From ac95601655e679217fc176a9b32f284c646d9b6f Mon Sep 17 00:00:00 2001
From: Anna Wellmann <a.wellmann@tu-braunschweig.de>
Date: Wed, 1 Sep 2021 16:07:15 +0200
Subject: [PATCH] Move grid calculation to KernelImp
---
.../VirtualFluids_GPU/Kernel/KernelImp.cpp | 21 +++++++++++++++-
src/gpu/VirtualFluids_GPU/Kernel/KernelImp.h | 2 ++
.../CumulantK17/CumulantK17Comp.cu | 20 +++-------------
.../CumulantK17chim/CumulantK17CompChim.cu | 20 +++-------------
.../CumulantK17CompChimSparse.cu | 24 ++++---------------
.../CumulantK17CompChimSparse.h | 1 -
6 files changed, 32 insertions(+), 56 deletions(-)
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.cpp b/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.cpp
index 3151e6bed..e6960cc18 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.cpp
+++ b/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.cpp
@@ -27,6 +27,25 @@ void KernelImp::setCheckParameterStrategy(std::shared_ptr<CheckParameterStrategy
this->checkStrategy = strategy;
}
+
+
KernelImp::KernelImp(std::shared_ptr<Parameter> para, int level) : para(para), level(level) {}
-KernelImp::KernelImp() {}
\ No newline at end of file
+KernelImp::KernelImp() {}
+
+std::unique_ptr<std::pair<dim3, dim3>> KernelImp::calcGridDimensions(unsigned int size_Mat, int numberOfThreads)
+{
+ int Grid = (size_Mat / numberOfThreads) + 1;
+ int Grid1, Grid2;
+ if (Grid > 512) {
+ Grid1 = 512;
+ Grid2 = (Grid / Grid1) + 1;
+ } else {
+ Grid1 = 1;
+ Grid2 = Grid;
+ }
+ dim3 grid(Grid1, Grid2);
+ dim3 threads(numberOfThreads, 1, 1);
+ std::pair<dim3, dim3> dimensions(grid, threads);
+ return std::make_unique<std::pair<dim3, dim3>>(dimensions);
+}
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.h b/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.h
index c5215dbdd..ca41116e4 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.h
+++ b/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.h
@@ -33,6 +33,8 @@ protected:
KernelGroup myKernelGroup;
vf::gpu::CudaGrid cudaGrid;
+
+ std::unique_ptr<std::pair<dim3, dim3>> calcGridDimensions(unsigned int size_Mat, int numberOfThreads);
};
#endif
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17/CumulantK17Comp.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17/CumulantK17Comp.cu
index 86704c614..1caf0340c 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17/CumulantK17Comp.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17/CumulantK17Comp.cu
@@ -10,23 +10,9 @@ std::shared_ptr<CumulantK17Comp> CumulantK17Comp::getNewInstance(std::shared_ptr
void CumulantK17Comp::run()
{
- int numberOfThreads = para->getParD(level)->numberofthreads;
- int size_Mat = para->getParD(level)->size_Mat_SP;
-
- int Grid = (size_Mat / numberOfThreads) + 1;
- int Grid1, Grid2;
- if (Grid>512)
- {
- Grid1 = 512;
- Grid2 = (Grid / Grid1) + 1;
- }
- else
- {
- Grid1 = 1;
- Grid2 = Grid;
- }
- dim3 grid(Grid1, Grid2);
- dim3 threads(numberOfThreads, 1, 1);
+ dim3 grid, threads;
+ std::tie(grid, threads) =
+ *calcGridDimensions(para->getParD(level)->size_Mat_SP, para->getParD(level)->numberofthreads);
LB_Kernel_CumulantK17Comp <<< grid, threads >>>(para->getParD(level)->omega,
para->getParD(level)->geoSP,
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chim/CumulantK17CompChim.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chim/CumulantK17CompChim.cu
index 38df36655..a31505f33 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chim/CumulantK17CompChim.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chim/CumulantK17CompChim.cu
@@ -10,23 +10,9 @@ std::shared_ptr<CumulantK17CompChim> CumulantK17CompChim::getNewInstance(std::sh
void CumulantK17CompChim::run()
{
- int numberOfThreads = para->getParD(level)->numberofthreads;
- int size_Mat = para->getParD(level)->size_Mat_SP;
-
- int Grid = (size_Mat / numberOfThreads) + 1;
- int Grid1, Grid2;
- if (Grid>512)
- {
- Grid1 = 512;
- Grid2 = (Grid / Grid1) + 1;
- }
- else
- {
- Grid1 = 1;
- Grid2 = Grid;
- }
- dim3 grid(Grid1, Grid2);
- dim3 threads(numberOfThreads, 1, 1);
+ dim3 grid, threads;
+ std::tie(grid, threads) =
+ *calcGridDimensions(para->getParD(level)->size_Mat_SP, para->getParD(level)->numberofthreads);
LB_Kernel_CumulantK17CompChim <<< grid, threads >>>(
para->getParD(level)->omega,
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chimSparse/CumulantK17CompChimSparse.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chimSparse/CumulantK17CompChimSparse.cu
index 3f493de19..b2ba5fdb1 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chimSparse/CumulantK17CompChimSparse.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chimSparse/CumulantK17CompChimSparse.cu
@@ -14,7 +14,8 @@ std::shared_ptr<CumulantK17CompChimSparse> CumulantK17CompChimSparse::getNewInst
void CumulantK17CompChimSparse::run()
{
dim3 grid, threads;
- std::tie(grid, threads) = *calcGridDimensions(para->getParD(level)->numberOfFluidNodes);
+ std::tie(grid, threads) =
+ *calcGridDimensions(para->getParD(level)->numberOfFluidNodes, para->getParD(level)->numberofthreads);
LB_Kernel_CumulantK17CompChimSparse <<< grid, threads >>>(
para->getParD(level)->omega,
@@ -35,7 +36,8 @@ void CumulantK17CompChimSparse::run()
void CumulantK17CompChimSparse::runOnIndices(const unsigned int *indices, unsigned int size_indices, int streamIndex)
{
dim3 grid, threads;
- std::tie(grid, threads) = *calcGridDimensions(para->getParD(level)->numberOfFluidNodes);
+ std::tie(grid, threads) =
+ *calcGridDimensions(para->getParD(level)->numberOfFluidNodes, para->getParD(level)->numberofthreads);
cudaStream_t stream = (streamIndex == -1) ? CU_STREAM_LEGACY : para->getStreamManager().getStream(streamIndex);
@@ -66,21 +68,3 @@ CumulantK17CompChimSparse::CumulantK17CompChimSparse(std::shared_ptr<Parameter>
myKernelGroup = BasicKernel;
}
-std::unique_ptr<std::pair<dim3, dim3>> CumulantK17CompChimSparse::calcGridDimensions(unsigned int size_Mat)
-{
- int numberOfThreads = para->getParD(level)->numberofthreads;
-
- int Grid = (size_Mat / numberOfThreads) + 1;
- int Grid1, Grid2;
- if (Grid > 512) {
- Grid1 = 512;
- Grid2 = (Grid / Grid1) + 1;
- } else {
- Grid1 = 1;
- Grid2 = Grid;
- }
- dim3 grid(Grid1, Grid2);
- dim3 threads(numberOfThreads, 1, 1);
- std::pair<dim3, dim3> dimensions(grid, threads);
- return std::make_unique<std::pair<dim3, dim3>>(dimensions);
-}
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chimSparse/CumulantK17CompChimSparse.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chimSparse/CumulantK17CompChimSparse.h
index 906c88033..8233e9b9b 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chimSparse/CumulantK17CompChimSparse.h
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chimSparse/CumulantK17CompChimSparse.h
@@ -13,7 +13,6 @@ public:
private:
CumulantK17CompChimSparse();
CumulantK17CompChimSparse(std::shared_ptr<Parameter> para, int level);
- std::unique_ptr<std::pair<dim3, dim3>> calcGridDimensions(unsigned int size_Mat);
};
#endif
--
GitLab