/* File produced by Kranc */ #define KRANC_C #include #include #include #include #include #include "cctk.h" #include "cctk_Arguments.h" #include "cctk_Parameters.h" #include "GenericFD.h" #include "Differencing.h" #include "cctk_Loop.h" #include "loopcontrol.h" #include "OpenCLRunTime.h" #include "vectors.h" /* Define macros used in calculations */ #define INITVALUE (42) #define ScalarINV(x) ((CCTK_REAL)1.0 / (x)) #define ScalarSQR(x) ((x) * (x)) #define ScalarCUB(x) ((x) * ScalarSQR(x)) #define ScalarQAD(x) (ScalarSQR(ScalarSQR(x))) #define INV(x) (kdiv(ToReal(1.0),x)) #define SQR(x) (kmul(x,x)) #define CUB(x) (kmul(x,SQR(x))) #define QAD(x) (SQR(SQR(x))) static void WT_CL_Gaussian_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[]) { DECLARE_CCTK_ARGUMENTS; DECLARE_CCTK_PARAMETERS; char const *const source = "\n" "/* Include user-supplied include files */\n" "\n" "/* Initialise finite differencing variables */\n" "ptrdiff_t const di = 1;\n" "ptrdiff_t const dj = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" "ptrdiff_t const dk = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" "ptrdiff_t const cdi = sizeof(CCTK_REAL) * di;\n" "ptrdiff_t const cdj = sizeof(CCTK_REAL) * dj;\n" "ptrdiff_t const cdk = sizeof(CCTK_REAL) * dk;\n" "CCTK_REAL_VEC const dx = ToReal(CCTK_DELTA_SPACE(0));\n" "CCTK_REAL_VEC const dy = ToReal(CCTK_DELTA_SPACE(1));\n" "CCTK_REAL_VEC const dz = ToReal(CCTK_DELTA_SPACE(2));\n" "CCTK_REAL_VEC const dt = ToReal(CCTK_DELTA_TIME);\n" "CCTK_REAL_VEC const t = ToReal(cctk_time);\n" "CCTK_REAL_VEC const dxi = INV(dx);\n" "CCTK_REAL_VEC const dyi = INV(dy);\n" "CCTK_REAL_VEC const dzi = INV(dz);\n" "CCTK_REAL_VEC const khalf = ToReal(0.5);\n" "CCTK_REAL_VEC const kthird = ToReal(1.0/3.0);\n" "CCTK_REAL_VEC const ktwothird = ToReal(2.0/3.0);\n" "CCTK_REAL_VEC const kfourthird = ToReal(4.0/3.0);\n" "CCTK_REAL_VEC const keightthird = ToReal(8.0/3.0);\n" "CCTK_REAL_VEC const hdxi = kmul(ToReal(0.5), dxi);\n" "CCTK_REAL_VEC const hdyi = kmul(ToReal(0.5), dyi);\n" "CCTK_REAL_VEC const hdzi = kmul(ToReal(0.5), dzi);\n" "\n" "/* Initialize predefined quantities */\n" "CCTK_REAL_VEC const p1o12dx = kdiv(ToReal(0.0833333333333333333333333333333),dx);\n" "CCTK_REAL_VEC const p1o12dy = kdiv(ToReal(0.0833333333333333333333333333333),dy);\n" "CCTK_REAL_VEC const p1o12dz = kdiv(ToReal(0.0833333333333333333333333333333),dz);\n" "CCTK_REAL_VEC const p1o144dxdy = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dy,dx));\n" "CCTK_REAL_VEC const p1o144dxdz = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dx));\n" "CCTK_REAL_VEC const p1o144dydz = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dy));\n" "CCTK_REAL_VEC const pm1o12dx2 = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dx,dx));\n" "CCTK_REAL_VEC const pm1o12dy2 = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dy,dy));\n" "CCTK_REAL_VEC const pm1o12dz2 = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dz,dz));\n" "\n" "/* Assign local copies of arrays functions */\n" "\n" "\n" "\n" "/* Calculate temporaries and arrays functions */\n" "\n" "/* Copy local copies back to grid functions */\n" "\n" "/* Loop over the grid points */\n" "#pragma omp parallel\n" "LC_LOOP3VEC(WT_CL_Gaussian,\n" " i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2],\n" " cctk_ash[0],cctk_ash[1],cctk_ash[2],\n" " CCTK_REAL_VEC_SIZE)\n" "{\n" " ptrdiff_t const index = di*i + dj*j + dk*k;\n" " \n" " /* Assign local copies of grid functions */\n" " \n" " CCTK_REAL_VEC rL = vec_load(r[index]);\n" " \n" " \n" " /* Include user supplied include files */\n" " \n" " /* Precompute derivatives */\n" " \n" " /* Calculate temporaries and grid functions */\n" " CCTK_REAL_VEC uL = \n" " kmul(kexp(kdiv(kmul(kmul(rL,rL),ToReal(-0.5)),ToReal(ScalarSQR(width)))),ToReal(amplitude));\n" " \n" " CCTK_REAL_VEC rhoL = ToReal(0);\n" " \n" " /* Copy local copies back to grid functions */\n" " vec_store_partial_prepare(i,lc_imin,lc_imax);\n" " vec_store_nta_partial(rho[index],rhoL);\n" " vec_store_nta_partial(u[index],uL);\n" "}\n" "LC_ENDLOOP3VEC(WT_CL_Gaussian);\n" "" ; char const *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}; OpenCLRunTime_CallKernel(cctkGH, CCTK_THORNSTRING, "WT_CL_Gaussian", sources, groups, NULL, NULL, NULL, -1, imin, imax, &kernel); } extern "C" void WT_CL_Gaussian(CCTK_ARGUMENTS) { DECLARE_CCTK_ARGUMENTS; DECLARE_CCTK_PARAMETERS; if (verbose > 1) { CCTK_VInfo(CCTK_THORNSTRING,"Entering WT_CL_Gaussian_Body"); } if (cctk_iteration % WT_CL_Gaussian_calc_every != WT_CL_Gaussian_calc_offset) { return; } const char *const groups[] = { "grid::coordinates", "ML_WaveToy_CL::WT_rho", "ML_WaveToy_CL::WT_u"}; GenericFD_AssertGroupStorage(cctkGH, "WT_CL_Gaussian", 3, groups); GenericFD_LoopOverEverything(cctkGH, WT_CL_Gaussian_Body); if (verbose > 1) { CCTK_VInfo(CCTK_THORNSTRING,"Leaving WT_CL_Gaussian_Body"); } }