vertex model Kernels

CUDA kernels and callers for 2D vertex models. More...

Functions
bool	gpu_vm_geometry (Dscalar2 *d_vertexPositions, int *d_cellVertexNum, int *d_cellVertices, int *d_vertexCellNeighbors, Dscalar2 *d_voroCur, Dscalar4 *d_voroLastNext, Dscalar2 *d_AreaPerimeter, int N, Index2D &n_idx, gpubox &Box)
	Call the kernel to calculate the area and perimeter of each cell.
bool	gpu_vm_get_cell_positions (Dscalar2 *d_cellPositions, Dscalar2 *d_vertexPositions, int *d_cellVertexNum, int *d_cellVertices, int N, Index2D &n_idx, gpubox &Box)
	Call the kernel to calculate the position of each cell from the position of its vertices.
bool	gpu_vm_test_edges_for_T1 (Dscalar2 *d_vertexPositions, int *d_vertexNeighbors, int *d_vertexEdgeFlips, int *d_vertexCellNeighbors, int *d_cellVertexNum, int *d_cellVertices, gpubox &Box, Dscalar T1THRESHOLD, int Nvertices, int vertexMax, int *d_grow, Index2D &n_idx)
	Call the kernel to test every edge for a T1 event, see if vertexMax needs to increase.
bool	gpu_vm_parse_multiple_flips (int *d_vertexEdgeFlips, int *d_vertexEdgeFlipsCurrent, int *d_vertexNeighbors, int *d_vertexCellNeighbors, int *d_cellVertexNum, int *d_cellVertices, int *d_finishedFlippingEdges, int *d_cellEdgeFlips, int4 *d_cellSets, Index2D &n_idx, int Ncells)
	determine whether any edges need to be flipped, and if we need to loop through the flipping routine, writing to d_finishedFlippingEdges the current state
bool	gpu_vm_flip_edges (int *d_vertexEdgeFlipsCurrent, Dscalar2 *d_vertexPositions, int *d_vertexNeighbors, int *d_vertexCellNeighbors, int *d_cellVertexNum, int *d_cellVertices, int *d_cellEdgeFlips, int4 *d_cellSets, gpubox &Box, Index2D &n_idx, int Nvertices, int Ncells)
	Call the kernel to flip at most one edge per cell, write to d_finishedFlippingEdges the current state.
bool	gpu_avm_force_sets (int *d_vertexCellNeighbors, Dscalar2 *d_voroCur, Dscalar4 *d_voroLastNext, Dscalar2 *d_AreaPerimeter, Dscalar2 *d_AreaPerimeterPreferences, Dscalar2 *d_vertexForceSets, int nForceSets, Dscalar KA, Dscalar KP)
	Call the kernel to calculate force sets.
bool	gpu_avm_sum_force_sets (Dscalar2 *d_vertexForceSets, Dscalar2 *d_vertexForces, int Nvertices)
	Call the kernel to sum up the force sets to get net force on each vertex.
__global__ void	vm_geometry_kernel (const Dscalar2 *__restrict__ d_vertexPositions, const int *__restrict__ d_cellVertexNum, const int *__restrict__ d_cellVertices, const int *__restrict__ d_vertexCellNeighbors, Dscalar2 *__restrict__ d_voroCur, Dscalar4 *__restrict__ d_voroLastNext, Dscalar2 *__restrict__ d_AreaPerimeter, int N, Index2D n_idx, gpubox Box)
__global__ void	vm_get_cell_positions_kernel (Dscalar2 *d_cellPositions, Dscalar2 *d_vertexPositions, int *d_nn, int *d_n, int N, Index2D n_idx, gpubox Box)
__global__ void	vm_simple_T1_test_kernel (Dscalar2 *d_vertexPositions, int *d_vertexNeighbors, int *d_vertexEdgeFlips, int *d_vertexCellNeighbors, int *d_cellVertexNum, gpubox Box, Dscalar T1THRESHOLD, int NvTimes3, int vertexMax, int *d_grow)
__global__ void	vm_one_T1_per_cell_per_vertex_kernel (int *__restrict__ d_vertexEdgeFlips, int *__restrict__ d_vertexEdgeFlipsCurrent, const int *__restrict__ d_vertexNeighbors, const int *__restrict__ d_vertexCellNeighbors, const int *__restrict__ d_cellVertexNum, const int *__restrict__ d_cellVertices, int *d_finishedFlippingEdges, int *d_cellEdgeFlips, int4 *d_cellSets, Index2D n_idx, int Ncells)
__global__ void	vm_flip_edges_kernel (int *d_vertexEdgeFlipsCurrent, Dscalar2 *d_vertexPositions, int *d_vertexNeighbors, int *d_vertexCellNeighbors, int *d_cellVertexNum, int *d_cellVertices, int *d_cellEdgeFlips, int4 *d_cellSets, gpubox Box, Index2D n_idx, int NvTimes3)
__global__ void	avm_force_sets_kernel (int *d_vertexCellNeighbors, Dscalar2 *d_voroCur, Dscalar4 *d_voroLastNext, Dscalar2 *d_AreaPerimeter, Dscalar2 *d_AreaPerimeterPreferences, Dscalar2 *d_vertexForceSets, int nForceSets, Dscalar KA, Dscalar KP)
__global__ void	avm_sum_force_sets_kernel (Dscalar2 *d_vertexForceSets, Dscalar2 *d_vertexForces, int N)

Detailed Description

CUDA kernels and callers for 2D vertex models.

Function Documentation

vm_geometry_kernel()

__global__ void vm_geometry_kernel	(	const Dscalar2 *__restrict__	d_vertexPositions,
		const int *__restrict__	d_cellVertexNum,
		const int *__restrict__	d_cellVertices,
		const int *__restrict__	d_vertexCellNeighbors,
		Dscalar2 *__restrict__	d_voroCur,
		Dscalar4 *__restrict__	d_voroLastNext,
		Dscalar2 *__restrict__	d_AreaPerimeter,
		int	N,
		Index2D	n_idx,
		gpubox	Box
	)

Since the cells are NOT guaranteed to be convex, the area of the cell must take into account any self-intersections. The strategy is the same as in the CPU branch.

vm_get_cell_positions_kernel()

__global__ void vm_get_cell_positions_kernel	(	Dscalar2 *	d_cellPositions,
		Dscalar2 *	d_vertexPositions,
		int *	d_nn,
		int *	d_n,
		int	N,
		Index2D	n_idx,
		gpubox	Box
	)

This function is being deprecated, but is still useful for calculating, e.g. the mean-squared displacement of the cells without transferring data back to the host

vm_simple_T1_test_kernel()

__global__ void vm_simple_T1_test_kernel	(	Dscalar2 *	d_vertexPositions,
		int *	d_vertexNeighbors,
		int *	d_vertexEdgeFlips,
		int *	d_vertexCellNeighbors,
		int *	d_cellVertexNum,
		gpubox	Box,
		Dscalar	T1THRESHOLD,
		int	NvTimes3,
		int	vertexMax,
		int *	d_grow
	)

Run through every pair of vertices (once), see if any T1 transitions should be done, and see if the cell-vertex list needs to grow

vm_one_T1_per_cell_per_vertex_kernel()

__global__ void vm_one_T1_per_cell_per_vertex_kernel	(	int *__restrict__	d_vertexEdgeFlips,
		int *__restrict__	d_vertexEdgeFlipsCurrent,
		const int *__restrict__	d_vertexNeighbors,
		const int *__restrict__	d_vertexCellNeighbors,
		const int *__restrict__	d_cellVertexNum,
		const int *__restrict__	d_cellVertices,
		int *	d_finishedFlippingEdges,
		int *	d_cellEdgeFlips,
		int4 *	d_cellSets,
		Index2D	n_idx,
		int	Ncells
	)

There will be severe topology mismatches if a cell is involved in more than one T1 transition simultaneously (due to incoherent updates of the cellVertices structure). So, go through the current list of edges that are marked to take part in a T1 transition and select one edge per cell to be flipped on this trip through the functions.

vm_flip_edges_kernel()

__global__ void vm_flip_edges_kernel	(	int *	d_vertexEdgeFlipsCurrent,
		Dscalar2 *	d_vertexPositions,
		int *	d_vertexNeighbors,
		int *	d_vertexCellNeighbors,
		int *	d_cellVertexNum,
		int *	d_cellVertices,
		int *	d_cellEdgeFlips,
		int4 *	d_cellSets,
		gpubox	Box,
		Index2D	n_idx,
		int	NvTimes3
	)

Flip any edge labeled for re-wiring in the vertexEdgeFlipsCurrent list

avm_force_sets_kernel()

__global__ void avm_force_sets_kernel	(	int *	d_vertexCellNeighbors,
		Dscalar2 *	d_voroCur,
		Dscalar4 *	d_voroLastNext,
		Dscalar2 *	d_AreaPerimeter,
		Dscalar2 *	d_AreaPerimeterPreferences,
		Dscalar2 *	d_vertexForceSets,
		int	nForceSets,
		Dscalar	KA,
		Dscalar	KP
	)

The force on a vertex has a contribution from how moving that vertex affects each of the neighboring cells...compute those force sets

avm_sum_force_sets_kernel()

__global__ void avm_sum_force_sets_kernel	(	Dscalar2 *	d_vertexForceSets,
		Dscalar2 *	d_vertexForces,
		int	N
	)

the force on a vertex is decomposable into the force contribution from each of its voronoi vertices... add 'em up!