Fix memory leak in gpu kernel modifications

src/gpu/VirtualFluids_GPU/GPU/CalcMac27.cu

@@ -42,14 +42,18 @@ extern "C" __global__ void LBCalcMac27( real* vxD,
 
    const unsigned int k = nx*(ny*z + y) + x; // Zugriff auf arrays im device
 
+
+   if(k >= size_Mat)
+      return;
+
+   if(!vf::gpu::isValidFluidNode(geoD[k]))
+      return;
+
    rhoD[k] = c0o1;
    vxD[k]  = c0o1;
    vyD[k]  = c0o1;
    vzD[k]  = c0o1;
 
-   if(!vf::gpu::isValidFluidNode(k, size_Mat, geoD[k]))
-      return;
-
    vf::gpu::DistributionWrapper distr_wrapper(distributions, size_Mat, isEvenTimestep, k, neighborX, neighborY, neighborZ);
    const auto& distribution = distr_wrapper.distribution;
 
@@ -266,15 +270,18 @@ extern "C" __global__ void LBCalcMacCompSP27(real *vxD, real *vyD, real *vzD, re
 {
     const unsigned k = vf::gpu::getNodeIndex();
 
+    if(k >= size_Mat)
+        return;
+
+    if (!vf::gpu::isValidFluidNode(geoD[k]))
+        return;
+
     pressD[k] = c0o1;
     rhoD[k]   = c0o1;
     vxD[k]    = c0o1;
     vyD[k]    = c0o1;
     vzD[k]    = c0o1;
 
-    if (!vf::gpu::isValidFluidNode(k, size_Mat, geoD[k]))
-        return;
-
     vf::gpu::DistributionWrapper distr_wrapper(distributions, size_Mat, isEvenTimestep, k, neighborX, neighborY,
                                                neighborZ);
     const auto &distribution = distr_wrapper.distribution;

src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/RunLBMKernel.cuh

@@ -32,9 +32,11 @@ template<typename KernelFunctor>
 __global__ void runKernel(KernelFunctor kernel, GPUKernelParameter kernelParameter)
 {
     const uint k = getNodeIndex();
-    const uint nodeType = kernelParameter.typeOfGridNode[k];
 
-    if (!isValidFluidNode(k, kernelParameter.size_Mat, nodeType))
+    if(k >= kernelParameter.size_Mat)
+        return;
+
+    if (!isValidFluidNode(kernelParameter.typeOfGridNode[k]))
         return;
 
     DistributionWrapper distributionWrapper {

src/gpu/VirtualFluids_GPU/Kernel/Utilities/DistributionHelper.cu

@@ -158,10 +158,9 @@ __device__ unsigned int getNodeIndex()
     return nx * (ny * z + y) + x;
 }
 
-__device__ bool isValidFluidNode(uint k, int size_Mat, uint nodeType)
+__device__ bool isValidFluidNode(uint nodeType)
 {
-    return (k < size_Mat) &&
-           (nodeType == GEO_FLUID || nodeType == GEO_PM_0 || nodeType == GEO_PM_1 || nodeType == GEO_PM_2);
+    return (nodeType == GEO_FLUID || nodeType == GEO_PM_0 || nodeType == GEO_PM_1 || nodeType == GEO_PM_2);
 }
 
 

src/gpu/VirtualFluids_GPU/Kernel/Utilities/DistributionHelper.cuh

@@ -90,7 +90,7 @@ struct DistributionWrapper
 
 __device__ unsigned int getNodeIndex();
 
-__device__ bool isValidFluidNode(uint k, int size_Mat, uint nodeType);
+__device__ bool isValidFluidNode(uint nodeType);
 
 }
 }