OpenFPM_pdata  4.1.0
Project that contain the implementation of distributed structures
amr_base_gpu_unit_tests.cu
1 /*
2  * amr_base_gpu_unit_tests.cu
3  *
4  * Created on: Aug 28, 2019
5  * Author: i-bird
6  */
7 
8 /*
9  * amr_base_unit_test.cpp
10  *
11  * Created on: Oct 5, 2017
12  * Author: i-bird
13  */
14 #define BOOST_TEST_DYN_LINK
15 #include <boost/test/unit_test.hpp>
16 
17 #include "Amr/grid_dist_amr.hpp"
18 #include "Point_test.hpp"
19 #include "Grid/tests/grid_dist_id_util_tests.hpp"
20 
22 {
23  template<typename grid_type, typename ite_type>
24  __device__ void operator()(grid_type & grid, ite_type itg, float spacing, Point<3,float> center)
25  {
26  GRID_ID_3_GLOBAL(itg);
27 
28  __shared__ bool is_block_empty;
29 
30  if (threadIdx.x == 0 && threadIdx.y == 0 && threadIdx.z == 0)
31  {is_block_empty = true;}
32 
33  grid.init();
34 
35  int offset = 0;
37  bool out = grid.getInsertBlockOffset(itg,key,blk,offset);
38 
39  auto blockId = grid.getBlockLinId(blk);
40 
41  const float x = keyg.get(0)*spacing - center.get(0);
42  const float y = keyg.get(1)*spacing - center.get(1);
43  const float z = keyg.get(2)*spacing - center.get(2);
44 
45  float radius = sqrt((float) (x*x + y*y + z*z));
46 
47  bool is_active = radius < 0.4 && radius > 0.3;
48 
49  if (is_active == true)
50  {is_block_empty = false;}
51 
52  __syncthreads();
53 
54  if (is_block_empty == false)
55  {
56  auto ec = grid.insertBlock(blockId);
57 
58  if ( is_active == true)
59  {
60  ec.template get<0>()[offset] = x+y+z;
61  ec.template get<grid_type::pMask>()[offset] = 1;
62  }
63  }
64 
65  __syncthreads();
66 
67  grid.flush_block_insert();
68  }
69 };
70 
71 
72 
73 BOOST_AUTO_TEST_SUITE( amr_grid_dist_id_test )
74 
75 
76 BOOST_AUTO_TEST_CASE( grid_dist_id_amr_gpu )
77 {
78  // Domain
79  Box<3,float> domain3({0.0,0.0,0.0},{1.0,1.0,1.0});
80 
81 
82  Ghost<3,long int> g(1);
83  sgrid_dist_amr_gpu<3,float,aggregate<float>> amr_g(domain3,g);
84 
85  size_t g_sz[3] = {4,4,4};
86 
87  size_t n_lvl = 6;
88 
89  amr_g.initLevels(n_lvl,g_sz);
90 
91  for (size_t i = 0 ; i < amr_g.getNLvl() ; i++)
92  {
93  // Fill the AMR with something
94 
95  size_t count = 0;
96 
97  auto it = amr_g.getGridIteratorGPU(i);
98  it.setGPUInsertBuffer(1);
99 
100  Point<3,float> center({0.5,0.5,0.5});
101 
102  it.launch(amr_launch_sparse(),it.getSpacing(0),center);
103  amr_g.getDistGrid(i).template flush<smax_<0>>(FLUSH_ON_DEVICE);
104 
105  amr_g.getDistGrid(i).template deviceToHost<0>();
106 
107  auto it2 = amr_g.getDistGrid(i).getDomainIterator();
108 
109  while (it2.isNext())
110  {
111  auto key = it2.get();
112  auto keyg = it2.getGKey(key);
113 
114  count++;
115 
116  ++it2;
117  }
118 
119  auto & v_cl = create_vcluster();
120 
121  v_cl.sum(count);
122  v_cl.execute();
123 
124  switch(i)
125  {
126  case 0:
127  BOOST_REQUIRE_EQUAL(count,0);
128  break;
129  case 1:
130  BOOST_REQUIRE_EQUAL(count,30);
131  break;
132  case 2:
133  BOOST_REQUIRE_EQUAL(count,282);
134  break;
135  case 3:
136  BOOST_REQUIRE_EQUAL(count,2192);
137  break;
138  case 4:
139  BOOST_REQUIRE_EQUAL(count,16890);
140  break;
141  case 5:
142  BOOST_REQUIRE_EQUAL(count,136992);
143  break;
144  }
145  }
146 
147  // Iterate across all the levels initialized
148 /* auto it = amr_g.getDomainIterator();
149 
150  size_t count = 0;
151 
152  while (it.isNext())
153  {
154  count++;
155 
156  ++it;
157  }
158 
159  Vcluster<> & v_cl = create_vcluster();
160 
161  v_cl.sum(count);
162  v_cl.execute();
163 
164  BOOST_REQUIRE_EQUAL(count,correct_result);
165 
166  auto itc = amr_g.getDomainIteratorCells();
167 
168  size_t count_c = 0;
169 
170  while (itc.isNext())
171  {
172  count_c++;
173 
174  ++itc;
175  }
176 
177  v_cl.sum(count_c);
178  v_cl.execute();
179 
180  auto it_level = amr_g.getDomainIteratorCells(3);
181 
182  while (it_level.isNext())
183  {
184  auto key = it_level.get();
185 
186  amr_g.template get<0>(3,key);
187 
188  ++it_level;
189  }
190 
191  BOOST_REQUIRE_EQUAL(count_c,correct_result_cell);*/
192 }
193 
194 BOOST_AUTO_TEST_CASE( grid_dist_id_amr_gpu_link_test )
195 {
196 // To uncomment (It does not work when we run the full suite for some weird reason)
197 
198 #if 0
199 
200  auto & v_cl = create_vcluster();
201 
202  // Domain
203  Box<2,float> domain({0.0,0.0},{1.0,1.0});
204 
205  Ghost<2,long int> g(1);
206  sgrid_dist_amr_gpu<2,float,aggregate<float>> amr_g(domain,g);
207 
208  size_t g_sz[2] = {17,17};
209 
210  size_t n_lvl = 3;
211 
212  amr_g.initLevels(n_lvl,g_sz);
213 
214  grid_key_dx<2> start({5,5});
215  grid_key_dx<2> start_lvl_dw({9,9});
216  grid_key_dx<2> stop_lvl_dw({12,12});
217  grid_key_dx<2> start_lvl_dw2({19,19});
218  grid_key_dx<2> stop_lvl_dw2({23,23});
219 
220  auto it = amr_g.getGridIterator(0,start,start);
221  auto it2 = amr_g.getGridIterator(1,start_lvl_dw,stop_lvl_dw);
222  auto it3 = amr_g.getGridIterator(2,start_lvl_dw2,stop_lvl_dw2);
223 // it.setGPUInsertBuffer(4);
224 
225  auto & lvl_0 = amr_g.getDistGrid(0);
226  auto & lvl_1 = amr_g.getDistGrid(1);
227  auto & lvl_2 = amr_g.getDistGrid(2);
228 
229  // Add points in level 0
230 
231  while (it.isNext())
232  {
233  auto key = it.get_dist();
234 
235  lvl_0.template insertFlush<0>(key) = 1.0;
236 
237  ++it;
238  }
239 
240  while (it2.isNext())
241  {
242  auto key = it2.get_dist();
243 
244  lvl_1.template insertFlush<0>(key) = 2.0;
245 
246  ++it2;
247  }
248 
249  while (it3.isNext())
250  {
251  auto key = it3.get_dist();
252 
253  lvl_2.template insertFlush<0>(key) = 3.0;
254 
255  ++it3;
256  }
257 
258  amr_g.hostToDevice<0>();
259  amr_g.tagBoundaries<NNStar<2>>();
260  amr_g.construct_level_connections();
261 
263 
264  auto & lvl_zero_d = amr_g.getDistGrid(0);
265  auto & lvl_one_d = amr_g.getDistGrid(1);
266  auto & lvl_two_d = amr_g.getDistGrid(2);
267 
268  // For each local grid
269 
270  for (int i = 0 ; i < lvl_zero_d.getN_loc_grid() ; i++)
271  {
272 
273  // Check
274  auto & lvl_zero = lvl_zero_d.get_loc_grid(i);
275  auto & lvl_one = lvl_one_d.get_loc_grid(i);
276  auto & lvl_two = lvl_two_d.get_loc_grid(i);
277 
278  auto & offs_dw_link = lvl_zero.getDownLinksOffsets();
279  auto & dw_links = lvl_zero.getDownLinks();
280 
281  BOOST_REQUIRE_EQUAL(offs_dw_link.size(),1);
282  BOOST_REQUIRE_EQUAL(dw_links.size(),4);
283 
284  auto & indexL0 = lvl_zero.private_get_blockMap().getIndexBuffer();
285  auto & indexL1 = lvl_one.private_get_blockMap().getIndexBuffer();
286  auto & indexL2 = lvl_two.private_get_blockMap().getIndexBuffer();
287 
288  auto & dataL0 = lvl_zero.private_get_blockMap().getDataBuffer();
289  auto & dataL1 = lvl_one.private_get_blockMap().getDataBuffer();
290  auto & dataL2 = lvl_two.private_get_blockMap().getDataBuffer();
291 
292  dw_links.template deviceToHost<0,1>();
293 
294  BOOST_REQUIRE_EQUAL(dataL1.template get<0>(dw_links.template get<0>(0))[dw_links.template get<1>(0)],2);
295  BOOST_REQUIRE_EQUAL(dataL1.template get<0>(dw_links.template get<0>(1))[dw_links.template get<1>(1)],2);
296  BOOST_REQUIRE_EQUAL(dataL1.template get<0>(dw_links.template get<0>(2))[dw_links.template get<1>(2)],2);
297  BOOST_REQUIRE_EQUAL(dataL1.template get<0>(dw_links.template get<0>(3))[dw_links.template get<1>(3)],2);
298 
299  auto & offs_dw_link_1 = lvl_one.getDownLinksOffsets();
300  auto & dw_links_1 = lvl_one.getDownLinks();
301 
302  BOOST_REQUIRE_EQUAL(offs_dw_link_1.size(),12);
303  BOOST_REQUIRE_EQUAL(dw_links_1.size(),9);
304 
305  dw_links_1.template deviceToHost<0,1>();
306 
307  BOOST_REQUIRE_EQUAL(dataL2.template get<0>(dw_links_1.template get<0>(0))[dw_links_1.template get<1>(0)],3);
308  BOOST_REQUIRE_EQUAL(dataL2.template get<0>(dw_links_1.template get<0>(1))[dw_links_1.template get<1>(1)],3);
309  BOOST_REQUIRE_EQUAL(dataL2.template get<0>(dw_links_1.template get<0>(2))[dw_links_1.template get<1>(2)],3);
310  BOOST_REQUIRE_EQUAL(dataL2.template get<0>(dw_links_1.template get<0>(3))[dw_links_1.template get<1>(3)],3);
311  BOOST_REQUIRE_EQUAL(dataL2.template get<0>(dw_links_1.template get<0>(4))[dw_links_1.template get<1>(4)],3);
312  BOOST_REQUIRE_EQUAL(dataL2.template get<0>(dw_links_1.template get<0>(5))[dw_links_1.template get<1>(5)],3);
313  BOOST_REQUIRE_EQUAL(dataL2.template get<0>(dw_links_1.template get<0>(6))[dw_links_1.template get<1>(6)],3);
314  BOOST_REQUIRE_EQUAL(dataL2.template get<0>(dw_links_1.template get<0>(7))[dw_links_1.template get<1>(7)],3);
315  BOOST_REQUIRE_EQUAL(dataL2.template get<0>(dw_links_1.template get<0>(8))[dw_links_1.template get<1>(8)],3);
316  }
317 
319 
320 #endif
321 }
322 
323 BOOST_AUTO_TEST_CASE( grid_dist_id_amr_gpu_link_test_more_dense )
324 {
325  // To uncomment (It does not work when we run the full suite for some weird reason)
326 
327  #if 0
328 
329  auto & v_cl = create_vcluster();
330 
331  // Domain
332  Box<2,float> domain({0.0,0.0},{1.0,1.0});
333 
334  Ghost<2,long int> g(1);
335  sgrid_dist_amr_gpu<2,float,aggregate<float>> amr_g(domain,g);
336 
337  size_t g_sz[2] = {17,17};
338 
339  size_t n_lvl = 3;
340 
341  amr_g.initLevels(n_lvl,g_sz);
342 
343  grid_key_dx<2> start({1,1});
344  grid_key_dx<2> stop({15,15});
345  grid_key_dx<2> start_lvl_dw({2,2});
346  grid_key_dx<2> stop_lvl_dw({31,31});
347  grid_key_dx<2> start_lvl_dw2({4,4});
348  grid_key_dx<2> stop_lvl_dw2({63,63});
349 
350  auto it = amr_g.getGridIterator(0,start,stop);
351  auto it2 = amr_g.getGridIterator(1,start_lvl_dw,stop_lvl_dw);
352  auto it3 = amr_g.getGridIterator(2,start_lvl_dw2,stop_lvl_dw2);
353 
354  auto & lvl_0 = amr_g.getDistGrid(0);
355  auto & lvl_1 = amr_g.getDistGrid(1);
356  auto & lvl_2 = amr_g.getDistGrid(2);
357 
358  // Add points in level 0
359 
360  while (it.isNext())
361  {
362  auto key = it.get_dist();
363 
364  lvl_0.template insertFlush<0>(key) = 1.0;
365 
366  ++it;
367  }
368 
369  while (it2.isNext())
370  {
371  auto key = it2.get_dist();
372 
373  lvl_1.template insertFlush<0>(key) = 2.0;
374 
375  ++it2;
376  }
377 
378  while (it3.isNext())
379  {
380  auto key = it3.get_dist();
381 
382  lvl_2.template insertFlush<0>(key) = 3.0;
383 
384  ++it3;
385  }
386 
387  amr_g.hostToDevice<0>();
388  amr_g.ghost_get<0>(RUN_ON_DEVICE);
389  amr_g.tagBoundaries<NNStar<2>>();
390  amr_g.ghost_get<0>(RUN_ON_DEVICE);
391  amr_g.construct_level_connections();
392  amr_g.deviceToHost<0>();
393  amr_g.write("TESTOUT");
394 
396 
397  auto & lvl_zero_d = amr_g.getDistGrid(0);
398  auto & lvl_one_d = amr_g.getDistGrid(1);
399  auto & lvl_two_d = amr_g.getDistGrid(2);
400 
401  // For each local grid
402 
403  size_t tot_dw_offs_12 = 0;
404  size_t tot_dw_lk_12 = 0;
405 
406  size_t tot_dw_offs_23 = 0;
407  size_t tot_dw_lk_23 = 0;
408 
409  size_t tot_up_offs_12 = 0;
410  size_t tot_up_lk_12 = 0;
411 
412  size_t tot_up_offs_23 = 0;
413  size_t tot_up_lk_23 = 0;
414 
415  for (int i = 0 ; i < lvl_zero_d.getN_loc_grid() ; i++)
416  {
417 
418  // Check
419  auto & lvl_zero = lvl_zero_d.get_loc_grid(i);
420  auto & lvl_one = lvl_one_d.get_loc_grid(i);
421  auto & lvl_two = lvl_two_d.get_loc_grid(i);
422 
423  auto & offs_dw_link = lvl_zero.getDownLinksOffsets();
424  auto & dw_links = lvl_zero.getDownLinks();
425 
426  auto & offs_up_link = lvl_one.getUpLinksOffsets();
427  auto & up_links = lvl_one.getUpLinks();
428 
429  tot_dw_offs_12 += offs_dw_link.size();
430  tot_dw_lk_12 += dw_links.size();
431 
432  tot_up_offs_12 += offs_up_link.size();
433  tot_up_lk_12 += up_links.size();
434 
435  auto & indexL0 = lvl_zero.private_get_blockMap().getIndexBuffer();
436  auto & indexL1 = lvl_one.private_get_blockMap().getIndexBuffer();
437  auto & indexL2 = lvl_two.private_get_blockMap().getIndexBuffer();
438 
439  auto & dataL0 = lvl_zero.private_get_blockMap().getDataBuffer();
440  auto & dataL1 = lvl_one.private_get_blockMap().getDataBuffer();
441  auto & dataL2 = lvl_two.private_get_blockMap().getDataBuffer();
442 
443  dw_links.template deviceToHost<0,1>();
444  up_links.template deviceToHost<0,1>();
445 
446  for (int i = 0 ; i < dw_links.size(); i++)
447  {
448  BOOST_REQUIRE_EQUAL(dataL1.template get<0>(dw_links.template get<0>(i))[dw_links.template get<1>(i)],2);
449  }
450 
451  for (int i = 0 ; i < up_links.size(); i++)
452  {
453  BOOST_REQUIRE_EQUAL(dataL0.template get<0>(up_links.template get<0>(i))[up_links.template get<1>(i)],1);
454  }
455 
456  auto & offs_dw_link_1 = lvl_one.getDownLinksOffsets();
457  auto & dw_links_1 = lvl_one.getDownLinks();
458 
459  auto & offs_up_link_1 = lvl_two.getUpLinksOffsets();
460  auto & up_links_1 = lvl_two.getUpLinks();
461 
462  tot_dw_offs_23 += offs_dw_link_1.size();
463  tot_dw_lk_23 += dw_links_1.size();
464 
465  tot_up_offs_23 += offs_up_link_1.size();
466  tot_up_lk_23 += up_links_1.size();
467 
468  dw_links_1.template deviceToHost<0,1>();
469  up_links_1.template deviceToHost<0,1>();
470 
471  for (int i = 0 ; i < dw_links_1.size(); i++)
472  {
473  BOOST_REQUIRE_EQUAL(dataL2.template get<0>(dw_links_1.template get<0>(i))[dw_links_1.template get<1>(i)],3);
474  }
475 
476  for (int i = 0 ; i < up_links_1.size(); i++)
477  {
478  BOOST_REQUIRE_EQUAL(dataL1.template get<0>(up_links_1.template get<0>(i))[up_links_1.template get<1>(i)],2);
479  }
480  }
481 
482  v_cl.sum(tot_dw_offs_12);
483  v_cl.sum(tot_dw_lk_12);
484 
485  v_cl.sum(tot_dw_offs_23);
486  v_cl.sum(tot_dw_lk_23);
487 
488  v_cl.sum(tot_up_offs_12);
489  v_cl.sum(tot_up_lk_12);
490 
491  v_cl.sum(tot_up_offs_23);
492  v_cl.sum(tot_up_lk_23);
493 
494  v_cl.execute();
495 
496  BOOST_REQUIRE_EQUAL(tot_dw_offs_12,56);
497  BOOST_REQUIRE_EQUAL(tot_dw_lk_12,56*4);
498 
499  BOOST_REQUIRE_EQUAL(tot_dw_offs_23,116);
500  BOOST_REQUIRE_EQUAL(tot_dw_lk_23,116*4);
501 
502  BOOST_REQUIRE_EQUAL(tot_up_offs_12,116);
503  BOOST_REQUIRE_EQUAL(tot_up_lk_12,116);
504 
505  BOOST_REQUIRE_EQUAL(tot_up_offs_23,236);
506  BOOST_REQUIRE_EQUAL(tot_up_lk_23,236);
507 
509 
510 #endif
511 }
512 
513 BOOST_AUTO_TEST_SUITE_END()
grid_key_dx is the key to access any element in the grid
Definition: grid_key.hpp:18
Definition: Ghost.hpp:39
This is a distributed grid.
__device__ __host__ const T & get(unsigned int i) const
Get coordinate.
Definition: Point.hpp:172