OpenFPM_pdata  3.0.0
Project that contain the implementation of distributed structures
SparseGridGpu_ker< dim, blockEdgeSize, AggregateBlockT, ct_params, indexT, layout_base, GridSmT, linearizer, BcT > Class Template Reference

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

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_integhostLayerToThreadsMapping
 
openfpm::vector_gpu_ker< aggregate< indexT >, memory_traits_intenn_blocks
 
openfpm::vector_gpu_ker< aggregate< indexT >, memory_traits_intebuffPnt
 
- 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
blockIdIndex of the center block
neighboursPosneighborhood blocks around blockId (if neighbourPos is null loadGhost)
sharedRegionPtrshared region

Definition at line 780 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
blockIdIndex of the center block
neighboursPosneighborhood blocks around blockId (if neighbourPos is null loadGhost)
sharedRegionPtrshared region

Definition at line 804 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
coordpoint
Returns
a reference to the data point

Definition at line 389 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
coordpoint
Returns
a reference to the data point

Definition at line 403 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
blockLinIdlinearized block index
Returns
the block coordinates

Definition at line 210 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
blockLinIdlinearized block index
Returns
the block coordinates

Definition at line 197 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 228 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
blockCoordblock coordinates
Returns
the linearized index

Definition at line 185 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 239 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
linearizedindex
Returns
the point coordinate

Definition at line 130 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
linearizedindex
Returns
the point coordinate

Definition at line 143 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 249 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
blockCoordblock coordinate
offsetin the block
Returns
the global coordinate position

Definition at line 93 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
pPoint the thread is processing
blkBlock id the point is falling into
offsetlinearized index within the block
Returns
true if the point is inactive

Definition at line 159 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
coordpoint coordinates
Returns
the linearized position

Definition at line 106 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
posdata block position
offsetinside the data block
positionwhere to move

Definition at line 264 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 460 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
pThe property to retrieve from global memory.
CoordTThe coordinate type.
SharedThe type of the shared memory region.
Parameters
coordThe coordinate of the block.
sharedRegionPtrThe pointer to the shared memory region.

Definition at line 480 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
CoordTThe coordinate type.
propsThe set of properties to retrieve from global memory.
Parameters
coordThe coordinate of the block.
sharedRegionPtrThe array of pointers to the shared memory regions, one for each property.

Definition at line 497 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
CoordTThe coordinate type.
SharedThe type of the shared memory region.
propsThe set of properties to retrieve from global memory.
Parameters
coordThe coordinate of the block.
sharedRegionPtrThe pointer to the shared memory region.

Definition at line 548 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
pThe property to retrieve from global memory.
CoordTThe coordinate type.
SharedThe type of the shared memory region.
Parameters
coordThe coordinate of the block.
sharedRegionPtrThe pointer to the shared memory region.

Definition at line 516 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
idirection
Returns
the linearized position

Definition at line 118 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
CoordTThe coordinate type.
SharedThe type of the shared memory region.
propsThe set of properties to retrieve from global memory.
Parameters
coordThe coordinate of the block.
sharedRegionPtrThe array of pointers to the shared memory regions, one for each property.

Definition at line 590 of file SparseGridGpu_ker.cuh.


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