Detailed Description

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT, typename ct_params, typename indexT, template< typename > class layout_base, typename GridSmT, typename linearizer, typename BcT>
class SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >

Definition at line 29 of file SparseGridGpu_ker.cuh.

Inheritance diagram for SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >:

Public Types
typedef AggregateBlockT	AggregateBlockType

typedef indexT	indexT_

typedef int	yes_has_check_device_pointer
	Indicate this structure has a function to check the device pointer.

Public Types inherited from BlockMapGpu_ker< AggregateBlockT, indexT, layout_base >
typedef AggregateBlockT	AggregateType

typedef InsertBlockWrapper< AggregateBlockT, pMask >	InsertBlockWrapperType

Public Member Functions
	SparseGridGpu_ker (const openfpm::vector_sparse_gpu_ker< AggregateBlockT, indexT, layout_base > &blockMap, linearizer &grid, GridSmT extendedBlockGeometry, unsigned int stencilSupportRadius, openfpm::vector_gpu_ker< aggregate< short int, short int >, memory_traits_inte > ghostLayerToThreadsMapping, openfpm::vector_gpu_ker< aggregate< indexT >, memory_traits_inte > nn_blocks, openfpm::vector_gpu_ker< aggregate< indexT >, memory_traits_inte > buffPnt, unsigned int ghostLayerSize, BcT &bck)
	constructor

template<typename CoordT >
__device__ __host__ grid_key_dx< dim, CoordT >	getGlobalCoord (const grid_key_dx< dim, CoordT > &blockCoord, unsigned int offset)
	Get the coordinate of the block and the offset id inside the block it give the global coordinate.

template<typename CoordT >
__device__ size_t	getLinId (const grid_key_dx< dim, CoordT > &coord) const
	Linearization of global coordinates.

__device__ unsigned int	size (unsigned int i)
	Size of the sparse grid in each direction.

__device__ grid_key_dx< dim, int >	getCoord (size_t linId) const
	The inversion of getLinId.

__device__ grid_key_dx< dim, int >	getCoord (size_t dataBlockId, unsigned offset) const
	The inversion of getLinId.

template<typename ite_type >
__device__ bool	getInsertBlockOffset (const ite_type &itd, const grid_key_dx< dim, int > &p, grid_key_dx< dim, int > &blk, int &offset)
	Given a point to insert, return the block-id and offset of that point.

template<typename CoordT >
__device__ size_t	getBlockLinId (CoordT blockCoord) const
	Linearization of block coordinates.

__device__ grid_key_dx< dim, int >	getBlockCoord (size_t blockLinId) const
	The inversion of getBlockLinId.

__device__ grid_key_dx< dim, int >	getBlockBaseCoord (size_t blockLinId) const
	Given a linearized block index it return the coordinated of the lower-left point in 2D or in general the origin of the block in global coordinates.

__device__ grid_key_dx< dim, int >	getNeighbour (grid_key_dx< dim, int > base, unsigned int dimension, char offset) const

constexpr __device__ unsigned int	getBlockSize () const
	Return the size of the block.

__device__ unsigned int	getEnlargedBlockSize () const
	Return the size of the block + ghost needed to apply the stencil.

template<typename NN_type , typename indexT2 >
__device__ block_offset< indexT2 >	getNNPoint (openfpm::sparse_index< unsigned int > pos, unsigned int offset, const grid_key_dx< dim, indexT2 > &mov)
	Get the neighborhood point in one direction.

__device__ unsigned int	posToEnlargedBlockPos (unsigned int pos) const

__device__ grid_key_dx< dim, int >	getCoordInEnlargedBlock (const unsigned int offset) const

__device__ unsigned int	getLinIdInEnlargedBlock (const unsigned int offset) const

template<typename Coordtype >
__device__ unsigned int	getNeighbourLinIdInEnlargedBlock (const grid_key_dx< dim, Coordtype > &base, grid_key_dx< dim, Coordtype > &offsets) const

template<typename Coordtype >
__device__ unsigned int	getNeighbourLinIdInEnlargedBlock (const grid_key_dx< dim, Coordtype > &base, unsigned int dimension, char offset) const

__device__ bool	getIfBoundaryElementInEnlargedBlock (const grid_key_dx< dim, int > coordInEnlargedBlock, char(&boundaryDirection)[dim])

template<unsigned int p, typename CoordT >
__device__ auto	get (const grid_key_dx< dim, CoordT > &coord) const -> ScalarTypeOf< AggregateBlockT, p >

template<typename CoordT >
__device__ void	get_sparse (const grid_key_dx< dim, CoordT > &coord, unsigned int &dataBlockPos, unsigned int &offset) const

template<unsigned int p, typename CoordT >
__device__ auto	get (const block_offset< CoordT > &coord) const -> decltype(std::declval< BlockMapGpu_ker< AggregateBlockT, indexT, layout_base > >().getblockMap().template get_ele< p >(coord.pos)[coord.off])
	Access the grid point.

template<unsigned int p, typename CoordT >
__device__ auto	get (const block_offset< CoordT > &coord) -> decltype(std::declval< BlockMapGpu_ker< AggregateBlockT, indexT, layout_base > >().getblockMap().template get_ele< p >(coord.pos)[coord.off])
	Access the grid point.

template<unsigned int p, typename CoordT >
__device__ auto	insert (const grid_key_dx< dim, CoordT > &coord) -> ScalarTypeOf< AggregateBlockT, p > &

template<typename CoordT >
__device__ unsigned int	getBlockId (const grid_key_dx< dim, CoordT > &coord)

template<typename CoordT >
__device__ unsigned int	getOffset (const grid_key_dx< dim, CoordT > &coord)

template<typename CoordT >
__device__ auto	getBlock (const grid_key_dx< dim, CoordT > &coord) -> decltype(BlockMapGpu_ker< AggregateBlockT, indexT, layout_base >::getBlock(0))

__device__ auto	getBlock (const unsigned int blockLinId) -> decltype(BlockMapGpu_ker< AggregateBlockT, indexT, layout_base >::getBlock(0))

template<unsigned int chunksPerBlocks = 1, typename CoordT >
__device__ auto	insertBlock (const grid_key_dx< dim, CoordT > &coord) -> decltype(BlockMapGpu_ker< AggregateBlockT, indexT, layout_base >::insertBlock(0))

template<unsigned int chunksPerBlocks = 1>
__device__ auto	insertBlock (const indexT blockLinId, const unsigned int stride=8192) -> decltype(BlockMapGpu_ker< AggregateBlockT, indexT, layout_base >::insertBlock(0))

__device__ auto	getPointBuffer () -> decltype(buffPnt) &
	Return the buffer of points.

template<unsigned int p, typename AggrWrapperT >
__device__ void	loadBlock (AggrWrapperT &block, ScalarTypeOf< AggregateBlockT, p > *sharedRegion)

template<unsigned int ... props, typename AggrWrapperT >
__device__ void	loadBlock (AggrWrapperT &block, void *sharedRegionPtr[sizeof...(props)])

template<unsigned int p, typename AggrWrapperT , typename CoordT >
__device__ void	loadGhostBlock (const AggrWrapperT &dataBlockLoad, const grid_key_dx< dim, CoordT > &coord, ScalarTypeOf< AggregateBlockT, p > *sharedRegion)

template<unsigned int p, typename AggrWrapperT >
__device__ void	loadGhostBlock (const AggrWrapperT &dataBlockLoad, const openfpm::sparse_index< unsigned int > blockLinId, ScalarTypeOf< AggregateBlockT, p > *sharedRegion)

template<unsigned int p, typename AggrWrapperT >
__device__ void	loadGhostBlock (const AggrWrapperT &dataBlockLoad, const openfpm::sparse_index< unsigned int > blockLinId, ScalarTypeOf< AggregateBlockT, p > sharedRegion, unsigned char mask)

template<unsigned int ... props, typename CoordT >
__device__ void	loadGhost (const grid_key_dx< dim, CoordT > &coord, const int neighboursPos, void sharedRegionPtr[sizeof...(props)])

template<unsigned int ... props>
__device__ void	loadGhost (const unsigned int blockLinId, const int neighboursPos, void sharedRegionPtr[sizeof...(props)])

template<unsigned int p, typename AggrWrapperT >
__device__ void	storeBlock (AggrWrapperT &block, ScalarTypeOf< AggregateBlockT, p > *sharedRegion)

template<unsigned int p, typename CoordT >
__device__ void	storeBlockInPlace (const grid_key_dx< dim, CoordT > &coord, ScalarTypeOf< AggregateBlockT, p > *sharedRegion)

template<unsigned int ... props, typename AggrWrapperT >
__device__ void	storeBlock (AggrWrapperT &block, void *sharedRegionPtr[sizeof...(props)])

template<unsigned int ... props, typename CoordT >
__device__ void	storeBlockInPlace (const grid_key_dx< dim, CoordT > &coord, void *sharedRegionPtr[sizeof...(props)])

template<unsigned int p, typename CoordT >
__device__ ScalarTypeOf< AggregateBlockT, p > &	get (grid_key_dx< dim, CoordT > coord, Box< dim, indexT > sharedMemBox, ScalarTypeOf< AggregateBlockT, p > *sharedRegion)

template<typename CoordT >
__device__ void	remove (const grid_key_dx< dim, CoordT > &coord)

template<typename keyIndexT >
__device__ bool	isPadding (grid_key_dx< dim, keyIndexT > coord) const

template<typename NNtype >
__device__ indexT	getNeighboursPos (const indexT blockId, const unsigned int offset)

Public Member Functions inherited from BlockMapGpu_ker< AggregateBlockT, indexT, layout_base >
	BlockMapGpu_ker (openfpm::vector_sparse_gpu_ker< AggregateBlockT, indexT, layout_base > blockMap)

__device__ auto	get (unsigned int linId) const -> ScalarTypeOf< AggregateBlockT, p >

__device__ auto	get (unsigned int blockId, unsigned int offset) const -> ScalarTypeOf< AggregateBlockT, p >

__device__ auto	getBlock (unsigned int blockId) -> decltype(blockMap.get(0))

__device__ ScalarTypeOf< AggregateBlockT, p > &	getReference (unsigned int linId)

__device__ ScalarTypeOf< AggregateBlockT, p > &	getReference (unsigned int blockId, unsigned int offset)

__device__ auto	insert (unsigned int linId) -> ScalarTypeOf< AggregateBlockT, p > &

__device__ auto	insert (unsigned int blockId, unsigned int offset) -> ScalarTypeOf< AggregateBlockT, p > &

__device__ auto	insertBlock (indexT blockId, unsigned int stride=8192) -> decltype(blockMap.insert(0))

__device__ openfpm::vector_sparse_gpu_ker< AggregateBlockT, indexT, layout_base > &	getblockMap ()

__device__ void	get_sparse (unsigned int linId, unsigned int &dataBlockPos, unsigned int &offset) const

__device__ void	init ()

__device__ void	flush_block_insert ()

__device__ ScalarTypeOf< AggregateBlockT, pMask >	getMask (unsigned int linId) const

__device__ void	remove (unsigned int linId)

__device__ void	remove (unsigned int blockId, unsigned int offset)

__device__ auto	getIndexBuffer () -> decltype(blockMap.getIndexBuffer())
	Return the index buffer for the sparse vector.

__device__ auto	getDataBuffer () -> decltype(blockMap.getDataBuffer())
	Return the data buffer for the sparse vector.

Static Public Member Functions
template<typename headers_type >
static __device__ int	unpack_headers (headers_type &headers, unsigned char *data, int ih, int sz_pack)


static constexpr __device__ unsigned int	getBlockEdgeSize ()
	Return the size of the block edge size.

template<typename BitMaskT >
static __device__ bool	isPadding (const BitMaskT &bitMask)

template<typename BitMaskT >
static __device__ void	setPadding (BitMaskT &bitMask)

template<typename BitMaskT >
static __device__ void	unsetPadding (BitMaskT &bitMask)

Static Public Member Functions inherited from BlockMapGpu_ker< AggregateBlockT, indexT, layout_base >
static __device__ __host__ bool	getBit (const BitMaskT &bitMask, unsigned char pos)

static __device__ __host__ void	setBit (BitMaskT &bitMask, unsigned char pos)

static __device__ __host__ void	unsetBit (BitMaskT &bitMask, unsigned char pos)

static __device__ unsigned int	getBlockId (unsigned int linId)

static __device__ unsigned int	getOffset (unsigned int linId)

static __device__ bool	exist (const BitMaskT &bitMask)

static __device__ void	setExist (BitMaskT &bitMask)

static __device__ void	unsetExist (BitMaskT &bitMask)

Data Fields
unsigned int	stencilSupportRadius

Static Public Attributes
static constexpr unsigned int	d = dim

static constexpr unsigned int	dims = dim

static constexpr unsigned int	blockEdgeSize_ = blockEdgeSize

Static Public Attributes inherited from BlockMapGpu_ker< AggregateBlockT, indexT, layout_base >
static const unsigned int	pMask

Protected Attributes
unsigned int	ghostLayerSize

openfpm::vector_gpu_ker< aggregate< short int, short int >, memory_traits_inte >	ghostLayerToThreadsMapping

openfpm::vector_gpu_ker< aggregate< indexT >, memory_traits_inte >	nn_blocks

openfpm::vector_gpu_ker< aggregate< indexT >, memory_traits_inte >	buffPnt

Protected Attributes inherited from BlockMapGpu_ker< AggregateBlockT, indexT, layout_base >
openfpm::vector_sparse_gpu_ker< AggregateBlockT, indexT, layout_base >	blockMap

Static Protected Attributes
static const unsigned char	PADDING_BIT = 1

static constexpr unsigned int	blockSize = BlockTypeOf<AggregateBlockT, 0>::size

Static Protected Attributes inherited from BlockMapGpu_ker< AggregateBlockT, indexT, layout_base >
static const unsigned char	EXIST_BIT

Private Member Functions
template<unsigned int p, typename AggrWrapperT , typename SharedPtrT >
__device__ void	__loadBlock (const AggrWrapperT &block, SharedPtrT sharedRegionPtr)

template<unsigned int p, unsigned int ... props, typename AggrWrapperT >
__device__ void	__loadBlock (const AggrWrapperT &block, void *sharedRegionPtr[sizeof...(props)+1])

template<unsigned int p, typename SharedPtrT >
__device__ void	__loadGhostNoNN (const unsigned int blockId, SharedPtrT *sharedRegionPtr)

template<unsigned int p, typename AggrWrapperT , typename SharedPtrT >
__device__ void	__loadGhostBlock (const AggrWrapperT &block, const openfpm::sparse_index< unsigned int > blockId, SharedPtrT *sharedRegionPtr)
	Load the ghost area in the shared region.

template<unsigned int p, typename AggrWrapperT , typename SharedPtrT >
__device__ void	__loadGhostBlock (const AggrWrapperT &block, const openfpm::sparse_index< unsigned int > blockId, SharedPtrT sharedRegionPtr, unsigned char maskPtr)
	Load the ghost area in the shared region.

template<unsigned int p, unsigned int ... props>
__device__ void	__loadGhost (const unsigned int blockId, const int neighboursPos, void sharedRegionPtr[sizeof...(props)+1])

template<unsigned int p, typename AggrWrapperT , typename SharedPtrT >
__device__ void	__storeBlock (AggrWrapperT &block, SharedPtrT sharedRegionPtr)

template<unsigned int p, unsigned int ... props, typename AggrWrapperT >
__device__ void	__storeBlock (AggrWrapperT &block, void *sharedRegionPtr[sizeof...(props)+1])

Private Attributes
linearizer	grid

GridSmT	blockWithGhostGrid

BcT	background
	background values

Member Typedef Documentation

◆ AggregateBlockType

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

typedef AggregateBlockT SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::AggregateBlockType

Definition at line 51 of file SparseGridGpu_ker.cuh.

◆ indexT_

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

typedef indexT SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::indexT_

Definition at line 52 of file SparseGridGpu_ker.cuh.

◆ yes_has_check_device_pointer

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

typedef int SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::yes_has_check_device_pointer

Indicate this structure has a function to check the device pointer.

Definition at line 55 of file SparseGridGpu_ker.cuh.

Constructor & Destructor Documentation

◆ SparseGridGpu_ker()

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::SparseGridGpu_ker	(	const openfpm::vector_sparse_gpu_ker< AggregateBlockT, indexT, layout_base > &	blockMap,
		linearizer &	grid,
		GridSmT	extendedBlockGeometry,
		unsigned int	stencilSupportRadius,
		openfpm::vector_gpu_ker< aggregate< short int, short int >, memory_traits_inte >	ghostLayerToThreadsMapping,
		openfpm::vector_gpu_ker< aggregate< indexT >, memory_traits_inte >	nn_blocks,
		openfpm::vector_gpu_ker< aggregate< indexT >, memory_traits_inte >	buffPnt,
		unsigned int	ghostLayerSize,
		BcT &	bck
	)

inline

constructor

this constructor is in general called in the function toKernel() of SparseGridGpu

Definition at line 64 of file SparseGridGpu_ker.cuh.

Member Function Documentation

◆ __loadBlock() [1/2]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<unsigned int p, typename AggrWrapperT , typename SharedPtrT >

__device__ void SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::__loadBlock	(	const AggrWrapperT &	block,
		SharedPtrT	sharedRegionPtr
	)

inlineprivate

Definition at line 743 of file SparseGridGpu_ker.cuh.

◆ __loadBlock() [2/2]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<unsigned int p, unsigned int ... props, typename AggrWrapperT >

__device__ void SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::__loadBlock	(	const AggrWrapperT &	block,
		void *	sharedRegionPtr[sizeof...(props)+1]
	)

inlineprivate

Definition at line 765 of file SparseGridGpu_ker.cuh.

◆ __loadGhost()

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<unsigned int p, unsigned int ... props>

__device__ void SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::__loadGhost	(	const unsigned int	blockId,
		const int *	neighboursPos,
		void *	sharedRegionPtr[sizeof...(props)+1]
	)

inlineprivate

Definition at line 862 of file SparseGridGpu_ker.cuh.

◆ __loadGhostBlock() [1/2]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<unsigned int p, typename AggrWrapperT , typename SharedPtrT >

__device__ void SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::__loadGhostBlock	(	const AggrWrapperT &	block,
		const openfpm::sparse_index< unsigned int >	blockId,
		SharedPtrT *	sharedRegionPtr
	)

inlineprivate

Load the ghost area in the shared region.

Parameters

blockId	Index of the center block
neighboursPos	neighborhood blocks around blockId (if neighbourPos is null loadGhost)
sharedRegionPtr	shared region

Definition at line 820 of file SparseGridGpu_ker.cuh.

◆ __loadGhostBlock() [2/2]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<unsigned int p, typename AggrWrapperT , typename SharedPtrT >

__device__ void SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::__loadGhostBlock	(	const AggrWrapperT &	block,
		const openfpm::sparse_index< unsigned int >	blockId,
		SharedPtrT *	sharedRegionPtr,
		unsigned char *	maskPtr
	)

inlineprivate

Load the ghost area in the shared region.

Parameters

blockId	Index of the center block
neighboursPos	neighborhood blocks around blockId (if neighbourPos is null loadGhost)
sharedRegionPtr	shared region

Definition at line 844 of file SparseGridGpu_ker.cuh.

◆ __loadGhostNoNN()

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<unsigned int p, typename SharedPtrT >

__device__ void SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::__loadGhostNoNN	(	const unsigned int	blockId,
		SharedPtrT *	sharedRegionPtr
	)

inlineprivate

Definition at line 782 of file SparseGridGpu_ker.cuh.

◆ __storeBlock() [1/2]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<unsigned int p, typename AggrWrapperT , typename SharedPtrT >

__device__ void SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::__storeBlock	(	AggrWrapperT &	block,
		SharedPtrT	sharedRegionPtr
	)

inlineprivate

Definition at line 877 of file SparseGridGpu_ker.cuh.

◆ __storeBlock() [2/2]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<unsigned int p, unsigned int ... props, typename AggrWrapperT >

__device__ void SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::__storeBlock	(	AggrWrapperT &	block,
		void *	sharedRegionPtr[sizeof...(props)+1]
	)

inlineprivate

Definition at line 900 of file SparseGridGpu_ker.cuh.

◆ get() [1/4]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<unsigned int p, typename CoordT >

__device__ auto SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::get ( const block_offset< CoordT > & coord ) -> decltype(std::declval<BlockMapGpu_ker<AggregateBlockT, indexT, layout_base>>().getblockMap().template get_ele<p>(coord.pos)[coord.off])

inline

Access the grid point.

Parameters

coord point

Returns: a reference to the data point

Definition at line 440 of file SparseGridGpu_ker.cuh.

◆ get() [2/4]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<unsigned int p, typename CoordT >

__device__ auto SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::get ( const block_offset< CoordT > & coord ) const -> decltype(std::declval<BlockMapGpu_ker<AggregateBlockT, indexT, layout_base>>().getblockMap().template get_ele<p>(coord.pos)[coord.off])

inline

Access the grid point.

Parameters

coord point

Returns: a reference to the data point

Definition at line 426 of file SparseGridGpu_ker.cuh.

◆ get() [3/4]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<unsigned int p, typename CoordT >

__device__ auto SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::get ( const grid_key_dx< dim, CoordT > & coord ) const -> ScalarTypeOf<AggregateBlockT, p>

inline

Definition at line 403 of file SparseGridGpu_ker.cuh.

◆ get() [4/4]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<unsigned int p, typename CoordT >

__device__ ScalarTypeOf< AggregateBlockT, p > & SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::get	(	grid_key_dx< dim, CoordT >	coord,
		Box< dim, indexT >	sharedMemBox,
		ScalarTypeOf< AggregateBlockT, p > *	sharedRegion
	)

inline

Definition at line 640 of file SparseGridGpu_ker.cuh.

◆ get_sparse()

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<typename CoordT >

__device__ void SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::get_sparse	(	const grid_key_dx< dim, CoordT > &	coord,
		unsigned int &	dataBlockPos,
		unsigned int &	offset
	)		const

inline

Definition at line 412 of file SparseGridGpu_ker.cuh.

◆ getBlock() [1/2]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<typename CoordT >

__device__ auto SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::getBlock ( const grid_key_dx< dim, CoordT > & coord ) -> decltype(BlockMapGpu_ker<AggregateBlockT, indexT, layout_base>::getBlock(0))

inline

Definition at line 467 of file SparseGridGpu_ker.cuh.

◆ getBlock() [2/2]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

__device__ auto SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::getBlock ( const unsigned int blockLinId ) -> decltype(BlockMapGpu_ker<AggregateBlockT, indexT, layout_base>::getBlock(0))

inline

Definition at line 473 of file SparseGridGpu_ker.cuh.

◆ getBlockBaseCoord()

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

__device__ grid_key_dx< dim, int > SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::getBlockBaseCoord ( size_t blockLinId ) const

inline

Given a linearized block index it return the coordinated of the lower-left point in 2D or in general the origin of the block in global coordinates.

Parameters

blockLinId linearized block index

Returns: the block coordinates

Definition at line 243 of file SparseGridGpu_ker.cuh.

◆ getBlockCoord()

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

__device__ grid_key_dx< dim, int > SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::getBlockCoord ( size_t blockLinId ) const

inline

The inversion of getBlockLinId.

Parameters

blockLinId linearized block index

Returns: the block coordinates

Definition at line 230 of file SparseGridGpu_ker.cuh.

◆ getBlockEdgeSize()

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

static constexpr __device__ unsigned int SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::getBlockEdgeSize ( )

inlinestaticconstexpr

Return the size of the block edge size.

Returns: the block edge size

Definition at line 261 of file SparseGridGpu_ker.cuh.

◆ getBlockId()

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<typename CoordT >

__device__ unsigned int SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::getBlockId ( const grid_key_dx< dim, CoordT > & coord )

inline

Definition at line 453 of file SparseGridGpu_ker.cuh.

◆ getBlockLinId()

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<typename CoordT >

__device__ size_t SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::getBlockLinId ( CoordT blockCoord ) const

inline

Linearization of block coordinates.

Parameters

blockCoord block coordinates

Returns: the linearized index

Definition at line 218 of file SparseGridGpu_ker.cuh.

◆ getBlockSize()

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

constexpr __device__ unsigned int SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::getBlockSize ( ) const

inlineconstexpr

Return the size of the block.

Returns: the block size

Definition at line 272 of file SparseGridGpu_ker.cuh.

◆ getCoord() [1/2]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

__device__ grid_key_dx< dim, int > SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::getCoord	(	size_t	dataBlockId,
		unsigned	offset
	)		const

inline

The inversion of getLinId.

Parameters

linearized index

Returns: the point coordinate

Definition at line 176 of file SparseGridGpu_ker.cuh.

◆ getCoord() [2/2]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

__device__ grid_key_dx< dim, int > SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::getCoord ( size_t linId ) const

inline

The inversion of getLinId.

Parameters

linearized index

Returns: the point coordinate

Definition at line 163 of file SparseGridGpu_ker.cuh.

◆ getCoordInEnlargedBlock()

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

__device__ grid_key_dx< dim, int > SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::getCoordInEnlargedBlock ( const unsigned int offset ) const

inline

Definition at line 341 of file SparseGridGpu_ker.cuh.

◆ getEnlargedBlockSize()

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

__device__ unsigned int SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::getEnlargedBlockSize ( ) const

inline

Return the size of the block + ghost needed to apply the stencil.

Returns: the block size + ghost

Definition at line 282 of file SparseGridGpu_ker.cuh.

◆ getGlobalCoord()

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<typename CoordT >

__device__ __host__ grid_key_dx< dim, CoordT > SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::getGlobalCoord	(	const grid_key_dx< dim, CoordT > &	blockCoord,
		unsigned int	offset
	)

inline

Get the coordinate of the block and the offset id inside the block it give the global coordinate.

Parameters

blockCoord	block coordinate
offset	in the block

Returns: the global coordinate position

Definition at line 126 of file SparseGridGpu_ker.cuh.

◆ getIfBoundaryElementInEnlargedBlock()

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

__device__ bool SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::getIfBoundaryElementInEnlargedBlock	(	const grid_key_dx< dim, int >	coordInEnlargedBlock,
		char(&)	boundaryDirection[dim]
	)

inline

Definition at line 375 of file SparseGridGpu_ker.cuh.

◆ getInsertBlockOffset()

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<typename ite_type >

__device__ bool SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::getInsertBlockOffset	(	const ite_type &	itd,
		const grid_key_dx< dim, int > &	p,
		grid_key_dx< dim, int > &	blk,
		int &	offset
	)

inline

Given a point to insert, return the block-id and offset of that point.

Parameters

p	Point the thread is processing
blk	Block id the point is falling into
offset	linearized index within the block

Returns: true if the point is inactive

Definition at line 192 of file SparseGridGpu_ker.cuh.

◆ getLinId()

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<typename CoordT >

__device__ size_t SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::getLinId ( const grid_key_dx< dim, CoordT > & coord ) const

inline

Linearization of global coordinates.

Parameters

coord point coordinates

Returns: the linearized position

Definition at line 139 of file SparseGridGpu_ker.cuh.

◆ getLinIdInEnlargedBlock()

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

__device__ unsigned int SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::getLinIdInEnlargedBlock ( const unsigned int offset ) const

inline

Definition at line 349 of file SparseGridGpu_ker.cuh.

◆ getNeighbour()

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

__device__ grid_key_dx< dim, int > SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::getNeighbour	(	grid_key_dx< dim, int >	base,
		unsigned int	dimension,
		char	offset
	)		const

inline

Definition at line 248 of file SparseGridGpu_ker.cuh.

◆ getNeighbourLinIdInEnlargedBlock() [1/2]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<typename Coordtype >

__device__ unsigned int SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::getNeighbourLinIdInEnlargedBlock	(	const grid_key_dx< dim, Coordtype > &	base,
		grid_key_dx< dim, Coordtype > &	offsets
	)		const

inline

Definition at line 360 of file SparseGridGpu_ker.cuh.

◆ getNeighbourLinIdInEnlargedBlock() [2/2]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<typename Coordtype >

__device__ unsigned int SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::getNeighbourLinIdInEnlargedBlock	(	const grid_key_dx< dim, Coordtype > &	base,
		unsigned int	dimension,
		char	offset
	)		const

inline

Definition at line 368 of file SparseGridGpu_ker.cuh.

◆ getNeighboursPos()

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<typename NNtype >

__device__ indexT SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::getNeighboursPos	(	const indexT	blockId,
		const unsigned int	offset
	)

inline

Definition at line 696 of file SparseGridGpu_ker.cuh.

◆ getNNPoint()

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<typename NN_type , typename indexT2 >

__device__ block_offset< indexT2 > SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::getNNPoint	(	openfpm::sparse_index< unsigned int >	pos,
		unsigned int	offset,
		const grid_key_dx< dim, indexT2 > &	mov
	)

inline

Get the neighborhood point in one direction.

Warning: This function assume you do not move more than one block edge size in one direction

Parameters

pos	data block position
offset	inside the data block
position	where to move

Definition at line 297 of file SparseGridGpu_ker.cuh.

◆ getOffset()

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<typename CoordT >

__device__ unsigned int SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::getOffset ( const grid_key_dx< dim, CoordT > & coord )

inline

Definition at line 460 of file SparseGridGpu_ker.cuh.

◆ getPointBuffer()

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

__device__ auto SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::getPointBuffer ( ) -> decltype(buffPnt) &

inline

Return the buffer of points.

Returns: the buffer of points

Definition at line 497 of file SparseGridGpu_ker.cuh.

◆ insert()

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<unsigned int p, typename CoordT >

__device__ auto SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::insert ( const grid_key_dx< dim, CoordT > & coord ) -> ScalarTypeOf<AggregateBlockT, p>&

inline

Definition at line 447 of file SparseGridGpu_ker.cuh.

◆ insertBlock() [1/2]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<unsigned int chunksPerBlocks = 1, typename CoordT >

__device__ auto SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::insertBlock ( const grid_key_dx< dim, CoordT > & coord ) -> decltype(BlockMapGpu_ker<AggregateBlockT, indexT, layout_base>::insertBlock(0))

inline

Definition at line 480 of file SparseGridGpu_ker.cuh.

◆ insertBlock() [2/2]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<unsigned int chunksPerBlocks = 1>

__device__ auto SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::insertBlock	(	const indexT	blockLinId,
		const unsigned int	stride = `8192`
	)		-> decltype(BlockMapGpu_ker<AggregateBlockT, indexT, layout_base>::insertBlock(0))

inline

Definition at line 487 of file SparseGridGpu_ker.cuh.

◆ isPadding() [1/2]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<typename BitMaskT >

static __device__ bool SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::isPadding ( const BitMaskT & bitMask )

inlinestatic

Definition at line 671 of file SparseGridGpu_ker.cuh.

◆ isPadding() [2/2]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<typename keyIndexT >

__device__ bool SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::isPadding ( grid_key_dx< dim, keyIndexT > coord ) const

inline

Definition at line 677 of file SparseGridGpu_ker.cuh.

◆ loadBlock() [1/2]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<unsigned int p, typename AggrWrapperT >

__device__ void SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::loadBlock	(	AggrWrapperT &	block,
		ScalarTypeOf< AggregateBlockT, p > *	sharedRegion
	)

inline

Load a data block into the inner part of a shared memory region. The given shared memory region should be shaped as a dim-dimensional array and sized so that it can contain the block plus the ghost layer around it.

Template Parameters

p	The property to retrieve from global memory.
CoordT	The coordinate type.
Shared	The type of the shared memory region.

Parameters

coord	The coordinate of the block.
sharedRegionPtr	The pointer to the shared memory region.

Definition at line 517 of file SparseGridGpu_ker.cuh.

◆ loadBlock() [2/2]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<unsigned int ... props, typename AggrWrapperT >

__device__ void SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::loadBlock	(	AggrWrapperT &	block,
		void *	sharedRegionPtr[sizeof...(props)]
	)

inline

Load a data block into the inner part of a shared memory region. The given shared memory region should be shaped as a dim-dimensional array and sized so that it can contain the block plus the ghost layer around it.

Template Parameters

CoordT	The coordinate type.
props	The set of properties to retrieve from global memory.

Parameters

coord	The coordinate of the block.
sharedRegionPtr	The array of pointers to the shared memory regions, one for each property.

Definition at line 534 of file SparseGridGpu_ker.cuh.

◆ loadGhost() [1/2]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<unsigned int ... props, typename CoordT >

__device__ void SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::loadGhost	(	const grid_key_dx< dim, CoordT > &	coord,
		const int *	neighboursPos,
		void *	sharedRegionPtr[sizeof...(props)]
	)

inline

Load the ghost layer of a data block into the boundary part of a shared memory region. The given shared memory region should be shaped as a dim-dimensional array and sized so that it can contain the block plus the ghost layer around it.

Template Parameters

CoordT	The coordinate type.
Shared	The type of the shared memory region.
props	The set of properties to retrieve from global memory.

Parameters

coord	The coordinate of the block.
sharedRegionPtr	The pointer to the shared memory region.

Definition at line 585 of file SparseGridGpu_ker.cuh.

◆ loadGhost() [2/2]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<unsigned int ... props>

__device__ void SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::loadGhost	(	const unsigned int	blockLinId,
		const int *	neighboursPos,
		void *	sharedRegionPtr[sizeof...(props)]
	)

inline

Definition at line 593 of file SparseGridGpu_ker.cuh.

◆ loadGhostBlock() [1/3]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<unsigned int p, typename AggrWrapperT , typename CoordT >

__device__ void SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::loadGhostBlock	(	const AggrWrapperT &	dataBlockLoad,
		const grid_key_dx< dim, CoordT > &	coord,
		ScalarTypeOf< AggregateBlockT, p > *	sharedRegion
	)

inline

Load the ghost layer of a data block into the boundary part of a shared memory region. The given shared memory region should be shaped as a dim-dimensional array and sized so that it can contain the block plus the ghost layer around it.

Template Parameters

p	The property to retrieve from global memory.
CoordT	The coordinate type.
Shared	The type of the shared memory region.

Parameters

coord	The coordinate of the block.
sharedRegionPtr	The pointer to the shared memory region.

Definition at line 553 of file SparseGridGpu_ker.cuh.

◆ loadGhostBlock() [2/3]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<unsigned int p, typename AggrWrapperT >

__device__ void SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::loadGhostBlock	(	const AggrWrapperT &	dataBlockLoad,
		const openfpm::sparse_index< unsigned int >	blockLinId,
		ScalarTypeOf< AggregateBlockT, p > *	sharedRegion
	)

inline

Definition at line 561 of file SparseGridGpu_ker.cuh.

◆ loadGhostBlock() [3/3]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<unsigned int p, typename AggrWrapperT >

__device__ void SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::loadGhostBlock	(	const AggrWrapperT &	dataBlockLoad,
		const openfpm::sparse_index< unsigned int >	blockLinId,
		ScalarTypeOf< AggregateBlockT, p > *	sharedRegion,
		unsigned char *	mask
	)

inline

Definition at line 568 of file SparseGridGpu_ker.cuh.

◆ posToEnlargedBlockPos()

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

__device__ unsigned int SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::posToEnlargedBlockPos ( unsigned int pos ) const

inline

Definition at line 329 of file SparseGridGpu_ker.cuh.

◆ remove()

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<typename CoordT >

__device__ void SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::remove ( const grid_key_dx< dim, CoordT > & coord )

inline

Definition at line 665 of file SparseGridGpu_ker.cuh.

◆ setPadding()

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<typename BitMaskT >

static __device__ void SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::setPadding ( BitMaskT & bitMask )

inlinestatic

Definition at line 684 of file SparseGridGpu_ker.cuh.

◆ size()

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

__device__ unsigned int SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::size ( unsigned int i )

inline

Size of the sparse grid in each direction.

Parameters

i direction

Returns: the linearized position

Definition at line 151 of file SparseGridGpu_ker.cuh.

◆ storeBlock() [1/2]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<unsigned int p, typename AggrWrapperT >

__device__ void SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::storeBlock	(	AggrWrapperT &	block,
		ScalarTypeOf< AggregateBlockT, p > *	sharedRegion
	)

inline

Definition at line 600 of file SparseGridGpu_ker.cuh.

◆ storeBlock() [2/2]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<unsigned int ... props, typename AggrWrapperT >

__device__ void SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::storeBlock	(	AggrWrapperT &	block,
		void *	sharedRegionPtr[sizeof...(props)]
	)

inline

Read a data block from the inner part of a shared memory region and store it in global memory. The given shared memory region should be shaped as a dim-dimensional array and sized so that it can contain the block plus the ghost layer around it.

Template Parameters

CoordT	The coordinate type.
Shared	The type of the shared memory region.
props	The set of properties to retrieve from global memory.

Parameters

coord	The coordinate of the block.
sharedRegionPtr	The array of pointers to the shared memory regions, one for each property.

Definition at line 627 of file SparseGridGpu_ker.cuh.

◆ storeBlockInPlace() [1/2]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<unsigned int p, typename CoordT >

__device__ void SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::storeBlockInPlace	(	const grid_key_dx< dim, CoordT > &	coord,
		ScalarTypeOf< AggregateBlockT, p > *	sharedRegion
	)

inline

Definition at line 608 of file SparseGridGpu_ker.cuh.

◆ storeBlockInPlace() [2/2]

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<unsigned int ... props, typename CoordT >

__device__ void SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::storeBlockInPlace	(	const grid_key_dx< dim, CoordT > &	coord,
		void *	sharedRegionPtr[sizeof...(props)]
	)

inline

Definition at line 633 of file SparseGridGpu_ker.cuh.

◆ unpack_headers()

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<typename headers_type >

static __device__ int SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::unpack_headers	(	headers_type &	headers,
		unsigned char *	data,
		int	ih,
		int	sz_pack
	)

inlinestatic

Parameters

Definition at line 90 of file SparseGridGpu_ker.cuh.

◆ unsetPadding()

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

template<typename BitMaskT >

static __device__ void SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::unsetPadding ( BitMaskT & bitMask )

inlinestatic

Definition at line 690 of file SparseGridGpu_ker.cuh.

Field Documentation

◆ background

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

BcT SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::background

private

background values

Definition at line 36 of file SparseGridGpu_ker.cuh.

◆ blockEdgeSize_

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

constexpr unsigned int SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::blockEdgeSize_ = blockEdgeSize

staticconstexpr

Definition at line 49 of file SparseGridGpu_ker.cuh.

◆ blockSize

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

constexpr unsigned int SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::blockSize = BlockTypeOf<AggregateBlockT, 0>::size

staticconstexprprotected

Definition at line 40 of file SparseGridGpu_ker.cuh.

◆ blockWithGhostGrid

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

GridSmT SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::blockWithGhostGrid

private

Definition at line 33 of file SparseGridGpu_ker.cuh.

◆ buffPnt

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

openfpm::vector_gpu_ker<aggregate<indexT>,memory_traits_inte> SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::buffPnt

protected

Definition at line 44 of file SparseGridGpu_ker.cuh.

◆ d

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

constexpr unsigned int SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::d = dim

staticconstexpr

Definition at line 47 of file SparseGridGpu_ker.cuh.

◆ dims

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

constexpr unsigned int SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::dims = dim

staticconstexpr

Definition at line 48 of file SparseGridGpu_ker.cuh.

◆ ghostLayerSize

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

unsigned int SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::ghostLayerSize

protected

Definition at line 41 of file SparseGridGpu_ker.cuh.

◆ ghostLayerToThreadsMapping

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

openfpm::vector_gpu_ker<aggregate<short int,short int>,memory_traits_inte> SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::ghostLayerToThreadsMapping

protected

Definition at line 42 of file SparseGridGpu_ker.cuh.

◆ grid

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

linearizer SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::grid

private

Definition at line 32 of file SparseGridGpu_ker.cuh.

◆ nn_blocks

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

openfpm::vector_gpu_ker<aggregate<indexT>,memory_traits_inte> SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::nn_blocks

protected

Definition at line 43 of file SparseGridGpu_ker.cuh.

◆ PADDING_BIT

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

const unsigned char SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::PADDING_BIT = 1

staticprotected

Definition at line 39 of file SparseGridGpu_ker.cuh.

◆ stencilSupportRadius

template<unsigned int dim, unsigned int blockEdgeSize, typename AggregateBlockT , typename ct_params , typename indexT , template< typename > class layout_base, typename GridSmT , typename linearizer , typename BcT >

unsigned int SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT >::stencilSupportRadius

Definition at line 50 of file SparseGridGpu_ker.cuh.

The documentation for this class was generated from the following file:

openfpm_data/src/SparseGridGpu/SparseGridGpu_ker.cuh

Detailed Description

Public Types

Public Member Functions

Static Public Member Functions

Data Fields

Static Public Attributes

Protected Attributes

Static Protected Attributes

Private Member Functions

Private Attributes

Member Typedef Documentation

◆ AggregateBlockType

◆ indexT_

◆ yes_has_check_device_pointer

Constructor & Destructor Documentation

◆ SparseGridGpu_ker()

Member Function Documentation

◆ __loadBlock() [1/2]

◆ __loadBlock() [2/2]

◆ __loadGhost()

◆ __loadGhostBlock() [1/2]

◆ __loadGhostBlock() [2/2]

◆ __loadGhostNoNN()

◆ __storeBlock() [1/2]

◆ __storeBlock() [2/2]

◆ get() [1/4]

◆ get() [2/4]

◆ get() [3/4]

◆ get() [4/4]

◆ get_sparse()

◆ getBlock() [1/2]

◆ getBlock() [2/2]

◆ getBlockBaseCoord()

◆ getBlockCoord()

◆ getBlockEdgeSize()

◆ getBlockId()

◆ getBlockLinId()

◆ getBlockSize()

◆ getCoord() [1/2]

◆ getCoord() [2/2]

◆ getCoordInEnlargedBlock()

◆ getEnlargedBlockSize()

◆ getGlobalCoord()

◆ getIfBoundaryElementInEnlargedBlock()

◆ getInsertBlockOffset()

◆ getLinId()

◆ getLinIdInEnlargedBlock()

◆ getNeighbour()

◆ getNeighbourLinIdInEnlargedBlock() [1/2]

◆ getNeighbourLinIdInEnlargedBlock() [2/2]

◆ getNeighboursPos()

◆ getNNPoint()

◆ getOffset()

◆ getPointBuffer()

◆ insert()

◆ insertBlock() [1/2]

◆ insertBlock() [2/2]

◆ isPadding() [1/2]

◆ isPadding() [2/2]

◆ loadBlock() [1/2]

◆ loadBlock() [2/2]

◆ loadGhost() [1/2]

◆ loadGhost() [2/2]

◆ loadGhostBlock() [1/3]

◆ loadGhostBlock() [2/3]

◆ loadGhostBlock() [3/3]

◆ posToEnlargedBlockPos()

◆ remove()

◆ setPadding()

◆ size()

◆ storeBlock() [1/2]

◆ storeBlock() [2/2]

◆ storeBlockInPlace() [1/2]

◆ storeBlockInPlace() [2/2]

◆ unpack_headers()

◆ unsetPadding()

Field Documentation

◆ background

◆ blockEdgeSize_

◆ blockSize