OpenFPM_pdata  4.1.0
Project that contain the implementation of distributed structures
sgrid_dist_id_gpu_unit_tests.cu
1 #include <ostream>
2 #define BOOST_TEST_DYN_LINK
3 
4 #include "config.h"
5 #include <boost/test/unit_test.hpp>
6 #include "Grid/grid_dist_id.hpp"
7 
8 
9 BOOST_AUTO_TEST_SUITE( sgrid_gpu_test_suite )
10 
11 template<unsigned int p>
13 {
14  template<typename SparseGridGpu_type, typename ite_type>
15  __device__ void operator()(SparseGridGpu_type & sg, ite_type & ite, float c)
16  {
17  GRID_ID_2_GLOBAL(ite);
18 
19  sg.init();
20 
21  if (inactive == false)
22  {sg.template insert<p>(key) = c + keyg.get(0) + keyg.get(1);}
23 
24  __syncthreads();
25 
26  sg.flush_block_insert();
27  }
28 };
29 
30 
31 
32 template<unsigned int p>
34 {
35  template<typename SparseGridGpu_type, typename ite_type>
36  __device__ void operator()(SparseGridGpu_type & sg, ite_type & ite, float c)
37  {
38  GRID_ID_3_GLOBAL(ite);
39 
40  sg.init();
41 
42  if (inactive == false)
43  {sg.template insert<p>(key) = c + keyg.get(0) + keyg.get(1) + keyg.get(2);}
44 
45  __syncthreads();
46 
47  sg.flush_block_insert();
48  }
49 };
50 
51 
52 BOOST_AUTO_TEST_CASE( sgrid_gpu_test_base )
53 {
54  size_t sz[2] = {17,17};
55  periodicity<2> bc = {PERIODIC,PERIODIC};
56 
57  Ghost<2,long int> g(1);
58 
59  Box<2,float> domain({0.0,0.0},{1.0,1.0});
60 
61  sgrid_dist_id_gpu<2,float,aggregate<float>> gdist(sz,domain,g,bc);
62 
63  gdist.template setBackgroundValue<0>(666);
64 
66 
67  Box<2,size_t> box({1,1},{1,1});
68  auto it = gdist.getGridIterator(box.getKP1(),box.getKP2());
69 
71 
72  float c = 5.0;
73 
74  gdist.template iterateGridGPU<insert_kernel2D<0>>(it,c);
75  gdist.template flush<smax_<0>>(flush_type::FLUSH_ON_DEVICE);
76 
77  gdist.template deviceToHost<0>();
78 
79  {
80  Box<2,size_t> box2({0,0},{15,15});
81 
82  auto it = gdist.getGridIterator(box2.getKP1(),box2.getKP2());
83 
84  while (it.isNext())
85  {
86  auto p = it.get_dist();
87  auto p2 = it.get();
88 
89  if (p2.get(0) == box.getLow(0) && p2.get(1) == box.getLow(1))
90  {
91  BOOST_REQUIRE_EQUAL(gdist.template get<0>(p), 7.0);
92  }
93  else
94  {
95  BOOST_REQUIRE_EQUAL(gdist.template get<0>(p), 666.0);
96  }
97 
98  ++it;
99  }
100  }
101 
102  //
103 
104  c = 3.0;
105 
106  Box<2,size_t> box3({3,3},{11,11});
107 
108  auto it3 = gdist.getGridIterator(box3.getKP1(),box3.getKP2());
109 
110  gdist.template iterateGridGPU<insert_kernel2D<0>>(it3,c);
111  gdist.template flush<smax_<0>>(flush_type::FLUSH_ON_DEVICE);
112  gdist.template deviceToHost<0>();
113 
114  {
115  Box<2,size_t> box2({0,0},{15,15});
116 
117  auto it = gdist.getGridIterator(box2.getKP1(),box2.getKP2());
118 
119  while (it.isNext())
120  {
121  auto p = it.get_dist();
122  auto p2 = it.get();
123 
124  Point<2,size_t> p2_ = p2.toPoint();
125 
126  if (box.isInside(p2_))
127  {
128  BOOST_REQUIRE_EQUAL(gdist.template get<0>(p), 7.0);
129  }
130  else if (box3.isInside(p2_))
131  {
132  float tst = c + p2.get(0) + p2.get(1);
133 
134  BOOST_REQUIRE_EQUAL(gdist.template get<0>(p), tst);
135  }
136  else
137  {
138  BOOST_REQUIRE_EQUAL(gdist.template get<0>(p), 666.0);
139  }
140 
141  ++it;
142  }
143  }
144 }
145 
146 
147 BOOST_AUTO_TEST_CASE( sgrid_gpu_test_output )
148 {
149  auto & v_cl = create_vcluster();
150 
151  if (v_cl.size() > 3){return;}
152 
153  size_t sz[2] = {17,17};
154  periodicity<2> bc = {PERIODIC,PERIODIC};
155 
156  Ghost<2,long int> g(1);
157 
158  Box<2,float> domain({0.0,0.0},{1.0,1.0});
159 
160  sgrid_dist_id_gpu<2,float,aggregate<float>> gdist(sz,domain,g,bc);
161 
162  gdist.template setBackgroundValue<0>(666);
163 
165 
166  Box<2,size_t> box({1,1},{15,15});
167  auto it = gdist.getGridIterator(box.getKP1(),box.getKP2());
168 
170 
171  float c = 5.0;
172 
173  gdist.template iterateGridGPU<insert_kernel2D<0>>(it,c);
174  gdist.template flush<smax_<0>>(flush_type::FLUSH_ON_DEVICE);
175 
176  gdist.template deviceToHost<0>();
177 
178  gdist.write("sgrid_gpu_output");
179 
180  std::string file_test("sgrid_gpu_output_" + std::to_string(v_cl.size()) + "_" + std::to_string(v_cl.rank()) + ".vtk");
181  std::string file("sgrid_gpu_output_" + std::to_string(v_cl.rank()) + ".vtk");
182 
183  #ifndef HAVE_OSX
184  bool test = compare(file,"test_data/" + file_test);
185 
186  BOOST_REQUIRE_EQUAL(true,test);
187  #endif
188 }
189 
190 
191 BOOST_AUTO_TEST_CASE( sgrid_gpu_test_save_and_load )
192 {
193  auto & v_cl = create_vcluster();
194 
195  if (v_cl.size() > 3){return;}
196 
197  size_t sz[2] = {370,370};
198  periodicity<2> bc = {PERIODIC,PERIODIC};
199 
200  Ghost<2,long int> g(1);
201 
202  Box<2,float> domain({0.0,0.0},{1.0,1.0});
203 
204  sgrid_dist_id_gpu<2,float,aggregate<float,float>> gdist(sz,domain,g,bc);
205 
206  gdist.template setBackgroundValue<0>(666);
207 
209 
210  Box<2,size_t> box({1,1},{350,350});
211  auto it = gdist.getGridIterator(box.getKP1(),box.getKP2());
212 
214 
215  typedef typename GetAddBlockType<decltype(gdist)>::type InsertBlockT;
216 
217  float c = 5.0;
218 
219  gdist.addPoints([] __device__ (int i, int j)
220  {
221  return true;
222  },
223  [c] __device__ (InsertBlockT & data, int i, int j)
224  {
225  data.template get<0>() = c + i + j;
226  data.template get<1>() = c + 1000 + i + j;
227  }
228  );
229 
230  gdist.template flush<smax_<0>,smax_<1>>(flush_type::FLUSH_ON_DEVICE);
231 
232  gdist.template deviceToHost<0>();
233 
234  gdist.save("sgrid_gpu_output_hdf5");
235 
236  // Now load
237 
238  sgrid_dist_id_gpu<2,float,aggregate<float,float>> gdist2(sz,domain,g,bc);
239 
240  gdist2.load("sgrid_gpu_output_hdf5");
241 
242  gdist2.template ghost_get<0,1>(RUN_ON_DEVICE);
243 
244  gdist2.deviceToHost<0,1>();
245  gdist.deviceToHost<0,1>();
246 
247  bool match = true;
248 
249 
250  auto it2 = gdist2.getDomainIterator();
251 
252  while (it2.isNext())
253  {
254  auto p = it2.get();
255 
256  auto key = it2.getGKey(p);
257 
258  auto p_xp1 = p.move(0,1);
259  auto p_xm1 = p.move(0,-1);
260  auto p_yp1 = p.move(1,1);
261  auto p_ym1 = p.move(1,-1);
262 
263  auto key_xp1 = key.move(0,1);
264  auto key_xm1 = key.move(0,-1);
265  auto key_yp1 = key.move(1,1);
266  auto key_ym1 = key.move(1,-1);
267 
268  if (box.isInside(key_xp1.toPoint()))
269  {
270  match &= gdist.template get<0>(p_xp1) == c + key_xp1.get(0) + key_xp1.get(1);
271 
272  if (match == false)
273  {
274  std::cout << gdist.template get<0>(p_xp1) << " " << c + key_xp1.get(0) + key_xp1.get(1) << std::endl;
275  break;
276  }
277  }
278 
279  if (box.isInside(key_xm1.toPoint()))
280  {
281  match &= gdist.template get<0>(p_xm1) == c + key_xm1.get(0) + key_xm1.get(1);
282 
283  if (match == false)
284  {
285  std::cout << gdist.template get<0>(p_xm1) << " " << c + key_xm1.get(0) + key_xm1.get(1) << std::endl;
286  break;
287  }
288  }
289 
290  if (box.isInside(key_yp1.toPoint()))
291  {
292  match &= gdist.template get<0>(p_yp1) == c + key_yp1.get(0) + key_yp1.get(1);
293 
294  if (match == false)
295  {
296  std::cout << gdist.template get<0>(p_yp1) << " " << c + key_yp1.get(0) + key_yp1.get(1) << std::endl;
297  break;
298  }
299  }
300 
301  if (box.isInside(key_ym1.toPoint()))
302  {
303  match &= gdist.template get<0>(p_ym1) == c + key_ym1.get(0) + key_ym1.get(1);
304 
305  if (match == false)
306  {
307  std::cout << gdist.template get<0>(p_ym1) << " " << c + key_ym1.get(0) + key_ym1.get(1) << std::endl;
308  break;
309  }
310  }
311 
312  ++it2;
313  }
314 
315 
316  BOOST_REQUIRE_EQUAL(match,true);
317 }
318 
319 void sgrid_ghost_get(size_t (& sz)[2],size_t (& sz2)[2])
320 {
321  periodicity<2> bc = {PERIODIC,PERIODIC};
322 
323  Ghost<2,long int> g(1);
324 
325  Box<2,float> domain({0.0,0.0},{1.0,1.0});
326 
327  sgrid_dist_id_gpu<2,float,aggregate<float>> gdist(sz,domain,g,bc);
328 
329  gdist.template setBackgroundValue<0>(666);
330 
332 
333  Box<2,size_t> box({1,1},{sz2[0],sz2[1]});
334  auto it = gdist.getGridIterator(box.getKP1(),box.getKP2());
335 
337 
338  float c = 5.0;
339 
340  gdist.template iterateGridGPU<insert_kernel2D<0>>(it,c);
341  gdist.template flush<smax_<0>>(flush_type::FLUSH_ON_DEVICE);
342 
343  gdist.template deviceToHost<0>();
344 
345  gdist.template ghost_get<0>(RUN_ON_DEVICE);
346 
347  gdist.template deviceToHost<0>();
348 
349  // Now we check that ghost is correct
350 
351  auto it2 = gdist.getDomainIterator();
352 
353  bool match = true;
354 
355  while (it2.isNext())
356  {
357  auto p = it2.get();
358 
359  auto key = it2.getGKey(p);
360 
361  auto p_xp1 = p.move(0,1);
362  auto p_xm1 = p.move(0,-1);
363  auto p_yp1 = p.move(1,1);
364  auto p_ym1 = p.move(1,-1);
365 
366  auto key_xp1 = key.move(0,1);
367  auto key_xm1 = key.move(0,-1);
368  auto key_yp1 = key.move(1,1);
369  auto key_ym1 = key.move(1,-1);
370 
371  if (box.isInside(key_xp1.toPoint()))
372  {
373  match &= gdist.template get<0>(p_xp1) == c + key_xp1.get(0) + key_xp1.get(1);
374 
375  if (match == false)
376  {
377  std::cout << gdist.template get<0>(p_xp1) << " " << c + key_xp1.get(0) + key_xp1.get(1) << std::endl;
378  break;
379  }
380  }
381 
382  if (box.isInside(key_xm1.toPoint()))
383  {
384  match &= gdist.template get<0>(p_xm1) == c + key_xm1.get(0) + key_xm1.get(1);
385 
386  if (match == false)
387  {
388  std::cout << gdist.template get<0>(p_xm1) << " " << c + key_xm1.get(0) + key_xm1.get(1) << std::endl;
389  break;
390  }
391  }
392 
393  if (box.isInside(key_yp1.toPoint()))
394  {
395  match &= gdist.template get<0>(p_yp1) == c + key_yp1.get(0) + key_yp1.get(1);
396 
397  if (match == false)
398  {
399  std::cout << gdist.template get<0>(p_yp1) << " " << c + key_yp1.get(0) + key_yp1.get(1) << std::endl;
400  break;
401  }
402  }
403 
404  if (box.isInside(key_ym1.toPoint()))
405  {
406  match &= gdist.template get<0>(p_ym1) == c + key_ym1.get(0) + key_ym1.get(1);
407 
408  if (match == false)
409  {
410  std::cout << gdist.template get<0>(p_ym1) << " " << c + key_ym1.get(0) + key_ym1.get(1) << std::endl;
411  break;
412  }
413  }
414 
415  ++it2;
416  }
417 
418  BOOST_REQUIRE_EQUAL(match,true);
419 }
420 
421 BOOST_AUTO_TEST_CASE( sgrid_gpu_test_ghost_get )
422 {
423  size_t sz[2] = {17,17};
424  size_t sz6[2] = {15,15};
425  sgrid_ghost_get(sz,sz6);
426 
427  return;
428 
429  size_t sz2[2] = {170,170};
430  size_t sz3[2] = {15,15};
431  sgrid_ghost_get(sz2,sz3);
432 
433  size_t sz4[2] = {168,168};
434  sgrid_ghost_get(sz2,sz4);
435 }
436 
437 BOOST_AUTO_TEST_CASE( sgrid_gpu_app_point_test_no_box )
438 {
439  size_t sz[3] = {75,75,75};
440  periodicity<3> bc = {PERIODIC,PERIODIC,PERIODIC};
441 
442  Ghost<3,long int> g(1);
443 
444  Box<3,float> domain({0.0,0.0,0.0},{1.0,1.0,1.0});
445 
446  sgrid_dist_id_gpu<3,float,aggregate<float,float,float,float>> gdist(sz,domain,g,bc);
447 
448  gdist.template setBackgroundValue<0>(666);
449  gdist.template setBackgroundValue<1>(666);
450  gdist.template setBackgroundValue<2>(666);
451  gdist.template setBackgroundValue<3>(666);
452 
454 
455  Box<3,size_t> box({1,1,1},{sz[0],sz[1],sz[2]});
456 
458 
459  float c = 5.0;
460 
461  typedef typename GetAddBlockType<decltype(gdist)>::type InsertBlockT;
462 
463  CudaMemory cmem;
464  cmem.allocate(sizeof(int));
465 
466  *(int *)cmem.getPointer() = 0.0;
467 
468  cmem.hostToDevice();
469 
470  int * cnt = (int *)cmem.getDevicePointer();
471 
472  gdist.addPoints([cnt] __device__ (int i, int j, int k)
473  {
474  atomicAdd(cnt,1);
475 
476  return true;
477  },
478  [c] __device__ (InsertBlockT & data, int i, int j, int k)
479  {
480  data.template get<0>() = c + i + j;
481  data.template get<1>() = c + 1000 + i + j;
482  }
483  );
484 
485  gdist.template flush<smax_<0>,smax_<1>>(flush_type::FLUSH_ON_DEVICE);
486  gdist.template ghost_get<0,1>(RUN_ON_DEVICE);
487 
488  cmem.deviceToHost();
489 
490  int cnt_host = *(int *)cmem.getPointer();
491 
492  auto & v_cl = create_vcluster();
493 
494  v_cl.sum(cnt_host);
495  v_cl.execute();
496 
497  BOOST_REQUIRE_EQUAL(cnt_host,75*75*75);
498 }
499 
500 
501 BOOST_AUTO_TEST_CASE( sgrid_gpu_app_point_test )
502 {
503  size_t sz[3] = {75,75,75};
504  periodicity<3> bc = {PERIODIC,PERIODIC,PERIODIC};
505 
506  Ghost<3,long int> g(1);
507 
508  Box<3,float> domain({0.0,0.0,0.0},{1.0,1.0,1.0});
509 
510  sgrid_dist_id_gpu<3,float,aggregate<float,float,float,float>> gdist(sz,domain,g,bc);
511 
512  gdist.template setBackgroundValue<0>(666);
513  gdist.template setBackgroundValue<1>(666);
514  gdist.template setBackgroundValue<2>(666);
515  gdist.template setBackgroundValue<3>(666);
516 
518 
519  Box<3,size_t> box({1,1,1},{sz[0],sz[1],sz[2]});
520 
522 
523  float c = 5.0;
524 
525  typedef typename GetAddBlockType<decltype(gdist)>::type InsertBlockT;
526 
527  CudaMemory cmem;
528  cmem.allocate(sizeof(int));
529  CudaMemory cmem_out;
530  cmem_out.allocate(sizeof(int));
531 
532  *(int *)cmem.getPointer() = 0.0;
533  *(int *)cmem_out.getPointer() = 0.0;
534 
535  cmem.hostToDevice();
536  cmem_out.hostToDevice();
537 
538  int * cnt = (int *)cmem.getDevicePointer();
539  int * cnt_out = (int *)cmem_out.getDevicePointer();
540 
541  Box<3,size_t> bx({23,23,23},{70,70,70});
542 
543  gdist.addPoints(bx.getKP1(),bx.getKP2(),
544  [cnt,cnt_out,bx] __device__ (int i, int j, int k)
545  {
546  Point<3,int> p({i,j,k});
547 
548  if (bx.isInside(p))
549  {atomicAdd(cnt,1);}
550  else
551  {
552  atomicAdd(cnt_out,1);
553  }
554 
555  return true;
556  },
557  [c] __device__ (InsertBlockT & data, int i, int j, int k)
558  {
559  data.template get<0>() = c + i + j;
560  data.template get<1>() = c + 1000 + i + j;
561  }
562  );
563 
564  gdist.template flush<smax_<0>,smax_<1>>(flush_type::FLUSH_ON_DEVICE);
565  gdist.template ghost_get<0,1>(RUN_ON_DEVICE);
566 
567  cmem.deviceToHost();
568  cmem_out.deviceToHost();
569 
570  int cnt_host = *(int *)cmem.getPointer();
571  int cnt_host_out = *(int *)cmem_out.getPointer();
572 
573  auto & v_cl = create_vcluster();
574 
575  v_cl.sum(cnt_host_out);
576  v_cl.sum(cnt_host);
577  v_cl.execute();
578 
579  BOOST_REQUIRE_EQUAL(cnt_host_out,0);
580  BOOST_REQUIRE_EQUAL(cnt_host,bx.getVolumeKey());
581 }
582 
583 
584 BOOST_AUTO_TEST_CASE( sgrid_gpu_test_conv2_test )
585 {
586  size_t sz[2] = {164,164};
587  periodicity<2> bc = {PERIODIC,PERIODIC};
588 
589  Ghost<2,long int> g(1);
590 
591  Box<2,float> domain({0.0,0.0},{1.0,1.0});
592 
593  sgrid_dist_id_gpu<2,float,aggregate<float,float,float,float>> gdist(sz,domain,g,bc);
594 
595  gdist.template setBackgroundValue<0>(666);
596  gdist.template setBackgroundValue<1>(666);
597  gdist.template setBackgroundValue<2>(666);
598  gdist.template setBackgroundValue<3>(666);
599 
601 
602  Box<2,size_t> box({1,1},{sz[0],sz[1]});
603 
605 
606  float c = 5.0;
607 
608  typedef typename GetAddBlockType<decltype(gdist)>::type InsertBlockT;
609 
610  gdist.addPoints(box.getKP1(),box.getKP2(),
611  [] __device__ (int i, int j)
612  {
613  return true;
614  },
615  [c] __device__ (InsertBlockT & data, int i, int j)
616  {
617  data.template get<0>() = c + i + j;
618  data.template get<1>() = c + 1000 + i + j;
619  }
620  );
621 
622  gdist.template flush<smax_<0>,smax_<1>>(flush_type::FLUSH_ON_DEVICE);
623  gdist.template ghost_get<0,1>(RUN_ON_DEVICE);
624 
625 
626  // Now run the convolution
627 
628  typedef typename GetCpBlockType<decltype(gdist),0,1>::type CpBlockType;
629 
630  gdist.template conv2<0,1,2,3,1>({2,2},{(int)sz[0]-2,(int)sz[1]-2},[] __device__ (float & u_out, float & v_out, CpBlockType & u, CpBlockType & v,int i, int j){
631  u_out = u(i+1,j) - u(i-1,j) + u(i,j+1) - u(i,j-1);
632  v_out = v(i+1,j) - v(i-1,j) + v(i,j+1) - v(i,j-1);
633  });
634 
635  gdist.deviceToHost<0,1,2,3>();
636 
637  // Now we check that ghost is correct
638 
639  auto it3 = gdist.getSubDomainIterator({2,2},{(int)sz[0]-2,(int)sz[1]-2});
640 
641  bool match = true;
642 
643  while (it3.isNext())
644  {
645  auto p = it3.get();
646 
647  auto p_xp1 = p.move(0,1);
648  auto p_xm1 = p.move(0,-1);
649  auto p_yp1 = p.move(1,1);
650  auto p_ym1 = p.move(1,-1);
651 
652  float sub1 = gdist.template get<2>(p);
653  float sub2 = gdist.template get<3>(p);
654 
655  if (sub1 != 4.0 || sub2 != 4.0)
656  {
657  std::cout << sub1 << " " << sub2 << std::endl;
658  std::cout << gdist.template get<0>(p_xp1) << " " << gdist.template get<0>(p_xm1) << std::endl;
659  std::cout << gdist.template get<1>(p_xp1) << " " << gdist.template get<1>(p_xm1) << std::endl;
660  match = false;
661  break;
662  }
663 
664  ++it3;
665  }
666 
667  BOOST_REQUIRE_EQUAL(match,true);
668 }
669 
670 
671 BOOST_AUTO_TEST_CASE( sgrid_gpu_test_conv2_test_3d )
672 {
673  #ifdef CUDA_ON_CPU
674  size_t sz[3] = {20,20,20};
675  #else
676  size_t sz[3] = {60,60,60};
677  #endif
678  periodicity<3> bc = {PERIODIC,PERIODIC,PERIODIC};
679 
680  Ghost<3,long int> g(1);
681 
682  Box<3,float> domain({0.0,0.0,0.0},{1.0,1.0,1.0});
683 
684  sgrid_dist_id_gpu<3,float,aggregate<float,float,float,float>> gdist(sz,domain,g,bc);
685 
686  gdist.template setBackgroundValue<0>(666);
687  gdist.template setBackgroundValue<1>(666);
688  gdist.template setBackgroundValue<2>(666);
689  gdist.template setBackgroundValue<3>(666);
690 
692 
693  Box<3,size_t> box({1,1,1},{sz[0]-1,sz[1]-1,sz[2]-1});
694 
696 
697  float c = 5.0;
698 
699  typedef typename GetAddBlockType<decltype(gdist)>::type InsertBlockT;
700 
701  gdist.addPoints(box.getKP1(),box.getKP2(),
702  [] __device__ (int i, int j, int k)
703  {
704  return true;
705  },
706  [c] __device__ (InsertBlockT & data, int i, int j, int k)
707  {
708  data.template get<0>() = c + i + j + k;
709  data.template get<1>() = c + 1000 + i + j + k;
710  }
711  );
712 
713  gdist.template flush<smax_<0>,smax_<1>>(flush_type::FLUSH_ON_DEVICE);
714 
715  gdist.template ghost_get<0,1>(RUN_ON_DEVICE);
716 
717  for (int i = 0 ; i < 10 ; i++)
718  {
719  gdist.template ghost_get<0,1>(RUN_ON_DEVICE);
720  }
721 
722  // Now run the convolution
723 
724  typedef typename GetCpBlockType<decltype(gdist),0,1>::type CpBlockType;
725 
726  gdist.template conv2<0,1,2,3,1>({2,2,2},{(int)sz[0]-2,(int)sz[1]-2,(int)sz[2]-2},[] __device__ (float & u_out, float & v_out, CpBlockType & u, CpBlockType & v,int i, int j, int k){
727  u_out = u(i+1,j,k) - u(i-1,j,k) + u(i,j+1,k) - u(i,j-1,k) + u(i,j,k+1) - u(i,j,k-1);
728  v_out = v(i+1,j,k) - v(i-1,j,k) + v(i,j+1,k) - v(i,j-1,k) + v(i,j,k+1) - v(i,j,k-1);
729  });
730 
731  gdist.deviceToHost<0,1,2,3>();
732 
733  // Now we check that ghost is correct
734 
735  auto it3 = gdist.getSubDomainIterator({2,2,2},{(int)sz[0]-2,(int)sz[1]-2,(int)sz[2]-2});
736 
737  bool match = true;
738 
739  while (it3.isNext())
740  {
741  auto p = it3.get();
742 
743  auto p_xp1 = p.move(0,1);
744  auto p_xm1 = p.move(0,-1);
745  auto p_yp1 = p.move(1,1);
746  auto p_ym1 = p.move(1,-1);
747  auto p_zp1 = p.move(2,1);
748  auto p_zm1 = p.move(2,-1);
749 
750  float sub1 = gdist.template get<2>(p);
751  float sub2 = gdist.template get<3>(p);
752 
753  if (sub1 != 6.0 || sub2 != 6.0)
754  {
755  std::cout << sub1 << " " << sub2 << std::endl;
756  std::cout << gdist.template get<0>(p_xp1) << " " << gdist.template get<0>(p_xm1) << std::endl;
757  std::cout << gdist.template get<1>(p_xp1) << " " << gdist.template get<1>(p_xm1) << std::endl;
758  match = false;
759  break;
760  }
761 
762  ++it3;
763  }
764 
765  BOOST_REQUIRE_EQUAL(match,true);
766 }
767 
768 BOOST_AUTO_TEST_CASE( sgrid_gpu_test_conv_cross_block_test_3d )
769 {
770  #ifdef CUDA_ON_CPU
771  size_t sz[3] = {20,20,20};
772  #else
773  size_t sz[3] = {60,60,60};
774  #endif
775  periodicity<3> bc = {PERIODIC,PERIODIC,PERIODIC};
776 
777  Ghost<3,long int> g(1);
778 
779  Box<3,float> domain({0.0,0.0,0.0},{1.0,1.0,1.0});
780 
781  sgrid_dist_id_gpu<3,float,aggregate<float,float,float,float>> gdist(sz,domain,g,bc);
782 
783  gdist.template setBackgroundValue<0>(666);
784  gdist.template setBackgroundValue<1>(666);
785  gdist.template setBackgroundValue<2>(666);
786  gdist.template setBackgroundValue<3>(666);
787 
789 
790  Box<3,size_t> box({1,1,1},{sz[0]-1,sz[1]-1,sz[2]-1});
791 
793 
794  float c = 5.0;
795 
796  typedef typename GetAddBlockType<decltype(gdist)>::type InsertBlockT;
797 
798  gdist.addPoints(box.getKP1(),box.getKP2(),
799  [] __device__ (int i, int j, int k)
800  {
801  return true;
802  },
803  [c] __device__ (InsertBlockT & data, int i, int j, int k)
804  {
805  data.template get<0>() = c + i + j + k;
806  data.template get<1>() = c + 1000 + i + j + k;
807  }
808  );
809 
810  gdist.template flush<smax_<0>,smax_<1>>(flush_type::FLUSH_ON_DEVICE);
811 
812  gdist.template ghost_get<0,1>(RUN_ON_DEVICE);
813 
814  for (int i = 0 ; i < 10 ; i++)
815  {
816  gdist.template ghost_get<0,1>(RUN_ON_DEVICE);
817  }
818 
819  // Now run the convolution
820 
821  typedef typename GetCpBlockType<decltype(gdist),0,1>::type CpBlockType;
822 
823  gdist.template conv_cross_b<0,1,1>({2,2,2},{(int)sz[0]-2,(int)sz[1]-2,(int)sz[2]-2},[] __device__ (CpBlockType & u, auto & block, int offset,int i, int j, int k){
824  return u(i+1,j,k) - u(i-1,j,k) + u(i,j+1,k) - u(i,j-1,k) + u(i,j,k+1) - u(i,j,k-1) + block.template get<0>()[offset];
825  });
826 
827  gdist.deviceToHost<0,1,2,3>();
828 
829  // Now we check that ghost is correct
830 
831  auto it3 = gdist.getSubDomainIterator({2,2,2},{(int)sz[0]-2,(int)sz[1]-2,(int)sz[2]-2});
832 
833  bool match = true;
834 
835  while (it3.isNext())
836  {
837  auto p = it3.get();
838 
839  auto p_xp1 = p.move(0,1);
840  auto p_xm1 = p.move(0,-1);
841  auto p_yp1 = p.move(1,1);
842  auto p_ym1 = p.move(1,-1);
843  auto p_zp1 = p.move(2,1);
844  auto p_zm1 = p.move(2,-1);
845 
846  float sub1 = gdist.template get<1>(p);
847 
848  if (sub1 != 6.0 + gdist.template get<0>(p))
849  {
850  std::cout << sub1 << std::endl;
851  std::cout << gdist.template get<0>(p_xp1) << " " << gdist.template get<0>(p_xm1) << std::endl;
852  std::cout << gdist.template get<1>(p_xp1) << " " << gdist.template get<1>(p_xm1) << std::endl;
853  match = false;
854  break;
855  }
856 
857  ++it3;
858  }
859 
860  BOOST_REQUIRE_EQUAL(match,true);
861 }
862 
863 BOOST_AUTO_TEST_CASE( sgrid_gpu_test_conv2_b_test_3d )
864 {
865  #ifdef CUDA_ON_CPU
866  size_t sz[3] = {20,20,20};
867  #else
868  size_t sz[3] = {60,60,60};
869  #endif
870  periodicity<3> bc = {PERIODIC,PERIODIC,PERIODIC};
871 
872  Ghost<3,long int> g(1);
873 
874  Box<3,float> domain({0.0,0.0,0.0},{1.0,1.0,1.0});
875 
876  sgrid_dist_id_gpu<3,float,aggregate<float,float,float,float>> gdist(sz,domain,g,bc);
877 
878  gdist.template setBackgroundValue<0>(666);
879  gdist.template setBackgroundValue<1>(666);
880  gdist.template setBackgroundValue<2>(666);
881  gdist.template setBackgroundValue<3>(666);
882 
884 
885  Box<3,size_t> box({1,1,1},{sz[0]-1,sz[1]-1,sz[2]-1});
886 
888 
889  float c = 5.0;
890 
891  typedef typename GetAddBlockType<decltype(gdist)>::type InsertBlockT;
892 
893  gdist.addPoints(box.getKP1(),box.getKP2(),
894  [] __device__ (int i, int j, int k)
895  {
896  return true;
897  },
898  [c] __device__ (InsertBlockT & data, int i, int j, int k)
899  {
900  data.template get<0>() = c + i + j + k;
901  data.template get<1>() = c + 1000 + i + j + k;
902  }
903  );
904 
905  gdist.template flush<smax_<0>,smax_<1>>(flush_type::FLUSH_ON_DEVICE);
906 
907  gdist.template ghost_get<0,1>(RUN_ON_DEVICE);
908 
909  for (int i = 0 ; i < 10 ; i++)
910  {
911  gdist.template ghost_get<0,1>(RUN_ON_DEVICE);
912  }
913 
914  // Now run the convolution
915 
916  typedef typename GetCpBlockType<decltype(gdist),0,1>::type CpBlockType;
917 
918  gdist.template conv2_b<0,1,2,3,1>({2,2,2},{(int)sz[0]-2,(int)sz[1]-2,(int)sz[2]-2},[] __device__ (float & u_out, float & v_out, CpBlockType & u, CpBlockType & v, auto & block, int offset,int i, int j, int k){
919  u_out = u(i+1,j,k) - u(i-1,j,k) + u(i,j+1,k) - u(i,j-1,k) + u(i,j,k+1) - u(i,j,k-1) + block.template get<0>()[offset];
920  v_out = v(i+1,j,k) - v(i-1,j,k) + v(i,j+1,k) - v(i,j-1,k) + v(i,j,k+1) - v(i,j,k-1) + block.template get<1>()[offset];
921  });
922 
923  gdist.deviceToHost<0,1,2,3>();
924 
925  // Now we check that ghost is correct
926 
927  auto it3 = gdist.getSubDomainIterator({2,2,2},{(int)sz[0]-2,(int)sz[1]-2,(int)sz[2]-2});
928 
929  bool match = true;
930 
931  while (it3.isNext())
932  {
933  auto p = it3.get();
934 
935  auto p_xp1 = p.move(0,1);
936  auto p_xm1 = p.move(0,-1);
937  auto p_yp1 = p.move(1,1);
938  auto p_ym1 = p.move(1,-1);
939  auto p_zp1 = p.move(2,1);
940  auto p_zm1 = p.move(2,-1);
941 
942  float sub1 = gdist.template get<2>(p) ;
943  float sub2 = gdist.template get<3>(p);
944 
945  if (sub1 != 6.0 + gdist.template get<0>(p) || sub2 != 6.0 + gdist.template get<1>(p))
946  {
947  std::cout << sub1 << " " << sub2 << std::endl;
948  std::cout << gdist.template get<0>(p_xp1) << " " << gdist.template get<0>(p_xm1) << std::endl;
949  std::cout << gdist.template get<1>(p_xp1) << " " << gdist.template get<1>(p_xm1) << std::endl;
950  match = false;
951  break;
952  }
953 
954  ++it3;
955  }
956 
957  BOOST_REQUIRE_EQUAL(match,true);
958 }
959 
960 BOOST_AUTO_TEST_CASE( sgrid_gpu_test_ghost_point_remove )
961 {
962  size_t sz[3] = {60,60,60};
963  periodicity<3> bc = {PERIODIC,PERIODIC,PERIODIC};
964 
965  Ghost<3,long int> g(1);
966 
967  Box<3,float> domain({0.0,0.0,0.0},{1.0,1.0,1.0});
968 
969  sgrid_dist_id_gpu<3,float,aggregate<float,float,float,float>> gdist(sz,domain,g,bc);
970 
971  gdist.template setBackgroundValue<0>(666);
972  gdist.template setBackgroundValue<1>(666);
973  gdist.template setBackgroundValue<2>(666);
974  gdist.template setBackgroundValue<3>(666);
975 
977 
978  Box<3,size_t> box({1,1,1},{sz[0]-1,sz[1]-1,sz[2]-1});
979 
981 
982  float c = 5.0;
983 
984  typedef typename GetAddBlockType<decltype(gdist)>::type InsertBlockT;
985 
986  gdist.addPoints(box.getKP1(),box.getKP2(),
987  [] __device__ (int i, int j, int k)
988  {
989  return true;
990  },
991  [c] __device__ (InsertBlockT & data, int i, int j, int k)
992  {
993  data.template get<0>() = c + i + j + k;
994  data.template get<1>() = c + 1000 + i + j + k;
995  }
996  );
997 
998  gdist.template flush<smax_<0>,smax_<1>>(flush_type::FLUSH_ON_DEVICE);
999 
1000  gdist.template ghost_get<0,1>(RUN_ON_DEVICE);
1001 
1002  // Remove the right side of the points
1003  Box<3,size_t> bxR({59,0,0},{59,59,59});
1004  gdist.removePoints(bxR);
1005 
1006  // Remove the right side of the points
1007  Box<3,size_t> bxT({0,0,59},{59,59,59});
1008  gdist.removePoints(bxT);
1009 
1010  // Remove the right side of the points
1011  Box<3,size_t> bxD({0,59,0},{59,59,59});
1012  gdist.removePoints(bxD);
1013 
1014  gdist.template ghost_get<0,1>(RUN_ON_DEVICE);
1015 
1016  for (int i = 0 ; i < 10 ; i++)
1017  {
1018  gdist.template ghost_get<0,1>(RUN_ON_DEVICE);
1019  }
1020 
1021  // Now run the convolution
1022 
1023  typedef typename GetCpBlockType<decltype(gdist),0,1>::type CpBlockType;
1024 
1025  gdist.template conv2<0,1,2,3,1>({2,2,2},{(int)sz[0]-3,(int)sz[1]-3,(int)sz[2]-3},[] __device__ (float & u_out, float & v_out, CpBlockType & u, CpBlockType & v,int i, int j, int k){
1026  u_out = u(i+1,j,k) - u(i-1,j,k) + u(i,j+1,k) - u(i,j-1,k) + u(i,j,k+1) - u(i,j,k-1);
1027  v_out = v(i+1,j,k) - v(i-1,j,k) + v(i,j+1,k) - v(i,j-1,k) + v(i,j,k+1) - v(i,j,k-1);
1028  });
1029 
1030  gdist.deviceToHost<0,1,2,3>();
1031 
1032  // Now we check that ghost is correct
1033 
1034  auto it3 = gdist.getSubDomainIterator({2,2,2},{(int)sz[0]-3,(int)sz[1]-3,(int)sz[2]-3});
1035 
1036  bool match = true;
1037 
1038  while (it3.isNext())
1039  {
1040  auto p = it3.get();
1041 
1042  auto p_xp1 = p.move(0,1);
1043  auto p_xm1 = p.move(0,-1);
1044  auto p_yp1 = p.move(1,1);
1045  auto p_ym1 = p.move(1,-1);
1046  auto p_zp1 = p.move(2,1);
1047  auto p_zm1 = p.move(2,-1);
1048 
1049  float sub1 = gdist.template get<2>(p);
1050  float sub2 = gdist.template get<3>(p);
1051 
1052  if (sub1 != 6.0 || sub2 != 6.0)
1053  {
1054  std::cout << sub1 << " " << sub2 << std::endl;
1055  std::cout << gdist.template get<0>(p_xp1) << " " << gdist.template get<0>(p_xm1) << std::endl;
1056  std::cout << gdist.template get<1>(p_xp1) << " " << gdist.template get<1>(p_xm1) << std::endl;
1057  match = false;
1058  break;
1059  }
1060 
1061  ++it3;
1062  }
1063 
1064  BOOST_REQUIRE_EQUAL(match,true);
1065 
1066  gdist.template deviceToHost<0,1,2,3>();
1067 
1068  auto it4 = gdist.getDomainGhostIterator();
1069  Box<3,long int> bin({0,0,0},{59,59,59});
1070 
1071  match = true;
1072 
1073  while (it4.isNext())
1074  {
1075  auto p = it4.get();
1076 
1077  // We have to check we have no point in the ghost area
1078  auto gkey = it4.getGKey(p);
1079 
1080  if (bin.isInside(gkey.toPoint()) == false)
1081  {match = false;}
1082 
1083  ++it4;
1084  }
1085 
1086  BOOST_REQUIRE_EQUAL(match,true);
1087 }
1088 
1089 BOOST_AUTO_TEST_CASE( sgrid_gpu_test_skip_labelling )
1090 {
1091  size_t sz[3] = {60,60,60};
1092  periodicity<3> bc = {PERIODIC,PERIODIC,PERIODIC};
1093 
1094  Ghost<3,long int> g(1);
1095 
1096  Box<3,float> domain({0.0,0.0,0.0},{1.0,1.0,1.0});
1097 
1098  sgrid_dist_id_gpu<3,float,aggregate<float,float,float,float>> gdist(sz,domain,g,bc);
1099 
1100  gdist.template setBackgroundValue<0>(666);
1101  gdist.template setBackgroundValue<1>(666);
1102  gdist.template setBackgroundValue<2>(666);
1103  gdist.template setBackgroundValue<3>(666);
1104 
1106 
1107  Box<3,size_t> box({1,1,1},{sz[0]-1,sz[1]-1,sz[2]-1});
1108 
1110 
1111  float c = 5.0;
1112 
1113  typedef typename GetAddBlockType<decltype(gdist)>::type InsertBlockT;
1114 
1115  gdist.addPoints(box.getKP1(),box.getKP2(),
1116  [] __device__ (int i, int j, int k)
1117  {
1118  return true;
1119  },
1120  [c] __device__ (InsertBlockT & data, int i, int j, int k)
1121  {
1122  data.template get<0>() = c + i + j + k;
1123  data.template get<1>() = c + 1000 + i + j + k;
1124  }
1125  );
1126 
1127  gdist.template flush<smax_<0>,smax_<1>>(flush_type::FLUSH_ON_DEVICE);
1128 
1129  gdist.template ghost_get<0,1>(RUN_ON_DEVICE);
1130 
1131  // Now run the convolution
1132 
1133  typedef typename GetCpBlockType<decltype(gdist),0,1>::type CpBlockType;
1134 
1135  gdist.template conv2<0,1,0,1,1>({0,0,0},{(int)sz[0]-1,(int)sz[1]-1,(int)sz[2]-1},[] __device__ (float & u_out, float & v_out, CpBlockType & u, CpBlockType & v,int i, int j, int k){
1136  u_out = 1*u(i,j,k);
1137  v_out = 1*v(i,j,k);
1138  });
1139 
1140  gdist.template ghost_get<0,1>(RUN_ON_DEVICE | SKIP_LABELLING);
1141 
1142  gdist.template conv2<0,1,0,1,1>({0,0,0},{(int)sz[0]-1,(int)sz[1]-1,(int)sz[2]-1},[] __device__ (float & u_out, float & v_out, CpBlockType & u, CpBlockType & v,int i, int j, int k){
1143  u_out = 5*u(i,j,k);
1144  v_out = 5*v(i,j,k);
1145  });
1146 
1147  gdist.template ghost_get<0,1>(RUN_ON_DEVICE | SKIP_LABELLING);
1148 
1149  gdist.template conv2<0,1,0,1,1>({0,0,0},{(int)sz[0]-1,(int)sz[1]-1,(int)sz[2]-1},[] __device__ (float & u_out, float & v_out, CpBlockType & u, CpBlockType & v,int i, int j, int k){
1150  u_out = 2*u(i,j,k);
1151  v_out = 2*v(i,j,k);
1152  });
1153 
1154  gdist.template ghost_get<0,1>(RUN_ON_DEVICE | SKIP_LABELLING);
1155 
1156  gdist.template conv2<0,1,2,3,1>({2,2,2},{(int)sz[0]-3,(int)sz[1]-3,(int)sz[2]-3},[] __device__ (float & u_out, float & v_out, CpBlockType & u, CpBlockType & v,int i, int j, int k){
1157  u_out = u(i+1,j,k) - u(i-1,j,k) + u(i,j+1,k) - u(i,j-1,k) + u(i,j,k+1) - u(i,j,k-1);
1158  v_out = v(i+1,j,k) - v(i-1,j,k) + v(i,j+1,k) - v(i,j-1,k) + v(i,j,k+1) - v(i,j,k-1);
1159  });
1160 
1161 
1162  gdist.deviceToHost<0,1,2,3>();
1163 
1164  // Now we check that ghost is correct
1165 
1166  auto it3 = gdist.getSubDomainIterator({2,2,2},{(int)sz[0]-3,(int)sz[1]-3,(int)sz[2]-3});
1167 
1168  bool match = true;
1169 
1170  while (it3.isNext())
1171  {
1172  auto p = it3.get();
1173 
1174  auto p_xp1 = p.move(0,1);
1175  auto p_xm1 = p.move(0,-1);
1176  auto p_yp1 = p.move(1,1);
1177  auto p_ym1 = p.move(1,-1);
1178  auto p_zp1 = p.move(2,1);
1179  auto p_zm1 = p.move(2,-1);
1180 
1181  float sub1 = gdist.template get<2>(p);
1182  float sub2 = gdist.template get<3>(p);
1183 
1184  if (sub1 != 6.0*10.0 || sub2 != 6.0*10.0)
1185  {
1186  std::cout << sub1 << " " << sub2 << std::endl;
1187  std::cout << gdist.template get<0>(p_xp1) << " " << gdist.template get<0>(p_xm1) << std::endl;
1188  std::cout << gdist.template get<1>(p_xp1) << " " << gdist.template get<1>(p_xm1) << std::endl;
1189  match = false;
1190  break;
1191  }
1192 
1193  ++it3;
1194  }
1195 
1196  BOOST_REQUIRE_EQUAL(match,true);
1197 }
1198 
1199 BOOST_AUTO_TEST_CASE( sgrid_gpu_test_conv_background )
1200 {
1201  size_t sz[3] = {60,60,60};
1202  periodicity<3> bc = {PERIODIC,PERIODIC,PERIODIC};
1203 
1204  Ghost<3,long int> g(1);
1205 
1206  Box<3,float> domain({0.0,0.0,0.0},{1.0,1.0,1.0});
1207 
1208  sgrid_dist_id_gpu<3,float,aggregate<float,float,float,float>> gdist(sz,domain,g,bc);
1209 
1210  gdist.template setBackgroundValue<0>(666);
1211  gdist.template setBackgroundValue<1>(666);
1212  gdist.template setBackgroundValue<2>(666);
1213  gdist.template setBackgroundValue<3>(666);
1214 
1216 
1217  Box<3,size_t> box({1,1,1},{sz[0]-1,sz[1]-1,sz[2]-1});
1218 
1220 
1221  float c = 5.0;
1222 
1223  typedef typename GetAddBlockType<decltype(gdist)>::type InsertBlockT;
1224 
1225  gdist.addPoints(box.getKP1(),box.getKP2(),
1226  [] __device__ (int i, int j, int k)
1227  {
1228  return (i == 30 && j == 30 && k == 30);
1229  },
1230  [c] __device__ (InsertBlockT & data, int i, int j, int k)
1231  {
1232  data.template get<0>() = 0;
1233  data.template get<1>() = 0;
1234  }
1235  );
1236 
1237  gdist.template flush<smax_<0>,smax_<1>>(flush_type::FLUSH_ON_DEVICE);
1238 
1239  gdist.template ghost_get<0,1>(RUN_ON_DEVICE);
1240 
1241  // Now run the convolution
1242 
1243  typedef typename GetCpBlockType<decltype(gdist),0,1>::type CpBlockType;
1244 
1245  gdist.template conv2<0,1,2,3,1>({0,0,0},{(int)sz[0]-1,(int)sz[1]-1,(int)sz[2]-1},[] __device__ (float & u_out, float & v_out, CpBlockType & u, CpBlockType & v,int i, int j, int k){
1246  u_out = u(i+1,j,k) + u(i,j+1,k) + u(i,j,k+1);
1247  v_out = v(i+1,j,k) + v(i,j+1,k) + v(i,j,k+1);
1248  });
1249 
1250  gdist.deviceToHost<0,1,2,3>();
1251 
1252  // Now we check that ghost is correct
1253 
1254  auto it3 = gdist.getDomainIterator();
1255 
1256  bool match = true;
1257 
1258  int count = 0;
1259 
1260  while (it3.isNext())
1261  {
1262  auto p = it3.get();
1263 
1264  float sub1 = gdist.template get<2>(p);
1265  float sub2 = gdist.template get<3>(p);
1266 
1267  if (sub1 != 3*666.0 || sub2 != 3*666.0)
1268  {
1269  std::cout << sub1 << " " << sub2 << std::endl;
1270  match = false;
1271  break;
1272  }
1273 
1274  count++;
1275 
1276  ++it3;
1277  }
1278 
1279  BOOST_REQUIRE(count == 0 || count == 1);
1280  BOOST_REQUIRE_EQUAL(match,true);
1281 }
1282 
1283 BOOST_AUTO_TEST_CASE( grid_dense_to_sparse_conversion )
1284 {
1285  Box<3,float> domain({0.0,0.0,0.0},{1.0,1.0,1.0});
1286 
1287  // grid size
1288  size_t sz[3];
1289  sz[0] = 32;
1290  sz[1] = 32;
1291  sz[2] = 32;
1292 
1293  // Ghost
1294  Ghost<3,long int> g(1);
1295 
1296  periodicity<3> pr = {PERIODIC,PERIODIC,PERIODIC};
1297 
1298  // Distributed grid with id decomposition
1299  grid_dist_id<3, float, aggregate<float,float>> g_dist(sz,domain,g,pr);
1300 
1301  auto it = g_dist.getDomainIterator();
1302 
1303  while (it.isNext())
1304  {
1305  auto p = it.get();
1306  auto gkey = it.getGKey(p);
1307 
1308  g_dist.template getProp<0>(p) = gkey.get(0) + gkey.get(1) + gkey.get(2);
1309  g_dist.template getProp<1>(p) = 3.0*gkey.get(0) + gkey.get(1) + gkey.get(2);
1310 
1311  ++it;
1312  }
1313 
1314  sgrid_dist_id_gpu<3,float,aggregate<float,float>> sgdist(g_dist.getDecomposition(),sz,g);
1315 
1316  while (it.isNext())
1317  {
1318  auto p = it.get();
1319  auto gkey = it.getGKey(p);
1320 
1321  sgdist.template insertFlush<0>(p) = g_dist.template get<0>(p);
1322 
1323  ++it;
1324  }
1325 
1326 
1327  bool check = true;
1328 
1329  while (it.isNext())
1330  {
1331  auto p = it.get();
1332  auto gkey = it.getGKey(p);
1333 
1334  check &= sgdist.template getProp<0>(p) == g_dist.template get<0>(p);
1335 
1336  ++it;
1337  }
1338 
1339  BOOST_REQUIRE_EQUAL(check,true);
1340 }
1341 
1342 BOOST_AUTO_TEST_SUITE_END()
virtual bool allocate(size_t sz)
allocate memory
Definition: CudaMemory.cu:38
virtual void hostToDevice()
Move memory from host to device.
Definition: CudaMemory.cu:508
virtual void * getPointer()
get a readable pointer with the data
Definition: CudaMemory.cu:352
This class implement the point shape in an N-dimensional space.
Definition: Point.hpp:27
Definition: Ghost.hpp:39
virtual void * getDevicePointer()
get a readable pointer with the data
Definition: CudaMemory.cu:497
This is a distributed grid.
KeyT const ValueT ValueT OffsetIteratorT OffsetIteratorT int
[in] The number of segments that comprise the sorting data
This class represent an N-dimensional box.
Definition: Box.hpp:60
get the type of the block
virtual void deviceToHost()
Move memory from device to host.
Definition: CudaMemory.cu:367
get the type of the insertBlock
Boundary conditions.
Definition: common.hpp:21