Run new kernel only on fluid nodes

src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chimSparse/CumulantK17CompChimSparse.cu

@@ -29,7 +29,7 @@ void CumulantK17CompChimSparse::run()
 	dim3 grid(Grid1, Grid2);
 	dim3 threads(numberOfThreads, 1, 1);
 
-	LB_Kernel_CumulantK17CompChimSparse<<<grid, threads>>>(
+	LB_Kernel_CumulantK17CompChimSparse <<< grid, threads >>>(
 		para->getParD(level)->omega,
 		para->getParD(level)->geoSP,
 		para->getParD(level)->neighborX_SP,
@@ -40,7 +40,9 @@ void CumulantK17CompChimSparse::run()
 		level,
 		para->getForcesDev(),
         para->getQuadricLimitersDev(),
-		para->getParD(level)->evenOrOdd);
+		para->getParD(level)->evenOrOdd,
+        para->getParD(level)->fluidNodeIndices,
+		para->getParD(level)->numberOfFluidNodes);
 	getLastCudaError("LB_Kernel_CumulantK17CompChim execution failed");
 }
 

src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chimSparse/CumulantK17CompChimSparse_Device.cu

@@ -50,7 +50,9 @@ extern "C" __global__ void LB_Kernel_CumulantK17CompChimSparse(
 	int level,
 	real* forces,
 	real* quadricLimiters,
-	bool isEvenTimestep)
+	bool isEvenTimestep,
+    const uint *fluidNodeIndices, 
+    uint numberOfFluidNodes)
 {
     //////////////////////////////////////////////////////////////////////////
     //! Cumulant K17 Kernel is based on \ref
@@ -70,17 +72,19 @@ extern "C" __global__ void LB_Kernel_CumulantK17CompChimSparse(
     const unsigned nx = blockDim.x;
     const unsigned ny = gridDim.x;
 
-    const unsigned k = nx * (ny * z + y) + x;
+    const unsigned k_thread = nx * (ny * z + y) + x;
 
     //////////////////////////////////////////////////////////////////////////
-    // run for all indices in size_Mat and fluid nodes
-    if ((k < size_Mat) && (typeOfGridNode[k] == GEO_FLUID)) {
+    // run for all indices in fluidNodeIndices
+    if (k_thread < numberOfFluidNodes) {
         //////////////////////////////////////////////////////////////////////////
         //! - Read distributions: style of reading and writing the distributions from/to stored arrays dependent on
         //! timestep is based on the esoteric twist algorithm \ref <a
         //! href="https://doi.org/10.3390/computation5020019"><b>[ M. Geier et al. (2017),
         //! DOI:10.3390/computation5020019 ]</b></a>
-        //!
+
+        const unsigned k = fluidNodeIndices[k_thread];
+
         Distributions27 dist;
         if (isEvenTimestep) {
             dist.f[dirE]    = &distributions[dirE * size_Mat];

src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chimSparse/CumulantK17CompChimSparse_Device.cuh

@@ -15,5 +15,7 @@ extern "C" __global__ void LB_Kernel_CumulantK17CompChimSparse(
 	int level,
 	real* forces,
 	real* quadricLimiters,
-	bool isEvenTimestep);
+	bool isEvenTimestep,
+	const uint* fluidNodeIndices,
+	uint numberOfFluidNodes);
 #endif