Add method to pass the indices to the kernel

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();
 
     //////////////////////////////////////////////////////////////////////////
 

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

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;

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

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

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

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