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 More...

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. More...

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

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

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

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

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. More...

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

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

__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. More...

__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. More...

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

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. More...

__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. More...

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. More...

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. More...

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)

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

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

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

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

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

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

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

template<unsigned int nChunksPerBlocks = 1>
__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. More...

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

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. More...

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 >
template<typename BitMaskT >
static __device__ __host__ bool	getBit (const BitMaskT &bitMask, unsigned char pos)

template<typename BitMaskT >
static __device__ __host__ void	setBit (BitMaskT &bitMask, unsigned char pos)

template<typename BitMaskT >
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)

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

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

template<typename BitMaskT >
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 = AggregateBlockT::max_prop_real - 1

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 = 0

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. More...

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. More...

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

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

◆ __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.

◆ get() [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__ 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() [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 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.

◆ 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 ( )

inlinestatic

Return the size of the block edge size.

Returns: the block edge size

Definition at line 261 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

inline

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 linId ) const

inline

The inversion of getLinId.

Parameters

linearized index

Returns: the point coordinate

Definition at line 163 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	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.

◆ 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.

◆ 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.

◆ 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.

◆ 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.

◆ 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()

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.

◆ loadGhostBlock()

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.

◆ 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()

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.

◆ 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.

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

Constructor & Destructor Documentation

◆ SparseGridGpu_ker()

Member Function Documentation

◆ __loadGhostBlock() [1/2]

◆ __loadGhostBlock() [2/2]

◆ get() [1/2]

◆ get() [2/2]

◆ getBlockBaseCoord()

◆ getBlockCoord()

◆ getBlockEdgeSize()

◆ getBlockLinId()

◆ getBlockSize()

◆ getCoord() [1/2]

◆ getCoord() [2/2]

◆ getEnlargedBlockSize()

◆ getGlobalCoord()

◆ getInsertBlockOffset()

◆ getLinId()

◆ getNNPoint()

◆ getPointBuffer()

◆ loadBlock() [1/2]

◆ loadBlock() [2/2]

◆ loadGhost()

◆ loadGhostBlock()

◆ size()

◆ storeBlock()

◆ unpack_headers()