PeriDyno 1.0.0
Loading...
Searching...
No Matches
VirtualSpatiallyAdaptiveStrategy.cu
Go to the documentation of this file.
1#include "VirtualSpatiallyAdaptiveStrategy.h"
2#include "Node.h"
3#include "ParticleSystem/Module/SummationDensity.h"
4#include "Topology/GridHash.h"
5
6#include <thrust/sort.h>
7
8
9namespace dyno
10{
11 __constant__ int diff_v[27][3] =
12 {
13 0, 0, 0,
14 0, 0, 1,
15 0, 1, 0,
16 1, 0, 0,
17 0, 0, -1,
18 0, -1, 0,
19 -1, 0, 0,
20 0, 1, 1,
21 0, 1, -1,
22 0, -1, 1,
23 0, -1, -1,
24 1, 0, 1,
25 1, 0, -1,
26 -1, 0, 1,
27 -1, 0, -1,
28 1, 1, 0,
29 1, -1, 0,
30 -1, 1, 0,
31 -1, -1, 0,
32 1, 1, 1,
33 1, 1, -1,
34 1, -1, 1,
35 -1, 1, 1,
36 1, -1, -1,
37 -1, 1, -1,
38 -1, -1, 1,
39 -1, -1, -1
40 };
41
42
43 __constant__ int diff_v_33[33][3] =
44 {
45 0, 0, 0,
46 0, 0, 1,
47 0, 1, 0,
48 1, 0, 0,
49 0, 0, -1,
50 0, -1, 0,
51 -1, 0, 0,
52 0, 1, 1,
53 0, 1, -1,
54 0, -1, 1,
55 0, -1, -1,
56 1, 0, 1,
57 1, 0, -1,
58 -1, 0, 1,
59 -1, 0, -1,
60 1, 1, 0,
61 1, -1, 0,
62 -1, 1, 0,
63 -1, -1, 0,
64 1, 1, 1,
65 1, 1, -1,
66 1, -1, 1,
67 -1, 1, 1,
68 1, -1, -1,
69 -1, 1, -1,
70 -1, -1, 1,
71 -1, -1, -1,
72 2, 0, 0,
73 -2, 0, 0,
74 0, 2, 0,
75 0, -2, 0,
76 0, 0, 2,
77 0, 0, -2
78
79 };
80
81
82 __constant__ int diff_v_125[125][3] =
83 {
84 -2, -2, -2,
85 -2, -2, -1,
86 -2, -2, 0,
87 -2, -2, 1,
88 -2, -2, 2,
89 -2, -1, -2,
90 -2, -1, -1,
91 -2, -1, 0,
92 -2, -1, 1,
93 -2, -1, 2,
94 -2, 0, -2,
95 -2, 0, -1,
96 -2, 0, 0,
97 -2, 0, 1,
98 -2, 0, 2,
99 -2, 1, -2,
100 -2, 1, -1,
101 -2, 1, 0,
102 -2, 1, 1,
103 -2, 1, 2,
104 -2, 2, -2,
105 -2, 2, -1,
106 -2, 2, 0,
107 -2, 2, 1,
108 -2, 2, 2,
109 -1, -2, -2,
110 -1, -2, -1,
111 -1, -2, 0,
112 -1, -2, 1,
113 -1, -2, 2,
114 -1, -1, -2,
115 -1, -1, -1,
116 -1, -1, 0,
117 -1, -1, 1,
118 -1, -1, 2,
119 -1, 0, -2,
120 -1, 0, -1,
121 -1, 0, 0,
122 -1, 0, 1,
123 -1, 0, 2,
124 -1, 1, -2,
125 -1, 1, -1,
126 -1, 1, 0,
127 -1, 1, 1,
128 -1, 1, 2,
129 -1, 2, -2,
130 -1, 2, -1,
131 -1, 2, 0,
132 -1, 2, 1,
133 -1, 2, 2,
134 0, -2, -2,
135 0, -2, -1,
136 0, -2, 0,
137 0, -2, 1,
138 0, -2, 2,
139 0, -1, -2,
140 0, -1, -1,
141 0, -1, 0,
142 0, -1, 1,
143 0, -1, 2,
144 0, 0, -2,
145 0, 0, -1,
146 0, 0, 0,
147 0, 0, 1,
148 0, 0, 2,
149 0, 1, -2,
150 0, 1, -1,
151 0, 1, 0,
152 0, 1, 1,
153 0, 1, 2,
154 0, 2, -2,
155 0, 2, -1,
156 0, 2, 0,
157 0, 2, 1,
158 0, 2, 2,
159 1, -2, -2,
160 1, -2, -1,
161 1, -2, 0,
162 1, -2, 1,
163 1, -2, 2,
164 1, -1, -2,
165 1, -1, -1,
166 1, -1, 0,
167 1, -1, 1,
168 1, -1, 2,
169 1, 0, -2,
170 1, 0, -1,
171 1, 0, 0,
172 1, 0, 1,
173 1, 0, 2,
174 1, 1, -2,
175 1, 1, -1,
176 1, 1, 0,
177 1, 1, 1,
178 1, 1, 2,
179 1, 2, -2,
180 1, 2, -1,
181 1, 2, 0,
182 1, 2, 1,
183 1, 2, 2,
184 2, -2, -2,
185 2, -2, -1,
186 2, -2, 0,
187 2, -2, 1,
188 2, -2, 2,
189 2, -1, -2,
190 2, -1, -1,
191 2, -1, 0,
192 2, -1, 1,
193 2, -1, 2,
194 2, 0, -2,
195 2, 0, -1,
196 2, 0, 0,
197 2, 0, 1,
198 2, 0, 2,
199 2, 1, -2,
200 2, 1, -1,
201 2, 1, 0,
202 2, 1, 1,
203 2, 1, 2,
204 2, 2, -2,
205 2, 2, -1,
206 2, 2, 0,
207 2, 2, 1,
208 2, 2, 2
209 };
210
211
212 IMPLEMENT_TCLASS(VirtualSpatiallyAdaptiveStrategy, TDataType)
213
214 template<typename TDataType>
215 VirtualSpatiallyAdaptiveStrategy<TDataType>::VirtualSpatiallyAdaptiveStrategy()
216 : VirtualParticleGenerator<TDataType>()
217 {
218
219 this->varSamplingDistance()->setValue(Real(0.005));
220 this->varRestDensity()->setValue(Real(1000));
221 gridSize = this->varSamplingDistance()->getData();
222
223
224 }
225
226 template<typename TDataType>
227 VirtualSpatiallyAdaptiveStrategy<TDataType>::~VirtualSpatiallyAdaptiveStrategy()
228 {
229
230 }
231
232 template<typename Real, typename Coord>
233 __global__ void AFV_VirtualPositionCompute(
234 DArray<Coord> pos,
235 DArray<Coord> vpos,
236 Coord loPoint,
237 int nx,
238 int ny,
239 int nz,
240 Real ds
241 )
242 {
243 int pId = threadIdx.x + (blockIdx.x * blockDim.x);
244 if (pId >= vpos.size()) return;
245 if (nx == 0) return;
246
247 vpos[pId][1] = loPoint[1] + Real(pId / (nx * nz)) * ds;
248 vpos[pId][0] = loPoint[0] + Real(pId % nx) * ds;
249 vpos[pId][2] = loPoint[2] + Real(pId % (nx * nz) / nx) * ds;
250 }
251
252 /*
253 *@brief Virtual particles' candinate points
254 * Every particle has 8 neighbors.
255 */
256 template<typename Real, typename Coord>
257 __global__ void AFV_AnchorNeighbor_8
258 (
259 DArray<Coord> anchorPoint,
260 DArray<Coord> pos,
261 Coord origin,
262 Real dh
263 )
264 {
265 int id = threadIdx.x + (blockIdx.x * blockDim.x);
266 if (id >= pos.size()) return;
267
268
269
270 Coord pos_ref = pos[id] - origin + Coord (dh/4.0);
271 //Coord pos_ref = pos[id] - origin;
272 Coord a(0);
273
274 a[0] = (Real)((int)(floor(pos_ref[0] / dh))) * dh;
275 a[1] = (Real)((int)(floor(pos_ref[1] / dh))) * dh;
276 a[2] = (Real)((int)(floor(pos_ref[2] / dh))) * dh;
277
278 anchorPoint[id * 8] = Coord(a[0], a[1], a[2]);
279 anchorPoint[id * 8 + 1] = Coord(a[0] + dh, a[1], a[2]);
280 anchorPoint[id * 8 + 2] = Coord(a[0] + dh, a[1] + dh, a[2]);
281 anchorPoint[id * 8 + 3] = Coord(a[0] + dh, a[1] + dh, a[2] + dh);
282 anchorPoint[id * 8 + 4] = Coord(a[0], a[1] + dh, a[2]);
283 anchorPoint[id * 8 + 5] = Coord(a[0], a[1] + dh, a[2] + dh);
284 anchorPoint[id * 8 + 6] = Coord(a[0], a[1], a[2] + dh);
285 anchorPoint[id * 8 + 7] = Coord(a[0] + dh, a[1], a[2] + dh);
286
287 }
288
289
290 /*
291 *@brief Virtual particles' candinate points
292 * Every particle has 27 neighbors.
293 */
294 template<typename Real, typename Coord>
295 __global__ void AFV_AnchorNeighbor_27
296 (
297 DArray<Coord> anchorPoint,
298 DArray<Coord> pos,
299 Coord origin,
300 Real dh
301 )
302 {
303 int id = threadIdx.x + (blockIdx.x * blockDim.x);
304 if (id >= pos.size()) return;
305
306 Coord pos_ref = pos[id] - origin + dh/2;
307 Coord a(0);
308
309 a[0] = (Real)((int)(floor(pos_ref[0] / dh))) * dh;
310 a[1] = (Real)((int)(floor(pos_ref[1] / dh))) * dh;
311 a[2] = (Real)((int)(floor(pos_ref[2] / dh))) * dh;
312
313
314 for (int i = 0; i < 27; i++)
315 {
316 anchorPoint[id * 27 + i] = Coord( a[0] + dh * Real(diff_v[i][0]),
317 a[1] + dh * Real(diff_v[i][1]),
318 a[2] + dh * Real(diff_v[i][2])
319 );
320 }
321
322 }
323
324 /*
325*@brief Virtual particles' candinate points
326* Every particle has 33 neighbors.
327*/
328 template<typename Real, typename Coord>
329 __global__ void AFV_AnchorNeighbor_33
330 (
331 DArray<Coord> anchorPoint,
332 DArray<Coord> pos,
333 Coord origin,
334 Real dh
335 )
336 {
337 int id = threadIdx.x + (blockIdx.x * blockDim.x);
338 if (id >= pos.size()) return;
339
340 //Coord pos_ref = pos[id] - origin + dh / 2;
341 Coord pos_ref = pos[id] - origin + dh / 2;
342 Coord a(0);
343
344 a[0] = (Real)((int)(floor(pos_ref[0] / dh))) * dh;
345 a[1] = (Real)((int)(floor(pos_ref[1] / dh))) * dh;
346 a[2] = (Real)((int)(floor(pos_ref[2] / dh))) * dh;
347
348
349 for (int i = 0; i < 33; i++)
350 {
351 anchorPoint[id * 33 + i] = Coord(a[0] + dh * Real(diff_v_33[i][0]),
352 a[1] + dh * Real(diff_v_33[i][1]),
353 a[2] + dh * Real(diff_v_33[i][2])
354 );
355 }
356
357 }
358
359 /*
360*@brief Virtual particles' candinate points
361* Every particle has 125 neighbors.
362*/
363 template<typename Real, typename Coord>
364 __global__ void AFV_AnchorNeighbor_125
365 (
366 DArray<Coord> anchorPoint,
367 DArray<Coord> pos,
368 Coord origin,
369 Real dh
370 )
371 {
372 int id = threadIdx.x + (blockIdx.x * blockDim.x);
373 if (id >= pos.size()) return;
374
375 Coord pos_ref = pos[id] - origin + dh / 2;
376 Coord a(0);
377
378 a[0] = (Real)((int)(floor(pos_ref[0] / dh))) * dh;
379 a[1] = (Real)((int)(floor(pos_ref[1] / dh))) * dh;
380 a[2] = (Real)((int)(floor(pos_ref[2] / dh))) * dh;
381
382
383 for (int i = 0; i < 125; i++)
384 {
385 anchorPoint[id * 125 + i] = Coord(a[0] + dh * Real(diff_v_125[i][0]),
386 a[1] + dh * Real(diff_v_125[i][1]),
387 a[2] + dh * Real(diff_v_125[i][2])
388 );
389 }
390
391 }
392
393
394 template< typename Coord>
395 __global__ void AFV_PositionToMortonCode_TEST
396 (
397 DArray<uint32_t> mortonCode,
398 DArray<Coord> gridPoint,
399 Real dh
400 )
401 {
402 int id = threadIdx.x + (blockIdx.x * blockDim.x);
403 if (id >= gridPoint.size()) return;
404
405 uint32_t x = (uint32_t)(gridPoint[id][0] / dh);
406 uint32_t y = (uint32_t)(gridPoint[id][1] / dh);
407 uint32_t z = (uint32_t)(gridPoint[id][2] / dh);
408
409 uint32_t xi = x;
410 uint32_t yi = y;
411 uint32_t zi = z;
412
413 uint32_t key = 0;
414
415
416 xi = (xi | (xi << 16)) & 0x030000FF;
417 xi = (xi | (xi << 8)) & 0x0300F00F;
418 xi = (xi | (xi << 4)) & 0x030C30C3;
419 xi = (xi | (xi << 2)) & 0x09249249;
420
421 yi = (yi | (yi << 16)) & 0x030000FF;
422 yi = (yi | (yi << 8)) & 0x0300F00F;
423 yi = (yi | (yi << 4)) & 0x030C30C3;
424 yi = (yi | (yi << 2)) & 0x09249249;
425
426 zi = (zi | (zi << 16)) & 0x030000FF;
427 zi = (zi | (zi << 8)) & 0x0300F00F;
428 zi = (zi | (zi << 4)) & 0x030C30C3;
429 zi = (zi | (zi << 2)) & 0x09249249;
430
431 key = xi | (yi << 1) | (zi << 2);
432
433 uint32_t xv = key & 0x09249249;
434 uint32_t yv = (key >> 1) & 0x09249249;
435 uint32_t zv = (key >> 2) & 0x09249249;
436
437 xv = ((xv >> 2) | xv) & 0x030C30C3;
438 xv = ((xv >> 4) | xv) & 0x0300F00F;
439 xv = ((xv >> 8) | xv) & 0x030000FF;
440 xv = ((xv >> 16) | xv) & 0x000003FF;
441
442 yv = ((yv >> 2) | yv) & 0x030C30C3;
443 yv = ((yv >> 4) | yv) & 0x0300F00F;
444 yv = ((yv >> 8) | yv) & 0x030000FF;
445 yv = ((yv >> 16) | yv) & 0x000003FF;
446
447 zv = ((zv >> 2) | zv) & 0x030C30C3;
448 zv = ((zv >> 4) | zv) & 0x0300F00F;
449 zv = ((zv >> 8) | zv) & 0x030000FF;
450 zv = ((zv >> 16) | zv) & 0x000003FF;
451
452 if(x!=xv || y!=yv || z!=zv)
453 printf("gridPoint: %d, %d, %d || key: %d || invert: %d, %d, %d \r\n", x, y, z, key, xv, yv, zv);
454
455 }
456
457 /*
458 *@brief Positions(3-dimension) convert to Morton Codes
459 */
460 template< typename Coord>
461 __global__ void AFV_PositionToMortonCode
462 (
463 DArray<uint32_t> mortonCode,
464 DArray<Coord> gridPoint,
465 Real dh
466 )
467 {
468 int id = threadIdx.x + (blockIdx.x * blockDim.x);
469 if (id >= gridPoint.size()) return;
470
471 uint32_t x = (uint32_t)(gridPoint[id][0] / dh + 100 * EPSILON);
472 uint32_t y = (uint32_t)(gridPoint[id][1] / dh + 100 * EPSILON);
473 uint32_t z = (uint32_t)(gridPoint[id][2] / dh + 100 * EPSILON);
474
475 uint32_t xi = x;
476 uint32_t yi = y;
477 uint32_t zi = z;
478
479 uint32_t key = 0;
480
481
482 xi = (xi | (xi << 16)) & 0x030000FF;
483 xi = (xi | (xi << 8)) & 0x0300F00F;
484 xi = (xi | (xi << 4)) & 0x030C30C3;
485 xi = (xi | (xi << 2)) & 0x09249249;
486
487 yi = (yi | (yi << 16)) & 0x030000FF;
488 yi = (yi | (yi << 8)) & 0x0300F00F;
489 yi = (yi | (yi << 4)) & 0x030C30C3;
490 yi = (yi | (yi << 2)) & 0x09249249;
491
492 zi = (zi | (zi << 16)) & 0x030000FF;
493 zi = (zi | (zi << 8)) & 0x0300F00F;
494 zi = (zi | (zi << 4)) & 0x030C30C3;
495 zi = (zi | (zi << 2)) & 0x09249249;
496
497 key = xi | (yi << 1) | (zi << 2);
498
499 mortonCode[id] = key;
500 }
501
502
503 /*
504 *@brief Morton Codes convert to Positions(3-dimension)
505 */
506 template< typename Coord>
507 __global__ void AFV_MortonCodeToPosition
508 (
509 DArray<Coord> gridPoint,
510 DArray<uint32_t> mortonCode,
511 Real dh
512 )
513 {
514 int id = threadIdx.x + (blockIdx.x * blockDim.x);
515 if (id >= mortonCode.size()) return;
516
517
518 uint32_t key = mortonCode[id];
519
520 uint32_t xv = key & 0x09249249;
521 uint32_t yv = (key >> 1) & 0x09249249;
522 uint32_t zv = (key >> 2) & 0x09249249;
523
524 xv = ((xv >> 2) | xv) & 0x030C30C3;
525 xv = ((xv >> 4) | xv) & 0x0300F00F;
526 xv = ((xv >> 8) | xv) & 0x030000FF;
527 xv = ((xv >> 16) | xv) & 0x000003FF;
528
529 yv = ((yv >> 2) | yv) & 0x030C30C3;
530 yv = ((yv >> 4) | yv) & 0x0300F00F;
531 yv = ((yv >> 8) | yv) & 0x030000FF;
532 yv = ((yv >> 16) | yv) & 0x000003FF;
533
534 zv = ((zv >> 2) | zv) & 0x030C30C3;
535 zv = ((zv >> 4) | zv) & 0x0300F00F;
536 zv = ((zv >> 8) | zv) & 0x030000FF;
537 zv = ((zv >> 16) | zv) & 0x000003FF;
538
539 Real x = (float)(xv)* dh;
540 Real y = (float)(yv)* dh;
541 Real z = (float)(zv)* dh;
542
543 gridPoint[id] = Coord(x, y, z);
544 }
545
546
547 /*
548 *@brief Search for adjacent elements with the same value
549 */
550 __global__ void AFV_RepeatedElementSearch
551 (
552 DArray<uint32_t> morton,
553 DArray<uint32_t> counter
554 )
555 {
556 int id = threadIdx.x + (blockIdx.x * blockDim.x);
557 if (id >= morton.size()) return;
558
559
560 if (id == 0 || morton[id] != morton[id - 1])
561 {
562 counter[id] = 1;
563 }
564 else
565 counter[id] = 0;
566 }
567
568 /*
569 *@brief Accumulate the number of non-repeating elements
570 */
571 __global__ void AFV_nonRepeatedElementsCal
572 (
573 DArray<uint32_t> non_repeated_elements,
574 DArray<uint32_t> post_elements,
575 DArray<uint32_t> counter
576 )
577 {
578 int id = threadIdx.x + (blockIdx.x * blockDim.x);
579 if (id >= post_elements.size()) return;
580
581 if (id == 0 || post_elements[id] != post_elements[id - 1])
582 {
583 non_repeated_elements[counter[id]] = post_elements[id];
584 }
585 }
586
587
588
589
590 template< typename Coord>
591 __global__ void AFV_CopyToVpos(
592 DArray<Coord> vpos,
593 Coord origin,
594 DArray<Coord> elements
595
596 )
597 {
598 int id = threadIdx.x + (blockIdx.x * blockDim.x);
599 if (id >= elements.size()) return;
600 vpos[id] = elements[id] + origin;
601 }
602
603
604 __global__ void AFV_ElementsPrint(
605 DArray<uint32_t> elements
606 )
607 {
608 int id = threadIdx.x + (blockIdx.x * blockDim.x);
609 if (id >= elements.size()) return;
610
611 if (id == 100)
612 {
613 printf("size: %d \r\n", elements.size());
614 for(int i = 0; i < elements.size(); i++)
615 printf("%d\t", elements[i]);
616 }
617 }
618
619
620 template< typename Real>
621 __global__ void AFV_SortMortonCodePrint
622 (
623 DArray<uint32_t> mortonCode,
624 DArray<uint32_t> counter,
625 Real dh
626 )
627 {
628 int id = threadIdx.x + (blockIdx.x * blockDim.x);
629 if (id >= mortonCode.size()) return;
630
631 if (id == 0)
632 {
633 for (int i = 0; i < mortonCode.size(); i++)
634 {
635 printf("%d , %d \t", counter[i], mortonCode[i]);
636 }
637 }
638 }
639
640
641
642 template<typename TDataType>
643 void VirtualSpatiallyAdaptiveStrategy<TDataType>::constrain()
644 {
645 cudaDeviceSynchronize();
646 gridSize = this->varSamplingDistance()->getData();
647 int num = this->inRPosition()->size();
648 if (num == 0) return;
649
650 if (num == 0) return;
651
652 int node_num = num * (int)(this->varCandidatePointCount()->getDataPtr()->currentKey());
653
654 if (m_anchorPoint.size() != node_num)
655 {
656 m_anchorPoint.resize(node_num);
657 }
658
659 if (m_anchorPointCodes.size() != node_num)
660 {
661 m_anchorPointCodes.resize(node_num);
662 }
663
664 if (m_nonRepeatedCount.size() != node_num)
665 {
666 m_nonRepeatedCount.resize(node_num);
667 }
668
669 Reduction<Coord> reduce;
670 Coord hiBound = reduce.maximum(this->inRPosition()->getData().begin(), this->inRPosition()->getData().size());
671 Coord loBound = reduce.minimum(this->inRPosition()->getData().begin(), this->inRPosition()->getData().size());
672
673 int padding = 2;
674 hiBound += Coord((padding + 1)* gridSize);
675 loBound -= Coord(padding * gridSize);
676
677 loBound[0] = (Real)((int)(floor(loBound[0] / gridSize)))*gridSize;
678 loBound[1] = (Real)((int)(floor(loBound[1] / gridSize)))*gridSize;
679 loBound[2] = (Real)((int)(floor(loBound[2] / gridSize)))*gridSize;
680
681 origin = Coord(loBound[0], loBound[1], loBound[2]);
682
683 if (this->varCandidatePointCount()->getValue() == CandidatePointCount::neighbors_8) {
684 cuExecute(this->inRPosition()->getData().size(),
685 AFV_AnchorNeighbor_8,
686 m_anchorPoint,
687 this->inRPosition()->getData(),
688 origin,
689 gridSize
690 );
691
692 }
693 else if (this->varCandidatePointCount()->getValue() == CandidatePointCount::neighbors_27) {
694 cuExecute(this->inRPosition()->getData().size(),
695 AFV_AnchorNeighbor_27,
696 m_anchorPoint,
697 this->inRPosition()->getData(),
698 origin,
699 gridSize
700 );
701 }
702 else if (this->varCandidatePointCount()->getValue() == CandidatePointCount::neighbors_33) {
703 cuExecute(this->inRPosition()->getData().size(),
704 AFV_AnchorNeighbor_33,
705 m_anchorPoint,
706 this->inRPosition()->getData(),
707 origin,
708 gridSize
709 );
710 }
711 else if (this->varCandidatePointCount()->getValue() == CandidatePointCount::neighbors_125) {
712 cuExecute(this->inRPosition()->getData().size(),
713 AFV_AnchorNeighbor_125,
714 m_anchorPoint,
715 this->inRPosition()->getData(),
716 origin,
717 gridSize
718 );
719 }
720 else
721 {
722 std::cout << "*VIRTUAL PARTICLE:: AdaptiveVirtualPosition ERROR!!!!!" << std::endl;
723 }
724
725 cuExecute(m_anchorPointCodes.size(),
726 AFV_PositionToMortonCode,
727 m_anchorPointCodes,
728 m_anchorPoint,
729 gridSize
730 );
731
732 thrust::sort(thrust::device, m_anchorPointCodes.begin(), m_anchorPointCodes.begin() + m_anchorPointCodes.size());
733
734 cuExecute(m_anchorPointCodes.size(),
735 AFV_RepeatedElementSearch,
736 m_anchorPointCodes,
737 m_nonRepeatedCount
738 );
739
740
741 int candidatePoint_num = thrust::reduce(thrust::device, m_nonRepeatedCount.begin(), m_nonRepeatedCount.begin() + m_nonRepeatedCount.size(), (int)0, thrust::plus<uint32_t>());
742
743 thrust::exclusive_scan(thrust::device, m_nonRepeatedCount.begin(), m_nonRepeatedCount.begin() + m_nonRepeatedCount.size(), m_nonRepeatedCount.begin());
744
745 m_candidateCodes.resize(candidatePoint_num);
746
747
748 cuExecute(m_anchorPointCodes.size(),
749 AFV_nonRepeatedElementsCal,
750 m_candidateCodes,
751 m_anchorPointCodes,
752 m_nonRepeatedCount
753 );
754
755 cuSynchronize();
756
757 m_virtual_position.resize(m_candidateCodes.size());
758
759 cuExecute(m_candidateCodes.size(),
760 AFV_MortonCodeToPosition,
761 m_virtual_position,
762 m_candidateCodes,
763 gridSize
764 );
765
766
767 if (this->outVirtualParticles()->isEmpty())
768 {
769 this->outVirtualParticles()->allocate();
770 }
771
772 this->outVirtualParticles()->resize(m_virtual_position.size());
773
774 cuExecute(m_virtual_position.size(),
775 AFV_CopyToVpos,
776 this->outVirtualParticles()->getData(),
777 origin,
778 m_virtual_position
779 );
780
781 std::cout << "*DUAL-ISPH::SpatiallyAdaptiveStrategy(S.C.)::Model_"<< this->varCandidatePointCount()->getDataPtr()->currentKey()
782 <<"::RealPoints:" << this->inRPosition()->size() << "||VirtualPoints:" << this->outVirtualParticles()->size() <<std::endl;
783 }
784
785
786
787 DEFINE_CLASS(VirtualSpatiallyAdaptiveStrategy);
788
789}
Generated by doxygen 1.13.2