diff options
Diffstat (limited to 'ML_WaveToy_CL/src/WT_CL_Standing.cc')
| -rw-r--r-- | ML_WaveToy_CL/src/WT_CL_Standing.cc | 91 |
1 files changed, 49 insertions, 42 deletions
diff --git a/ML_WaveToy_CL/src/WT_CL_Standing.cc b/ML_WaveToy_CL/src/WT_CL_Standing.cc index 3f79b33..6d590b8 100644 --- a/ML_WaveToy_CL/src/WT_CL_Standing.cc +++ b/ML_WaveToy_CL/src/WT_CL_Standing.cc @@ -28,49 +28,49 @@ #define CUB(x) (kmul(x,SQR(x))) #define QAD(x) (SQR(SQR(x))) -static void WT_CL_Standing_Body(cGH const * restrict const cctkGH, int const dir, int const face, CCTK_REAL const normal[3], CCTK_REAL const tangentA[3], CCTK_REAL const tangentB[3], int const imin[3], int const imax[3], int const n_subblock_gfs, CCTK_REAL * restrict const subblock_gfs[]) +static void WT_CL_Standing_Body(const cGH* restrict const cctkGH, const int dir, const int face, const CCTK_REAL normal[3], const CCTK_REAL tangentA[3], const CCTK_REAL tangentB[3], const int imin[3], const int imax[3], const int n_subblock_gfs, CCTK_REAL* restrict const subblock_gfs[]) { DECLARE_CCTK_ARGUMENTS; DECLARE_CCTK_PARAMETERS; - char const *const source = + const char* const source = "\n" "/* Include user-supplied include files */\n" "\n" "/* Initialise finite differencing variables */\n" - "ptrdiff_t /*const*/ di CCTK_ATTRIBUTE_UNUSED = 1;\n" - "ptrdiff_t /*const*/ dj CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" - "ptrdiff_t /*const*/ dk CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" - "ptrdiff_t /*const*/ cdi CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * di;\n" - "ptrdiff_t /*const*/ cdj CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dj;\n" - "ptrdiff_t /*const*/ cdk CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dk;\n" - "CCTK_REAL_VEC /*const*/ dx CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(0));\n" - "CCTK_REAL_VEC /*const*/ dy CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(1));\n" - "CCTK_REAL_VEC /*const*/ dz CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(2));\n" - "CCTK_REAL_VEC /*const*/ dt CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_TIME);\n" - "CCTK_REAL_VEC /*const*/ t CCTK_ATTRIBUTE_UNUSED = ToReal(cctk_time);\n" - "CCTK_REAL_VEC /*const*/ dxi CCTK_ATTRIBUTE_UNUSED = INV(dx);\n" - "CCTK_REAL_VEC /*const*/ dyi CCTK_ATTRIBUTE_UNUSED = INV(dy);\n" - "CCTK_REAL_VEC /*const*/ dzi CCTK_ATTRIBUTE_UNUSED = INV(dz);\n" - "CCTK_REAL_VEC /*const*/ khalf CCTK_ATTRIBUTE_UNUSED = ToReal(0.5);\n" - "CCTK_REAL_VEC /*const*/ kthird CCTK_ATTRIBUTE_UNUSED = ToReal(1.0/3.0);\n" - "CCTK_REAL_VEC /*const*/ ktwothird CCTK_ATTRIBUTE_UNUSED = ToReal(2.0/3.0);\n" - "CCTK_REAL_VEC /*const*/ kfourthird CCTK_ATTRIBUTE_UNUSED = ToReal(4.0/3.0);\n" - "CCTK_REAL_VEC /*const*/ keightthird CCTK_ATTRIBUTE_UNUSED = ToReal(8.0/3.0);\n" - "CCTK_REAL_VEC /*const*/ hdxi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dxi);\n" - "CCTK_REAL_VEC /*const*/ hdyi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dyi);\n" - "CCTK_REAL_VEC /*const*/ hdzi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dzi);\n" + "const ptrdiff_t di CCTK_ATTRIBUTE_UNUSED = 1;\n" + "const ptrdiff_t dj CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" + "const ptrdiff_t dk CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" + "const ptrdiff_t cdi CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * di;\n" + "const ptrdiff_t cdj CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dj;\n" + "const ptrdiff_t cdk CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dk;\n" + "const CCTK_REAL_VEC dx CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(0));\n" + "const CCTK_REAL_VEC dy CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(1));\n" + "const CCTK_REAL_VEC dz CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(2));\n" + "const CCTK_REAL_VEC dt CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_TIME);\n" + "const CCTK_REAL_VEC t CCTK_ATTRIBUTE_UNUSED = ToReal(cctk_time);\n" + "const CCTK_REAL_VEC dxi CCTK_ATTRIBUTE_UNUSED = INV(dx);\n" + "const CCTK_REAL_VEC dyi CCTK_ATTRIBUTE_UNUSED = INV(dy);\n" + "const CCTK_REAL_VEC dzi CCTK_ATTRIBUTE_UNUSED = INV(dz);\n" + "const CCTK_REAL_VEC khalf CCTK_ATTRIBUTE_UNUSED = ToReal(0.5);\n" + "const CCTK_REAL_VEC kthird CCTK_ATTRIBUTE_UNUSED = ToReal(1.0/3.0);\n" + "const CCTK_REAL_VEC ktwothird CCTK_ATTRIBUTE_UNUSED = ToReal(2.0/3.0);\n" + "const CCTK_REAL_VEC kfourthird CCTK_ATTRIBUTE_UNUSED = ToReal(4.0/3.0);\n" + "const CCTK_REAL_VEC keightthird CCTK_ATTRIBUTE_UNUSED = ToReal(8.0/3.0);\n" + "const CCTK_REAL_VEC hdxi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dxi);\n" + "const CCTK_REAL_VEC hdyi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dyi);\n" + "const CCTK_REAL_VEC hdzi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dzi);\n" "\n" "/* Initialize predefined quantities */\n" - "CCTK_REAL_VEC /*const*/ p1o12dx CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dx);\n" - "CCTK_REAL_VEC /*const*/ p1o12dy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dy);\n" - "CCTK_REAL_VEC /*const*/ p1o12dz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dz);\n" - "CCTK_REAL_VEC /*const*/ p1o144dxdy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dy,dx));\n" - "CCTK_REAL_VEC /*const*/ p1o144dxdz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dx));\n" - "CCTK_REAL_VEC /*const*/ p1o144dydz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dy));\n" - "CCTK_REAL_VEC /*const*/ pm1o12dx2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dx,dx));\n" - "CCTK_REAL_VEC /*const*/ pm1o12dy2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dy,dy));\n" - "CCTK_REAL_VEC /*const*/ pm1o12dz2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dz,dz));\n" + "const CCTK_REAL_VEC p1o12dx CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dx);\n" + "const CCTK_REAL_VEC p1o12dy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dy);\n" + "const CCTK_REAL_VEC p1o12dz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dz);\n" + "const CCTK_REAL_VEC p1o144dxdy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dy,dx));\n" + "const CCTK_REAL_VEC p1o144dxdz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dx));\n" + "const CCTK_REAL_VEC p1o144dydz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dy));\n" + "const CCTK_REAL_VEC pm1o12dx2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dx,dx));\n" + "const CCTK_REAL_VEC pm1o12dy2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dy,dy));\n" + "const CCTK_REAL_VEC pm1o12dz2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dz,dz));\n" "\n" "/* Assign local copies of arrays functions */\n" "\n" @@ -81,13 +81,20 @@ static void WT_CL_Standing_Body(cGH const * restrict const cctkGH, int const dir "/* Copy local copies back to grid functions */\n" "\n" "/* Loop over the grid points */\n" - "#pragma omp parallel\n" + "const int imin0=imin[0];\n" + "const int imin1=imin[1];\n" + "const int imin2=imin[2];\n" + "const int imax0=imax[0];\n" + "const int imax1=imax[1];\n" + "const int imax2=imax[2];\n" + "#pragma omp parallel // reduction(+: vec_iter_counter, vec_op_counter, vec_mem_counter)\n" "CCTK_LOOP3STR(WT_CL_Standing,\n" - " i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2],\n" + " i,j,k, imin0,imin1,imin2, imax0,imax1,imax2,\n" " cctk_ash[0],cctk_ash[1],cctk_ash[2],\n" - " kimin,kimax, CCTK_REAL_VEC_SIZE)\n" + " vecimin,vecimax, CCTK_REAL_VEC_SIZE)\n" "{\n" - " ptrdiff_t /*const*/ index CCTK_ATTRIBUTE_UNUSED = di*i + dj*j + dk*k;\n" + " const ptrdiff_t index CCTK_ATTRIBUTE_UNUSED = di*i + dj*j + dk*k;\n" + " // vec_iter_counter+=CCTK_REAL_VEC_SIZE;\n" " \n" " /* Assign local copies of grid functions */\n" " \n" @@ -105,13 +112,13 @@ static void WT_CL_Standing_Body(cGH const * restrict const cctkGH, int const dir " ToReal(3.14159265358979323846264338328*ScalarINV(width));\n" " \n" " CCTK_REAL_VEC CCTK_ATTRIBUTE_UNUSED omega = \n" - " ksqrt(kmul(kmul(kvec,kvec),ToReal(3.)));\n" + " ksqrt(kmul(kmul(kvec,kvec),ToReal(3)));\n" " \n" " CCTK_REAL_VEC CCTK_ATTRIBUTE_UNUSED uL = \n" " kmul(kcos(kmul(xL,kvec)),kmul(kcos(kmul(yL,kvec)),kmul(kcos(kmul(zL,kvec)),kmul(kcos(kmul(omega,t)),ToReal(amplitude)))));\n" " \n" " CCTK_REAL_VEC CCTK_ATTRIBUTE_UNUSED rhoL = \n" - " kmul(omega,kmul(kcos(kmul(xL,kvec)),kmul(kcos(kmul(yL,kvec)),kmul(kcos(kmul(zL,kvec)),kmul(ksin(kmul(omega,t)),ToReal(-1.*amplitude))))));\n" + " kneg(kmul(omega,kmul(kcos(kmul(xL,kvec)),kmul(kcos(kmul(yL,kvec)),kmul(kcos(kmul(zL,kvec)),kmul(ksin(kmul(omega,t)),ToReal(amplitude)))))));\n" " \n" " /* Copy local copies back to grid functions */\n" " vec_store_partial_prepare(i,lc_imin,lc_imax);\n" @@ -122,14 +129,14 @@ static void WT_CL_Standing_Body(cGH const * restrict const cctkGH, int const dir "" ; - char const *const groups[] = { + const char* const groups[] = { "grid::coordinates", "ML_WaveToy_CL::WT_rho", "ML_WaveToy_CL::WT_u", NULL}; static struct OpenCLKernel *kernel = NULL; - char const *const sources[] = {differencing, source, NULL}; + const char* const sources[] = {differencing, source, NULL}; OpenCLRunTime_CallKernel(cctkGH, CCTK_THORNSTRING, "WT_CL_Standing", sources, groups, NULL, NULL, NULL, -1, imin, imax, &kernel); @@ -152,7 +159,7 @@ extern "C" void WT_CL_Standing(CCTK_ARGUMENTS) return; } - const char *const groups[] = { + const char* const groups[] = { "grid::coordinates", "ML_WaveToy_CL::WT_rho", "ML_WaveToy_CL::WT_u"}; |