Move grid calculation to KernelImp

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);
+}

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

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,

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,

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);
-}

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