diff --git a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp
index 442dcb6f962e4d0240f0cd93f49b332191d5a317..9cb31dac883590a5e9782c55f2e7df1948e7b645 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 c66d6afd40e4261ce0a6800c6239071c81c95179..0ce9383b9dc30e9dcdca4586f6d06c2593dbf027 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 4b54277ddd405eb619191895065af9bb3b780063..ceb1638adbfb3abf38dae61714a69356b0bb8361 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 5e4c5aa08e37e88008da13466bfeed6893ec94f6..cd06a7c94b2ffce4dcbb20d4f564fa82a56321f2 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 08b71d42e79564d9eac887289a1ae36824095c46..0986e7b9759a69cf6df6c9624a3e67fcb8716100 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 87b4714f2a815c369de759b33ecbd69bc49742ef..6c3555c8f74762de4f9d57053fe0ec10c959893d 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 ad53cd75bdf3282e1b63102620d19bf0c88f718c..f1645cfaec3ca28b0856bffc0ecb45319746216e 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