1/*
2 * grid_dist_id_iterator_sub.hpp
3 *
4 * Created on: Feb 4, 2015
5 * Author: Pietro Incardona
6 */
7
8#ifndef GRID_DIST_ID_ITERATOR_HPP_
9#define GRID_DIST_ID_ITERATOR_HPP_
10
11#define FREE 1
12#define FIXED 2
13#define ITERATION_ISOLATION 4
14
15#include "Grid/grid_dist_key.hpp"
16#include "VCluster/VCluster.hpp"
17#include "util/GBoxes.hpp"
18
19#ifdef __NVCC__
20#include "SparseGridGpu/encap_num.hpp"
21#endif
22
23template<unsigned int dim>
24struct launch_insert_sparse_lambda_call
25{
26 template<typename ec_type, typename lambda_t,typename coord_type>
27 __device__ inline static void call(ec_type & ec,lambda_t f, coord_type coord)
28 {
29 printf("Not implemented in this direction \n");
30 }
31
32 template<typename ite_type>
33 __device__ inline static bool set_keys(grid_key_dx<3,int> & key, grid_key_dx<3,int> & keyg, ite_type & itg)
34 {
35 return false;
36 }
37};
38
39template<>
40struct launch_insert_sparse_lambda_call<3>
41{
42 template<typename grid_type, typename lambda_t1, typename lambda_t2,typename itd_type, typename coord_type>
43 __device__ inline static void call(grid_type & grid,
44 lambda_t1 f1, lambda_t2 f2,
45 unsigned int blockId,
46 itd_type itd,
47 coord_type & key,
48 coord_type & keyg,unsigned int offset, bool & is_block_empty,
49 bool is_in)
50 {
51#ifdef __NVCC__
52
53 bool is_active = false;
54 if (is_in == true)
55 {is_active = f1(keyg.get(0),keyg.get(1),keyg.get(2));}
56
57 if (is_active == true)
58 {is_block_empty = false;}
59
60 __syncthreads();
61
62 if (is_block_empty == false)
63 {
64 auto ec = grid.insertBlock(blockId);
65 enc_num<decltype(grid.insertBlock(blockId))> ecn(ec,offset);
66
67 if ( is_active == true)
68 {
69 f2(ecn,keyg.get(0),keyg.get(1),keyg.get(2));
70 ec.template get<grid_type::pMask>()[offset] = 1;
71 }
72 }
73
74#endif
75 }
76
77 template<typename ite_type>
78 __device__ inline static bool set_keys(grid_key_dx<3,int> & key, grid_key_dx<3,int> & keyg, ite_type & itg)
79 {
80#ifdef __NVCC__
81
82 key.set_d(0,threadIdx.x + blockIdx.x * blockDim.x + itg.start.get(0));
83 key.set_d(1,threadIdx.y + blockIdx.y * blockDim.y + itg.start.get(1));
84 key.set_d(2,threadIdx.z + blockIdx.z * blockDim.z + itg.start.get(2));
85
86 keyg.set_d(0,key.get(0) + itg.origin.get(0));
87 keyg.set_d(1,key.get(1) + itg.origin.get(1));
88 keyg.set_d(2,key.get(2) + itg.origin.get(2));
89
90 if (key.get(0) > itg.stop.get(0) || key.get(1) > itg.stop.get(1) || key.get(2) > itg.stop.get(2) ||
91 key.get(0) < itg.start_base.get(0) || key.get(1) < itg.start_base.get(1) || key.get(2) < itg.start_base.get(2))
92 {return true;}
93#endif
94 return false;
95 }
96};
97
98template<>
99struct launch_insert_sparse_lambda_call<2>
100{
101 template<typename grid_type, typename lambda_t1, typename lambda_t2,typename itd_type, typename coord_type>
102 __device__ inline static void call(grid_type & grid,
103 lambda_t1 f1, lambda_t2 f2,
104 unsigned int blockId,
105 itd_type itd,
106 coord_type & key,
107 coord_type & keyg,unsigned int offset, bool & is_block_empty,
108 bool is_in)
109 {
110#ifdef __NVCC__
111
112 bool is_active = false;
113 if (is_in == true)
114 {is_active = f1(keyg.get(0),keyg.get(1));}
115
116 if (is_active == true)
117 {is_block_empty = false;}
118
119 __syncthreads();
120
121 if (is_block_empty == false)
122 {
123 auto ec = grid.insertBlock(blockId);
124 enc_num<decltype(grid.insertBlock(blockId))> ecn(ec,offset);
125
126 if ( is_active == true)
127 {
128 f2(ecn,keyg.get(0),keyg.get(1));
129 ec.template get<grid_type::pMask>()[offset] = 1;
130 }
131 }
132
133#endif
134 }
135
136 template<typename ite_type>
137 __device__ inline static bool set_keys(grid_key_dx<2,int> & key, grid_key_dx<2,int> & keyg, ite_type & itg)
138 {
139#ifdef __NVCC__
140 key.set_d(0,threadIdx.x + blockIdx.x * blockDim.x + itg.start.get(0));
141 key.set_d(1,threadIdx.y + blockIdx.y * blockDim.y + itg.start.get(1));
142
143 keyg.set_d(0,key.get(0) + itg.origin.get(0));
144 keyg.set_d(1,key.get(1) + itg.origin.get(1));
145
146 if (key.get(0) > itg.stop.get(0) || key.get(1) > itg.stop.get(1) ||
147 key.get(0) < itg.start_base.get(0) || key.get(1) < itg.start_base.get(1))
148 {return true;}
149#endif
150 return false;
151 }
152};
153
154struct launch_insert_sparse
155{
156 template<typename grid_type, typename ite_type, typename lambda_f1, typename lambda_f2>
157 __device__ void operator()(grid_type & grid, ite_type itg, bool & is_block_empty, lambda_f1 f1, lambda_f2 f2)
158 {
159#ifdef __NVCC__
160
161 grid_key_dx<grid_type::dims,int> key;
162 grid_key_dx<grid_type::dims,int> keyg;
163
164 bool not_active = launch_insert_sparse_lambda_call<grid_type::dims>::set_keys(key,keyg,itg);
165
166 if (threadIdx.x == 0 && threadIdx.y == 0 && threadIdx.z == 0)
167 {is_block_empty = true;}
168
169 grid.init();
170
171 int offset = 0;
172 grid_key_dx<grid_type::dims,int> blk;
173 bool out = grid.template getInsertBlockOffset<ite_type>(itg,key,blk,offset);
174
175 auto blockId = grid.getBlockLinId(blk);
176
177 launch_insert_sparse_lambda_call<grid_type::dims>::call(grid,f1,f2,blockId,itg,key,keyg,offset,is_block_empty,!not_active);
178
179 __syncthreads();
180
181 grid.flush_block_insert();
182#endif
183 }
184};
185
186template<bool is_free>
187struct selvg
188{
189 template<typename a_it_type, typename gdb_ext_type, typename gList_type>
190 static inline void call(a_it_type & a_it, gdb_ext_type & gdb_ext, gList_type & gList, size_t & g_c)
191 {
192 if (gdb_ext.get(g_c).Dbox.isValid() == false)
193 {g_c++;}
194 else
195 {
196 a_it.reinitialize(gList.get(g_c).getIterator(gdb_ext.get(g_c).Dbox.getKP1(),gdb_ext.get(g_c).Dbox.getKP2()));
197 if (a_it.isNext() == false) {g_c++;}
198 }
199 }
200};
201
202template<>
203struct selvg<false>
204{
205 template<typename a_it_type, typename gdb_ext_type, typename gList_type>
206 static inline void call(a_it_type & a_it, gdb_ext_type & gdb_ext, gList_type & gList, size_t & g_c)
207 {
208 // Full iterator (no subset)
209 a_it.reinitialize(gList.get(g_c).getIterator());
210 if (a_it.isNext() == false) {g_c++;}
211 }
212};
213
214/*! \brief Distributed grid iterator
215 *
216 * Iterator across the local elements of the distributed grid
217 *
218 * \tparam dim dimensionality of the grid
219 * \tparam device_grid type of basic grid
220 * \tparam stencil it inject the code to calculate stencil offset
221 * \tparam sub_iterator it indicate the sub-iterator type of the device_grid
222 *
223 */
224template<unsigned int dim, typename device_grid, typename device_sub_it, int impl, typename stencil = no_stencil >
225class grid_dist_iterator
226{
227 //! grid list counter
228 size_t g_c;
229
230 //! List of the grids we are going to iterate
231 const openfpm::vector<device_grid> & gList;
232
233 //! Extension of each grid: domain and ghost + domain
234 const openfpm::vector<GBoxes<device_grid::dims>> & gdb_ext;
235
236 //! Actual iterator
237 device_sub_it a_it;
238
239 //! stop point (is the grid size)
240 grid_key_dx<dim> stop;
241
242 /*! \brief from g_c increment g_c until you find a valid grid
243 *
244 */
245 void selectValidGrid()
246 {
247 do
248 {
249 if (impl == FREE)
250 {
251 // When the grid has size 0 potentially all the other informations are garbage
252 while (g_c < gList.size() && (gList.get(g_c).size() == 0 || gdb_ext.get(g_c).Dbox.isValid() == false ) ) g_c++;
253 }
254 else
255 {
256 // When the grid has size 0 potentially all the other informations are garbage
257 while (g_c < gList.size() && (gList.get(g_c).size() == 0 || gdb_ext.get(g_c).GDbox.isValid() == false) ) g_c++;
258 }
259
260 // get the next grid iterator
261 if (g_c < gList.size())
262 {
263 selvg<impl == FREE>::call(a_it,gdb_ext,gList,g_c);
264 }
265 } while (g_c < gList.size() && a_it.isNext() == false);
266
267 }
268
269 public:
270
271 /*! \brief Constructor of the distributed grid iterator
272 *
273 * \param gk std::vector of the local grid
274 * \param gdb_ext set of local subdomains
275 * \param stop end point
276 *
277 */
278 grid_dist_iterator(const openfpm::vector<device_grid> & gk,
279 const openfpm::vector<GBoxes<device_grid::dims>> & gdb_ext,
280 const grid_key_dx<dim> & stop)
281 :g_c(0),gList(gk),gdb_ext(gdb_ext),stop(stop)
282 {
283 // Initialize the current iterator
284 // with the first grid
285 selectValidGrid();
286 }
287
288
289 /*! \brief Constructor of the distributed grid iterator with
290 * stencil support
291 *
292 * \param gk std::vector of the local grid
293 * \param gdb_ext set of local subdomains
294 * \param stop end point
295 * \param stencil_pnt stencil points
296 *
297 */
298 grid_dist_iterator(openfpm::vector<device_grid> & gk,
299 const openfpm::vector<GBoxes<device_grid::dims>> & gdb_ext,
300 const grid_key_dx<dim> & stop,
301 const grid_key_dx<dim> (& stencil_pnt)[stencil::nsp])
302 :g_c(0),gList(gk),gdb_ext(gdb_ext),a_it(stencil_pnt),stop(stop)
303 {
304 // Initialize the current iterator
305 // with the first grid
306 selectValidGrid();
307 }
308
309 //! Copy constructor
310 grid_dist_iterator(const grid_dist_iterator<dim,device_grid,device_sub_it,impl,stencil> & g)
311 :g_c(g.g_c),gList(g.gList),gdb_ext(g.gdb_ext),a_it(g.a_it),stop(g.stop)
312 {}
313
314 //! Copy constructor
315 grid_dist_iterator(grid_dist_iterator<dim,device_grid,device_sub_it,impl,stencil> && g)
316 :g_c(g.g_c),gList(g.gList),gdb_ext(g.gdb_ext),a_it(g.a_it),stop(g.stop)
317 {}
318
319 //! Destructor
320 ~grid_dist_iterator()
321 {
322 }
323
324 /*! \brief Get the next element
325 *
326 * \return the next grid_key
327 *
328 */
329 inline grid_dist_iterator<dim,device_grid,device_sub_it,impl,stencil> & operator++()
330 {
331 ++a_it;
332
333 // check if a_it is at the end
334
335 if (a_it.isNext() == true)
336 return *this;
337 else
338 {
339 // switch to the new grid
340 g_c++;
341
342 selectValidGrid();
343 }
344
345 return *this;
346 }
347
348 /*! \brief Check if there is the next element
349 *
350 * \return true if there is the next, false otherwise
351 *
352 */
353 inline bool isNext() const
354 {
355 // If there are no other grid stop
356
357 if (g_c >= gList.size())
358 {return false;}
359
360 return true;
361 }
362
363 /*! \brief Get the actual key
364 *
365 * \return the actual key
366 *
367 */
368 inline grid_dist_key_dx<dim, typename device_grid::base_key> get() const
369 {
370 return grid_dist_key_dx<dim,typename device_grid::base_key>(g_c,a_it.get());
371 }
372
373 /*! \brief it return the stop point of the iterator
374 *
375 * The stop point of the iterator is just the grid size
376 *
377 * \return the stop point
378 *
379 */
380 inline grid_key_dx<dim> getStop() const
381 {
382 return stop;
383 }
384
385 /*! \brief it return the start point of the iterator
386 *
387 * The start point of the iterator is the point with all coordinates zeros
388 *
389 * \return the start point
390 *
391 */
392 inline grid_key_dx<dim> getStart() const
393 {
394 grid_key_dx<dim> start;
395
396 start.zero();
397
398 return start;
399 }
400
401 /*! \brief Get the boxes
402 *
403 * Get the boxes that define the local grids
404 *
405 * \return Vector of local boxes
406 *
407 */
408 inline const openfpm::vector<GBoxes<device_grid::dims>> & getGBoxes()
409 {
410 return gdb_ext;
411 }
412
413 /*! \brief Convert a g_dist_key_dx into a global key
414 *
415 * \see grid_dist_key_dx
416 * \see grid_dist_iterator
417 *
418 * \param k key position in local coordinates
419 *
420 * \return the global position in the grid
421 *
422 */
423 inline grid_key_dx<dim> getGKey(const grid_dist_key_dx<dim,typename device_grid::base_key> & k)
424 {
425 // Get the sub-domain id
426 size_t sub_id = k.getSub();
427
428 auto k_glob = k.getKey();
429
430 // shift
431 auto k_glob2 = k_glob + gdb_ext.get(sub_id).origin;
432
433 return k_glob2;
434 }
435
436 /*! \brief Return the stencil point offset
437 *
438 * \tparam id
439 *
440 * \return linearized distributed key
441 *
442 */
443 template<unsigned int id> inline grid_dist_lin_dx getStencil()
444 {
445 return grid_dist_lin_dx(g_c,a_it.template getStencil<id>());
446 }
447};
448
449
450
451#endif /* GRID_DIST_ID_ITERATOR_SUB_HPP_ */
452