OpenFPM_pdata  4.1.0
Project that contain the implementation of distributed structures
map_vector_sparse_cuda_ker_unit_tests.cu
1 /*
2  * map_vector_sparse_cuda_ker_unit_tests.cuh
3  *
4  * Created on: Jan 24, 2019
5  * Author: i-bird
6  */
7 
8 #ifndef MAP_VECTOR_SPARSE_CUDA_KER_UNIT_TESTS_CUH_
9 #define MAP_VECTOR_SPARSE_CUDA_KER_UNIT_TESTS_CUH_
10 
11 #define BOOST_TEST_DYN_LINK
12 #include <boost/test/unit_test.hpp>
13 #include "Vector/map_vector_sparse.hpp"
14 #include "Vector/map_vector.hpp"
15 #include <map>
16 
17 template<typename vd_type>
18 __global__ void test_insert_sparse(vd_type vd_insert)
19 {
20  vd_insert.init();
21 
22  int p = blockIdx.x*blockDim.x + threadIdx.x;
23 
24  p *= 2;
25 
26  auto ie = vd_insert.insert(10000 - p);
27  ie.template get<0>() = p + 100;
28  ie.template get<1>() = p + 10100;
29  ie.template get<2>() = p + 20100;
30  vd_insert.flush_block_insert();
31 }
32 
33 //template<unsigned int blockLength, typename vd_type>
34 //__global__ void test_insert_sparse_block(vd_type vd_insert)
35 //{
36 // vd_insert.init();
37 //
38 // int p = blockIdx.x*blockDim.x + threadIdx.x;
39 //
40 // p *= 2;
41 //
42 // auto ie = vd_insert.insert(10000 - p);
43 // ie.template get<0>() = p + 100;
44 // for (unsigned int i = 0 ; i < blockLength ; i++)
45 // {
46 // ie.template get<1>()[i] = p + 10100 + i;
47 // }
48 //
49 // vd_insert.flush_block_insert();
50 //}
51 
52 template<typename vd_type>
53 __global__ void test_remove_sparse(vd_type vd_insert)
54 {
55  vd_insert.init();
56 
57  int p = blockIdx.x*blockDim.x + threadIdx.x;
58 
59  p *= 2;
60 
61  vd_insert.remove(10000 - p);
62  vd_insert.flush_block_remove();
63 }
64 
65 template<typename vd_type>
66 __global__ void test_insert_sparse2(vd_type vd_insert)
67 {
68  vd_insert.init();
69 
70  int p = blockIdx.x*blockDim.x + threadIdx.x;
71 
72  p *= 2;
73 
74  auto ie = vd_insert.insert(9000 - p);
75  ie.template get<0>() = p + 3000;
76  ie.template get<1>() = p + 13000;
77  ie.template get<2>() = p + 23000;
78 
79  vd_insert.flush_block_insert();
80 }
81 
82 template<typename vd_type>
83 __global__ void test_remove_sparse2(vd_type vd_insert)
84 {
85  vd_insert.init();
86 
87  int p = blockIdx.x*blockDim.x + threadIdx.x;
88 
89  p *= 2;
90 
91  vd_insert.remove(9000 - p);
92 
93  vd_insert.flush_block_remove();
94 }
95 
96 template<typename vd_type>
97 __global__ void test_insert_sparse2_inc(vd_type vd_insert)
98 {
99  vd_insert.init_ins_inc();
100 
101  int p = blockIdx.x*blockDim.x + threadIdx.x;
102 
103  p *= 2;
104 
105  auto ie = vd_insert.insert(9000 - p);
106  ie.template get<0>() = p + 3000;
107  ie.template get<1>() = p + 13000;
108  ie.template get<2>() = p + 23000;
109 
110  vd_insert.flush_block_insert();
111 }
112 
113 template<typename vd_type>
114 __global__ void test_remove_sparse2_inc(vd_type vd_insert)
115 {
116  vd_insert.init_rem_inc();
117 
118  int p = blockIdx.x*blockDim.x + threadIdx.x;
119 
120  p *= 2;
121 
122  vd_insert.remove(9000 - p);
123 
124  vd_insert.flush_block_remove();
125 }
126 
127 template<typename vd_type>
128 __global__ void test_insert_sparse3(vd_type vd_insert)
129 {
130  vd_insert.init();
131 
132  int p = blockIdx.x*blockDim.x + threadIdx.x;
133 
134  p *= 2;
135 
136  auto ie = vd_insert.insert(p);
137  ie.template get<0>() = 5;
138  ie.template get<1>() = 1;
139  ie.template get<2>() = 1;
140 
141  vd_insert.flush_block_insert();
142 }
143 
144 template<typename vd_type>
145 __global__ void test_remove_sparse3(vd_type vd_insert)
146 {
147  vd_insert.init();
148 
149  int p = blockIdx.x*blockDim.x + threadIdx.x;
150 
151  p *= 2;
152 
153  vd_insert.remove(p);
154 
155  vd_insert.flush_block_remove();
156 }
157 
158 template<typename vd_sparse_type, typename vector_out_type>
159 __global__ void test_sparse_get_test(vd_sparse_type vd_test, vector_out_type output)
160 {
161  int p = blockIdx.x*blockDim.x + threadIdx.x;
162  int i = blockIdx.x*blockDim.x + threadIdx.x;
163 
164  p *= 2;
165 
166  int v;
167 
168  output.template get<0>(i) = vd_test.template get<0>(10000 - p,v);
169  output.template get<1>(i) = vd_test.template get_ele<1>(v);
170  output.template get<2>(i) = vd_test.template get_ele<2>(v);
171 }
172 
173 //template<unsigned int blockLength, typename vd_sparse_type, typename vector_out_type>
174 //__global__ void test_sparse_get_test_block(vd_sparse_type vd_test, vector_out_type output)
175 //{
176 // int p = blockIdx.x*blockDim.x + threadIdx.x;
177 // int i = blockIdx.x*blockDim.x + threadIdx.x;
178 //
179 // p *= 2;
180 //
181 // int v;
182 //
183 // output.template get<0>(i) = vd_test.template get<0>(10000 - p,v);
184 // for (int j=0; j<blockLength; ++j)
185 // {
186 // output.template get<1>(i)[j] = vd_test.template get_ele<1>(v)[j];
187 // }
188 //}
189 
190 BOOST_AUTO_TEST_SUITE( vector_cuda_sparse )
191 
192 BOOST_AUTO_TEST_CASE( vector_sparse_cuda_gpu )
193 {
194  openfpm::vector_sparse_gpu<aggregate<size_t,float,double>> vs;
195 
196  vs.template setBackground<0>(17);
197 
198  vs.template setBackground<1>(18);
199 
200  vs.template setBackground<2>(19);
201 
202  vs.setGPUInsertBuffer(10,1024);
203 
204  // we launch a kernel to insert data
205  CUDA_LAUNCH_DIM3(test_insert_sparse,10,100,vs.toKernel());
206 
207  mgpu::ofp_context_t ctx;
208  vs.flush<sadd_<0>,smin_<1>,smax_<2> >(ctx,flush_type::FLUSH_ON_DEVICE);
209 
210  vs.setGPUInsertBuffer(10,1024);
211  CUDA_LAUNCH_DIM3(test_insert_sparse2,10,100,vs.toKernel());
212 
213  vs.flush<sadd_<0>,smin_<1>,smax_<2> >(ctx,flush_type::FLUSH_ON_DEVICE);
214 
215  vs.setGPUInsertBuffer(4000,512);
216  CUDA_LAUNCH_DIM3(test_insert_sparse3,4000,256,vs.toKernel());
217 
218  vs.flush<sadd_<0>,smin_<1>,smax_<2> >(ctx,flush_type::FLUSH_ON_DEVICE);
219 
220  openfpm::vector_gpu<aggregate<size_t,float,double>> output;
221 
222  output.resize(1500);
223 
224  CUDA_LAUNCH_DIM3(test_sparse_get_test,10,150,vs.toKernel(),output.toKernel());
225 
226  output.template deviceToHost<0,1,2>();
227  vs.template deviceToHost<0,1,2>();
228 
229  bool match = true;
230  for (size_t i = 0 ; i < output.size() ; i++)
231  {
232  match &= output.template get<0>(i) == vs.template get<0>(10000 - 2*i);
233  match &= output.template get<1>(i) == vs.template get<1>(10000 - 2*i);
234  match &= output.template get<2>(i) == vs.template get<2>(10000 - 2*i);
235  }
236 
237  BOOST_REQUIRE_EQUAL(match,true);
238 
239  vs.clear();
240 
241  CUDA_LAUNCH_DIM3(test_sparse_get_test,10,150,vs.toKernel(),output.toKernel());
242 
243  output.template deviceToHost<0,1,2>();
244  vs.template deviceToHost<0,1,2>();
245 
246  match = true;
247  for (size_t i = 0 ; i < output.size() ; i++)
248  {
249  match &= output.template get<0>(i) == 17;
250  match &= output.template get<1>(i) == 18;
251  match &= output.template get<2>(i) == 19;
252 
253 
254  if (match == false){break;}
255  }
256 
257  BOOST_REQUIRE_EQUAL(match,true);
258 }
259 
260 
261 BOOST_AUTO_TEST_CASE( vector_sparse_cuda_gpu_incremental_add )
262 {
263  openfpm::vector_sparse_gpu<aggregate<size_t,float,double>> vs;
264 
265  vs.template setBackground<0>(17);
266 
267  vs.template setBackground<1>(18);
268 
269  vs.template setBackground<2>(19);
270 
271  vs.setGPUInsertBuffer(10,1024);
272 
273  auto aaa = vs.toKernel();
274 // auto nindex = aaa.private_get_vct_nadd_index();
275 
276  // we launch a kernel to insert data
277  CUDA_LAUNCH_DIM3(test_insert_sparse,10,100,aaa);
278 
279 /* for (int i = 0 ; i < nindex.size() ; i++)
280  {
281  std::cout << aaa.template get<0>(i) << std::endl;
282  }*/
283 
284  CUDA_LAUNCH_DIM3(test_insert_sparse2_inc,10,100,vs.toKernel());
285  CUDA_LAUNCH_DIM3(test_insert_sparse2_inc,10,100,vs.toKernel());
286  CUDA_LAUNCH_DIM3(test_insert_sparse2_inc,10,100,vs.toKernel());
287 
288  mgpu::ofp_context_t ctx;
289 
290  vs.flush<sadd_<0>,smin_<1>,smax_<2> >(ctx,flush_type::FLUSH_ON_DEVICE);
291 
292  vs.template deviceToHost<0,1,2>();
293 
294  BOOST_REQUIRE_EQUAL(vs.size(),1500);
295 
296  bool match = true;
297  for (size_t i = 500 ; i < 1000 ; i++)
298  {
299  match &= vs.template get<0>(9000 - 2*i) == 3*(2*i + 3000);
300  match &= vs.template get<1>(9000 - 2*i) == 2*i + 13000;
301  match &= vs.template get<2>(9000 - 2*i) == 2*i + 23000;
302 
303  if (match == false)
304  {
305  std::cout << i << " " << vs.template get<0>(9000 - 2*i) << "!=" << 3*(2*i + 3000) << " " << vs.template get<1>(9000 - 2*i) << "!=" << 2*i + 13000 << " " << vs.template get<2>(9000 - 2*i) << "!=" << 2*i + 23000 << std::endl;
306  break;
307  }
308  }
309 
310  for (size_t i = 0 ; i < 500 ; i++)
311  {
312  match &= vs.template get<0>(9000 - 2*i) == 3*(2*i + 3000) + 2*i + 1100;
313  match &= vs.template get<1>(9000 - 2*i) == 2*i + 11100;
314  match &= vs.template get<2>(9000 - 2*i) == 2*i + 23000;
315 
316  if (match == false)
317  {
318  std::cout << i << " " << vs.template get<0>(9000 - 2*i) << "!=" << 3*(2*i + 3000) << " " << vs.template get<1>(9000 - 2*i) << "!=" << 2*i + 13000 << " " << vs.template get<2>(9000 - 2*i) << "!=" << 2*i + 23000 << std::endl;
319  break;
320  }
321  }
322 
323  for (size_t i = 0 ; i < 500 ; i++)
324  {
325  match &= vs.template get<0>(10000 - 2*i) == 2*i + 100;
326  match &= vs.template get<1>(10000 - 2*i) == 2*i + 10100;
327  match &= vs.template get<2>(10000 - 2*i) == 2*i + 20100;
328 
329  if (match == false)
330  {
331  std::cout << i << " " << vs.template get<0>(9000 - 2*i) << "!=" << 3*(2*i + 3000) << " " << vs.template get<1>(9000 - 2*i) << "!=" << 2*i + 13000 << " " << vs.template get<2>(9000 - 2*i) << "!=" << 2*i + 23000 << std::endl;
332  break;
333  }
334  }
335 
336  BOOST_REQUIRE_EQUAL(match,true);
337 }
338 
339 BOOST_AUTO_TEST_CASE( vector_sparse_cuda_gpu_get )
340 {
341  openfpm::vector_sparse_gpu<aggregate<size_t,float,double>> vs;
342 
343  vs.template setBackground<0>(0);
344 
345  vs.template setBackground<1>(0);
346 
347  vs.template setBackground<2>(0);
348 
349  vs.setGPUInsertBuffer(10,1024);
350 
351  // we launch a kernel to insert data
352  CUDA_LAUNCH_DIM3(test_insert_sparse,10,100,vs.toKernel());
353 
354 
355  mgpu::ofp_context_t ctx;
356  vs.flush<sadd_<0>,smin_<1>,smax_<2> >(ctx,flush_type::FLUSH_ON_DEVICE);
357 
358  vs.template deviceToHost<0,1,2>();
359 
360  bool match = true;
361  for (size_t i = 0 ; i < 1000 ; i++)
362  {
363  match &= vs.template get<0>(10000 - 2*i) == 2*i + 100;
364  match &= vs.template get<1>(10000 - 2*i) == 2*i + 10100;
365  match &= vs.template get<2>(10000 - 2*i) == 2*i + 20100;
366  }
367 
368  BOOST_REQUIRE_EQUAL(match,true);
369 
370  vs.setGPUInsertBuffer(10,1024);
371  CUDA_LAUNCH_DIM3(test_insert_sparse2,10,100,vs.toKernel());
372 
373  vs.flush<sadd_<0>,smin_<1>,smax_<2> >(ctx,flush_type::FLUSH_ON_DEVICE);
374 
375  vs.template deviceToHost<0,1,2>();
376 
377  BOOST_REQUIRE_EQUAL(vs.size(),1500);
378 
379  match = true;
380  for (size_t i = 500 ; i < 1000 ; i++)
381  {
382  match &= vs.template get<0>(9000 - 2*i) == 2*i + 3000;
383  match &= vs.template get<1>(9000 - 2*i) == 2*i + 13000;
384  match &= vs.template get<2>(9000 - 2*i) == 2*i + 23000;
385  }
386 
387  for (size_t i = 0 ; i < 500 ; i++)
388  {
389  match &= vs.template get<0>(9000 - 2*i) == 2*i + 3000 + 2*i + 1100;
390  match &= vs.template get<1>(9000 - 2*i) == 2*i + 11100;
391  match &= vs.template get<2>(9000 - 2*i) == 2*i + 23000;
392 
393  if (match == false)
394  {
395  std::cout << 0 << " " << vs.template get<0>(9000 - 2*i) << " " << 2*i + 3000 + 2*i + 1100 << std::endl;
396  std::cout << 1 << " " << vs.template get<1>(9000 - 2*i) << " " << 2*i + 11100 << std::endl;
397  std::cout << 2 << " " << vs.template get<2>(9000 - 2*i) << " " << 2*i + 23000 << std::endl;
398  }
399  }
400 
401  for (size_t i = 0 ; i < 500 ; i++)
402  {
403  match &= vs.template get<0>(10000 - 2*i) == 2*i + 100;
404  match &= vs.template get<1>(10000 - 2*i) == 2*i + 10100;
405  match &= vs.template get<2>(10000 - 2*i) == 2*i + 20100;
406  }
407 
408  BOOST_REQUIRE_EQUAL(match,true);
409 
410  vs.setGPUInsertBuffer(4000,512);
411  CUDA_LAUNCH_DIM3(test_insert_sparse3,4000,256,vs.toKernel());
412 
413  vs.flush<sadd_<0>,smin_<1>,smax_<2> >(ctx,flush_type::FLUSH_ON_DEVICE);
414  vs.template deviceToHost<0,1,2>();
415 
416  for (size_t i = 0 ; i <= 3500 ; i++)
417  {
418  match &= vs.template get<0>(2*i) == 5;
419  match &= vs.template get<1>(2*i) == 1;
420  match &= vs.template get<2>(2*i) == 1;
421  }
422 
423  BOOST_REQUIRE_EQUAL(match,true);
424 
425  for (size_t i = 3501 ; i <= 4000 ; i++)
426  {
427  match &= vs.template get<0>(2*i) == 5 - 2*i + 3000 + 9000;
428  match &= vs.template get<1>(2*i) == 1;
429  match &= vs.template get<2>(2*i) == 23000 + 9000 - 2*i;
430 
431  if (match == false)
432  {
433  std::cout << i << " " << vs.template get<0>(2*i) << " " << 5 - 2*i + 3000 + 9000 << std::endl;
434  std::cout << i << " " << vs.template get<1>(2*i) << " " << 1 << std::endl;
435  std::cout << i << " " << vs.template get<2>(2*i) << " " << 23000 + 9000 - 2*i << std::endl;
436  }
437  }
438 
439  BOOST_REQUIRE_EQUAL(match,true);
440 
441  for (size_t i = 4001 ; i <= 4500 ; i++)
442  {
443  match &= vs.template get<0>(2*i) == 5 - 2*i + 1100 - 2*i + 3000 + 18000;
444  match &= vs.template get<1>(2*i) == 1;
445  match &= vs.template get<2>(2*i) == 23000 + 9000 - 2*i;
446  }
447 
448  BOOST_REQUIRE_EQUAL(match,true);
449 
450  for (size_t i = 4501 ; i <= 5000 ; i++)
451  {
452  match &= vs.template get<0>(2*i) == 5 - 2*i + 1100 + 9000;
453  match &= vs.template get<1>(2*i) == 1;
454  match &= vs.template get<2>(2*i) == 21100 + 9000 - 2*i;
455  }
456 
457  BOOST_REQUIRE_EQUAL(match,true);
458 }
459 
460 BOOST_AUTO_TEST_CASE( vector_sparse_cuda_gpu_special_function )
461 {
462  openfpm::vector_sparse_gpu<aggregate<size_t,float,double>> vs;
463 
464  vs.template setBackground<0>(17);
465 
466  vs.template setBackground<1>(18);
467 
468  vs.template setBackground<2>(19);
469 
470  vs.setGPUInsertBuffer(10,1024);
471 
472  // we launch a kernel to insert data
473  CUDA_LAUNCH_DIM3(test_insert_sparse,10,100,vs.toKernel());
474  CUDA_LAUNCH_DIM3(test_insert_sparse2_inc,10,100,vs.toKernel());
475  CUDA_LAUNCH_DIM3(test_insert_sparse2_inc,10,100,vs.toKernel());
476  CUDA_LAUNCH_DIM3(test_insert_sparse2_inc,10,100,vs.toKernel());
477 
478  mgpu::ofp_context_t ctx;
479 
480  vs.flush<sstart_<0>>(ctx,flush_type::FLUSH_ON_DEVICE);
481 
482  vs.template deviceToHost<0>();
483 
484  BOOST_REQUIRE_EQUAL(vs.size(),1500);
485 
486  bool match = true;
487 
488  size_t count = 0;
489  for (size_t i = 999 ; i >= 500 ; i--)
490  {
491  match &= vs.template get<0>(9000 - 2*i) == count;
492  count += 3;
493  }
494 
495  for (long int i = 499 ; i >= 0 ; i--)
496  {
497  match &= vs.template get<0>(9000 - 2*i) == count;
498  count += 4;
499  }
500 
501  for (long int i = 499 ; i >= 0 ; i--)
502  {
503  match &= vs.template get<0>(10000 - 2*i) == count;
504  count += 1;
505  }
506 
507  BOOST_REQUIRE_EQUAL(match,true);
508 }
509 
511 
512 void check_lines(openfpm::vector_sparse_gpu<aggregate<size_t,float,double>> & vs,
513  bool s1,
514  bool s2,
515  bool s3,
516  bool s4)
517 {
518  bool match = true;
519 
520  for (size_t i = 0 ; i <= 3500 ; i++)
521  {
522  if (s1 == true)
523  {
524  match &= vs.template get<0>(2*i) == 5;
525  match &= vs.template get<1>(2*i) == 1;
526  match &= vs.template get<2>(2*i) == 1;
527  }
528  else
529  {
530  match &= vs.template get<0>(2*i) == 17;
531  match &= vs.template get<1>(2*i) == 18;
532  match &= vs.template get<2>(2*i) == 19;
533  }
534  }
535 
536  BOOST_REQUIRE_EQUAL(match,true);
537 
538  for (size_t i = 3501 ; i <= 4000 ; i++)
539  {
540  if (s2 == true)
541  {
542  match &= vs.template get<0>(2*i) == 5 - 2*i + 3000 + 9000;
543  match &= vs.template get<1>(2*i) == 1;
544  match &= vs.template get<2>(2*i) == 23000 + 9000 - 2*i;
545  }
546  else
547  {
548  match &= vs.template get<0>(2*i) == 17;
549  match &= vs.template get<1>(2*i) == 18;
550  match &= vs.template get<2>(2*i) == 19;
551  }
552  }
553 
554  BOOST_REQUIRE_EQUAL(match,true);
555 
556  for (size_t i = 4001 ; i <= 4500 ; i++)
557  {
558  if (s3 == true)
559  {
560  match &= vs.template get<0>(2*i) == 5 - 2*i + 1100 - 2*i + 3000 + 18000;
561  match &= vs.template get<1>(2*i) == 1;
562  match &= vs.template get<2>(2*i) == 23000 + 9000 - 2*i;
563  }
564  else
565  {
566  match &= vs.template get<0>(2*i) == 17;
567  match &= vs.template get<1>(2*i) == 18;
568  match &= vs.template get<2>(2*i) == 19;
569  }
570  }
571 
572  BOOST_REQUIRE_EQUAL(match,true);
573 
574  for (size_t i = 4501 ; i <= 5000 ; i++)
575  {
576  if (s4 == true)
577  {
578  match &= vs.template get<0>(2*i) == 5 - 2*i + 1100 + 9000;
579  match &= vs.template get<1>(2*i) == 1;
580  match &= vs.template get<2>(2*i) == 21100 + 9000 - 2*i;
581  }
582  else
583  {
584  match &= vs.template get<0>(2*i) == 17;
585  match &= vs.template get<1>(2*i) == 18;
586  match &= vs.template get<2>(2*i) == 19;
587  }
588  }
589 }
590 
591 
592 BOOST_AUTO_TEST_CASE( vector_sparse_cuda_gpu_remove )
593 {
594  openfpm::vector_sparse_gpu<aggregate<size_t,float,double>> vs;
595 
596  vs.template setBackground<0>(17);
597 
598  vs.template setBackground<1>(18);
599 
600  vs.template setBackground<2>(19);
601 
602  vs.setGPUInsertBuffer(10,1024);
603 
604  // we launch a kernel to insert data
605  CUDA_LAUNCH_DIM3(test_insert_sparse,10,100,vs.toKernel());
606 
607  mgpu::ofp_context_t ctx;
608  vs.flush<sadd_<0>,smin_<1>,smax_<2> >(ctx,flush_type::FLUSH_ON_DEVICE);
609 
610  vs.setGPUInsertBuffer(10,1024);
611  CUDA_LAUNCH_DIM3(test_insert_sparse2,10,100,vs.toKernel());
612 
613  vs.flush<sadd_<0>,smin_<1>,smax_<2> >(ctx,flush_type::FLUSH_ON_DEVICE);
614 
615  vs.setGPUInsertBuffer(4000,512);
616  CUDA_LAUNCH_DIM3(test_insert_sparse3,4000,256,vs.toKernel());
617 
618  vs.flush<sadd_<0>,smin_<1>,smax_<2> >(ctx,flush_type::FLUSH_ON_DEVICE);
619 
620  // we launch a kernel to insert data
621  vs.setGPURemoveBuffer(10,1024);
622  CUDA_LAUNCH_DIM3(test_remove_sparse,10,100,vs.toKernel());
623 
624  size_t sz = vs.size();
625 
626  vs.flush_remove(ctx,flush_type::FLUSH_ON_DEVICE);
627  vs.template deviceToHost<0,1,2>();
628 
629  BOOST_REQUIRE_EQUAL(vs.size(),sz - 1000);
630 
631  check_lines(vs,
632  true,
633  true,
634  false,
635  false);
636 
637  // we launch a kernel to insert data
638  vs.setGPURemoveBuffer(10,1024);
639  CUDA_LAUNCH_DIM3(test_remove_sparse2,10,100,vs.toKernel());
640 
641  vs.flush_remove(ctx,flush_type::FLUSH_ON_DEVICE);
642 
643  BOOST_REQUIRE_EQUAL(vs.size(),sz - 1500);
644 
645  vs.template deviceToHost<0,1,2>();
646 
647  check_lines(vs,
648  true,
649  false,
650  false,
651  false);
652 
653  vs.setGPURemoveBuffer(4000,512);
654  CUDA_LAUNCH_DIM3(test_remove_sparse3,4000,256,vs.toKernel());
655 
656  vs.flush_remove(ctx,flush_type::FLUSH_ON_DEVICE);
657 
658  BOOST_REQUIRE_EQUAL(vs.size(),0);
659 
660  vs.template deviceToHost<0,1,2>();
661 
662  BOOST_REQUIRE_EQUAL(vs.size(),0);
663 
664  check_lines(vs,
665  false,
666  false,
667  false,
668  false);
669 }
670 
671 BOOST_AUTO_TEST_CASE( vector_sparse_cuda_gpu_remove_incremental )
672 {
673  openfpm::vector_sparse_gpu<aggregate<size_t,float,double>> vs;
674 
675  vs.template setBackground<0>(17);
676 
677  vs.template setBackground<1>(18);
678 
679  vs.template setBackground<2>(19);
680 
681  vs.setGPUInsertBuffer(10,1024);
682 
683  // we launch a kernel to insert data
684  CUDA_LAUNCH_DIM3(test_insert_sparse,10,100,vs.toKernel());
685  CUDA_LAUNCH_DIM3(test_insert_sparse2_inc,10,100,vs.toKernel());
686  CUDA_LAUNCH_DIM3(test_insert_sparse2_inc,10,100,vs.toKernel());
687  CUDA_LAUNCH_DIM3(test_insert_sparse2_inc,10,100,vs.toKernel());
688 
689  mgpu::ofp_context_t ctx;
690 
691  vs.flush<sadd_<0>,sadd_<1>,sadd_<2>>(ctx,flush_type::FLUSH_ON_DEVICE);
692 
693  // we launch a kernel to insert data
694  vs.setGPURemoveBuffer(10,1024);
695  CUDA_LAUNCH_DIM3(test_remove_sparse,10,100,vs.toKernel());
696  CUDA_LAUNCH_DIM3(test_remove_sparse2_inc,10,99,vs.toKernel());
697  CUDA_LAUNCH_DIM3(test_remove_sparse2_inc,10,99,vs.toKernel());
698  CUDA_LAUNCH_DIM3(test_remove_sparse2_inc,10,99,vs.toKernel());
699 
700  vs.flush_remove(ctx,flush_type::FLUSH_ON_DEVICE);
701 
702  BOOST_REQUIRE_EQUAL(vs.size(),10);
703 
704  vs.template deviceToHost<0,1,2>();
705 
706  BOOST_REQUIRE_EQUAL(vs.template get<0>(7020),14940);
707  BOOST_REQUIRE_EQUAL(vs.template get<0>(7018),14946);
708  BOOST_REQUIRE_EQUAL(vs.template get<0>(7016),14952);
709  BOOST_REQUIRE_EQUAL(vs.template get<0>(7014),14958);
710  BOOST_REQUIRE_EQUAL(vs.template get<0>(7012),14964);
711  BOOST_REQUIRE_EQUAL(vs.template get<0>(7010),14970);
712  BOOST_REQUIRE_EQUAL(vs.template get<0>(7008),14976);
713  BOOST_REQUIRE_EQUAL(vs.template get<0>(7006),14982);
714  BOOST_REQUIRE_EQUAL(vs.template get<0>(7004),14988);
715  BOOST_REQUIRE_EQUAL(vs.template get<0>(7002),14994);
716 
717  BOOST_REQUIRE_EQUAL(vs.template get<1>(7020),44940);
718  BOOST_REQUIRE_EQUAL(vs.template get<1>(7018),44946);
719  BOOST_REQUIRE_EQUAL(vs.template get<1>(7016),44952);
720  BOOST_REQUIRE_EQUAL(vs.template get<1>(7014),44958);
721  BOOST_REQUIRE_EQUAL(vs.template get<1>(7012),44964);
722  BOOST_REQUIRE_EQUAL(vs.template get<1>(7010),44970);
723  BOOST_REQUIRE_EQUAL(vs.template get<1>(7008),44976);
724  BOOST_REQUIRE_EQUAL(vs.template get<1>(7006),44982);
725  BOOST_REQUIRE_EQUAL(vs.template get<1>(7004),44988);
726  BOOST_REQUIRE_EQUAL(vs.template get<1>(7002),44994);
727 
728  BOOST_REQUIRE_EQUAL(vs.template get<2>(7020),74940);
729  BOOST_REQUIRE_EQUAL(vs.template get<2>(7018),74946);
730  BOOST_REQUIRE_EQUAL(vs.template get<2>(7016),74952);
731  BOOST_REQUIRE_EQUAL(vs.template get<2>(7014),74958);
732  BOOST_REQUIRE_EQUAL(vs.template get<2>(7012),74964);
733  BOOST_REQUIRE_EQUAL(vs.template get<2>(7010),74970);
734  BOOST_REQUIRE_EQUAL(vs.template get<2>(7008),74976);
735  BOOST_REQUIRE_EQUAL(vs.template get<2>(7006),74982);
736  BOOST_REQUIRE_EQUAL(vs.template get<2>(7004),74988);
737  BOOST_REQUIRE_EQUAL(vs.template get<2>(7002),74994);
738 }
739 
740 
741 BOOST_AUTO_TEST_SUITE_END()
742 
743 #endif /* MAP_VECTOR_SPARSE_CUDA_KER_UNIT_TESTS_CUH_ */
openfpm::vector_sparse::flush_remove
void flush_remove(mgpu::ofp_context_t &context, flush_type opt=FLUSH_ON_HOST)
merge the added element to the main data array
Definition: map_vector_sparse.hpp:1812
smax_
Definition: map_vector_sparse_cuda_kernels.cuh:174
sadd_
Definition: map_vector_sparse_cuda_kernels.cuh:91
smin_
Definition: map_vector_sparse_cuda_kernels.cuh:249
openfpm::vector_sparse::setGPURemoveBuffer
void setGPURemoveBuffer(int nblock, int nslot)
set the gpu remove buffer for every block
Definition: map_vector_sparse.hpp:1936
openfpm::vector_sparse::setGPUInsertBuffer
void setGPUInsertBuffer(int nblock, int nslot)
set the gpu insert buffer for every block
Definition: map_vector_sparse.hpp:1893
openfpm::vector_sparse::size
size_t size()
Return how many element you have in this map.
Definition: map_vector_sparse.hpp:1830
openfpm::vector_sparse::toKernel
vector_sparse_gpu_ker< T, Ti, layout_base > toKernel()
toKernel function transform this structure into one that can be used on GPU
Definition: map_vector_sparse.hpp:1874
openfpm::vector_sparse::flush
void flush(mgpu::ofp_context_t &context, flush_type opt=FLUSH_ON_HOST)
merge the added element to the main data array
Definition: map_vector_sparse.hpp:1794
openfpm::vector_sparse::clear
void clear()
Clear all from all the elements.
Definition: map_vector_sparse.hpp:1962
aggregate
aggregate of properties, from a list of object if create a struct that follow the OPENFPM native stru...
Definition: aggregate.hpp:214
openfpm::vector_sparse
Definition: map_vector_sparse.hpp:790
openfpm::vector
Implementation of 1-D std::vector like structure.
Definition: map_vector.hpp:202