OpenFPM_pdata  4.1.0
Project that contain the implementation of distributed structures
SparseGridGpu_testKernels.cuh
1 //
2 // Created by tommaso on 15/8/19.
3 //
4 
5 #ifndef OPENFPM_PDATA_SPARSEGRIDGPU_TESTKERNELS_CUH
6 #define OPENFPM_PDATA_SPARSEGRIDGPU_TESTKERNELS_CUH
7 
9 
10 template<unsigned int dim, unsigned int p_src, unsigned int p_dst>
11 struct BoundaryStencilSetX
12 {
13  // This is an example of a boundary stencil setting the value to the same value as the x coordinate
14 
15  typedef NNStar<dim> stencil_type;
16 
17  static constexpr unsigned int supportRadius = 1;
18 
19  template<typename SparseGridT, typename DataBlockWrapperT>
20  static inline __device__ void stencil(
21  SparseGridT & sparseGrid,
22  const unsigned int dataBlockId,
23  openfpm::sparse_index<unsigned int> dataBlockIdPos,
24  unsigned int offset,
25  grid_key_dx<dim, int> & pointCoord,
26  DataBlockWrapperT & dataBlockLoad,
27  DataBlockWrapperT & dataBlockStore,
28  unsigned char curMask)
29  {
30  if (curMask == mask_sparse::EXIST_AND_PADDING)
31  {
32  dataBlockStore.template get<p_dst>()[offset] = pointCoord.get(0);
33  }
34  }
35 };
36 
37 template<unsigned int dim, unsigned int p_src, unsigned int p_dst, typename ScalarT = float>
38 struct BoundaryStencilSetXRescaled
39 {
40  // This is an example of a boundary stencil setting the value to the same value as the x coordinate
41 
42  typedef NNStar<dim> stencil_type;
43 
44  static constexpr unsigned int supportRadius = 1;
45 
46  template<typename SparseGridT, typename DataBlockWrapperT>
47  static inline __device__ void stencil(
48  SparseGridT & sparseGrid,
49  const unsigned int dataBlockId,
50  openfpm::sparse_index<unsigned int> dataBlockIdPos,
51  unsigned int offset,
52  grid_key_dx<dim, int> & pointCoord,
53  DataBlockWrapperT & dataBlockLoad,
54  DataBlockWrapperT & dataBlockStore,
55  unsigned char curMask,
56  ScalarT minX, ScalarT maxX, ScalarT minValue, ScalarT maxValue)
57  {
58  if (curMask == mask_sparse::EXIST_AND_PADDING)
59  {
60  const ScalarT x = pointCoord.get(0);
61 
62  auto value = maxValue * (x - minX) / (maxX - minX - 1);
63  if (x < minX)
64  {
65  value = minValue;
66  }
67  else if (x > maxX)
68  {
69  value = maxValue;
70  }
71  dataBlockStore.template get<p_dst>()[offset] = value;
72  }
73  }
74 
75  template <typename SparseGridT, typename CtxT>
76  static inline void __host__ flush(SparseGridT & sparseGrid, CtxT & ctx)
77  {
78  // No flush
79  }
80 };
81 
83 
84 template<unsigned int p, typename SparseGridType, typename ValueType>
85 __global__ void insertSphere(SparseGridType sparseGrid, grid_key_dx<2,int> start, float r1, float r2, ValueType value)
86 {
87  constexpr unsigned int pMask = SparseGridType::pMask;
88  typedef BlockTypeOf<typename SparseGridType::AggregateType, p> BlockT;
89  typedef BlockTypeOf<typename SparseGridType::AggregateType, pMask> MaskBlockT;
90 
91  typedef typename SparseGridType::indexT_ idType;
92 
93  grid_key_dx<2,idType> blk({
94  blockIdx.x + start.get(0) / sparseGrid.getBlockEdgeSize(),
95  blockIdx.y + start.get(1) / sparseGrid.getBlockEdgeSize()
96  });
97 
98  unsigned int offset = threadIdx.x;
99 
100  __shared__ bool is_block_empty;
101 
102  if (threadIdx.x == 0 && threadIdx.y == 0)
103  {is_block_empty = true;}
104 
105  sparseGrid.init();
106 
107  auto blockId = sparseGrid.getBlockLinId(blk);
108 
109  grid_key_dx<2,idType> keyg;
110  keyg = sparseGrid.getGlobalCoord(blk,offset);
111 
112  float radius = sqrt( (float)
113  (keyg.get(0) - (start.get(0) + gridDim.x/2*SparseGridType::blockEdgeSize_))
114  * (keyg.get(0) - (start.get(0) + gridDim.x/2*SparseGridType::blockEdgeSize_))
115  + (keyg.get(1) - (start.get(1) + gridDim.y/2*SparseGridType::blockEdgeSize_))
116  * (keyg.get(1) - (start.get(1) + gridDim.y/2*SparseGridType::blockEdgeSize_)) );
117 
118  bool is_active = radius < r1 && radius > r2;
119 
120  if (is_active == true)
121  {
122  is_block_empty = false;
123  }
124 
125  __syncthreads();
126 
127  if (is_block_empty == false)
128  {
129  auto ec = sparseGrid.insertBlock(blockId);
130 
131  if ( is_active == true)
132  {
133  ec.template get<p>()[offset] = value;
134  BlockMapGpu_ker<>::setExist(ec.template get<pMask>()[offset]);
135  }
136  }
137 
138  __syncthreads();
139 
140  sparseGrid.flush_block_insert();
141 }
142 
143 template<unsigned int p, typename SparseGridType, typename ValueType>
144 __global__ void insertSphere3D(SparseGridType sparseGrid, grid_key_dx<3,int> start, float r1, float r2, ValueType value)
145 {
146  constexpr unsigned int pMask = SparseGridType::pMask;
147  typedef BlockTypeOf<typename SparseGridType::AggregateType, p> BlockT;
148  typedef BlockTypeOf<typename SparseGridType::AggregateType, pMask> MaskBlockT;
149 
150  typedef typename SparseGridType::indexT_ idType;
151 
152  grid_key_dx<3,idType> blk({
153  blockIdx.x + start.get(0) / sparseGrid.getBlockEdgeSize(),
154  blockIdx.y + start.get(1) / sparseGrid.getBlockEdgeSize(),
155  blockIdx.z + start.get(2) / sparseGrid.getBlockEdgeSize()});
156 
157  unsigned int offset = threadIdx.x;
158 
159  __shared__ bool is_block_empty;
160 
161  if (threadIdx.x == 0 && threadIdx.y == 0 && threadIdx.z == 0)
162  {is_block_empty = true;}
163 
164  sparseGrid.init();
165 
166  auto blockId = sparseGrid.getBlockLinId(blk);
167 
168  grid_key_dx<3,idType> keyg;
169  keyg = sparseGrid.getGlobalCoord(blk,offset);
170 
171  const long int x = (long int)keyg.get(0) - (start.get(0) + gridDim.x / 2 * SparseGridType::blockEdgeSize_);
172  const long int y = (long int)keyg.get(1) - (start.get(1) + gridDim.y / 2 * SparseGridType::blockEdgeSize_);
173  const long int z = (long int)keyg.get(2) - (start.get(2) + gridDim.z / 2 * SparseGridType::blockEdgeSize_);
174 
175  float radius = sqrt((float) (x*x + y*y + z*z));
176 
177  bool is_active = radius < r1 && radius >= r2;
178 
179  if (is_active == true)
180  {
181  is_block_empty = false;
182  }
183 
184  __syncthreads();
185 
186  if (is_block_empty == false)
187  {
188  auto ec = sparseGrid.insertBlock(blockId);
189 
190  if ( is_active == true)
191  {
192  ec.template get<p>()[offset] = value;
193  BlockMapGpu_ker<>::setExist(ec.template get<pMask>()[offset]);
194  }
195  }
196 
197  __syncthreads();
198 
199  sparseGrid.flush_block_insert();
200 }
201 
202 template<unsigned int p, typename SparseGridType, typename ValueType>
203 __global__ void insertSphere3D_radius(SparseGridType sparseGrid, grid_key_dx<3,int> start, float r1, float r2, ValueType value)
204 {
205  constexpr unsigned int pMask = SparseGridType::pMask;
206  typedef BlockTypeOf<typename SparseGridType::AggregateType, p> BlockT;
207  typedef BlockTypeOf<typename SparseGridType::AggregateType, pMask> MaskBlockT;
208 
209  typedef typename SparseGridType::indexT_ idType;
210 
211  grid_key_dx<3,idType> blk({
212  blockIdx.x + start.get(0) / sparseGrid.getBlockEdgeSize(),
213  blockIdx.y + start.get(1) / sparseGrid.getBlockEdgeSize(),
214  blockIdx.z + start.get(2) / sparseGrid.getBlockEdgeSize()});
215 
216  unsigned int offset = threadIdx.x;
217 
218  __shared__ bool is_block_empty;
219 
220  if (threadIdx.x == 0 && threadIdx.y == 0 && threadIdx.z == 0)
221  {is_block_empty = true;}
222 
223  sparseGrid.init();
224 
225  auto blockId = sparseGrid.getBlockLinId(blk);
226 
227  grid_key_dx<3,idType> keyg;
228  keyg = sparseGrid.getGlobalCoord(blk,offset);
229 
230  const long int x = (long int)keyg.get(0) - (start.get(0) + gridDim.x / 2 * SparseGridType::blockEdgeSize_);
231  const long int y = (long int)keyg.get(1) - (start.get(1) + gridDim.y / 2 * SparseGridType::blockEdgeSize_);
232  const long int z = (long int)keyg.get(2) - (start.get(2) + gridDim.z / 2 * SparseGridType::blockEdgeSize_);
233 
234  float radius = sqrt((float) (x*x + y*y + z*z));
235 
236  bool is_active = radius < r1 && radius > r2;
237 
238  if (is_active == true)
239  {
240  is_block_empty = false;
241  }
242 
243  __syncthreads();
244 
245  if (is_block_empty == false)
246  {
247  auto ec = sparseGrid.insertBlock(blockId);
248 
249  if ( is_active == true)
250  {
251  ec.template get<p>()[offset] = x+y+z;
252  BlockMapGpu_ker<>::setExist(ec.template get<pMask>()[offset]);
253  }
254  }
255 
256  __syncthreads();
257 
258  sparseGrid.flush_block_insert();
259 }
260 
261 template<unsigned int p, typename SparseGridType, typename ValueType>
262 __global__ void insertSphere3D_radiusV(SparseGridType sparseGrid, grid_key_dx<3,int> start, float r1, float r2, ValueType value)
263 {
264  constexpr unsigned int pMask = SparseGridType::pMask;
265  typedef BlockTypeOf<typename SparseGridType::AggregateType, p> BlockT;
266  typedef BlockTypeOf<typename SparseGridType::AggregateType, pMask> MaskBlockT;
267 
268  typedef typename SparseGridType::indexT_ idType;
269 
270  grid_key_dx<3,idType> blk({
271  blockIdx.x + start.get(0) / sparseGrid.getBlockEdgeSize(),
272  blockIdx.y + start.get(1) / sparseGrid.getBlockEdgeSize(),
273  blockIdx.z + start.get(2) / sparseGrid.getBlockEdgeSize()});
274 
275  unsigned int offset = threadIdx.x;
276 
277  __shared__ bool is_block_empty;
278 
279  if (threadIdx.x == 0 && threadIdx.y == 0 && threadIdx.z == 0)
280  {is_block_empty = true;}
281 
282  sparseGrid.init();
283 
284  auto blockId = sparseGrid.getBlockLinId(blk);
285 
286  grid_key_dx<3,idType> keyg;
287  keyg = sparseGrid.getGlobalCoord(blk,offset);
288 
289  const long int x = (long int)keyg.get(0) - (start.get(0) + gridDim.x / 2 * SparseGridType::blockEdgeSize_);
290  const long int y = (long int)keyg.get(1) - (start.get(1) + gridDim.y / 2 * SparseGridType::blockEdgeSize_);
291  const long int z = (long int)keyg.get(2) - (start.get(2) + gridDim.z / 2 * SparseGridType::blockEdgeSize_);
292 
293  float radius = sqrt((float) (x*x + y*y + z*z));
294 
295  bool is_active = radius < r1 && radius > r2;
296 
297  if (is_active == true)
298  {
299  is_block_empty = false;
300  }
301 
302  __syncthreads();
303 
304  if (is_block_empty == false)
305  {
306  auto ec = sparseGrid.insertBlock(blockId);
307 
308  if ( is_active == true)
309  {
310  ec.template get<p>()[offset] = x+y+z;
311  ec.template get<p+1>()[0][offset] = x;
312  ec.template get<p+1>()[1][offset] = y;
313  ec.template get<p+1>()[2][offset] = z;
314  BlockMapGpu_ker<>::setExist(ec.template get<pMask>()[offset]);
315  }
316  }
317 
318  __syncthreads();
319 
320  sparseGrid.flush_block_insert();
321 }
322 
323 template<unsigned int p, typename SparseGridType, typename ValueType>
324 __global__ void removeSphere3D_even_radiusV(SparseGridType sparseGrid, grid_key_dx<3,int> start, float r1, float r2, ValueType value)
325 {
326  constexpr unsigned int pMask = SparseGridType::pMask;
327  typedef BlockTypeOf<typename SparseGridType::AggregateType, p> BlockT;
328  typedef BlockTypeOf<typename SparseGridType::AggregateType, pMask> MaskBlockT;
329 
330  typedef typename SparseGridType::indexT_ idType;
331 
332  grid_key_dx<3,idType> blk({
333  blockIdx.x + start.get(0) / sparseGrid.getBlockEdgeSize(),
334  blockIdx.y + start.get(1) / sparseGrid.getBlockEdgeSize(),
335  blockIdx.z + start.get(2) / sparseGrid.getBlockEdgeSize()});
336 
337  unsigned int offset = threadIdx.x;
338 
339  auto blockId = sparseGrid.getBlockLinId(blk);
340 
341  grid_key_dx<3,idType> keyg;
342  keyg = sparseGrid.getGlobalCoord(blk,offset);
343 
344  const long int x = (long int)keyg.get(0) - (start.get(0) + gridDim.x / 2 * SparseGridType::blockEdgeSize_);
345  const long int y = (long int)keyg.get(1) - (start.get(1) + gridDim.y / 2 * SparseGridType::blockEdgeSize_);
346  const long int z = (long int)keyg.get(2) - (start.get(2) + gridDim.z / 2 * SparseGridType::blockEdgeSize_);
347 
348  float radius = sqrt((float) (x*x + y*y + z*z));
349 
350  bool is_active = radius < r1 && radius > r2 && (keyg.get(0) + keyg.get(1) + keyg.get(2)) % 2 == 0;
351 
352  if (is_active == true)
353  {
354  sparseGrid.remove(keyg);
355  }
356 }
357 
358 #endif //OPENFPM_PDATA_SPARSEGRIDGPU_TESTKERNELS_CUH
grid_key_dx< dim, int >
grid_key_dx::get
__device__ __host__ index_type get(index_type i) const
Get the i index.
Definition: grid_key.hpp:503
cub::int
KeyT const ValueT ValueT OffsetIteratorT OffsetIteratorT int
[in] The number of segments that comprise the sorting data
Definition: dispatch_radix_sort.cuh:331
openfpm::sparse_index
Definition: map_vector_sparse_cuda_ker.cuh:33
BlockMapGpu_ker
Definition: BlockMapGpu_ker.cuh:71