milupHPC documentation
  • include
  • sph
sph.cuh
Go to the documentation of this file.
1
15#ifndef MILUPHPC_SPH_CUH
16#define MILUPHPC_SPH_CUH
17
18#include "../particles.cuh"
19#include "../subdomain_key_tree/subdomain.cuh"
20#include "../parameter.h"
21#include "../helper.cuh"
22#include "../materials/material.cuh"
23#include <float.h>
24
25#include <boost/mpi.hpp>
26#include <assert.h>
27
29namespace SPH {
30
34 void exchangeParticleEntry(SubDomainKeyTree *subDomainKeyTree, real *entry, real *toSend, integer *sendLengths,
35 integer *receiveLengths, integer numParticlesLocal);
36
38 namespace Kernel {
39
61 __global__ void fixedRadiusNN_bruteForce(Tree *tree, Particles *particles, integer *interactions,
62 integer numParticlesLocal, integer numParticles, integer numNodes);
63
93 __global__ void fixedRadiusNN(Tree *tree, Particles *particles, integer *interactions, real radius,
94 integer numParticlesLocal, integer numParticles, integer numNodes);
95
120 __global__ void fixedRadiusNN_withinBox(Tree *tree, Particles *particles, integer *interactions,
121 integer numParticlesLocal, integer numParticles, integer numNodes);
122
137 __global__ void fixedRadiusNN_sharedMemory(Tree *tree, Particles *particles, integer *interactions,
138 integer numParticlesLocal, integer numParticles, integer numNodes);
139
156 __global__ void fixedRadiusNN_variableSML(Material *materials, Tree *tree, Particles *particles, integer *interactions,
157 integer numParticlesLocal, integer numParticles,
158 integer numNodes);
159
173 __device__ void redoNeighborSearch(Tree *tree, Particles *particles, int particleId,
174 int *interactions, real radius, integer numParticles, integer numNodes);
175
190 __global__ void compTheta(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles,
191 DomainList *lowestDomainList, Curve::Type curveType);
192
211 __global__ void symbolicForce(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles,
212 DomainList *lowestDomainList, integer *sendIndices, real searchRadius,
213 integer n, integer m, integer relevantIndex,
214 Curve::Type curveType);
215
216 __global__ void symbolicForce_test(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles,
217 DomainList *lowestDomainList, integer *sendIndices, real searchRadius,
218 integer n, integer m, integer relevantProc, integer relevantIndicesCounter,
219 integer relevantIndexOld, Curve::Type curveType);
220
221 __global__ void symbolicForce_test2(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles,
222 DomainList *domainList, integer *sendIndices, real searchRadius,
223 integer n, integer m, integer relevantProc, integer relevantIndicesCounter,
224 Curve::Type curveType);
225
241 __global__ void collectSendIndices(Tree *tree, Particles *particles, integer *sendIndices,
242 integer *particles2Send, integer *particlesCount,
243 integer n, integer length, Curve::Type curveType);
244
245 __global__ void collectSendIndices_test2(Tree *tree, Particles *particles, integer *sendIndices,
246 integer *particles2Send, integer *particlesCount,
247 integer numParticlesLocal, integer numParticles,
248 integer treeIndex, int currentProc, Curve::Type curveType);
249
253 __global__ void particles2Send(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles,
254 DomainList *domainList, DomainList *lowestDomainList, integer maxLevel,
255 integer *toSend, integer *sendCount, integer *alreadyInserted,
256 integer insertOffset,
257 integer numParticlesLocal, integer numParticles, integer numNodes, real radius,
258 Curve::Type curveType = Curve::lebesgue);
259
263 __global__ void collectSendIndicesBackup(integer *toSend, integer *toSendCollected, integer count);
264
268 __global__ void collectSendEntriesBackup(SubDomainKeyTree *subDomainKeyTree, real *entry, real *toSend,
269 integer *sendIndices, integer *sendCount, integer totalSendCount,
270 integer insertOffset);
271
286 __global__ void insertReceivedParticles(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles,
287 DomainList *domainList, DomainList *lowestDomainList, int n, int m);
288
289
293 __global__ void calculateCentersOfMass(Tree *tree, Particles *particles, integer level);
294
310 __global__ void determineSearchRadii(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles,
311 DomainList *domainList, DomainList *lowestDomainList, real *searchRadii,
312 int n, int m, Curve::Type curveType);
313
324 __global__ void info(Tree *tree, Particles *particles, Helper *helper,
325 integer numParticlesLocal, integer numParticles, integer numNodes);
326
328 namespace Launch {
329
335 real fixedRadiusNN_bruteForce(Tree *tree, Particles *particles, integer *interactions,
336 integer numParticlesLocal, integer numParticles, integer numNodes);
337
343 real fixedRadiusNN(Tree *tree, Particles *particles, integer *interactions, real radius,
344 integer numParticlesLocal, integer numParticles, integer numNodes);
345
351 real fixedRadiusNN_sharedMemory(Tree *tree, Particles *particles, integer *interactions,
352 integer numParticlesLocal, integer numParticles, integer numNodes);
353
359 real fixedRadiusNN_withinBox(Tree *tree, Particles *particles, integer *interactions,
360 integer numParticlesLocal, integer numParticles, integer numNodes);
361
367 real fixedRadiusNN_variableSML(Material *materials, Tree *tree, Particles *particles, integer *interactions,
368 integer numParticlesLocal, integer numParticles,
369 integer numNodes);
370
376 real compTheta(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles,
377 DomainList *lowestDomainList, Curve::Type curveType);
378
384 real symbolicForce(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles,
385 DomainList *lowestDomainList, integer *sendIndices, real searchRadius,
386 integer n, integer m, integer relevantIndex,
387 Curve::Type curveType);
388
389 real symbolicForce_test(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles,
390 DomainList *lowestDomainList, integer *sendIndices, real searchRadius,
391 integer n, integer m, integer relevantProc, integer relevantIndicesCounter,
392 integer relevantIndexOld, Curve::Type curveType);
393
394 real symbolicForce_test2(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles,
395 DomainList *domainList, integer *sendIndices, real searchRadius,
396 integer n, integer m, integer relevantProc, integer relevantIndicesCounter,
397 Curve::Type curveType);
398
404 real collectSendIndices(Tree *tree, Particles *particles, integer *sendIndices,
405 integer *particles2Send, integer *particlesCount,
406 integer n, integer length, Curve::Type curveType);
407
408 real collectSendIndices_test2(Tree *tree, Particles *particles, integer *sendIndices,
409 integer *particles2Send, integer *particlesCount,
410 integer numParticlesLocal, integer numParticles,
411 integer treeIndex, int currentProc, Curve::Type curveType);
412
418 real particles2Send(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles,
419 DomainList *domainList, DomainList *lowestDomainList, integer maxLevel,
420 integer *toSend, integer *sendCount, integer *alreadyInserted,
421 integer insertOffset,
422 integer numParticlesLocal, integer numParticles, integer numNodes, real radius,
423 Curve::Type curveType = Curve::lebesgue);
424
430 real collectSendIndicesBackup(integer *toSend, integer *toSendCollected, integer count);
431
437 real collectSendEntriesBackup(SubDomainKeyTree *subDomainKeyTree, real *entry, real *toSend, integer *sendIndices,
438 integer *sendCount, integer totalSendCount, integer insertOffset);
439
445 real insertReceivedParticles(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles,
446 DomainList *domainList, DomainList *lowestDomainList, int n, int m);
447
453 real info(Tree *tree, Particles *particles, Helper *helper,
454 integer numParticlesLocal, integer numParticles, integer numNodes);
455
461 real calculateCentersOfMass(Tree *tree, Particles *particles, integer level);
462
468 real determineSearchRadii(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles,
469 DomainList *domainList, DomainList *lowestDomainList, real *searchRadii,
470 int n, int m, Curve::Type curveType);
471 }
472 }
473
474}
475
476#endif //MILUPHPC_SPH_CUH
DomainList
Definition: subdomain.cuh:494
Helper
Definition: helper.cuh:24
Material
Material parameters.
Definition: material.cuh:88
Particles
Particle(s) class based on SoA (Structur of Arrays).
Definition: particles.cuh:50
SubDomainKeyTree
SubDomainKeyTree class handling rank, number of processes and ranges.
Definition: subdomain.cuh:62
Tree
Tree class.
Definition: tree.cuh:140
DomainListNS::Kernel::lowestDomainList
__global__ void lowestDomainList(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles, DomainList *domainList, DomainList *lowestDomainList, integer n, integer m)
Kernel to create the lowest domain list.
Definition: subdomain.cu:2039
Kernel
Definition: device_rhs.cuh:7
ProfilerIds::Time::tree
const char *const tree
Definition: h5profiler.h:57
ProfilerIds::numParticlesLocal
const char *const numParticlesLocal
Definition: h5profiler.h:30
ProfilerIds::numParticles
const char *const numParticles
Definition: h5profiler.h:29
SPH::Kernel::Launch::symbolicForce_test2
real symbolicForce_test2(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles, DomainList *domainList, integer *sendIndices, real searchRadius, integer n, integer m, integer relevantProc, integer relevantIndicesCounter, Curve::Type curveType)
Definition: sph.cu:3163
SPH::Kernel::Launch::symbolicForce_test
real symbolicForce_test(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles, DomainList *lowestDomainList, integer *sendIndices, real searchRadius, integer n, integer m, integer relevantProc, integer relevantIndicesCounter, integer relevantIndexOld, Curve::Type curveType)
Definition: sph.cu:3153
SPH::Kernel::Launch::fixedRadiusNN_sharedMemory
real fixedRadiusNN_sharedMemory(Tree *tree, Particles *particles, integer *interactions, integer numParticlesLocal, integer numParticles, integer numNodes)
Wrapper for SPH::Kernel::fixedRadiusNN_sharedMemory().
Definition: sph.cu:3113
SPH::Kernel::Launch::compTheta
real compTheta(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles, DomainList *lowestDomainList, Curve::Type curveType)
Wrapper for SPH::Kernel::compTheta().
Definition: sph.cu:3137
SPH::Kernel::Launch::collectSendEntriesBackup
real collectSendEntriesBackup(SubDomainKeyTree *subDomainKeyTree, real *entry, real *toSend, integer *sendIndices, integer *sendCount, integer totalSendCount, integer insertOffset)
Wrapper for SPH::Kernel::collectSendEntriesBackup().
Definition: sph.cu:3209
SPH::Kernel::Launch::insertReceivedParticles
real insertReceivedParticles(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles, DomainList *domainList, DomainList *lowestDomainList, int n, int m)
Wrapper for SPH::Kernel::insertReceivedParticles().
Definition: sph.cu:3216
SPH::Kernel::Launch::collectSendIndices_test2
real collectSendIndices_test2(Tree *tree, Particles *particles, integer *sendIndices, integer *particles2Send, integer *particlesCount, integer numParticlesLocal, integer numParticles, integer treeIndex, int currentProc, Curve::Type curveType)
Definition: sph.cu:3180
SPH::Kernel::Launch::collectSendIndices
real collectSendIndices(Tree *tree, Particles *particles, integer *sendIndices, integer *particles2Send, integer *particlesCount, integer n, integer length, Curve::Type curveType)
Wrapper for SPH::Kernel::collectSendIndices().
Definition: sph.cu:3172
SPH::Kernel::Launch::fixedRadiusNN_bruteForce
real fixedRadiusNN_bruteForce(Tree *tree, Particles *particles, integer *interactions, integer numParticlesLocal, integer numParticles, integer numNodes)
Wrapper for SPH::Kernel::fixedRadiusNN_bruteForce().
Definition: sph.cu:3087
SPH::Kernel::Launch::collectSendIndicesBackup
real collectSendIndicesBackup(integer *toSend, integer *toSendCollected, integer count)
Wrapper for SPH::Kernel::collectSendIndicesBackup().
Definition: sph.cu:3203
SPH::Kernel::Launch::fixedRadiusNN
real fixedRadiusNN(Tree *tree, Particles *particles, integer *interactions, real radius, integer numParticlesLocal, integer numParticles, integer numNodes)
Wrapper for SPH::Kernel::fixedRadiusNN().
Definition: sph.cu:3094
SPH::Kernel::Launch::fixedRadiusNN_withinBox
real fixedRadiusNN_withinBox(Tree *tree, Particles *particles, integer *interactions, integer numParticlesLocal, integer numParticles, integer numNodes)
Wrapper for SPH::Kernel::fixedRadiusNN_withinBox().
Definition: sph.cu:3103
SPH::Kernel::Launch::particles2Send
real particles2Send(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles, DomainList *domainList, DomainList *lowestDomainList, integer maxLevel, integer *toSend, integer *sendCount, integer *alreadyInserted, integer insertOffset, integer numParticlesLocal, integer numParticles, integer numNodes, real radius, Curve::Type curveType=Curve::lebesgue)
Wrapper for SPH::Kernel::particles2Send().
Definition: sph.cu:3190
SPH::Kernel::Launch::symbolicForce
real symbolicForce(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles, DomainList *lowestDomainList, integer *sendIndices, real searchRadius, integer n, integer m, integer relevantIndex, Curve::Type curveType)
Wrapper for SPH::Kernel::symbolicForce().
Definition: sph.cu:3144
SPH::Kernel::Launch::info
real info(Tree *tree, Particles *particles, Helper *helper, integer numParticlesLocal, integer numParticles, integer numNodes)
Wrapper for SPH::Kernel::info().
Definition: sph.cu:3223
SPH::Kernel::Launch::determineSearchRadii
real determineSearchRadii(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles, DomainList *domainList, DomainList *lowestDomainList, real *searchRadii, int n, int m, Curve::Type curveType)
Wrapper for SPH::Kernel::determineSearchRadii().
Definition: sph.cu:3236
SPH::Kernel::Launch::fixedRadiusNN_variableSML
real fixedRadiusNN_variableSML(Material *materials, Tree *tree, Particles *particles, integer *interactions, integer numParticlesLocal, integer numParticles, integer numNodes)
Wrapper for SPH::Kernel::fixedRadiusNN_variableSML().
Definition: sph.cu:3128
SPH::Kernel::Launch::calculateCentersOfMass
real calculateCentersOfMass(Tree *tree, Particles *particles, integer level)
Wrapper for SPH::Kernel::calculateCentersOfMass().
Definition: sph.cu:3230
SPH::Kernel::redoNeighborSearch
__device__ void redoNeighborSearch(Tree *tree, Particles *particles, int particleId, int *interactions, real radius, integer numParticles, integer numNodes)
Redo the neighbor search (FRNN).
Definition: sph.cu:963
SPH::Kernel::compTheta
__global__ void compTheta(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles, DomainList *lowestDomainList, Curve::Type curveType)
Find the relevant (lowest) domain list nodes as preparation for finding particles to be exchanged bet...
Definition: sph.cu:1080
SPH::Kernel::calculateCentersOfMass
__global__ void calculateCentersOfMass(Tree *tree, Particles *particles, integer level)
Definition: sph.cu:2551
SPH::Kernel::symbolicForce
__global__ void symbolicForce(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles, DomainList *lowestDomainList, integer *sendIndices, real searchRadius, integer n, integer m, integer relevantIndex, Curve::Type curveType)
Find the particles that need to be exchanged between processes to grant correctness of SPH forces.
Definition: sph.cu:1128
SPH::Kernel::symbolicForce_test2
__global__ void symbolicForce_test2(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles, DomainList *domainList, integer *sendIndices, real searchRadius, integer n, integer m, integer relevantProc, integer relevantIndicesCounter, Curve::Type curveType)
Definition: sph.cu:1401
SPH::Kernel::collectSendIndices_test2
__global__ void collectSendIndices_test2(Tree *tree, Particles *particles, integer *sendIndices, integer *particles2Send, integer *particlesCount, integer numParticlesLocal, integer numParticles, integer treeIndex, int currentProc, Curve::Type curveType)
Definition: sph.cu:1550
SPH::Kernel::insertReceivedParticles
__global__ void insertReceivedParticles(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles, DomainList *domainList, DomainList *lowestDomainList, int n, int m)
Insert the received particles into the local tree.
Definition: sph.cu:1875
SPH::Kernel::collectSendIndicesBackup
__global__ void collectSendIndicesBackup(integer *toSend, integer *toSendCollected, integer count)
Definition: sph.cu:1776
SPH::Kernel::fixedRadiusNN_variableSML
__global__ void fixedRadiusNN_variableSML(Material *materials, Tree *tree, Particles *particles, integer *interactions, integer numParticlesLocal, integer numParticles, integer numNodes)
Fixed-radius near neighbor search for iteratively finding appropriate smoothing length.
Definition: sph.cu:739
SPH::Kernel::particles2Send
__global__ void particles2Send(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles, DomainList *domainList, DomainList *lowestDomainList, integer maxLevel, integer *toSend, integer *sendCount, integer *alreadyInserted, integer insertOffset, integer numParticlesLocal, integer numParticles, integer numNodes, real radius, Curve::Type curveType=Curve::lebesgue)
Definition: sph.cu:1575
SPH::Kernel::info
__global__ void info(Tree *tree, Particles *particles, Helper *helper, integer numParticlesLocal, integer numParticles, integer numNodes)
Info/Debug kernel.
Definition: sph.cu:3062
SPH::Kernel::collectSendIndices
__global__ void collectSendIndices(Tree *tree, Particles *particles, integer *sendIndices, integer *particles2Send, integer *particlesCount, integer n, integer length, Curve::Type curveType)
Collect the found particles into contiguous memory in order to facilitate sending via MPI.
Definition: sph.cu:1525
SPH::Kernel::fixedRadiusNN
__global__ void fixedRadiusNN(Tree *tree, Particles *particles, integer *interactions, real radius, integer numParticlesLocal, integer numParticles, integer numNodes)
Fixed-radius near neighbor search (default method via explicit stack).
Definition: sph.cu:93
SPH::Kernel::fixedRadiusNN_withinBox
__global__ void fixedRadiusNN_withinBox(Tree *tree, Particles *particles, integer *interactions, integer numParticlesLocal, integer numParticles, integer numNodes)
Fixed-radius near neighbor search (nested stack method).
Definition: sph.cu:283
SPH::Kernel::determineSearchRadii
__global__ void determineSearchRadii(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles, DomainList *domainList, DomainList *lowestDomainList, real *searchRadii, int n, int m, Curve::Type curveType)
Determine the search radius needed for SPH::Kernel::symbolicForce().
Definition: sph.cu:2619
SPH::Kernel::symbolicForce_test
__global__ void symbolicForce_test(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles, DomainList *lowestDomainList, integer *sendIndices, real searchRadius, integer n, integer m, integer relevantProc, integer relevantIndicesCounter, integer relevantIndexOld, Curve::Type curveType)
Definition: sph.cu:1271
SPH::Kernel::collectSendEntriesBackup
__global__ void collectSendEntriesBackup(SubDomainKeyTree *subDomainKeyTree, real *entry, real *toSend, integer *sendIndices, integer *sendCount, integer totalSendCount, integer insertOffset)
Definition: sph.cu:1789
SPH::Kernel::fixedRadiusNN_sharedMemory
__global__ void fixedRadiusNN_sharedMemory(Tree *tree, Particles *particles, integer *interactions, integer numParticlesLocal, integer numParticles, integer numNodes)
Fixed-radius near neighbor search (brute-force method).
Definition: sph.cu:592
SPH::Kernel::fixedRadiusNN_bruteForce
__global__ void fixedRadiusNN_bruteForce(Tree *tree, Particles *particles, integer *interactions, integer numParticlesLocal, integer numParticles, integer numNodes)
Fixed-radius near neighbor search (brute-force method).
Definition: sph.cu:36
SPH
SPH related functions and kernels.
Definition: density.cuh:23
SPH::exchangeParticleEntry
void exchangeParticleEntry(SubDomainKeyTree *subDomainKeyTree, real *entry, real *toSend, integer *sendLengths, integer *receiveLengths, integer numParticlesLocal)
Definition: sph.cu:12
real
double real
Definition: parameter.h:15
integer
int integer
Definition: parameter.h:17
Curve::Type
Type
Definition: parameter.h:206
Curve::lebesgue
@ lebesgue
Definition: parameter.h:207

milupHPC - include/sph/sph.cuh Source File
Generated on Wed Aug 31 2022 12:16:52 by Doxygen 1.9.3