aaron/ispc
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

+1

Browse Source
This commit is contained in:
Evghenii
2014-02-02 19:01:10 +01:00
parent 92e69c8197
commit 6d034596d3
2 changed files with 28 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files

14
examples_ptx/nbody/nbody.cu
View File

@@ -39,6 +39,16 @@ void closeNbody()
closeNbody___export<<<1,1>>>(); closeNbody___export<<<1,1>>>();
} }
static inline __device__
real rsqrt_real(real r2)
{
#if 1
return r2> (real)0.0 ? (real)1.0/sqrt(r2) : 0;
#else
return r2> (real)0.0 ? rsqrt((float)r2) : 0;
#endif
}
__global__ __global__
@@ -75,7 +85,7 @@ void computeForces(
const real dy = jposy - iposy; const real dy = jposy - iposy;
const real dz = jposz - iposz; const real dz = jposz - iposz;
const real r2 = dx*dx + dy*dy + dz*dz; const real r2 = dx*dx + dy*dy + dz*dz;
const real rinv = r2; // > 0.0 ? rsqrt((float)r2) : 0; const real rinv = rsqrt_real(r2);
const real mrinv = -jmass * rinv; const real mrinv = -jmass * rinv;
const real mrinv3 = mrinv * rinv*rinv; const real mrinv3 = mrinv * rinv*rinv;
iaccx += mrinv3 * dx; iaccx += mrinv3 * dx;
@@ -118,7 +128,7 @@ void computeForces(
const real dy = jposy - iposy; const real dy = jposy - iposy;
const real dz = jposz - iposz; const real dz = jposz - iposz;
const real r2 = dx*dx + dy*dy + dz*dz; const real r2 = dx*dx + dy*dy + dz*dz;
const real rinv = r2 > 0.0 ? 1.0/sqrt(r2) : 0; const real rinv = rsqrt_real(r2);
const real mrinv = -jmass * rinv; const real mrinv = -jmass * rinv;
const real mrinv3 = mrinv * rinv*rinv; const real mrinv3 = mrinv * rinv*rinv;
iaccx += mrinv3 * dx; iaccx += mrinv3 * dx;

18
examples_ptx/nbody/nbody.ispc
View File

@@ -26,6 +26,16 @@ void closeNbody()
delete gpotList; delete gpotList;
} }
static inline
real rsqrt_real(real r2)
{
#if 1
return r2> (real)0.0 ? (real)1.0d0/sqrt(r2) : 0;
#else
return r2> (real)0.0 ? rsqrt((float)r2) : 0;
#endif
}
uniform int nn = programCount; uniform int nn = programCount;
@@ -64,7 +74,7 @@ void computeForces(
const real dy = jposy - iposy; const real dy = jposy - iposy;
const real dz = jposz - iposz; const real dz = jposz - iposz;
const real r2 = dx*dx + dy*dy + dz*dz; const real r2 = dx*dx + dy*dy + dz*dz;
const real rinv = r2> 0.0d ? 1.0/sqrt(r2) : 0; const real rinv = rsqrt_real(r2);
const real mrinv = -jmass * rinv; const real mrinv = -jmass * rinv;
const real mrinv3 = mrinv * rinv*rinv; const real mrinv3 = mrinv * rinv*rinv;
@@ -104,7 +114,11 @@ void computeForces(
const real dy = jposy - iposy; const real dy = jposy - iposy;
const real dz = jposz - iposz; const real dz = jposz - iposz;
const real r2 = dx*dx + dy*dy + dz*dz; const real r2 = dx*dx + dy*dy + dz*dz;
const real rinv = r2 > 0.0d ? 1.0/sqrt(r2) : 0; #if 0 /* faster */
const real rinv = r2> (real)0.0 ? (real)1.0d0/sqrt(r2) : 0;
#else
const real rinv = rsqrt_real(r2);
#endif
const real mrinv = -jmass * rinv; const real mrinv = -jmass * rinv;
const real mrinv3 = mrinv * rinv*rinv; const real mrinv3 = mrinv * rinv*rinv;