spatial_cell Namespace Reference

Namespaces
namespace	Transfer

Classes
struct	Population
class	SpatialCell

Typedefs
typedef std::array< unsigned int, 3 >	velocity_cell_indices_t
typedef std::array< vmesh::LocalID, 3 >	velocity_block_indices_t

Functions
void	update_velocity_block_content_lists (dccrg::Dccrg< SpatialCell, dccrg::Cartesian_Geometry > &mpiGrid, const vector< CellID > &cells, const uint popID)
void	adjust_velocity_blocks_in_cells (dccrg::Dccrg< spatial_cell::SpatialCell, dccrg::Cartesian_Geometry > &mpiGrid, const vector< CellID > &cellsToAdjust, const uint popID)
void	update_velocity_block_content_lists (dccrg::Dccrg< spatial_cell::SpatialCell, dccrg::Cartesian_Geometry > &mpiGrid, const vector< CellID > &cells, const uint popID)
	spatial_cell::update_velocity_block_content_lists Finds blocks above the sparsity threshold
void	clear_maps_caller (const uint nCells, const size_t largestSizePower, gpuStream_t stream, const size_t offset)
void	batch_adjust_blocks_caller (dccrg::Dccrg< spatial_cell::SpatialCell, dccrg::Cartesian_Geometry > &mpiGrid, const vector< CellID > &cellsToAdjust, const uint cellOffset, uint &out_largestBlocksToChange, uint &out_largestBlocksBeforeOrAfter, const uint popID)
void	extract_to_replace_caller (Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > **input_maps, split::SplitVector< Hashinator::hash_pair< vmesh::GlobalID, vmesh::LocalID > > **output_vecs, vmesh::LocalID *output_sizes, vmesh::VelocityMesh **rule_meshes, Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > **rule_maps, split::SplitVector< vmesh::GlobalID > **rule_vectors, const uint nCells, gpuStream_t stream)
void	extract_to_delete_or_move_caller (Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > **input_maps, split::SplitVector< Hashinator::hash_pair< vmesh::GlobalID, vmesh::LocalID > > **output_vecs, vmesh::LocalID *output_sizes, vmesh::VelocityMesh **rule_meshes, Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > **rule_maps, split::SplitVector< vmesh::GlobalID > **rule_vectors, const uint nCells, gpuStream_t stream)
void	extract_to_add_caller (Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > **input_maps, split::SplitVector< vmesh::GlobalID > **output_vecs, vmesh::LocalID *output_sizes, vmesh::VelocityMesh **rule_meshes, Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > **rule_maps, split::SplitVector< vmesh::GlobalID > **rule_vectors, const uint nCells, gpuStream_t stream)
static __global__ void	resize_vmesh_ondevice_kernel (vmesh::VelocityMesh *vmesh, vmesh::LocalID nBlocks)
static __global__ void	__launch_bounds__ (WID3, 4) population_scale_kernel(vmesh
static __global__ void	resize_and_empty_kernel (vmesh::VelocityMesh *vmesh, vmesh::VelocityBlockContainer *blockContainer, vmesh::LocalID newSize)
template<typename fileReal>
__global__ void	add_blocks_from_buffer_kernel (const vmesh::VelocityMesh *vmesh, vmesh::VelocityBlockContainer *blockContainer, const vmesh::LocalID startLID, const vmesh::GlobalID *gpuInitBlocks, const fileReal *gpuInitBuffer, const uint nBlocks)

Detailed Description

Spatial cell class for Vlasiator that supports a variable number of velocity blocks.

Typedef Documentation

velocity_block_indices_t

typedef std::array< vmesh::LocalID, 3 > spatial_cell::velocity_block_indices_t

Defines the indices of a velocity block in the velocity grid. Indices start from 0 and the first value is the index in x direction. Note: these are the (i,j,k) indices of the block. Valid values are ([0,vx_length[,[0,vy_length[,[0,vz_length[).

Definition at line 230 of file spatial_cell_cpu.hpp.

velocity_cell_indices_t

typedef std::array< unsigned int, 3 > spatial_cell::velocity_cell_indices_t

Defines the indices of a velocity cell in a velocity block. Indices start from 0 and the first value is the index in x direction. Note: these are the (i,j,k) indices of the cell within the block. Valid values are ([0,WID[,[0,WID[,[0,WID[).

Definition at line 225 of file spatial_cell_cpu.hpp.

Function Documentation

__launch_bounds__()

__global__ void spatial_cell::__launch_bounds__ ( WID3 , 4  )

static

GPU kernel for scaling a particle population

GPU kernel for adding a particle population to another with a scaling factor This kernel increments existing blocks, creates new ones if the block does not exist, and so is not parallel-safe.

GPU kernel for replacing one particle population with another

Definition at line 69 of file spatial_cell_gpu.hpp.

add_blocks_from_buffer_kernel()

template<typename fileReal>

__global__ void spatial_cell::add_blocks_from_buffer_kernel	(	const vmesh::VelocityMesh *	vmesh,
		vmesh::VelocityBlockContainer *	blockContainer,
		const vmesh::LocalID	startLID,
		const vmesh::GlobalID *	gpuInitBlocks,
		const fileReal *	gpuInitBuffer,
		const uint	nBlocks )

GPU kernel for populating block data and parameters based on list of globalIDs and avgs.

Definition at line 510 of file spatial_cell_gpu.hpp.

Here is the call graph for this function:

adjust_velocity_blocks_in_cells()

void spatial_cell::adjust_velocity_blocks_in_cells	(	dccrg::Dccrg< spatial_cell::SpatialCell, dccrg::Cartesian_Geometry > &	mpiGrid,
		const vector< CellID > &	cellsToAdjust,
		const uint	popID )

Now extract vectors to be used in actual block adjustment. Previous kernels may have added dummy (to be added) entries to velocity_block_with_content_map with LID=vmesh->invalidLocalID() Or non-content blocks which should be retained (with correct LIDs).

Note: These batch operations always include deletion of not-needed blocks.

Rules used in extracting keys or elements from hashmaps: These are provided with the value of nBlocksAfterAdjust as the threshold argument by the kernel. To this end, rule_meshes, rule_maps, and rule_vectors pointer buffers are provided to the kernels.

Definition at line 55 of file block_adjust_cpu.cpp.

Here is the call graph for this function:

Here is the caller graph for this function:

batch_adjust_blocks_caller()

void spatial_cell::batch_adjust_blocks_caller	(	dccrg::Dccrg< spatial_cell::SpatialCell, dccrg::Cartesian_Geometry > &	mpiGrid,
		const vector< CellID > &	cellsToAdjust,
		const uint	cellOffset,
		uint &	out_largestBlocksToChange,
		uint &	out_largestBlocksBeforeOrAfter,
		const uint	popID )

Definition at line 701 of file block_adjust_gpu.cpp.

Here is the call graph for this function:

Here is the caller graph for this function:

clear_maps_caller()

void spatial_cell::clear_maps_caller	(	const uint	nCells,
		const size_t	largestSizePower,
		gpuStream_t	stream,
		const size_t	offset )

Definition at line 682 of file block_adjust_gpu.cpp.

Here is the call graph for this function:

Here is the caller graph for this function:

extract_to_add_caller()

void spatial_cell::extract_to_add_caller	(	Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > **	input_maps,
		split::SplitVector< vmesh::GlobalID > **	output_vecs,
		vmesh::LocalID *	output_sizes,
		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 898 of file block_adjust_gpu.cpp.

Here is the call graph for this function:

extract_to_delete_or_move_caller()

void spatial_cell::extract_to_delete_or_move_caller	(	Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > **	input_maps,
		split::SplitVector< Hashinator::hash_pair< vmesh::GlobalID, vmesh::LocalID > > **	output_vecs,
		vmesh::LocalID *	output_sizes,
		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 862 of file block_adjust_gpu.cpp.

Here is the call graph for this function:

Here is the caller graph for this function:

extract_to_replace_caller()

void spatial_cell::extract_to_replace_caller	(	Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > **	input_maps,
		split::SplitVector< Hashinator::hash_pair< vmesh::GlobalID, vmesh::LocalID > > **	output_vecs,
		vmesh::LocalID *	output_sizes,
		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 824 of file block_adjust_gpu.cpp.

Here is the call graph for this function:

Here is the caller graph for this function:

resize_and_empty_kernel()

__global__ void spatial_cell::resize_and_empty_kernel ( vmesh::VelocityMesh * vmesh, vmesh::VelocityBlockContainer * blockContainer, vmesh::LocalID newSize )

static

GPU kernel for clearing the vmesh globalToLocalMap and setting the localToGlobalMap size to newSize. If the maps do not have sufficient capacity, raises a reallocation flag and returns early. This allows operating on-device without page faults.

Definition at line 215 of file spatial_cell_gpu.hpp.

Here is the call graph for this function:

Here is the caller graph for this function:

resize_vmesh_ondevice_kernel()

__global__ void spatial_cell::resize_vmesh_ondevice_kernel ( vmesh::VelocityMesh * vmesh, vmesh::LocalID nBlocks )

static

GPU mini-kernel for resizing a vmesh on-device

Definition at line 61 of file spatial_cell_gpu.hpp.

Here is the caller graph for this function:

update_velocity_block_content_lists() [1/2]

void spatial_cell::update_velocity_block_content_lists	(	dccrg::Dccrg< spatial_cell::SpatialCell, dccrg::Cartesian_Geometry > &	mpiGrid,
		const vector< CellID > &	cells,
		const uint	popID )

spatial_cell::update_velocity_block_content_lists Finds blocks above the sparsity threshold

Bulk call over listed cells of spatial grid Prepares the content / no-content velocity block lists for all requested cells, for the requested popID

Definition at line 37 of file block_adjust_gpu.cpp.

Here is the call graph for this function:

update_velocity_block_content_lists() [2/2]

void spatial_cell::update_velocity_block_content_lists	(	dccrg::Dccrg< SpatialCell, dccrg::Cartesian_Geometry > &	mpiGrid,
		const vector< CellID > &	cells,
		const uint	popID )

Bulk call over listed cells of spatial grid Prepares the content / no-content velocity block lists for all requested cells, for the requested popID

Definition at line 33 of file block_adjust_cpu.cpp.

Here is the caller graph for this function: