From c0c3cd42e320f4b98d0c5a6a68a62ba0e57b1e9b Mon Sep 17 00:00:00 2001
From: Anna Wellmann <a.wellmann@tu-bs.de>
Date: Wed, 21 Dec 2022 08:17:29 +0000
Subject: [PATCH] Add coordinate as parameter to sponge kernel
---
.../CumulantK17Sponge/CumulantK17Sponge.cu | 194 +++++++++---------
.../CumulantK17Sponge_Device.cu | 35 ++--
.../CumulantK17Sponge_Device.cuh | 3 +
3 files changed, 122 insertions(+), 110 deletions(-)
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Sponge/CumulantK17Sponge.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Sponge/CumulantK17Sponge.cu
index 7a313e783..28c7cafd4 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Sponge/CumulantK17Sponge.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Sponge/CumulantK17Sponge.cu
@@ -6,7 +6,7 @@
#include <cuda.h>
-template<TurbulenceModel turbulenceModel>
+template<TurbulenceModel turbulenceModel>
std::shared_ptr< CumulantK17Sponge<turbulenceModel> > CumulantK17Sponge<turbulenceModel>::getNewInstance(std::shared_ptr<Parameter> para, int level)
{
return std::shared_ptr<CumulantK17Sponge<turbulenceModel> >(new CumulantK17Sponge<turbulenceModel>(para,level));
@@ -15,23 +15,25 @@ std::shared_ptr< CumulantK17Sponge<turbulenceModel> > CumulantK17Sponge<turbulen
template<TurbulenceModel turbulenceModel>
void CumulantK17Sponge<turbulenceModel>::run()
{
- LB_Kernel_CumulantK17Sponge < turbulenceModel, false, false > <<< cudaGrid.grid, cudaGrid.threads >>>( para->getParD(level)->omega,
- para->getParD(level)->typeOfGridNode,
- para->getParD(level)->neighborX, para->getParD(level)->neighborY, para->getParD(level)->neighborZ,
- para->getParD(level)->distributions.f[0],
- para->getParD(level)->rho,
- para->getParD(level)->velocityX, para->getParD(level)->velocityY, para->getParD(level)->velocityZ,
- para->getParD(level)->turbViscosity,
- para->getSGSConstant(),
- (unsigned long)para->getParD(level)->numberOfNodes,
- level,
- para->getIsBodyForce(),
- para->getForcesDev(),
- para->getParD(level)->forceX_SP, para->getParD(level)->forceY_SP, para->getParD(level)->forceZ_SP,
- para->getQuadricLimitersDev(),
- para->getParD(level)->isEvenTimestep,
- para->getParD(level)->taggedFluidNodeIndices[CollisionTemplate::Default],
- para->getParD(level)->numberOfTaggedFluidNodes[CollisionTemplate::Default]);
+ LB_Kernel_CumulantK17Sponge < turbulenceModel, false, false > <<< cudaGrid.grid, cudaGrid.threads >>>(
+ para->getParD(level)->omega,
+ para->getParD(level)->typeOfGridNode,
+ para->getParD(level)->neighborX, para->getParD(level)->neighborY, para->getParD(level)->neighborZ,
+ para->getParD(level)->distributions.f[0],
+ para->getParD(level)->rho,
+ para->getParD(level)->velocityX, para->getParD(level)->velocityY, para->getParD(level)->velocityZ,
+ para->getParD(level)->turbViscosity,
+ para->getSGSConstant(),
+ (unsigned long)para->getParD(level)->numberOfNodes,
+ level,
+ para->getIsBodyForce(),
+ para->getForcesDev(),
+ para->getParD(level)->forceX_SP, para->getParD(level)->forceY_SP, para->getParD(level)->forceZ_SP,
+ para->getParD(level)->coordinateX, para->getParD(level)->coordinateY, para->getParD(level)->coordinateZ,
+ para->getQuadricLimitersDev(),
+ para->getParD(level)->isEvenTimestep,
+ para->getParD(level)->taggedFluidNodeIndices[CollisionTemplate::Default],
+ para->getParD(level)->numberOfTaggedFluidNodes[CollisionTemplate::Default]);
getLastCudaError("LB_Kernel_CumulantK17Sponge execution failed");
}
@@ -40,91 +42,95 @@ template<TurbulenceModel turbulenceModel>
void CumulantK17Sponge<turbulenceModel>::runOnIndices( const unsigned int *indices, unsigned int size_indices, CollisionTemplate collisionTemplate, CudaStreamIndex streamIndex )
{
cudaStream_t stream = para->getStreamManager()->getStream(streamIndex);
-
+
switch (collisionTemplate)
{
case CollisionTemplate::Default:
- LB_Kernel_CumulantK17Sponge < turbulenceModel, false, false > <<< cudaGrid.grid, cudaGrid.threads, 0, stream >>>(
- para->getParD(level)->omega,
- para->getParD(level)->typeOfGridNode,
- para->getParD(level)->neighborX, para->getParD(level)->neighborY, para->getParD(level)->neighborZ,
- para->getParD(level)->distributions.f[0],
- para->getParD(level)->rho,
- para->getParD(level)->velocityX, para->getParD(level)->velocityY, para->getParD(level)->velocityZ,
- para->getParD(level)->turbViscosity,
- para->getSGSConstant(),
- (unsigned long)para->getParD(level)->numberOfNodes,
- level,
- para->getIsBodyForce(),
- para->getForcesDev(),
- para->getParD(level)->forceX_SP, para->getParD(level)->forceY_SP, para->getParD(level)->forceZ_SP,
- para->getQuadricLimitersDev(),
- para->getParD(level)->isEvenTimestep,
- indices,
- size_indices);
+ LB_Kernel_CumulantK17Sponge < turbulenceModel, false, false > <<< cudaGrid.grid, cudaGrid.threads, 0, stream >>>(
+ para->getParD(level)->omega,
+ para->getParD(level)->typeOfGridNode,
+ para->getParD(level)->neighborX, para->getParD(level)->neighborY, para->getParD(level)->neighborZ,
+ para->getParD(level)->distributions.f[0],
+ para->getParD(level)->rho,
+ para->getParD(level)->velocityX, para->getParD(level)->velocityY, para->getParD(level)->velocityZ,
+ para->getParD(level)->turbViscosity,
+ para->getSGSConstant(),
+ (unsigned long)para->getParD(level)->numberOfNodes,
+ level,
+ para->getIsBodyForce(),
+ para->getForcesDev(),
+ para->getParD(level)->forceX_SP, para->getParD(level)->forceY_SP, para->getParD(level)->forceZ_SP,
+ para->getParD(level)->coordinateX, para->getParD(level)->coordinateY, para->getParD(level)->coordinateZ,
+ para->getQuadricLimitersDev(),
+ para->getParD(level)->isEvenTimestep,
+ indices,
+ size_indices);
break;
-
+
case CollisionTemplate::WriteMacroVars:
- LB_Kernel_CumulantK17Sponge < turbulenceModel, true, false > <<< cudaGrid.grid, cudaGrid.threads, 0, stream >>>(
- para->getParD(level)->omega,
- para->getParD(level)->typeOfGridNode,
- para->getParD(level)->neighborX, para->getParD(level)->neighborY, para->getParD(level)->neighborZ,
- para->getParD(level)->distributions.f[0],
- para->getParD(level)->rho,
- para->getParD(level)->velocityX, para->getParD(level)->velocityY, para->getParD(level)->velocityZ,
- para->getParD(level)->turbViscosity,
- para->getSGSConstant(),
- (unsigned long)para->getParD(level)->numberOfNodes,
- level,
- para->getIsBodyForce(),
- para->getForcesDev(),
- para->getParD(level)->forceX_SP, para->getParD(level)->forceY_SP, para->getParD(level)->forceZ_SP,
- para->getQuadricLimitersDev(),
- para->getParD(level)->isEvenTimestep,
- indices,
- size_indices);
+ LB_Kernel_CumulantK17Sponge < turbulenceModel, true, false > <<< cudaGrid.grid, cudaGrid.threads, 0, stream >>>(
+ para->getParD(level)->omega,
+ para->getParD(level)->typeOfGridNode,
+ para->getParD(level)->neighborX, para->getParD(level)->neighborY, para->getParD(level)->neighborZ,
+ para->getParD(level)->distributions.f[0],
+ para->getParD(level)->rho,
+ para->getParD(level)->velocityX, para->getParD(level)->velocityY, para->getParD(level)->velocityZ,
+ para->getParD(level)->turbViscosity,
+ para->getSGSConstant(),
+ (unsigned long)para->getParD(level)->numberOfNodes,
+ level,
+ para->getIsBodyForce(),
+ para->getForcesDev(),
+ para->getParD(level)->forceX_SP, para->getParD(level)->forceY_SP, para->getParD(level)->forceZ_SP,
+ para->getParD(level)->coordinateX, para->getParD(level)->coordinateY, para->getParD(level)->coordinateZ,
+ para->getQuadricLimitersDev(),
+ para->getParD(level)->isEvenTimestep,
+ indices,
+ size_indices);
break;
-
+
case CollisionTemplate::Border:
case CollisionTemplate::AllFeatures:
- LB_Kernel_CumulantK17Sponge < turbulenceModel, true, true > <<< cudaGrid.grid, cudaGrid.threads, 0, stream >>>(
- para->getParD(level)->omega,
- para->getParD(level)->typeOfGridNode,
- para->getParD(level)->neighborX, para->getParD(level)->neighborY, para->getParD(level)->neighborZ,
- para->getParD(level)->distributions.f[0],
- para->getParD(level)->rho,
- para->getParD(level)->velocityX, para->getParD(level)->velocityY, para->getParD(level)->velocityZ,
- para->getParD(level)->turbViscosity,
- para->getSGSConstant(),
- (unsigned long)para->getParD(level)->numberOfNodes,
- level,
- para->getIsBodyForce(),
- para->getForcesDev(),
- para->getParD(level)->forceX_SP, para->getParD(level)->forceY_SP, para->getParD(level)->forceZ_SP,
- para->getQuadricLimitersDev(),
- para->getParD(level)->isEvenTimestep,
- indices,
- size_indices);
+ LB_Kernel_CumulantK17Sponge < turbulenceModel, true, true > <<< cudaGrid.grid, cudaGrid.threads, 0, stream >>>(
+ para->getParD(level)->omega,
+ para->getParD(level)->typeOfGridNode,
+ para->getParD(level)->neighborX, para->getParD(level)->neighborY, para->getParD(level)->neighborZ,
+ para->getParD(level)->distributions.f[0],
+ para->getParD(level)->rho,
+ para->getParD(level)->velocityX, para->getParD(level)->velocityY, para->getParD(level)->velocityZ,
+ para->getParD(level)->turbViscosity,
+ para->getSGSConstant(),
+ (unsigned long)para->getParD(level)->numberOfNodes,
+ level,
+ para->getIsBodyForce(),
+ para->getForcesDev(),
+ para->getParD(level)->forceX_SP, para->getParD(level)->forceY_SP, para->getParD(level)->forceZ_SP,
+ para->getParD(level)->coordinateX, para->getParD(level)->coordinateY, para->getParD(level)->coordinateZ,
+ para->getQuadricLimitersDev(),
+ para->getParD(level)->isEvenTimestep,
+ indices,
+ size_indices);
break;
case CollisionTemplate::ApplyBodyForce:
- LB_Kernel_CumulantK17Sponge < turbulenceModel, false, true > <<< cudaGrid.grid, cudaGrid.threads, 0, stream >>>(
- para->getParD(level)->omega,
- para->getParD(level)->typeOfGridNode,
- para->getParD(level)->neighborX, para->getParD(level)->neighborY, para->getParD(level)->neighborZ,
- para->getParD(level)->distributions.f[0],
- para->getParD(level)->rho,
- para->getParD(level)->velocityX, para->getParD(level)->velocityY, para->getParD(level)->velocityZ,
- para->getParD(level)->turbViscosity,
- para->getSGSConstant(),
- (unsigned long)para->getParD(level)->numberOfNodes,
- level,
- para->getIsBodyForce(),
- para->getForcesDev(),
- para->getParD(level)->forceX_SP, para->getParD(level)->forceY_SP, para->getParD(level)->forceZ_SP,
- para->getQuadricLimitersDev(),
- para->getParD(level)->isEvenTimestep,
- indices,
- size_indices);
+ LB_Kernel_CumulantK17Sponge < turbulenceModel, false, true > <<< cudaGrid.grid, cudaGrid.threads, 0, stream >>>(
+ para->getParD(level)->omega,
+ para->getParD(level)->typeOfGridNode,
+ para->getParD(level)->neighborX, para->getParD(level)->neighborY, para->getParD(level)->neighborZ,
+ para->getParD(level)->distributions.f[0],
+ para->getParD(level)->rho,
+ para->getParD(level)->velocityX, para->getParD(level)->velocityY, para->getParD(level)->velocityZ,
+ para->getParD(level)->turbViscosity,
+ para->getSGSConstant(),
+ (unsigned long)para->getParD(level)->numberOfNodes,
+ level,
+ para->getIsBodyForce(),
+ para->getForcesDev(),
+ para->getParD(level)->forceX_SP, para->getParD(level)->forceY_SP, para->getParD(level)->forceZ_SP,
+ para->getParD(level)->coordinateX, para->getParD(level)->coordinateY, para->getParD(level)->coordinateZ,
+ para->getQuadricLimitersDev(),
+ para->getParD(level)->isEvenTimestep,
+ indices,
+ size_indices);
break;
default:
throw std::runtime_error("Invalid CollisionTemplate in CumulantK17Sponge::runOnIndices()");
@@ -146,7 +152,7 @@ CumulantK17Sponge<turbulenceModel>::CumulantK17Sponge(std::shared_ptr<Parameter>
this->cudaGrid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes);
this->kernelUsesFluidNodeIndices = true;
-
+
VF_LOG_INFO("Using turbulence model: {}", turbulenceModel);
}
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Sponge/CumulantK17Sponge_Device.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Sponge/CumulantK17Sponge_Device.cu
index df5495432..83c31acdb 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Sponge/CumulantK17Sponge_Device.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Sponge/CumulantK17Sponge_Device.cu
@@ -76,6 +76,9 @@ __global__ void LB_Kernel_CumulantK17Sponge(
real* bodyForceX,
real* bodyForceY,
real* bodyForceZ,
+ real* coordX,
+ real* coordY,
+ real* coordZ,
real* quadricLimiters,
bool isEvenTimestep,
const uint *fluidNodeIndices,
@@ -695,34 +698,34 @@ __global__ void LB_Kernel_CumulantK17Sponge(
(dist.f[DIR_MMM])[k_MMM] = f_PPP;
}
-template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::AMD, true, true > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* bodyForceZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
+template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::AMD, true, true > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* bodyForceZ, real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
-template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::Smagorinsky, true, true > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* bodyForceZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
+template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::Smagorinsky, true, true > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
-template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::QR, true, true > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* bodyForceZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
+template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::QR, true, true > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
-template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::None, true, true > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* bodyForceZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
+template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::None, true, true > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
-template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::AMD, true, false > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* bodyForceZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
+template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::AMD, true, false > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
-template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::Smagorinsky, true, false > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* bodyForceZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
+template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::Smagorinsky, true, false > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
-template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::QR, true, false > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* bodyForceZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
+template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::QR, true, false > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
-template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::None, true, false > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* bodyForceZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
+template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::None, true, false > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
-template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::AMD, false, true > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* bodyForceZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
+template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::AMD, false, true > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
-template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::Smagorinsky, false, true > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* bodyForceZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
+template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::Smagorinsky, false, true > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
-template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::QR, false, true > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* bodyForceZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
+template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::QR, false, true > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
-template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::None, false, true > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* bodyForceZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
+template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::None, false, true > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
-template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::AMD, false, false > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* bodyForceZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
+template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::AMD, false, false > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
-template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::Smagorinsky, false, false > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* bodyForceZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
+template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::Smagorinsky, false, false > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
-template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::QR, false, false > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* bodyForceZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
+template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::QR, false, false > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
-template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::None, false, false > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* bodyForceZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
\ No newline at end of file
+template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::None, false, false > ( real omega_in, uint* typeOfGridNode, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, bool bodyForce, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes);
\ No newline at end of file
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Sponge/CumulantK17Sponge_Device.cuh b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Sponge/CumulantK17Sponge_Device.cuh
index 28da79958..d119fa439 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Sponge/CumulantK17Sponge_Device.cuh
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Sponge/CumulantK17Sponge_Device.cuh
@@ -24,6 +24,9 @@ template< TurbulenceModel turbulenceModel, bool writeMacroscopicVariables, bool
real* bodyForceX,
real* bodyForceY,
real* bodyForceZ,
+ real* coordX,
+ real* coordY,
+ real* coordZ,
real* quadricLimiters,
bool isEvenTimestep,
const uint *fluidNodeIndices,
--
GitLab