From f3e10ded2adf42c130248a6c946a286b19117b2a Mon Sep 17 00:00:00 2001 From: "LEGOLAS\\lenz" <lenz@irmb.tu-bs.de> Date: Wed, 4 Dec 2019 16:14:36 +0100 Subject: [PATCH] renames kernels --- .../CumulantAA2016/CumulantAA2016CompSP27.cu | 53 ----------------- .../CumulantAA2016/CumulantAA2016CompSP27.h | 17 ------ .../CumulantAA2016CompSP27_Device.cuh | 18 ------ .../CumulantAA2016CompBulkSP27.h | 17 ------ .../CumulantF3/CumulantF3CompSP27.h | 17 ------ .../CumulantF32018/CumulantF32018CompSP27.h | 16 ------ .../CumulantK15/CumulantK15Comp.cu | 53 +++++++++++++++++ .../CumulantK15/CumulantK15Comp.h | 16 ++++++ .../CumulantK15Comp_Device.cu} | 2 +- .../CumulantK15/CumulantK15Comp_Device.cuh | 17 ++++++ .../package.include | 0 .../CumulantK15Bulk/CumulantK15BulkComp.cu | 56 ++++++++++++++++++ .../CumulantK15Bulk/CumulantK15BulkComp.h | 17 ++++++ .../CumulantK15BulkComp_Device.cu} | 2 +- .../CumulantK15BulkComp_Device.cuh} | 2 +- .../package.include | 0 .../CumulantK15SpongeComp.cu | 57 +++++++++++++++++++ .../CumulantK15Sponge/CumulantK15SpongeComp.h | 17 ++++++ .../CumulantK15SpongeComp_Device.cu} | 2 +- .../CumulantK15SpongeComp_Device.cuh} | 2 +- .../package.include | 0 .../CumulantK17/CumulantK17Comp.cu | 53 +++++++++++++++++ .../CumulantK17/CumulantK17Comp.h | 17 ++++++ .../CumulantK17Comp_Device.cu} | 2 +- .../CumulantK17/CumulantK17Comp_Device.cuh | 18 ++++++ .../package.include | 0 .../CumulantK17BulkComp.cu} | 18 +++--- .../CumulantK17Bulk/CumulantK17BulkComp.h | 17 ++++++ .../CumulantK17BulkComp_Device.cu} | 2 +- .../CumulantK17BulkComp_Device.cuh} | 2 +- .../package.include | 0 .../CumulantK18Comp.cu} | 16 +++--- .../CumulantK18/CumulantK18Comp.h | 17 ++++++ .../CumulantK18Comp_Device.cu} | 4 +- .../CumulantK18Comp_Device.cuh} | 2 +- .../package.include | 0 .../CumulantK20Comp.cu} | 16 +++--- .../CumulantK20/CumulantK20Comp.h | 16 ++++++ .../CumulantK20Comp_Device.cu} | 4 +- .../CumulantK20Comp_Device.cuh} | 2 +- .../package.include | 0 .../CumulantOne/CumulantOneCompSP27.cu | 53 ----------------- .../CumulantOne/CumulantOneCompSP27.h | 16 ------ .../CumulantOneCompSP27_Device.cuh | 17 ------ .../CumulantOneCompBulkSP27.cu | 56 ------------------ .../CumulantOneBulk/CumulantOneCompBulkSP27.h | 17 ------ .../CumulantOneCompSpongeSP27.cu | 57 ------------------- .../CumulantOneCompSpongeSP27.h | 17 ------ .../CumulantK15Incomp.cu} | 18 +++--- .../CumulantK15/CumulantK15Incomp.h | 16 ++++++ .../CumulantK15Incomp_Device.cu} | 2 +- .../CumulantK15Incomp_Device.cuh} | 2 +- .../package.include | 0 .../CumulantOne/CumulantOneIncompSP27.h | 16 ------ .../KernelFactory/KernelFactoryImp.cpp | 52 ++++++++--------- .../Kernel/Utilities/KernelType.h | 16 +++--- .../Mapper/KernelMapper/KernelMapper.cpp | 16 +++--- src/VirtualFluids_GPU/Parameter/Parameter.cpp | 4 +- targets/apps/LBM/TGV_3D/TGV_3D.cpp | 4 +- 59 files changed, 482 insertions(+), 484 deletions(-) delete mode 100644 src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016/CumulantAA2016CompSP27.cu delete mode 100644 src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016/CumulantAA2016CompSP27.h delete mode 100644 src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016/CumulantAA2016CompSP27_Device.cuh delete mode 100644 src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016Bulk/CumulantAA2016CompBulkSP27.h delete mode 100644 src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF3/CumulantF3CompSP27.h delete mode 100644 src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF32018/CumulantF32018CompSP27.h create mode 100644 src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15/CumulantK15Comp.cu create mode 100644 src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15/CumulantK15Comp.h rename src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/{CumulantOne/CumulantOneCompSP27_Device.cu => CumulantK15/CumulantK15Comp_Device.cu} (99%) create mode 100644 src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15/CumulantK15Comp_Device.cuh rename src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/{CumulantAA2016 => CumulantK15}/package.include (100%) create mode 100644 src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Bulk/CumulantK15BulkComp.cu create mode 100644 src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Bulk/CumulantK15BulkComp.h rename src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/{CumulantOneBulk/CumulantOneCompBulkSP27_Device.cu => CumulantK15Bulk/CumulantK15BulkComp_Device.cu} (99%) rename src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/{CumulantOneBulk/CumulantOneCompBulkSP27_Device.cuh => CumulantK15Bulk/CumulantK15BulkComp_Device.cuh} (81%) rename src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/{CumulantAA2016Bulk => CumulantK15Bulk}/package.include (100%) create mode 100644 src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Sponge/CumulantK15SpongeComp.cu create mode 100644 src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Sponge/CumulantK15SpongeComp.h rename src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/{CumulantOneSponge/CumulantOneCompSpongeSP27_Device.cu => CumulantK15Sponge/CumulantK15SpongeComp_Device.cu} (99%) rename src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/{CumulantOneSponge/CumulantOneCompSpongeSP27_Device.cuh => CumulantK15Sponge/CumulantK15SpongeComp_Device.cuh} (81%) rename src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/{CumulantF3 => CumulantK15Sponge}/package.include (100%) create mode 100644 src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17/CumulantK17Comp.cu create mode 100644 src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17/CumulantK17Comp.h rename src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/{CumulantAA2016/CumulantAA2016CompSP27_Device.cu => CumulantK17/CumulantK17Comp_Device.cu} (99%) create mode 100644 src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17/CumulantK17Comp_Device.cuh rename src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/{CumulantF32018 => CumulantK17}/package.include (100%) rename src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/{CumulantAA2016Bulk/CumulantAA2016CompBulkSP27.cu => CumulantK17Bulk/CumulantK17BulkComp.cu} (57%) create mode 100644 src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17Bulk/CumulantK17BulkComp.h rename src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/{CumulantAA2016Bulk/CumulantAA2016CompBulkSP27_Device.cu => CumulantK17Bulk/CumulantK17BulkComp_Device.cu} (99%) rename src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/{CumulantAA2016Bulk/CumulantAA2016CompBulkSP27_Device.cuh => CumulantK17Bulk/CumulantK17BulkComp_Device.cuh} (82%) rename src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/{CumulantOne => CumulantK17Bulk}/package.include (100%) rename src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/{CumulantF3/CumulantF3CompSP27.cu => CumulantK18/CumulantK18Comp.cu} (67%) create mode 100644 src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK18/CumulantK18Comp.h rename src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/{CumulantF3/CumulantF3CompSP27_Device.cu => CumulantK18/CumulantK18Comp_Device.cu} (99%) rename src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/{CumulantF3/CumulantF3CompSP27_Device.cuh => CumulantK18/CumulantK18Comp_Device.cuh} (88%) rename src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/{CumulantOneBulk => CumulantK18}/package.include (100%) rename src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/{CumulantF32018/CumulantF32018CompSP27.cu => CumulantK20/CumulantK20Comp.cu} (65%) create mode 100644 src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK20/CumulantK20Comp.h rename src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/{CumulantF32018/CumulantF32018CompSP27_Device.cu => CumulantK20/CumulantK20Comp_Device.cu} (99%) rename src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/{CumulantF32018/CumulantF32018CompSP27_Device.cuh => CumulantK20/CumulantK20Comp_Device.cuh} (88%) rename src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/{CumulantOneSponge => CumulantK20}/package.include (100%) delete mode 100644 src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOne/CumulantOneCompSP27.cu delete mode 100644 src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOne/CumulantOneCompSP27.h delete mode 100644 src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOne/CumulantOneCompSP27_Device.cuh delete mode 100644 src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneBulk/CumulantOneCompBulkSP27.cu delete mode 100644 src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneBulk/CumulantOneCompBulkSP27.h delete mode 100644 src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneSponge/CumulantOneCompSpongeSP27.cu delete mode 100644 src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneSponge/CumulantOneCompSpongeSP27.h rename src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/{CumulantOne/CumulantOneIncompSP27.cu => CumulantK15/CumulantK15Incomp.cu} (58%) create mode 100644 src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantK15/CumulantK15Incomp.h rename src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/{CumulantOne/CumulantOneIncompSP27_Device.cu => CumulantK15/CumulantK15Incomp_Device.cu} (99%) rename src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/{CumulantOne/CumulantOneIncompSP27_Device.cuh => CumulantK15/CumulantK15Incomp_Device.cuh} (80%) rename src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/{CumulantOne => CumulantK15}/package.include (100%) delete mode 100644 src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantOne/CumulantOneIncompSP27.h diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016/CumulantAA2016CompSP27.cu b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016/CumulantAA2016CompSP27.cu deleted file mode 100644 index b7c5c67cb..000000000 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016/CumulantAA2016CompSP27.cu +++ /dev/null @@ -1,53 +0,0 @@ -#include "CumulantAA2016CompSP27.h" - -#include "Parameter/Parameter.h" -#include "CumulantAA2016CompSP27_Device.cuh" - -std::shared_ptr<CumulantAA2016CompSP27> CumulantAA2016CompSP27::getNewInstance(std::shared_ptr<Parameter> para, int level) -{ - return std::shared_ptr<CumulantAA2016CompSP27>(new CumulantAA2016CompSP27(para,level)); -} - -void CumulantAA2016CompSP27::run() -{ - int numberOfThreads = para->getParD(level)->numberofthreads; - int size_Mat = para->getParD(level)->size_Mat_SP; - - int Grid = (size_Mat / numberOfThreads) + 1; - int Grid1, Grid2; - if (Grid>512) - { - Grid1 = 512; - Grid2 = (Grid / Grid1) + 1; - } - else - { - Grid1 = 1; - Grid2 = Grid; - } - dim3 grid(Grid1, Grid2); - dim3 threads(numberOfThreads, 1, 1); - - LB_Kernel_Kum_AA2016_Comp_SP_27 << < grid, threads >> >( para->getParD(level)->omega, - para->getParD(level)->geoSP, - para->getParD(level)->neighborX_SP, - para->getParD(level)->neighborY_SP, - para->getParD(level)->neighborZ_SP, - para->getParD(level)->d0SP.f[0], - para->getParD(level)->size_Mat_SP, - level, - para->getForcesDev(), - para->getQuadricLimitersDev(), - para->getParD(level)->evenOrOdd); - getLastCudaError("LB_Kernel_Kum_AA2016_Comp_SP_27 execution failed"); -} - -CumulantAA2016CompSP27::CumulantAA2016CompSP27(std::shared_ptr<Parameter> para, int level) -{ - this->para = para; - this->level = level; - - myPreProcessorTypes.push_back(InitCompSP27); - - myKernelGroup = BasicKernel; -} \ No newline at end of file diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016/CumulantAA2016CompSP27.h b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016/CumulantAA2016CompSP27.h deleted file mode 100644 index f6e3dd40e..000000000 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016/CumulantAA2016CompSP27.h +++ /dev/null @@ -1,17 +0,0 @@ -#ifndef CUMULANT_AA2016_COMP_SP27_H -#define CUMULANT_AA2016_COMP_SP27_H - -#include "Kernel\KernelImp.h" - -class CumulantAA2016CompSP27 : public KernelImp -{ -public: - static std::shared_ptr<CumulantAA2016CompSP27> getNewInstance(std::shared_ptr< Parameter> para, int level); - void run(); - -private: - CumulantAA2016CompSP27(); - CumulantAA2016CompSP27(std::shared_ptr< Parameter> para, int level); -}; - -#endif diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016/CumulantAA2016CompSP27_Device.cuh b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016/CumulantAA2016CompSP27_Device.cuh deleted file mode 100644 index fdc68216f..000000000 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016/CumulantAA2016CompSP27_Device.cuh +++ /dev/null @@ -1,18 +0,0 @@ -#ifndef LB_Kernel_Kum_AA2016_Comp_SP_27_H -#define LB_Kernel_Kum_AA2016_Comp_SP_27_H - -#include <DataTypes.h> -#include <curand.h> - -extern "C" __global__ void LB_Kernel_Kum_AA2016_Comp_SP_27( real omega, - unsigned int* bcMatD, - unsigned int* neighborX, - unsigned int* neighborY, - unsigned int* neighborZ, - real* DDStart, - int size_Mat, - int level, - real* forces, - real* quadricLimiters, - bool EvenOrOdd); -#endif diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016Bulk/CumulantAA2016CompBulkSP27.h b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016Bulk/CumulantAA2016CompBulkSP27.h deleted file mode 100644 index 0d21d261d..000000000 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016Bulk/CumulantAA2016CompBulkSP27.h +++ /dev/null @@ -1,17 +0,0 @@ -#ifndef CUMULANT_AA2016_COMP_BULK_SP27_H -#define CUMULANT_AA2016_COMP_BULK_SP27_H - -#include "Kernel\KernelImp.h" - -class CumulantAA2016CompBulkSP27 : public KernelImp -{ -public: - static std::shared_ptr<CumulantAA2016CompBulkSP27> getNewInstance(std::shared_ptr< Parameter> para, int level); - void run(); - -private: - CumulantAA2016CompBulkSP27(); - CumulantAA2016CompBulkSP27(std::shared_ptr< Parameter> para, int level); - -}; -#endif \ No newline at end of file diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF3/CumulantF3CompSP27.h b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF3/CumulantF3CompSP27.h deleted file mode 100644 index eaaf6a8bb..000000000 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF3/CumulantF3CompSP27.h +++ /dev/null @@ -1,17 +0,0 @@ -#ifndef CUMULANT_F3_COMP_SP27_H -#define CUMULANT_F3_COMP_SP27_H - -#include "Kernel\KernelImp.h" - -class CumulantF3CompSP27 : public KernelImp -{ -public: - static std::shared_ptr< CumulantF3CompSP27> getNewInstance(std::shared_ptr< Parameter> para, int level); - void run(); - -private: - CumulantF3CompSP27(); - CumulantF3CompSP27(std::shared_ptr< Parameter> para, int level); -}; - -#endif \ No newline at end of file diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF32018/CumulantF32018CompSP27.h b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF32018/CumulantF32018CompSP27.h deleted file mode 100644 index aa40bf336..000000000 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF32018/CumulantF32018CompSP27.h +++ /dev/null @@ -1,16 +0,0 @@ -#ifndef CUMULANT_F32018_COMP_SP27_H -#define CUMULANT_F32018_COMP_SP27_H - -#include "Kernel\KernelImp.h" - -class CumulantF32018CompSP27 : public KernelImp -{ -public: - static std::shared_ptr< CumulantF32018CompSP27> getNewInstance(std::shared_ptr< Parameter> para, int level); - void run(); - -private: - CumulantF32018CompSP27(); - CumulantF32018CompSP27(std::shared_ptr< Parameter> para, int level); -}; -#endif \ No newline at end of file diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15/CumulantK15Comp.cu b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15/CumulantK15Comp.cu new file mode 100644 index 000000000..01cc11cbb --- /dev/null +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15/CumulantK15Comp.cu @@ -0,0 +1,53 @@ +#include "CumulantK15Comp.h" + +#include "CumulantK15Comp_Device.cuh" + +#include "Parameter/Parameter.h" + +std::shared_ptr<CumulantK15Comp> CumulantK15Comp::getNewInstance(std::shared_ptr<Parameter> para, int level) +{ + return std::shared_ptr<CumulantK15Comp>(new CumulantK15Comp(para, level)); +} + +void CumulantK15Comp::run() +{ + int numberOfThreads = para->getParD(level)->numberofthreads; + int size_Mat = para->getParD(level)->size_Mat_SP; + + int Grid = (size_Mat / numberOfThreads) + 1; + int Grid1, Grid2; + if (Grid > 512) + { + Grid1 = 512; + Grid2 = (Grid / Grid1) + 1; + } + else + { + Grid1 = 1; + Grid2 = Grid; + } + dim3 grid(Grid1, Grid2, 1); + dim3 threads(numberOfThreads, 1, 1); + + LB_Kernel_CumulantK15Comp <<< grid, threads >>>(para->getParD(level)->omega, + para->getParD(level)->geoSP, + para->getParD(level)->neighborX_SP, + para->getParD(level)->neighborY_SP, + para->getParD(level)->neighborZ_SP, + para->getParD(level)->d0SP.f[0], + size_Mat, + level, + para->getForcesDev(), + para->getParD(level)->evenOrOdd); + getLastCudaError("LB_Kernel_CumulantK15Comp execution failed"); +} + +CumulantK15Comp::CumulantK15Comp(std::shared_ptr<Parameter> para, int level) +{ + this->para = para; + this->level = level; + + myPreProcessorTypes.push_back(InitCompSP27); + + myKernelGroup = BasicKernel; +} \ No newline at end of file diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15/CumulantK15Comp.h b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15/CumulantK15Comp.h new file mode 100644 index 000000000..9e65a76b4 --- /dev/null +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15/CumulantK15Comp.h @@ -0,0 +1,16 @@ +#ifndef CUMULANT_ONE_COMP_SP27_H +#define CUMULANT_ONE_COMP_SP27_H + +#include "Kernel\KernelImp.h" + +class CumulantK15Comp : public KernelImp +{ +public: + static std::shared_ptr< CumulantK15Comp> getNewInstance(std::shared_ptr< Parameter> para, int level); + void run(); + +private: + CumulantK15Comp(); + CumulantK15Comp(std::shared_ptr< Parameter> para, int level); +}; +#endif \ No newline at end of file diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOne/CumulantOneCompSP27_Device.cu b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15/CumulantK15Comp_Device.cu similarity index 99% rename from src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOne/CumulantOneCompSP27_Device.cu rename to src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15/CumulantK15Comp_Device.cu index aad1ac081..183d83a6c 100644 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOne/CumulantOneCompSP27_Device.cu +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15/CumulantK15Comp_Device.cu @@ -3,7 +3,7 @@ #include "Core/RealConstants.h" #include "math.h" -extern "C" __global__ void LB_Kernel_Kum_One_Comp_SP_27(real omega, +extern "C" __global__ void LB_Kernel_CumulantK15Comp(real omega, unsigned int* bcMatD, unsigned int* neighborX, unsigned int* neighborY, diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15/CumulantK15Comp_Device.cuh b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15/CumulantK15Comp_Device.cuh new file mode 100644 index 000000000..79c4ac7a8 --- /dev/null +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15/CumulantK15Comp_Device.cuh @@ -0,0 +1,17 @@ +#ifndef LB_Kernel_Kum_ONE_COMP_SP_27_H +#define LB_Kernel_Kum_ONE_COMP_SP_27_H + +#include <DataTypes.h> +#include <curand.h> + +extern "C" __global__ void LB_Kernel_CumulantK15Comp( real omega, + unsigned int* bcMatD, + unsigned int* neighborX, + unsigned int* neighborY, + unsigned int* neighborZ, + real* DDStart, + int size_Mat, + int level, + real* forces, + bool EvenOrOdd); +#endif \ No newline at end of file diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016/package.include b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15/package.include similarity index 100% rename from src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016/package.include rename to src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15/package.include diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Bulk/CumulantK15BulkComp.cu b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Bulk/CumulantK15BulkComp.cu new file mode 100644 index 000000000..5d32fb19d --- /dev/null +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Bulk/CumulantK15BulkComp.cu @@ -0,0 +1,56 @@ +#include "CumulantK15BulkComp.h" + +#include "CumulantK15BulkComp_Device.cuh" +#include "Parameter\Parameter.h" + +std::shared_ptr<CumulantK15BulkComp> CumulantK15BulkComp::getNewInstance(std::shared_ptr<Parameter> para, int level) +{ + return std::shared_ptr<CumulantK15BulkComp>(new CumulantK15BulkComp(para, level)); +} + +void CumulantK15BulkComp::run() +{ + int size_Mat = para->getParD(level)->size_Mat_SP; + int numberOfThreads = para->getParD(level)->numberofthreads; + + int Grid = (size_Mat / numberOfThreads) + 1; + int Grid1, Grid2; + if (Grid>512) + { + Grid1 = 512; + Grid2 = (Grid / Grid1) + 1; + } + else + { + Grid1 = 1; + Grid2 = Grid; + } + dim3 grid(Grid1, Grid2); + dim3 threads(numberOfThreads, 1, 1); + + LB_Kernel_CumulantK15BulkComp <<< grid, threads >>>(para->getParD(level)->omega, + para->getParD(level)->geoSP, + para->getParD(level)->neighborX_SP, + para->getParD(level)->neighborY_SP, + para->getParD(level)->neighborZ_SP, + para->getParD(level)->d0SP.f[0], + para->getParD(level)->size_Mat_SP, + level, + para->getForcesDev(), + para->getParD(level)->evenOrOdd); + getLastCudaError("LB_Kernel_CumulantK15BulkComp execution failed"); +} + +CumulantK15BulkComp::CumulantK15BulkComp(std::shared_ptr<Parameter> para, int level) +{ + this->para = para; + this->level = level; + + myPreProcessorTypes.push_back(InitCompSP27); + + myKernelGroup = BasicKernel; +} + +CumulantK15BulkComp::CumulantK15BulkComp() +{ +} diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Bulk/CumulantK15BulkComp.h b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Bulk/CumulantK15BulkComp.h new file mode 100644 index 000000000..855b3380c --- /dev/null +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Bulk/CumulantK15BulkComp.h @@ -0,0 +1,17 @@ +#ifndef CUMULANT_ONE_COMP_BULK_SP27_H +#define CUMULANT_ONE_COMP_BULK_SP27_H + +#include "Kernel\KernelImp.h" + +class CumulantK15BulkComp : public KernelImp +{ +public: + static std::shared_ptr<CumulantK15BulkComp> getNewInstance(std::shared_ptr< Parameter> para, int level); + void run(); + +private: + CumulantK15BulkComp(); + CumulantK15BulkComp(std::shared_ptr< Parameter> para, int level); +}; + +#endif \ No newline at end of file diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneBulk/CumulantOneCompBulkSP27_Device.cu b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Bulk/CumulantK15BulkComp_Device.cu similarity index 99% rename from src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneBulk/CumulantOneCompBulkSP27_Device.cu rename to src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Bulk/CumulantK15BulkComp_Device.cu index f9c0f1949..dfda29010 100644 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneBulk/CumulantOneCompBulkSP27_Device.cu +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Bulk/CumulantK15BulkComp_Device.cu @@ -3,7 +3,7 @@ #include "Core/RealConstants.h" #include "math.h" -extern "C" __global__ void LB_Kernel_Cum_One_Comp_Bulk_SP_27(real omega, +extern "C" __global__ void LB_Kernel_CumulantK15BulkComp(real omega, unsigned int* bcMatD, unsigned int* neighborX, unsigned int* neighborY, diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneBulk/CumulantOneCompBulkSP27_Device.cuh b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Bulk/CumulantK15BulkComp_Device.cuh similarity index 81% rename from src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneBulk/CumulantOneCompBulkSP27_Device.cuh rename to src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Bulk/CumulantK15BulkComp_Device.cuh index f44daef00..a1a3aca96 100644 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneBulk/CumulantOneCompBulkSP27_Device.cuh +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Bulk/CumulantK15BulkComp_Device.cuh @@ -5,7 +5,7 @@ #include <curand.h> -extern "C" __global__ void LB_Kernel_Cum_One_Comp_Bulk_SP_27(real omega, +extern "C" __global__ void LB_Kernel_CumulantK15BulkComp(real omega, unsigned int* bcMatD, unsigned int* neighborX, unsigned int* neighborY, diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016Bulk/package.include b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Bulk/package.include similarity index 100% rename from src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016Bulk/package.include rename to src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Bulk/package.include diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Sponge/CumulantK15SpongeComp.cu b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Sponge/CumulantK15SpongeComp.cu new file mode 100644 index 000000000..79274e5c1 --- /dev/null +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Sponge/CumulantK15SpongeComp.cu @@ -0,0 +1,57 @@ +#include "CumulantK15SpongeComp.h" + +#include "CumulantK15SpongeComp_Device.cuh" +#include "Parameter\Parameter.h" + +std::shared_ptr<CumulantK15SpongeComp> CumulantK15SpongeComp::getNewInstance(std::shared_ptr<Parameter> para, int level) +{ + return std::shared_ptr<CumulantK15SpongeComp>(new CumulantK15SpongeComp(para, level)); +} + +void CumulantK15SpongeComp::run() +{ + int size_Mat = para->getParD(level)->size_Mat_SP; + int numberOfThreads = para->getParD(level)->numberofthreads; + + int Grid = (size_Mat / numberOfThreads) + 1; + int Grid1, Grid2; + if (Grid>512) + { + Grid1 = 512; + Grid2 = (Grid / Grid1) + 1; + } + else + { + Grid1 = 1; + Grid2 = Grid; + } + dim3 grid(Grid1, Grid2); + dim3 threads(numberOfThreads, 1, 1); + + LB_Kernel_CumulantK15SpongeComp <<< grid, threads >>>( para->getParD(level)->omega, + para->getParD(level)->geoSP, + para->getParD(level)->neighborX_SP, + para->getParD(level)->neighborY_SP, + para->getParD(level)->neighborZ_SP, + para->getParD(level)->coordX_SP, + para->getParD(level)->coordY_SP, + para->getParD(level)->coordZ_SP, + para->getParD(level)->d0SP.f[0], + para->getParD(level)->size_Mat_SP, + para->getParD(level)->evenOrOdd); + getLastCudaError("LB_Kernel_CumulantK15SpongeComp execution failed"); +} + +CumulantK15SpongeComp::CumulantK15SpongeComp(std::shared_ptr<Parameter> para, int level) +{ + this->para = para; + this->level = level; + + myPreProcessorTypes.push_back(InitCompSP27); + + myKernelGroup = BasicKernel; +} + +CumulantK15SpongeComp::CumulantK15SpongeComp() +{ +} diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Sponge/CumulantK15SpongeComp.h b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Sponge/CumulantK15SpongeComp.h new file mode 100644 index 000000000..8581fdd6e --- /dev/null +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Sponge/CumulantK15SpongeComp.h @@ -0,0 +1,17 @@ +#ifndef CUMULANT_ONE_COMP_SPONGE_SP27_H +#define CUMULANT_ONE_COMP_SPONGE_SP27_H + +#include "Kernel\KernelImp.h" + +class CumulantK15SpongeComp : public KernelImp +{ +public: + static std::shared_ptr<CumulantK15SpongeComp> getNewInstance(std::shared_ptr< Parameter> para, int level); + void run(); + +private: + CumulantK15SpongeComp(); + CumulantK15SpongeComp(std::shared_ptr< Parameter> para, int level); + +}; +#endif \ No newline at end of file diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneSponge/CumulantOneCompSpongeSP27_Device.cu b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Sponge/CumulantK15SpongeComp_Device.cu similarity index 99% rename from src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneSponge/CumulantOneCompSpongeSP27_Device.cu rename to src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Sponge/CumulantK15SpongeComp_Device.cu index 64c3e49a2..3a916bb89 100644 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneSponge/CumulantOneCompSpongeSP27_Device.cu +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Sponge/CumulantK15SpongeComp_Device.cu @@ -3,7 +3,7 @@ #include "Core/RealConstants.h" #include "math.h" -extern "C" __global__ void LB_Kernel_Cum_One_Comp_Sponge_SP_27(real omegaIn, +extern "C" __global__ void LB_Kernel_CumulantK15SpongeComp(real omegaIn, unsigned int* bcMatD, unsigned int* neighborX, unsigned int* neighborY, diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneSponge/CumulantOneCompSpongeSP27_Device.cuh b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Sponge/CumulantK15SpongeComp_Device.cuh similarity index 81% rename from src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneSponge/CumulantOneCompSpongeSP27_Device.cuh rename to src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Sponge/CumulantK15SpongeComp_Device.cuh index 87cf77963..b884e835f 100644 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneSponge/CumulantOneCompSpongeSP27_Device.cuh +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Sponge/CumulantK15SpongeComp_Device.cuh @@ -4,7 +4,7 @@ #include <DataTypes.h> #include <curand.h> -extern "C" __global__ void LB_Kernel_Cum_One_Comp_Sponge_SP_27(real omega, +extern "C" __global__ void LB_Kernel_CumulantK15SpongeComp(real omega, unsigned int* bcMatD, unsigned int* neighborX, unsigned int* neighborY, diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF3/package.include b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Sponge/package.include similarity index 100% rename from src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF3/package.include rename to src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Sponge/package.include diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17/CumulantK17Comp.cu b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17/CumulantK17Comp.cu new file mode 100644 index 000000000..86704c614 --- /dev/null +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17/CumulantK17Comp.cu @@ -0,0 +1,53 @@ +#include "CumulantK17Comp.h" + +#include "Parameter/Parameter.h" +#include "CumulantK17Comp_Device.cuh" + +std::shared_ptr<CumulantK17Comp> CumulantK17Comp::getNewInstance(std::shared_ptr<Parameter> para, int level) +{ + return std::shared_ptr<CumulantK17Comp>(new CumulantK17Comp(para,level)); +} + +void CumulantK17Comp::run() +{ + int numberOfThreads = para->getParD(level)->numberofthreads; + int size_Mat = para->getParD(level)->size_Mat_SP; + + int Grid = (size_Mat / numberOfThreads) + 1; + int Grid1, Grid2; + if (Grid>512) + { + Grid1 = 512; + Grid2 = (Grid / Grid1) + 1; + } + else + { + Grid1 = 1; + Grid2 = Grid; + } + dim3 grid(Grid1, Grid2); + dim3 threads(numberOfThreads, 1, 1); + + LB_Kernel_CumulantK17Comp <<< grid, threads >>>(para->getParD(level)->omega, + para->getParD(level)->geoSP, + para->getParD(level)->neighborX_SP, + para->getParD(level)->neighborY_SP, + para->getParD(level)->neighborZ_SP, + para->getParD(level)->d0SP.f[0], + para->getParD(level)->size_Mat_SP, + level, + para->getForcesDev(), + para->getQuadricLimitersDev(), + para->getParD(level)->evenOrOdd); + getLastCudaError("LB_Kernel_CumulantK17Comp execution failed"); +} + +CumulantK17Comp::CumulantK17Comp(std::shared_ptr<Parameter> para, int level) +{ + this->para = para; + this->level = level; + + myPreProcessorTypes.push_back(InitCompSP27); + + myKernelGroup = BasicKernel; +} \ No newline at end of file diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17/CumulantK17Comp.h b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17/CumulantK17Comp.h new file mode 100644 index 000000000..675827493 --- /dev/null +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17/CumulantK17Comp.h @@ -0,0 +1,17 @@ +#ifndef CUMULANT_AA2016_COMP_SP27_H +#define CUMULANT_AA2016_COMP_SP27_H + +#include "Kernel\KernelImp.h" + +class CumulantK17Comp : public KernelImp +{ +public: + static std::shared_ptr<CumulantK17Comp> getNewInstance(std::shared_ptr< Parameter> para, int level); + void run(); + +private: + CumulantK17Comp(); + CumulantK17Comp(std::shared_ptr< Parameter> para, int level); +}; + +#endif diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016/CumulantAA2016CompSP27_Device.cu b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17/CumulantK17Comp_Device.cu similarity index 99% rename from src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016/CumulantAA2016CompSP27_Device.cu rename to src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17/CumulantK17Comp_Device.cu index e8d8d2f38..9d5f512bd 100644 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016/CumulantAA2016CompSP27_Device.cu +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17/CumulantK17Comp_Device.cu @@ -4,7 +4,7 @@ #include "math.h" -extern "C" __global__ void LB_Kernel_Kum_AA2016_Comp_SP_27(real omega, +extern "C" __global__ void LB_Kernel_CumulantK17Comp(real omega, unsigned int* bcMatD, unsigned int* neighborX, unsigned int* neighborY, diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17/CumulantK17Comp_Device.cuh b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17/CumulantK17Comp_Device.cuh new file mode 100644 index 000000000..f0379ac75 --- /dev/null +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17/CumulantK17Comp_Device.cuh @@ -0,0 +1,18 @@ +#ifndef LB_Kernel_Kum_AA2016_Comp_SP_27_H +#define LB_Kernel_Kum_AA2016_Comp_SP_27_H + +#include <DataTypes.h> +#include <curand.h> + +extern "C" __global__ void LB_Kernel_CumulantK17Comp( real omega, + unsigned int* bcMatD, + unsigned int* neighborX, + unsigned int* neighborY, + unsigned int* neighborZ, + real* DDStart, + int size_Mat, + int level, + real* forces, + real* quadricLimiters, + bool EvenOrOdd); +#endif diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF32018/package.include b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17/package.include similarity index 100% rename from src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF32018/package.include rename to src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17/package.include diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016Bulk/CumulantAA2016CompBulkSP27.cu b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17Bulk/CumulantK17BulkComp.cu similarity index 57% rename from src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016Bulk/CumulantAA2016CompBulkSP27.cu rename to src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17Bulk/CumulantK17BulkComp.cu index 868b6b763..8a3432c86 100644 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016Bulk/CumulantAA2016CompBulkSP27.cu +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17Bulk/CumulantK17BulkComp.cu @@ -1,14 +1,14 @@ -#include "CumulantAA2016CompBulkSP27.h" +#include "CumulantK17BulkComp.h" -#include "CumulantAA2016CompBulkSP27_Device.cuh" +#include "CumulantK17BulkComp_Device.cuh" #include "Parameter\Parameter.h" -std::shared_ptr<CumulantAA2016CompBulkSP27> CumulantAA2016CompBulkSP27::getNewInstance(std::shared_ptr<Parameter> para, int level) +std::shared_ptr<CumulantK17BulkComp> CumulantK17BulkComp::getNewInstance(std::shared_ptr<Parameter> para, int level) { - return std::shared_ptr<CumulantAA2016CompBulkSP27>(new CumulantAA2016CompBulkSP27(para, level)); + return std::shared_ptr<CumulantK17BulkComp>(new CumulantK17BulkComp(para, level)); } -void CumulantAA2016CompBulkSP27::run() +void CumulantK17BulkComp::run() { int size_Array = para->getParD(level)->size_Array_SP; int numberOfThreads = para->getParD(level)->numberofthreads; @@ -17,7 +17,7 @@ void CumulantAA2016CompBulkSP27::run() dim3 grid(Grid, 1, 1); dim3 threads(numberOfThreads, 1, 1); - LB_Kernel_Cum_AA2016_Comp_Bulk_SP_27 << < grid, threads >> >( para->getParD(level)->omega, + LB_Kernel_CumulantK17BulkComp << < grid, threads >> >( para->getParD(level)->omega, para->getParD(level)->geoSP, para->getParD(level)->neighborX_SP, para->getParD(level)->neighborY_SP, @@ -28,10 +28,10 @@ void CumulantAA2016CompBulkSP27::run() para->getForcesDev(), para->getQuadricLimitersDev(), para->getParD(level)->evenOrOdd); - getLastCudaError("LB_Kernel_Kum_AA2016_Comp_Bulk_SP_27 execution failed"); + getLastCudaError("LB_Kernel_CumulantK17BulkComp execution failed"); } -CumulantAA2016CompBulkSP27::CumulantAA2016CompBulkSP27(std::shared_ptr<Parameter> para, int level) +CumulantK17BulkComp::CumulantK17BulkComp(std::shared_ptr<Parameter> para, int level) { this->para = para; this->level = level; @@ -41,6 +41,6 @@ CumulantAA2016CompBulkSP27::CumulantAA2016CompBulkSP27(std::shared_ptr<Parameter myKernelGroup = BasicKernel; } -CumulantAA2016CompBulkSP27::CumulantAA2016CompBulkSP27() +CumulantK17BulkComp::CumulantK17BulkComp() { } diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17Bulk/CumulantK17BulkComp.h b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17Bulk/CumulantK17BulkComp.h new file mode 100644 index 000000000..9e24c45e8 --- /dev/null +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17Bulk/CumulantK17BulkComp.h @@ -0,0 +1,17 @@ +#ifndef CUMULANT_AA2016_COMP_BULK_SP27_H +#define CUMULANT_AA2016_COMP_BULK_SP27_H + +#include "Kernel\KernelImp.h" + +class CumulantK17BulkComp : public KernelImp +{ +public: + static std::shared_ptr<CumulantK17BulkComp> getNewInstance(std::shared_ptr< Parameter> para, int level); + void run(); + +private: + CumulantK17BulkComp(); + CumulantK17BulkComp(std::shared_ptr< Parameter> para, int level); + +}; +#endif \ No newline at end of file diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016Bulk/CumulantAA2016CompBulkSP27_Device.cu b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17Bulk/CumulantK17BulkComp_Device.cu similarity index 99% rename from src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016Bulk/CumulantAA2016CompBulkSP27_Device.cu rename to src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17Bulk/CumulantK17BulkComp_Device.cu index 3336f0554..653fd9f7b 100644 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016Bulk/CumulantAA2016CompBulkSP27_Device.cu +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17Bulk/CumulantK17BulkComp_Device.cu @@ -3,7 +3,7 @@ #include "Core/RealConstants.h" #include "math.h" -extern "C" __global__ void LB_Kernel_Cum_AA2016_Comp_Bulk_SP_27(real omega, +extern "C" __global__ void LB_Kernel_CumulantK17BulkComp(real omega, unsigned int* bcMatD, unsigned int* neighborX, unsigned int* neighborY, diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016Bulk/CumulantAA2016CompBulkSP27_Device.cuh b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17Bulk/CumulantK17BulkComp_Device.cuh similarity index 82% rename from src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016Bulk/CumulantAA2016CompBulkSP27_Device.cuh rename to src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17Bulk/CumulantK17BulkComp_Device.cuh index 129ecb340..6f0744715 100644 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016Bulk/CumulantAA2016CompBulkSP27_Device.cuh +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17Bulk/CumulantK17BulkComp_Device.cuh @@ -4,7 +4,7 @@ #include <DataTypes.h> #include <curand.h> -extern "C" __global__ void LB_Kernel_Cum_AA2016_Comp_Bulk_SP_27(real omega, +extern "C" __global__ void LB_Kernel_CumulantK17BulkComp(real omega, unsigned int* bcMatD, unsigned int* neighborX, unsigned int* neighborY, diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOne/package.include b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17Bulk/package.include similarity index 100% rename from src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOne/package.include rename to src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17Bulk/package.include diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF3/CumulantF3CompSP27.cu b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK18/CumulantK18Comp.cu similarity index 67% rename from src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF3/CumulantF3CompSP27.cu rename to src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK18/CumulantK18Comp.cu index 907e9e686..f5b9bbe63 100644 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF3/CumulantF3CompSP27.cu +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK18/CumulantK18Comp.cu @@ -1,15 +1,15 @@ -#include "CumulantF3CompSP27.h" +#include "CumulantK18Comp.h" -#include "CumulantF3CompSP27_Device.cuh" +#include "CumulantK18Comp_Device.cuh" #include "Parameter/Parameter.h" -std::shared_ptr<CumulantF3CompSP27> CumulantF3CompSP27::getNewInstance(std::shared_ptr<Parameter> para, int level) +std::shared_ptr<CumulantK18Comp> CumulantK18Comp::getNewInstance(std::shared_ptr<Parameter> para, int level) { - return std::shared_ptr<CumulantF3CompSP27>(new CumulantF3CompSP27(para, level)); + return std::shared_ptr<CumulantK18Comp>(new CumulantK18Comp(para, level)); } -void CumulantF3CompSP27::run() +void CumulantK18Comp::run() { int numberOfThreads = para->getParD(level)->numberofthreads; int size_Mat = para->getParD(level)->size_Mat_SP; @@ -29,7 +29,7 @@ void CumulantF3CompSP27::run() dim3 grid(Grid1, Grid2); dim3 threads(numberOfThreads, 1, 1); - LB_Kernel_Cumulant_D3Q27F3 << < grid, threads >> >( para->getParD(level)->omega, + LB_Kernel_CumulantK18Comp << < grid, threads >> >( para->getParD(level)->omega, para->getParD(level)->geoSP, para->getParD(level)->neighborX_SP, para->getParD(level)->neighborY_SP, @@ -41,10 +41,10 @@ void CumulantF3CompSP27::run() para->getForcesDev(), para->getQuadricLimitersDev(), para->getParD(level)->evenOrOdd); - getLastCudaError("LB_Kernel_Cumulant_D3Q27F3 execution failed"); + getLastCudaError("LB_Kernel_CumulantK18Comp execution failed"); } -CumulantF3CompSP27::CumulantF3CompSP27(std::shared_ptr<Parameter> para, int level) +CumulantK18Comp::CumulantK18Comp(std::shared_ptr<Parameter> para, int level) { this->para = para; this->level = level; diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK18/CumulantK18Comp.h b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK18/CumulantK18Comp.h new file mode 100644 index 000000000..e8612c86c --- /dev/null +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK18/CumulantK18Comp.h @@ -0,0 +1,17 @@ +#ifndef CUMULANT_F3_COMP_SP27_H +#define CUMULANT_F3_COMP_SP27_H + +#include "Kernel\KernelImp.h" + +class CumulantK18Comp : public KernelImp +{ +public: + static std::shared_ptr< CumulantK18Comp> getNewInstance(std::shared_ptr< Parameter> para, int level); + void run(); + +private: + CumulantK18Comp(); + CumulantK18Comp(std::shared_ptr< Parameter> para, int level); +}; + +#endif \ No newline at end of file diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF3/CumulantF3CompSP27_Device.cu b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK18/CumulantK18Comp_Device.cu similarity index 99% rename from src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF3/CumulantF3CompSP27_Device.cu rename to src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK18/CumulantK18Comp_Device.cu index 4dd6490d6..f882ecc2a 100644 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF3/CumulantF3CompSP27_Device.cu +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK18/CumulantK18Comp_Device.cu @@ -1,11 +1,9 @@ -#include "CumulantF3CompSP27_Device.cuh" - #include "LBM/LB.h" #include "LBM/D3Q27.h" #include "Core/RealConstants.h" #include "math.h" -extern "C" __global__ void LB_Kernel_Cumulant_D3Q27F3( +extern "C" __global__ void LB_Kernel_CumulantK18Comp( real omega, unsigned int* bcMatD, unsigned int* neighborX, diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF3/CumulantF3CompSP27_Device.cuh b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK18/CumulantK18Comp_Device.cuh similarity index 88% rename from src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF3/CumulantF3CompSP27_Device.cuh rename to src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK18/CumulantK18Comp_Device.cuh index 04fa719b5..45bdcf958 100644 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF3/CumulantF3CompSP27_Device.cuh +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK18/CumulantK18Comp_Device.cuh @@ -4,7 +4,7 @@ #include <DataTypes.h> #include <curand.h> -extern "C" __global__ void LB_Kernel_Cumulant_D3Q27F3( real omega, +extern "C" __global__ void LB_Kernel_CumulantK18Comp( real omega, unsigned int* bcMatD, unsigned int* neighborX, unsigned int* neighborY, diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneBulk/package.include b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK18/package.include similarity index 100% rename from src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneBulk/package.include rename to src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK18/package.include diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF32018/CumulantF32018CompSP27.cu b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK20/CumulantK20Comp.cu similarity index 65% rename from src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF32018/CumulantF32018CompSP27.cu rename to src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK20/CumulantK20Comp.cu index afa545011..086a8e205 100644 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF32018/CumulantF32018CompSP27.cu +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK20/CumulantK20Comp.cu @@ -1,15 +1,15 @@ -#include "CumulantF32018CompSP27.h" +#include "CumulantK20Comp.h" -#include "CumulantF32018CompSP27_Device.cuh" +#include "CumulantK20Comp_Device.cuh" #include "Parameter/Parameter.h" -std::shared_ptr<CumulantF32018CompSP27> CumulantF32018CompSP27::getNewInstance(std::shared_ptr<Parameter> para, int level) +std::shared_ptr<CumulantK20Comp> CumulantK20Comp::getNewInstance(std::shared_ptr<Parameter> para, int level) { - return std::shared_ptr<CumulantF32018CompSP27>(new CumulantF32018CompSP27(para, level)); + return std::shared_ptr<CumulantK20Comp>(new CumulantK20Comp(para, level)); } -void CumulantF32018CompSP27::run() +void CumulantK20Comp::run() { int numberOfThreads = para->getParD(level)->numberofthreads; int size_Mat = para->getParD(level)->size_Mat_SP; @@ -29,7 +29,7 @@ void CumulantF32018CompSP27::run() dim3 grid(Grid1, Grid2); dim3 threads(numberOfThreads, 1, 1); - LB_Kernel_Cumulant_D3Q27F3_2018 << < grid, threads >> >( para->getParD(level)->omega, + LB_Kernel_CumulantK20Comp << < grid, threads >> >( para->getParD(level)->omega, para->getParD(level)->geoSP, para->getParD(level)->neighborX_SP, para->getParD(level)->neighborY_SP, @@ -41,10 +41,10 @@ void CumulantF32018CompSP27::run() para->getForcesDev(), para->getQuadricLimitersDev(), para->getParD(level)->evenOrOdd); - getLastCudaError("LB_Kernel_Cumulant_D3Q27F3 execution failed"); + getLastCudaError("LB_Kernel_CumulantK20Comp execution failed"); } -CumulantF32018CompSP27::CumulantF32018CompSP27(std::shared_ptr<Parameter> para, int level) +CumulantK20Comp::CumulantK20Comp(std::shared_ptr<Parameter> para, int level) { this->para = para; this->level = level; diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK20/CumulantK20Comp.h b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK20/CumulantK20Comp.h new file mode 100644 index 000000000..f7f9ce8a2 --- /dev/null +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK20/CumulantK20Comp.h @@ -0,0 +1,16 @@ +#ifndef CUMULANT_F32018_COMP_SP27_H +#define CUMULANT_F32018_COMP_SP27_H + +#include "Kernel\KernelImp.h" + +class CumulantK20Comp : public KernelImp +{ +public: + static std::shared_ptr< CumulantK20Comp> getNewInstance(std::shared_ptr< Parameter> para, int level); + void run(); + +private: + CumulantK20Comp(); + CumulantK20Comp(std::shared_ptr< Parameter> para, int level); +}; +#endif \ No newline at end of file diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF32018/CumulantF32018CompSP27_Device.cu b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK20/CumulantK20Comp_Device.cu similarity index 99% rename from src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF32018/CumulantF32018CompSP27_Device.cu rename to src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK20/CumulantK20Comp_Device.cu index 463ac1b6d..3324d52e9 100644 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF32018/CumulantF32018CompSP27_Device.cu +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK20/CumulantK20Comp_Device.cu @@ -1,11 +1,9 @@ -#include "CumulantF32018CompSP27_Device.cuh" - #include "LBM/LB.h" #include "LBM/D3Q27.h" #include "Core/RealConstants.h" #include "math.h" -extern "C" __global__ void LB_Kernel_Cumulant_D3Q27F3_2018( +extern "C" __global__ void LB_Kernel_CumulantK20Comp( real omega, unsigned int* bcMatD, unsigned int* neighborX, diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF32018/CumulantF32018CompSP27_Device.cuh b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK20/CumulantK20Comp_Device.cuh similarity index 88% rename from src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF32018/CumulantF32018CompSP27_Device.cuh rename to src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK20/CumulantK20Comp_Device.cuh index 503432b2f..19647055b 100644 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF32018/CumulantF32018CompSP27_Device.cuh +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK20/CumulantK20Comp_Device.cuh @@ -4,7 +4,7 @@ #include <DataTypes.h> #include <curand.h> -extern "C" __global__ void LB_Kernel_Cumulant_D3Q27F3_2018( real omega, +extern "C" __global__ void LB_Kernel_CumulantK20Comp( real omega, unsigned int* bcMatD, unsigned int* neighborX, unsigned int* neighborY, diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneSponge/package.include b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK20/package.include similarity index 100% rename from src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneSponge/package.include rename to src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK20/package.include diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOne/CumulantOneCompSP27.cu b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOne/CumulantOneCompSP27.cu deleted file mode 100644 index bb8bdc5c6..000000000 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOne/CumulantOneCompSP27.cu +++ /dev/null @@ -1,53 +0,0 @@ -#include "CumulantOneCompSP27.h" - -#include "CumulantOneCompSP27_Device.cuh" - -#include "Parameter/Parameter.h" - -std::shared_ptr<CumulantOneCompSP27> CumulantOneCompSP27::getNewInstance(std::shared_ptr<Parameter> para, int level) -{ - return std::shared_ptr<CumulantOneCompSP27>(new CumulantOneCompSP27(para, level)); -} - -void CumulantOneCompSP27::run() -{ - int numberOfThreads = para->getParD(level)->numberofthreads; - int size_Mat = para->getParD(level)->size_Mat_SP; - - int Grid = (size_Mat / numberOfThreads) + 1; - int Grid1, Grid2; - if (Grid > 512) - { - Grid1 = 512; - Grid2 = (Grid / Grid1) + 1; - } - else - { - Grid1 = 1; - Grid2 = Grid; - } - dim3 grid(Grid1, Grid2, 1); - dim3 threads(numberOfThreads, 1, 1); - - LB_Kernel_Kum_One_Comp_SP_27 << < grid, threads >> >( para->getParD(level)->omega, - para->getParD(level)->geoSP, - para->getParD(level)->neighborX_SP, - para->getParD(level)->neighborY_SP, - para->getParD(level)->neighborZ_SP, - para->getParD(level)->d0SP.f[0], - size_Mat, - level, - para->getForcesDev(), - para->getParD(level)->evenOrOdd); - getLastCudaError("LB_Kernel_Kum_New_Comp_SP_27 execution failed"); -} - -CumulantOneCompSP27::CumulantOneCompSP27(std::shared_ptr<Parameter> para, int level) -{ - this->para = para; - this->level = level; - - myPreProcessorTypes.push_back(InitCompSP27); - - myKernelGroup = BasicKernel; -} \ No newline at end of file diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOne/CumulantOneCompSP27.h b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOne/CumulantOneCompSP27.h deleted file mode 100644 index d5c175510..000000000 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOne/CumulantOneCompSP27.h +++ /dev/null @@ -1,16 +0,0 @@ -#ifndef CUMULANT_ONE_COMP_SP27_H -#define CUMULANT_ONE_COMP_SP27_H - -#include "Kernel\KernelImp.h" - -class CumulantOneCompSP27 : public KernelImp -{ -public: - static std::shared_ptr< CumulantOneCompSP27> getNewInstance(std::shared_ptr< Parameter> para, int level); - void run(); - -private: - CumulantOneCompSP27(); - CumulantOneCompSP27(std::shared_ptr< Parameter> para, int level); -}; -#endif \ No newline at end of file diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOne/CumulantOneCompSP27_Device.cuh b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOne/CumulantOneCompSP27_Device.cuh deleted file mode 100644 index 028a52251..000000000 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOne/CumulantOneCompSP27_Device.cuh +++ /dev/null @@ -1,17 +0,0 @@ -#ifndef LB_Kernel_Kum_ONE_COMP_SP_27_H -#define LB_Kernel_Kum_ONE_COMP_SP_27_H - -#include <DataTypes.h> -#include <curand.h> - -extern "C" __global__ void LB_Kernel_Kum_One_Comp_SP_27( real omega, - unsigned int* bcMatD, - unsigned int* neighborX, - unsigned int* neighborY, - unsigned int* neighborZ, - real* DDStart, - int size_Mat, - int level, - real* forces, - bool EvenOrOdd); -#endif \ No newline at end of file diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneBulk/CumulantOneCompBulkSP27.cu b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneBulk/CumulantOneCompBulkSP27.cu deleted file mode 100644 index 796b5c2a2..000000000 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneBulk/CumulantOneCompBulkSP27.cu +++ /dev/null @@ -1,56 +0,0 @@ -#include "CumulantOneCompBulkSP27.h" - -#include "CumulantOneCompBulkSP27_Device.cuh" -#include "Parameter\Parameter.h" - -std::shared_ptr<CumulantOneCompBulkSP27> CumulantOneCompBulkSP27::getNewInstance(std::shared_ptr<Parameter> para, int level) -{ - return std::shared_ptr<CumulantOneCompBulkSP27>(new CumulantOneCompBulkSP27(para, level)); -} - -void CumulantOneCompBulkSP27::run() -{ - int size_Mat = para->getParD(level)->size_Mat_SP; - int numberOfThreads = para->getParD(level)->numberofthreads; - - int Grid = (size_Mat / numberOfThreads) + 1; - int Grid1, Grid2; - if (Grid>512) - { - Grid1 = 512; - Grid2 = (Grid / Grid1) + 1; - } - else - { - Grid1 = 1; - Grid2 = Grid; - } - dim3 grid(Grid1, Grid2); - dim3 threads(numberOfThreads, 1, 1); - - LB_Kernel_Cum_One_Comp_Bulk_SP_27 << < grid, threads >> >(para->getParD(level)->omega, - para->getParD(level)->geoSP, - para->getParD(level)->neighborX_SP, - para->getParD(level)->neighborY_SP, - para->getParD(level)->neighborZ_SP, - para->getParD(level)->d0SP.f[0], - para->getParD(level)->size_Mat_SP, - level, - para->getForcesDev(), - para->getParD(level)->evenOrOdd); - getLastCudaError("LB_Kernel_Kum_New_Comp_Bulk_SP_27 execution failed"); -} - -CumulantOneCompBulkSP27::CumulantOneCompBulkSP27(std::shared_ptr<Parameter> para, int level) -{ - this->para = para; - this->level = level; - - myPreProcessorTypes.push_back(InitCompSP27); - - myKernelGroup = BasicKernel; -} - -CumulantOneCompBulkSP27::CumulantOneCompBulkSP27() -{ -} diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneBulk/CumulantOneCompBulkSP27.h b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneBulk/CumulantOneCompBulkSP27.h deleted file mode 100644 index 8c0873edd..000000000 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneBulk/CumulantOneCompBulkSP27.h +++ /dev/null @@ -1,17 +0,0 @@ -#ifndef CUMULANT_ONE_COMP_BULK_SP27_H -#define CUMULANT_ONE_COMP_BULK_SP27_H - -#include "Kernel\KernelImp.h" - -class CumulantOneCompBulkSP27 : public KernelImp -{ -public: - static std::shared_ptr<CumulantOneCompBulkSP27> getNewInstance(std::shared_ptr< Parameter> para, int level); - void run(); - -private: - CumulantOneCompBulkSP27(); - CumulantOneCompBulkSP27(std::shared_ptr< Parameter> para, int level); -}; - -#endif \ No newline at end of file diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneSponge/CumulantOneCompSpongeSP27.cu b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneSponge/CumulantOneCompSpongeSP27.cu deleted file mode 100644 index 1ea1bdba9..000000000 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneSponge/CumulantOneCompSpongeSP27.cu +++ /dev/null @@ -1,57 +0,0 @@ -#include "CumulantOneCompSpongeSP27.h" - -#include "CumulantOneCompSpongeSP27_Device.cuh" -#include "Parameter\Parameter.h" - -std::shared_ptr<CumulantOneCompSpongeSP27> CumulantOneCompSpongeSP27::getNewInstance(std::shared_ptr<Parameter> para, int level) -{ - return std::shared_ptr<CumulantOneCompSpongeSP27>(new CumulantOneCompSpongeSP27(para, level)); -} - -void CumulantOneCompSpongeSP27::run() -{ - int size_Mat = para->getParD(level)->size_Mat_SP; - int numberOfThreads = para->getParD(level)->numberofthreads; - - int Grid = (size_Mat / numberOfThreads) + 1; - int Grid1, Grid2; - if (Grid>512) - { - Grid1 = 512; - Grid2 = (Grid / Grid1) + 1; - } - else - { - Grid1 = 1; - Grid2 = Grid; - } - dim3 grid(Grid1, Grid2); - dim3 threads(numberOfThreads, 1, 1); - - LB_Kernel_Cum_One_Comp_Sponge_SP_27 << < grid, threads >> >( para->getParD(level)->omega, - para->getParD(level)->geoSP, - para->getParD(level)->neighborX_SP, - para->getParD(level)->neighborY_SP, - para->getParD(level)->neighborZ_SP, - para->getParD(level)->coordX_SP, - para->getParD(level)->coordY_SP, - para->getParD(level)->coordZ_SP, - para->getParD(level)->d0SP.f[0], - para->getParD(level)->size_Mat_SP, - para->getParD(level)->evenOrOdd); - getLastCudaError("LB_Kernel_Kum_New_Comp_Sponge_SP_27 execution failed"); -} - -CumulantOneCompSpongeSP27::CumulantOneCompSpongeSP27(std::shared_ptr<Parameter> para, int level) -{ - this->para = para; - this->level = level; - - myPreProcessorTypes.push_back(InitCompSP27); - - myKernelGroup = BasicKernel; -} - -CumulantOneCompSpongeSP27::CumulantOneCompSpongeSP27() -{ -} diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneSponge/CumulantOneCompSpongeSP27.h b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneSponge/CumulantOneCompSpongeSP27.h deleted file mode 100644 index 6fb9f0bad..000000000 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneSponge/CumulantOneCompSpongeSP27.h +++ /dev/null @@ -1,17 +0,0 @@ -#ifndef CUMULANT_ONE_COMP_SPONGE_SP27_H -#define CUMULANT_ONE_COMP_SPONGE_SP27_H - -#include "Kernel\KernelImp.h" - -class CumulantOneCompSpongeSP27 : public KernelImp -{ -public: - static std::shared_ptr<CumulantOneCompSpongeSP27> getNewInstance(std::shared_ptr< Parameter> para, int level); - void run(); - -private: - CumulantOneCompSpongeSP27(); - CumulantOneCompSpongeSP27(std::shared_ptr< Parameter> para, int level); - -}; -#endif \ No newline at end of file diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantOne/CumulantOneIncompSP27.cu b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantK15/CumulantK15Incomp.cu similarity index 58% rename from src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantOne/CumulantOneIncompSP27.cu rename to src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantK15/CumulantK15Incomp.cu index 10f80a78f..d6609827d 100644 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantOne/CumulantOneIncompSP27.cu +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantK15/CumulantK15Incomp.cu @@ -1,14 +1,14 @@ -#include "CumulantOneIncompSP27.h" +#include "CumulantK15Incomp.h" -#include "CumulantOneIncompSP27_Device.cuh" +#include "CumulantK15Incomp_Device.cuh" #include "Parameter\Parameter.h" -std::shared_ptr<CumulantOneIncompSP27> CumulantOneIncompSP27::getNewInstance(std::shared_ptr<Parameter> para, int level) +std::shared_ptr<CumulantK15Incomp> CumulantK15Incomp::getNewInstance(std::shared_ptr<Parameter> para, int level) { - return std::shared_ptr<CumulantOneIncompSP27>(new CumulantOneIncompSP27(para, level)); + return std::shared_ptr<CumulantK15Incomp>(new CumulantK15Incomp(para, level)); } -void CumulantOneIncompSP27::run() +void CumulantK15Incomp::run() { int size_Mat = para->getParD(level)->size_Mat_SP; int numberOfThreads = para->getParD(level)->numberofthreads; @@ -28,7 +28,7 @@ void CumulantOneIncompSP27::run() dim3 grid(Grid1, Grid2); dim3 threads(numberOfThreads, 1, 1); - LB_Kernel_Cum_One_Incomp_SP_27 << < grid, threads >> >( para->getParD(level)->omega, + LB_Kernel_CumulantK15Incomp <<< grid, threads >>>( para->getParD(level)->omega, para->getParD(level)->geoSP, para->getParD(level)->neighborX_SP, para->getParD(level)->neighborY_SP, @@ -36,10 +36,10 @@ void CumulantOneIncompSP27::run() para->getParD(level)->d0SP.f[0], para->getParD(level)->size_Mat_SP, para->getParD(level)->evenOrOdd); - getLastCudaError("LB_Kernel_Kum_New_SP_27 execution failed"); + getLastCudaError("LB_Kernel_CumulantK15Incomp execution failed"); } -CumulantOneIncompSP27::CumulantOneIncompSP27(std::shared_ptr<Parameter> para, int level) +CumulantK15Incomp::CumulantK15Incomp(std::shared_ptr<Parameter> para, int level) { this->para = para; this->level = level; @@ -49,6 +49,6 @@ CumulantOneIncompSP27::CumulantOneIncompSP27(std::shared_ptr<Parameter> para, in myKernelGroup = BasicKernel; } -CumulantOneIncompSP27::CumulantOneIncompSP27() +CumulantK15Incomp::CumulantK15Incomp() { } diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantK15/CumulantK15Incomp.h b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantK15/CumulantK15Incomp.h new file mode 100644 index 000000000..eee935847 --- /dev/null +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantK15/CumulantK15Incomp.h @@ -0,0 +1,16 @@ +#ifndef CUMULANT_ONE_INCOMP_SP27_H +#define CUMULANT_ONE_INCOMP_SP27_H + +#include "Kernel\KernelImp.h" + +class CumulantK15Incomp : public KernelImp +{ +public: + static std::shared_ptr<CumulantK15Incomp> getNewInstance(std::shared_ptr< Parameter> para, int level); + void run(); + +private: + CumulantK15Incomp(); + CumulantK15Incomp(std::shared_ptr< Parameter> para, int level); +}; +#endif \ No newline at end of file diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantOne/CumulantOneIncompSP27_Device.cu b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantK15/CumulantK15Incomp_Device.cu similarity index 99% rename from src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantOne/CumulantOneIncompSP27_Device.cu rename to src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantK15/CumulantK15Incomp_Device.cu index 054c243d9..829a9e0c3 100644 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantOne/CumulantOneIncompSP27_Device.cu +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantK15/CumulantK15Incomp_Device.cu @@ -3,7 +3,7 @@ #include "Core/RealConstants.h" #include "math.h" -extern "C" __global__ void LB_Kernel_Cum_One_Incomp_SP_27(real omega, +extern "C" __global__ void LB_Kernel_CumulantK15Incomp(real omega, unsigned int* bcMatD, unsigned int* neighborX, unsigned int* neighborY, diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantOne/CumulantOneIncompSP27_Device.cuh b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantK15/CumulantK15Incomp_Device.cuh similarity index 80% rename from src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantOne/CumulantOneIncompSP27_Device.cuh rename to src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantK15/CumulantK15Incomp_Device.cuh index d56d6d558..beb4f50be 100644 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantOne/CumulantOneIncompSP27_Device.cuh +++ b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantK15/CumulantK15Incomp_Device.cuh @@ -4,7 +4,7 @@ #include <DataTypes.h> #include <curand.h> -extern "C" __global__ void LB_Kernel_Cum_One_Incomp_SP_27(real s9, +extern "C" __global__ void LB_Kernel_CumulantK15Incomp(real s9, unsigned int* bcMatD, unsigned int* neighborX, unsigned int* neighborY, diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantOne/package.include b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantK15/package.include similarity index 100% rename from src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantOne/package.include rename to src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantK15/package.include diff --git a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantOne/CumulantOneIncompSP27.h b/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantOne/CumulantOneIncompSP27.h deleted file mode 100644 index 70a3e9b88..000000000 --- a/src/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantOne/CumulantOneIncompSP27.h +++ /dev/null @@ -1,16 +0,0 @@ -#ifndef CUMULANT_ONE_INCOMP_SP27_H -#define CUMULANT_ONE_INCOMP_SP27_H - -#include "Kernel\KernelImp.h" - -class CumulantOneIncompSP27 : public KernelImp -{ -public: - static std::shared_ptr<CumulantOneIncompSP27> getNewInstance(std::shared_ptr< Parameter> para, int level); - void run(); - -private: - CumulantOneIncompSP27(); - CumulantOneIncompSP27(std::shared_ptr< Parameter> para, int level); -}; -#endif \ No newline at end of file diff --git a/src/VirtualFluids_GPU/Kernel/Utilities/KernelFactory/KernelFactoryImp.cpp b/src/VirtualFluids_GPU/Kernel/Utilities/KernelFactory/KernelFactoryImp.cpp index 11347f229..04945ccbc 100644 --- a/src/VirtualFluids_GPU/Kernel/Utilities/KernelFactory/KernelFactoryImp.cpp +++ b/src/VirtualFluids_GPU/Kernel/Utilities/KernelFactory/KernelFactoryImp.cpp @@ -6,14 +6,14 @@ #include "Kernel/Kernels/BasicKernels/Advection/Compressible/BGKPlus/BGKPlusCompSP27.h" #include "Kernel/Kernels/BasicKernels/Advection/Compressible/Cascade/CascadeCompSP27.h" #include "Kernel/Kernels/BasicKernels/Advection/Compressible/Cumulant/CumulantCompSP27.h" -#include "Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016/CumulantAA2016CompSP27.h" -#include "Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAA2016Bulk/CumulantAA2016CompBulkSP27.h" +#include "Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17/CumulantK17Comp.h" +#include "Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK17Bulk/CumulantK17BulkComp.h" #include "Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAll4/CumulantAll4CompSP27.h" -#include "Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF3/CumulantF3CompSP27.h" -#include "Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantF32018/CumulantF32018CompSP27.h" -#include "Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOne/CumulantOneCompSP27.h" -#include "Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneBulk/CumulantOneCompBulkSP27.h" -#include "Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantOneSponge/CumulantOneCompSpongeSP27.h" +#include "Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK18/CumulantK18Comp.h" +#include "Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK20/CumulantK20Comp.h" +#include "Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15/CumulantK15Comp.h" +#include "Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Bulk/CumulantK15BulkComp.h" +#include "Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantK15Sponge/CumulantK15SpongeComp.h" #include "Kernel/Kernels/BasicKernels/Advection/Compressible/MRT/MRTCompSP27.h" #include "Kernel/Kernels/BasicKernels/Advection/Incompressible/BGK/BGKIncompSP27.h" @@ -21,7 +21,7 @@ #include "Kernel/Kernels/BasicKernels/Advection/Incompressible/Cascade/CascadeIncompSP27.h" #include "Kernel/Kernels/BasicKernels/Advection/Incompressible/Cumulant1hSP27/Cumulant1hIncompSP27.h" #include "Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantIsoSP27/CumulantIsoIncompSP27.h" -#include "Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantOne/CumulantOneIncompSP27.h" +#include "Kernel/Kernels/BasicKernels/Advection/Incompressible/CumulantK15/CumulantK15Incomp.h" #include "Kernel/Kernels/BasicKernels/Advection/Incompressible/MRT/MRTIncompSP27.h" #include "Kernel/Kernels/BasicKernels/AdvectionDiffusion/Compressible/Mod27/ADComp27/ADComp27.h" @@ -109,36 +109,36 @@ std::shared_ptr<Kernel> KernelFactoryImp::makeKernel(std::shared_ptr<Parameter> newKernel = CumulantCompSP27::getNewInstance(para, level); checkStrategy = AdvecCompStrategy::getInstance(); break; - case LB_CumulantAA2016CompSP27: - newKernel = CumulantAA2016CompSP27::getNewInstance(para, level); + case LB_CumulantK17Comp: + newKernel = CumulantK17Comp::getNewInstance(para, level); checkStrategy = AdvecCompStrategy::getInstance(); break; - case LB_CumulantAA2016CompBulkSP27: - newKernel = CumulantAA2016CompBulkSP27::getNewInstance(para, level); + case LB_CumulantK17BulkComp: + newKernel = CumulantK17BulkComp::getNewInstance(para, level); checkStrategy = AdvecCompStrategy::getInstance(); break; case LB_CumulantAll4CompSP27: newKernel = CumulantAll4CompSP27::getNewInstance(para, level); checkStrategy = AdvecCompStrategy::getInstance(); break; - case LB_CumulantF3CompSP27: - newKernel = CumulantF3CompSP27::getNewInstance(para, level); + case LB_CumulantK18Comp: + newKernel = CumulantK18Comp::getNewInstance(para, level); checkStrategy = AdvecCompStrategy::getInstance(); break; - case LB_CumulantF32018CompSP27: - newKernel = CumulantF32018CompSP27::getNewInstance(para, level); + case LB_CumulantK20Comp: + newKernel = CumulantK20Comp::getNewInstance(para, level); checkStrategy = AdvecCompStrategy::getInstance(); break; - case LB_CumulantOneCompSP27: - newKernel = CumulantOneCompSP27::getNewInstance(para, level); + case LB_CumulantK15Comp: + newKernel = CumulantK15Comp::getNewInstance(para, level); checkStrategy = AdvecCompStrategy::getInstance(); break; - case LB_CumulantOneCompBulkSP27: - newKernel = CumulantOneCompSP27::getNewInstance(para, level); + case LB_CumulantK15BulkComp: + newKernel = CumulantK15Comp::getNewInstance(para, level); checkStrategy = AdvecCompStrategy::getInstance(); break; - case LB_CumulantOneCompSpongeSP27: - newKernel = CumulantOneCompSpongeSP27::getNewInstance(para, level); + case LB_CumulantK15SpongeComp: + newKernel = CumulantK15SpongeComp::getNewInstance(para, level); checkStrategy = AdvecCompStrategy::getInstance(); break; case LB_MRTCompSP27: @@ -167,8 +167,8 @@ std::shared_ptr<Kernel> KernelFactoryImp::makeKernel(std::shared_ptr<Parameter> newKernel = CumulantIsoIncompSP27::getNewInstance(para, level); checkStrategy = AdvecIncompStrategy::getInstance(); break; - case LB_CumulantOneIncompSP27: - newKernel = CumulantOneIncompSP27::getNewInstance(para, level); + case LB_CumulantK15Incomp: + newKernel = CumulantK15Incomp::getNewInstance(para, level); checkStrategy = AdvecIncompStrategy::getInstance(); break; case LB_MRTIncompSP27: @@ -208,7 +208,7 @@ std::shared_ptr<Kernel> KernelFactoryImp::makeKernel(std::shared_ptr<Parameter> return newKernel; } else - throw std::exception("Kernelfactory does not know the KernelType."); + throw std::exception("KernelFactory does not know the KernelType."); } @@ -245,7 +245,7 @@ std::shared_ptr<ADKernel> KernelFactoryImp::makeAdvDifKernel(std::shared_ptr<Par return newKernel; } else - throw std::exception("Kernelfactory does not know the KernelType."); + throw std::exception("KernelFactory does not know the KernelType."); } KernelFactoryImp::KernelFactoryImp() diff --git a/src/VirtualFluids_GPU/Kernel/Utilities/KernelType.h b/src/VirtualFluids_GPU/Kernel/Utilities/KernelType.h index 6d7b04477..e1d14920f 100644 --- a/src/VirtualFluids_GPU/Kernel/Utilities/KernelType.h +++ b/src/VirtualFluids_GPU/Kernel/Utilities/KernelType.h @@ -7,14 +7,14 @@ enum KernelType { LB_BGKPlusCompSP27, LB_CascadeCompSP27, LB_CumulantCompSP27, - LB_CumulantAA2016CompSP27, - LB_CumulantAA2016CompBulkSP27, + LB_CumulantK17Comp, + LB_CumulantK17BulkComp, LB_CumulantAll4CompSP27, - LB_CumulantF3CompSP27, - LB_CumulantF32018CompSP27, - LB_CumulantOneCompSP27, - LB_CumulantOneCompBulkSP27, - LB_CumulantOneCompSpongeSP27, + LB_CumulantK18Comp, + LB_CumulantK20Comp, + LB_CumulantK15Comp, + LB_CumulantK15BulkComp, + LB_CumulantK15SpongeComp, LB_MRTCompSP27, LB_BGKIncompSP27, @@ -22,7 +22,7 @@ enum KernelType { LB_CascadeIncompSP27, LB_Cumulant1hIncompSP27, LB_CumulantIsoIncompSP27, - LB_CumulantOneIncompSP27, + LB_CumulantK15Incomp, LB_MRTIncompSP27, LB_PMCumulantOneCompSP27, diff --git a/src/VirtualFluids_GPU/Kernel/Utilities/Mapper/KernelMapper/KernelMapper.cpp b/src/VirtualFluids_GPU/Kernel/Utilities/Mapper/KernelMapper/KernelMapper.cpp index 26aaf3fdd..fd9a9d317 100644 --- a/src/VirtualFluids_GPU/Kernel/Utilities/Mapper/KernelMapper/KernelMapper.cpp +++ b/src/VirtualFluids_GPU/Kernel/Utilities/Mapper/KernelMapper/KernelMapper.cpp @@ -24,21 +24,21 @@ KernelMapper::KernelMapper() myEnumMapper.addEnum(LB_BGKPlusCompSP27, "BGKPlusCompSP27"); myEnumMapper.addEnum(LB_CascadeCompSP27, "CascadeCompSP27"); myEnumMapper.addEnum(LB_CumulantCompSP27, "CumulantCompSP27"); - myEnumMapper.addEnum(LB_CumulantAA2016CompSP27, "CumulantAA2016CompSP27"); - myEnumMapper.addEnum(LB_CumulantAA2016CompBulkSP27, "CumulantAA2016CompBulkSP27"); + myEnumMapper.addEnum(LB_CumulantK17Comp, "CumulantK17Comp"); + myEnumMapper.addEnum(LB_CumulantK17BulkComp, "CumulantK17BulkComp"); myEnumMapper.addEnum(LB_CumulantAll4CompSP27, "CumulantAll4CompSP27"); - myEnumMapper.addEnum(LB_CumulantF3CompSP27, "CumulantF3CompSP27"); - myEnumMapper.addEnum(LB_CumulantF32018CompSP27, "CumulantF32018CompSP27"); - myEnumMapper.addEnum(LB_CumulantOneCompSP27, "CumulantOneCompSP27"); - myEnumMapper.addEnum(LB_CumulantOneCompBulkSP27, "CumulantOneCompBulkSP27"); - myEnumMapper.addEnum(LB_CumulantOneCompSpongeSP27, "CumulantOneCompSpongeSP27"); + myEnumMapper.addEnum(LB_CumulantK18Comp, "CumulantK18Comp"); + myEnumMapper.addEnum(LB_CumulantK20Comp, "CumulantK20Comp"); + myEnumMapper.addEnum(LB_CumulantK15Comp, "CumulantK15Comp"); + myEnumMapper.addEnum(LB_CumulantK15BulkComp, "CumulantK15BulkComp"); + myEnumMapper.addEnum(LB_CumulantK15SpongeComp, "CumulantK15SpongeComp"); myEnumMapper.addEnum(LB_MRTCompSP27, "MRTCompSP27"); myEnumMapper.addEnum(LB_BGKIncompSP27, "BGKIncompSP27"); myEnumMapper.addEnum(LB_BGKPlusIncompSP27, "BGKPlusIncompSP27"); myEnumMapper.addEnum(LB_CascadeIncompSP27, "CascadeIncompSP27"); myEnumMapper.addEnum(LB_Cumulant1hIncompSP27, "Cumulant1hIncompSP27"); myEnumMapper.addEnum(LB_CumulantIsoIncompSP27, "CumulantIsoIncompSP27"); - myEnumMapper.addEnum(LB_CumulantOneIncompSP27, "CumulantOneIncompSP27"); + myEnumMapper.addEnum(LB_CumulantK15Incomp, "CumulantK15Incomp"); myEnumMapper.addEnum(LB_MRTIncompSP27, "MRTIncompSP27"); myEnumMapper.addEnum(LB_PMCumulantOneCompSP27, "PMCumulantOneCompSP27"); myEnumMapper.addEnum(LB_WaleCumulantAA2016CompSP27, "WaleCumulantAA2016CompSP27"); diff --git a/src/VirtualFluids_GPU/Parameter/Parameter.cpp b/src/VirtualFluids_GPU/Parameter/Parameter.cpp index 1711dd342..b24ef75cf 100644 --- a/src/VirtualFluids_GPU/Parameter/Parameter.cpp +++ b/src/VirtualFluids_GPU/Parameter/Parameter.cpp @@ -504,7 +504,7 @@ Parameter::Parameter(SPtr<ConfigData> configData, Communicator* comm) if (configData->isMainKernelInConfigFile()) this->setMainKernel(kernelMapper->getEnum(configData->getMainKernel())); else - this->setMainKernel(kernelMapper->getEnum("CumulantOneCompSP27")); + this->setMainKernel(kernelMapper->getEnum("CumulantK15Comp")); if (configData->isMultiKernelOnInConfigFile()) this->setMultiKernelOn(configData->getMultiKernelOn()); @@ -537,7 +537,7 @@ Parameter::Parameter(SPtr<ConfigData> configData, Communicator* comm) std::vector<KernelType> tmp; for (int i = 0; i < this->getMaxLevel()+1; i++) { - tmp.push_back(kernelMapper->getEnum("CumulantOneCompSP27")); + tmp.push_back(kernelMapper->getEnum("CumulantK15Comp")); } this->setMultiKernel(tmp); } diff --git a/targets/apps/LBM/TGV_3D/TGV_3D.cpp b/targets/apps/LBM/TGV_3D/TGV_3D.cpp index 3915b77ce..6481947a1 100644 --- a/targets/apps/LBM/TGV_3D/TGV_3D.cpp +++ b/targets/apps/LBM/TGV_3D/TGV_3D.cpp @@ -92,7 +92,7 @@ uint gpuIndex = 0; bool useLimiter = false; -std::string kernel( "CumulantAA2016CompSP27" ); +std::string kernel( "CumulantK17Comp" ); std::string path("F:/Work/Computations/out/TaylorGreen3DNew/"); //LEGOLAS //std::string path("E:/DrivenCavity/results/"); //TESLA03 @@ -181,6 +181,8 @@ void multipleLevel(const std::string& configPath) _path << path; _path << kernel; + if( useLimiter ) _path << "_Limiter"; + path = _path.str(); ////////////////////////////////////////////////////////////////////////// -- GitLab