From 27ad2040e537ea8e724a4e98a8e9a090ccd30441 Mon Sep 17 00:00:00 2001
From: "TESLA03\\Master" <a.wellmann@tu-bs.de>
Date: Wed, 21 Jul 2021 14:30:34 +0200
Subject: [PATCH] Run new kernel only on fluid nodes
---
.../CumulantK17CompChimSparse.cu | 6 ++++--
.../CumulantK17CompChimSparse_Device.cu | 14 +++++++++-----
.../CumulantK17CompChimSparse_Device.cuh | 4 +++-
3 files changed, 16 insertions(+), 8 deletions(-)
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 0c197e53e..26b85c03a 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 0acac9085..a520ebe25 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 719264237..e750a08c2 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
--
GitLab