diff --git a/src/gpu/VirtualFluids_GPU/GPU/GPU_Kernels.cuh b/src/gpu/VirtualFluids_GPU/GPU/GPU_Kernels.cuh index 877390c822b4828b0007249be524d2534a2482f0..74ebf3bea73c221207d3dda7a6a2f29de083ffde 100644 --- a/src/gpu/VirtualFluids_GPU/GPU/GPU_Kernels.cuh +++ b/src/gpu/VirtualFluids_GPU/GPU/GPU_Kernels.cuh @@ -64,16 +64,17 @@ __global__ void LB_Kernel_Kum_New_Comp_SRT_SP_27( real* forces, bool EvenOrOdd); -__global__ void LB_Kernel_Cumulant_D3Q27All4(real omega, - unsigned int* bcMatD, - unsigned int* neighborX, - unsigned int* neighborY, - unsigned int* neighborZ, - real* DDStart, - unsigned long long numberOfLBnodes, - int level, - real* forces, - bool EvenOrOdd); +__global__ void K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants_Device( + real omega, + unsigned int* bcMatD, + unsigned int* neighborX, + unsigned int* neighborY, + unsigned int* neighborZ, + real* DDStart, + unsigned long long numberOfLBnodes, + int level, + real* forces, + bool EvenOrOdd); __global__ void LB_Kernel_Kum_AA2016_Comp_Bulk_SP_27(real omega, diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B15CompressibleNavierStokesBGKplus/B15CompressibleNavierStokesBGKplus.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B15CompressibleNavierStokesBGKplus/B15CompressibleNavierStokesBGKplus.cu new file mode 100644 index 0000000000000000000000000000000000000000..46d3c6edc413781826d8dc5147e9f46c3ae5a128 --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B15CompressibleNavierStokesBGKplus/B15CompressibleNavierStokesBGKplus.cu @@ -0,0 +1,40 @@ +#include "B15CompressibleNavierStokesBGKplus.h" + +#include "B15CompressibleNavierStokesBGKplus_Device.cuh" +#include "Parameter/Parameter.h" +#include "cuda/CudaGrid.h" + +std::shared_ptr<B15CompressibleNavierStokesBGKplus> B15CompressibleNavierStokesBGKplus::getNewInstance(std::shared_ptr<Parameter> para, int level) +{ + return std::shared_ptr<B15CompressibleNavierStokesBGKplus>(new B15CompressibleNavierStokesBGKplus(para, level)); +} + +void B15CompressibleNavierStokesBGKplus::run() +{ + vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes); + + B15CompressibleNavierStokesBGKplus_Device<<<grid.grid, grid.threads>>>( + para->getParD(level)->omega, + para->getParD(level)->typeOfGridNode, + para->getParD(level)->neighborX, + para->getParD(level)->neighborY, + para->getParD(level)->neighborZ, + para->getParD(level)->distributions.f[0], + para->getParD(level)->numberOfNodes, + para->getParD(level)->isEvenTimestep); + getLastCudaError("B15CompressibleNavierStokesBGKplus_Device execution failed"); +} + +B15CompressibleNavierStokesBGKplus::B15CompressibleNavierStokesBGKplus(std::shared_ptr<Parameter> para, int level) +{ + this->para = para; + this->level = level; + + myPreProcessorTypes.push_back(InitCompSP27); + + +} + +B15CompressibleNavierStokesBGKplus::B15CompressibleNavierStokesBGKplus() +{ +} diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B15CompressibleNavierStokesBGKplus/B15CompressibleNavierStokesBGKplus.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B15CompressibleNavierStokesBGKplus/B15CompressibleNavierStokesBGKplus.h new file mode 100644 index 0000000000000000000000000000000000000000..0347e30dc1aec97bcdd6e21eb58cf929084987cc --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B15CompressibleNavierStokesBGKplus/B15CompressibleNavierStokesBGKplus.h @@ -0,0 +1,17 @@ +#ifndef B15CompressibleNavierStokesBGKplus_H +#define B15CompressibleNavierStokesBGKplus_H + +#include "Kernel/KernelImp.h" + +class B15CompressibleNavierStokesBGKplus : public KernelImp +{ +public: + static std::shared_ptr<B15CompressibleNavierStokesBGKplus> getNewInstance(std::shared_ptr< Parameter> para, int level); + void run(); + +private: + B15CompressibleNavierStokesBGKplus(); + B15CompressibleNavierStokesBGKplus(std::shared_ptr< Parameter> para, int level); +}; + +#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKPlus/BGKPlusCompSP27_Device.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B15CompressibleNavierStokesBGKplus/B15CompressibleNavierStokesBGKplus_Device.cu similarity index 99% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKPlus/BGKPlusCompSP27_Device.cu rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B15CompressibleNavierStokesBGKplus/B15CompressibleNavierStokesBGKplus_Device.cu index 638210bd2da8ebf30bda603a3f6d70c19468193e..b21213bf58fd786a4006faf485f18e5c18694e2e 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKPlus/BGKPlusCompSP27_Device.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B15CompressibleNavierStokesBGKplus/B15CompressibleNavierStokesBGKplus_Device.cu @@ -6,7 +6,7 @@ using namespace vf::basics::constant; using namespace vf::lbm::dir; #include "math.h" -__global__ void LB_Kernel_BGK_Plus_Comp_SP_27( +__global__ void B15CompressibleNavierStokesBGKplus_Device( real omega, unsigned int* bcMatD, unsigned int* neighborX, diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B15CompressibleNavierStokesBGKplus/B15CompressibleNavierStokesBGKplus_Device.cuh b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B15CompressibleNavierStokesBGKplus/B15CompressibleNavierStokesBGKplus_Device.cuh new file mode 100644 index 0000000000000000000000000000000000000000..2a4775b95d0b4d371fadb332a08868fda0fafd37 --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B15CompressibleNavierStokesBGKplus/B15CompressibleNavierStokesBGKplus_Device.cuh @@ -0,0 +1,16 @@ +#ifndef B15CompressibleNavierStokesBGKplus_Device_H +#define B15CompressibleNavierStokesBGKplus_Device_H + +#include <DataTypes.h> +#include <curand.h> + +__global__ void B15CompressibleNavierStokesBGKplus_Device( + real omega, + unsigned int* bcMatD, + unsigned int* neighborX, + unsigned int* neighborY, + unsigned int* neighborZ, + real* DDStart, + int size_Mat, + bool EvenOrOdd); +#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKUnified/BGKUnified.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B15CompressibleNavierStokesBGKplusUnified/B15CompressibleNavierStokesBGKplusUnified.cu similarity index 83% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKUnified/BGKUnified.cu rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B15CompressibleNavierStokesBGKplusUnified/B15CompressibleNavierStokesBGKplusUnified.cu index 0a5ac6cf7a1b6564a61d0150b187b10b584222b8..bafd4477dc611b77f17c896dc85832b71acdccc0 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKUnified/BGKUnified.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B15CompressibleNavierStokesBGKplusUnified/B15CompressibleNavierStokesBGKplusUnified.cu @@ -1,4 +1,4 @@ -#include "BGKUnified.h" +#include "B15CompressibleNavierStokesBGKplusUnified.h" #include <stdexcept> @@ -15,7 +15,7 @@ namespace gpu { -BGKUnified::BGKUnified(std::shared_ptr<Parameter> para, int level) +B15CompressibleNavierStokesBGKplusUnified::B15CompressibleNavierStokesBGKplusUnified(std::shared_ptr<Parameter> para, int level) : KernelImp(para, level) { #ifndef BUILD_CUDA_LTO @@ -30,7 +30,7 @@ BGKUnified::BGKUnified(std::shared_ptr<Parameter> para, int level) } -void BGKUnified::run() +void B15CompressibleNavierStokesBGKplusUnified::run() { GPUKernelParameter kernelParameter{ para->getParD(level)->omega, diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B15CompressibleNavierStokesBGKplusUnified/B15CompressibleNavierStokesBGKplusUnified.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B15CompressibleNavierStokesBGKplusUnified/B15CompressibleNavierStokesBGKplusUnified.h new file mode 100644 index 0000000000000000000000000000000000000000..ffc8ce41d13369b78d19ff0d55bab988696cdb5d --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B15CompressibleNavierStokesBGKplusUnified/B15CompressibleNavierStokesBGKplusUnified.h @@ -0,0 +1,22 @@ +#ifndef B15CompressibleNavierStokesBGKplusUnified_H +#define B15CompressibleNavierStokesBGKplusUnified_H + +#include "Kernel/KernelImp.h" + +namespace vf +{ +namespace gpu +{ + +class B15CompressibleNavierStokesBGKplusUnified : public KernelImp +{ +public: + B15CompressibleNavierStokesBGKplusUnified(std::shared_ptr<Parameter> para, int level); + + void run(); +}; + +} +} + +#endif diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGK/BGKCompSP27.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B92CompressibleNavierStokes/B92CompressibleNavierStokes.cu similarity index 57% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGK/BGKCompSP27.cu rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B92CompressibleNavierStokes/B92CompressibleNavierStokes.cu index 4aef26b7dd31435b2dadceb78ac1e0b7ebedf029..b708caf38990e563ab65caec26a68294c8589361 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGK/BGKCompSP27.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B92CompressibleNavierStokes/B92CompressibleNavierStokes.cu @@ -1,19 +1,19 @@ -#include "BGKCompSP27.h" +#include "B92CompressibleNavierStokes.h" -#include "BGKCompSP27_Device.cuh" +#include "B92CompressibleNavierStokes_Device.cuh" #include "Parameter/Parameter.h" #include "cuda/CudaGrid.h" -std::shared_ptr<BGKCompSP27> BGKCompSP27::getNewInstance(std::shared_ptr<Parameter> para, int level) +std::shared_ptr<B92CompressibleNavierStokes> B92CompressibleNavierStokes::getNewInstance(std::shared_ptr<Parameter> para, int level) { - return std::shared_ptr<BGKCompSP27>(new BGKCompSP27(para, level)); + return std::shared_ptr<B92CompressibleNavierStokes>(new B92CompressibleNavierStokes(para, level)); } -void BGKCompSP27::run() +void B92CompressibleNavierStokes::run() { vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes); - LB_Kernel_BGK_Comp_SP_27<<< grid.grid, grid.threads >>>( + B92CompressibleNavierStokes_Device<<< grid.grid, grid.threads >>>( para->getParD(level)->omega, para->getParD(level)->typeOfGridNode, para->getParD(level)->neighborX, @@ -25,7 +25,7 @@ void BGKCompSP27::run() getLastCudaError("LB_Kernel_BGK_Comp_SP_27 execution failed"); } -BGKCompSP27::BGKCompSP27(std::shared_ptr<Parameter> para, int level) +B92CompressibleNavierStokes::B92CompressibleNavierStokes(std::shared_ptr<Parameter> para, int level) { this->para = para; this->level = level; @@ -34,6 +34,6 @@ BGKCompSP27::BGKCompSP27(std::shared_ptr<Parameter> para, int level) } -BGKCompSP27::BGKCompSP27() +B92CompressibleNavierStokes::B92CompressibleNavierStokes() { } diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B92CompressibleNavierStokes/B92CompressibleNavierStokes.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B92CompressibleNavierStokes/B92CompressibleNavierStokes.h new file mode 100644 index 0000000000000000000000000000000000000000..8b7ba23b33c6c78c1a87113a4e61748da42d400d --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B92CompressibleNavierStokes/B92CompressibleNavierStokes.h @@ -0,0 +1,16 @@ +#ifndef B92CompressibleNavierStokes_H +#define B92CompressibleNavierStokes_H + +#include "Kernel/KernelImp.h" + +class B92CompressibleNavierStokes : public KernelImp +{ +public: + static std::shared_ptr<B92CompressibleNavierStokes> getNewInstance(std::shared_ptr< Parameter> para, int level); + void run(); + +private: + B92CompressibleNavierStokes(); + B92CompressibleNavierStokes(std::shared_ptr< Parameter> para, int level); +}; +#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGK/BGKCompSP27_Device.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B92CompressibleNavierStokes/B92CompressibleNavierStokes_Device.cu similarity index 99% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGK/BGKCompSP27_Device.cu rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B92CompressibleNavierStokes/B92CompressibleNavierStokes_Device.cu index 44add98d9607642531a6068021d0a4e831cb3d4e..d032936a96f77304071d014590e35512c1ca7771 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGK/BGKCompSP27_Device.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B92CompressibleNavierStokes/B92CompressibleNavierStokes_Device.cu @@ -7,7 +7,7 @@ using namespace vf::lbm::dir; #include "math.h" -__global__ void LB_Kernel_BGK_Comp_SP_27( real omega, +__global__ void B92CompressibleNavierStokes_Device( real omega, unsigned int* bcMatD, unsigned int* neighborX, unsigned int* neighborY, diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B92CompressibleNavierStokes/B92CompressibleNavierStokes_Device.cuh b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B92CompressibleNavierStokes/B92CompressibleNavierStokes_Device.cuh new file mode 100644 index 0000000000000000000000000000000000000000..4e0450452ddeca71cfd7c76074988159e51fa78d --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B92CompressibleNavierStokes/B92CompressibleNavierStokes_Device.cuh @@ -0,0 +1,17 @@ +#ifndef B92CompressibleNavierStokes_Device_H +#define B92CompressibleNavierStokes_Device_H + +#include <DataTypes.h> +#include <curand.h> + +__global__ void B92CompressibleNavierStokes_Device( + real omega, + unsigned int* bcMatD, + unsigned int* neighborX, + unsigned int* neighborY, + unsigned int* neighborZ, + real* DDStart, + int size_Mat, + bool EvenOrOdd); + +#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGK/BGKCompSP27.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGK/BGKCompSP27.h deleted file mode 100644 index 4163d87fa5105a737138cd5c5e67fe5b01edaeed..0000000000000000000000000000000000000000 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGK/BGKCompSP27.h +++ /dev/null @@ -1,16 +0,0 @@ -#ifndef BGK_COMP_SP27_H -#define BGK_COMP_SP27_H - -#include "Kernel/KernelImp.h" - -class BGKCompSP27 : public KernelImp -{ -public: - static std::shared_ptr<BGKCompSP27> getNewInstance(std::shared_ptr< Parameter> para, int level); - void run(); - -private: - BGKCompSP27(); - BGKCompSP27(std::shared_ptr< Parameter> para, int level); -}; -#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGK/BGKCompSP27_Device.cuh b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGK/BGKCompSP27_Device.cuh deleted file mode 100644 index 59a5240862ed92a9ea3e9187c503ee9233da7e5a..0000000000000000000000000000000000000000 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGK/BGKCompSP27_Device.cuh +++ /dev/null @@ -1,16 +0,0 @@ -#ifndef LB_Kernel_BGK_COMP_SP_27_H -#define LB_Kernel_BGK_COMP_SP_27_H - -#include <DataTypes.h> -#include <curand.h> - -__global__ void LB_Kernel_BGK_Comp_SP_27( real omega, - unsigned int* bcMatD, - unsigned int* neighborX, - unsigned int* neighborY, - unsigned int* neighborZ, - real* DDStart, - int size_Mat, - bool EvenOrOdd); - -#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKPlus/BGKPlusCompSP27.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKPlus/BGKPlusCompSP27.cu deleted file mode 100644 index 00aaf3c27f16a5d53e7aee225214f05bd62a541a..0000000000000000000000000000000000000000 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKPlus/BGKPlusCompSP27.cu +++ /dev/null @@ -1,40 +0,0 @@ -#include "BGKPlusCompSP27.h" - -#include "BGKPlusCompSP27_Device.cuh" -#include "Parameter/Parameter.h" -#include "cuda/CudaGrid.h" - -std::shared_ptr<BGKPlusCompSP27> BGKPlusCompSP27::getNewInstance(std::shared_ptr<Parameter> para, int level) -{ - return std::shared_ptr<BGKPlusCompSP27>(new BGKPlusCompSP27(para, level)); -} - -void BGKPlusCompSP27::run() -{ - vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes); - - LB_Kernel_BGK_Plus_Comp_SP_27 <<< grid.grid, grid.threads >>>( - para->getParD(level)->omega, - para->getParD(level)->typeOfGridNode, - para->getParD(level)->neighborX, - para->getParD(level)->neighborY, - para->getParD(level)->neighborZ, - para->getParD(level)->distributions.f[0], - para->getParD(level)->numberOfNodes, - para->getParD(level)->isEvenTimestep); - getLastCudaError("LB_Kernel_BGK_Plus_Comp_SP_27 execution failed"); -} - -BGKPlusCompSP27::BGKPlusCompSP27(std::shared_ptr<Parameter> para, int level) -{ - this->para = para; - this->level = level; - - myPreProcessorTypes.push_back(InitCompSP27); - - -} - -BGKPlusCompSP27::BGKPlusCompSP27() -{ -} diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKPlus/BGKPlusCompSP27.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKPlus/BGKPlusCompSP27.h deleted file mode 100644 index da40eb8aeee5d19646a66b7a0f8f8da498b85a2a..0000000000000000000000000000000000000000 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKPlus/BGKPlusCompSP27.h +++ /dev/null @@ -1,17 +0,0 @@ -#ifndef BGK_PLUS_COMP_SP27_H -#define BGK_PLUS_COMP_SP27_H - -#include "Kernel/KernelImp.h" - -class BGKPlusCompSP27 : public KernelImp -{ -public: - static std::shared_ptr<BGKPlusCompSP27> getNewInstance(std::shared_ptr< Parameter> para, int level); - void run(); - -private: - BGKPlusCompSP27(); - BGKPlusCompSP27(std::shared_ptr< Parameter> para, int level); -}; - -#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKPlus/BGKPlusCompSP27_Device.cuh b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKPlus/BGKPlusCompSP27_Device.cuh deleted file mode 100644 index 9e991ffa4b16e0df78fe23f6ee5a1e0678919cd7..0000000000000000000000000000000000000000 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKPlus/BGKPlusCompSP27_Device.cuh +++ /dev/null @@ -1,15 +0,0 @@ -#ifndef LB_Kernel_BGK_PLUS_Comp_SP_27_H -#define LB_Kernel_BGK_PLUS_Comp_SP_27_H - -#include <DataTypes.h> -#include <curand.h> - -__global__ void LB_Kernel_BGK_Plus_Comp_SP_27( real omega, - unsigned int* bcMatD, - unsigned int* neighborX, - unsigned int* neighborY, - unsigned int* neighborZ, - real* DDStart, - int size_Mat, - bool EvenOrOdd); -#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKUnified/BGKUnified.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKUnified/BGKUnified.h deleted file mode 100644 index 762eaaa5935bd01fa6ae002521a40e45cd239dfd..0000000000000000000000000000000000000000 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKUnified/BGKUnified.h +++ /dev/null @@ -1,22 +0,0 @@ -#ifndef GPU_BKGUnified_H -#define GPU_BKGUnified_H - -#include "Kernel/KernelImp.h" - -namespace vf -{ -namespace gpu -{ - -class BGKUnified : public KernelImp -{ -public: - BGKUnified(std::shared_ptr<Parameter> para, int level); - - void run(); -}; - -} -} - -#endif diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cascade/CascadeCompSP27.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/C06CompressibleNavierStokes/C06CompressibleNavierStokes.cu similarity index 54% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cascade/CascadeCompSP27.cu rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/C06CompressibleNavierStokes/C06CompressibleNavierStokes.cu index 664b46fcebd277b0c93300d86b2171edf4f91b2a..ca943474683ca9b5517ef97626ad2449b74119f2 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cascade/CascadeCompSP27.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/C06CompressibleNavierStokes/C06CompressibleNavierStokes.cu @@ -1,19 +1,19 @@ -#include "CascadeCompSP27.h" +#include "C06CompressibleNavierStokes.h" -#include "CascadeCompSP27_Device.cuh" +#include "C06CompressibleNavierStokes_Device.cuh" #include "Parameter/Parameter.h" #include "cuda/CudaGrid.h" -std::shared_ptr<CascadeCompSP27> CascadeCompSP27::getNewInstance(std::shared_ptr<Parameter> para, int level) +std::shared_ptr<C06CompressibleNavierStokes> C06CompressibleNavierStokes::getNewInstance(std::shared_ptr<Parameter> para, int level) { - return std::shared_ptr<CascadeCompSP27>(new CascadeCompSP27(para, level)); + return std::shared_ptr<C06CompressibleNavierStokes>(new C06CompressibleNavierStokes(para, level)); } -void CascadeCompSP27::run() +void C06CompressibleNavierStokes::run() { vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes); - LB_Kernel_Cascade_Comp_SP_27 <<< grid.grid, grid.threads >>>( + C06CompressibleNavierStokes_Device <<< grid.grid, grid.threads >>>( para->getParD(level)->omega, para->getParD(level)->typeOfGridNode, para->getParD(level)->neighborX, @@ -25,7 +25,7 @@ void CascadeCompSP27::run() getLastCudaError("LB_Kernel_Cascade_Comp_SP_27 execution failed"); } -CascadeCompSP27::CascadeCompSP27(std::shared_ptr<Parameter> para, int level) +C06CompressibleNavierStokes::C06CompressibleNavierStokes(std::shared_ptr<Parameter> para, int level) { this->para = para; this->level = level; @@ -35,6 +35,6 @@ CascadeCompSP27::CascadeCompSP27(std::shared_ptr<Parameter> para, int level) } -CascadeCompSP27::CascadeCompSP27() +C06CompressibleNavierStokes::C06CompressibleNavierStokes() { } diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/C06CompressibleNavierStokes/C06CompressibleNavierStokes.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/C06CompressibleNavierStokes/C06CompressibleNavierStokes.h new file mode 100644 index 0000000000000000000000000000000000000000..8004ad52ea2471ed6319fc373175e509dedef5ff --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/C06CompressibleNavierStokes/C06CompressibleNavierStokes.h @@ -0,0 +1,16 @@ +#ifndef C06CompressibleNavierStokes_H +#define C06CompressibleNavierStokes_H + +#include "Kernel/KernelImp.h" + +class C06CompressibleNavierStokes : public KernelImp +{ +public: + static std::shared_ptr<C06CompressibleNavierStokes> getNewInstance(std::shared_ptr< Parameter> para, int level); + void run(); + +private: + C06CompressibleNavierStokes(); + C06CompressibleNavierStokes(std::shared_ptr< Parameter> para, int level); +}; +#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cascade/CascadeCompSP27_Device.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/C06CompressibleNavierStokes/C06CompressibleNavierStokes_Device.cu similarity index 99% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cascade/CascadeCompSP27_Device.cu rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/C06CompressibleNavierStokes/C06CompressibleNavierStokes_Device.cu index 6bd4415c7edcb5c9954874c074801b865cce3efe..6cf0fd639523d4cef8290f860f03b4bc75e3252c 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cascade/CascadeCompSP27_Device.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/C06CompressibleNavierStokes/C06CompressibleNavierStokes_Device.cu @@ -6,7 +6,7 @@ using namespace vf::basics::constant; using namespace vf::lbm::dir; #include "math.h" -__global__ void LB_Kernel_Cascade_Comp_SP_27(real omega, +__global__ void C06CompressibleNavierStokes_Device(real omega, unsigned int* bcMatD, unsigned int* neighborX, unsigned int* neighborY, diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cascade/CascadeCompSP27_Device.cuh b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/C06CompressibleNavierStokes/C06CompressibleNavierStokes_Device.cuh similarity index 57% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cascade/CascadeCompSP27_Device.cuh rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/C06CompressibleNavierStokes/C06CompressibleNavierStokes_Device.cuh index 8a49bd02af3ab420b42bc257e8668dd3ff9eca2c..fb8fc657a03c89efbad879ca1c9b4e493acabd86 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cascade/CascadeCompSP27_Device.cuh +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/C06CompressibleNavierStokes/C06CompressibleNavierStokes_Device.cuh @@ -1,10 +1,11 @@ -#ifndef LB_KERNEL_CASCADE_COMP_SP_27_H -#define LB_KERNEL_CASCADE_COMP_SP_27_H +#ifndef C06CompressibleNavierStokes_Device_H +#define C06CompressibleNavierStokes_Device_H #include <DataTypes.h> #include <curand.h> -__global__ void LB_Kernel_Cascade_Comp_SP_27(real s9, +__global__ void C06CompressibleNavierStokes_Device( + real s9, unsigned int* bcMatD, unsigned int* neighborX, unsigned int* neighborY, diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cascade/CascadeCompSP27.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cascade/CascadeCompSP27.h deleted file mode 100644 index 900f41884c5f19f050c33a0e6141010e0b72b46f..0000000000000000000000000000000000000000 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cascade/CascadeCompSP27.h +++ /dev/null @@ -1,16 +0,0 @@ -#ifndef CASCADE_COMP_SP27_H -#define CASCADE_COMP_SP27_H - -#include "Kernel/KernelImp.h" - -class CascadeCompSP27 : public KernelImp -{ -public: - static std::shared_ptr<CascadeCompSP27> getNewInstance(std::shared_ptr< Parameter> para, int level); - void run(); - -private: - CascadeCompSP27(); - CascadeCompSP27(std::shared_ptr< Parameter> para, int level); -}; -#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cumulant/CumulantCompSP27.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cumulant/CumulantCompSP27.h deleted file mode 100644 index 0f9317f12b9438f770cd60355c41dc13599f1ad2..0000000000000000000000000000000000000000 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cumulant/CumulantCompSP27.h +++ /dev/null @@ -1,17 +0,0 @@ -#ifndef CUMULANT_COMP_SP27_H -#define CUMULANT_COMP_SP27_H - -#include "Kernel/KernelImp.h" - -class CumulantCompSP27 : public KernelImp -{ -public: - static std::shared_ptr<CumulantCompSP27> getNewInstance(std::shared_ptr< Parameter> para, int level); - void run(); - -private: - CumulantCompSP27(); - CumulantCompSP27(std::shared_ptr< Parameter> para, int level); - -}; -#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantAll4/CumulantAll4CompSP27.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantAll4/CumulantAll4CompSP27.cu deleted file mode 100644 index c8aad41b87ef39514f6cf5abc8b8bff42a869346..0000000000000000000000000000000000000000 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantAll4/CumulantAll4CompSP27.cu +++ /dev/null @@ -1,39 +0,0 @@ -#include "CumulantAll4CompSP27.h" - -#include "CumulantAll4CompSP27_Device.cuh" -#include "Parameter/Parameter.h" -#include "cuda/CudaGrid.h" - -std::shared_ptr<CumulantAll4CompSP27> CumulantAll4CompSP27::getNewInstance(std::shared_ptr<Parameter> para, int level) -{ - return std::shared_ptr<CumulantAll4CompSP27>(new CumulantAll4CompSP27(para, level)); -} - -void CumulantAll4CompSP27::run() -{ - vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes); - - LB_Kernel_Cumulant_D3Q27All4 <<< grid.grid, grid.threads >>>( - para->getParD(level)->omega, - para->getParD(level)->typeOfGridNode, - para->getParD(level)->neighborX, - para->getParD(level)->neighborY, - para->getParD(level)->neighborZ, - para->getParD(level)->distributions.f[0], - para->getParD(level)->numberOfNodes, - level, - para->getForcesDev(), - para->getQuadricLimitersDev(), - para->getParD(level)->isEvenTimestep); - getLastCudaError("LB_Kernel_Cumulant_D3Q27All4 execution failed"); -} - -CumulantAll4CompSP27::CumulantAll4CompSP27(std::shared_ptr<Parameter> para, int level) -{ - this->para = para; - this->level = level; - - myPreProcessorTypes.push_back(InitCompSP27); - - -} \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantAll4/CumulantAll4CompSP27.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantAll4/CumulantAll4CompSP27.h deleted file mode 100644 index e36d51bc471ff38a4f213f415c9e86bd677add32..0000000000000000000000000000000000000000 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantAll4/CumulantAll4CompSP27.h +++ /dev/null @@ -1,16 +0,0 @@ -#ifndef CUMULANT_ALL4_COMP_SP27_H -#define CUMULANT_ALL4_COMP_SP27_H - -#include "Kernel/KernelImp.h" - -class CumulantAll4CompSP27 : public KernelImp -{ -public: - static std::shared_ptr<CumulantAll4CompSP27> getNewInstance(std::shared_ptr< Parameter> para, int level); - void run(); - -private: - CumulantAll4CompSP27(); - CumulantAll4CompSP27(std::shared_ptr< Parameter> para, int level); -}; -#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantAll4/CumulantAll4CompSP27_Device.cuh b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantAll4/CumulantAll4CompSP27_Device.cuh deleted file mode 100644 index 5f23194d561d106cf2493c36199444f8da15efd7..0000000000000000000000000000000000000000 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantAll4/CumulantAll4CompSP27_Device.cuh +++ /dev/null @@ -1,18 +0,0 @@ -#ifndef LB_Kernel_Cumulant_D3Q27All4_H -#define LB_Kernel_Cumulant_D3Q27All4_H - -#include <DataTypes.h> -#include <curand.h> - -__global__ void LB_Kernel_Cumulant_D3Q27All4( 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 \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK18/CumulantK18Comp.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK18/CumulantK18Comp.h deleted file mode 100644 index 6cffff12643d066d167d19bf9ab2f0c450979d3f..0000000000000000000000000000000000000000 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK18/CumulantK18Comp.h +++ /dev/null @@ -1,17 +0,0 @@ -#ifndef CUMULANT_K18_COMP_H -#define CUMULANT_K18_COMP_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/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK18/CumulantK18Comp_Device.cuh b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK18/CumulantK18Comp_Device.cuh deleted file mode 100644 index 60a15145e3c117cadb7485f2899ba768b10eb0c1..0000000000000000000000000000000000000000 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK18/CumulantK18Comp_Device.cuh +++ /dev/null @@ -1,20 +0,0 @@ -#ifndef LB_Kernel_Cumulant_K18_H -#define LB_Kernel_Cumulant_K18_H - -#include <DataTypes.h> -#include <curand.h> - -__global__ void LB_Kernel_CumulantK18Comp( real omega, - unsigned int* bcMatD, - unsigned int* neighborX, - unsigned int* neighborY, - unsigned int* neighborZ, - real* DDStart, - real* F3, - int size_Mat, - int level, - real* forces, - real* quadricLimiters, - bool EvenOrOdd); - -#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK20/CumulantK20Comp.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK20/CumulantK20Comp.h deleted file mode 100644 index 24745df47ed75eb3369f41068d49d27702eee461..0000000000000000000000000000000000000000 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK20/CumulantK20Comp.h +++ /dev/null @@ -1,16 +0,0 @@ -#ifndef CUMULANT_K20_COMP_H -#define CUMULANT_K20_COMP_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/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cumulant/CumulantCompSP27.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K08CompressibleNavierStokes/K08CompressibleNavierStokes.cu similarity index 54% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cumulant/CumulantCompSP27.cu rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K08CompressibleNavierStokes/K08CompressibleNavierStokes.cu index 218623b7c51099717f6aaa6f375a82516e0c0dae..d56b9f2ecea4704be7c5c5d6d843492e74312f3e 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cumulant/CumulantCompSP27.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K08CompressibleNavierStokes/K08CompressibleNavierStokes.cu @@ -1,19 +1,19 @@ -#include "CumulantCompSP27.h" +#include "K08CompressibleNavierStokes.h" -#include "CumulantCompSP27_Device.cuh" +#include "K08CompressibleNavierStokes_Device.cuh" #include "Parameter/Parameter.h" #include "cuda/CudaGrid.h" -std::shared_ptr<CumulantCompSP27> CumulantCompSP27::getNewInstance(std::shared_ptr<Parameter> para, int level) +std::shared_ptr<K08CompressibleNavierStokes> K08CompressibleNavierStokes::getNewInstance(std::shared_ptr<Parameter> para, int level) { - return std::shared_ptr<CumulantCompSP27>(new CumulantCompSP27(para, level)); + return std::shared_ptr<K08CompressibleNavierStokes>(new K08CompressibleNavierStokes(para, level)); } -void CumulantCompSP27::run() +void K08CompressibleNavierStokes::run() { vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes); - LB_Kernel_Cum_Comp_SP_27 <<< grid.grid, grid.threads >>>( + K08CompressibleNavierStokes_Device <<< grid.grid, grid.threads >>>( para->getParD(level)->omega, para->getParD(level)->typeOfGridNode, para->getParD(level)->neighborX, @@ -26,7 +26,7 @@ void CumulantCompSP27::run() } -CumulantCompSP27::CumulantCompSP27(std::shared_ptr<Parameter> para, int level) +K08CompressibleNavierStokes::K08CompressibleNavierStokes(std::shared_ptr<Parameter> para, int level) { this->para = para; this->level = level; @@ -36,6 +36,6 @@ CumulantCompSP27::CumulantCompSP27(std::shared_ptr<Parameter> para, int level) } -CumulantCompSP27::CumulantCompSP27() +K08CompressibleNavierStokes::K08CompressibleNavierStokes() { } \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K08CompressibleNavierStokes/K08CompressibleNavierStokes.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K08CompressibleNavierStokes/K08CompressibleNavierStokes.h new file mode 100644 index 0000000000000000000000000000000000000000..2f91a459818173775aeb0b966ec6677fe7803b48 --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K08CompressibleNavierStokes/K08CompressibleNavierStokes.h @@ -0,0 +1,17 @@ +#ifndef K08CompressibleNavierStokes_H +#define K08CompressibleNavierStokes_H + +#include "Kernel/KernelImp.h" + +class K08CompressibleNavierStokes : public KernelImp +{ +public: + static std::shared_ptr<K08CompressibleNavierStokes> getNewInstance(std::shared_ptr< Parameter> para, int level); + void run(); + +private: + K08CompressibleNavierStokes(); + K08CompressibleNavierStokes(std::shared_ptr< Parameter> para, int level); + +}; +#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cumulant/CumulantCompSP27_Device.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K08CompressibleNavierStokes/K08CompressibleNavierStokes_Device.cu similarity index 99% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cumulant/CumulantCompSP27_Device.cu rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K08CompressibleNavierStokes/K08CompressibleNavierStokes_Device.cu index 6ab3385b86611614eceeb0018f6beef73031711c..021ec25f3d70802402bb2d96ca6a462e60774982 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cumulant/CumulantCompSP27_Device.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K08CompressibleNavierStokes/K08CompressibleNavierStokes_Device.cu @@ -6,7 +6,7 @@ using namespace vf::basics::constant; using namespace vf::lbm::dir; #include "math.h" -__global__ void LB_Kernel_Cum_Comp_SP_27(real omega, +__global__ void K08CompressibleNavierStokes_Device(real omega, unsigned int* bcMatD, unsigned int* neighborX, unsigned int* neighborY, diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cumulant/CumulantCompSP27_Device.cuh b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K08CompressibleNavierStokes/K08CompressibleNavierStokes_Device.cuh similarity index 57% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cumulant/CumulantCompSP27_Device.cuh rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K08CompressibleNavierStokes/K08CompressibleNavierStokes_Device.cuh index cf6a926698e4082383b04c5f1e2d886c6dca6380..b2a6d02b6ab40e8dc89238b207c7f274ab7fe828 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cumulant/CumulantCompSP27_Device.cuh +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K08CompressibleNavierStokes/K08CompressibleNavierStokes_Device.cuh @@ -1,10 +1,11 @@ -#ifndef LB_KERNEL_CUM_COMP_SP_27_H -#define LB_KERNEL_CUM_COMP_SP_27_H +#ifndef K08CompressibleNavierStokes_Device_H +#define K08CompressibleNavierStokes_Device_H #include <DataTypes.h> #include <curand.h> -__global__ void LB_Kernel_Cum_Comp_SP_27(real s9, +__global__ void K08CompressibleNavierStokes_Device( + real s9, unsigned int* bcMatD, unsigned int* neighborX, unsigned int* neighborY, diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K15CompressibleNavierStokes/K15CompressibleNavierStokes_Device.cuh b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K15CompressibleNavierStokes/K15CompressibleNavierStokes_Device.cuh index 77d42493ead97135314a55c72d294f3520f39a73..0dfdb3488136575921eb169cde4364d9de1507f5 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K15CompressibleNavierStokes/K15CompressibleNavierStokes_Device.cuh +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K15CompressibleNavierStokes/K15CompressibleNavierStokes_Device.cuh @@ -4,14 +4,15 @@ #include <DataTypes.h> #include <curand.h> -__global__ void K15CompressibleNavierStokes_Device( 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); +__global__ void K15CompressibleNavierStokes_Device( + 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/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants.cu new file mode 100644 index 0000000000000000000000000000000000000000..1957ae332acaeed5e29584079c670bf7ddcfadd2 --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants.cu @@ -0,0 +1,39 @@ +#include "K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants.h" + +#include "K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants_Device.cuh" +#include "Parameter/Parameter.h" +#include "cuda/CudaGrid.h" + +std::shared_ptr<K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants> K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants::getNewInstance(std::shared_ptr<Parameter> para, int level) +{ + return std::shared_ptr<K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants>(new K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants(para, level)); +} + +void K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants::run() +{ + vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes); + + K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants_Device <<< grid.grid, grid.threads >>>( + para->getParD(level)->omega, + para->getParD(level)->typeOfGridNode, + para->getParD(level)->neighborX, + para->getParD(level)->neighborY, + para->getParD(level)->neighborZ, + para->getParD(level)->distributions.f[0], + para->getParD(level)->numberOfNodes, + level, + para->getForcesDev(), + para->getQuadricLimitersDev(), + para->getParD(level)->isEvenTimestep); + getLastCudaError("LB_Kernel_Cumulant_D3Q27All4 execution failed"); +} + +K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants::K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants(std::shared_ptr<Parameter> para, int level) +{ + this->para = para; + this->level = level; + + myPreProcessorTypes.push_back(InitCompSP27); + + +} \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants.h new file mode 100644 index 0000000000000000000000000000000000000000..cfc30852016375a4c6512f1ed1e35f069d5e366a --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants.h @@ -0,0 +1,16 @@ +#ifndef K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants_H +#define K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants_H + +#include "Kernel/KernelImp.h" + +class K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants : public KernelImp +{ +public: + static std::shared_ptr<K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants> getNewInstance(std::shared_ptr< Parameter> para, int level); + void run(); + +private: + K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants(); + K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants(std::shared_ptr< Parameter> para, int level); +}; +#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantAll4/CumulantAll4CompSP27_Device.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants_Device.cu similarity index 99% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantAll4/CumulantAll4CompSP27_Device.cu rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants_Device.cu index 7a5e39d6f1f95f5b34bb38f1514ab728f477c34b..4919f56979ef004d3d74049feadc43895ad7cf28 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantAll4/CumulantAll4CompSP27_Device.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants_Device.cu @@ -1,4 +1,4 @@ -#include "CumulantAll4CompSP27_Device.cuh" +#include "K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants_Device.cuh" #include "LBM/LB.h" #include "lbm/constants/D3Q27.h" @@ -8,17 +8,18 @@ using namespace vf::basics::constant; using namespace vf::lbm::dir; #include "math.h" -__global__ void LB_Kernel_Cumulant_D3Q27All4( 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) +__global__ void K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants_Device( + 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) { //////////////////////////////////////////////////////////////////////////////// const unsigned x = threadIdx.x; // Globaler x-Index diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants_Device.cuh b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants_Device.cuh new file mode 100644 index 0000000000000000000000000000000000000000..df7f6a09676740503aaed1f4b16b0b8db96624c8 --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants_Device.cuh @@ -0,0 +1,19 @@ +#ifndef K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants_Device_H +#define K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants_Device_H + +#include <DataTypes.h> +#include <curand.h> + +__global__ void K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants_Device( + 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 \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK18/CumulantK18Comp.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K18CompressibleNavierStokes/K18CompressibleNavierStokes.cu similarity index 61% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK18/CumulantK18Comp.cu rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K18CompressibleNavierStokes/K18CompressibleNavierStokes.cu index 15d3509e735faa08b97d0876600c30876829c35f..801c7d90576324559873afa81e2d5f543219d5ae 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK18/CumulantK18Comp.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K18CompressibleNavierStokes/K18CompressibleNavierStokes.cu @@ -1,19 +1,19 @@ -#include "CumulantK18Comp.h" +#include "K18CompressibleNavierStokes.h" -#include "CumulantK18Comp_Device.cuh" +#include "K18CompressibleNavierStokes_Device.cuh" #include "Parameter/Parameter.h" #include "cuda/CudaGrid.h" -std::shared_ptr<CumulantK18Comp> CumulantK18Comp::getNewInstance(std::shared_ptr<Parameter> para, int level) +std::shared_ptr<K18CompressibleNavierStokes> K18CompressibleNavierStokes::getNewInstance(std::shared_ptr<Parameter> para, int level) { - return std::shared_ptr<CumulantK18Comp>(new CumulantK18Comp(para, level)); + return std::shared_ptr<K18CompressibleNavierStokes>(new K18CompressibleNavierStokes(para, level)); } -void CumulantK18Comp::run() +void K18CompressibleNavierStokes::run() { vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes); - LB_Kernel_CumulantK18Comp <<< grid.grid, grid.threads >>>( + K18CompressibleNavierStokes_Device <<< grid.grid, grid.threads >>>( para->getParD(level)->omega, para->getParD(level)->typeOfGridNode, para->getParD(level)->neighborX, @@ -29,7 +29,7 @@ void CumulantK18Comp::run() getLastCudaError("LB_Kernel_CumulantK18Comp execution failed"); } -CumulantK18Comp::CumulantK18Comp(std::shared_ptr<Parameter> para, int level) +K18CompressibleNavierStokes::K18CompressibleNavierStokes(std::shared_ptr<Parameter> para, int level) { this->para = para; this->level = level; diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K18CompressibleNavierStokes/K18CompressibleNavierStokes.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K18CompressibleNavierStokes/K18CompressibleNavierStokes.h new file mode 100644 index 0000000000000000000000000000000000000000..d75920c78efe48ad586c4c92f775c3868195e9b8 --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K18CompressibleNavierStokes/K18CompressibleNavierStokes.h @@ -0,0 +1,17 @@ +#ifndef K18CompressibleNavierStokes_H +#define K18CompressibleNavierStokes_H + +#include "Kernel/KernelImp.h" + +class K18CompressibleNavierStokes : public KernelImp +{ +public: + static std::shared_ptr< K18CompressibleNavierStokes> getNewInstance(std::shared_ptr< Parameter> para, int level); + void run(); + +private: + K18CompressibleNavierStokes(); + K18CompressibleNavierStokes(std::shared_ptr< Parameter> para, int level); +}; + +#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK18/CumulantK18Comp_Device.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K18CompressibleNavierStokes/K18CompressibleNavierStokes_Device.cu similarity index 99% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK18/CumulantK18Comp_Device.cu rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K18CompressibleNavierStokes/K18CompressibleNavierStokes_Device.cu index c585c19aaca870e97fec63b0cd1742f7aad32556..169dedd3660ba9fc3c5e15bdfb60941bad2856c3 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK18/CumulantK18Comp_Device.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K18CompressibleNavierStokes/K18CompressibleNavierStokes_Device.cu @@ -6,7 +6,7 @@ using namespace vf::basics::constant; using namespace vf::lbm::dir; #include "math.h" -__global__ void LB_Kernel_CumulantK18Comp( +__global__ void K18CompressibleNavierStokes_Device( real omega, unsigned int* bcMatD, unsigned int* neighborX, diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K18CompressibleNavierStokes/K18CompressibleNavierStokes_Device.cuh b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K18CompressibleNavierStokes/K18CompressibleNavierStokes_Device.cuh new file mode 100644 index 0000000000000000000000000000000000000000..601071a73c8ea2e0e9d3d0e1b3ba69f921fd821d --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K18CompressibleNavierStokes/K18CompressibleNavierStokes_Device.cuh @@ -0,0 +1,21 @@ +#ifndef K18CompressibleNavierStokes_Device_H +#define K18CompressibleNavierStokes_Device_H + +#include <DataTypes.h> +#include <curand.h> + +__global__ void K18CompressibleNavierStokes_Device( + real omega, + unsigned int* bcMatD, + unsigned int* neighborX, + unsigned int* neighborY, + unsigned int* neighborZ, + real* DDStart, + real* F3, + int size_Mat, + int level, + real* forces, + real* quadricLimiters, + bool EvenOrOdd); + +#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK20/CumulantK20Comp.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K20CompressibleNavierStokes/K20CompressibleNavierStokes.cu similarity index 61% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK20/CumulantK20Comp.cu rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K20CompressibleNavierStokes/K20CompressibleNavierStokes.cu index 8181cdb690b5813c368eddadb9cda58a7d749302..b47f8b633d349a5d128e4bbeed1d642261e6ea3b 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK20/CumulantK20Comp.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K20CompressibleNavierStokes/K20CompressibleNavierStokes.cu @@ -1,19 +1,19 @@ -#include "CumulantK20Comp.h" +#include "K20CompressibleNavierStokes.h" -#include "CumulantK20Comp_Device.cuh" +#include "K20CompressibleNavierStokes_Device.cuh" #include "Parameter/Parameter.h" #include "cuda/CudaGrid.h" -std::shared_ptr<CumulantK20Comp> CumulantK20Comp::getNewInstance(std::shared_ptr<Parameter> para, int level) +std::shared_ptr<K20CompressibleNavierStokes> K20CompressibleNavierStokes::getNewInstance(std::shared_ptr<Parameter> para, int level) { - return std::shared_ptr<CumulantK20Comp>(new CumulantK20Comp(para, level)); + return std::shared_ptr<K20CompressibleNavierStokes>(new K20CompressibleNavierStokes(para, level)); } -void CumulantK20Comp::run() +void K20CompressibleNavierStokes::run() { vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes); - LB_Kernel_CumulantK20Comp <<< grid.grid, grid.threads >>>( + K20CompressibleNavierStokes_Device <<< grid.grid, grid.threads >>>( para->getParD(level)->omega, para->getParD(level)->typeOfGridNode, para->getParD(level)->neighborX, @@ -29,7 +29,7 @@ void CumulantK20Comp::run() getLastCudaError("LB_Kernel_CumulantK20Comp execution failed"); } -CumulantK20Comp::CumulantK20Comp(std::shared_ptr<Parameter> para, int level) +K20CompressibleNavierStokes::K20CompressibleNavierStokes(std::shared_ptr<Parameter> para, int level) { this->para = para; this->level = level; diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K20CompressibleNavierStokes/K20CompressibleNavierStokes.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K20CompressibleNavierStokes/K20CompressibleNavierStokes.h new file mode 100644 index 0000000000000000000000000000000000000000..fc6d4e72acf330a482e80c145a669bda74c0dd8a --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K20CompressibleNavierStokes/K20CompressibleNavierStokes.h @@ -0,0 +1,16 @@ +#ifndef K20CompressibleNavierStokes_H +#define K20CompressibleNavierStokes_H + +#include "Kernel/KernelImp.h" + +class K20CompressibleNavierStokes : public KernelImp +{ +public: + static std::shared_ptr< K20CompressibleNavierStokes> getNewInstance(std::shared_ptr< Parameter> para, int level); + void run(); + +private: + K20CompressibleNavierStokes(); + K20CompressibleNavierStokes(std::shared_ptr< Parameter> para, int level); +}; +#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK20/CumulantK20Comp_Device.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K20CompressibleNavierStokes/K20CompressibleNavierStokes_Device.cu similarity index 99% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK20/CumulantK20Comp_Device.cu rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K20CompressibleNavierStokes/K20CompressibleNavierStokes_Device.cu index 0a26eff29c3624bd78ef7dd4a8f675b8cb1a99d3..d722baa512dea2d7cc01ebf4f986f3239345fb40 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK20/CumulantK20Comp_Device.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K20CompressibleNavierStokes/K20CompressibleNavierStokes_Device.cu @@ -6,7 +6,7 @@ using namespace vf::basics::constant; using namespace vf::lbm::dir; #include "math.h" -__global__ void LB_Kernel_CumulantK20Comp( +__global__ void K20CompressibleNavierStokes_Device( real omega, unsigned int* bcMatD, unsigned int* neighborX, diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK20/CumulantK20Comp_Device.cuh b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K20CompressibleNavierStokes/K20CompressibleNavierStokes_Device.cuh similarity index 75% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK20/CumulantK20Comp_Device.cuh rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K20CompressibleNavierStokes/K20CompressibleNavierStokes_Device.cuh index 17691f621b5a46d29556d71304195f4a346a7ec6..b86fc38428b964304772cf283b310d6056be2682 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK20/CumulantK20Comp_Device.cuh +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K20CompressibleNavierStokes/K20CompressibleNavierStokes_Device.cuh @@ -1,10 +1,10 @@ -#ifndef LB_Kernel_Cumulant_K20_H -#define LB_Kernel_Cumulant_K20_H +#ifndef K20CompressibleNavierStokes_Device_H +#define K20CompressibleNavierStokes_Device_H #include <DataTypes.h> #include <curand.h> -__global__ void LB_Kernel_CumulantK20Comp( real omega, +__global__ void K20CompressibleNavierStokes_Device( real omega, unsigned int* bcMatD, unsigned int* neighborX, unsigned int* neighborY, diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/MRT/MRTCompSP27.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/M02CompressibleNavierStokes/M02CompressibleNavierStokes.cu similarity index 57% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/MRT/MRTCompSP27.cu rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/M02CompressibleNavierStokes/M02CompressibleNavierStokes.cu index 6e11bd97a2e76cca3983a83f785a2435d40f594b..bc8d4e11f1efd9325c005846164b24a8459242ce 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/MRT/MRTCompSP27.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/M02CompressibleNavierStokes/M02CompressibleNavierStokes.cu @@ -1,19 +1,19 @@ -#include "MRTCompSP27.h" +#include "M02CompressibleNavierStokes.h" -#include "MRTCompSP27_Device.cuh" +#include "M02CompressibleNavierStokes_Device.cuh" #include "Parameter/Parameter.h" #include "cuda/CudaGrid.h" -std::shared_ptr<MRTCompSP27> MRTCompSP27::getNewInstance(std::shared_ptr<Parameter> para, int level) +std::shared_ptr<M02CompressibleNavierStokes> M02CompressibleNavierStokes::getNewInstance(std::shared_ptr<Parameter> para, int level) { - return std::shared_ptr<MRTCompSP27>(new MRTCompSP27(para, level)); + return std::shared_ptr<M02CompressibleNavierStokes>(new M02CompressibleNavierStokes(para, level)); } -void MRTCompSP27::run() +void M02CompressibleNavierStokes::run() { vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes); - LB_Kernel_MRT_Comp_SP_27 <<< grid.grid, grid.threads >>>( + M02CompressibleNavierStokes_Device <<< grid.grid, grid.threads >>>( para->getParD(level)->omega, para->getParD(level)->typeOfGridNode, para->getParD(level)->neighborX, @@ -25,16 +25,15 @@ void MRTCompSP27::run() getLastCudaError("LB_Kernel_MRT_Comp_SP_27 execution failed"); } -MRTCompSP27::MRTCompSP27(std::shared_ptr<Parameter> para, int level) +M02CompressibleNavierStokes::M02CompressibleNavierStokes(std::shared_ptr<Parameter> para, int level) { this->para = para; this->level = level; myPreProcessorTypes.push_back(InitCompSP27); - } -MRTCompSP27::MRTCompSP27() +M02CompressibleNavierStokes::M02CompressibleNavierStokes() { } diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/M02CompressibleNavierStokes/M02CompressibleNavierStokes.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/M02CompressibleNavierStokes/M02CompressibleNavierStokes.h new file mode 100644 index 0000000000000000000000000000000000000000..b18eb3dfb7f144e664d9dd2236b5a76a1c63dfa2 --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/M02CompressibleNavierStokes/M02CompressibleNavierStokes.h @@ -0,0 +1,17 @@ +#ifndef M02CompressibleNavierStokes_H +#define M02CompressibleNavierStokes_H + +#include "Kernel/KernelImp.h" + + +class M02CompressibleNavierStokes : public KernelImp +{ +public: + static std::shared_ptr<M02CompressibleNavierStokes> getNewInstance(std::shared_ptr< Parameter> para, int level); + void run(); + +private: + M02CompressibleNavierStokes(); + M02CompressibleNavierStokes(std::shared_ptr< Parameter> para, int level); +}; +#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/MRT/MRTCompSP27_Device.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/M02CompressibleNavierStokes/M02CompressibleNavierStokes_Device.cu similarity index 99% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/MRT/MRTCompSP27_Device.cu rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/M02CompressibleNavierStokes/M02CompressibleNavierStokes_Device.cu index 41b24a349da75586be4fb818c9eb3f194a2447fd..ae0f67fdb44884ffbcdf555fd48b275b64499c85 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/MRT/MRTCompSP27_Device.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/M02CompressibleNavierStokes/M02CompressibleNavierStokes_Device.cu @@ -6,7 +6,7 @@ using namespace vf::basics::constant; using namespace vf::lbm::dir; #include "math.h" -__global__ void LB_Kernel_MRT_Comp_SP_27(real omega, +__global__ void M02CompressibleNavierStokes_Device(real omega, unsigned int* bcMatD, unsigned int* neighborX, unsigned int* neighborY, diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/MRT/MRTCompSP27_Device.cuh b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/M02CompressibleNavierStokes/M02CompressibleNavierStokes_Device.cuh similarity index 56% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/MRT/MRTCompSP27_Device.cuh rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/M02CompressibleNavierStokes/M02CompressibleNavierStokes_Device.cuh index 947ce68259432efe87af84fd9986916e62521397..06542a20c7c40367195fda9bdb3acac339e62f0e 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/MRT/MRTCompSP27_Device.cuh +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/M02CompressibleNavierStokes/M02CompressibleNavierStokes_Device.cuh @@ -1,10 +1,10 @@ -#ifndef LB_KERNEL_MRT_COMP_SP_27_H -#define LB_KERNEL_MRT_COMP_SP_27_H +#ifndef M02CompressibleNavierStokes_Device_H +#define M02CompressibleNavierStokes_Device_H #include <DataTypes.h> #include <curand.h> -__global__ void LB_Kernel_MRT_Comp_SP_27(real omega, +__global__ void M02CompressibleNavierStokes_Device(real omega, unsigned int* bcMatD, unsigned int* neighborX, unsigned int* neighborY, diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/MRT/MRTCompSP27.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/MRT/MRTCompSP27.h deleted file mode 100644 index 07aba55fd4bad868c919284ebf11846ba8f9ebec..0000000000000000000000000000000000000000 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/MRT/MRTCompSP27.h +++ /dev/null @@ -1,17 +0,0 @@ -#ifndef MRT_COMP_SP27_H -#define MRT_COMP_SP27_H - -#include "Kernel/KernelImp.h" - - -class MRTCompSP27 : public KernelImp -{ -public: - static std::shared_ptr<MRTCompSP27> getNewInstance(std::shared_ptr< Parameter> para, int level); - void run(); - -private: - MRTCompSP27(); - MRTCompSP27(std::shared_ptr< Parameter> para, int level); -}; -#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGKPlus/BGKPlusIncompSP27.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/B15IncompressibleNavierStokesBGKplus/B15IncompressibleNavierStokesBGKplus.cu similarity index 50% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGKPlus/BGKPlusIncompSP27.cu rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/B15IncompressibleNavierStokesBGKplus/B15IncompressibleNavierStokesBGKplus.cu index 84a55b89d68f9a1e18c5114f8088a7dee24a4cd1..3bc4fcbecfe656cf1f1b451ab52e1a2f192908f1 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGKPlus/BGKPlusIncompSP27.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/B15IncompressibleNavierStokesBGKplus/B15IncompressibleNavierStokesBGKplus.cu @@ -1,19 +1,19 @@ -#include "BGKPlusIncompSP27.h" +#include "B15IncompressibleNavierStokesBGKplus.h" -#include "BGKPlusIncompSP27_Device.cuh" +#include "B15IncompressibleNavierStokesBGKplus_Device.cuh" #include "Parameter/Parameter.h" #include "cuda/CudaGrid.h" -std::shared_ptr<BGKPlusIncompSP27> BGKPlusIncompSP27::getNewInstance(std::shared_ptr<Parameter> para, int level) +std::shared_ptr<B15IncompressibleNavierStokesBGKplus> B15IncompressibleNavierStokesBGKplus::getNewInstance(std::shared_ptr<Parameter> para, int level) { - return std::shared_ptr<BGKPlusIncompSP27>(new BGKPlusIncompSP27(para, level)); + return std::shared_ptr<B15IncompressibleNavierStokesBGKplus>(new B15IncompressibleNavierStokesBGKplus(para, level)); } -void BGKPlusIncompSP27::run() +void B15IncompressibleNavierStokesBGKplus::run() { vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes); - LB_Kernel_BGK_Plus_Incomp_SP_27 <<< grid.grid, grid.threads >>>( + B15IncompressibleNavierStokesBGKplus_Device <<< grid.grid, grid.threads >>>( para->getParD(level)->omega, para->getParD(level)->typeOfGridNode, para->getParD(level)->neighborX, @@ -25,7 +25,7 @@ void BGKPlusIncompSP27::run() getLastCudaError("LB_Kernel_BGK_Plus_Incomp_SP_27 execution failed"); } -BGKPlusIncompSP27::BGKPlusIncompSP27(std::shared_ptr<Parameter> para, int level) +B15IncompressibleNavierStokesBGKplus::B15IncompressibleNavierStokesBGKplus(std::shared_ptr<Parameter> para, int level) { this->para = para; this->level = level; @@ -35,6 +35,6 @@ BGKPlusIncompSP27::BGKPlusIncompSP27(std::shared_ptr<Parameter> para, int level) } -BGKPlusIncompSP27::BGKPlusIncompSP27() +B15IncompressibleNavierStokesBGKplus::B15IncompressibleNavierStokesBGKplus() { } diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/B15IncompressibleNavierStokesBGKplus/B15IncompressibleNavierStokesBGKplus.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/B15IncompressibleNavierStokesBGKplus/B15IncompressibleNavierStokesBGKplus.h new file mode 100644 index 0000000000000000000000000000000000000000..2426d32cb0262c90aa32e8eaeea40ad9e273b75d --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/B15IncompressibleNavierStokesBGKplus/B15IncompressibleNavierStokesBGKplus.h @@ -0,0 +1,17 @@ +#ifndef B15IncompressibleNavierStokesBGKplus_H +#define B15IncompressibleNavierStokesBGKplus_H + +#include "Kernel/KernelImp.h" + +class B15IncompressibleNavierStokesBGKplus : public KernelImp +{ +public: + static std::shared_ptr<B15IncompressibleNavierStokesBGKplus> getNewInstance(std::shared_ptr< Parameter> para, int level); + void run(); + +private: + B15IncompressibleNavierStokesBGKplus(); + B15IncompressibleNavierStokesBGKplus(std::shared_ptr< Parameter> para, int level); + +}; +#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGKPlus/BGKPlusIncompSP27_Device.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/B15IncompressibleNavierStokesBGKplus/B15IncompressibleNavierStokesBGKplus_Device.cu similarity index 99% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGKPlus/BGKPlusIncompSP27_Device.cu rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/B15IncompressibleNavierStokesBGKplus/B15IncompressibleNavierStokesBGKplus_Device.cu index 1fee181619070e3bc30370a1bc32b949a3af9a2a..49764bc07bc910911dff0668105252f75d062205 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGKPlus/BGKPlusIncompSP27_Device.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/B15IncompressibleNavierStokesBGKplus/B15IncompressibleNavierStokesBGKplus_Device.cu @@ -6,7 +6,7 @@ using namespace vf::basics::constant; using namespace vf::lbm::dir; #include "math.h" -__global__ void LB_Kernel_BGK_Plus_Incomp_SP_27(real omega, +__global__ void B15IncompressibleNavierStokesBGKplus_Device(real omega, unsigned int* bcMatD, unsigned int* neighborX, unsigned int* neighborY, diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/B15IncompressibleNavierStokesBGKplus/B15IncompressibleNavierStokesBGKplus_Device.cuh b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/B15IncompressibleNavierStokesBGKplus/B15IncompressibleNavierStokesBGKplus_Device.cuh new file mode 100644 index 0000000000000000000000000000000000000000..1c812f936541deaba06060fda0367dc60f231e63 --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/B15IncompressibleNavierStokesBGKplus/B15IncompressibleNavierStokesBGKplus_Device.cuh @@ -0,0 +1,17 @@ +#ifndef B15IncompressibleNavierStokesBGKplus_Device_H +#define B15IncompressibleNavierStokesBGKplus_Device_H + +#include <DataTypes.h> +#include <curand.h> + +__global__ void B15IncompressibleNavierStokesBGKplus_Device( + real omega, + unsigned int* bcMatD, + unsigned int* neighborX, + unsigned int* neighborY, + unsigned int* neighborZ, + real* DDStart, + int size_Mat, + bool EvenOrOdd); + +#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGK/BGKIncompSP27.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/B92IncompressibleNavierStokes/B92IncompressibleNavierStokes.cu similarity index 57% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGK/BGKIncompSP27.cu rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/B92IncompressibleNavierStokes/B92IncompressibleNavierStokes.cu index 39bd1f3491d0d70e4734d04ef8a2d6e38cdc6448..730eb540739ffb7461d2b5cd458d4e790ce9d369 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGK/BGKIncompSP27.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/B92IncompressibleNavierStokes/B92IncompressibleNavierStokes.cu @@ -1,19 +1,19 @@ -#include "BGKIncompSP27.h" +#include "B92IncompressibleNavierStokes.h" -#include "BGKIncompSP27_Device.cuh" +#include "B92IncompressibleNavierStokes_Device.cuh" #include "Parameter/Parameter.h" #include "cuda/CudaGrid.h" -std::shared_ptr<BGKIncompSP27> BGKIncompSP27::getNewInstance(std::shared_ptr<Parameter> para, int level) +std::shared_ptr<B92IncompressibleNavierStokes> B92IncompressibleNavierStokes::getNewInstance(std::shared_ptr<Parameter> para, int level) { - return std::shared_ptr<BGKIncompSP27>(new BGKIncompSP27(para, level)); + return std::shared_ptr<B92IncompressibleNavierStokes>(new B92IncompressibleNavierStokes(para, level)); } -void BGKIncompSP27::run() +void B92IncompressibleNavierStokes::run() { vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes); - LB_Kernel_BGK_Incomp_SP_27 <<< grid.grid, grid.threads >>>( + B92IncompressibleNavierStokes_Device <<< grid.grid, grid.threads >>>( para->getParD(level)->omega, para->getParD(level)->typeOfGridNode, para->getParD(level)->neighborX, @@ -25,7 +25,7 @@ void BGKIncompSP27::run() getLastCudaError("LB_Kernel_BGK_Incomp_SP_27 execution failed"); } -BGKIncompSP27::BGKIncompSP27(std::shared_ptr<Parameter> para, int level) +B92IncompressibleNavierStokes::B92IncompressibleNavierStokes(std::shared_ptr<Parameter> para, int level) { this->para = para; this->level = level; @@ -35,6 +35,6 @@ BGKIncompSP27::BGKIncompSP27(std::shared_ptr<Parameter> para, int level) } -BGKIncompSP27::BGKIncompSP27() +B92IncompressibleNavierStokes::B92IncompressibleNavierStokes() { } diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/B92IncompressibleNavierStokes/B92IncompressibleNavierStokes.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/B92IncompressibleNavierStokes/B92IncompressibleNavierStokes.h new file mode 100644 index 0000000000000000000000000000000000000000..782b764460b85b6525d8c3ec89ef25e333eadd40 --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/B92IncompressibleNavierStokes/B92IncompressibleNavierStokes.h @@ -0,0 +1,17 @@ +#ifndef B92IncompressibleNavierStokes_H +#define B92IncompressibleNavierStokes_H + +#include "Kernel/KernelImp.h" + + +class B92IncompressibleNavierStokes : public KernelImp +{ +public: + static std::shared_ptr<B92IncompressibleNavierStokes> getNewInstance(std::shared_ptr< Parameter> para, int level); + void run(); + +private: + B92IncompressibleNavierStokes(); + B92IncompressibleNavierStokes(std::shared_ptr< Parameter> para, int level); +}; +#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGK/BGKIncompSP27_Device.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/B92IncompressibleNavierStokes/B92IncompressibleNavierStokes_Device.cu similarity index 99% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGK/BGKIncompSP27_Device.cu rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/B92IncompressibleNavierStokes/B92IncompressibleNavierStokes_Device.cu index 6cfde7648bdfe8808cb39dd1b80d6537bb3c3280..1d343eb3020816fe3114538a5362162e4a86169e 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGK/BGKIncompSP27_Device.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/B92IncompressibleNavierStokes/B92IncompressibleNavierStokes_Device.cu @@ -6,7 +6,7 @@ using namespace vf::basics::constant; using namespace vf::lbm::dir; #include "math.h" -__global__ void LB_Kernel_BGK_Incomp_SP_27(real omega, +__global__ void B92IncompressibleNavierStokes_Device(real omega, unsigned int* bcMatD, unsigned int* neighborX, unsigned int* neighborY, diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGK/BGKIncompSP27_Device.cuh b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/B92IncompressibleNavierStokes/B92IncompressibleNavierStokes_Device.cuh similarity index 80% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGK/BGKIncompSP27_Device.cuh rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/B92IncompressibleNavierStokes/B92IncompressibleNavierStokes_Device.cuh index f1a90b45238a2df4d93860d7e77cb1242b9fbd90..8b81ebf323cd27faaa60f4b5f75d113a17a2b5e8 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGK/BGKIncompSP27_Device.cuh +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/B92IncompressibleNavierStokes/B92IncompressibleNavierStokes_Device.cuh @@ -4,7 +4,8 @@ #include <DataTypes.h> #include <curand.h> -__global__ void LB_Kernel_BGK_Incomp_SP_27(real omega, +__global__ void B92IncompressibleNavierStokes_Device( + real omega, unsigned int* bcMatD, unsigned int* neighborX, unsigned int* neighborY, diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGK/BGKIncompSP27.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGK/BGKIncompSP27.h deleted file mode 100644 index 04308f6919fd15495761d929da06a583457ba0c0..0000000000000000000000000000000000000000 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGK/BGKIncompSP27.h +++ /dev/null @@ -1,17 +0,0 @@ -#ifndef BGK_INCOMP_SP27_H -#define BGK_INCOMP_SP27_H - -#include "Kernel/KernelImp.h" - - -class BGKIncompSP27 : public KernelImp -{ -public: - static std::shared_ptr<BGKIncompSP27> getNewInstance(std::shared_ptr< Parameter> para, int level); - void run(); - -private: - BGKIncompSP27(); - BGKIncompSP27(std::shared_ptr< Parameter> para, int level); -}; -#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGKPlus/BGKPlusIncompSP27.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGKPlus/BGKPlusIncompSP27.h deleted file mode 100644 index c6a1dcaca1ea267738e4a9cec735273ce4ccdb9c..0000000000000000000000000000000000000000 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGKPlus/BGKPlusIncompSP27.h +++ /dev/null @@ -1,17 +0,0 @@ -#ifndef BGK_PLUS_INCOMP_SP27_H -#define BGK_PLUS_INCOMP_SP27_H - -#include "Kernel/KernelImp.h" - -class BGKPlusIncompSP27 : public KernelImp -{ -public: - static std::shared_ptr<BGKPlusIncompSP27> getNewInstance(std::shared_ptr< Parameter> para, int level); - void run(); - -private: - BGKPlusIncompSP27(); - BGKPlusIncompSP27(std::shared_ptr< Parameter> para, int level); - -}; -#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGKPlus/BGKPlusIncompSP27_Device.cuh b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGKPlus/BGKPlusIncompSP27_Device.cuh deleted file mode 100644 index 7f85f3ca29d4d8d7620e2503df9947fd7e42fe8f..0000000000000000000000000000000000000000 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGKPlus/BGKPlusIncompSP27_Device.cuh +++ /dev/null @@ -1,16 +0,0 @@ -#ifndef LB_KERNEL_BGK_PLUS_INCOMP_SP_27_H -#define LB_KERNEL_BGK_PLUS_INCOMP_SP_27_H - -#include <DataTypes.h> -#include <curand.h> - -__global__ void LB_Kernel_BGK_Plus_Incomp_SP_27(real omega, - unsigned int* bcMatD, - unsigned int* neighborX, - unsigned int* neighborY, - unsigned int* neighborZ, - real* DDStart, - int size_Mat, - bool EvenOrOdd); - -#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cascade/CascadeIncompSP27.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/C06IncompressibleNavierStokes/C06IncompressibleNavierStokes.cu similarity index 57% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cascade/CascadeIncompSP27.cu rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/C06IncompressibleNavierStokes/C06IncompressibleNavierStokes.cu index b060137f2d505886ee02a4b72e372ce8b4d48a78..eebe0587bffd55749ed16cd5bec976cb100cb457 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cascade/CascadeIncompSP27.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/C06IncompressibleNavierStokes/C06IncompressibleNavierStokes.cu @@ -1,19 +1,19 @@ -#include "CascadeIncompSP27.h" +#include "C06IncompressibleNavierStokes.h" #include "CascadeIncompSP27_Device.cuh" #include "Parameter/Parameter.h" #include "cuda/CudaGrid.h" -std::shared_ptr<CascadeIncompSP27> CascadeIncompSP27::getNewInstance(std::shared_ptr<Parameter> para, int level) +std::shared_ptr<C06IncompressibleNavierStokes> C06IncompressibleNavierStokes::getNewInstance(std::shared_ptr<Parameter> para, int level) { - return std::shared_ptr<CascadeIncompSP27>(new CascadeIncompSP27(para, level)); + return std::shared_ptr<C06IncompressibleNavierStokes>(new C06IncompressibleNavierStokes(para, level)); } -void CascadeIncompSP27::run() +void C06IncompressibleNavierStokes::run() { vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes); - LB_Kernel_Cascade_Incomp_SP_27 <<< grid.grid, grid.threads >>>( + C06IncompressibleNavierStokes_Device <<< grid.grid, grid.threads >>>( para->getParD(level)->omega, para->getParD(level)->typeOfGridNode, para->getParD(level)->neighborX, @@ -25,7 +25,7 @@ void CascadeIncompSP27::run() getLastCudaError("LB_Kernel_Cascade_Incomp_SP_27 execution failed"); } -CascadeIncompSP27::CascadeIncompSP27(std::shared_ptr<Parameter> para, int level) +C06IncompressibleNavierStokes::C06IncompressibleNavierStokes(std::shared_ptr<Parameter> para, int level) { this->para = para; this->level = level; @@ -35,6 +35,6 @@ CascadeIncompSP27::CascadeIncompSP27(std::shared_ptr<Parameter> para, int level) } -CascadeIncompSP27::CascadeIncompSP27() +C06IncompressibleNavierStokes::C06IncompressibleNavierStokes() { } diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/C06IncompressibleNavierStokes/C06IncompressibleNavierStokes.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/C06IncompressibleNavierStokes/C06IncompressibleNavierStokes.h new file mode 100644 index 0000000000000000000000000000000000000000..292b37ee75e1f8098845888776089d448e133332 --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/C06IncompressibleNavierStokes/C06IncompressibleNavierStokes.h @@ -0,0 +1,17 @@ +#ifndef C06IncompressibleNavierStokes_H +#define C06IncompressibleNavierStokes_H + +#include "Kernel/KernelImp.h" + +class C06IncompressibleNavierStokes : public KernelImp +{ +public: + static std::shared_ptr<C06IncompressibleNavierStokes> getNewInstance(std::shared_ptr< Parameter> para, int level); + void run(); + +private: + C06IncompressibleNavierStokes(); + C06IncompressibleNavierStokes(std::shared_ptr< Parameter> para, int level); + +}; +#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cascade/CascadeIncompSP27_Device.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/C06IncompressibleNavierStokes/CascadeIncompSP27_Device.cu similarity index 99% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cascade/CascadeIncompSP27_Device.cu rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/C06IncompressibleNavierStokes/CascadeIncompSP27_Device.cu index 0346f12cf609c355aedd4743dd7f971f444d5fe9..3f51dd4044887803a1b6b5f2470639ce4537b8d2 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cascade/CascadeIncompSP27_Device.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/C06IncompressibleNavierStokes/CascadeIncompSP27_Device.cu @@ -6,7 +6,8 @@ using namespace vf::basics::constant; using namespace vf::lbm::dir; #include "math.h" -__global__ void LB_Kernel_Cascade_Incomp_SP_27(real omega, +__global__ void C06IncompressibleNavierStokes_Device( + real omega, unsigned int* bcMatD, unsigned int* neighborX, unsigned int* neighborY, diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cascade/CascadeIncompSP27_Device.cuh b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/C06IncompressibleNavierStokes/CascadeIncompSP27_Device.cuh similarity index 56% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cascade/CascadeIncompSP27_Device.cuh rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/C06IncompressibleNavierStokes/CascadeIncompSP27_Device.cuh index a531fa7bd64b9782f43f800b29fe666504612f1a..daebc3345411812700a8776da1b1f5917dff2f59 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cascade/CascadeIncompSP27_Device.cuh +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/C06IncompressibleNavierStokes/CascadeIncompSP27_Device.cuh @@ -1,10 +1,11 @@ -#ifndef LB_KERNEL_CASCADE_INCOMP_SP_27_H -#define LB_KERNEL_CASCADE_INCOMP_SP_27_H +#ifndef C06IncompressibleNavierStokes_Device_H +#define C06IncompressibleNavierStokes_Device_H #include <DataTypes.h> #include <curand.h> -__global__ void LB_Kernel_Cascade_Incomp_SP_27(real s9, +__global__ void C06IncompressibleNavierStokes_Device( + real s9, unsigned int* bcMatD, unsigned int* neighborX, unsigned int* neighborY, diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cascade/CascadeIncompSP27.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cascade/CascadeIncompSP27.h deleted file mode 100644 index 4ebb0c50a642a943defb7a42ce96cba72d03ddf4..0000000000000000000000000000000000000000 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cascade/CascadeIncompSP27.h +++ /dev/null @@ -1,17 +0,0 @@ -#ifndef CASCADE_INCOMP_P27_H -#define CASCADE_INCOMP_P27_H - -#include "Kernel/KernelImp.h" - -class CascadeIncompSP27 : public KernelImp -{ -public: - static std::shared_ptr<CascadeIncompSP27> getNewInstance(std::shared_ptr< Parameter> para, int level); - void run(); - -private: - CascadeIncompSP27(); - CascadeIncompSP27(std::shared_ptr< Parameter> para, int level); - -}; -#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cumulant1hSP27/Cumulant1hIncompSP27.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cumulant1hSP27/Cumulant1hIncompSP27.cu deleted file mode 100644 index 2cade430786b17567c47264f0638dba259b3192d..0000000000000000000000000000000000000000 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cumulant1hSP27/Cumulant1hIncompSP27.cu +++ /dev/null @@ -1,45 +0,0 @@ -#include "Cumulant1hIncompSP27.h" - -#include "Cumulant1hIncompSP27_Device.cuh" -#include "Parameter/Parameter.h" -#include "cuda/CudaGrid.h" - -std::shared_ptr<Cumulant1hIncompSP27> Cumulant1hIncompSP27::getNewInstance(std::shared_ptr<Parameter> para, int level) -{ - return std::shared_ptr<Cumulant1hIncompSP27>(new Cumulant1hIncompSP27(para, level)); -} - -void Cumulant1hIncompSP27::run() -{ - vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes); - - LB_Kernel_Cum_1h_Incomp_SP_27 <<< grid.grid, grid.threads >>>( - para->getParD(level)->omega, - para->getParD(level)->deltaPhi, - para->getAngularVelocity(), - para->getParD(level)->typeOfGridNode, - para->getParD(level)->neighborX, - para->getParD(level)->neighborY, - para->getParD(level)->neighborZ, - para->getParD(level)->coordinateX, - para->getParD(level)->coordinateY, - para->getParD(level)->coordinateZ, - para->getParD(level)->distributions.f[0], - para->getParD(level)->numberOfNodes, - para->getParD(level)->isEvenTimestep); - getLastCudaError("LB_Kernel_Cum_1h_Incomp_SP_27 execution failed"); -} - -Cumulant1hIncompSP27::Cumulant1hIncompSP27(std::shared_ptr<Parameter> para, int level) -{ - this->para = para; - this->level = level; - - myPreProcessorTypes.push_back(InitSP27); - - -} - -Cumulant1hIncompSP27::Cumulant1hIncompSP27() -{ -} diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cumulant1hSP27/Cumulant1hIncompSP27.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cumulant1hSP27/Cumulant1hIncompSP27.h deleted file mode 100644 index c6086649dc577afd6a00b684769858deb7d85e20..0000000000000000000000000000000000000000 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cumulant1hSP27/Cumulant1hIncompSP27.h +++ /dev/null @@ -1,17 +0,0 @@ -#ifndef CUMULANT_1H_INCOMP_SP27_H -#define CUMULANT_1H_INCOMP_SP27_H - -#include "Kernel/KernelImp.h" - -class Cumulant1hIncompSP27 : public KernelImp -{ -public: - static std::shared_ptr<Cumulant1hIncompSP27> getNewInstance(std::shared_ptr< Parameter> para, int level); - void run(); - -private: - Cumulant1hIncompSP27(); - Cumulant1hIncompSP27(std::shared_ptr< Parameter> para, int level); - -}; -#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantIsoSP27/CumulantIsoIncompSP27.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantIsoSP27/CumulantIsoIncompSP27.h deleted file mode 100644 index 99d2230be843d9acd6ff6e95283e8530a52acbef..0000000000000000000000000000000000000000 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantIsoSP27/CumulantIsoIncompSP27.h +++ /dev/null @@ -1,18 +0,0 @@ -#ifndef CUMULANT_ISO_INCOMP_SP27_H -#define CUMULANT_ISO_INCOMP_SP27_H - -#include "Kernel/KernelImp.h" - -class CumulantIsoIncompSP27 : public KernelImp -{ -public: - static std::shared_ptr<CumulantIsoIncompSP27> getNewInstance(std::shared_ptr< Parameter> para, int level); - void run(); - -private: - CumulantIsoIncompSP27(); - CumulantIsoIncompSP27(std::shared_ptr< Parameter> para, int level); - -}; - -#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantK15/CumulantK15Incomp.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantK15/CumulantK15Incomp.h deleted file mode 100644 index c08737b4331df171efb72594d487da3ea346f8ca..0000000000000000000000000000000000000000 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantK15/CumulantK15Incomp.h +++ /dev/null @@ -1,16 +0,0 @@ -#ifndef CUMULANT_K15_INCOMP_H -#define CUMULANT_K15_INCOMP_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/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantK15/CumulantK15Incomp_Device.cuh b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantK15/CumulantK15Incomp_Device.cuh deleted file mode 100644 index f2b5063f9db6d55b9efb547c0c05d450463e0509..0000000000000000000000000000000000000000 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantK15/CumulantK15Incomp_Device.cuh +++ /dev/null @@ -1,16 +0,0 @@ -#ifndef LB_KERNEL_CUMULANT_K15_IMCOMP_H -#define LB_KERNEL_CUMULANT_K15_IMCOMP_H - -#include <DataTypes.h> -#include <curand.h> - -__global__ void LB_Kernel_CumulantK15Incomp(real s9, - unsigned int* bcMatD, - unsigned int* neighborX, - unsigned int* neighborY, - unsigned int* neighborZ, - real* DDStart, - int size_Mat, - bool EvenOrOdd); - -#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantK15/CumulantK15Incomp.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokes/K15IncompressibleNavierStokes.cu similarity index 53% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantK15/CumulantK15Incomp.cu rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokes/K15IncompressibleNavierStokes.cu index c597924193d859a35dddaa7b37a56e21d265ceba..bfbf84b4073a87cbbf833ea63ae9d457ac5965f4 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantK15/CumulantK15Incomp.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokes/K15IncompressibleNavierStokes.cu @@ -1,19 +1,19 @@ -#include "CumulantK15Incomp.h" +#include "K15IncompressibleNavierStokes.h" -#include "CumulantK15Incomp_Device.cuh" +#include "K15IncompressibleNavierStokes_Device.cuh" #include "Parameter/Parameter.h" #include "cuda/CudaGrid.h" -std::shared_ptr<CumulantK15Incomp> CumulantK15Incomp::getNewInstance(std::shared_ptr<Parameter> para, int level) +std::shared_ptr<K15IncompressibleNavierStokes> K15IncompressibleNavierStokes::getNewInstance(std::shared_ptr<Parameter> para, int level) { - return std::shared_ptr<CumulantK15Incomp>(new CumulantK15Incomp(para, level)); + return std::shared_ptr<K15IncompressibleNavierStokes>(new K15IncompressibleNavierStokes(para, level)); } -void CumulantK15Incomp::run() +void K15IncompressibleNavierStokes::run() { vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes); - LB_Kernel_CumulantK15Incomp <<< grid.grid, grid.threads >>>( + K15IncompressibleNavierStokes_Device <<< grid.grid, grid.threads >>>( para->getParD(level)->omega, para->getParD(level)->typeOfGridNode, para->getParD(level)->neighborX, @@ -25,7 +25,7 @@ void CumulantK15Incomp::run() getLastCudaError("LB_Kernel_CumulantK15Incomp execution failed"); } -CumulantK15Incomp::CumulantK15Incomp(std::shared_ptr<Parameter> para, int level) +K15IncompressibleNavierStokes::K15IncompressibleNavierStokes(std::shared_ptr<Parameter> para, int level) { this->para = para; this->level = level; @@ -35,6 +35,6 @@ CumulantK15Incomp::CumulantK15Incomp(std::shared_ptr<Parameter> para, int level) } -CumulantK15Incomp::CumulantK15Incomp() +K15IncompressibleNavierStokes::K15IncompressibleNavierStokes() { } diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokes/K15IncompressibleNavierStokes.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokes/K15IncompressibleNavierStokes.h new file mode 100644 index 0000000000000000000000000000000000000000..40d06dfaef95cdd71ab1d11fbfc356087345ce91 --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokes/K15IncompressibleNavierStokes.h @@ -0,0 +1,16 @@ +#ifndef K15IncompressibleNavierStokes_H +#define K15IncompressibleNavierStokes_H + +#include "Kernel/KernelImp.h" + +class K15IncompressibleNavierStokes : public KernelImp +{ +public: + static std::shared_ptr<K15IncompressibleNavierStokes> getNewInstance(std::shared_ptr< Parameter> para, int level); + void run(); + +private: + K15IncompressibleNavierStokes(); + K15IncompressibleNavierStokes(std::shared_ptr< Parameter> para, int level); +}; +#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantK15/CumulantK15Incomp_Device.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokes/K15IncompressibleNavierStokes_Device.cu similarity index 99% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantK15/CumulantK15Incomp_Device.cu rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokes/K15IncompressibleNavierStokes_Device.cu index 9fcfeaa97d97b9bfcf2ad0227464cbcabbc28f44..5b666eb73c7a19d472d26f2047b04462f83d0d23 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantK15/CumulantK15Incomp_Device.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokes/K15IncompressibleNavierStokes_Device.cu @@ -6,7 +6,8 @@ using namespace vf::basics::constant; using namespace vf::lbm::dir; #include "math.h" -__global__ void LB_Kernel_CumulantK15Incomp(real omega, +__global__ void K15IncompressibleNavierStokes_Device( + real omega, unsigned int* bcMatD, unsigned int* neighborX, unsigned int* neighborY, diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokes/K15IncompressibleNavierStokes_Device.cuh b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokes/K15IncompressibleNavierStokes_Device.cuh new file mode 100644 index 0000000000000000000000000000000000000000..9252b17df43b405e7098359d3c5ff7406cc90759 --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokes/K15IncompressibleNavierStokes_Device.cuh @@ -0,0 +1,17 @@ +#ifndef K15IncompressibleNavierStokes_Device_H +#define K15IncompressibleNavierStokes_Device_H + +#include <DataTypes.h> +#include <curand.h> + +__global__ void K15IncompressibleNavierStokes_Device( + real s9, + unsigned int* bcMatD, + unsigned int* neighborX, + unsigned int* neighborY, + unsigned int* neighborZ, + real* DDStart, + int size_Mat, + bool EvenOrOdd); + +#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantIsoSP27/CumulantIsoIncompSP27.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokesIsoCheck/K15IncompressibleNavierStokesIsoCheck.cu similarity index 55% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantIsoSP27/CumulantIsoIncompSP27.cu rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokesIsoCheck/K15IncompressibleNavierStokesIsoCheck.cu index 840067da7f34a4415b1b14458ae0fc8d316e366d..e975bf9df0dce6ff1ebd3d463c3061013b526354 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantIsoSP27/CumulantIsoIncompSP27.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokesIsoCheck/K15IncompressibleNavierStokesIsoCheck.cu @@ -1,19 +1,19 @@ -#include "CumulantIsoIncompSP27.h" +#include "K15IncompressibleNavierStokesIsoCheck.h" -#include "CumulantIsoIncompSP27_Device.cuh" +#include "K15IncompressibleNavierStokesIsoCheck_Device.cuh" #include "Parameter/Parameter.h" #include "cuda/CudaGrid.h" -std::shared_ptr<CumulantIsoIncompSP27> CumulantIsoIncompSP27::getNewInstance(std::shared_ptr<Parameter> para, int level) +std::shared_ptr<K15IncompressibleNavierStokesIsoCheck> K15IncompressibleNavierStokesIsoCheck::getNewInstance(std::shared_ptr<Parameter> para, int level) { - return std::shared_ptr<CumulantIsoIncompSP27>(new CumulantIsoIncompSP27(para, level)); + return std::shared_ptr<K15IncompressibleNavierStokesIsoCheck>(new K15IncompressibleNavierStokesIsoCheck(para, level)); } -void CumulantIsoIncompSP27::run() +void K15IncompressibleNavierStokesIsoCheck::run() { vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes); - LB_Kernel_Cum_IsoTest_Incomp_SP_27 <<< grid.grid, grid.threads >>>( + K15IncompressibleNavierStokesIsoCheck_Device <<< grid.grid, grid.threads >>>( para->getParD(level)->omega, para->getParD(level)->typeOfGridNode, para->getParD(level)->neighborX, @@ -28,7 +28,7 @@ void CumulantIsoIncompSP27::run() getLastCudaError("LB_Kernel_Cum_IsoTest_Incomp_SP_27 execution failed"); } -CumulantIsoIncompSP27::CumulantIsoIncompSP27(std::shared_ptr<Parameter> para, int level) +K15IncompressibleNavierStokesIsoCheck::K15IncompressibleNavierStokesIsoCheck(std::shared_ptr<Parameter> para, int level) { this->para = para; this->level = level; @@ -38,6 +38,6 @@ CumulantIsoIncompSP27::CumulantIsoIncompSP27(std::shared_ptr<Parameter> para, in } -CumulantIsoIncompSP27::CumulantIsoIncompSP27() +K15IncompressibleNavierStokesIsoCheck::K15IncompressibleNavierStokesIsoCheck() { } diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokesIsoCheck/K15IncompressibleNavierStokesIsoCheck.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokesIsoCheck/K15IncompressibleNavierStokesIsoCheck.h new file mode 100644 index 0000000000000000000000000000000000000000..c2871b389c421b191cb65b36fd8cab3c6c35cb2c --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokesIsoCheck/K15IncompressibleNavierStokesIsoCheck.h @@ -0,0 +1,18 @@ +#ifndef K15IncompressibleNavierStokesIsoCheck_H +#define K15IncompressibleNavierStokesIsoCheck_H + +#include "Kernel/KernelImp.h" + +class K15IncompressibleNavierStokesIsoCheck : public KernelImp +{ +public: + static std::shared_ptr<K15IncompressibleNavierStokesIsoCheck> getNewInstance(std::shared_ptr< Parameter> para, int level); + void run(); + +private: + K15IncompressibleNavierStokesIsoCheck(); + K15IncompressibleNavierStokesIsoCheck(std::shared_ptr< Parameter> para, int level); + +}; + +#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantIsoSP27/CumulantIsoIncompSP27_Device.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokesIsoCheck/K15IncompressibleNavierStokesIsoCheck_Device.cu similarity index 99% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantIsoSP27/CumulantIsoIncompSP27_Device.cu rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokesIsoCheck/K15IncompressibleNavierStokesIsoCheck_Device.cu index 623d3c2c26b2c4d8fdfdaa718ca92662dfe5b548..6c91d9bfc4cf2c39b812a493986da29b45165407 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantIsoSP27/CumulantIsoIncompSP27_Device.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokesIsoCheck/K15IncompressibleNavierStokesIsoCheck_Device.cu @@ -6,7 +6,8 @@ using namespace vf::basics::constant; using namespace vf::lbm::dir; #include "math.h" -__global__ void LB_Kernel_Cum_IsoTest_Incomp_SP_27(real omega, +__global__ void K15IncompressibleNavierStokesIsoCheck_Device( + real omega, unsigned int* bcMatD, unsigned int* neighborX, unsigned int* neighborY, diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantIsoSP27/CumulantIsoIncompSP27_Device.cuh b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokesIsoCheck/K15IncompressibleNavierStokesIsoCheck_Device.cuh similarity index 56% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantIsoSP27/CumulantIsoIncompSP27_Device.cuh rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokesIsoCheck/K15IncompressibleNavierStokesIsoCheck_Device.cuh index 57dc7180ce7900333b4071db409a03ac847dd641..b392245582cec7a098a9d7bf2583d75a578b6993 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantIsoSP27/CumulantIsoIncompSP27_Device.cuh +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokesIsoCheck/K15IncompressibleNavierStokesIsoCheck_Device.cuh @@ -1,10 +1,11 @@ -#ifndef LB_KERNEL_CUM_ISO_TEST_INCOMP_SP_27 -#define LB_KERNEL_CUM_ISO_TEST_INCOMP_SP_27 +#ifndef K15IncompressibleNavierStokesIsoCheck_Device_27 +#define K15IncompressibleNavierStokesIsoCheck_Device_27 #include <DataTypes.h> #include <curand.h> -__global__ void LB_Kernel_Cum_IsoTest_Incomp_SP_27(real omega, +__global__ void K15IncompressibleNavierStokesIsoCheck_Device( + real omega, unsigned int* bcMatD, unsigned int* neighborX, unsigned int* neighborY, diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokesRotatingVelocityField/K15IncompressibleNavierStokesRotatingVelocityField.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokesRotatingVelocityField/K15IncompressibleNavierStokesRotatingVelocityField.cu new file mode 100644 index 0000000000000000000000000000000000000000..e405246db6d7ef92d86f23fd36231dfe6ea7d831 --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokesRotatingVelocityField/K15IncompressibleNavierStokesRotatingVelocityField.cu @@ -0,0 +1,45 @@ +#include "K15IncompressibleNavierStokesRotatingVelocityField.h" + +#include "K15IncompressibleNavierStokesRotatingVelocityField_Device.cuh" +#include "Parameter/Parameter.h" +#include "cuda/CudaGrid.h" + +std::shared_ptr<K15IncompressibleNavierStokesRotatingVelocityField> K15IncompressibleNavierStokesRotatingVelocityField::getNewInstance(std::shared_ptr<Parameter> para, int level) +{ + return std::shared_ptr<K15IncompressibleNavierStokesRotatingVelocityField>(new K15IncompressibleNavierStokesRotatingVelocityField(para, level)); +} + +void K15IncompressibleNavierStokesRotatingVelocityField::run() +{ + vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes); + + K15IncompressibleNavierStokesRotatingVelocityField_Device <<< grid.grid, grid.threads >>>( + para->getParD(level)->omega, + para->getParD(level)->deltaPhi, + para->getAngularVelocity(), + para->getParD(level)->typeOfGridNode, + para->getParD(level)->neighborX, + para->getParD(level)->neighborY, + para->getParD(level)->neighborZ, + para->getParD(level)->coordinateX, + para->getParD(level)->coordinateY, + para->getParD(level)->coordinateZ, + para->getParD(level)->distributions.f[0], + para->getParD(level)->numberOfNodes, + para->getParD(level)->isEvenTimestep); + getLastCudaError("K15IncompressibleNavierStokesRotatingVelocityField_Device execution failed"); +} + +K15IncompressibleNavierStokesRotatingVelocityField::K15IncompressibleNavierStokesRotatingVelocityField(std::shared_ptr<Parameter> para, int level) +{ + this->para = para; + this->level = level; + + myPreProcessorTypes.push_back(InitSP27); + + +} + +K15IncompressibleNavierStokesRotatingVelocityField::K15IncompressibleNavierStokesRotatingVelocityField() +{ +} diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokesRotatingVelocityField/K15IncompressibleNavierStokesRotatingVelocityField.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokesRotatingVelocityField/K15IncompressibleNavierStokesRotatingVelocityField.h new file mode 100644 index 0000000000000000000000000000000000000000..94858833254f2e2b40c0ae3ef00eaa0789ded88b --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokesRotatingVelocityField/K15IncompressibleNavierStokesRotatingVelocityField.h @@ -0,0 +1,18 @@ +#ifndef K15IncompressibleNavierStokesRotatingVelocityField_H +#define K15IncompressibleNavierStokesRotatingVelocityField_H + +#include "Kernel/KernelImp.h" + +class K15IncompressibleNavierStokesRotatingVelocityField : public KernelImp +{ +public: + static std::shared_ptr<K15IncompressibleNavierStokesRotatingVelocityField> getNewInstance(std::shared_ptr< Parameter> para, int level); + void run(); + +private: + K15IncompressibleNavierStokesRotatingVelocityField(); + K15IncompressibleNavierStokesRotatingVelocityField(std::shared_ptr< Parameter> para, int level); + +}; + +#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cumulant1hSP27/Cumulant1hIncompSP27_Device.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokesRotatingVelocityField/K15IncompressibleNavierStokesRotatingVelocityField_Device.cu similarity index 99% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cumulant1hSP27/Cumulant1hIncompSP27_Device.cu rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokesRotatingVelocityField/K15IncompressibleNavierStokesRotatingVelocityField_Device.cu index ac88396b42483e3178869be585124fb031d099ec..082a5dc3882a498838d174e0c7b24ae134dca622 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cumulant1hSP27/Cumulant1hIncompSP27_Device.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokesRotatingVelocityField/K15IncompressibleNavierStokesRotatingVelocityField_Device.cu @@ -6,7 +6,8 @@ using namespace vf::basics::constant; using namespace vf::lbm::dir; #include "math.h" -__global__ void LB_Kernel_Cum_1h_Incomp_SP_27(real omega, +__global__ void K15IncompressibleNavierStokesRotatingVelocityField_Device( + real omega, real deltaPhi, real angularVelocity, unsigned int* bcMatD, diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cumulant1hSP27/Cumulant1hIncompSP27_Device.cuh b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokesRotatingVelocityField/K15IncompressibleNavierStokesRotatingVelocityField_Device.cuh similarity index 56% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cumulant1hSP27/Cumulant1hIncompSP27_Device.cuh rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokesRotatingVelocityField/K15IncompressibleNavierStokesRotatingVelocityField_Device.cuh index 0679e770f244ad2e2c59c9cacf9d3524a640a42e..b313fcb8f2e7aa4544249ce26f6ada3cf85c7a79 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cumulant1hSP27/Cumulant1hIncompSP27_Device.cuh +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokesRotatingVelocityField/K15IncompressibleNavierStokesRotatingVelocityField_Device.cuh @@ -1,10 +1,11 @@ -#ifndef LB_KERNEL_CUM_1H_INCOMP_SP_27_H -#define LB_KERNEL_CUM_1H_INCOMP_SP_27_H +#ifndef K15IncompressibleNavierStokesRotatingVelocityField_Device_H +#define K15IncompressibleNavierStokesRotatingVelocityField_Device_H #include <DataTypes.h> #include <curand.h> -__global__ void LB_Kernel_Cum_1h_Incomp_SP_27(real omega, +__global__ void K15IncompressibleNavierStokesRotatingVelocityField_Device( + real omega, real deltaPhi, real angularVelocity, unsigned int* bcMatD, diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/MRT/MRTIncompSP27.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/M02IncompressibleNavierStokes/M02IncompressibleNavierStokes.cu similarity index 57% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/MRT/MRTIncompSP27.cu rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/M02IncompressibleNavierStokes/M02IncompressibleNavierStokes.cu index daa90091fe092a98741d0764e2327f3ce4c9d2bc..71660b82ed7b439f912679d8bd4617e6b6e2a012 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/MRT/MRTIncompSP27.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/M02IncompressibleNavierStokes/M02IncompressibleNavierStokes.cu @@ -1,19 +1,19 @@ -#include "MRTIncompSP27.h" +#include "M02IncompressibleNavierStokes.h" -#include "MRTIncompSP27_Device.cuh" +#include "M02IncompressibleNavierStokes_Device.cuh" #include "Parameter/Parameter.h" #include "cuda/CudaGrid.h" -std::shared_ptr<MRTIncompSP27> MRTIncompSP27::getNewInstance(std::shared_ptr<Parameter> para, int level) +std::shared_ptr<M02IncompressibleNavierStokes> M02IncompressibleNavierStokes::getNewInstance(std::shared_ptr<Parameter> para, int level) { - return std::shared_ptr<MRTIncompSP27>(new MRTIncompSP27(para, level)); + return std::shared_ptr<M02IncompressibleNavierStokes>(new M02IncompressibleNavierStokes(para, level)); } -void MRTIncompSP27::run() +void M02IncompressibleNavierStokes::run() { vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes); - LB_Kernel_MRT_Incomp_SP_27 <<< grid.grid, grid.threads >>>( + M02IncompressibleNavierStokes_Device <<< grid.grid, grid.threads >>>( para->getParD(level)->omega, para->getParD(level)->typeOfGridNode, para->getParD(level)->neighborX, @@ -25,7 +25,7 @@ void MRTIncompSP27::run() getLastCudaError("LB_Kernel_MRT_Incomp_SP_27 execution failed"); } -MRTIncompSP27::MRTIncompSP27(std::shared_ptr<Parameter> para, int level) +M02IncompressibleNavierStokes::M02IncompressibleNavierStokes(std::shared_ptr<Parameter> para, int level) { this->para = para; this->level = level; @@ -35,6 +35,6 @@ MRTIncompSP27::MRTIncompSP27(std::shared_ptr<Parameter> para, int level) } -MRTIncompSP27::MRTIncompSP27() +M02IncompressibleNavierStokes::M02IncompressibleNavierStokes() { } diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/M02IncompressibleNavierStokes/M02IncompressibleNavierStokes.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/M02IncompressibleNavierStokes/M02IncompressibleNavierStokes.h new file mode 100644 index 0000000000000000000000000000000000000000..dde5846114b84abb6c276f155e3f583f7de0ecd6 --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/M02IncompressibleNavierStokes/M02IncompressibleNavierStokes.h @@ -0,0 +1,16 @@ +#ifndef M02IncompressibleNavierStokes_H +#define M02IncompressibleNavierStokes_H + +#include "Kernel/KernelImp.h" + +class M02IncompressibleNavierStokes : public KernelImp +{ +public: + static std::shared_ptr<M02IncompressibleNavierStokes> getNewInstance(std::shared_ptr<Parameter> para, int level); + void run(); + +private: + M02IncompressibleNavierStokes(); + M02IncompressibleNavierStokes(std::shared_ptr< Parameter> para, int level); +}; +#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/MRT/MRTIncompSP27_Device.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/M02IncompressibleNavierStokes/M02IncompressibleNavierStokes_Device.cu similarity index 99% rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/MRT/MRTIncompSP27_Device.cu rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/M02IncompressibleNavierStokes/M02IncompressibleNavierStokes_Device.cu index 2fbc7d64d5f10a2c6313647e508fbb2192dd435b..4c94c960e0bf9521c558ede0c00076f2bfe82ce4 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/MRT/MRTIncompSP27_Device.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/M02IncompressibleNavierStokes/M02IncompressibleNavierStokes_Device.cu @@ -6,7 +6,8 @@ using namespace vf::basics::constant; using namespace vf::lbm::dir; #include "math.h" -__global__ void LB_Kernel_MRT_Incomp_SP_27(real omega, +__global__ void M02IncompressibleNavierStokes_Device( + real omega, unsigned int* bcMatD, unsigned int* neighborX, unsigned int* neighborY, diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/M02IncompressibleNavierStokes/M02IncompressibleNavierStokes_Device.cuh b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/M02IncompressibleNavierStokes/M02IncompressibleNavierStokes_Device.cuh new file mode 100644 index 0000000000000000000000000000000000000000..c62c73b10fcab18f8e4c7b3189dc10204a60de28 --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/M02IncompressibleNavierStokes/M02IncompressibleNavierStokes_Device.cuh @@ -0,0 +1,17 @@ +#ifndef M02IncompressibleNavierStokes_Device_H +#define M02IncompressibleNavierStokes_Device_H + +#include <DataTypes.h> +#include <curand.h> + +__global__ void M02IncompressibleNavierStokes_Device( + real omega, + unsigned int* bcMatD, + unsigned int* neighborX, + unsigned int* neighborY, + unsigned int* neighborZ, + real* DDStart, + int size_Mat, + bool EvenOrOdd); + +#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/MRT/MRTIncompSP27.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/MRT/MRTIncompSP27.h deleted file mode 100644 index dd467be23b888d0968520f706b3e2dd3da6c3fd7..0000000000000000000000000000000000000000 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/MRT/MRTIncompSP27.h +++ /dev/null @@ -1,16 +0,0 @@ -#ifndef MRT_INCOMP_SP27_H -#define MRT_INCOMP_SP27_H - -#include "Kernel/KernelImp.h" - -class MRTIncompSP27 : public KernelImp -{ -public: - static std::shared_ptr<MRTIncompSP27> getNewInstance(std::shared_ptr<Parameter> para, int level); - void run(); - -private: - MRTIncompSP27(); - MRTIncompSP27(std::shared_ptr< Parameter> para, int level); -}; -#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/MRT/MRTIncompSP27_Device.cuh b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/MRT/MRTIncompSP27_Device.cuh deleted file mode 100644 index d3a9fcea7c1a53e4084acf8dc5f1f815d0da967d..0000000000000000000000000000000000000000 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/MRT/MRTIncompSP27_Device.cuh +++ /dev/null @@ -1,16 +0,0 @@ -#ifndef LB_KERNEL_MRT_INCOMP_SP_27_H -#define LB_KERNEL_MRT_INCOMP_SP_27_H - -#include <DataTypes.h> -#include <curand.h> - -__global__ void LB_Kernel_MRT_Incomp_SP_27(real omega, - unsigned int* bcMatD, - unsigned int* neighborX, - unsigned int* neighborY, - unsigned int* neighborZ, - real* DDStart, - int size_Mat, - bool EvenOrOdd); - -#endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Utilities/KernelFactory/KernelFactoryImp.cpp b/src/gpu/VirtualFluids_GPU/Kernel/Utilities/KernelFactory/KernelFactoryImp.cpp index 694251a71e2562cdc66cfe478d27dbc4256aa456..beb0cc930c9cca91168981c9ec2e9e46a906442d 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Utilities/KernelFactory/KernelFactoryImp.cpp +++ b/src/gpu/VirtualFluids_GPU/Kernel/Utilities/KernelFactory/KernelFactoryImp.cpp @@ -7,12 +7,12 @@ #include "Kernel/Utilities/KernelTypes.h" //LBM kernel (compressible) -#include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGK/BGKCompSP27.h" -#include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKUnified/BGKUnified.h" -#include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKPlus/BGKPlusCompSP27.h" -#include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/MRT/MRTCompSP27.h" -#include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cascade/CascadeCompSP27.h" -#include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cumulant/CumulantCompSP27.h" +#include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B92CompressibleNavierStokes/B92CompressibleNavierStokes.h" +#include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B15CompressibleNavierStokesBGKplusUnified/B15CompressibleNavierStokesBGKplusUnified.h" +#include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B15CompressibleNavierStokesBGKplus/B15CompressibleNavierStokesBGKplus.h" +#include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/M02CompressibleNavierStokes/M02CompressibleNavierStokes.h" +#include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/C06CompressibleNavierStokes/C06CompressibleNavierStokes.h" +#include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K08CompressibleNavierStokes/K08CompressibleNavierStokes.h" #include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K15CompressibleNavierStokes/K15CompressibleNavierStokes.h" #include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K15CompressibleNavierStokesBulkViscosity/K15CompressibleNavierStokesBulkViscosity.h" #include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K15CompressibleNavierStokesSponge/K15CompressibleNavierStokesSponge.h" @@ -21,18 +21,18 @@ #include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K17CompressibleNavierStokesChimeraLegacy/K17CompressibleNavierStokesChimeraLegacy.h" #include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K17CompressibleNavierStokesBulkViscosity/K17CompressibleNavierStokesBulkViscosity.h" #include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K17CompressibleNavierStokesUnified/K17CompressibleNavierStokesUnified.h" -#include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantAll4/CumulantAll4CompSP27.h" -#include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK18/CumulantK18Comp.h" -#include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK20/CumulantK20Comp.h" +#include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants.h" +#include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K18CompressibleNavierStokes/K18CompressibleNavierStokes.h" +#include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K20CompressibleNavierStokes/K20CompressibleNavierStokes.h" //LBM kernel (inkompressible) -#include "Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGK/BGKIncompSP27.h" -#include "Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGKPlus/BGKPlusIncompSP27.h" -#include "Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cascade/CascadeIncompSP27.h" -#include "Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cumulant1hSP27/Cumulant1hIncompSP27.h" -#include "Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantIsoSP27/CumulantIsoIncompSP27.h" -#include "Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantK15/CumulantK15Incomp.h" -#include "Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/MRT/MRTIncompSP27.h" +#include "Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/B92IncompressibleNavierStokes/B92IncompressibleNavierStokes.h" +#include "Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/B15IncompressibleNavierStokesBGKplus/B15IncompressibleNavierStokesBGKplus.h" +#include "Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/C06IncompressibleNavierStokes/C06IncompressibleNavierStokes.h" +#include "Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokesRotatingVelocityField/K15IncompressibleNavierStokesRotatingVelocityField.h" +//#include "Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokesIsoTest/K15IncompressibleNavierStokesIsoTest.h" +#include "Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokes/K15IncompressibleNavierStokes.h" +#include "Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/M02IncompressibleNavierStokes/M02IncompressibleNavierStokes.h" //advection diffusion kernel (compressible) #include "Kernel/Kernels/BasicKernels/AdvectionDiffusion/Compressible/Mod27/ADComp27/ADComp27.h" @@ -101,22 +101,22 @@ std::shared_ptr<Kernel> KernelFactoryImp::makeKernel(std::shared_ptr<Parameter> std::shared_ptr<CheckParameterStrategy> checkStrategy; if (kernel == CollisionKernel::Compressible::BGK) { - newKernel = BGKCompSP27::getNewInstance(para, level); // compressible + newKernel = B92CompressibleNavierStokes::getNewInstance(para, level); // compressible checkStrategy = FluidFlowCompStrategy::getInstance(); // || } else if (kernel == CollisionKernel::Compressible::BGKUnified) { // \/ - newKernel = std::make_shared<vf::gpu::BGKUnified>(para, level); + newKernel = std::make_shared<vf::gpu::B15CompressibleNavierStokesBGKplusUnified>(para, level); checkStrategy = FluidFlowCompStrategy::getInstance(); } else if (kernel == CollisionKernel::Compressible::BGKPlus) { - newKernel = BGKPlusCompSP27::getNewInstance(para, level); + newKernel = B15CompressibleNavierStokesBGKplus::getNewInstance(para, level); checkStrategy = FluidFlowCompStrategy::getInstance(); } else if (kernel == CollisionKernel::Compressible::MRT) { - newKernel = MRTCompSP27::getNewInstance(para, level); + newKernel = M02CompressibleNavierStokes::getNewInstance(para, level); checkStrategy = FluidFlowCompStrategy::getInstance(); } else if (kernel == CollisionKernel::Compressible::Cascade) { - newKernel = CascadeCompSP27::getNewInstance(para, level); + newKernel = C06CompressibleNavierStokes::getNewInstance(para, level); checkStrategy = FluidFlowCompStrategy::getInstance(); } else if (kernel == CollisionKernel::Compressible::CumulantClassic) { - newKernel = CumulantCompSP27::getNewInstance(para, level); + newKernel = K08CompressibleNavierStokes::getNewInstance(para, level); checkStrategy = FluidFlowCompStrategy::getInstance(); } else if (kernel == CollisionKernel::Compressible::CumulantK15Unified) { newKernel = std::make_shared<vf::gpu::K15CompressibleNavierStokesUnified>(para, level); @@ -151,13 +151,13 @@ std::shared_ptr<Kernel> KernelFactoryImp::makeKernel(std::shared_ptr<Parameter> } checkStrategy = FluidFlowCompStrategy::getInstance(); } else if (kernel == CollisionKernel::Compressible::CumulantAll4SP27) { - newKernel = CumulantAll4CompSP27::getNewInstance(para, level); + newKernel = K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants::getNewInstance(para, level); checkStrategy = FluidFlowCompStrategy::getInstance(); } else if (kernel == CollisionKernel::Compressible::CumulantK18) { - newKernel = CumulantK18Comp::getNewInstance(para, level); + newKernel = K18CompressibleNavierStokes::getNewInstance(para, level); checkStrategy = FluidFlowCompStrategy::getInstance(); } else if (kernel == CollisionKernel::Compressible::CumulantK20) { - newKernel = CumulantK20Comp::getNewInstance(para, level); + newKernel = K20CompressibleNavierStokes::getNewInstance(para, level); checkStrategy = FluidFlowCompStrategy::getInstance(); } else if (kernel == CollisionKernel::Compressible::K15CompressibleNavierStokes) { newKernel = K15CompressibleNavierStokes::getNewInstance(para, level); @@ -170,25 +170,25 @@ std::shared_ptr<Kernel> KernelFactoryImp::makeKernel(std::shared_ptr<Parameter> checkStrategy = FluidFlowCompStrategy::getInstance(); // compressible } //=============== else if ( kernel == CollisionKernel::Incompressible::BGK) { // incompressible - newKernel = BGKIncompSP27::getNewInstance(para, level); // || + newKernel = B92IncompressibleNavierStokes::getNewInstance(para, level); // || checkStrategy = FluidFlowIncompStrategy::getInstance(); // \/ } else if (kernel == CollisionKernel::Incompressible::BGKPlus) { - newKernel = BGKPlusIncompSP27::getNewInstance(para, level); + newKernel = B15IncompressibleNavierStokesBGKplus::getNewInstance(para, level); checkStrategy = FluidFlowIncompStrategy::getInstance(); } else if (kernel == CollisionKernel::Incompressible::MRT) { - newKernel = MRTIncompSP27::getNewInstance(para, level); + newKernel = M02IncompressibleNavierStokes::getNewInstance(para, level); checkStrategy = FluidFlowIncompStrategy::getInstance(); } else if (kernel == CollisionKernel::Incompressible::Cascade) { - newKernel = CascadeIncompSP27::getNewInstance(para, level); + newKernel = C06IncompressibleNavierStokes::getNewInstance(para, level); checkStrategy = FluidFlowIncompStrategy::getInstance(); } else if (kernel == CollisionKernel::Incompressible::Cumulant1h) { - newKernel = Cumulant1hIncompSP27::getNewInstance(para, level); - checkStrategy = FluidFlowIncompStrategy::getInstance(); - } else if (kernel == CollisionKernel::Incompressible::CumulantIsometric) { - newKernel = CumulantIsoIncompSP27::getNewInstance(para, level); + newKernel = K15IncompressibleNavierStokesRotatingVelocityField::getNewInstance(para, level); checkStrategy = FluidFlowIncompStrategy::getInstance(); + //} else if (kernel == CollisionKernel::Incompressible::CumulantIsometric) { + // newKernel = K15IncompressibleNavierStokesIsoTest::getNewInstance(para, level); + // checkStrategy = FluidFlowIncompStrategy::getInstance(); } else if (kernel == CollisionKernel::Incompressible::CumulantK15) { // /\ // - newKernel = CumulantK15Incomp::getNewInstance(para, level); // || + newKernel = K15IncompressibleNavierStokes::getNewInstance(para, level); // || checkStrategy = FluidFlowIncompStrategy::getInstance(); // incompressible } //=============== else if (kernel == CollisionKernel::PorousMedia::CumulantOne) { // porous media