#include "../grid.h"
#include "../object_wrapper.h"
#include "../memoryallocation.h"
#include "gpu_1d_ppm_nonuniform.hpp"
#include "gpu_trans_map_amr.hpp"
#include "cpu_trans_pencils.hpp"
#include "../arch/gpu_base.hpp"

Include dependency graph for gpu_trans_map_amr.cpp:

Functions
__device__ bool	check_skip_blocks (const Realf __restrict__ const pencilBlockData, const uint centerOffset)
__global__ void	__launch_bounds__ (WID3) translation_kernel(const uint dimension
	switch (dimension)
	for (uint thisBlockIndex=startingBlockIndex+blockIdx.x;thisBlockIndex< nAllBlocks;thisBlockIndex+=blockIndexIncrement)
__global__ void	__launch_bounds__ (GPUTHREADS *WARPSPERBLOCK) gather_union_of_blocks_kernel(Hashinator
bool	trans_map_1d_amr (const dccrg::Dccrg< SpatialCell, dccrg::Cartesian_Geometry > &mpiGrid, const vector< CellID > &localPropagatedCells, const vector< CellID > &remoteTargetCells, std::vector< uint > &nPencilsLB, const uint dimension, const Realf dt, const uint popID)
int	get_sibling_index (dccrg::Dccrg< SpatialCell, dccrg::Cartesian_Geometry > &mpiGrid, const CellID &cellid)
static __global__ void	remote_increment_kernel (Realf blockData, Realf neighborData, vmesh::LocalID nBlocks)
void	update_remote_mapping_contribution_amr (dccrg::Dccrg< SpatialCell, dccrg::Cartesian_Geometry > &mpiGrid, const uint dimension, int direction, const uint popID)

Variables
__global__ void const Realf	dt
__global__ void const Realf const uint *__restrict__	pencilLengths
__global__ void const Realf const uint __restrict__ const uint __restrict__	pencilStarts
__global__ void const Realf const uint __restrict__ const uint __restrict__ const vmesh::GlobalID *__restrict__	allBlocks
__global__ void const Realf const uint __restrict__ const uint __restrict__ const vmesh::GlobalID *__restrict__ const uint	nAllBlocks
__global__ void const Realf const uint __restrict__ const uint __restrict__ const vmesh::GlobalID *__restrict__ const uint const uint	nPencils
__global__ void const Realf const uint __restrict__ const uint __restrict__ const vmesh::GlobalID *__restrict__ const uint const uint const uint	sumOfLengths
__global__ void const Realf const uint __restrict__ const uint __restrict__ const vmesh::GlobalID *__restrict__ const uint const uint const uint const Realf	threshold
__global__ void const Realf const uint __restrict__ const uint __restrict__ const vmesh::GlobalID __restrict__ const uint const uint const uint const Realf const vmesh::VelocityMesh __restrict__ const *	dev_allPencilsMeshes
__global__ void const Realf const uint __restrict__ const uint __restrict__ const vmesh::GlobalID __restrict__ const uint const uint const uint const Realf const vmesh::VelocityMesh __restrict__ const vmesh::VelocityBlockContainer **	dev_allPencilsContainers
__global__ void const Realf const uint __restrict__ const uint __restrict__ const vmesh::GlobalID __restrict__ const uint const uint const uint const Realf const vmesh::VelocityMesh __restrict__ const vmesh::VelocityBlockContainer Realf **	pencilBlockData
__global__ void const Realf const uint __restrict__ const uint __restrict__ const vmesh::GlobalID __restrict__ const uint const uint const uint const Realf const vmesh::VelocityMesh __restrict__ const vmesh::VelocityBlockContainer Realf Realf **	dev_blockDataOrdered
__global__ void const Realf const uint __restrict__ const uint __restrict__ const vmesh::GlobalID __restrict__ const uint const uint const uint const Realf const vmesh::VelocityMesh __restrict__ const vmesh::VelocityBlockContainer Realf Realf const Realf *__restrict__	pencilDZ
__global__ void const Realf const uint __restrict__ const uint __restrict__ const vmesh::GlobalID __restrict__ const uint const uint const uint const Realf const vmesh::VelocityMesh __restrict__ const vmesh::VelocityBlockContainer Realf Realf const Realf __restrict__ const Realf __restrict__	pencilRatios
__global__ void const Realf const uint __restrict__ const uint __restrict__ const vmesh::GlobalID __restrict__ const uint const uint const uint const Realf const vmesh::VelocityMesh __restrict__ const vmesh::VelocityBlockContainer Realf Realf const Realf __restrict__ const Realf __restrict__ uint *	pencilBlocksCount
__global__ void const Realf const uint __restrict__ const uint __restrict__ const vmesh::GlobalID __restrict__ const uint const uint const uint const Realf const vmesh::VelocityMesh __restrict__ const vmesh::VelocityBlockContainer Realf Realf const Realf __restrict__ const Realf __restrict__ uint uint *	dev_pencilsInBin
__global__ void const Realf const uint __restrict__ const uint __restrict__ const vmesh::GlobalID __restrict__ const uint const uint const uint const Realf const vmesh::VelocityMesh __restrict__ const vmesh::VelocityBlockContainer Realf Realf const Realf __restrict__ const Realf __restrict__ uint uint uint *	dev_binStart
__global__ void const Realf const uint __restrict__ const uint __restrict__ const vmesh::GlobalID __restrict__ const uint const uint const uint const Realf const vmesh::VelocityMesh __restrict__ const vmesh::VelocityBlockContainer Realf Realf const Realf __restrict__ const Realf __restrict__ uint uint uint uint *	dev_binSize
__global__ void const Realf const uint __restrict__ const uint __restrict__ const vmesh::GlobalID __restrict__ const uint const uint const uint const Realf const vmesh::VelocityMesh __restrict__ const vmesh::VelocityBlockContainer Realf Realf const Realf __restrict__ const Realf __restrict__ uint uint uint uint const uint	numberOfBins
const uint	blockIndexIncrement = gridDim.y*gridDim.x
Realf *	pencilOrderedSource = dev_blockDataOrdered[blockIdx.y]
const vmesh::LocalID	ti = (threadIdx.z)blockDim.xblockDim.y + threadIdx.y*blockDim.x + threadIdx.x
uint	vz_index
const uint	pencilBlockDataOffset = (blockIdx.x * sumOfLengths) + (blockIdx.y * sumOfLengths * gridDim.x)
const uint	pencilOrderedSourceOffset = blockIdx.x * sumOfLengths * WID3
const uint	pencilBlocksCountOffset = (blockIdx.x * nPencils) + (blockIdx.y * nPencils * gridDim.x)
const vmesh::VelocityMesh *__restrict__	randovmesh = dev_allPencilsMeshes[0]
const Realf	dvz = randovmesh->getCellSize()[dimension]
const Realf	vz_min = randovmesh->getMeshMinLimits()[dimension]

Function Documentation

◆ __launch_bounds__() [1/2]

__global__ void __launch_bounds__ ( GPUTHREADS * WARPSPERBLOCK )

Definition at line 340 of file gpu_trans_map_amr.cpp.

◆ __launch_bounds__() [2/2]

__global__ void __launch_bounds__ ( WID3 ) const

◆ check_skip_blocks()

__device__ bool check_skip_blocks	(	const Realf __restrict__ const	pencilBlockData,
		const uint	centerOffset )

inline

Definition at line 39 of file gpu_trans_map_amr.cpp.

Here is the caller graph for this function:

◆ for()

for ( )

Definition at line 133 of file gpu_trans_map_amr.cpp.

Here is the call graph for this function:

◆ get_sibling_index()

int get_sibling_index	(	dccrg::Dccrg< SpatialCell, dccrg::Cartesian_Geometry > &	mpiGrid,
		const CellID &	cellid )

Definition at line 612 of file gpu_trans_map_amr.cpp.

Here is the caller graph for this function:

◆ remote_increment_kernel()

__global__ void remote_increment_kernel	(	Realf *	blockData,
		Realf *	neighborData,
		vmesh::LocalID	nBlocks )

static

GPU kernel for incrementing one set of blocks with another used for accruing remote neighbor contributions

Definition at line 644 of file gpu_trans_map_amr.cpp.

Here is the caller graph for this function:

◆ switch()

switch ( dimension )

Definition at line 109 of file gpu_trans_map_amr.cpp.

◆ trans_map_1d_amr()

bool trans_map_1d_amr	(	const dccrg::Dccrg< SpatialCell, dccrg::Cartesian_Geometry > &	mpiGrid,
		const vector< CellID > &	localPropagatedCells,
		const vector< CellID > &	remoteTargetCells,
		std::vector< uint > &	nPencilsLB,
		const uint	dimension,
		const Realf	dt,
		const uint	popID )

Definition at line 372 of file gpu_trans_map_amr.cpp.

Here is the call graph for this function:

◆ update_remote_mapping_contribution_amr()

void update_remote_mapping_contribution_amr	(	dccrg::Dccrg< SpatialCell, dccrg::Cartesian_Geometry > &	mpiGrid,
		const uint	dimension,
		int	direction,
		const uint	popID )

GPUTODO: First attempts at using unified memory for remote neighbours Should probably: a) Switch to using pure device buffers. Some MPI-CUDA implementations do not work well with UB buffers. b) transition to re-using buffers and ensuring size is suitable? If that path is taken, it should also check for any local cells not on process boundary (anymore, due to LB) and free the buffers from those cells... so gets tricky.

Definition at line 673 of file gpu_trans_map_amr.cpp.

Here is the call graph for this function:

Variable Documentation

◆ allBlocks

__global__ void const Realf const uint* __restrict__ const uint* __restrict__ const vmesh::GlobalID* __restrict__ allBlocks

Definition at line 80 of file gpu_trans_map_amr.cpp.

◆ blockIndexIncrement

const uint blockIndexIncrement = gridDim.y*gridDim.x

Definition at line 101 of file gpu_trans_map_amr.cpp.

◆ dev_allPencilsContainers

__global__ void const Realf const uint* __restrict__ const uint* __restrict__ const vmesh::GlobalID* __restrict__ const uint const uint const uint const Realf const vmesh::VelocityMesh* __restrict__ const vmesh::VelocityBlockContainer* * dev_allPencilsContainers

Definition at line 86 of file gpu_trans_map_amr.cpp.

◆ dev_allPencilsMeshes

__global__ void const Realf const uint* __restrict__ const uint* __restrict__ const vmesh::GlobalID* __restrict__ const uint const uint const uint const Realf const vmesh::VelocityMesh* __restrict__ const* dev_allPencilsMeshes

Definition at line 85 of file gpu_trans_map_amr.cpp.

◆ dev_binSize

__global__ void const Realf const uint* __restrict__ const uint* __restrict__ const vmesh::GlobalID* __restrict__ const uint const uint const uint const Realf const vmesh::VelocityMesh* __restrict__ const vmesh::VelocityBlockContainer Realf Realf const Realf* __restrict__ const Realf* __restrict__ uint uint uint uint* dev_binSize

Definition at line 94 of file gpu_trans_map_amr.cpp.

◆ dev_binStart

__global__ void const Realf const uint* __restrict__ const uint* __restrict__ const vmesh::GlobalID* __restrict__ const uint const uint const uint const Realf const vmesh::VelocityMesh* __restrict__ const vmesh::VelocityBlockContainer Realf Realf const Realf* __restrict__ const Realf* __restrict__ uint uint uint* dev_binStart

Definition at line 93 of file gpu_trans_map_amr.cpp.

◆ dev_blockDataOrdered

__global__ void const Realf const uint* __restrict__ const uint* __restrict__ const vmesh::GlobalID* __restrict__ const uint const uint const uint const Realf const vmesh::VelocityMesh* __restrict__ const vmesh::VelocityBlockContainer Realf Realf** dev_blockDataOrdered

Definition at line 88 of file gpu_trans_map_amr.cpp.

◆ dev_pencilsInBin

__global__ void const Realf const uint* __restrict__ const uint* __restrict__ const vmesh::GlobalID* __restrict__ const uint const uint const uint const Realf const vmesh::VelocityMesh* __restrict__ const vmesh::VelocityBlockContainer Realf Realf const Realf* __restrict__ const Realf* __restrict__ uint uint* dev_pencilsInBin

Definition at line 92 of file gpu_trans_map_amr.cpp.

◆ dt

__global__ void const Realf dt

Definition at line 77 of file gpu_trans_map_amr.cpp.

◆ dvz

const Realf dvz = randovmesh->getCellSize()[dimension]

Definition at line 129 of file gpu_trans_map_amr.cpp.

◆ nAllBlocks

__global__ void const Realf const uint* __restrict__ const uint* __restrict__ const vmesh::GlobalID* __restrict__ const uint nAllBlocks

Definition at line 81 of file gpu_trans_map_amr.cpp.

◆ nPencils

__global__ void const Realf const uint* __restrict__ const uint* __restrict__ const vmesh::GlobalID* __restrict__ const uint const uint nPencils

Definition at line 82 of file gpu_trans_map_amr.cpp.

◆ numberOfBins

__global__ void const Realf const uint* __restrict__ const uint* __restrict__ const vmesh::GlobalID* __restrict__ const uint const uint const uint const Realf const vmesh::VelocityMesh* __restrict__ const vmesh::VelocityBlockContainer Realf Realf const Realf* __restrict__ const Realf* __restrict__ uint uint uint uint const uint numberOfBins

Initial value:

{
   
   const uint startingBlockIndex = blockIdx.y*gridDim.x

Definition at line 95 of file gpu_trans_map_amr.cpp.

◆ pencilBlockData

__global__ void const Realf const uint* __restrict__ const uint* __restrict__ const vmesh::GlobalID* __restrict__ const uint const uint const uint const Realf const vmesh::VelocityMesh* __restrict__ const vmesh::VelocityBlockContainer Realf** pencilBlockData

Definition at line 87 of file gpu_trans_map_amr.cpp.

◆ pencilBlockDataOffset

const uint pencilBlockDataOffset = (blockIdx.x * sumOfLengths) + (blockIdx.y * sumOfLengths * gridDim.x)

Definition at line 125 of file gpu_trans_map_amr.cpp.

◆ pencilBlocksCount

__global__ void const Realf const uint* __restrict__ const uint* __restrict__ const vmesh::GlobalID* __restrict__ const uint const uint const uint const Realf const vmesh::VelocityMesh* __restrict__ const vmesh::VelocityBlockContainer Realf Realf const Realf* __restrict__ const Realf* __restrict__ uint* pencilBlocksCount

Definition at line 91 of file gpu_trans_map_amr.cpp.

◆ pencilBlocksCountOffset

const uint pencilBlocksCountOffset = (blockIdx.x * nPencils) + (blockIdx.y * nPencils * gridDim.x)

Definition at line 127 of file gpu_trans_map_amr.cpp.

◆ pencilDZ

__global__ void const Realf const uint* __restrict__ const uint* __restrict__ const vmesh::GlobalID* __restrict__ const uint const uint const uint const Realf const vmesh::VelocityMesh* __restrict__ const vmesh::VelocityBlockContainer Realf Realf const Realf* __restrict__ pencilDZ

Definition at line 89 of file gpu_trans_map_amr.cpp.

◆ pencilLengths

__global__ void const Realf const uint* __restrict__ pencilLengths

Definition at line 78 of file gpu_trans_map_amr.cpp.

◆ pencilOrderedSource

Realf* pencilOrderedSource = dev_blockDataOrdered[blockIdx.y]

Definition at line 102 of file gpu_trans_map_amr.cpp.

◆ pencilOrderedSourceOffset

const uint pencilOrderedSourceOffset = blockIdx.x * sumOfLengths * WID3

Definition at line 126 of file gpu_trans_map_amr.cpp.

◆ pencilRatios

__global__ void const Realf const uint* __restrict__ const uint* __restrict__ const vmesh::GlobalID* __restrict__ const uint const uint const uint const Realf const vmesh::VelocityMesh* __restrict__ const vmesh::VelocityBlockContainer Realf Realf const Realf* __restrict__ const Realf* __restrict__ pencilRatios

Definition at line 90 of file gpu_trans_map_amr.cpp.

◆ pencilStarts

__global__ void const Realf const uint* __restrict__ const uint* __restrict__ pencilStarts

Definition at line 79 of file gpu_trans_map_amr.cpp.

◆ randovmesh

const vmesh::VelocityMesh* __restrict__ randovmesh = dev_allPencilsMeshes[0]

Definition at line 128 of file gpu_trans_map_amr.cpp.

◆ sumOfLengths

__global__ void const Realf const uint* __restrict__ const uint* __restrict__ const vmesh::GlobalID* __restrict__ const uint const uint const uint sumOfLengths

Definition at line 83 of file gpu_trans_map_amr.cpp.

◆ threshold

__global__ void const Realf const uint* __restrict__ const uint* __restrict__ const vmesh::GlobalID* __restrict__ const uint const uint const uint const Realf threshold

Definition at line 84 of file gpu_trans_map_amr.cpp.

◆ ti

const vmesh::LocalID ti = (threadIdx.z)*blockDim.x*blockDim.y + threadIdx.y*blockDim.x + threadIdx.x

Definition at line 105 of file gpu_trans_map_amr.cpp.

◆ vz_index

uint vz_index

Definition at line 108 of file gpu_trans_map_amr.cpp.

◆ vz_min

const Realf vz_min = randovmesh->getMeshMinLimits()[dimension]

Definition at line 130 of file gpu_trans_map_amr.cpp.

Functions

Variables

Function Documentation

◆ __launch_bounds__() [1/2]

◆ __launch_bounds__() [2/2]

◆ check_skip_blocks()

◆ for()

◆ get_sibling_index()

◆ remote_increment_kernel()

◆ switch()

◆ trans_map_1d_amr()

◆ update_remote_mapping_contribution_amr()

Variable Documentation

◆ allBlocks

◆ blockIndexIncrement

◆ dev_allPencilsContainers

◆ dev_allPencilsMeshes

◆ dev_binSize

◆ dev_binStart

◆ dev_blockDataOrdered

◆ dev_pencilsInBin

◆ dt

◆ dvz

◆ nAllBlocks

◆ nPencils

◆ numberOfBins

◆ pencilBlockData

◆ pencilBlockDataOffset

◆ pencilBlocksCount

◆ pencilBlocksCountOffset

◆ pencilDZ

◆ pencilLengths

◆ pencilOrderedSource

◆ pencilOrderedSourceOffset

◆ pencilRatios

◆ pencilStarts

◆ randovmesh

◆ sumOfLengths

◆ threshold

◆ ti

◆ vz_index

◆ vz_min