OpenFPM_pdata  4.1.0
Project that contain the implementation of distributed structures
 
Loading...
Searching...
No Matches
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
10template<unsigned int dim, unsigned int p_src, unsigned int p_dst>
11struct 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
37template<unsigned int dim, unsigned int p_src, unsigned int p_dst, typename ScalarT = float>
38struct 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
84template<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
143template<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
202template<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
261template<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
323template<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
BlockMapGpu_ker
Definition BlockMapGpu_ker.cuh:72
grid_key_dx
grid_key_dx is the key to access any element in the grid
Definition grid_key.hpp:19
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:336
openfpm::sparse_index
Definition map_vector_sparse_cuda_ker.cuh:34