From 1ed15dfd360a4c431994185e35730416b8535e2b Mon Sep 17 00:00:00 2001 From: Martin Schoenherr <m.schoenherr@tu-braunschweig.de> Date: Fri, 7 Jul 2023 08:40:10 +0200 Subject: [PATCH] change the kernel names --- src/gpu/VirtualFluids_GPU/GPU/GPU_Kernels.cuh | 21 +++--- .../B15CompressibleNavierStokesBGKplus.cu | 40 +++++++++++ .../B15CompressibleNavierStokesBGKplus.h | 17 +++++ ...CompressibleNavierStokesBGKplus_Device.cu} | 2 +- ...CompressibleNavierStokesBGKplus_Device.cuh | 16 +++++ ...CompressibleNavierStokesBGKplusUnified.cu} | 6 +- ...15CompressibleNavierStokesBGKplusUnified.h | 22 ++++++ .../B92CompressibleNavierStokes.cu} | 16 ++--- .../B92CompressibleNavierStokes.h | 16 +++++ .../B92CompressibleNavierStokes_Device.cu} | 2 +- .../B92CompressibleNavierStokes_Device.cuh | 17 +++++ .../FluidFlow/Compressible/BGK/BGKCompSP27.h | 16 ----- .../Compressible/BGK/BGKCompSP27_Device.cuh | 16 ----- .../Compressible/BGKPlus/BGKPlusCompSP27.cu | 40 ----------- .../Compressible/BGKPlus/BGKPlusCompSP27.h | 17 ----- .../BGKPlus/BGKPlusCompSP27_Device.cuh | 15 ---- .../Compressible/BGKUnified/BGKUnified.h | 22 ------ .../C06CompressibleNavierStokes.cu} | 16 ++--- .../C06CompressibleNavierStokes.h | 16 +++++ .../C06CompressibleNavierStokes_Device.cu} | 2 +- .../C06CompressibleNavierStokes_Device.cuh} | 7 +- .../Compressible/Cascade/CascadeCompSP27.h | 16 ----- .../Compressible/Cumulant/CumulantCompSP27.h | 17 ----- .../CumulantAll4/CumulantAll4CompSP27.cu | 39 ----------- .../CumulantAll4/CumulantAll4CompSP27.h | 16 ----- .../CumulantAll4CompSP27_Device.cuh | 18 ----- .../CumulantK18/CumulantK18Comp.h | 17 ----- .../CumulantK18/CumulantK18Comp_Device.cuh | 20 ------ .../CumulantK20/CumulantK20Comp.h | 16 ----- .../K08CompressibleNavierStokes.cu} | 16 ++--- .../K08CompressibleNavierStokes.h | 17 +++++ .../K08CompressibleNavierStokes_Device.cu} | 2 +- .../K08CompressibleNavierStokes_Device.cuh} | 7 +- .../K15CompressibleNavierStokes_Device.cuh | 21 +++--- ...erStokesSecondDerivatesFrom5thCumulants.cu | 39 +++++++++++ ...ierStokesSecondDerivatesFrom5thCumulants.h | 16 +++++ ...SecondDerivatesFrom5thCumulants_Device.cu} | 25 +++---- ...SecondDerivatesFrom5thCumulants_Device.cuh | 19 ++++++ .../K18CompressibleNavierStokes.cu} | 14 ++-- .../K18CompressibleNavierStokes.h | 17 +++++ .../K18CompressibleNavierStokes_Device.cu} | 2 +- .../K18CompressibleNavierStokes_Device.cuh | 21 ++++++ .../K20CompressibleNavierStokes.cu} | 14 ++-- .../K20CompressibleNavierStokes.h | 16 +++++ .../K20CompressibleNavierStokes_Device.cu} | 2 +- .../K20CompressibleNavierStokes_Device.cuh} | 6 +- .../M02CompressibleNavierStokes.cu} | 17 +++-- .../M02CompressibleNavierStokes.h | 17 +++++ .../M02CompressibleNavierStokes_Device.cu} | 2 +- .../M02CompressibleNavierStokes_Device.cuh} | 6 +- .../FluidFlow/Compressible/MRT/MRTCompSP27.h | 17 ----- .../B15IncompressibleNavierStokesBGKplus.cu} | 16 ++--- .../B15IncompressibleNavierStokesBGKplus.h | 17 +++++ ...compressibleNavierStokesBGKplus_Device.cu} | 2 +- ...compressibleNavierStokesBGKplus_Device.cuh | 17 +++++ .../B92IncompressibleNavierStokes.cu} | 16 ++--- .../B92IncompressibleNavierStokes.h | 17 +++++ .../B92IncompressibleNavierStokes_Device.cu} | 2 +- .../B92IncompressibleNavierStokes_Device.cuh} | 3 +- .../Incompressible/BGK/BGKIncompSP27.h | 17 ----- .../BGKPlus/BGKPlusIncompSP27.h | 17 ----- .../BGKPlus/BGKPlusIncompSP27_Device.cuh | 16 ----- .../C06IncompressibleNavierStokes.cu} | 14 ++-- .../C06IncompressibleNavierStokes.h | 17 +++++ .../CascadeIncompSP27_Device.cu | 3 +- .../CascadeIncompSP27_Device.cuh | 7 +- .../Cascade/CascadeIncompSP27.h | 17 ----- .../Cumulant1hSP27/Cumulant1hIncompSP27.cu | 45 ------------ .../Cumulant1hSP27/Cumulant1hIncompSP27.h | 17 ----- .../CumulantK15/CumulantK15Incomp.h | 16 ----- .../CumulantK15/CumulantK15Incomp_Device.cuh | 16 ----- .../K15IncompressibleNavierStokes.cu} | 16 ++--- .../K15IncompressibleNavierStokes.h | 16 +++++ .../K15IncompressibleNavierStokes_Device.cu} | 3 +- .../K15IncompressibleNavierStokes_Device.cuh | 17 +++++ ...ssibleNavierStokesRotatingVelocityField.cu | 45 ++++++++++++ ...essibleNavierStokesRotatingVelocityField.h | 18 +++++ ...vierStokesRotatingVelocityField_Device.cu} | 3 +- ...ierStokesRotatingVelocityField_Device.cuh} | 7 +- .../M02IncompressibleNavierStokes.cu} | 16 ++--- .../M02IncompressibleNavierStokes.h | 16 +++++ .../M02IncompressibleNavierStokes_Device.cu} | 3 +- .../M02IncompressibleNavierStokes_Device.cuh | 17 +++++ .../Incompressible/MRT/MRTIncompSP27.h | 16 ----- .../MRT/MRTIncompSP27_Device.cuh | 16 ----- .../KernelFactory/KernelFactoryImp.cpp | 68 +++++++++---------- 86 files changed, 701 insertions(+), 682 deletions(-) create mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B15CompressibleNavierStokesBGKplus/B15CompressibleNavierStokesBGKplus.cu create mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B15CompressibleNavierStokesBGKplus/B15CompressibleNavierStokesBGKplus.h rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/{BGKPlus/BGKPlusCompSP27_Device.cu => B15CompressibleNavierStokesBGKplus/B15CompressibleNavierStokesBGKplus_Device.cu} (99%) create mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B15CompressibleNavierStokesBGKplus/B15CompressibleNavierStokesBGKplus_Device.cuh rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/{BGKUnified/BGKUnified.cu => B15CompressibleNavierStokesBGKplusUnified/B15CompressibleNavierStokesBGKplusUnified.cu} (83%) create mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B15CompressibleNavierStokesBGKplusUnified/B15CompressibleNavierStokesBGKplusUnified.h rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/{BGK/BGKCompSP27.cu => B92CompressibleNavierStokes/B92CompressibleNavierStokes.cu} (57%) create mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B92CompressibleNavierStokes/B92CompressibleNavierStokes.h rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/{BGK/BGKCompSP27_Device.cu => B92CompressibleNavierStokes/B92CompressibleNavierStokes_Device.cu} (99%) create mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/B92CompressibleNavierStokes/B92CompressibleNavierStokes_Device.cuh delete mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGK/BGKCompSP27.h delete mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGK/BGKCompSP27_Device.cuh delete mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKPlus/BGKPlusCompSP27.cu delete mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKPlus/BGKPlusCompSP27.h delete mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKPlus/BGKPlusCompSP27_Device.cuh delete mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKUnified/BGKUnified.h rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/{Cascade/CascadeCompSP27.cu => C06CompressibleNavierStokes/C06CompressibleNavierStokes.cu} (54%) create mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/C06CompressibleNavierStokes/C06CompressibleNavierStokes.h rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/{Cascade/CascadeCompSP27_Device.cu => C06CompressibleNavierStokes/C06CompressibleNavierStokes_Device.cu} (99%) rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/{Cascade/CascadeCompSP27_Device.cuh => C06CompressibleNavierStokes/C06CompressibleNavierStokes_Device.cuh} (57%) delete mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cascade/CascadeCompSP27.h delete mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/Cumulant/CumulantCompSP27.h delete mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantAll4/CumulantAll4CompSP27.cu delete mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantAll4/CumulantAll4CompSP27.h delete mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantAll4/CumulantAll4CompSP27_Device.cuh delete mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK18/CumulantK18Comp.h delete mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK18/CumulantK18Comp_Device.cuh delete mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK20/CumulantK20Comp.h rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/{Cumulant/CumulantCompSP27.cu => K08CompressibleNavierStokes/K08CompressibleNavierStokes.cu} (54%) create mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K08CompressibleNavierStokes/K08CompressibleNavierStokes.h rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/{Cumulant/CumulantCompSP27_Device.cu => K08CompressibleNavierStokes/K08CompressibleNavierStokes_Device.cu} (99%) rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/{Cumulant/CumulantCompSP27_Device.cuh => K08CompressibleNavierStokes/K08CompressibleNavierStokes_Device.cuh} (57%) create mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants.cu create mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants.h rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/{CumulantAll4/CumulantAll4CompSP27_Device.cu => K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants_Device.cu} (99%) create mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants/K17CompressibleNavierStokesSecondDerivatesFrom5thCumulants_Device.cuh rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/{CumulantK18/CumulantK18Comp.cu => K18CompressibleNavierStokes/K18CompressibleNavierStokes.cu} (61%) create mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K18CompressibleNavierStokes/K18CompressibleNavierStokes.h rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/{CumulantK18/CumulantK18Comp_Device.cu => K18CompressibleNavierStokes/K18CompressibleNavierStokes_Device.cu} (99%) create mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K18CompressibleNavierStokes/K18CompressibleNavierStokes_Device.cuh rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/{CumulantK20/CumulantK20Comp.cu => K20CompressibleNavierStokes/K20CompressibleNavierStokes.cu} (61%) create mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/K20CompressibleNavierStokes/K20CompressibleNavierStokes.h rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/{CumulantK20/CumulantK20Comp_Device.cu => K20CompressibleNavierStokes/K20CompressibleNavierStokes_Device.cu} (99%) rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/{CumulantK20/CumulantK20Comp_Device.cuh => K20CompressibleNavierStokes/K20CompressibleNavierStokes_Device.cuh} (75%) rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/{MRT/MRTCompSP27.cu => M02CompressibleNavierStokes/M02CompressibleNavierStokes.cu} (57%) create mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/M02CompressibleNavierStokes/M02CompressibleNavierStokes.h rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/{MRT/MRTCompSP27_Device.cu => M02CompressibleNavierStokes/M02CompressibleNavierStokes_Device.cu} (99%) rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/{MRT/MRTCompSP27_Device.cuh => M02CompressibleNavierStokes/M02CompressibleNavierStokes_Device.cuh} (56%) delete mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/MRT/MRTCompSP27.h rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/{BGKPlus/BGKPlusIncompSP27.cu => B15IncompressibleNavierStokesBGKplus/B15IncompressibleNavierStokesBGKplus.cu} (50%) create mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/B15IncompressibleNavierStokesBGKplus/B15IncompressibleNavierStokesBGKplus.h rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/{BGKPlus/BGKPlusIncompSP27_Device.cu => B15IncompressibleNavierStokesBGKplus/B15IncompressibleNavierStokesBGKplus_Device.cu} (99%) create mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/B15IncompressibleNavierStokesBGKplus/B15IncompressibleNavierStokesBGKplus_Device.cuh rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/{BGK/BGKIncompSP27.cu => B92IncompressibleNavierStokes/B92IncompressibleNavierStokes.cu} (57%) create mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/B92IncompressibleNavierStokes/B92IncompressibleNavierStokes.h rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/{BGK/BGKIncompSP27_Device.cu => B92IncompressibleNavierStokes/B92IncompressibleNavierStokes_Device.cu} (99%) rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/{BGK/BGKIncompSP27_Device.cuh => B92IncompressibleNavierStokes/B92IncompressibleNavierStokes_Device.cuh} (80%) delete mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGK/BGKIncompSP27.h delete mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGKPlus/BGKPlusIncompSP27.h delete mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/BGKPlus/BGKPlusIncompSP27_Device.cuh rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/{Cascade/CascadeIncompSP27.cu => C06IncompressibleNavierStokes/C06IncompressibleNavierStokes.cu} (57%) create mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/C06IncompressibleNavierStokes/C06IncompressibleNavierStokes.h rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/{Cascade => C06IncompressibleNavierStokes}/CascadeIncompSP27_Device.cu (99%) rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/{Cascade => C06IncompressibleNavierStokes}/CascadeIncompSP27_Device.cuh (56%) delete mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cascade/CascadeIncompSP27.h delete mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cumulant1hSP27/Cumulant1hIncompSP27.cu delete mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/Cumulant1hSP27/Cumulant1hIncompSP27.h delete mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantK15/CumulantK15Incomp.h delete mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantK15/CumulantK15Incomp_Device.cuh rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/{CumulantK15/CumulantK15Incomp.cu => K15IncompressibleNavierStokes/K15IncompressibleNavierStokes.cu} (53%) create mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokes/K15IncompressibleNavierStokes.h rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/{CumulantK15/CumulantK15Incomp_Device.cu => K15IncompressibleNavierStokes/K15IncompressibleNavierStokes_Device.cu} (99%) create mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokes/K15IncompressibleNavierStokes_Device.cuh create mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokesRotatingVelocityField/K15IncompressibleNavierStokesRotatingVelocityField.cu create mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokesRotatingVelocityField/K15IncompressibleNavierStokesRotatingVelocityField.h rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/{Cumulant1hSP27/Cumulant1hIncompSP27_Device.cu => K15IncompressibleNavierStokesRotatingVelocityField/K15IncompressibleNavierStokesRotatingVelocityField_Device.cu} (99%) rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/{Cumulant1hSP27/Cumulant1hIncompSP27_Device.cuh => K15IncompressibleNavierStokesRotatingVelocityField/K15IncompressibleNavierStokesRotatingVelocityField_Device.cuh} (56%) rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/{MRT/MRTIncompSP27.cu => M02IncompressibleNavierStokes/M02IncompressibleNavierStokes.cu} (57%) create mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/M02IncompressibleNavierStokes/M02IncompressibleNavierStokes.h rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/{MRT/MRTIncompSP27_Device.cu => M02IncompressibleNavierStokes/M02IncompressibleNavierStokes_Device.cu} (99%) create mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/M02IncompressibleNavierStokes/M02IncompressibleNavierStokes_Device.cuh delete mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/MRT/MRTIncompSP27.h delete mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/MRT/MRTIncompSP27_Device.cuh diff --git a/src/gpu/VirtualFluids_GPU/GPU/GPU_Kernels.cuh b/src/gpu/VirtualFluids_GPU/GPU/GPU_Kernels.cuh index 877390c82..74ebf3bea 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 000000000..46d3c6edc --- /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 000000000..0347e30dc --- /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 638210bd2..b21213bf5 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 000000000..2a4775b95 --- /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 0a5ac6cf7..bafd4477d 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 000000000..ffc8ce41d --- /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 4aef26b7d..b708caf38 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 000000000..8b7ba23b3 --- /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 44add98d9..d032936a9 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 000000000..4e0450452 --- /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 4163d87fa..000000000 --- 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 59a524086..000000000 --- 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 00aaf3c27..000000000 --- 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 da40eb8ae..000000000 --- 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 9e991ffa4..000000000 --- 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 762eaaa59..000000000 --- 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 664b46fce..ca9434746 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 000000000..8004ad52e --- /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 6bd4415c7..6cf0fd639 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 8a49bd02a..fb8fc657a 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 900f41884..000000000 --- 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 0f9317f12..000000000 --- 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 c8aad41b8..000000000 --- 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 e36d51bc4..000000000 --- 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 5f23194d5..000000000 --- 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 6cffff126..000000000 --- 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 60a15145e..000000000 --- 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 24745df47..000000000 --- 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 218623b7c..d56b9f2ec 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 000000000..2f91a4598 --- /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 6ab3385b8..021ec25f3 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 cf6a92669..b2a6d02b6 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 77d42493e..0dfdb3488 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 000000000..1957ae332 --- /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 000000000..cfc308520 --- /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 7a5e39d6f..4919f5697 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 000000000..df7f6a096 --- /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 15d3509e7..801c7d905 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 000000000..d75920c78 --- /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 c585c19aa..169dedd36 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 000000000..601071a73 --- /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 8181cdb69..b47f8b633 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 000000000..fc6d4e72a --- /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 0a26eff29..d722baa51 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 17691f621..b86fc3842 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 6e11bd97a..bc8d4e11f 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 000000000..b18eb3dfb --- /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 41b24a349..ae0f67fdb 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 947ce6825..06542a20c 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 07aba55fd..000000000 --- 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 84a55b89d..3bc4fcbec 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 000000000..2426d32cb --- /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 1fee18161..49764bc07 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 000000000..1c812f936 --- /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 39bd1f349..730eb5407 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 000000000..782b76446 --- /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 6cfde7648..1d343eb30 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 f1a90b452..8b81ebf32 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 04308f691..000000000 --- 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 c6a1dcaca..000000000 --- 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 7f85f3ca2..000000000 --- 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 b060137f2..eebe0587b 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 000000000..292b37ee7 --- /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 0346f12cf..3f51dd404 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 a531fa7bd..daebc3345 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 4ebb0c50a..000000000 --- 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 2cade4307..000000000 --- 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 c6086649d..000000000 --- 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/CumulantK15/CumulantK15Incomp.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/CumulantK15/CumulantK15Incomp.h deleted file mode 100644 index c08737b43..000000000 --- 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 f2b5063f9..000000000 --- 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 c59792419..bfbf84b40 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 000000000..40d06dfae --- /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 9fcfeaa97..5b666eb73 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 000000000..9252b17df --- /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/K15IncompressibleNavierStokesRotatingVelocityField/K15IncompressibleNavierStokesRotatingVelocityField.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Incompressible/K15IncompressibleNavierStokesRotatingVelocityField/K15IncompressibleNavierStokesRotatingVelocityField.cu new file mode 100644 index 000000000..e405246db --- /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 000000000..948588332 --- /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 ac88396b4..082a5dc38 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 0679e770f..b313fcb8f 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 daa90091f..71660b82e 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 000000000..dde584611 --- /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 2fbc7d64d..4c94c960e 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 000000000..c62c73b10 --- /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 dd467be23..000000000 --- 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 d3a9fcea7..000000000 --- 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 694251a71..beb0cc930 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 -- GitLab