milupHPC documentation
  • src
particles.cu
Go to the documentation of this file.
1#include "../include/particles.cuh"
2#include "../include/cuda_utils/cuda_launcher.cuh"
3
4CUDA_CALLABLE_MEMBER Particles::Particles() {
5
6}
7
8#if DIM == 1
9
10CUDA_CALLABLE_MEMBER Particles::Particles(integer *numParticles, integer *numNodes, real *mass, real *x, real *vx,
11 real *ax, integer *level, idInteger *uid, integer *materialId,
12 real *sml, integer *nnl, integer *noi, real *e, real *dedt, real *cs,
13 real *rho, real *p) :
14 numParticles(numParticles), numNodes(numNodes), mass(mass), x(x), vx(vx),
15 ax(ax), level(level), uid(uid), materialId(materialId), sml(sml),
16 nnl(nnl), noi(noi), e(e), dedt(dedt), cs(cs), rho(rho), p(p) {
17
18}
19
20CUDA_CALLABLE_MEMBER void Particles::set(integer *numParticles, integer *numNodes, real *mass, real *x, real *vx,
21 real *ax, integer *level, idInteger *uid, integer *materialId,
22 real *sml, integer *nnl, integer *noi, real *e, real *dedt,
23 real *cs, real *rho, real *p) {
24
25 this->numParticles = numParticles;
26 this->numNodes = numNodes;
27 this->mass = mass;
28 this->x = x;
29 this->vx = vx;
30 this->ax = ax;
31 this->level = level;
32 this->uid = uid;
33 this->materialId = materialId;
34 this->sml = sml;
35 this->nnl = nnl;
36 this->noi = noi;
37 this->e = e;
38 this->dedt = dedt;
39 this->cs = cs;
40 this->rho = rho;
41 this->p = p;
42
43}
44
45#elif DIM == 2
46
47CUDA_CALLABLE_MEMBER Particles::Particles(integer *numParticles, integer *numNodes, real *mass, real *x, real *y,
48 real *vx, real *vy, real *ax, real *ay, idInteger *uid,
49 integer *materialId, real *sml, integer *nnl, integer *noi,
50 real *e, real *dedt, real *cs, real *rho, real *p) :
51 numParticles(numParticles), numNodes(numNodes),
52 mass(mass), x(x), y(y), vx(vx), vy(vy), ax(ax), ay(ay),
53 uid(uid), materialId(materialId), sml(sml), nnl(nnl), noi(noi),
54 e(e), dedt(dedt), cs(cs), rho(rho), p(p) {
55
56}
57
58CUDA_CALLABLE_MEMBER void Particles::set(integer *numParticles, integer *numNodes, real *mass, real *x, real *y,
59 real *vx, real *vy, real *ax, real *ay, integer *level, idInteger *uid,
60 integer *materialId, real *sml, integer *nnl, integer *noi, real *e,
61 real *dedt, real *cs, real *rho, real *p) {
62
63 this->numParticles = numParticles;
64 this->numNodes = numNodes;
65 this->mass = mass;
66 this->x = x;
67 this->y = y;
68 this->vx = vx;
69 this->vy = vy;
70 this->ax = ax;
71 this->ay = ay;
72 this->level = level;
73 this->uid = uid;
74 this->materialId = materialId;
75 this->sml = sml;
76 this->nnl = nnl;
77 this->noi = noi;
78 this->e = e;
79 this->dedt = dedt;
80 this->cs = cs;
81 this->rho = rho;
82 this->p = p;
83}
84
85#else
86
87CUDA_CALLABLE_MEMBER Particles::Particles(integer *numParticles, integer *numNodes, real *mass, real *x, real *y,
88 real *z, real *vx, real *vy, real *vz, real *ax, real *ay, real *az,
89 idInteger *uid, integer *materialId, real *sml, integer *nnl, integer *noi,
90 real *e, real *dedt, real *cs, real *rho, real *p) :
91 numParticles(numParticles), numNodes(numNodes), mass(mass), x(x), y(y), z(z),
92 vx(vx), vy(vy), vz(vz), ax(ax), ay(ay), az(az), uid(uid),
93 materialId(materialId), sml(sml), nnl(nnl), noi(noi), e(e), dedt(dedt),
94 cs(cs), rho(rho), p(p) {
95
96}
97
98CUDA_CALLABLE_MEMBER void Particles::set(integer *numParticles, integer *numNodes, real *mass, real *x, real *y,
99 real *z, real *vx, real *vy, real *vz, real *ax, real *ay, real *az,
100 integer *level, idInteger *uid, integer *materialId, real *sml, integer *nnl,
101 integer *noi, real *e, real *dedt, real *cs, real *rho, real *p) {
102
103 this->numParticles = numParticles;
104 this->numNodes = numNodes;
105 this->mass = mass;
106 this->x = x;
107 this->y = y;
108 this->z = z;
109 this->vx = vx;
110 this->vy = vy;
111 this->vz = vz;
112 this->ax = ax;
113 this->ay = ay;
114 this->az = az;
115 this->level = level;
116 this->uid = uid;
117 this->materialId = materialId;
118 this->sml = sml;
119 this->nnl = nnl;
120 this->noi = noi;
121 this->e = e;
122 this->dedt = dedt;
123 this->cs = cs;
124 this->rho = rho;
125 this->p = p;
126}
127#endif
128
129#if DIM == 1
130 CUDA_CALLABLE_MEMBER void Particles::setGravity(real *g_ax) {
131 this->g_ax = g_ax;
132 }
133#elif DIM == 2
134 CUDA_CALLABLE_MEMBER void Particles::setGravity(real *g_ax, real *g_ay) {
135 this->g_ax = g_ax;
136 this->g_ay = g_ay;
137 }
138#else
139 CUDA_CALLABLE_MEMBER void Particles::setGravity(real *g_ax, real *g_ay, real *g_az) {
140 this->g_ax = g_ax;
141 this->g_ay = g_ay;
142 this->g_az = g_az;
143 }
144#endif
145
146#if DIM == 1
147 CUDA_CALLABLE_MEMBER void Particles::setLeapfrog(real *ax_old, real *g_ax_old) {
148 this->ax_old = ax_old;
149 this->g_ax_old = g_ax_old;
150 }
151#elif DIM == 2
152 CUDA_CALLABLE_MEMBER void Particles::setLeapfrog(real *ax_old, real *ay_old, real *g_ax_old, real *g_ay_old) {
153 this->ax_old = ax_old;
154 this->ay_old = ay_old;
155 this->g_ax_old = g_ax_old;
156 this->g_ay_old = g_ay_old;
157 }
158#else
159
160 CUDA_CALLABLE_MEMBER void Particles::setLeapfrog(real *ax_old, real *ay_old, real *az_old, real *g_ax_old, real *g_ay_old,
161 real *g_az_old) {
162 this->ax_old = ax_old;
163 this->ay_old = ay_old;
164 this->az_old = az_old;
165 this->g_ax_old = g_ax_old;
166 this->g_ay_old = g_ay_old;
167 this->g_az_old = g_az_old;
168 }
169#endif
170
171CUDA_CALLABLE_MEMBER void Particles::setU(real *u) {
172 this->u = u;
173}
174
175CUDA_CALLABLE_MEMBER void Particles::setNodeType(integer *nodeType) {
176 this->nodeType = nodeType;
177}
178
179CUDA_CALLABLE_MEMBER void Particles::setArtificialViscosity(real *muijmax) {
180 this->muijmax = muijmax;
181}
182
183#if BALSARA_SWITCH
184 CUDA_CALLABLE_MEMBER void Particles::setDivCurl(real *divv, real *curlv) {
185 this->divv = divv;
186 this->curlv = curlv;
187 }
188#endif
189
190//#if INTEGRATE_DENSITY
191 CUDA_CALLABLE_MEMBER void Particles::setIntegrateDensity(real *drhodt) {
192 this->drhodt = drhodt;
193 }
194//#endif
195#if VARIABLE_SML || INTEGRATE_SML
196 CUDA_CALLABLE_MEMBER void Particles::setVariableSML(real *dsmldt) {
197 this->dsmldt = dsmldt;
198 }
199#endif
200#if SML_CORRECTION
201 CUDA_CALLABLE_MEMBER void Particles::setSMLCorrection(real *sml_omega) {
202 this->sml_omega = sml_omega;
203 }
204#endif
205#if NAVIER_STOKES
206 CUDA_CALLABLE_MEMBER void Particles::setNavierStokes(real *Tshear, real *eta) {
207 this->Tshear = Tshear;
208 this->eta = eta;
209 }
210#endif
211#if SOLID
212 CUDA_CALLABLE_MEMBER void Particles::setSolid(real *S, real *dSdt, real *localStrain) {
213 this->S = S;
214 this->dSdt = dSdt;
215 this->localStrain = localStrain;
216 }
217#endif
218#if SOLID || NAVIER_STOKES
219 CUDA_CALLABLE_MEMBER void Particles::setSolidNavierStokes(real *sigma) {
220 this->sigma = sigma;
221 }
222#endif
223#if ARTIFICIAL_STRESS
224 CUDA_CALLABLE_MEMBER void Particles::setArtificialStress(real *R) {
225 this->R = R;
226 }
227#endif
228#if POROSITY
229CUDA_CALLABLE_MEMBER void Particles::setPorosity(real *pold, real *alpha_jutzi, real *alpha_jutzi_old, real *dalphadt,
230 real *dalphadp, real *dp, real *dalphadrho, real *f, real *delpdelrho,
231 real *delpdele, real *cs_old, real *alpha_epspor, real *dalpha_epspordt,
232 real *epsilon_v, real *depsilon_vdt) {
233 this->pold = pold;
234 this->alpha_jutzi = alpha_jutzi;
235 this->alpha_jutzi_old = alpha_jutzi_old;
236 this->dalphadt = dalphadt;
237 this->dalphadp = dalphadp;
238 this->dp = dp;
239 this->dalphadrho = dalphadrho;
240 this->f = f;
241 this->delpdelrho = delpdelrho;
242 this->delpdele = delpdele;
243 this->cs_old = cs_old;
244 this->alpha_epspor = alpha_epspor;
245 this->dalpha_epspordt = dalpha_epspordt;
246 this->epsilon_v = epsilon_v;
247 this->depsilon_vdt = depsilon_vdt;
248 }
249#endif
250#if ZERO_CONSISTENCY
251 CUDA_CALLABLE_MEMBER void Particles::setZeroConsistency(real *shepardCorrection) {
252 this->shepardCorrection = shepardCorrection;
253 }
254#endif
255#if LINEAR_CONSISTENCY
256 CUDA_CALLABLE_MEMBER void Particles::setLinearConsistency(real *tensorialCorrectionMatrix) {
257 this->tensorialCorrectionMatrix = tensorialCorrectionMatrix;
258 }
259#endif
260#if FRAGMENTATION
261 CUDA_CALLABLE_MEMBER void Particles::setFragmentation(real *d, real *damage_total, real *dddt, integer *numFlaws,
262 integer *maxNumFlaws, integer *numActiveFlaws, real *flaws) {
263 this->d = d;
264 this->damage_total = damage_total;
265 this->dddt = dddt;
266 this->numFlaws = numFlaws;
267 this->maxNumFlaws = maxNumFlaws;
268 this->numActiveFlaws = numActiveFlaws;
269 this->flaws = flaws;
270 }
271#if PALPHA_POROSITY
272 CUDA_CALLABLE_MEMBER void Particles::setPalphaPorosity(real *damage_porjutzi, real *ddamage_porjutzidt) {
273 this->damage_porjutzi = damage_porjutzi;
274 this->ddamage_porjutzidt = ddamage_porjutzidt;
275 }
276#endif
277#endif
278
279CUDA_CALLABLE_MEMBER void Particles::reset(integer index) {
280 level[index] = -1;
281 nodeType[index] = 0;
282 x[index] = 0.;
283#if DIM > 1
284 y[index] = 0.;
285#if DIM == 3
286 z[index] = 0.;
287#endif
288#endif
289 mass[index] = 0.;
290}
291
292CUDA_CALLABLE_MEMBER real Particles::distance(integer index_1, integer index_2) {
293
294 real dx;
295 if (x[index_1] < x[index_2]) {
296 dx = x[index_2] - x[index_1];
297 }
298 else if (x[index_1] > x[index_2]) {
299 dx = x[index_1] - x[index_2];
300 }
301 else {
302 dx = 0.f;
303 }
304#if DIM > 1
305 real dy;
306 if (y[index_1] < y[index_2]) {
307 dy = y[index_2] - y[index_1];
308 }
309 else if (y[index_1] > y[index_2]) {
310 dy = y[index_1] - y[index_2];
311 }
312 else {
313 dy = 0.f;
314 }
315#if DIM == 3
316 real dz;
317 if (z[index_1] < z[index_2]) {
318 dz = z[index_2] - z[index_1];
319 }
320 else if (z[index_1] > z[index_2]) {
321 dz = z[index_1] - z[index_2];
322 }
323 else {
324 dz = 0.f;
325 }
326#endif
327#endif
328
329#if DIM == 1
330 return sqrt(dx*dx);
331#elif DIM == 2
332 return sqrt(dx*dx + dy*dy);
333#else
334 return sqrt(dx*dx + dy*dy + dz*dz);
335#endif
336
337}
338
339CUDA_CALLABLE_MEMBER real Particles::weightedEntry(integer index, Entry::Name entry) {
340 switch (entry) {
341 case Entry::x: {
342 return x[index] * mass[index];
343 } break;
344#if DIM > 1
345 case Entry::y: {
346 return y[index] * mass[index];
347 } break;
348#if DIM == 3
349 case Entry::z: {
350 return z[index] * mass[index];
351 }
352#endif
353#endif
354 default: {
355 printf("Entry not available!\n");
356 return (real)0;
357 }
358 }
359}
360
361CUDA_CALLABLE_MEMBER Particles::~Particles() {
362
363}
364
365
366namespace ParticlesNS {
367
368 namespace Kernel {
369
370 /*__global__ void check4nans(Particles *particles, integer n) {
371 int bodyIndex = threadIdx.x + blockDim.x * blockIdx.x;
372 int stride = blockDim.x * gridDim.x;
373 int offset = 0;
374
375 while ((bodyIndex + offset) < n) {
376 if (std::isnan(particles->x[bodyIndex + offset]) || std::isnan(particles->mass[bodyIndex + offset])
377#if DIM > 1
378 || std::isnan(particles->y[bodyIndex + offset])
379#if DIM == 3
380 || std::isnan(particles->z[bodyIndex + offset])
381#endif
382#endif
383 ) {
384#if DIM == 1
385 printf("NAN for index: %i (%f) %f\n", bodyIndex + offset,
386 particles->x[bodyIndex + offset],
387 particles->mass[bodyIndex + offset]);
388#elif DIM == 2
389 printf("NAN for index: %i (%f, %f) %f\n", bodyIndex + offset,
390 particles->x[bodyIndex + offset],
391 particles->y[bodyIndex + offset],
392 particles->mass[bodyIndex + offset]);
393#else
394 printf("NAN for index: %i (%f, %f, %f) %f\n", bodyIndex + offset,
395 particles->x[bodyIndex + offset],
396 particles->y[bodyIndex + offset],
397 particles->z[bodyIndex + offset],
398 particles->mass[bodyIndex + offset]);
399#endif
400 assert(0);
401
402
403 }
404
405 if (particles->mass[bodyIndex + offset] == 0. || particles->sml[bodyIndex + offset] == 0.) {
406#if DIM == 1
407 printf("ATTENTION for index: %i (%f) %f\n", bodyIndex + offset,
408 particles->x[bodyIndex + offset],
409 particles->mass[bodyIndex + offset]);
410#elif DIM == 2
411 printf("ATTENTION for index: %i (%f, %f) %f\n", bodyIndex + offset,
412 particles->x[bodyIndex + offset],
413 particles->y[bodyIndex + offset],
414 particles->mass[bodyIndex + offset]);
415#else
416 printf("ATTENTION for index: %i (%e, %e, %e) %e sml = %e\n", bodyIndex + offset,
417 particles->x[bodyIndex + offset],
418 particles->y[bodyIndex + offset],
419 particles->z[bodyIndex + offset],
420 particles->mass[bodyIndex + offset],
421 particles->sml[bodyIndex + offset]);
422#endif
423 assert(0);
424 }
425
426
427
428 offset += stride;
429 }
430 }*/
431
432 __global__ void check4nans(Particles *particles, integer n) {
433 int bodyIndex = threadIdx.x + blockDim.x * blockIdx.x;
434 int stride = blockDim.x * gridDim.x;
435 int offset = 0;
436
437 while ((bodyIndex + offset) < n) {
438
439 if (particles->level[bodyIndex + offset] == -1) {
440 printf("level[%i] = %i!\n", bodyIndex + offset, particles->level[bodyIndex + offset]);
441 assert(0);
442 }
443 /*
444 if (std::isnan(particles->x[bodyIndex + offset]) || std::isnan(particles->mass[bodyIndex + offset])
445 #if DIM > 1
446 || std::isnan(particles->y[bodyIndex + offset])
447 #if DIM == 3
448 || std::isnan(particles->z[bodyIndex + offset])
449#endif
450#endif
451 ) {
452#if DIM == 1
453 printf("NAN for index: %i (%f) %f\n", bodyIndex + offset,
454 particles->x[bodyIndex + offset],
455 particles->mass[bodyIndex + offset]);
456#elif DIM == 2
457 printf("NAN for index: %i (%f, %f) %f\n", bodyIndex + offset,
458 particles->x[bodyIndex + offset],
459 particles->y[bodyIndex + offset],
460 particles->mass[bodyIndex + offset]);
461#else
462 printf("NAN for index: %i (%f, %f, %f) %f\n", bodyIndex + offset,
463 particles->x[bodyIndex + offset],
464 particles->y[bodyIndex + offset],
465 particles->z[bodyIndex + offset],
466 particles->mass[bodyIndex + offset]);
467#endif
468 assert(0);
469
470
471 }
472
473 if (particles->mass[bodyIndex + offset] == 0. || particles->sml[bodyIndex + offset] == 0.) {
474#if DIM == 1
475 printf("ATTENTION for index: %i (%f) %f\n", bodyIndex + offset,
476 particles->x[bodyIndex + offset],
477 particles->mass[bodyIndex + offset]);
478#elif DIM == 2
479 printf("ATTENTION for index: %i (%f, %f) %f\n", bodyIndex + offset,
480 particles->x[bodyIndex + offset],
481 particles->y[bodyIndex + offset],
482 particles->mass[bodyIndex + offset]);
483#else
484 printf("ATTENTION for index: %i (%e, %e, %e) %e sml = %e\n", bodyIndex + offset,
485 particles->x[bodyIndex + offset],
486 particles->y[bodyIndex + offset],
487 particles->z[bodyIndex + offset],
488 particles->mass[bodyIndex + offset],
489 particles->sml[bodyIndex + offset]);
490#endif
491 assert(0);
492 }
493
494 if (particles->x[bodyIndex + offset] > 1.e250
495 #if DIM > 1
496 || particles->y[bodyIndex + offset] > 1.e250
497 #if DIM == 3
498 || particles->z[bodyIndex + offset] > 1.e250
499#endif
500#endif
501 ) {
502 printf("HUGE entry for index: %i (%e, %e, %e) %e\n", bodyIndex + offset,
503 particles->x[bodyIndex + offset],
504 particles->y[bodyIndex + offset],
505 particles->z[bodyIndex + offset],
506 particles->mass[bodyIndex + offset]);
507 assert(0);
508 }
509
510 //if (bodyIndex + offset == 128121) {
511 // printf("INFO for index: %i (%e, %e, %e) %e sml = %e\n", bodyIndex + offset,
512 // particles->x[bodyIndex + offset],
513 // particles->y[bodyIndex + offset],
514 // particles->z[bodyIndex + offset],
515 // particles->mass[bodyIndex + offset],
516 // particles->sml[bodyIndex + offset]);
517 //}
518
519 //if (particles->sml[bodyIndex + offset] < 1.e-20) {
520 // printf("sml = %e\n", particles->sml[bodyIndex + offset]);
521 // assert(0);
522 //}
523 */
524 offset += stride;
525 }
526 }
527
528 __global__ void info(Particles *particles, integer n, integer m, integer k) {
529 int bodyIndex = threadIdx.x + blockDim.x * blockIdx.x;
530 int stride = blockDim.x * gridDim.x;
531 int offset = 0;
532
533#if DIM == 1
534 printf("not implemented yet for DIM == 1...\n");
535#elif DIM == 2
536 printf("not implemented yet for DIM == 2...\n");
537#else
538 while ((bodyIndex + offset) < n) {
539 //if ((bodyIndex + offset) % 100 == 0) {
540 printf("x[%i] = (%f, %f, %f) v = (%f, %f, %f) a = (%f, %f, %f) mass = %f\n", bodyIndex + offset,
541 particles->x[bodyIndex + offset], particles->y[bodyIndex + offset], particles->z[bodyIndex + offset],
542 particles->vx[bodyIndex + offset], particles->vy[bodyIndex + offset], particles->vz[bodyIndex + offset],
543 particles->ax[bodyIndex + offset], particles->ay[bodyIndex + offset], particles->az[bodyIndex + offset],
544 particles->mass[bodyIndex + offset]);
545 //}
546 offset += stride;
547 }
548
549 offset = m;
550 while ((bodyIndex + offset) < k && (bodyIndex + offset) > m) {
551 //if ((bodyIndex + offset) % 100 == 0) {
552 printf("x[%i] = (%f, %f, %f) mass = %f\n", bodyIndex + offset, particles->x[bodyIndex + offset],
553 particles->y[bodyIndex + offset],
554 particles->z[bodyIndex + offset],
555 particles->mass[bodyIndex + offset]);
556 //}
557 offset += stride;
558 }
559#endif
560
561 }
562
563 real Launch::check4nans(Particles *particles, integer n) {
564 ExecutionPolicy executionPolicy;
565 return cuda::launch(true, executionPolicy, ::ParticlesNS::Kernel::check4nans, particles, n);
566 }
567
568 real Launch::info(Particles *particles, integer n, integer m, integer k) {
569 ExecutionPolicy executionPolicy;
570 return cuda::launch(true, executionPolicy, ::ParticlesNS::Kernel::info, particles, n, m, k);
571 }
572
573#if DIM == 1
574
575 __global__ void set(Particles *particles, integer *numParticles, integer *numNodes, real *mass, real *x,
576 real *vx, real *ax, integer *level, idInteger *uid, integer *materialId,
577 real *sml, integer *nnl, integer *noi, real *e, real *dedt, real *cs, real *rho, real *p) {
578
579 particles->set(numParticles, numNodes, mass, x, vx, ax, level, uid, materialId, sml, nnl, noi, e,
580 dedt, cs, rho, p);
581
582 }
583
584 void Launch::set(Particles *particles, integer *numParticles, integer *numNodes, real *mass, real *x, real *vx,
585 real *ax, integer *level, idInteger *uid, integer *materialId, real *sml,
586 integer *nnl, integer *noi, real *e, real *dedt, real *cs, real *rho, real *p) {
587
588 ExecutionPolicy executionPolicy(1, 1);
589 cuda::launch(false, executionPolicy, ::ParticlesNS::Kernel::set, particles, numParticles, numNodes, mass,
590 x, vx, ax, level, uid, materialId, sml, nnl, noi, e, dedt, cs, rho, p);
591
592 }
593
594#elif DIM == 2
595
596 __global__ void set(Particles *particles, integer *numParticles, integer *numNodes, real *mass, real *x,
597 real *y, real *vx, real *vy, real *ax, real *ay, integer *level,
598 idInteger *uid, integer *materialId, real *sml, integer *nnl, integer *noi, real *e,
599 real *dedt, real *cs, real *rho, real *p) {
600
601 particles->set(numParticles, numNodes, mass, x, y, vx, vy, ax, ay, level, uid, materialId, sml,
602 nnl, noi, e, dedt, cs, rho, p);
603
604 }
605
606 void Launch::set(Particles *particles, integer *numParticles, integer *numNodes, real *mass, real *x, real *y,
607 real *vx, real *vy, real *ax, real *ay, integer *level, idInteger *uid,
608 integer *materialId, real *sml, integer *nnl, integer *noi, real *e, real *dedt, real *cs,
609 real *rho, real *p) {
610
611 ExecutionPolicy executionPolicy(1, 1);
612 cuda::launch(false, executionPolicy, ::ParticlesNS::Kernel::set, particles, numParticles, numNodes,
613 mass, x, y, vx, vy, ax, ay, level, uid, materialId, sml, nnl, noi, e, dedt,
614 cs, rho, p);
615
616 }
617
618
619#else
620
621 __global__ void set(Particles *particles, integer *numParticles, integer *numNodes, real *mass, real *x,
622 real *y, real *z, real *vx, real *vy, real *vz, real *ax, real *ay, real *az,
623 integer *level, idInteger *uid, integer *materialId, real *sml, integer *nnl,
624 integer *noi, real *e, real *dedt, real *cs, real *rho, real *p) {
625
626 particles->set(numParticles, numNodes, mass, x, y, z, vx, vy, vz, ax, ay, az, level, uid,
627 materialId, sml, nnl, noi, e, dedt, cs, rho, p);
628
629 }
630
631 void Launch::set(Particles *particles, integer *numParticles, integer *numNodes, real *mass, real *x, real *y,
632 real *z, real *vx, real *vy, real *vz, real *ax, real *ay, real *az, integer *level,
633 idInteger *uid, integer *materialId, real *sml, integer *nnl, integer *noi, real *e,
634 real *dedt, real *cs, real *rho, real *p) {
635
636 ExecutionPolicy executionPolicy(1, 1);
637 cuda::launch(false, executionPolicy, ::ParticlesNS::Kernel::set, particles, numParticles, numNodes,
638 mass, x, y, z, vx, vy, vz, ax, ay, az, level, uid, materialId, sml,
639 nnl, noi, e, dedt, cs, rho, p);
640 //setKernel<<<1, 1>>>(particles, count, mass, x, y, z, vx, vy, vz, ax, ay, az);
641
642 }
643
644#endif
645
646#if DIM == 1
647 __global__ void setGravity(Particles *particles, real *g_ax) {
648 particles->setGravity(g_ax);
649 }
650
651 namespace Launch {
652 void setGravity(Particles *particles, real *g_ax) {
653 ExecutionPolicy executionPolicy(1, 1);
654 cuda::launch(false, executionPolicy, ::ParticlesNS::Kernel::setGravity, particles, g_ax);
655 }
656 }
657
658#elif DIM == 2
659 __global__ void setGravity(Particles *particles, real *g_ax, real *g_ay) {
660 particles->setGravity(g_ax, g_ay);
661 }
662
663 namespace Launch {
664 void setGravity(Particles *particles, real *g_ax, real *g_ay) {
665 ExecutionPolicy executionPolicy(1, 1);
666 cuda::launch(false, executionPolicy, ::ParticlesNS::Kernel::setGravity, particles,
667 g_ax, g_ay);
668 }
669 }
670#else
671 __global__ void setGravity(Particles *particles, real *g_ax, real *g_ay, real *g_az) {
672 particles->setGravity(g_ax, g_ay, g_az);
673 }
674
675 namespace Launch {
676 void setGravity(Particles *particles, real *g_ax, real *g_ay, real *g_az) {
677 ExecutionPolicy executionPolicy(1, 1);
678 cuda::launch(false, executionPolicy, ::ParticlesNS::Kernel::setGravity, particles,
679 g_ax, g_ay, g_az);
680 }
681 }
682#endif
683
684#if DIM == 1
685 __global__ void setLeapfrog(Particles *particles, real *ax_old, real *g_ax_old) {
686 particles->setLeapfrog(ax_old, g_ax_old);
687 }
688
689 namespace Launch {
690 void setLeapfrog(Particles *particles, real *ax_old, real *g_ax_old) {
691 ExecutionPolicy executionPolicy(1, 1);
692 cuda::launch(false, executionPolicy, ::ParticlesNS::Kernel::setLeapfrog, particles,
693 ax_old, g_ax_old);
694 }
695 }
696
697#elif DIM == 2
698 __global__ void setLeapfrog(Particles *particles, real *ax_old, real *ay_old, real *g_ax_old, real *g_ay_old) {
699 particles->setLeapfrog(ax_old, ay_old, g_ax_old, g_ay_old);
700 }
701
702 namespace Launch {
703 void setLeapfrog(Particles *particles, real *ax_old, real *ay_old, real *g_ax_old, real *g_ay_old) {
704 ExecutionPolicy executionPolicy(1, 1);
705 cuda::launch(false, executionPolicy, ::ParticlesNS::Kernel::setLeapfrog, particles,
706 ax_old, ay_old, g_ax_old, g_ay_old);
707 }
708 }
709#else
710 __global__ void setLeapfrog(Particles *particles, real *ax_old, real *ay_old, real *az_old, real *g_ax_old,
711 real *g_ay_old, real *g_az_old) {
712 particles->setLeapfrog(ax_old, ay_old, az_old, g_ax_old, g_ay_old, g_az_old);
713 }
714
715 namespace Launch {
716 void setLeapfrog(Particles *particles, real *ax_old, real *ay_old, real *az_old, real *g_ax_old,
717 real *g_ay_old, real *g_az_old) {
718 ExecutionPolicy executionPolicy(1, 1);
719 cuda::launch(false, executionPolicy, ::ParticlesNS::Kernel::setLeapfrog, particles,
720 ax_old, ay_old, az_old, g_ax_old, g_ay_old, g_az_old);
721 }
722 }
723#endif
724
725 __global__ void setU(Particles *particles, real *u) {
726 particles->setU(u);
727 }
728
729 namespace Launch {
730 void setU(Particles *particles, real *u) {
731 ExecutionPolicy executionPolicy(1, 1);
732 cuda::launch(false, executionPolicy, ::ParticlesNS::Kernel::setU, particles, u);
733 }
734 }
735
736 __global__ void setNodeType(Particles *particles, integer *nodeType) {
737 particles->setNodeType(nodeType);
738 }
739 namespace Launch {
740 void setNodeType(Particles *particles, integer *nodeType) {
741 ExecutionPolicy executionPolicy(1, 1);
742 cuda::launch(false, executionPolicy, ::ParticlesNS::Kernel::setNodeType, particles, nodeType);
743 }
744 }
745
746 __global__ void setArtificialViscosity(Particles *particles, real *muijmax) {
747 particles->setArtificialViscosity(muijmax);
748 }
749
750 namespace Launch {
751 void setArtificialViscosity(Particles *particles, real *muijmax) {
752 ExecutionPolicy executionPolicy(1, 1);
753 cuda::launch(false, executionPolicy, ::ParticlesNS::Kernel::setArtificialViscosity, particles,
754 muijmax);
755 }
756 }
757
758#if BALSARA_SWITCH
759 __global__ void setDivCurl(Particles *particles, real *divv, real *curlv) {
760 particles->setDivCurl(divv, curlv);
761 }
762 namespace Launch {
763 void setDivCurl(Particles *particles, real *divv, real *curlv) {
764 ExecutionPolicy executionPolicy(1, 1);
765 cuda::launch(false, executionPolicy, ::ParticlesNS::Kernel::setDivCurl, particles, divv, curlv);
766 }
767 }
768#endif
769
770//#if INTEGRATE_DENSITY
771 __global__ void setIntegrateDensity(Particles *particles, real *drhodt) {
772 particles->setIntegrateDensity(drhodt);
773 }
774 void Launch::setIntegrateDensity(Particles *particles, real *drhodt) {
775 ExecutionPolicy executionPolicy(1, 1);
776 cuda::launch(false, executionPolicy, ::ParticlesNS::Kernel::setIntegrateDensity, particles, drhodt);
777 }
778//#endif
779#if VARIABLE_SML || INTEGRATE_SML
780 __global__ void setVariableSML(Particles *particles, real *dsmldt) {
781 particles->setVariableSML(dsmldt);
782 }
783 void Launch::setVariableSML(Particles *particles, real *dsmldt) {
784 ExecutionPolicy executionPolicy(1, 1);
785 cuda::launch(false, executionPolicy, ::ParticlesNS::Kernel::setVariableSML, particles, dsmldt);
786 }
787#endif
788#if SML_CORRECTION
789 __global__ void setSMLCorrection(Particles *particles, real *sml_omega) {
790 particles->setSMLCorrection(sml_omega);
791 }
792
793 void Launch::setSMLCorrection(Particles *particles, real *sml_omega) {
794 ExecutionPolicy executionPolicy(1, 1);
795 cuda::launch(false, executionPolicy, ::ParticlesNS::Kernel::setSMLCorrection, particles, sml_omega);
796 }
797#endif
798#if NAVIER_STOKES
799 __global__ void setNavierStokes(Particles *particles, real *Tshear, real *eta) {
800 particles->setNavierStokes(Tshear, eta);
801 }
802 namespace Launch {
803 void setNavierStokes(Particles *particles, real *Tshear, real *eta) {
804 ExecutionPolicy executionPolicy(1, 1);
805 cuda::launch(false, executionPolicy, ::ParticlesNS::Kernel::setNavierStokes, particles, Tshear, eta);
806 }
807 }
808#endif
809#if SOLID
810 __global__ void setSolid(Particles *particles, real *S, real *dSdt, real *localStrain) {
811 particles->setSolid(S, dSdt, localStrain);
812 }
813 void Launch::setSolid(Particles *particles, real *S, real *dSdt, real *localStrain) {
814 ExecutionPolicy executionPolicy(1, 1);
815 cuda::launch(false, executionPolicy, ::ParticlesNS::Kernel::setSolid, particles, S, dSdt, localStrain);
816 }
817#endif
818#if SOLID || NAVIER_STOKES
819 __global__ void setSolidNavierStokes(Particles *particles, real *sigma) {
820 particles->setSolidNavierStokes(sigma);
821 }
822 void Launch::setSolidNavierStokes(Particles *particles, real *sigma) {
823 ExecutionPolicy executionPolicy(1, 1);
824 cuda::launch(false, executionPolicy, ::ParticlesNS::Kernel::setSolidNavierStokes, particles, sigma);
825 }
826#endif
827#if ARTIFICIAL_STRESS
828 __global__ void setArtificialStress(Particles *particles, real *R) {
829 particles->setArtificialStress(R);
830 }
831 void Launch::setArtificialStress(Particles *particles, real *R) {
832 ExecutionPolicy executionPolicy(1, 1);
833 cuda::launch(false, executionPolicy, ::ParticlesNS::Kernel::setArtificialStress, particles, R);
834 }
835#endif
836#if POROSITY
837 __global__ void setPorosity(Particles *particles, real *pold, real *alpha_jutzi, real *alpha_jutzi_old, real *dalphadt,
838 real *dalphadp, real *dp, real *dalphadrho, real *f, real *delpdelrho,
839 real *delpdele, real *cs_old, real *alpha_epspor, real *dalpha_epspordt,
840 real *epsilon_v, real *depsilon_vdt) {
841 particles->setPorosity(pold, alpha_jutzi, alpha_jutzi_old, dalphadt,
842 dalphadp, dp, dalphadrho, f, delpdelrho,
843 delpdele, cs_old, alpha_epspor, dalpha_epspordt,
844 epsilon_v, depsilon_vdt);
845 }
846 void Launch::setPorosity(Particles *particles, real *pold, real *alpha_jutzi, real *alpha_jutzi_old, real *dalphadt,
847 real *dalphadp, real *dp, real *dalphadrho, real *f, real *delpdelrho,
848 real *delpdele, real *cs_old, real *alpha_epspor, real *dalpha_epspordt,
849 real *epsilon_v, real *depsilon_vdt) {
850 ExecutionPolicy executionPolicy(1, 1);
851 cuda::launch(false, executionPolicy, ::ParticlesNS::Kernel::setPorosity, particles, pold, alpha_jutzi,
852 alpha_jutzi_old, dalphadt, dalphadp, dp, dalphadrho, f, delpdelrho, delpdele, cs_old,
853 alpha_epspor, dalpha_epspordt, epsilon_v, depsilon_vdt);
854 }
855
856#endif
857#if ZERO_CONSISTENCY
858 __global__ void setZeroConsistency(Particles *particles, real *shepardCorrection) {
859 particles->setZeroConsistency(shepardCorrection);
860 }
861 void Launch::setZeroConsistency(Particles *particles, real *shepardCorrection) {
862 ExecutionPolicy executionPolicy(1, 1);
863 cuda::launch(false, executionPolicy, ::ParticlesNS::Kernel::setZeroConsistency, particles,
864 shepardCorrection);
865 }
866#endif
867#if LINEAR_CONSISTENCY
868 __global__ void setLinearConsistency(Particles *particles, real *tensorialCorrectionMatrix) {
869 particles->setLinearConsistency(tensorialCorrectionMatrix);
870 }
871 void Launch::setLinearConsistency(Particles *particles, real *tensorialCorrectionMatrix) {
872 ExecutionPolicy executionPolicy(1, 1);
873 cuda::launch(false, executionPolicy, ::ParticlesNS::Kernel::setLinearConsistency, particles,
874 tensorialCorrectionMatrix);
875 }
876#endif
877#if FRAGMENTATION
878 __global__ void setFragmentation(Particles *particles, real *d, real *damage_total, real *dddt, integer *numFlaws,
879 integer *maxNumFlaws, integer *numActiveFlaws, real *flaws) {
880 particles->setFragmentation(d, damage_total, dddt, numFlaws, maxNumFlaws, numActiveFlaws, flaws);
881 }
882 void Launch::setFragmentation(Particles *particles, real *d, real *damage_total, real *dddt, integer *numFlaws,
883 integer *maxNumFlaws, integer *numActiveFlaws, real *flaws) {
884 ExecutionPolicy executionPolicy(1, 1);
885 cuda::launch(false, executionPolicy, ::ParticlesNS::Kernel::setFragmentation, particles, d, damage_total,
886 dddt, numFlaws, maxNumFlaws, numActiveFlaws, flaws);
887 }
888#if PALPHA_POROSITY
889 __global__ void setPalphaPorosity(Particles *particles, real *damage_porjutzi, real *ddamage_porjutzidt) {
890 particles->setPalphaPorosity(damage_porjutzi, ddamage_porjutzidt);
891 }
892 void Launch::setPalphaPorosity(Particles *particles, real *damage_porjutzi, real *ddamage_porjutzidt) {
893 ExecutionPolicy executionPolicy(1, 1);
894 cuda::launch(false, executionPolicy, ::ParticlesNS::Kernel::setPalphaPorosity, particles, damage_porjutzi,
895 ddamage_porjutzidt);
896 }
897#endif
898#endif
899
900
901
902 __global__ void test(Particles *particles) {
903
904 int bodyIndex = threadIdx.x + blockDim.x * blockIdx.x;
905 int stride = blockDim.x * gridDim.x;
906 int offset = 0;
907
908 if (bodyIndex == 0) {
909 printf("device: numParticles = %i\n", *particles->numParticles);
910 }
911
912#if DIM == 1
913 printf("not implemented yet for DIM == 1...\n");
914#elif DIM == 2
915 printf("not implemented yet for DIM == 2...\n");
916#else
917 while ((bodyIndex + offset) < *particles->numParticles) {
918 if ((bodyIndex + offset) % 10000 == 0) {
919 printf("device: x[%i] = (%f, %f, %f)\n", bodyIndex + offset, particles->x[bodyIndex + offset],
920 particles->y[bodyIndex + offset],
921 particles->z[bodyIndex + offset]);
922 }
923 offset += stride;
924 }
925#endif
926
927
928 }
929
930 real Launch::test(Particles *particles, bool time) {
931 ExecutionPolicy executionPolicy;
932 return cuda::launch(time, executionPolicy, ::ParticlesNS::Kernel::test, particles);
933 //testKernel<<<256, 256>>>(particles);
934 }
935 }
936
937}
938
939CUDA_CALLABLE_MEMBER IntegratedParticles::IntegratedParticles() {
940
941}
942
943#if DIM == 1
944
945CUDA_CALLABLE_MEMBER IntegratedParticles::IntegratedParticles(idInteger *uid, real *rho, real *e, real *dedt, real *p,
946 real *cs, real *x, real *vx, real *ax) :
947 uid(uid), rho(rho), e(e), dedt(dedt), p(p), cs(cs),
948 x(x), vx(vx), ax(ax) {
949
950}
951
952CUDA_CALLABLE_MEMBER void IntegratedParticles::set(idInteger *uid, real *rho, real *e, real *dedt, real *p, real *cs,
953 real *x, real *vx, real *ax) {
954 this->uid = uid;
955 this->rho = rho;
956 this->e = e;
957 this->dedt = dedt;
958 this->p = p;
959 this->cs = cs;
960 this->x = x;
961 this->vx = vx;
962 this->ax = ax;
963}
964
965#elif DIM == 2
966
967CUDA_CALLABLE_MEMBER IntegratedParticles::IntegratedParticles(idInteger *uid, real *rho, real *e, real *dedt, real *p,
968 real *cs, real *x, real *y, real *vx, real *vy, real *ax,
969 real *ay) : uid(uid), rho(rho), e(e), dedt(dedt),
970 p(p), cs(cs), x(x), y(y), vx(vx), vy(vy), ax(ax),
971 ay(ay) {
972
973}
974
975CUDA_CALLABLE_MEMBER void IntegratedParticles::set(idInteger *uid, real *rho, real *e, real *dedt, real *p,
976 real *cs, real *x, real *y, real *vx, real *vy, real *ax,
977 real *ay) {
978 this->uid = uid;
979 this->rho = rho;
980 this->e = e;
981 this->dedt = dedt;
982 this->p = p;
983 this->cs = cs;
984 this->x = x;
985 this->y = y;
986 this->vx = vx;
987 this->vy = vy;
988 this->ax = ax;
989 this->ay = ay;
990}
991
992#else
993
994CUDA_CALLABLE_MEMBER IntegratedParticles::IntegratedParticles(idInteger *uid, real *rho, real *e, real *dedt, real *p,
995 real *cs, real *x, real *y, real *z, real *vx, real *vy,
996 real *vz, real *ax, real *ay, real *az) :
997 uid(uid), rho(rho), e(e), dedt(dedt), p(p), cs(cs),
998 x(x), y(y), z(z), vx(vx), vy(vy), vz(vz), ax(ax), ay(ay),
999 az(az) {
1000
1001}
1002
1003CUDA_CALLABLE_MEMBER void IntegratedParticles::set(idInteger *uid, real *rho, real *e, real *dedt, real *p,
1004 real *cs, real *x, real *y, real *z, real *vx, real *vy,
1005 real *vz, real *ax, real *ay, real *az) {
1006 this->uid = uid;
1007 this->rho = rho;
1008 this->e = e;
1009 this->dedt = dedt;
1010 this->p = p;
1011 this->cs = cs;
1012 this->x = x;
1013 this->y = y;
1014 this->z = z;
1015 this->vx = vx;
1016 this->vy = vy;
1017 this->vz = vz;
1018 this->ax = ax;
1019 this->ay = ay;
1020 this->az = az;
1021}
1022
1023#endif
1024
1025//CUDA_CALLABLE_MEMBER void IntegratedParticles::setLevel(integer *level) {
1026// this->level = level;
1027//}
1028
1029//CUDA_CALLABLE_MEMBER void IntegratedParticles::setNodeType(integer *nodeType) {
1030// this->nodeType = nodeType;
1031//}
1032
1033CUDA_CALLABLE_MEMBER void IntegratedParticles::setSML(real *sml) {
1034 this->sml = sml;
1035}
1036
1037//#if INTEGRATE_DENSITY
1038CUDA_CALLABLE_MEMBER void IntegratedParticles::setIntegrateDensity(real *drhodt) {
1039 this->drhodt = drhodt;
1040}
1041//#endif
1042
1043#if VARIABLE_SML || INTEGRATE_SML
1044CUDA_CALLABLE_MEMBER void IntegratedParticles::setIntegrateSML(real *dsmldt) {
1045 this->dsmldt = dsmldt;
1046}
1047#endif
1048
1049CUDA_CALLABLE_MEMBER void IntegratedParticles::reset(integer index) {
1050
1051 //TODO: what to reset?
1052 uid[index] = 0;
1053 x[index] = 0.;
1054 vx[index] = 0.;
1055 ax[index] = 0.;
1056#if DIM > 1
1057 y[index] = 0.;
1058 vy[index] = 0.;
1059 ay[index] = 0.;
1060#if DIM == 3
1061 z[index] = 0.;
1062 vz[index] = 0.;
1063 az[index] = 0.;
1064#endif
1065#endif
1066
1067#if SPH_SIM
1068 e[index] = 0.;
1069 dedt[index] = 0.;
1070 p[index] = 0.;
1071 rho[index] = 0.;
1072 drhodt[index] = 0.;
1073 cs[index] = 0.;
1074 sml[index] = 0.;
1075#endif
1076}
1077
1078CUDA_CALLABLE_MEMBER IntegratedParticles::~IntegratedParticles() {
1079
1080}
1081
1082namespace IntegratedParticlesNS {
1083
1084 namespace Kernel {
1085
1086#if DIM == 1
1087
1088 __global__ void set(IntegratedParticles *integratedParticles, idInteger *uid, real *rho, real *e, real *dedt,
1089 real *p, real *cs, real *x, real *vx, real *ax) {
1090 integratedParticles->set(uid, rho, e, dedt, p, cs, x, vx, ax);
1091 }
1092
1093 void Launch::set(IntegratedParticles *integratedParticles, idInteger *uid, real *rho, real *e, real *dedt,
1094 real *p, real *cs, real *x, real *vx, real *ax) {
1095 ExecutionPolicy executionPolicy(1, 1);
1096 cuda::launch(false, executionPolicy, ::IntegratedParticlesNS::Kernel::set, integratedParticles, uid,
1097 rho, e, dedt, p, cs, x, vx, ax);
1098 }
1099
1100#elif DIM == 2
1101
1102 __global__ void set(IntegratedParticles *integratedParticles, idInteger *uid, real *rho, real *e, real *dedt,
1103 real *p, real *cs, real *x, real *y, real *vx, real *vy, real *ax, real *ay) {
1104 integratedParticles->set(uid, rho, e, dedt, p, cs, x, y, vx, vy, ax, ay);
1105 }
1106
1107 void Launch::set(IntegratedParticles *integratedParticles, idInteger *uid, real *rho, real *e, real *dedt,
1108 real *p, real *cs, real *x, real *y, real *vx, real *vy, real *ax, real *ay) {
1109 ExecutionPolicy executionPolicy(1, 1);
1110 cuda::launch(false, executionPolicy, ::IntegratedParticlesNS::Kernel::set, integratedParticles, uid,
1111 rho, e, dedt, p, cs, x, y, vx, vy, ax, ay);
1112 }
1113
1114#else
1115
1116 __global__ void set(IntegratedParticles *integratedParticles, idInteger *uid, real *rho, real *e, real *dedt,
1117 real *p, real *cs, real *x, real *y, real *z, real *vx, real *vy, real *vz, real *ax,
1118 real *ay, real *az) {
1119 integratedParticles->set(uid, rho, e, dedt, p, cs, x, y, z, vx, vy, vz, ax, ay, az);
1120 }
1121
1122 void Launch::set(IntegratedParticles *integratedParticles, idInteger *uid, real *rho, real *e, real *dedt,
1123 real *p, real *cs, real *x, real *y, real *z, real *vx, real *vy, real *vz, real *ax,
1124 real *ay, real *az) {
1125 ExecutionPolicy executionPolicy(1, 1);
1126 cuda::launch(false, executionPolicy, ::IntegratedParticlesNS::Kernel::set, integratedParticles, uid,
1127 rho, e, dedt, p, cs, x, y, z, vx, vy, vz, ax, ay, az);
1128 }
1129
1130#endif
1131
1132 //__global__ void setLevel(IntegratedParticles *integratedParticles, integer *level) {
1133 // integratedParticles->setLevel(level);
1134 //}
1135 //namespace Launch {
1136 // void setLevel(IntegratedParticles *integratedParticles, integer *level) {
1137 // ExecutionPolicy executionPolicy(1, 1);
1138 // cuda::launch(false, executionPolicy, ::IntegratedParticlesNS::Kernel::setLevel,
1139 // integratedParticles, level);
1140 // }
1141 //}
1142
1143 //__global__ void setNodeType(IntegratedParticles *integratedParticles, integer *nodeType) {
1144 // integratedParticles->setNodeType(nodeType);
1145 //}
1146 //namespace Launch {
1147 // void setNodeType(IntegratedParticles *integratedParticles, integer *nodeType) {
1148 // ExecutionPolicy executionPolicy(1, 1);
1149 // cuda::launch(false, executionPolicy, ::IntegratedParticlesNS::Kernel::setLevel,
1150 // integratedParticles, level);
1151 // }
1152 //}
1153
1154 __global__ void setSML(IntegratedParticles *integratedParticles, real *sml) {
1155 integratedParticles->setSML(sml);
1156 }
1157
1158 namespace Launch {
1159 void setSML(IntegratedParticles *integratedParticles, real *sml) {
1160 ExecutionPolicy executionPolicy(1, 1);
1161 cuda::launch(false, executionPolicy, ::IntegratedParticlesNS::Kernel::setSML,
1162 integratedParticles, sml);
1163 }
1164 }
1165
1166//#if INTEGRATE_DENSITY
1167 __global__ void setIntegrateDensity(IntegratedParticles *integratedParticles, real *drhodt) {
1168 integratedParticles->setIntegrateDensity(drhodt);
1169 }
1170
1171 namespace Launch {
1172
1173 void setIntegrateDensity(IntegratedParticles *integratedParticles, real *drhodt) {
1174 ExecutionPolicy executionPolicy(1, 1);
1175 cuda::launch(false, executionPolicy, ::IntegratedParticlesNS::Kernel::setIntegrateDensity,
1176 integratedParticles, drhodt);
1177 }
1178
1179 }
1180//#endif
1181
1182#if VARIABLE_SML || INTEGRATE_SML
1183 __global__ void setIntegrateSML(IntegratedParticles *integratedParticles, real *dsmldt) {
1184 integratedParticles->setIntegrateSML(dsmldt);
1185 }
1186
1187 namespace Launch {
1188 void setIntegrateSML(IntegratedParticles *integratedParticles, real *dsmldt) {
1189 ExecutionPolicy executionPolicy(1, 1);
1190 cuda::launch(false, executionPolicy, ::IntegratedParticlesNS::Kernel::setIntegrateSML,
1191 integratedParticles, dsmldt);
1192 }
1193 }
1194#endif
1195 }
1196}
ExecutionPolicy
Execution policy/instruction for CUDA kernel execution.
Definition: cuda_launcher.cuh:33
IntegratedParticles
Definition: particles.cuh:979
IntegratedParticles::cs
real * cs
Definition: particles.cuh:1011
IntegratedParticles::y
real * y
Definition: particles.cuh:992
IntegratedParticles::ay
real * ay
Definition: particles.cuh:994
IntegratedParticles::rho
real * rho
Definition: particles.cuh:1007
IntegratedParticles::setSML
CUDA_CALLABLE_MEMBER void setSML(real *sml)
Definition: particles.cu:1033
IntegratedParticles::ax
real * ax
Definition: particles.cuh:989
IntegratedParticles::drhodt
real * drhodt
Definition: particles.cuh:1016
IntegratedParticles::vx
real * vx
Definition: particles.cuh:988
IntegratedParticles::IntegratedParticles
CUDA_CALLABLE_MEMBER IntegratedParticles()
Definition: particles.cu:939
IntegratedParticles::p
real * p
Definition: particles.cuh:1010
IntegratedParticles::~IntegratedParticles
CUDA_CALLABLE_MEMBER ~IntegratedParticles()
Definition: particles.cu:1078
IntegratedParticles::reset
CUDA_CALLABLE_MEMBER void reset(integer index)
Definition: particles.cu:1049
IntegratedParticles::uid
idInteger * uid
unique identifier
Definition: particles.cuh:984
IntegratedParticles::az
real * az
Definition: particles.cuh:999
IntegratedParticles::vz
real * vz
Definition: particles.cuh:998
IntegratedParticles::x
real * x
Definition: particles.cuh:987
IntegratedParticles::dedt
real * dedt
Definition: particles.cuh:1009
IntegratedParticles::sml
real * sml
Definition: particles.cuh:1013
IntegratedParticles::z
real * z
Definition: particles.cuh:997
IntegratedParticles::e
real * e
Definition: particles.cuh:1008
IntegratedParticles::vy
real * vy
Definition: particles.cuh:993
IntegratedParticles::setIntegrateDensity
CUDA_CALLABLE_MEMBER void setIntegrateDensity(real *drhodt)
Definition: particles.cu:1038
IntegratedParticles::set
CUDA_CALLABLE_MEMBER void set(idInteger *uid, real *rho, real *e, real *dedt, real *p, real *cs, real *x, real *y, real *z, real *vx, real *vy, real *vz, real *ax, real *ay, real *az)
Definition: particles.cu:1003
Particles
Particle(s) class based on SoA (Structur of Arrays).
Definition: particles.cuh:50
Particles::g_az_old
real * g_az_old
Definition: particles.cuh:97
Particles::noi
integer * noi
(pointer to) number of interactions (array)
Definition: particles.cuh:118
Particles::e
real * e
(pointer to) internal energy (array)
Definition: particles.cuh:121
Particles::ax_old
real * ax_old
Definition: particles.cuh:67
Particles::weightedEntry
CUDA_CALLABLE_MEMBER real weightedEntry(integer index, Entry::Name entry)
Definition: particles.cu:339
Particles::materialId
integer * materialId
(pointer to) material identifier (array)
Definition: particles.cuh:111
Particles::x
real * x
(pointer to) x position (array)
Definition: particles.cuh:62
Particles::rho
real * rho
(pointer to) density (array)
Definition: particles.cuh:133
Particles::level
integer * level
(pointer to) level of the (pseudo-)particles
Definition: particles.cuh:106
Particles::uid
idInteger * uid
(pointer to) unique identifier (array)
Definition: particles.cuh:109
Particles::drhodt
real * drhodt
(pointer to) time derivative of density (array)
Definition: particles.cuh:148
Particles::set
CUDA_CALLABLE_MEMBER void set(integer *numParticles, integer *numNodes, real *mass, real *x, real *y, real *z, real *vx, real *vy, real *vz, real *ax, real *ay, real *az, integer *level, idInteger *uid, integer *materialId, real *sml, integer *nnl, integer *noi, real *e, real *dedt, real *cs, real *rho, real *p)
Setter (DIM = 3) assigning variables to pointer members.
Definition: particles.cu:98
Particles::g_ax_old
real * g_ax_old
Definition: particles.cuh:89
Particles::g_ay
real * g_ay
(pointer to) y gravitational acceleration (array)
Definition: particles.cuh:92
Particles::ay
real * ay
(pointer to) y acceleration (array)
Definition: particles.cuh:74
Particles::nodeType
integer * nodeType
(pointer to) node type
Definition: particles.cuh:103
Particles::nnl
integer * nnl
(pointer to) near(est) neighbor list (array)
Definition: particles.cuh:116
Particles::az
real * az
(pointer to) z acceleration (array)
Definition: particles.cuh:82
Particles::ay_old
real * ay_old
Definition: particles.cuh:75
Particles::p
real * p
(pointer to) pressure (array)
Definition: particles.cuh:135
Particles::cs
real * cs
(pointer to) sound of speed (array)
Definition: particles.cuh:129
Particles::az_old
real * az_old
Definition: particles.cuh:83
Particles::y
real * y
(pointer to) y position (array)
Definition: particles.cuh:70
Particles::g_ay_old
real * g_ay_old
Definition: particles.cuh:93
Particles::ax
real * ax
(pointer to) x acceleration (array)
Definition: particles.cuh:66
Particles::reset
CUDA_CALLABLE_MEMBER void reset(integer index)
Reset (specific) entries.
Definition: particles.cu:279
Particles::u
real * u
energy (kinetic + gravitational for now)
Definition: particles.cuh:126
Particles::g_ax
real * g_ax
(pointer to) x gravitational acceleration (array)
Definition: particles.cuh:88
Particles::Particles
CUDA_CALLABLE_MEMBER Particles()
Definition: particles.cu:4
Particles::~Particles
CUDA_CALLABLE_MEMBER ~Particles()
Definition: particles.cu:361
Particles::sml
real * sml
(pointer to) smoothing length (array)
Definition: particles.cuh:113
Particles::dedt
real * dedt
(pointer to) time derivative of internal energy (array)
Definition: particles.cuh:123
Particles::mass
real * mass
(pointer to) mass (array)
Definition: particles.cuh:60
Particles::z
real * z
(pointer to) z position (array)
Definition: particles.cuh:78
Particles::numParticles
integer * numParticles
number of particles
Definition: particles.cuh:55
Particles::numNodes
integer * numNodes
number of nodes
Definition: particles.cuh:57
Particles::vz
real * vz
(pointer to) z velocity (array)
Definition: particles.cuh:80
Particles::distance
CUDA_CALLABLE_MEMBER real distance(integer index_1, integer index_2)
Distance of two particles.
Definition: particles.cu:292
Particles::setLeapfrog
CUDA_CALLABLE_MEMBER void setLeapfrog(real *ax_old, real *ay_old, real *az_old, real *g_ax_old, real *g_ay_old, real *g_az_old)
Definition: particles.cu:160
Particles::muijmax
real * muijmax
(pointer) to max(mu_ij) (array) needed for artificial viscosity and determining timestp
Definition: particles.cuh:138
Particles::setGravity
CUDA_CALLABLE_MEMBER void setGravity(real *g_ax, real *g_ay, real *g_az)
Constructor (DIM = 3) assigning gravitational acceleration to member variables.
Definition: particles.cu:139
Particles::g_az
real * g_az
(pointer to) z gravitational acceleration (array)
Definition: particles.cuh:96
Particles::vx
real * vx
(pointer to) x velocity (array)
Definition: particles.cuh:64
Particles::setNodeType
CUDA_CALLABLE_MEMBER void setNodeType(integer *nodeType)
Setter for node type.
Definition: particles.cu:175
Particles::vy
real * vy
(pointer to) y velocity (array)
Definition: particles.cuh:72
Particles::setArtificialViscosity
CUDA_CALLABLE_MEMBER void setArtificialViscosity(real *muijmax)
Setter for artificial viscosity (entry).
Definition: particles.cu:179
Particles::setIntegrateDensity
CUDA_CALLABLE_MEMBER void setIntegrateDensity(real *drhodt)
Setter in dependence of INTEGRATE_DENSITY.
Definition: particles.cu:191
Particles::setU
CUDA_CALLABLE_MEMBER void setU(real *u)
Setter for energy.
Definition: particles.cu:171
CUDA_CALLABLE_MEMBER
#define CUDA_CALLABLE_MEMBER
Definition: cuda_utilities.cuh:30
IntegratedParticlesNS::Kernel::Launch::setSML
void setSML(IntegratedParticles *integratedParticles, real *sml)
Definition: particles.cu:1159
IntegratedParticlesNS::Kernel::Launch::setIntegrateDensity
void setIntegrateDensity(IntegratedParticles *integratedParticles, real *drhodt)
Definition: particles.cu:1173
IntegratedParticlesNS::Kernel::Launch::set
void set(IntegratedParticles *integratedParticles, idInteger *uid, real *rho, real *e, real *dedt, real *p, real *cs, real *x, real *y, real *z, real *vx, real *vy, real *vz, real *ax, real *ay, real *az)
Definition: particles.cu:1122
IntegratedParticlesNS::Kernel::set
__global__ void set(IntegratedParticles *integratedParticles, idInteger *uid, real *rho, real *e, real *dedt, real *p, real *cs, real *x, real *y, real *z, real *vx, real *vy, real *vz, real *ax, real *ay, real *az)
Definition: particles.cu:1116
IntegratedParticlesNS::Kernel::setSML
__global__ void setSML(IntegratedParticles *integratedParticles, real *sml)
Definition: particles.cu:1154
IntegratedParticlesNS::Kernel::setIntegrateDensity
__global__ void setIntegrateDensity(IntegratedParticles *integratedParticles, real *drhodt)
Definition: particles.cu:1167
IntegratedParticlesNS
Definition: particles.cuh:1083
Kernel
Definition: device_rhs.cuh:7
MaterialNS::Kernel::info
__global__ void info(Material *material)
Debug kernel giving information about material(s).
Definition: material.cu:24
ParticlesNS::Kernel::Launch::setU
void setU(Particles *particles, real *u)
Definition: particles.cu:730
ParticlesNS::Kernel::Launch::setIntegrateDensity
void setIntegrateDensity(Particles *particles, real *drhodt)
Definition: particles.cu:774
ParticlesNS::Kernel::Launch::setGravity
void setGravity(Particles *particles, real *g_ax, real *g_ay, real *g_az)
Definition: particles.cu:676
ParticlesNS::Kernel::Launch::setLeapfrog
void setLeapfrog(Particles *particles, real *ax_old, real *ay_old, real *az_old, real *g_ax_old, real *g_ay_old, real *g_az_old)
Definition: particles.cu:716
ParticlesNS::Kernel::Launch::set
void set(Particles *particles, integer *numParticles, integer *numNodes, real *mass, real *x, real *y, real *z, real *vx, real *vy, real *vz, real *ax, real *ay, real *az, integer *level, idInteger *uid, integer *materialId, real *sml, integer *nnl, integer *noi, real *e, real *dedt, real *cs, real *rho, real *p)
Definition: particles.cu:631
ParticlesNS::Kernel::Launch::test
real test(Particles *particles, bool time=false)
Definition: particles.cu:930
ParticlesNS::Kernel::Launch::setNodeType
void setNodeType(Particles *particles, integer *nodeType)
Definition: particles.cu:740
ParticlesNS::Kernel::Launch::info
real info(Particles *particles, integer n, integer m, integer k)
Definition: particles.cu:568
ParticlesNS::Kernel::Launch::check4nans
real check4nans(Particles *particles, integer n)
Definition: particles.cu:563
ParticlesNS::Kernel::Launch::setArtificialViscosity
void setArtificialViscosity(Particles *particles, real *muijmax)
Definition: particles.cu:751
ParticlesNS::Kernel::set
__global__ void set(Particles *particles, integer *numParticles, integer *numNodes, real *mass, real *x, real *y, real *z, real *vx, real *vy, real *vz, real *ax, real *ay, real *az, integer *level, idInteger *uid, integer *materialId, real *sml, integer *nnl, integer *noi, real *e, real *dedt, real *cs, real *rho, real *p)
Definition: particles.cu:621
ParticlesNS::Kernel::setArtificialViscosity
__global__ void setArtificialViscosity(Particles *particles, real *muijmax)
Definition: particles.cu:746
ParticlesNS::Kernel::setNodeType
__global__ void setNodeType(Particles *particles, integer *nodeType)
Definition: particles.cu:736
ParticlesNS::Kernel::setIntegrateDensity
__global__ void setIntegrateDensity(Particles *particles, real *drhodt)
Definition: particles.cu:771
ParticlesNS::Kernel::setGravity
__global__ void setGravity(Particles *particles, real *g_ax, real *g_ay, real *g_az)
Definition: particles.cu:671
ParticlesNS::Kernel::info
__global__ void info(Particles *particles, integer n, integer m, integer k)
Definition: particles.cu:528
ParticlesNS::Kernel::setLeapfrog
__global__ void setLeapfrog(Particles *particles, real *ax_old, real *ay_old, real *az_old, real *g_ax_old, real *g_ay_old, real *g_az_old)
Definition: particles.cu:710
ParticlesNS::Kernel::check4nans
__global__ void check4nans(Particles *particles, integer n)
Definition: particles.cu:432
ParticlesNS::Kernel::setU
__global__ void setU(Particles *particles, real *u)
Definition: particles.cu:725
ParticlesNS::Kernel::test
__global__ void test(Particles *particles)
Definition: particles.cu:902
ParticlesNS
Particle class related functions and kernels.
Definition: particles.cuh:552
ProfilerIds::numParticles
const char *const numParticles
Definition: h5profiler.h:29
cuda::math::sqrt
__device__ real sqrt(real a)
Square root of a floating point value.
Definition: cuda_utilities.cu:456
cuda::set
void set(T *d_var, T val, std::size_t count=1)
Set device memory to a specific value.
Definition: cuda_runtime.h:56
cuda::launch
real launch(bool timeKernel, const ExecutionPolicy &policy, void(*f)(Arguments...), Arguments... args)
CUDA execution wrapper function.
Definition: cuda_launcher.cuh:114
real
double real
Definition: parameter.h:15
integer
int integer
Definition: parameter.h:17
idInteger
int idInteger
Definition: parameter.h:19
Entry::Name
Name
Definition: parameter.h:255
Entry::y
@ y
Definition: parameter.h:258
Entry::x
@ x
Definition: parameter.h:256
Entry::z
@ z
Definition: parameter.h:260

milupHPC - src/particles.cu Source File
Generated on Wed Aug 31 2022 12:16:53 by Doxygen 1.9.3