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

patch for stencil example

Browse Source
This commit is contained in:
Evghenii
2014-01-13 08:50:21 +01:00
parent 8190869714
commit 8f17468fa3
3 changed files with 11 additions and 5 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
examples_ptx/stencil/Makefile_gpu
View File

@@ -2,7 +2,7 @@ PROG=stencil
ISPC_SRC=stencil.ispc ISPC_SRC=stencil.ispc
CU_SRC=stencil.cu CU_SRC=stencil.cu
CXX_SRC=stencil.cpp stencil_serial.cpp CXX_SRC=stencil.cpp stencil_serial.cpp
PTXCC_REGMAX=128 PTXCC_REGMAX=92
LLVM_GPU=1 LLVM_GPU=1
NVVM_GPU=1 NVVM_GPU=1

4
examples_ptx/stencil/stencil.cu
View File

@@ -22,7 +22,7 @@ stencil_step( int x0, int x1,
const int x = xb + programIndex; const int x = xb + programIndex;
int index = (z * Nxy) + (y * Nx) + x; int index = (z * Nxy) + (y * Nx) + x;
#define A_cur(x, y, z) __ldg(&Ain[index + (x) + ((y) * Nx) + ((z) * Nxy)]) #define A_cur(x, y, z) *(&Ain[index + (x) + ((y) * Nx) + ((z) * Nxy)])
#define A_next(x, y, z) Aout[index + (x) + ((y) * Nx) + ((z) * Nxy)] #define A_next(x, y, z) Aout[index + (x) + ((y) * Nx) + ((z) * Nxy)]
double div = double div =
coef0 * A_cur(0, 0, 0) + coef0 * A_cur(0, 0, 0) +
@@ -38,7 +38,7 @@ stencil_step( int x0, int x1,
if (x < x1) if (x < x1)
A_next(0, 0, 0) = 2.0 * A_cur(0, 0, 0) - A_next(0, 0, 0) + A_next(0, 0, 0) = 2.0 * A_cur(0, 0, 0) - A_next(0, 0, 0) +
__ldg(&vsq[index]) * div; *(&vsq[index]) * div;
} }
} }

10
examples_ptx/stencil/stencil.ispc
View File

@@ -100,9 +100,15 @@ stencil_step(uniform int x0, uniform int x1,
} }
} }
#ifdef __NVPTX__
#define SPANX 32 #define SPANX 32
#define SPANY 2 #define SPANY 4
#define SPANZ 4 #define SPANZ 8
#else
#define SPANX 64
#define SPANY 4
#define SPANZ 8
#endif
static task void static task void
stencil_step_task(uniform int x0, uniform int x1, stencil_step_task(uniform int x0, uniform int x1,