OpenFPM_pdata  4.1.0
Project that contain the implementation of distributed structures
 
Loading...
Searching...
No Matches
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
21struct amr_launch_sparse
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;
36 grid_key_dx<3,int> blk;
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
73BOOST_AUTO_TEST_SUITE( amr_grid_dist_id_test )
74
75
76BOOST_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
194BOOST_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
323BOOST_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
513BOOST_AUTO_TEST_SUITE_END()
Box
This class represent an N-dimensional box.
Definition Box.hpp:61
Point
This class implement the point shape in an N-dimensional space.
Definition Point.hpp:28
Point::get
__device__ __host__ const T & get(unsigned int i) const
Get coordinate.
Definition Point.hpp:172
grid_dist_id
This is a distributed grid.
Definition grid_dist_id.hpp:278
grid_key_dx
grid_key_dx is the key to access any element in the grid
Definition grid_key.hpp:19
amr_launch_sparse
Definition amr_base_gpu_unit_tests.cu:22