Skip to content
Snippets Groups Projects
Commit b10b3a90 authored by Hkorb's avatar Hkorb
Browse files

check for zero denominator in AMD

parent 380e1b4a
No related branches found
No related tags found
1 merge request!170Kernel templetization and efficiency improvements
...@@ -5,6 +5,7 @@ ...@@ -5,6 +5,7 @@
#include <cuda_runtime.h> #include <cuda_runtime.h>
#include <helper_cuda.h> #include <helper_cuda.h>
#include "LBM/LB.h" #include "LBM/LB.h"
#include "Kernel/Utilities/DistributionHelper.cuh"
using namespace vf::lbm::constant; using namespace vf::lbm::constant;
...@@ -31,15 +32,7 @@ __global__ void calcAMD(real* vx, ...@@ -31,15 +32,7 @@ __global__ void calcAMD(real* vx,
uint size_Mat, uint size_Mat,
real SGSConstant) real SGSConstant)
{ {
const uint k = vf::gpu::getNodeIndex();
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;
if(k >= size_Mat) return; if(k >= size_Mat) return;
if(typeOfGridNode[k] != GEO_FLUID) return; if(typeOfGridNode[k] != GEO_FLUID) return;
...@@ -69,7 +62,7 @@ __global__ void calcAMD(real* vx, ...@@ -69,7 +62,7 @@ __global__ void calcAMD(real* vx,
(dvxdx*dvzdx + dvxdy*dvzdy + dvxdz*dvzdz) * (dvxdz+dvzdx) + (dvxdx*dvzdx + dvxdy*dvzdy + dvxdz*dvzdz) * (dvxdz+dvzdx) +
(dvydx*dvzdx + dvydy*dvzdy + dvydz*dvzdz) * (dvydz+dvzdy); (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) void calcTurbulentViscosityAMD(Parameter* para, int level)
......
0% Loading or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment