From 6d48710148fb3b65798aed25aa475328209c32f3 Mon Sep 17 00:00:00 2001
From: peters <peters@irmb.tu-bs.de>
Date: Thu, 10 Jun 2021 11:48:54 +0200
Subject: [PATCH] More generic renaming
---
.../Compressible/BGKUnified/BGKUnified.cu | 11 ++++++-----
.../CumulantK15Unified/CumulantK15Unified.cu | 12 ++++++------
.../CumulantK17Unified/CumulantK17Unified.cu | 11 ++++++-----
.../FluidFlow/Compressible/CumulantKernel.cu | 14 --------------
.../{CumulantKernel.cuh => RunLBMKernel.cuh} | 10 +++++-----
src/lbm/BGK.cpp | 2 +-
src/lbm/BGK.h | 4 ++--
src/lbm/CumulantChimera.cpp | 2 +-
src/lbm/CumulantChimera.h | 4 ++--
...umulantChimeraParameter.h => KernelParameter.h} | 2 +-
10 files changed, 30 insertions(+), 42 deletions(-)
delete mode 100644 src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantKernel.cu
rename src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/{CumulantKernel.cuh => RunLBMKernel.cuh} (76%)
rename src/lbm/{CumulantChimeraParameter.h => KernelParameter.h} (91%)
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 a99bb716b..bd3e08c85 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 7aecc31a6..90f6dad8e 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 0ebb95702..40398fe9c 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 945ab2474..000000000
--- 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 9c595a2ec..65305549e 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 8680436e5..407807cf5 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 a2320eca3..2c82f5bd4 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 87b89c94f..77133d1b6 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 35e96d7e7..e8740c7d3 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 6cb820c57..b60e43439 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;
--
GitLab