#include "gpu_acc_map.hpp"
#include "../spatial_cells/block_adjust_gpu.hpp"

Include dependency graph for gpu_acc_map.cpp:

Macros
#define	LOG_BANKS 4
#define	BANK_OFFSET(n)
#define	ACCELERATION_KERNEl_MIN_BLOCKS (REGISTERS_PER_MP/64)/(WID3)

Functions
__global__ void	prefill_probe_kernel (vmesh::VelocityMesh *__restrict__ vmeshes, vmesh::LocalID dev_probeCubeData, const uint flatExtent, const size_t Dacc, const size_t Dother, const vmesh::LocalID invalidLID, split::SplitVector< vmesh::GlobalID > lists_with_replace_new, split::SplitVector< Hashinator::hash_pair< vmesh::GlobalID, vmesh::LocalID > > lists_delete, split::SplitVector< Hashinator::hash_pair< vmesh::GlobalID, vmesh::LocalID > > lists_to_replace, split::SplitVector< Hashinator::hash_pair< vmesh::GlobalID, vmesh::LocalID > > lists_with_replace_old, split::SplitVector< vmesh::GlobalID > **dev_vbwcl_vec, const uint cumulativeOffset, const size_t gpu_probeStride)
	GPU kernel which fills the target probe cube with the invalid value for vmesh::LocalID.
__global__ void	fill_VBC_zero_kernel (vmesh::VelocityBlockContainer **blockContainers, const uint cumulativeOffset)
__global__ void	fill_probe_ordered (vmesh::VelocityMesh *__restrict__ vmeshes, vmesh::LocalID dev_probeCubeData, const uint flatExtent, const uint *__restrict__ gpu_block_indices_to_probe, const uint cumulativeOffset, const size_t gpu_probeStride)
	GPU kernel for taking the contents of a vmesh and placign the existing velocity blocks in a probe cube for further reduction by a flattening kernel.
__global__ void	flatten_probe_cube (vmesh::LocalID *dev_probeCubeData, const vmesh::LocalID Dacc, const vmesh::LocalID Dother, const size_t flatExtent, const vmesh::LocalID invalidLID, const size_t gpu_probeStride)
	GPU kernel which flattens the probe cube into two reduction results (counters): how many columns and how many blocks exist per columnset.
__global__ void	scan_probe (vmesh::VelocityMesh *__restrict__ vmeshes, vmesh::LocalID dev_probeCubeData, const vmesh::LocalID Dacc, const vmesh::LocalID Dother, const size_t flatExtent, vmesh::LocalID dev_numCols, vmesh::LocalID dev_numColSets, vmesh::LocalID dev_resizeSuccess, ColumnOffsets dev_columnOffsetData, const uint cumulativeOffset, const size_t gpu_probeStride)
	GPU kernel which performs exclusive prefix scans of the flattened probe cube, providing cumulative sums up to each index for use in further kernels.
__global__ void	build_column_offsets (vmesh::VelocityMesh *__restrict__ vmeshes, vmesh::LocalID dev_probeCubeData, const vmesh::LocalID D0, const vmesh::LocalID D1, const vmesh::LocalID D2, const int dimension, const size_t flatExtent, const vmesh::LocalID invalidLID, ColumnOffsets dev_columnOffsetData, split::SplitVector< vmesh::GlobalID > *dev_vbwcl_vec, const uint cumulativeOffset, const size_t gpu_probeStride)
	GPU kernel for building the columns and columnSets for each spatial cell, and storing offsets and lengths into the ColumnOffsets struct.
__global__ void	__launch_bounds__ (WID3) reorder_blocks_by_dimension_kernel(vmesh
	GPU kernel for reading block data in from the spatial cell VelocityBlockContainers, transposing each block so the acceleration direction is the k-index, and storing each block into ordered buffers which will then be fed to the acceleration kernel.
__global__ void	__launch_bounds__ (GPUTHREADS, 4) evaluate_column_extents_kernel(const uint dimension
	GPU Kernel for evaluating the constructed columns, which blocks need to be deleted from a velocity mesh, and which ones need to be added to it.
	if (setIndex< columnData->setColumnOffsets.size())
__global__ void	__launch_bounds__ (WID3, ACCELERATION_KERNEl_MIN_BLOCKS) acceleration_kernel(vmesh
	GPU kernel for main task of semi-Lagrangian acceleration. Reads data in from buffer, performs polynomial reconstruction and does piecewise integration of contribution, and stores results directly into VelocityBlockContainer.
__host__ bool	gpu_acc_map_1d (dccrg::Dccrg< spatial_cell::SpatialCell, dccrg::Cartesian_Geometry > &mpiGrid, vector< CellID > &launchCells, const uint popID, const uint dimension, const int Dacc, const int Dother, const size_t cumulativeOffset)
	This function performs the semi-Lagrangian acceleration for a provided list of spatial cells, for one popID, for one dimension. See gpu_acc_semilag.cpp for Information on the calling structure.

Variables
__global__ void vmesh::VelocityMesh **__restrict__	vmeshes
__global__ void vmesh::VelocityMesh *__restrict__ ColumnOffsets	dev_columnOffsetData
__global__ void vmesh::VelocityMesh __restrict__ ColumnOffsets split::SplitVector< vmesh::GlobalID >	lists_with_replace_new
__global__ void vmesh::VelocityMesh __restrict__ ColumnOffsets split::SplitVector< vmesh::GlobalID > Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID >	allMaps
__global__ void vmesh::VelocityMesh *__restrict__ ColumnOffsets split::SplitVector< vmesh::GlobalID > Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > const uint __restrict__	gpu_block_indices_to_id
__global__ void vmesh::VelocityMesh *__restrict__ ColumnOffsets split::SplitVector< vmesh::GlobalID > Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > const uint __restrict__ const Realf *	dev_intersections
__global__ void vmesh::VelocityMesh *__restrict__ ColumnOffsets split::SplitVector< vmesh::GlobalID > Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > const uint __restrict__ const Realf const int	bailout_velocity_space_wall_margin
__global__ void vmesh::VelocityMesh *__restrict__ ColumnOffsets split::SplitVector< vmesh::GlobalID > Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > const uint __restrict__ const Realf const int const int	max_v_length
__global__ void vmesh::VelocityMesh *__restrict__ ColumnOffsets split::SplitVector< vmesh::GlobalID > Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > const uint __restrict__ const Realf const int const int const Realf	v_min
__global__ void vmesh::VelocityMesh *__restrict__ ColumnOffsets split::SplitVector< vmesh::GlobalID > Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > const uint __restrict__ const Realf const int const int const Realf const Realf	dv
__global__ void vmesh::VelocityMesh *__restrict__ ColumnOffsets split::SplitVector< vmesh::GlobalID > Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > const uint __restrict__ const Realf const int const int const Realf const Realf vmesh::LocalID *	dev_resizeSuccess
__global__ void vmesh::VelocityMesh *__restrict__ ColumnOffsets split::SplitVector< vmesh::GlobalID > Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > const uint __restrict__ const Realf const int const int const Realf const Realf vmesh::LocalID vmesh::LocalID *	dev_overflownElements
__global__ void vmesh::VelocityMesh *__restrict__ ColumnOffsets split::SplitVector< vmesh::GlobalID > Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > const uint __restrict__ const Realf const int const int const Realf const Realf vmesh::LocalID vmesh::LocalID const uint	cumulativeOffset
const uint	setIndex = blockIdx.x
const uint	ti = threadIdx.x
const uint	parallelOffsetIndex = blockIdx.y
const uint	cellOffset = parallelOffsetIndex + cumulativeOffset
const Realf	intersection = dev_intersections[cellOffset*4+0]
const Realf	intersection_di = dev_intersections[cellOffset*4+1]
const Realf	intersection_dj = dev_intersections[cellOffset*4+2]
const Realf	intersection_dk = dev_intersections[cellOffset*4+3]
ColumnOffsets *	columnData = dev_columnOffsetData + parallelOffsetIndex
Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > *	dev_map_require = allMaps[2*cellOffset]
Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > *	dev_map_remove = allMaps[2*cellOffset+1]
split::SplitVector< vmesh::GlobalID > *	list_with_replace_new = lists_with_replace_new[cellOffset]
const vmesh::VelocityMesh *__restrict__	vmesh = vmeshes[cellOffset]
__shared__ int	isTargetBlock [MAX_BLOCKS_PER_DIM]
__shared__ int	isSourceBlock [MAX_BLOCKS_PER_DIM]

Macro Definition Documentation

◆ ACCELERATION_KERNEl_MIN_BLOCKS

#define ACCELERATION_KERNEl_MIN_BLOCKS (REGISTERS_PER_MP/64)/(WID3)

Definition at line 927 of file gpu_acc_map.cpp.

◆ BANK_OFFSET

#define BANK_OFFSET ( n )

Value:

0

Definition at line 39 of file gpu_acc_map.cpp.

◆ LOG_BANKS

#define LOG_BANKS 4

Definition at line 33 of file gpu_acc_map.cpp.

Function Documentation

◆ __launch_bounds__() [1/3]

__global__ void __launch_bounds__	(	GPUTHREADS	,
		4	) const

GPU Kernel for evaluating the constructed columns, which blocks need to be deleted from a velocity mesh, and which ones need to be added to it.

For each column, evaluates which blocks are source blocks, which are target blocks, and which ones are both. Then goes through this list and pushes values into two sets and one vector.

TODO: This kernel could probaly be streamlined, but it is fast to execute so hasn't been a priority.

Note: max_v_length, v_min, dv could all be queried from any random vmesh instead of passing from host.

Parameters

dimension	direction of acceleration
vmeshes	pointer to buffer of pointers to vmeshes of all cells on this process
dev_columnOffsetData	pointer to a buffer of ColumnOffsets structs (see arch/gpu_base.hpp) where per-cell information on all built columns and columnSets is stored
lists_with_replace_new	pointer to buffer of pointers to splitvectors, where newly added blocks are listed.
allMaps	pointer to buffer of pointers to maps, two per spatial cell, used for gathering deleted and required blocks.
gpu_block_indices_to_probe	3-element array used for converting block GID to location within probe cube
dev_intersections	pointer to buffer of per-cell intersection values pre-calculated for this acceleration direction.
bailout_velocity_space_wall_margin	integer for how close to v-space walls we allow blocks to be created before bailout
max_v_length	maximum extent of v-space for this direction of acceleration
v_min	start of v-space for this direction of acceleration
dv	size of v for each velocity cell in this direction of acceleration
dev_resizeSuccess	pointer to buffer where the amount by which the list_with_replace_new capacity was exceeded
dev_overflownElements	pointer to buffer for bailout flag for touching v-space walls
cumulativeOffset	the current cumulative offset at which to index the aforementioned buffers, using the grid index of this kernel on top of this provided offset.

◆ __launch_bounds__() [2/3]

__global__ void __launch_bounds__ ( WID3 )

GPU kernel for reading block data in from the spatial cell VelocityBlockContainers, transposing each block so the acceleration direction is the k-index, and storing each block into ordered buffers which will then be fed to the acceleration kernel.

Reads block data in from spatial cells and places it in the ordered container. Each block is in column-based order. Works column-per-column and adds the necessary one empty block at each end

Block contents are transposed so that values 0...WID3-1 (i.e. each block) are so that every WID3 consequtive elements form one "slice" perpendicular to the direction of acceleration. This is achieved by swapping the indexing order for two dimensions (if needed), in accordance with historic acceleration solver tradition.

Example for WID=4: Z ^ Y 60 61 62 63 | / 56 57 58 59 | / 52 53 54 55 |/ 48 49 50 51 *-----> X 44 45 46 47 40 41 42 43 36 37 38 39 32 33 34 35 28 29 30 31 24 25 26 27 20 21 22 23 16 17 18 19 12 13 1 15 8 9 10 11 4 5 6 7 0 1 2 3

now the first 16 elements of the reodered buffer would correspond to source indices as follows:

acceleration along X: (swap X and Z directions) 0 16 32 48 4 20 36 52 8 24 40 12 28 44 60

acceleration along Y: (swap Y and Z directions) 0 1 2 3 16 17 18 19 32 33 34 35 48 49 50 51

acceleration along Z: (no swaps) 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15

This way the acceleration kernel can access cells from the input buffer such that there is a constant offset WID2 between cells along the same acceleration direction.

Parameters

blockContainers	pointer to buffer of pointers to VelocityBlockContainers of all cells on this process
dev_blockDataOrdered	pointer to buffer of pointers, which point to allocated temporary arrays used for storing the phase-space data to be propagated.
gpu_cell_indices_to_id	buffer of 3 values used for converting block indices between directions
dev_vbwcl_vec	pointer to buffer of pointers to SplitVectors, re-cast to use as an (ordered) LID list
dev_columnOffsetData	pointer to a buffer of ColumnOffsets structs (see arch/gpu_base.hpp) where per-cell information on all built columns and columnSets is stored
dev_nColumns	pointer to buffer of data indicating how many columns each cell has
cumulativeOffset	the current cumulative offset at which to index the aforementioned buffers, using the grid index of this kernel on top of this provided offset.

GPU kernel for calculating second velocity moments from provided velocity block containers

Definition at line 665 of file gpu_acc_map.cpp.

◆ __launch_bounds__() [3/3]

__global__ void __launch_bounds__	(	WID3	,
		ACCELERATION_KERNEl_MIN_BLOCKS	)

GPU kernel for main task of semi-Lagrangian acceleration. Reads data in from buffer, performs polynomial reconstruction and does piecewise integration of contribution, and stores results directly into VelocityBlockContainer.

Note: v_min, dv could all be queried from any random vmesh instead of passing from host.

Parameters

vmeshes	pointer to buffer of pointers to vmeshes of all cells on this process
blockContainers	pointer to buffer of pointers to VelocityBlockContainers of all cells on this process
dev_blockDataOrdered	pointer to buffer of pointers, which point to allocated temporary arrays used for storing the phase-space data to be propagated.
gpu_cell_indices_to_id	buffer of 3 values used for converting block indices between directions
gpu_block_indices_to_probe	3-element array used for converting propageted block indexes bacl to GID
dev_columnOffsetData	pointer to a buffer of ColumnOffsets structs (see arch/gpu_base.hpp) where per-cell information on all built columns and columnSets is stored
dev_intersections	pointer to buffer of per-cell intersection values pre-calculated for this acceleration direction.
v_min	start of v-space for this direction of acceleration
i_dv	unity per size of v for each velocity cell in this direction of acceleration
dv	size of v for each velocity cell in this direction of acceleration
dev_minvalues	pointer to buffer of stored sparsity minValues for each cell, used by slope limiters
invalidLID	value to be used as the value for invalid vmesh::LocalID
cumulativeOffset	the current cumulative offset at which to index the aforementioned buffers, using the grid index of this kernel on top of this provided offset.

Definition at line 953 of file gpu_acc_map.cpp.

Here is the call graph for this function:

◆ build_column_offsets()

__global__ void build_column_offsets	(	vmesh::VelocityMesh **__restrict__	vmeshes,
		vmesh::LocalID *	dev_probeCubeData,
		const vmesh::LocalID	D0,
		const vmesh::LocalID	D1,
		const vmesh::LocalID	D2,
		const int	dimension,
		const size_t	flatExtent,
		const vmesh::LocalID	invalidLID,
		ColumnOffsets *	dev_columnOffsetData,
		split::SplitVector< vmesh::GlobalID > **	dev_vbwcl_vec,
		const uint	cumulativeOffset,
		const size_t	gpu_probeStride )

GPU kernel for building the columns and columnSets for each spatial cell, and storing offsets and lengths into the ColumnOffsets struct.

Utilizing the previously computed offsets (with the probe cube), this parallel kernel builds the offsets required for columns. It read the contents of the probe cube and outputs the stored GIDs and LIDs into provided buffers.

Parameters

vmeshes	pointer to buffer of pointers to vmeshes of all cells on this process
dev_probeCubeData	pointer to buffer of pointers, which point to allocated temporary arrays, which are recast and used as the probe cube.
D0	Vmesh maximal extents in X-direction
D1	Vmesh maximal extents in Y-direction
D2	Vmesh maximal extents in Z-direction
dimension	direction of acceleration
flatExtent	the product of the non-accelerated dimensions of the probe cube, i.e. the extent of it when it's flattend, rounded up to provide memory alignment.
invalidLID	value to be used as the value for invalid vmesh::LocalID
dev_columnOffsetData	pointer to a buffer of ColumnOffsets structs (see arch/gpu_base.hpp) where per-cell information on all built columns and columnSets is stored
dev_vbwcl_vec	pointer to buffer of pointers to SplitVectors, re-cast to use as an (ordered) LID list
cumulativeOffset	the current cumulative offset at which to index the aforementioned buffers, using the grid index of this kernel on top of this provided offset.

Definition at line 486 of file gpu_acc_map.cpp.

Here is the caller graph for this function:

◆ fill_probe_ordered()

__global__ void fill_probe_ordered	(	vmesh::VelocityMesh **__restrict__	vmeshes,
		vmesh::LocalID *	dev_probeCubeData,
		const uint	flatExtent,
		const uint *__restrict__	gpu_block_indices_to_probe,
		const uint	cumulativeOffset,
		const size_t	gpu_probeStride )

GPU kernel for taking the contents of a vmesh and placign the existing velocity blocks in a probe cube for further reduction by a flattening kernel.

One option for a probe cube would be to reduce (with __ballot_sync) along the direction of propagation (k) to get the number of blocks in a column. However, this does not have an obvious way to support gathering columnsets (several columns at one set of perpendicular i,j indices).

Thus, instead, for following analysis of the probe cube, we want the warp/wavefront to read a dimension not propagating along, because then the wavefront can loop over the dimension to propagate along (size Dacc), and each thread processes one potential columnSet. To even better parallelize and simplify, the two non-propagated dimensions (i,j) are merged int a dimension of size Dother (so it isn't an actual cube).

Since we read the data (existing blocks) in LID order, writes will be jumbled anyhow, so we can take this opportunity to make future accesses to the probe cube efficient by writing in a smart order.

TODO: ensure the first and second dimensions are powers of two for optimized reads? Then stepping will be based on array edge sizes instead of D0/1/2.

Parameters

vmeshes	pointer to buffer of pointers to vmeshes of all cells on this process
dev_probeCubeData	pointer to buffer of pointers, which point to allocated temporary arrays, which are recast and used as the probe cube.
flatExtent	the product of the non-accelerated dimensions of the probe cube, i.e. the extent of it when it's flattend, rounded up to provide memory alignment.
gpu_block_indices_to_probe	3-element array used for converting block GID to location within probe cube
cumulativeOffset	the current cumulative offset at which to index the aforementioned buffers, using the grid index of this kernel on top of this provided offset.

Definition at line 163 of file gpu_acc_map.cpp.

Here is the caller graph for this function:

◆ fill_VBC_zero_kernel()

__global__ void fill_VBC_zero_kernel	(	vmesh::VelocityBlockContainer **	blockContainers,
		const uint	cumulativeOffset )

GPU kernel which fills the target block container with zeroes

Parameters

blockContainers	pointer to buffer of pointers to per-cell velocity block containers
cumulativeOffset	the current cumulative offset at which to index the aforementioned buffers, using the grid index of this kernel on top of this provided offset.

Definition at line 114 of file gpu_acc_map.cpp.

Here is the call graph for this function:

Here is the caller graph for this function:

◆ flatten_probe_cube()

__global__ void flatten_probe_cube	(	vmesh::LocalID *	dev_probeCubeData,
		const vmesh::LocalID	Dacc,
		const vmesh::LocalID	Dother,
		const size_t	flatExtent,
		const vmesh::LocalID	invalidLID,
		const size_t	gpu_probeStride )

GPU kernel which flattens the probe cube into two reduction results (counters): how many columns and how many blocks exist per columnset.

In the probe cube, there's the dimension of acceleration (size Dacc) and the other two dimensions, merged (size Dother). For each index in the other two dimensions (i,j), we have a position in transcerse v-space associated with the possibility for constructing columns. Thus, that position is now termed a potential column position or potColumn.

This kernel loops over Dacc to find: (1) How many acceleration columns were found for each potColumn (2) How many blocks were found for each potColumn

Parameters

dev_probeCubeData	pointer to buffer of pointers, which point to allocated temporary arrays, which are recast and used as the probe cube.
Dacc	Extent of probe cube in accelerated dimension
Dother	Sum of extents of probe cube in transverse dimensions
flatExtent	the product of the non-accelerated dimensions of the probe cube, i.e. the extent of it when it's flattend, rounded up to provide memory alignment.
invalidLID	value to be used as the value for invalid vmesh::LocalID

Definition at line 218 of file gpu_acc_map.cpp.

Here is the caller graph for this function:

◆ gpu_acc_map_1d()

__host__ bool gpu_acc_map_1d	(	dccrg::Dccrg< spatial_cell::SpatialCell, dccrg::Cartesian_Geometry > &	mpiGrid,
		vector< CellID > &	launchCells,
		const uint	popID,
		const uint	dimension,
		const int	Dacc,
		const int	Dother,
		const size_t	cumulativeOffset )

This function performs the semi-Lagrangian acceleration for a provided list of spatial cells, for one popID, for one dimension. See gpu_acc_semilag.cpp for Information on the calling structure.

First, kernels are called to construct the existing columns and columnsets for the sparse velocity space data contained in this cell. A parallel launch and probe cubes are used for this, along with some unified memory column containers. Data from the current spatial cell is then read into an intermediate buffer in column-aligned order.

Then, utilizing pre-calculated SLICE-3D intersections, a kernel is launched to evaluate the extents of velocity space after each column(set) has been accelerated as requested. Block adjustment functions are called to batch-update the velocity space of the cell in question to match the new requirements (adding and removing cells as needed). Then, the velocity block container is cleared.

Finally the acceleration kernel itself is called. It reads the aligned velocity space data from the intermediate buffer, advects the columns according to the SLICE-3D intersections, and stores the resultant phase space density back into the velocity block container.

Parameters

mpiGrid	DCCRG container of spatial cells
launchCells	vector of cells for which to perform acceleration in this "chunk"
popID	ID of the accelerated particle species.
dimension	Velocity dimension for acceleration (VX, VY, or VZ)
Dacc	Maximal velocity block extent in accelerated dimension
Dother	Product of maximal velocity block extents in non-accelerated dimensions
cumulativeOffset	running counter for offset of indexing of launchCells into device buffers

New merged kernel approach without sorting for columns

First, we generate a "probe cube". It started off as an actual cube, but the transverse dimensions are considered as one. One dimension is that of the current acceleration (Dacc), and the other dimension is the product of the other two maximal velocity block domain extents (Dother). The "flattened" version is one where data is gathered over the acceleration direction into a single reduced value.

The flattened version of the probe cube must store: 1) how many columns per potential column position (potColumn) 2) how many blocks per potColumn 3) cumulative offset into columns per potColumn 4) cumulative offset into columnSets per potColumn 5) cumulative offset into blocks per potColumn

For reductions, each slice of the flattened array should have a size a multiple of 2*MAX_BLOCKSIZE:

For the gathered LIDlist, we re-use the allocation of spatial_cell->dev_velocity_block_with_content_list, It which contains variables of type vmesh::GlobalID (which should be the same as vmesh::LocalID, uint_32t). To ensure this static_cast is safe, we verify the sizes.

Use block adjustment callers / lambda rules for extracting required map contents, building up vectors to use for parallel adjustment.

Definition at line 1212 of file gpu_acc_map.cpp.

Here is the call graph for this function:

Here is the caller graph for this function:

◆ if()

if ( setIndex< columnData->setColumnOffsets. size() )

Definition at line 783 of file gpu_acc_map.cpp.

Here is the call graph for this function:

Here is the caller graph for this function:

◆ prefill_probe_kernel()

__global__ void prefill_probe_kernel	(	vmesh::VelocityMesh **__restrict__	vmeshes,
		vmesh::LocalID *	dev_probeCubeData,
		const uint	flatExtent,
		const size_t	Dacc,
		const size_t	Dother,
		const vmesh::LocalID	invalidLID,
		split::SplitVector< vmesh::GlobalID > **	lists_with_replace_new,
		split::SplitVector< Hashinator::hash_pair< vmesh::GlobalID, vmesh::LocalID > > **	lists_delete,
		split::SplitVector< Hashinator::hash_pair< vmesh::GlobalID, vmesh::LocalID > > **	lists_to_replace,
		split::SplitVector< Hashinator::hash_pair< vmesh::GlobalID, vmesh::LocalID > > **	lists_with_replace_old,
		split::SplitVector< vmesh::GlobalID > **	dev_vbwcl_vec,
		const uint	cumulativeOffset,
		const size_t	gpu_probeStride )

GPU kernel which fills the target probe cube with the invalid value for vmesh::LocalID.

Also clears provided vectors which are needed as empty for future kernels. Clearing on-device prevents page faulting of unified memory bookkeeping data. Resizes the per-cell velocity blocks with content list vector as it'll be re-used for a LID list.

Parameters

vmeshes	pointer to buffer of pointers to vmeshes of all cells on this process
dev_probeCubeData	pointer to buffer of pointers, which point to allocated temporary arrays, which are recast and used as the probe cube.
flatExtent	the product of the non-accelerated dimensions of the probe cube, i.e. the extent of it when it's flattend, rounded up to provide memory alignment.
nTot	number of total entries in probe cube
invalidLID	value to be used as the value for invalid vmesh::LocalID
lists_with_replace_new	pointer to buffer of pointers to SplitVectors to be cleared
lists_delete	pointer to buffer of pointers to SplitVectors to be cleared
lists_to_replace	pointer to buffer of pointers to SplitVectors to be cleared
lists_with_replace_old	pointer to buffer of pointers to SplitVectors to be cleared
dev_vbwcl_vec	pointer to buffer of pointers to SplitVectors, re-sized and re-cast to use as a LID list
cumulativeOffset	the current cumulative offset at which to index the aforementioned buffers, using the grid index of this kernel on top of this provided offset.

Definition at line 64 of file gpu_acc_map.cpp.

Here is the caller graph for this function:

◆ scan_probe()

__global__ void scan_probe	(	vmesh::VelocityMesh **__restrict__	vmeshes,
		vmesh::LocalID *	dev_probeCubeData,
		const vmesh::LocalID	Dacc,
		const vmesh::LocalID	Dother,
		const size_t	flatExtent,
		vmesh::LocalID *	dev_numCols,
		vmesh::LocalID *	dev_numColSets,
		vmesh::LocalID *	dev_resizeSuccess,
		ColumnOffsets *	dev_columnOffsetData,
		const uint	cumulativeOffset,
		const size_t	gpu_probeStride )

GPU kernel which performs exclusive prefix scans of the flattened probe cube, providing cumulative sums up to each index for use in further kernels.

This scan produces:

(1) the cumulative sum of columns up to the beginning of each potColumn (2) the cumulative sum of column sets up to the beginning of each potColumn (3) the cumulative sum of blocks up to the beginning of each potColumn

This is not a fully optimized scan as it uses only a single block in order to perform it all in one kernel. As the fixed-size shared memory buffer gets overwritten each cycle, we store the actual prefix scan results into the third and fourth entries in the probeFlattened buffer and keep track of the accumulated offset to the prefix.

The count of columns to be evaluated is estimated to remain somewhat small-ish, so there should not be all that many loops of the cycle to deal with.

Parameters

vmeshes	pointer to buffer of pointers to vmeshes of all cells on this process
dev_probeCubeData	pointer to buffer of pointers, which point to allocated temporary arrays, which are recast and used as the probe cube.
Dacc	Extent of probe cube in accelerated dimension
Dother	Sum of extents of probe cube in transverse dimensions
flatExtent	the product of the non-accelerated dimensions of the probe cube, i.e. the extent of it when it's flattend, rounded up to provide memory alignment.
dev_numCols	buffer for storing the total count of columns per cell (or probe cube)
dev_numColSets	buffer for storing the total count of columnSets per cell (or probe cube)
dev_resizeSuccess	buffer for storing a flag, if the columnData container re-sizing to match the counts of columns and columnSets was suffessful or not. If failed, re-sizing needs to occur on host.
dev_columnOffsetData	pointer to a buffer of ColumnOffsets structs (see arch/gpu_base.hpp) where per-cell information on all built columns and columnSets is stored
cumulativeOffset	the current cumulative offset at which to index the aforementioned buffers, using the grid index of this kernel on top of this provided offset.

Definition at line 306 of file gpu_acc_map.cpp.

Here is the call graph for this function:

Here is the caller graph for this function:

Variable Documentation

◆ allMaps

__global__ void vmesh::VelocityMesh** __restrict__ ColumnOffsets split::SplitVector<vmesh::GlobalID> Hashinator::Hashmap<vmesh::GlobalID,vmesh::LocalID>* * allMaps

Definition at line 751 of file gpu_acc_map.cpp.

◆ bailout_velocity_space_wall_margin

__global__ void vmesh::VelocityMesh** __restrict__ ColumnOffsets split::SplitVector<vmesh::GlobalID> Hashinator::Hashmap<vmesh::GlobalID,vmesh::LocalID> const uint* __restrict__ const Realf const int bailout_velocity_space_wall_margin

Definition at line 754 of file gpu_acc_map.cpp.

◆ cellOffset

const uint cellOffset = parallelOffsetIndex + cumulativeOffset

Definition at line 767 of file gpu_acc_map.cpp.

◆ columnData

ColumnOffsets* columnData = dev_columnOffsetData + parallelOffsetIndex

Definition at line 773 of file gpu_acc_map.cpp.

◆ cumulativeOffset

__global__ void vmesh::VelocityMesh** __restrict__ ColumnOffsets split::SplitVector<vmesh::GlobalID> Hashinator::Hashmap<vmesh::GlobalID,vmesh::LocalID> const uint* __restrict__ const Realf const int const int const Realf const Realf vmesh::LocalID vmesh::LocalID const uint cumulativeOffset

Initial value:

{

const uint warpSize = blockDim.x

Definition at line 760 of file gpu_acc_map.cpp.

◆ dev_columnOffsetData

__global__ void vmesh::VelocityMesh** __restrict__ ColumnOffsets* dev_columnOffsetData

Definition at line 749 of file gpu_acc_map.cpp.

◆ dev_intersections

__global__ void vmesh::VelocityMesh** __restrict__ ColumnOffsets split::SplitVector<vmesh::GlobalID> Hashinator::Hashmap<vmesh::GlobalID,vmesh::LocalID> const uint* __restrict__ const Realf* dev_intersections

Definition at line 753 of file gpu_acc_map.cpp.

◆ dev_map_remove

Hashinator::Hashmap<vmesh::GlobalID,vmesh::LocalID>* dev_map_remove = allMaps[2*cellOffset+1]

Definition at line 776 of file gpu_acc_map.cpp.

◆ dev_map_require

Hashinator::Hashmap<vmesh::GlobalID,vmesh::LocalID>* dev_map_require = allMaps[2*cellOffset]

Definition at line 775 of file gpu_acc_map.cpp.

◆ dev_overflownElements

__global__ void vmesh::VelocityMesh** __restrict__ ColumnOffsets split::SplitVector<vmesh::GlobalID> Hashinator::Hashmap<vmesh::GlobalID,vmesh::LocalID> const uint* __restrict__ const Realf const int const int const Realf const Realf vmesh::LocalID vmesh::LocalID* dev_overflownElements

Definition at line 759 of file gpu_acc_map.cpp.

◆ dev_resizeSuccess

__global__ void vmesh::VelocityMesh** __restrict__ ColumnOffsets split::SplitVector<vmesh::GlobalID> Hashinator::Hashmap<vmesh::GlobalID,vmesh::LocalID> const uint* __restrict__ const Realf const int const int const Realf const Realf vmesh::LocalID* dev_resizeSuccess

Definition at line 758 of file gpu_acc_map.cpp.

◆ dv

__global__ void vmesh::VelocityMesh** __restrict__ ColumnOffsets split::SplitVector<vmesh::GlobalID> Hashinator::Hashmap<vmesh::GlobalID,vmesh::LocalID> const uint* __restrict__ const Realf const int const int const Realf const Realf dv

Definition at line 757 of file gpu_acc_map.cpp.

◆ gpu_block_indices_to_id

__global__ void vmesh::VelocityMesh** __restrict__ ColumnOffsets split::SplitVector<vmesh::GlobalID> Hashinator::Hashmap<vmesh::GlobalID,vmesh::LocalID> const uint* __restrict__ gpu_block_indices_to_id

Definition at line 752 of file gpu_acc_map.cpp.

◆ intersection

const Realf intersection = dev_intersections[cellOffset*4+0]

Definition at line 769 of file gpu_acc_map.cpp.

◆ intersection_di

const Realf intersection_di = dev_intersections[cellOffset*4+1]

Definition at line 770 of file gpu_acc_map.cpp.

◆ intersection_dj

const Realf intersection_dj = dev_intersections[cellOffset*4+2]

Definition at line 771 of file gpu_acc_map.cpp.

◆ intersection_dk

const Realf intersection_dk = dev_intersections[cellOffset*4+3]

Definition at line 772 of file gpu_acc_map.cpp.

◆ isSourceBlock

__shared__ int isSourceBlock[MAX_BLOCKS_PER_DIM]

Definition at line 781 of file gpu_acc_map.cpp.

◆ isTargetBlock

__shared__ int isTargetBlock[MAX_BLOCKS_PER_DIM]

Definition at line 780 of file gpu_acc_map.cpp.

◆ list_with_replace_new

split::SplitVector<vmesh::GlobalID>* list_with_replace_new = lists_with_replace_new[cellOffset]

Definition at line 777 of file gpu_acc_map.cpp.

◆ lists_with_replace_new

__global__ void vmesh::VelocityMesh** __restrict__ ColumnOffsets split::SplitVector<vmesh::GlobalID>* * lists_with_replace_new

Definition at line 750 of file gpu_acc_map.cpp.

◆ max_v_length

__global__ void vmesh::VelocityMesh** __restrict__ ColumnOffsets split::SplitVector<vmesh::GlobalID> Hashinator::Hashmap<vmesh::GlobalID,vmesh::LocalID> const uint* __restrict__ const Realf const int const int max_v_length

Definition at line 755 of file gpu_acc_map.cpp.

◆ parallelOffsetIndex

const uint parallelOffsetIndex = blockIdx.y

Definition at line 766 of file gpu_acc_map.cpp.

◆ setIndex

const uint setIndex = blockIdx.x

Definition at line 763 of file gpu_acc_map.cpp.

◆ ti

const uint ti = threadIdx.x

Definition at line 764 of file gpu_acc_map.cpp.

◆ v_min

__global__ void vmesh::VelocityMesh** __restrict__ ColumnOffsets split::SplitVector<vmesh::GlobalID> Hashinator::Hashmap<vmesh::GlobalID,vmesh::LocalID> const uint* __restrict__ const Realf const int const int const Realf v_min

Definition at line 756 of file gpu_acc_map.cpp.

◆ vmesh

const vmesh::VelocityMesh* __restrict__ vmesh = vmeshes[cellOffset]

Definition at line 778 of file gpu_acc_map.cpp.

◆ vmeshes

__global__ void vmesh::VelocityMesh** __restrict__ vmeshes

Definition at line 748 of file gpu_acc_map.cpp.

Macros

Functions

Variables

Macro Definition Documentation

◆ ACCELERATION_KERNEl_MIN_BLOCKS

◆ BANK_OFFSET

◆ LOG_BANKS

Function Documentation

◆ __launch_bounds__() [1/3]

◆ __launch_bounds__() [2/3]

◆ __launch_bounds__() [3/3]

◆ build_column_offsets()

◆ fill_probe_ordered()

◆ fill_VBC_zero_kernel()

◆ flatten_probe_cube()

◆ gpu_acc_map_1d()

◆ if()

◆ prefill_probe_kernel()

◆ scan_probe()

Variable Documentation

◆ allMaps

◆ bailout_velocity_space_wall_margin

◆ cellOffset

◆ columnData

◆ cumulativeOffset

◆ dev_columnOffsetData

◆ dev_intersections

◆ dev_map_remove

◆ dev_map_require

◆ dev_overflownElements

◆ dev_resizeSuccess

◆ dv

◆ gpu_block_indices_to_id

◆ intersection

◆ intersection_di

◆ intersection_dj

◆ intersection_dk

◆ isSourceBlock

◆ isTargetBlock

◆ list_with_replace_new

◆ lists_with_replace_new

◆ max_v_length

◆ parallelOffsetIndex

◆ setIndex

◆ ti

◆ v_min

◆ vmesh

◆ vmeshes