From b10b3a90bd5904e382215aec93b7367b13858bb6 Mon Sep 17 00:00:00 2001 From: Henry <henry.korb@geo.uu.se> Date: Thu, 1 Sep 2022 10:55:55 +0200 Subject: [PATCH] check for zero denominator in AMD --- src/gpu/VirtualFluids_GPU/GPU/TurbulentViscosity.cu | 13 +++---------- 1 file changed, 3 insertions(+), 10 deletions(-) diff --git a/src/gpu/VirtualFluids_GPU/GPU/TurbulentViscosity.cu b/src/gpu/VirtualFluids_GPU/GPU/TurbulentViscosity.cu index ed7a199e4..d3faab442 100644 --- a/src/gpu/VirtualFluids_GPU/GPU/TurbulentViscosity.cu +++ b/src/gpu/VirtualFluids_GPU/GPU/TurbulentViscosity.cu @@ -5,6 +5,7 @@ #include <cuda_runtime.h> #include <helper_cuda.h> #include "LBM/LB.h" +#include "Kernel/Utilities/DistributionHelper.cuh" using namespace vf::lbm::constant; @@ -31,15 +32,7 @@ __global__ void calcAMD(real* vx, uint size_Mat, real SGSConstant) { - - const uint x = threadIdx.x; - const uint y = blockIdx.x; - const uint z = blockIdx.y; - - const uint nx = blockDim.x; - const uint ny = gridDim.x; - - const uint k = nx*(ny*z + y) + x; + const uint k = vf::gpu::getNodeIndex(); if(k >= size_Mat) return; if(typeOfGridNode[k] != GEO_FLUID) return; @@ -69,7 +62,7 @@ __global__ void calcAMD(real* vx, (dvxdx*dvzdx + dvxdy*dvzdy + dvxdz*dvzdz) * (dvxdz+dvzdx) + (dvydx*dvzdx + dvydy*dvzdy + dvydz*dvzdz) * (dvydz+dvzdy); - turbulentViscosity[k] = max(c0o1,-SGSConstant*enumerator)/denominator; + turbulentViscosity[k] = denominator != c0o1 ? max(c0o1,-SGSConstant*enumerator)/denominator : c0o1; } void calcTurbulentViscosityAMD(Parameter* para, int level) -- GitLab