diff --git a/examples_cuda/stencil/stencil.cu b/examples_cuda/stencil/stencil.cu index 1533505b..9240c708 100644 --- a/examples_cuda/stencil/stencil.cu +++ b/examples_cuda/stencil/stencil.cu @@ -31,31 +31,31 @@ stencil_step( int x0, int x1, const int x = xb + programIndex; #endif - int index = (z * Nxy) + (y * Nx) + x; -#define A_cur(x, y, z) Ain[index + (x) + ((y) * Nx) + ((z) * Nxy)] + int index = (z * Nxy) + (y * Nx) + x; +#define A_cur(x, y, z) __ldg(&Ain[index + (x) + ((y) * Nx) + ((z) * Nxy)]) #define A_next(x, y, z) Aout[index + (x) + ((y) * Nx) + ((z) * Nxy)] - double div = - coef0 * A_cur(0, 0, 0) + - coef1 * (A_cur(+1, 0, 0) + A_cur(-1, 0, 0) + - A_cur(0, +1, 0) + A_cur(0, -1, 0) + - A_cur(0, 0, +1) + A_cur(0, 0, -1)) + - coef2 * (A_cur(+2, 0, 0) + A_cur(-2, 0, 0) + - A_cur(0, +2, 0) + A_cur(0, -2, 0) + - A_cur(0, 0, +2) + A_cur(0, 0, -2)) + - coef3 * (A_cur(+3, 0, 0) + A_cur(-3, 0, 0) + - A_cur(0, +3, 0) + A_cur(0, -3, 0) + - A_cur(0, 0, +3) + A_cur(0, 0, -3)); + double div = + coef0 * A_cur(0, 0, 0) + + coef1 * (A_cur(+1, 0, 0) + A_cur(-1, 0, 0) + + A_cur(0, +1, 0) + A_cur(0, -1, 0) + + A_cur(0, 0, +1) + A_cur(0, 0, -1)) + + coef2 * (A_cur(+2, 0, 0) + A_cur(-2, 0, 0) + + A_cur(0, +2, 0) + A_cur(0, -2, 0) + + A_cur(0, 0, +2) + A_cur(0, 0, -2)) + + coef3 * (A_cur(+3, 0, 0) + A_cur(-3, 0, 0) + + A_cur(0, +3, 0) + A_cur(0, -3, 0) + + A_cur(0, 0, +3) + A_cur(0, 0, -3)); - if (x < x1) - A_next(0, 0, 0) = 2.0 * A_cur(0, 0, 0) - A_next(0, 0, 0) + - vsq[index] * div; - } + if (x < x1) + A_next(0, 0, 0) = 2.0 * A_cur(0, 0, 0) - A_next(0, 0, 0) + + __ldg(&vsq[index]) * div; + } } #define SPANX 32 -#define SPANY 8 -#define SPANZ 8 +#define SPANY 2 +#define SPANZ 4 __global__ void stencil_step_task( int x0, int x1, diff --git a/examples_cuda/stencil/stencil.ispc b/examples_cuda/stencil/stencil.ispc index 72c28ef6..c4746868 100644 --- a/examples_cuda/stencil/stencil.ispc +++ b/examples_cuda/stencil/stencil.ispc @@ -40,8 +40,44 @@ stencil_step(uniform int x0, uniform int x1, uniform const double Ain[], uniform double Aout[]) { const uniform int Nxy = Nx * Ny; +#if 0 +#define VER1 +#endif + +#ifdef VER1 + const uniform int x1o = 1; + const uniform int x2o = 2; + const uniform int x3o = 3; + const uniform int y1o = Nx; + const uniform int y2o = Nx*2; + const uniform int y3o = Nx*3; + const uniform int z1o = Nxy; + const uniform int z2o = Nxy*2; + const uniform int z3o = Nxy*3; +#endif + foreach (z = z0 ... z1, y = y0 ... y1, x = x0 ... x1) { +#ifdef VER1 + + int index = (z * Nxy) + (y * Nx) + x; +#define A_cur(x, y, z) Ain [index + (x) + (y) + (z)] +#define A_next(x, y, z) Aout[index + (x) + (y) + (z)] + double div = coef[0] * A_cur(0, 0, 0) + + coef[1] * (A_cur(+x1o, 0, 0) + A_cur(-x1o, 0, 0) + + A_cur(0, +y1o, 0) + A_cur(0, -y1o, 0) + + A_cur(0, 0, +z1o) + A_cur(0, 0, -z1o)) + + coef[2] * (A_cur(+x2o, 0, 0) + A_cur(-x2o, 0, 0) + + A_cur(0, +y2o, 0) + A_cur(0, -y2o, 0) + + A_cur(0, 0, +z2o) + A_cur(0, 0, -z2o)) + + coef[3] * (A_cur(+x3o, 0, 0) + A_cur(-x3o, 0, 0) + + A_cur(0, +y3o, 0) + A_cur(0, -y3o, 0) + + A_cur(0, 0, +z3o) + A_cur(0, 0, -z3o)); + A_next(0, 0, 0) = 2.0d0 * A_cur(0, 0, 0) - A_next(0, 0, 0) + + vsq[index] * div; + +#else + int index = (z * Nxy) + (y * Nx) + x; #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)] @@ -59,12 +95,14 @@ stencil_step(uniform int x0, uniform int x1, A_next(0, 0, 0) = 2.0 * A_cur(0, 0, 0) - A_next(0, 0, 0) + vsq[index] * div; +#endif + } } #define SPANX 32 -#define SPANY 8 -#define SPANZ 8 +#define SPANY 2 +#define SPANZ 4 static task void stencil_step_task(uniform int x0, uniform int x1, diff --git a/examples_cuda/stencil/stencil_cu.cpp b/examples_cuda/stencil/stencil_cu.cpp index a4674f59..de5f1854 100644 --- a/examples_cuda/stencil/stencil_cu.cpp +++ b/examples_cuda/stencil/stencil_cu.cpp @@ -129,7 +129,7 @@ int main() { // the minimum time of three runs. // double minTimeISPCTasks = 1e30; - const bool print_log = false; + bool print_log = true; const int nreg = 128; for (int i = 0; i < 3; ++i) { reset_and_start_timer(); @@ -152,6 +152,7 @@ int main() { &x0, &x1, &y0, &y1, &z0, &z1, &Nx, &Ny, &Nz, &d_coeff, &d_vsq, &d_Aispc0, &d_Aispc1}; double dt = 1e3*CUDALaunch(NULL, func_name, func_args, print_log, nreg); + print_log = false; minTimeISPCTasks = std::min(minTimeISPCTasks, dt); } memcpyD2H(Aispc[1], d_Aispc1, bufsize);