OpenFPM_pdata  4.1.0
Project that contain the implementation of distributed structures
 
Loading...
Searching...
No Matches
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
17template<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
52template<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
65template<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
82template<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
96template<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
113template<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
127template<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
144template<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
158template<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
190BOOST_AUTO_TEST_SUITE( vector_cuda_sparse )
191
192BOOST_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 gpu::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
261BOOST_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 gpu::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
339BOOST_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 gpu::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
460BOOST_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 gpu::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
512void 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
592BOOST_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 gpu::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
671BOOST_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 gpu::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
741BOOST_AUTO_TEST_SUITE_END()
742
743#endif /* MAP_VECTOR_SPARSE_CUDA_KER_UNIT_TESTS_CUH_ */
openfpm::vector_sparse
Definition map_vector_sparse.hpp:789
openfpm::vector_sparse::flush_remove
void flush_remove(gpu::ofp_context_t &context, flush_type opt=FLUSH_ON_HOST)
merge the added element to the main data array
Definition map_vector_sparse.hpp:1811
openfpm::vector_sparse::setGPURemoveBuffer
void setGPURemoveBuffer(int nblock, int nslot)
set the gpu remove buffer for every block
Definition map_vector_sparse.hpp:1935
openfpm::vector_sparse::clear
void clear()
Clear all from all the elements.
Definition map_vector_sparse.hpp:1961
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:1873
openfpm::vector_sparse::size
size_t size()
Return how many element you have in this map.
Definition map_vector_sparse.hpp:1829
openfpm::vector_sparse::flush
void flush(gpu::ofp_context_t &context, flush_type opt=FLUSH_ON_HOST)
merge the added element to the main data array
Definition map_vector_sparse.hpp:1793
openfpm::vector_sparse::setGPUInsertBuffer
void setGPUInsertBuffer(int nblock, int nslot)
set the gpu insert buffer for every block
Definition map_vector_sparse.hpp:1892
openfpm::vector
Implementation of 1-D std::vector like structure.
Definition map_vector.hpp:203
aggregate
aggregate of properties, from a list of object if create a struct that follow the OPENFPM native stru...
Definition aggregate.hpp:215
gpu::ofp_context_t
Definition ofp_context.hpp:303
sadd_
Definition map_vector_sparse_cuda_kernels.cuh:107
smax_
Definition map_vector_sparse_cuda_kernels.cuh:192
smin_
Definition map_vector_sparse_cuda_kernels.cuh:269