Commit bad52672 authored by incardon's avatar incardon

Multi level test

parent 02ef67f8
......@@ -152,6 +152,23 @@ public:
#endif // __NVCC__
}
inline __device__ void get_sparse(unsigned int linId, unsigned int & dataBlockPos , unsigned int & offset)
{
#ifdef __NVCC__
typedef BlockTypeOf<AggregateBlockT, pMask> BlockT;
unsigned int blockId = linId / BlockT::size;
offset = linId % BlockT::size;
const auto sid = blockMap.get_sparse(blockId);
dataBlockPos = sid.id;
#else // __NVCC__
std::cout << __FILE__ << ":" << __LINE__ << " error: you are supposed to compile this file with nvcc, if you want to use it with gpu" << std::endl;
#endif // __NVCC__
}
inline static __device__ unsigned int getBlockId(unsigned int linId)
{
#ifdef __NVCC__
......@@ -252,6 +269,7 @@ inline __device__ auto BlockMapGpu_ker<AggregateBlockT, indexT, layout_base>
#endif // __NVCC__
}
template<typename AggregateBlockT, typename indexT, template<typename> class layout_base>
template<unsigned int p>
inline __device__ auto BlockMapGpu_ker<AggregateBlockT, indexT, layout_base>
......
......@@ -459,6 +459,12 @@ private:
//! the set of all sub-set to pack
mutable openfpm::vector_gpu<Box<dim,int>> pack_subs;
//! links of the padding points with real points of a coarse sparsegrid
openfpm::vector_gpu<aggregate<size_t>> links_up;
//! links of the padding points with real points of a finer sparsegrid
openfpm::vector_gpu<aggregate<size_t>> link_dw;
protected:
static constexpr unsigned int blockSize = BlockTypeOf<AggregateBlockT, 0>::size;
typedef AggregateBlockT AggregateInternalT;
......@@ -1076,6 +1082,53 @@ public:
return BlockMapGpu<AggregateInternalT, threadBlockSize, indexT, layout_base>::template insert<p>(gridGeometry.LinId(coord));
}
/*! \brief construct link between levels
*
* \praram grid_up grid level up
* \param grid_dw grid level down
*
*/
void construct_link(self & grid_up, self & grid_dw, mgpu::ofp_context_t &context)
{
// Here it is crucial to use "auto &" as the type, as we need to be sure to pass the reference to the actual buffers!
auto & indexBuffer = BlockMapGpu<AggregateInternalT, threadBlockSize, indexT, layout_base>::blockMap.getIndexBuffer();
auto & dataBuffer = BlockMapGpu<AggregateInternalT, threadBlockSize, indexT, layout_base>::blockMap.getDataBuffer();
ite_gpu<1> ite;
ite.wthr.x = indexBuffer.size();
ite.wthr.y = 1;
ite.wthr.z = 1;
ite.thr.x = getBlockSize();
ite.thr.y = 1;
ite.thr.z = 1;
openfpm::vector_gpu<aggregate<unsigned int>> output;
output.resize(indexBuffer.size() + 1);
CUDA_LAUNCH((SparseGridGpuKernels::link_construct<dim,
BlockMapGpu<AggregateInternalT, threadBlockSize, indexT, layout_base>::pMask,
blockSize>),ite,grid_up.toKernel(),this->toKernel(),output.toKernel());
openfpm::scan((unsigned int *)output.template getDeviceBuffer<0>(),output.size(),(unsigned int *)output.template getDeviceBuffer<0>(),context);
output.template deviceToHost<0>(output.size()-1,output.size()-1);
unsigned int np_lup = output.template get<0>(output.size()-1);
links_up.resize(np_lup);
CUDA_LAUNCH((SparseGridGpuKernels::link_construct_insert<dim,
BlockMapGpu<AggregateInternalT, threadBlockSize, indexT, layout_base>::pMask,
blockSize>),ite,grid_up.toKernel(),this->toKernel(),output.toKernel(),links_up.toKernel());
}
/*! \Brief Before inser any element you have to call this function to initialize the insert buffer
*
* \param nBlock number of blocks the insert buffer has
......
......@@ -349,6 +349,15 @@ public:
return BlockMapGpu_ker<AggregateBlockT, indexT, layout_base>::template get<p>(grid.LinId(coord));
}
// Data management methods
template<typename CoordT>
inline __device__ void
get_sparse(const grid_key_dx<dim, CoordT> & coord, unsigned int & dataBlockPos, unsigned int & offset) const
{
return BlockMapGpu_ker<AggregateBlockT, indexT, layout_base>::get_sparse(grid.LinId(coord),dataBlockPos,offset);
}
/*! \brief Access the grid point
*
* \param coord point
......
......@@ -91,6 +91,76 @@ namespace SparseGridGpuKernels
sparseGrid.storeBlock<pMask>(dataBlock, enlargedBlock);
}
/*! \brief construct the link between 2 sparse grid
*
*
*/
template<unsigned int dim, unsigned int pMask, unsigned int chunk_size , typename SparseGridType, typename outputType>
__global__ void link_construct(SparseGridType grid_up, SparseGridType grid_cu, outputType out)
{
const unsigned int dataBlockPos = blockIdx.x;
const unsigned int offset = threadIdx.x;
auto & indexBuffer = grid_cu.getIndexBuffer();
auto & dataBuffer = grid_cu.getDataBuffer();
// if the point is a padding
if (dataBuffer.template get <pMask>(dataBlockPos)[offset] & 0x2)
{
auto id = indexBuffer.template get<0>(dataBlockPos);
grid_key_dx<dim> pos = grid_cu.getCoord(id*chunk_size + offset);
for (int i = 0 ; i < dim ; i++)
{pos.set_d(i,pos.get(i) / 2);}
if (grid_up.template get<pMask>(pos)[offset] == 0x1)
{
atomic_add(&out.template get<0>(dataBlockPos),1);
}
}
}
/*! \brief construct the link between 2 sparse grid
*
*
*/
template<unsigned int dim, unsigned int pMask, unsigned int chunk_size, typename SparseGridType, typename scanType, typename outputType>
__global__ void link_construct_insert(SparseGridType grid_up, SparseGridType grid_cu, scanType scan, outputType out)
{
const unsigned int dataBlockPos = blockIdx.x;
const unsigned int offset = threadIdx.x;
auto & indexBuffer = grid_cu.getIndexBuffer();
auto & dataBuffer = grid_cu.getDataBuffer();
auto & dataBuffer_up = grid_up.getDataBuffer();
__shared__ int cnt;
cnt = 0;
__syncthreads();
// if the point is a padding
if (dataBuffer.template get <pMask>(dataBlockPos)[offset] & 0x2)
{
auto id = indexBuffer.template get<0>(dataBlockPos);
grid_key_dx<dim> pos = grid_cu.getCoord(id*chunk_size + offset);
for (int i = 0 ; i < dim ; i++)
{pos.set_d(i,pos.get(i) / 2);}
unsigned int dataBlockPos_up;
unsigned int offset_up;
grid_up.get_sparse(pos,dataBlockPos_up,offset_up);
if (dataBuffer_up.template get<pMask>(dataBlockPos_up) == 0x1)
{
int c = atomicAdd(&cnt,1);
out.template get<0>(scan.template get<0>(dataBlockPos) + c) = dataBlockPos_up * chunk_size + offset_up;
}
}
}
/*! \brief find the neighborhood of each chunk
*
* \param indexBuffer Chunk indec buffer
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment