subdomain.cuh

Go to the documentation of this file.

 
#ifndef MILUPHPC_DOMAIN_CUH
#define MILUPHPC_DOMAIN_CUH
 
#include "../parameter.h"
#include "../cuda_utils/cuda_runtime.h"
#include "../cuda_utils/cuda_utilities.cuh"
#include "../helper_handler.h"
#include "tree.cuh"
 
//TODO: necessary/reasonable to have Key class?
//class Key {
    //keyType *keys;
    //integer *maxLevel;
    //CUDA_CALLABLE_MEMBER Key();
    //CUDA_CALLABLE_MEMBER Key(keyType *keys, integer *maxLevel);
    //CUDA_CALLABLE_MEMBER ~Key();
    //CUDA_CALLABLE_MEMBER void set(keyType *keys, integer *maxLevel);*/
//};
 
// Forward declaration of DomainList class
class DomainList;
// Forward declaration of SubDomainKeyTree class
class SubDomainKeyTree;
 
namespace KeyNS {
 
    CUDA_CALLABLE_MEMBER void key2Char(keyType key, integer maxLevel, char *keyAsChar);
 
    CUDA_CALLABLE_MEMBER integer key2proc(keyType key, SubDomainKeyTree *subDomainKeyTree/*, Curve::Type curveType=Curve::lebesgue*/);
}
 
class SubDomainKeyTree {
 
public:
    integer rank;
    integer numProcesses;
    keyType *range;
    //keyType *hilbertRange;
    integer *procParticleCounter;
 
    CUDA_CALLABLE_MEMBER SubDomainKeyTree();
    CUDA_CALLABLE_MEMBER SubDomainKeyTree(integer rank, integer numProcesses, keyType *range,
                                          integer *procParticleCounter);
    CUDA_CALLABLE_MEMBER ~SubDomainKeyTree();
 
    CUDA_CALLABLE_MEMBER void set(integer rank, integer numProcesses, keyType *range, integer *procParticleCounter);
 
    CUDA_CALLABLE_MEMBER integer key2proc(keyType key/*, Curve::Type curveType=Curve::lebesgue*/);
 
    CUDA_CALLABLE_MEMBER bool isDomainListNode(keyType key, integer maxLevel, integer level,
                                               Curve::Type curveType=Curve::lebesgue);
 
};
 
namespace SubDomainKeyTreeNS {
 
    namespace Kernel {
 
        __global__ void set(SubDomainKeyTree *subDomainKeyTree, integer rank, integer numProcesses, keyType *range,
                            integer *procParticleCounter);
 
        __global__ void test(SubDomainKeyTree *subDomainKeyTree);
 
        __global__ void buildDomainTree(Tree *tree, Particles *particles, DomainList *domainList, integer n, integer m);
 
        __global__ void buildDomainTree(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles, DomainList *domainList, integer n, integer m, integer level);
 
        __global__ void getParticleKeys(SubDomainKeyTree *subDomainKeyTree, Tree *tree,
                                                        Particles *particles, keyType *keys, integer maxLevel,
                                                        integer n, Curve::Type curveType = Curve::lebesgue);
 
        __global__ void particlesPerProcess(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles,
                                            integer n, integer m, Curve::Type curveType=Curve::lebesgue);
 
        __global__ void markParticlesProcess(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles,
                                             integer n, integer m, integer *sortArray,
                                             Curve::Type curveType=Curve::lebesgue);
 
        __global__ void zeroDomainListNodes(Particles *particles, DomainList *domainList,
                                            DomainList *lowestDomainList);
 
         template <typename T>
        __global__ void prepareLowestDomainExchange(Particles *particles, DomainList *lowestDomainList,
                                                    T *buffer, Entry::Name entry);
 
        template <typename T>
        __global__ void updateLowestDomainListNodes(Particles *particles, DomainList *lowestDomainList,
                                                    T *buffer, Entry::Name entry);
 
        __global__ void compLowestDomainListNodes(Tree *tree, Particles *particles, DomainList *lowestDomainList);
 
        __global__ void compLocalPseudoParticles(Tree *tree, Particles *particles, DomainList *domainList, int n);
 
        __global__ void compDomainListPseudoParticlesPerLevel(Tree *tree, Particles *particles, DomainList *domainList,
                                                              DomainList *lowestDomainList, int n, int level);
 
        __global__ void compDomainListPseudoParticles(Tree *tree, Particles *particles, DomainList *domainList,
                                                      DomainList *lowestDomainList, int n);
 
        __global__ void repairTree(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles,
                                   DomainList *domainList, DomainList *lowestDomainList,
                                   int n, int m, Curve::Type curveType);
 
        // notes:
        // - using Helper::keyTypeBuffer as keyHistRanges
        __global__ void createKeyHistRanges(Helper *helper, integer bins);
 
        // notes:
        // - using Helper::keyTypeBuffer as keyHistRanges
        // - using Helper::integerBuffer as keyHistCounts
        __global__ void keyHistCounter(Tree *tree, Particles *particles, SubDomainKeyTree *subDomainKeyTree,
                                       Helper *helper,
                /*keyType *keyHistRanges, integer *keyHistCounts,*/ int bins, int n,
                                       Curve::Type curveType=Curve::lebesgue);
 
        // notes:
        // - using Helper::keyTypeBuffer as keyHistRanges
        // - using Helper::integerBuffer as keyHistCounts
        __global__ void calculateNewRange(SubDomainKeyTree *subDomainKeyTree, Helper *helper,
                /*keyType *keyHistRanges, integer *keyHistCounts,*/ int bins, int n,
                                          Curve::Type curveType=Curve::lebesgue);
 
 
        namespace Launch {
 
            void set(SubDomainKeyTree *subDomainKeyTree, integer rank, integer numProcesses, keyType *range,
                     integer *procParticleCounter);
 
            void test(SubDomainKeyTree *subDomainKeyTree);
 
            real buildDomainTree(Tree *tree, Particles *particles, DomainList *domainList, integer n, integer m);
 
            real buildDomainTree(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles, DomainList *domainList, integer n, integer m, integer level);
 
            real getParticleKeys(SubDomainKeyTree *subDomainKeyTree, Tree *tree,
                                 Particles *particles, keyType *keys, integer maxLevel,
                                 integer n, Curve::Type curveType = Curve::lebesgue);
 
            real particlesPerProcess(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles,
                                                integer n, integer m, Curve::Type curveType=Curve::lebesgue);
 
            real markParticlesProcess(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles,
                                                 integer n, integer m, integer *sortArray,
                                                 Curve::Type curveType=Curve::lebesgue);
 
            real zeroDomainListNodes(Particles *particles, DomainList *domainList, DomainList *lowestDomainList);
 
            template <typename T>
            real prepareLowestDomainExchange(Particles *particles, DomainList *lowestDomainList,
                                             T *buffer, Entry::Name entry);
 
 
            template <typename T>
            real updateLowestDomainListNodes(Particles *particles, DomainList *lowestDomainList,
                                             T *buffer, Entry::Name entry);
 
            real compLowestDomainListNodes(Tree *tree, Particles *particles, DomainList *lowestDomainList);
 
            real compLocalPseudoParticles(Tree *tree, Particles *particles, DomainList *domainList, int n);
 
            real compDomainListPseudoParticlesPerLevel(Tree *tree, Particles *particles, DomainList *domainList,
                                                       DomainList *lowestDomainList, int n, int level);
 
            real compDomainListPseudoParticles(Tree *tree, Particles *particles, DomainList *domainList,
                                               DomainList *lowestDomainList, int n);
 
            real repairTree(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles,
                            DomainList *domainList, DomainList *lowestDomainList,
                            int n, int m, Curve::Type curveType);
 
            real createKeyHistRanges(Helper *helper, integer bins);
 
            real keyHistCounter(Tree *tree, Particles *particles, SubDomainKeyTree *subDomainKeyTree,
                                Helper *helper, int bins, int n, Curve::Type curveType=Curve::lebesgue);
 
            real calculateNewRange(SubDomainKeyTree *subDomainKeyTree, Helper *helper, int bins, int n,
                                   Curve::Type curveType=Curve::lebesgue);
        }
 
    }
 
}
 
class DomainList {
 
public:
 
    integer *domainListIndices;
    integer *domainListLevels;
    integer *domainListIndex;
    integer *domainListCounter;
    keyType *domainListKeys;
    keyType *sortedDomainListKeys;
    integer *relevantDomainListIndices;
    integer *relevantDomainListLevels;
    integer *relevantDomainListProcess;
 
    integer *relevantDomainListOriginalIndex;
 
    real *borders;
 
    CUDA_CALLABLE_MEMBER DomainList();
 
    CUDA_CALLABLE_MEMBER DomainList(integer *domainListIndices, integer *domainListLevels, integer *domainListIndex,
                                    integer *domainListCounter, keyType *domainListKeys, keyType *sortedDomainListKeys,
                                    integer *relevantDomainListIndices, integer *relevantDomainListLevels,
                                    integer *relevantDomainListProcess);
    CUDA_CALLABLE_MEMBER void set(integer *domainListIndices, integer *domainListLevels, integer *domainListIndex,
                                  integer *domainListCounter, keyType *domainListKeys, keyType *sortedDomainListKeys,
                                  integer *relevantDomainListIndices, integer *relevantDomainListLevels,
                                  integer *relevantDomainListProcess);
 
    CUDA_CALLABLE_MEMBER void setBorders(real *borders, integer *relevantDomainListOriginalIndex);
 
    CUDA_CALLABLE_MEMBER ~DomainList();
};
 
namespace DomainListNS {
 
    namespace Kernel {
 
        __global__ void set(DomainList *domainList, integer *domainListIndices, integer *domainListLevels,
                            integer *domainListIndex, integer *domainListCounter, keyType *domainListKeys,
                            keyType *sortedDomainListKeys, integer *relevantDomainListIndices,
                            integer *relevantDomainListLevels, integer *relevantDomainListProcess);
 
        __global__ void setBorders(DomainList *domainList, real *borders, integer *relevantDomainListOriginalIndex);
 
        __global__ void info(Particles *particles, DomainList *domainList);
 
        __global__ void info(Particles *particles, DomainList *domainList, DomainList *lowestDomainList);
 
        __global__ void createDomainList(SubDomainKeyTree *subDomainKeyTree, DomainList *domainList,
                                         integer maxLevel, Curve::Type curveType = Curve::lebesgue);
 
        __global__ void lowestDomainList(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles,
                                         DomainList *domainList, DomainList *lowestDomainList, integer n, integer m);
 
 
        namespace Launch {
 
            void set(DomainList *domainList, integer *domainListIndices, integer *domainListLevels,
                     integer *domainListIndex, integer *domainListCounter, keyType *domainListKeys,
                     keyType *sortedDomainListKeys, integer *relevantDomainListIndices,
                     integer *relevantDomainListLevels, integer *relevantDomainListProcess);
 
            void setBorders(DomainList *domainList, real *borders, integer *relevantDomainListOriginalIndex);
 
            real info(Particles *particles, DomainList *domainList);
 
            real info(Particles *particles, DomainList *domainList, DomainList *lowestDomainList);
 
            real createDomainList(SubDomainKeyTree *subDomainKeyTree, DomainList *domainList,
                                  integer maxLevel, Curve::Type curveType = Curve::lebesgue);
 
            real lowestDomainList(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles,
                                  DomainList *domainList, DomainList *lowestDomainList, integer n, integer m);
 
        }
    }
 
}
 
namespace ParticlesNS {
 
    __device__ bool applySphericalCriterion(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles,
                                            real d, int index);
 
    __device__ bool applyCubicCriterion(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles,
                                        real d, int index);
 
    namespace Kernel {
        __global__ void mark2remove(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles,
                                    int *particles2remove, int *counter, int criterion, real d,
                                    int numParticles);
 
        namespace Launch {
            real mark2remove(SubDomainKeyTree *subDomainKeyTree, Tree *tree, Particles *particles,
                             int *particles2remove, int *counter, int criterion, real d,
                             int numParticles);
        }
    }
}
 
namespace CudaUtils {
    namespace Kernel {
        template<typename T, typename U>
        __global__ void markDuplicatesTemp(Tree *tree, DomainList *domainList, T *array, U *entry1, U *entry2, U *entry3, integer *duplicateCounter, integer *child, int length);
 
        template <typename T, unsigned int blockSize>
        __global__ void reduceBlockwise(T *array, T *outputData, int n);
 
        template <typename T, unsigned int blockSize>
        __global__ void blockReduction(const T *indata, T *outdata);
 
        namespace Launch {
            template<typename T, typename U>
            real markDuplicatesTemp(Tree *tree, DomainList *domainList, T *array, U *entry1, U *entry2, U *entry3, integer *duplicateCounter, integer *child, int length);
 
            template <typename T, unsigned int blockSize>
            real reduceBlockwise(T *array, T *outputData, int n);
 
            template <typename T, unsigned int blockSize>
            real blockReduction(const T *indata, T *outdata);
        }
    }
}
 
namespace Physics {
    namespace Kernel {
 
        template <unsigned int blockSize>
        __global__ void calculateAngularMomentumBlockwise(Particles *particles, real *outputData, int n);
 
        template <unsigned int blockSize>
        __global__ void sumAngularMomentum(const real *indata, real *outdata);
 
        __global__ void kineticEnergy(Particles *particles, int n);
 
        namespace Launch {
            template <unsigned int blockSize>
            real calculateAngularMomentumBlockwise(Particles *particles, real *outputData, int n);
 
            template <unsigned int blockSize>
            real sumAngularMomentum(const real *indata, real *outdata);
 
            real kineticEnergy(Particles *particles, int n);
        }
    }
}
 
#endif //MILUPHPC_DOMAIN_CUH