Commit c1dbff50 authored by Steffen Christgau's avatar Steffen Christgau Committed by Matthias Rüster
Browse files

minor CUDA code tunings

- replace function call to powf(x, 2.0f) explicit arithmetics with x*x via macro
- avoid implicit conversion between double and float
parent 1088dd07
......@@ -47,7 +47,7 @@ __global__ void runWaveUpdateKernel( KernelData data ) {
float hh = data.h[ij] - data.cR1[ij] * ( data.fM[ij] - data.fM[data.le(ij)] + data.fN[ij] * data.cR6[j] - data.fN[data.dn(ij)]*data.cR6[j-1] );
absH = fabs(hh);
absH = fabsf(hh);
if( absH < dp.sshZeroThreshold ) {
hh = 0.f;
......@@ -87,6 +87,8 @@ __global__ void runFluxUpdateKernel( KernelData data ) {
}
#define SQR(x) ((x) * (x))
__global__ void runWaveBoundaryKernel( KernelData data ) {
KernelData& dt = data;
......@@ -97,43 +99,43 @@ __global__ void runWaveBoundaryKernel( KernelData data ) {
if( id <= dp.nI-1 ) {
ij = dt.idx(id,1);
dt.h[ij] = sqrtf( powf(dt.fN[ij],2.0f) + 0.25f*powf((dt.fM[ij] + dt.fM[dt.le(ij)]),2.0f) )*dt.cB1[id-1];
dt.h[ij] = sqrtf( SQR(dt.fN[ij]) + 0.25f*SQR((dt.fM[ij] + dt.fM[dt.le(ij)])) )*dt.cB1[id-1];
if( dt.fN[ij] > 0 ) dt.h[ij] = -dt.h[ij];
}
if( id <= dp.nI-1 ) {
ij = dt.idx(id,dp.nJ);
dt.h[ij] = sqrtf( powf(dt.fN[dt.dn(ij)],2.0f) + 0.25f*powf((dt.fM[ij] + dt.fM[dt.dn(ij)]),2.0f) )*dt.cB3[id-1];
dt.h[ij] = sqrtf( SQR(dt.fN[dt.dn(ij)]) + 0.25f*SQR((dt.fM[ij] + dt.fM[dt.dn(ij)])) )*dt.cB3[id-1];
if( dt.fN[dt.dn(ij)] < 0 ) dt.h[ij] = -dt.h[ij];
}
if( id <= dp.nJ-1 ) {
ij = dt.idx(1,id);
dt.h[ij] = sqrtf( powf(dt.fM[ij],2.0f) + 0.25f*powf((dt.fN[ij] + dt.fN[dt.dn(ij)]),2.0f) )*dt.cB2[id-1];
dt.h[ij] = sqrtf( SQR(dt.fM[ij]) + 0.25f*SQR((dt.fN[ij] + dt.fN[dt.dn(ij)])) )*dt.cB2[id-1];
if( dt.fM[ij] > 0 ) dt.h[ij] = -dt.h[ij];
}
if( id <= dp.nJ-1 ) {
ij = dt.idx(dp.nI,id);
dt.h[ij] = sqrtf( powf(dt.fM[dt.le(ij)],2.0f) + 0.25f*powf((dt.fN[ij] + dt.fN[dt.dn(ij)]),2.0f) )*dt.cB4[id-1];
dt.h[ij] = sqrtf( SQR(dt.fM[dt.le(ij)]) + 0.25f*SQR((dt.fN[ij] + dt.fN[dt.dn(ij)])) )*dt.cB4[id-1];
if( dt.fM[dt.le(ij)] < 0 ) dt.h[ij] = -dt.h[ij];
}
if( id == 2 ) {
ij = dt.idx(1,1);
dt.h[ij] = sqrtf( powf(dt.fM[ij],2.0f) + powf(dt.fN[ij],2.0f) )*dt.cB1[0];
dt.h[ij] = sqrtf( SQR(dt.fM[ij]) + SQR(dt.fN[ij]) )*dt.cB1[0];
if( dt.fN[ij] > 0 ) dt.h[ij] = -dt.h[ij];
ij = dt.idx(dp.nI,1);
dt.h[ij] = sqrtf( powf(dt.fM[dt.le(ij)],2.0f) + powf(dt.fN[ij],2.0f) )*dt.cB1[dp.nI-1];
dt.h[ij] = sqrtf( SQR(dt.fM[dt.le(ij)]) + SQR(dt.fN[ij]) )*dt.cB1[dp.nI-1];
if( dt.fN[ij] > 0 ) dt.h[ij] = -dt.h[ij];
ij = dt.idx(1,dp.nJ);
dt.h[ij] = sqrtf( powf(dt.fM[ij],2.0f) + powf(dt.fN[dt.dn(ij)],2.0f) )*dt.cB3[0];
dt.h[ij] = sqrtf( SQR(dt.fM[ij]) + SQR(dt.fN[dt.dn(ij)]) )*dt.cB3[0];
if( dt.fN[dt.dn(ij)] < 0 ) dt.h[ij] = -dt.h[ij];
ij = dt.idx(dp.nI,dp.nJ);
dt.h[ij] = sqrtf( powf(dt.fM[dt.le(ij)],2.0f) + powf(dt.fN[dt.dn(ij)],2.0f) )*dt.cB3[dp.nI-1];
dt.h[ij] = sqrtf( SQR(dt.fM[dt.le(ij)]) + SQR(dt.fN[dt.dn(ij)]) )*dt.cB3[dp.nI-1];
if( dt.fN[dt.dn(ij)] < 0 ) dt.h[ij] = -dt.h[ij];
}
}
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment