spatial_cell_gpu_kernels.hpp File Reference

This graph shows which files directly or indirectly include this file:

Go to the source code of this file.

Functions
__global__ void	__launch_bounds__ (WID3, 4) update_velocity_block_content_lists_kernel(const vmesh
__global__ void	update_velocity_halo_kernel (const vmesh::VelocityMesh *__restrict__ vmesh, const vmesh::LocalID velocity_block_with_content_list_size, const vmesh::GlobalID *__restrict__ velocity_block_with_content_list_data, Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > *dev_velocity_block_with_content_map, Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > *dev_velocity_block_with_no_content_map)
__global__ void	update_neighbour_halo_kernel (const vmesh::VelocityMesh *__restrict__ vmesh, const uint neighbour_count, const vmesh::GlobalID *__restrict__ const *dev_neigh_vbwcls, const vmesh::LocalID *__restrict__ dev_neigh_Nvbwcls, Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > *dev_velocity_block_with_content_map, Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > *dev_velocity_block_with_no_content_map)
__global__ void	update_vmesh_and_blockparameters_kernel (vmesh::VelocityMesh *dev_vmesh, vmesh::VelocityBlockContainer *dev_blockContainer, const vmesh::LocalID nLIDs)
__global__ void	resize_vbc_kernel_pre (vmesh::VelocityMesh *vmesh, vmesh::VelocityBlockContainer *blockContainer, const split::SplitVector< vmesh::GlobalID > *__restrict__ list_with_replace_new, const split::SplitVector< Hashinator::hash_pair< vmesh::GlobalID, vmesh::LocalID > > *__restrict__ list_delete, const split::SplitVector< Hashinator::hash_pair< vmesh::GlobalID, vmesh::LocalID > > *__restrict__ list_to_replace, const split::SplitVector< Hashinator::hash_pair< vmesh::GlobalID, vmesh::LocalID > > *__restrict__ list_with_replace_old, vmesh::LocalID *returnLID, Realf *gpu_rhoLossAdjust)
__global__ void	resize_vbc_kernel_post (vmesh::VelocityMesh *vmesh, vmesh::VelocityBlockContainer *blockContainer, const vmesh::LocalID nBlocksAfterAdjust)

Function Documentation

__launch_bounds__()

__global__ void __launch_bounds__	(	WID3	,
		4	) const

Note: these kernels are called only when adjusting blocks for a single spatial cell. Usually, these are performed as batch operations, found in block_adjust_gpu.cpp and block_adjust_gpu_kernels.hpp GPU kernel for identifying which blocks have relevant content

GPU kernel for updating blocks based on generated lists

Definition at line 38 of file spatial_cell_gpu_kernels.hpp.

Here is the call graph for this function:

resize_vbc_kernel_post()

__global__ void resize_vbc_kernel_post	(	vmesh::VelocityMesh *	vmesh,
		vmesh::VelocityBlockContainer *	blockContainer,
		const vmesh::LocalID	nBlocksAfterAdjust )

Mini-kernel for adjusting vmesh and VBC size on-device aftewards (shrink only)

Definition at line 341 of file spatial_cell_gpu_kernels.hpp.

Here is the call graph for this function:

Here is the caller graph for this function:

resize_vbc_kernel_pre()

__global__ void resize_vbc_kernel_pre	(	vmesh::VelocityMesh *	vmesh,
		vmesh::VelocityBlockContainer *	blockContainer,
		const split::SplitVector< vmesh::GlobalID > *__restrict__	list_with_replace_new,
		const split::SplitVector< Hashinator::hash_pair< vmesh::GlobalID, vmesh::LocalID > > *__restrict__	list_delete,
		const split::SplitVector< Hashinator::hash_pair< vmesh::GlobalID, vmesh::LocalID > > *__restrict__	list_to_replace,
		const split::SplitVector< Hashinator::hash_pair< vmesh::GlobalID, vmesh::LocalID > > *__restrict__	list_with_replace_old,
		vmesh::LocalID *	returnLID,
		Realf *	gpu_rhoLossAdjust )

Mini-kernel for checking list sizes and attempting to adjust vmesh and VBC size on-device

Definition at line 305 of file spatial_cell_gpu_kernels.hpp.

Here is the call graph for this function:

Here is the caller graph for this function:

update_neighbour_halo_kernel()

__global__ void update_neighbour_halo_kernel	(	const vmesh::VelocityMesh *__restrict__	vmesh,
		const uint	neighbour_count,
		const vmesh::GlobalID *__restrict__ const *	dev_neigh_vbwcls,
		const vmesh::LocalID *__restrict__	dev_neigh_Nvbwcls,
		Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > *	dev_velocity_block_with_content_map,
		Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > *	dev_velocity_block_with_no_content_map )

Gpu Kernel to quickly gather the spatial halo of neighbour content blocks

Definition at line 198 of file spatial_cell_gpu_kernels.hpp.

Here is the call graph for this function:

update_velocity_halo_kernel()

__global__ void update_velocity_halo_kernel	(	const vmesh::VelocityMesh *__restrict__	vmesh,
		const vmesh::LocalID	velocity_block_with_content_list_size,
		const vmesh::GlobalID *__restrict__	velocity_block_with_content_list_data,
		Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > *	dev_velocity_block_with_content_map,
		Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > *	dev_velocity_block_with_no_content_map )

Gpu Kernel to quickly gather the v-space halo of local content blocks Halo of 1 in each direction adds up to 26 neighbours. For NVIDIA/CUDA, we dan do 26 neighbours and 32 threads per warp in a single block. For AMD/HIP, we dan do 13 neighbours and 64 threads per warp in a single block, meaning two loops per cell. In either case, we launch blocks equal to velocity_block_with_content_list_size

Definition at line 123 of file spatial_cell_gpu_kernels.hpp.

Here is the call graph for this function:

Here is the caller graph for this function:

update_vmesh_and_blockparameters_kernel()

__global__ void update_vmesh_and_blockparameters_kernel	(	vmesh::VelocityMesh *	dev_vmesh,
		vmesh::VelocityBlockContainer *	dev_blockContainer,
		const vmesh::LocalID	nLIDs )

GPU kernel for resizing vmesh GlobalToLocalMap and blockcontainer based on vmesh localToGlobalMap contents (and size). Also populates block parameters. Assumes capacity is already large enough.

Definition at line 262 of file spatial_cell_gpu_kernels.hpp.

Here is the call graph for this function: