From 1007a1167628e212a5a306371e38032c9a93eca9 Mon Sep 17 00:00:00 2001
From: Anna Wellmann <a.wellmann@tu-braunschweig.de>
Date: Thu, 5 Aug 2021 11:04:08 +0200
Subject: [PATCH] Add method to pass the indices to the kernel
---
.../Calculation/UpdateGrid27.cpp | 16 ++++-
.../Calculation/UpdateGrid27.h | 2 +-
src/gpu/VirtualFluids_GPU/Kernel/Kernel.h | 1 +
.../VirtualFluids_GPU/Kernel/KernelImp.cpp | 9 ++-
src/gpu/VirtualFluids_GPU/Kernel/KernelImp.h | 1 +
.../CumulantK17CompChimSparse.cu | 61 +++++++++++++------
.../CumulantK17CompChimSparse.h | 2 +
7 files changed, 68 insertions(+), 24 deletions(-)
diff --git a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp
index 442dcb6f9..9cb31dac8 100644
--- a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp
+++ b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp
@@ -26,7 +26,11 @@ void updateGrid27(Parameter* para,
//////////////////////////////////////////////////////////////////////////
- collision(para, pm, level, t, kernels);
+ if (para->useStreams)
+ collision(para, pm, level, t, kernels, para->getParD(level)->fluidNodeIndices,
+ para->getParD(level)->numberOfFluidNodes);
+ else
+ collision(para, pm, level, t, kernels);
//////////////////////////////////////////////////////////////////////////
@@ -61,9 +65,15 @@ void updateGrid27(Parameter* para,
}
}
-void collision(Parameter* para, std::vector<std::shared_ptr<PorousMedia>>& pm, int level, unsigned int t, std::vector < SPtr< Kernel>>& kernels)
+void collision(Parameter* para, std::vector<std::shared_ptr<PorousMedia>>& pm, int level, unsigned int t, std::vector < SPtr< Kernel>>& kernels, uint* fluidNodeIndices, uint numberOfFluidNodes)
{
- kernels.at(level)->run();
+ if (para->useStreams)
+ if (fluidNodeIndices != nullptr && numberOfFluidNodes != 0)
+ kernels.at(level)->runOnIndices(fluidNodeIndices, numberOfFluidNodes);
+ else
+ std::cout << "in collision: fluidNodeIndices or numberOfFluidNodes not definded" << std::endl; // better use logger
+ else
+ kernels.at(level)->run();
//////////////////////////////////////////////////////////////////////////
diff --git a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h
index c66d6afd4..0ce9383b9 100644
--- a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h
+++ b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h
@@ -18,7 +18,7 @@ extern "C" void updateGrid27(Parameter* para,
unsigned int t,
std::vector < SPtr< Kernel>>& kernels);
-extern "C" void collision(Parameter* para, std::vector<std::shared_ptr<PorousMedia>>& pm, int level, unsigned int t, std::vector < SPtr< Kernel>>& kernels);
+extern "C" void collision(Parameter *para, std::vector<std::shared_ptr<PorousMedia>> &pm, int level, unsigned int t, std::vector<SPtr<Kernel>> &kernels, uint *fluidNodeIndices = nullptr, uint numberOfFluidNodes = 0);
extern "C" void collisionPorousMedia(Parameter* para, std::vector<std::shared_ptr<PorousMedia>>& pm, int level);
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernel.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernel.h
index 4b54277dd..ceb1638ad 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernel.h
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernel.h
@@ -13,6 +13,7 @@ class Kernel
public:
virtual ~Kernel() = default;
virtual void run() = 0;
+ virtual void runOnIndices(const unsigned int *indices, unsigned int size_indices) = 0;
virtual bool checkParameter() = 0;
virtual std::vector<PreProcessorType> getPreProcessorTypes() = 0;
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.cpp b/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.cpp
index 5e4c5aa08..cd06a7c94 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.cpp
+++ b/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.cpp
@@ -2,8 +2,13 @@
#include "Kernel/Utilities/CheckParameterStrategy/CheckParameterStrategy.h"
-bool KernelImp::checkParameter()
-{
+
+void KernelImp::runOnIndices(const unsigned int *indices, unsigned int size_indices)
+{
+ printf("Method not implemented for this Kernel \n");
+}
+
+bool KernelImp::checkParameter() {
return checkStrategy->checkParameter(para);
}
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.h b/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.h
index 08b71d42e..0986e7b97 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.h
+++ b/src/gpu/VirtualFluids_GPU/Kernel/KernelImp.h
@@ -14,6 +14,7 @@ class KernelImp : public Kernel
{
public:
virtual void run() = 0;
+ virtual void runOnIndices(const unsigned int *indices, unsigned int size_indices);
bool checkParameter();
std::vector<PreProcessorType> getPreProcessorTypes();
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 87b4714f2..6c3555c8f 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
@@ -11,23 +11,8 @@ std::shared_ptr<CumulantK17CompChimSparse> CumulantK17CompChimSparse::getNewInst
void CumulantK17CompChimSparse::run()
{
- int numberOfThreads = para->getParD(level)->numberofthreads;
- int size_Mat = para->getParD(level)->numberOfFluidNodes;
-
- 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)->numberOfFluidNodes);
LB_Kernel_CumulantK17CompChimSparse <<< grid, threads >>>(
para->getParD(level)->omega,
@@ -45,6 +30,27 @@ void CumulantK17CompChimSparse::run()
getLastCudaError("LB_Kernel_CumulantK17CompChim execution failed");
}
+void CumulantK17CompChimSparse::runOnIndices(const unsigned int *indices, unsigned int size_indices)
+{
+ dim3 grid, threads;
+ std::tie(grid, threads) = *calcGridDimensions(para->getParD(level)->numberOfFluidNodes);
+
+ LB_Kernel_CumulantK17CompChimSparse<<<grid, threads, 0, para->getStream(0)>>>(
+ para->getParD(level)->omega,
+ para->getParD(level)->neighborX_SP,
+ para->getParD(level)->neighborY_SP,
+ para->getParD(level)->neighborZ_SP,
+ para->getParD(level)->d0SP.f[0],
+ para->getParD(level)->size_Mat_SP,
+ level,
+ para->getForcesDev(),
+ para->getQuadricLimitersDev(),
+ para->getParD(level)->evenOrOdd,
+ indices,
+ size_indices);
+ getLastCudaError("LB_Kernel_CumulantK17CompChim execution failed");
+}
+
CumulantK17CompChimSparse::CumulantK17CompChimSparse(std::shared_ptr<Parameter> para, int level)
{
this->para = para;
@@ -53,4 +59,23 @@ CumulantK17CompChimSparse::CumulantK17CompChimSparse(std::shared_ptr<Parameter>
myPreProcessorTypes.push_back(InitCompSP27);
myKernelGroup = BasicKernel;
-}
\ No newline at end of file
+}
+
+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 ad53cd75b..f1645cfae 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
@@ -8,10 +8,12 @@ class CumulantK17CompChimSparse : public KernelImp
public:
static std::shared_ptr<CumulantK17CompChimSparse> getNewInstance(std::shared_ptr<Parameter> para, int level);
void run();
+ void runOnIndices(const unsigned int *indices, unsigned int size_indices) override;
private:
CumulantK17CompChimSparse();
CumulantK17CompChimSparse(std::shared_ptr<Parameter> para, int level);
+ std::unique_ptr<std::pair<dim3, dim3>> calcGridDimensions(unsigned int size_Mat);
};
#endif
--
GitLab