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 0c197e53eedf68bac078c972048cebba4540df4e..26b85c03ab86c9a8b3431c14986083996ee3f931 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 @@ -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"); } diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chimSparse/CumulantK17CompChimSparse_Device.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chimSparse/CumulantK17CompChimSparse_Device.cu index 0acac9085ae60a7320dc4038010b3548b9521a14..a520ebe25a32bfbc2595edd45d67a123ce1ade73 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chimSparse/CumulantK17CompChimSparse_Device.cu +++ b/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]; diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chimSparse/CumulantK17CompChimSparse_Device.cuh b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chimSparse/CumulantK17CompChimSparse_Device.cuh index 71926423726f9ffa6b9fd3bf54d000be89b0d244..e750a08c28c8b03440fbaf1ae8e4ba4e66d7761a 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chimSparse/CumulantK17CompChimSparse_Device.cuh +++ b/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