block_adjust_gpu_kernels.hpp File Reference

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

Go to the source code of this file.

Macros
#define	warpsPerBlockBatchContent   WID3/GPUTHREADS

Functions
__global__ void	__launch_bounds__ (WID3, WID3S_PER_MP) batch_update_velocity_block_content_lists_kernel(const vmesh
__global__ void	__launch_bounds__ (Hashinator::defaults::MAX_BLOCKSIZE, FULLBLOCKS_PER_MP) batch_reset_all_to_empty(Hashinator
template<typename Rule, typename ELEMENT, bool FIRSTONLY = false>
__global__ void	__launch_bounds__ (Hashinator::defaults::MAX_BLOCKSIZE, FULLBLOCKS_PER_MP) extract_GIDs_kernel(const Hashinator
template<typename Rule, typename ELEMENT, bool FIRSTONLY = false>
void	extract_GIDs_kernel_launcher (Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > **input_maps, split::SplitVector< ELEMENT > **output_vecs, vmesh::LocalID *output_sizes, Rule rule, vmesh::VelocityMesh **rule_meshes, Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > **rule_maps, split::SplitVector< vmesh::GlobalID > **rule_vectors, const uint nCells, gpuStream_t stream)
template<typename Rule>
__global__ void	__launch_bounds__ (Hashinator::defaults::MAX_BLOCKSIZE, FULLBLOCKS_PER_MP) extract_overflown_kernel(vmesh
template<typename Rule>
void	clean_tombstones_launcher (vmesh::VelocityMesh **vmeshes, split::SplitVector< Hashinator::hash_pair< vmesh::GlobalID, vmesh::LocalID > > **overflown_elements, vmesh::LocalID *output_sizes, Rule rule, const uint nCells, gpuStream_t stream)
__global__ void	__launch_bounds__ (GPUTHREADS, WARPS_PER_MP) batch_insert_kernel(vmesh
__global__ void	batch_update_velocity_halo_kernel (const vmesh::VelocityMesh *__restrict__ const *vmeshes, const split::SplitVector< vmesh::GlobalID > *__restrict__ const *velocity_block_with_content_lists, Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > **allMaps, const uint warpsPerBlockBatchHalo)
__global__ void	__launch_bounds__ (GPUTHREADS *WARPSPERBLOCK, FULLBLOCKS_PER_MP) batch_update_neighbour_halo_kernel(const vmesh
__global__ void	batch_resize_vbc_kernel_pre (vmesh::VelocityMesh **vmeshes, vmesh::VelocityBlockContainer **blockContainers, split::SplitVector< vmesh::GlobalID > **dev_list_with_replace_new, split::SplitVector< Hashinator::hash_pair< vmesh::GlobalID, vmesh::LocalID > > **dev_list_delete, split::SplitVector< Hashinator::hash_pair< vmesh::GlobalID, vmesh::LocalID > > **dev_list_to_replace, split::SplitVector< Hashinator::hash_pair< vmesh::GlobalID, vmesh::LocalID > > **dev_list_with_replace_old, vmesh::LocalID *dev_nBefore, vmesh::LocalID *dev_nAfter, vmesh::LocalID *dev_nBlocksToChange, vmesh::LocalID *dev_resizeSuccess, Real *dev_rhoLossAdjust)
__global__ void	batch_resize_vbc_kernel_post (vmesh::VelocityMesh **vmeshes, vmesh::VelocityBlockContainer **blockContainers, vmesh::LocalID *dev_nAfter)

Macro Definition Documentation

warpsPerBlockBatchContent

#define warpsPerBlockBatchContent   WID3/GPUTHREADS

Function Documentation

__launch_bounds__() [1/6]

__global__ void __launch_bounds__	(	GPUTHREADS *	WARPSPERBLOCK,
		FULLBLOCKS_PER_MP	) const

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

Definition at line 768 of file block_adjust_gpu_kernels.hpp.

__launch_bounds__() [2/6]

__global__ void __launch_bounds__	(	GPUTHREADS	,
		WARPS_PER_MP	)

Definition at line 488 of file block_adjust_gpu_kernels.hpp.

__launch_bounds__() [3/6]

__global__ void __launch_bounds__	(	Hashinator::defaults::MAX_BLOCKSIZE	,
		FULLBLOCKS_PER_MP	)

Definition at line 165 of file block_adjust_gpu_kernels.hpp.

__launch_bounds__() [4/6]

template<typename Rule, typename ELEMENT, bool FIRSTONLY = false>

__global__ void __launch_bounds__	(	Hashinator::defaults::MAX_BLOCKSIZE	,
		FULLBLOCKS_PER_MP	) const

Definition at line 212 of file block_adjust_gpu_kernels.hpp.

Here is the call graph for this function:

__launch_bounds__() [5/6]

template<typename Rule>

__global__ void __launch_bounds__	(	Hashinator::defaults::MAX_BLOCKSIZE	,
		FULLBLOCKS_PER_MP	)

Definition at line 358 of file block_adjust_gpu_kernels.hpp.

Here is the call graph for this function:

__launch_bounds__() [6/6]

__global__ void __launch_bounds__	(	WID3	,
		WID3S_PER_MP	) const

GPU kernel for identifying which blocks have relevant content

GPU kernel for updating blocks based on generated lists

GPU kernel for batch-scaling particle populations

Definition at line 43 of file block_adjust_gpu_kernels.hpp.

Here is the call graph for this function:

batch_resize_vbc_kernel_post()

__global__ void batch_resize_vbc_kernel_post	(	vmesh::VelocityMesh **	vmeshes,
		vmesh::VelocityBlockContainer **	blockContainers,
		vmesh::LocalID *	dev_nAfter )

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

Definition at line 879 of file block_adjust_gpu_kernels.hpp.

Here is the call graph for this function:

Here is the caller graph for this function:

batch_resize_vbc_kernel_pre()

__global__ void batch_resize_vbc_kernel_pre	(	vmesh::VelocityMesh **	vmeshes,
		vmesh::VelocityBlockContainer **	blockContainers,
		split::SplitVector< vmesh::GlobalID > **	dev_list_with_replace_new,
		split::SplitVector< Hashinator::hash_pair< vmesh::GlobalID, vmesh::LocalID > > **	dev_list_delete,
		split::SplitVector< Hashinator::hash_pair< vmesh::GlobalID, vmesh::LocalID > > **	dev_list_to_replace,
		split::SplitVector< Hashinator::hash_pair< vmesh::GlobalID, vmesh::LocalID > > **	dev_list_with_replace_old,
		vmesh::LocalID *	dev_nBefore,
		vmesh::LocalID *	dev_nAfter,
		vmesh::LocalID *	dev_nBlocksToChange,
		vmesh::LocalID *	dev_resizeSuccess,
		Real *	dev_rhoLossAdjust )

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

Definition at line 828 of file block_adjust_gpu_kernels.hpp.

Here is the call graph for this function:

Here is the caller graph for this function:

batch_update_velocity_halo_kernel()

__global__ void batch_update_velocity_halo_kernel	(	const vmesh::VelocityMesh *__restrict__ const *	vmeshes,
		const split::SplitVector< vmesh::GlobalID > *__restrict__ const *	velocity_block_with_content_lists,
		Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > **	allMaps,
		const uint	warpsPerBlockBatchHalo )

Gpu Kernel to quickly gather the v-space halo of local content blocks Halo of 1 in each direction adds up to 26 neighbours. This kernel does not use warp accessors so always does all 26 neighbors in a single block.

Definition at line 627 of file block_adjust_gpu_kernels.hpp.

clean_tombstones_launcher()

template<typename Rule>

void clean_tombstones_launcher	(	vmesh::VelocityMesh **	vmeshes,
		split::SplitVector< Hashinator::hash_pair< vmesh::GlobalID, vmesh::LocalID > > **	overflown_elements,
		vmesh::LocalID *	output_sizes,
		Rule	rule,
		const uint	nCells,
		gpuStream_t	stream )

Definition at line 467 of file block_adjust_gpu_kernels.hpp.

extract_GIDs_kernel_launcher()

template<typename Rule, typename ELEMENT, bool FIRSTONLY = false>

void extract_GIDs_kernel_launcher	(	Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > **	input_maps,
		split::SplitVector< ELEMENT > **	output_vecs,
		vmesh::LocalID *	output_sizes,
		Rule	rule,
		vmesh::VelocityMesh **	rule_meshes,
		Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > **	rule_maps,
		split::SplitVector< vmesh::GlobalID > **	rule_vectors,
		const uint	nCells,
		gpuStream_t	stream )

Definition at line 328 of file block_adjust_gpu_kernels.hpp.

Here is the caller graph for this function: