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/BGKUnified/BGKUnified.cu
index a99bb716b6e7cd08719aad3c781863db47367757..bd3e08c859163a0852498e148f7e1380ce7c3bdd 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKUnified/BGKUnified.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/BGKUnified/BGKUnified.cu
@@ -1,9 +1,10 @@
#include "BGKUnified.h"
+#include <stdexcept>
+
#include "Parameter/Parameter.h"
-#include "../CumulantKernel.cuh"
+#include "../RunLBMKernel.cuh"
#include "Kernel/Utilities/CudaGrid.h"
-#include <stdexcept>
#include <lbm/BGK.h>
@@ -15,7 +16,7 @@ std::shared_ptr<BGKUnified> BGKUnified::getNewInstance(std::shared_ptr<Parameter
void BGKUnified::run()
{
- vf::gpu::LBMKernelParameter kernelParameter{ para->getParD(level)->omega,
+ vf::gpu::GPUKernelParameter kernelParameter{ para->getParD(level)->omega,
para->getParD(level)->geoSP,
para->getParD(level)->neighborX_SP,
para->getParD(level)->neighborY_SP,
@@ -25,11 +26,11 @@ void BGKUnified::run()
nullptr, /* forces not used in bgk kernel */
para->getParD(level)->evenOrOdd };
- auto lambda = [] __device__(vf::lbm::CumulantChimeraParameter parameter) {
+ auto lambda = [] __device__(vf::lbm::KernelParameter parameter) {
return vf::lbm::bgk(parameter);
};
- vf::gpu::cumulantKernel<<<cudaGrid.grid, cudaGrid.threads>>>(lambda, kernelParameter);
+ vf::gpu::runKernel<<<cudaGrid.grid, cudaGrid.threads>>>(lambda, kernelParameter);
getLastCudaError("LB_Kernel_BGKUnified execution failed");
}
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK15Unified/CumulantK15Unified.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK15Unified/CumulantK15Unified.cu
index 7aecc31a6ed1e40f47abd72e8d02a93b261145c3..90f6dad8eea81d52957dc854f9cfd0a3282556f1 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK15Unified/CumulantK15Unified.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK15Unified/CumulantK15Unified.cu
@@ -1,10 +1,10 @@
#include "CumulantK15Unified.h"
-#include "../CumulantKernel.cuh"
+#include <stdexcept>
-#include "Parameter/Parameter.h"
+#include "../RunLBMKernel.cuh"
-#include <lbm/Distribution27.h>
+#include "Parameter/Parameter.h"
#include <lbm/CumulantChimera.h>
@@ -15,7 +15,7 @@ std::shared_ptr<CumulantK15Unified> CumulantK15Unified::getNewInstance(std::shar
void CumulantK15Unified::run()
{
- vf::gpu::LBMKernelParameter kernelParameter{ para->getParD(level)->omega,
+ vf::gpu::GPUKernelParameter kernelParameter{ para->getParD(level)->omega,
para->getParD(level)->geoSP,
para->getParD(level)->neighborX_SP,
para->getParD(level)->neighborY_SP,
@@ -25,11 +25,11 @@ void CumulantK15Unified::run()
para->getParD(level)->forcing,
para->getParD(level)->evenOrOdd };
- auto lambda = [] __device__(vf::lbm::CumulantChimeraParameter parameter) {
+ auto lambda = [] __device__(vf::lbm::KernelParameter parameter) {
return vf::lbm::cumulantChimera(parameter, vf::lbm::setRelaxationRatesK15);
};
- vf::gpu::cumulantKernel<<<cudaGrid.grid, cudaGrid.threads>>>(lambda, kernelParameter);
+ vf::gpu::runKernel<<<cudaGrid.grid, cudaGrid.threads>>>(lambda, kernelParameter);
getLastCudaError("LB_Kernel_CumulantK15Comp execution failed");
}
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Unified/CumulantK17Unified.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Unified/CumulantK17Unified.cu
index 0ebb95702e63e4caa499c00cae33470c86d0ab93..40398fe9cbd4856deb61c33d5ca20d5efddbe3b4 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Unified/CumulantK17Unified.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Unified/CumulantK17Unified.cu
@@ -1,9 +1,10 @@
#include "CumulantK17Unified.h"
+#include <stdexcept>
+
#include "Parameter/Parameter.h"
-#include "../CumulantKernel.cuh"
+#include "../RunLBMKernel.cuh"
#include "Kernel/Utilities/CudaGrid.h"
-#include <stdexcept>
#include <lbm/CumulantChimera.h>
@@ -15,7 +16,7 @@ std::shared_ptr<CumulantK17Unified> CumulantK17Unified::getNewInstance(std::shar
void CumulantK17Unified::run()
{
- vf::gpu::LBMKernelParameter kernelParameter{ para->getParD(level)->omega,
+ vf::gpu::GPUKernelParameter kernelParameter{ para->getParD(level)->omega,
para->getParD(level)->geoSP,
para->getParD(level)->neighborX_SP,
para->getParD(level)->neighborY_SP,
@@ -25,11 +26,11 @@ void CumulantK17Unified::run()
para->getParD(level)->forcing,
para->getParD(level)->evenOrOdd };
- auto lambda = [] __device__(vf::lbm::CumulantChimeraParameter parameter) {
+ auto lambda = [] __device__(vf::lbm::KernelParameter parameter) {
return vf::lbm::cumulantChimera(parameter, vf::lbm::setRelaxationRatesK17);
};
- vf::gpu::cumulantKernel<<<cudaGrid.grid, cudaGrid.threads>>>(lambda, kernelParameter);
+ vf::gpu::runKernel<<<cudaGrid.grid, cudaGrid.threads>>>(lambda, kernelParameter);
getLastCudaError("LB_Kernel_CumulantK17Unified execution failed");
}
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantKernel.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantKernel.cu
deleted file mode 100644
index 945ab247400e5733e2f662d684f483af6e90e981..0000000000000000000000000000000000000000
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantKernel.cu
+++ /dev/null
@@ -1,14 +0,0 @@
-#include "CumulantKernel.cuh"
-
-
-#include "Kernel/Utilities/DistributionHelper.cuh"
-
-
-namespace vf
-{
-namespace gpu
-{
-
-
-}
-}
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantKernel.cuh b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/RunLBMKernel.cuh
similarity index 76%
rename from src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantKernel.cuh
rename to src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/RunLBMKernel.cuh
index 9c595a2ec17ef34a8e4d75863266854935f20716..65305549e6ed3bf41432e70175caef5165aacf8e 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantKernel.cuh
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/RunLBMKernel.cuh
@@ -6,7 +6,7 @@
#include <cuda_runtime.h>
#include <lbm/Distribution27.h>
-#include <lbm/CumulantChimeraParameter.h>
+#include <lbm/KernelParameter.h>
#include "Kernel/Utilities/DistributionHelper.cuh"
@@ -16,7 +16,7 @@ namespace gpu
{
-struct LBMKernelParameter
+struct GPUKernelParameter
{
real omega;
unsigned int* typeOfGridNode;
@@ -30,7 +30,7 @@ struct LBMKernelParameter
};
template<typename KernelFunctor>
-__global__ void cumulantKernel(KernelFunctor kernel, LBMKernelParameter kernelParameter)
+__global__ void runKernel(KernelFunctor kernel, GPUKernelParameter kernelParameter)
{
const uint k = vf::gpu::getNodeIndex();
const uint nodeType = kernelParameter.typeOfGridNode[k];
@@ -48,8 +48,8 @@ __global__ void cumulantKernel(KernelFunctor kernel, LBMKernelParameter kernelPa
kernelParameter.neighborZ
};
- lbm::CumulantChimeraParameter chimeraParameter {distributionWrapper.distribution, kernelParameter.omega, kernelParameter.forces};
- kernel(chimeraParameter);
+ lbm::KernelParameter parameter {distributionWrapper.distribution, kernelParameter.omega, kernelParameter.forces};
+ kernel(parameter);
distributionWrapper.write();
}
diff --git a/src/lbm/BGK.cpp b/src/lbm/BGK.cpp
index 8680436e5a042d713df78eeabcbe8d35d61bcefd..407807cf52dae2526e3782c72565852ce64a34d5 100644
--- a/src/lbm/BGK.cpp
+++ b/src/lbm/BGK.cpp
@@ -20,7 +20,7 @@ using namespace constant;
-__host__ __device__ void bgk(CumulantChimeraParameter parameter)
+__host__ __device__ void bgk(KernelParameter parameter)
{
auto& distribution = parameter.distribution;
const auto omega = parameter.omega;
diff --git a/src/lbm/BGK.h b/src/lbm/BGK.h
index a2320eca357d0b3aee3bb2236536290c2a5a954b..2c82f5bd445ee008954add02fd0d6d6093364e90 100644
--- a/src/lbm/BGK.h
+++ b/src/lbm/BGK.h
@@ -10,14 +10,14 @@
#include <basics/Core/DataTypes.h>
-#include "CumulantChimeraParameter.h"
+#include "KernelParameter.h"
namespace vf
{
namespace lbm
{
-__host__ __device__ void bgk(CumulantChimeraParameter parameter);
+__host__ __device__ void bgk(KernelParameter parameter);
}
}
diff --git a/src/lbm/CumulantChimera.cpp b/src/lbm/CumulantChimera.cpp
index 87b89c94f6a767ca91dc3ce952cb332738ad676a..77133d1b698767b9569c91ac1b7d68a68dfeb58b 100644
--- a/src/lbm/CumulantChimera.cpp
+++ b/src/lbm/CumulantChimera.cpp
@@ -71,7 +71,7 @@ __host__ __device__ void setRelaxationRatesK15(real omega, real &OxxPyyPzz, real
//! and \ref
//! <a href="https://doi.org/10.1016/j.jcp.2017.07.004"><b>[ M. Geier et al. (2017), DOI:10.1016/j.jcp.2017.07.004 ]</b></a>
//////////////////////////////////////////////////////////////////////////
-__host__ __device__ void cumulantChimera(CumulantChimeraParameter parameter, RelaxationRatesFunctor setRelaxationRates)
+__host__ __device__ void cumulantChimera(KernelParameter parameter, RelaxationRatesFunctor setRelaxationRates)
{
auto& distribution = parameter.distribution;
const auto omega = parameter.omega;
diff --git a/src/lbm/CumulantChimera.h b/src/lbm/CumulantChimera.h
index 35e96d7e75026a1a23b20f6ed1c2c76f533f82d4..e8740c7d3f5b988a6fdc5c3b16ab6a90e0a28b83 100644
--- a/src/lbm/CumulantChimera.h
+++ b/src/lbm/CumulantChimera.h
@@ -10,7 +10,7 @@
#include <basics/Core/DataTypes.h>
-#include "CumulantChimeraParameter.h"
+#include "KernelParameter.h"
namespace vf
{
@@ -27,7 +27,7 @@ using RelaxationRatesFunctor = void(*)(real omega, real &OxxPyyPzz, real &OxyyPx
real &O4, real &O5, real &O6);
-__host__ __device__ void cumulantChimera(CumulantChimeraParameter parameter, RelaxationRatesFunctor setRelaxationRates);
+__host__ __device__ void cumulantChimera(KernelParameter parameter, RelaxationRatesFunctor setRelaxationRates);
}
}
diff --git a/src/lbm/CumulantChimeraParameter.h b/src/lbm/KernelParameter.h
similarity index 91%
rename from src/lbm/CumulantChimeraParameter.h
rename to src/lbm/KernelParameter.h
index 6cb820c57ce4a857eaec9c49560936fb80ec3b24..b60e43439dda3f2c592b67b6f9491d010891a32c 100644
--- a/src/lbm/CumulantChimeraParameter.h
+++ b/src/lbm/KernelParameter.h
@@ -18,7 +18,7 @@ namespace lbm
{
-struct CumulantChimeraParameter
+struct KernelParameter
{
Distribution27& distribution;
real omega;