diff options
Diffstat (limited to 'Examples/EMScript/src/EM_initial.cc')
| -rw-r--r-- | Examples/EMScript/src/EM_initial.cc | 298 |
1 files changed, 183 insertions, 115 deletions
diff --git a/Examples/EMScript/src/EM_initial.cc b/Examples/EMScript/src/EM_initial.cc index ddbd59e..34e72f1 100644 --- a/Examples/EMScript/src/EM_initial.cc +++ b/Examples/EMScript/src/EM_initial.cc @@ -14,13 +14,15 @@ #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 QAD(x) (SQR(SQR(x))) -#define INV(x) ((1.0) / (x)) -#define SQR(x) ((x) * (x)) -#define CUB(x) ((x) * (x) * (x)) +#define INV(x) (kdiv(ToReal(1.0),x)) +#define SQR(x) (kmul(x,x)) +#define CUB(x) (kmul(x,SQR(x))) extern "C" void EM_initial_SelectBCs(CCTK_ARGUMENTS) { @@ -28,12 +30,12 @@ extern "C" void EM_initial_SelectBCs(CCTK_ARGUMENTS) DECLARE_CCTK_PARAMETERS; CCTK_INT ierr = 0; - ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "EMScript::B_group","flat"); + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "My_New_Implementation::B_group","flat"); if (ierr < 0) - CCTK_WARN(1, "Failed to register flat BC for EMScript::B_group."); - ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "EMScript::El_group","flat"); + CCTK_WARN(1, "Failed to register flat BC for My_New_Implementation::B_group."); + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "My_New_Implementation::El_group","flat"); if (ierr < 0) - CCTK_WARN(1, "Failed to register flat BC for EMScript::El_group."); + CCTK_WARN(1, "Failed to register flat BC for My_New_Implementation::El_group."); return; } @@ -42,116 +44,182 @@ static void EM_initial_Body(cGH const * restrict const cctkGH, int const dir, in DECLARE_CCTK_ARGUMENTS; DECLARE_CCTK_PARAMETERS; + char const * const source = + "\n" + "/* Declare finite differencing variables */\n" + "\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 = kmul(INV(dx),ToReal(0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const p1o12dy = kmul(INV(dy),ToReal(0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const p1o12dz = kmul(INV(dz),ToReal(0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const p1o144dxdy = kmul(INV(dx),kmul(INV(dy),ToReal(0.00694444444444444444444444444444)));\n" + "CCTK_REAL_VEC const p1o144dxdz = kmul(INV(dx),kmul(INV(dz),ToReal(0.00694444444444444444444444444444)));\n" + "CCTK_REAL_VEC const p1o144dydz = kmul(INV(dy),kmul(INV(dz),ToReal(0.00694444444444444444444444444444)));\n" + "CCTK_REAL_VEC const p1o2dx = kmul(INV(dx),ToReal(0.5));\n" + "CCTK_REAL_VEC const p1o2dy = kmul(INV(dy),ToReal(0.5));\n" + "CCTK_REAL_VEC const p1o2dz = kmul(INV(dz),ToReal(0.5));\n" + "CCTK_REAL_VEC const p1o4dxdy = kmul(INV(dx),kmul(INV(dy),ToReal(0.25)));\n" + "CCTK_REAL_VEC const p1o4dxdz = kmul(INV(dx),kmul(INV(dz),ToReal(0.25)));\n" + "CCTK_REAL_VEC const p1o4dydz = kmul(INV(dy),kmul(INV(dz),ToReal(0.25)));\n" + "CCTK_REAL_VEC const p1odx2 = INV(SQR(dx));\n" + "CCTK_REAL_VEC const p1ody2 = INV(SQR(dy));\n" + "CCTK_REAL_VEC const p1odz2 = INV(SQR(dz));\n" + "CCTK_REAL_VEC const pm1o12dx2 = kmul(INV(SQR(dx)),ToReal(-0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const pm1o12dy2 = kmul(INV(SQR(dy)),ToReal(-0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const pm1o12dz2 = kmul(INV(SQR(dz)),ToReal(-0.0833333333333333333333333333333));\n" + "\n" + "/* Jacobian variable pointers */\n" + "bool const use_jacobian = (!CCTK_IsFunctionAliased(\"MultiPatch_GetMap\") || MultiPatch_GetMap(cctkGH) != jacobian_identity_map)\n" + " && strlen(jacobian_group) > 0;\n" + "if (use_jacobian && strlen(jacobian_derivative_group) == 0)\n" + "{\n" + " CCTK_WARN (1, \"GenericFD::jacobian_group and GenericFD::jacobian_derivative_group must both be set to valid group names\");\n" + "}\n" + "\n" + "CCTK_REAL const *restrict jacobian_ptrs[9];\n" + "if (use_jacobian) GenericFD_GroupDataPointers(cctkGH, jacobian_group,\n" + " 9, jacobian_ptrs);\n" + "\n" + "CCTK_REAL const *restrict const J11 = use_jacobian ? jacobian_ptrs[0] : 0;\n" + "CCTK_REAL const *restrict const J12 = use_jacobian ? jacobian_ptrs[1] : 0;\n" + "CCTK_REAL const *restrict const J13 = use_jacobian ? jacobian_ptrs[2] : 0;\n" + "CCTK_REAL const *restrict const J21 = use_jacobian ? jacobian_ptrs[3] : 0;\n" + "CCTK_REAL const *restrict const J22 = use_jacobian ? jacobian_ptrs[4] : 0;\n" + "CCTK_REAL const *restrict const J23 = use_jacobian ? jacobian_ptrs[5] : 0;\n" + "CCTK_REAL const *restrict const J31 = use_jacobian ? jacobian_ptrs[6] : 0;\n" + "CCTK_REAL const *restrict const J32 = use_jacobian ? jacobian_ptrs[7] : 0;\n" + "CCTK_REAL const *restrict const J33 = use_jacobian ? jacobian_ptrs[8] : 0;\n" + "\n" + "CCTK_REAL const *restrict jacobian_derivative_ptrs[18];\n" + "if (use_jacobian) GenericFD_GroupDataPointers(cctkGH, jacobian_derivative_group,\n" + " 18, jacobian_derivative_ptrs);\n" + "\n" + "CCTK_REAL const *restrict const dJ111 = use_jacobian ? jacobian_derivative_ptrs[0] : 0;\n" + "CCTK_REAL const *restrict const dJ112 = use_jacobian ? jacobian_derivative_ptrs[1] : 0;\n" + "CCTK_REAL const *restrict const dJ113 = use_jacobian ? jacobian_derivative_ptrs[2] : 0;\n" + "CCTK_REAL const *restrict const dJ122 = use_jacobian ? jacobian_derivative_ptrs[3] : 0;\n" + "CCTK_REAL const *restrict const dJ123 = use_jacobian ? jacobian_derivative_ptrs[4] : 0;\n" + "CCTK_REAL const *restrict const dJ133 = use_jacobian ? jacobian_derivative_ptrs[5] : 0;\n" + "CCTK_REAL const *restrict const dJ211 = use_jacobian ? jacobian_derivative_ptrs[6] : 0;\n" + "CCTK_REAL const *restrict const dJ212 = use_jacobian ? jacobian_derivative_ptrs[7] : 0;\n" + "CCTK_REAL const *restrict const dJ213 = use_jacobian ? jacobian_derivative_ptrs[8] : 0;\n" + "CCTK_REAL const *restrict const dJ222 = use_jacobian ? jacobian_derivative_ptrs[9] : 0;\n" + "CCTK_REAL const *restrict const dJ223 = use_jacobian ? jacobian_derivative_ptrs[10] : 0;\n" + "CCTK_REAL const *restrict const dJ233 = use_jacobian ? jacobian_derivative_ptrs[11] : 0;\n" + "CCTK_REAL const *restrict const dJ311 = use_jacobian ? jacobian_derivative_ptrs[12] : 0;\n" + "CCTK_REAL const *restrict const dJ312 = use_jacobian ? jacobian_derivative_ptrs[13] : 0;\n" + "CCTK_REAL const *restrict const dJ313 = use_jacobian ? jacobian_derivative_ptrs[14] : 0;\n" + "CCTK_REAL const *restrict const dJ322 = use_jacobian ? jacobian_derivative_ptrs[15] : 0;\n" + "CCTK_REAL const *restrict const dJ323 = use_jacobian ? jacobian_derivative_ptrs[16] : 0;\n" + "CCTK_REAL const *restrict const dJ333 = use_jacobian ? jacobian_derivative_ptrs[17] : 0;\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 (EM_initial,\n" + " i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2],\n" + " cctk_lsh[0],cctk_lsh[1],cctk_lsh[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 xL = vec_load(x[index]);\n" + " CCTK_REAL_VEC yL = vec_load(y[index]);\n" + " \n" + " \n" + " \n" + " /* Include user supplied include files */\n" + " \n" + " /* Precompute derivatives */\n" + " \n" + " switch(fdOrder)\n" + " {\n" + " case 2:\n" + " break;\n" + " \n" + " case 4:\n" + " break;\n" + " }\n" + " \n" + " /* Calculate temporaries and grid functions */\n" + " CCTK_REAL_VEC sigma = ToReal(1);\n" + " \n" + " CCTK_REAL_VEC csetemp0 = kadd(xL,yL);\n" + " \n" + " CCTK_REAL_VEC El1L = \n" + " kmul(sigma,kcos(kmul(csetemp0,kmul(Pi,ToReal(2)))));\n" + " \n" + " CCTK_REAL_VEC csetemp1 = ksub(ToReal(1),sigma);\n" + " \n" + " CCTK_REAL_VEC csetemp2 = kmul(xL,kmul(Pi,ToReal(2)));\n" + " \n" + " CCTK_REAL_VEC El2L = \n" + " knmadd(sigma,kcos(kmul(csetemp0,kmul(Pi,ToReal(2)))),kmul(csetemp1,kcos(csetemp2)));\n" + " \n" + " CCTK_REAL_VEC El3L = ToReal(0);\n" + " \n" + " CCTK_REAL_VEC B1L = ToReal(0);\n" + " \n" + " CCTK_REAL_VEC B2L = ToReal(0);\n" + " \n" + " CCTK_REAL_VEC B3L = \n" + " kmadd(csetemp1,kcos(csetemp2),kmul(sigma,kcos(kmul(csetemp0,kmul(Pi,ToReal(2))))));\n" + " \n" + " /* Copy local copies back to grid functions */\n" + " vec_store_nta_partial(B1[index],B1L);\n" + " vec_store_nta_partial(B2[index],B2L);\n" + " vec_store_nta_partial(B3[index],B3L);\n" + " vec_store_nta_partial(El1[index],El1L);\n" + " vec_store_nta_partial(El2[index],El2L);\n" + " vec_store_nta_partial(El3[index],El3L);\n" + "}\n" + "LC_ENDLOOP3VEC (EM_initial);\n" + "" + ; - /* Declare finite differencing variables */ + char const * const groups[] = {"My_New_Implementation::B_group","My_New_Implementation::El_group","grid::coordinates",NULL}; - /* Include user-supplied include files */ + static struct OpenCLKernel * kernel = NULL; + char const * const sources[] = {differencing, source, NULL}; + OpenCLRunTime_CallKernel (cctkGH, CCTK_THORNSTRING, "EM_initial", + sources, groups, NULL, NULL, NULL, -1, + imin, imax, &kernel); - /* Initialise finite differencing variables */ - ptrdiff_t const di = 1; - ptrdiff_t const dj = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0); - ptrdiff_t const dk = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0); - ptrdiff_t const cdi = sizeof(CCTK_REAL) * di; - ptrdiff_t const cdj = sizeof(CCTK_REAL) * dj; - ptrdiff_t const cdk = sizeof(CCTK_REAL) * dk; - CCTK_REAL const dx = ToReal(CCTK_DELTA_SPACE(0)); - CCTK_REAL const dy = ToReal(CCTK_DELTA_SPACE(1)); - CCTK_REAL const dz = ToReal(CCTK_DELTA_SPACE(2)); - CCTK_REAL const dt = ToReal(CCTK_DELTA_TIME); - CCTK_REAL const t = ToReal(cctk_time); - CCTK_REAL const dxi = INV(dx); - CCTK_REAL const dyi = INV(dy); - CCTK_REAL const dzi = INV(dz); - CCTK_REAL const khalf = 0.5; - CCTK_REAL const kthird = 1/3.0; - CCTK_REAL const ktwothird = 2.0/3.0; - CCTK_REAL const kfourthird = 4.0/3.0; - CCTK_REAL const keightthird = 8.0/3.0; - CCTK_REAL const hdxi = 0.5 * dxi; - CCTK_REAL const hdyi = 0.5 * dyi; - CCTK_REAL const hdzi = 0.5 * dzi; - - /* Initialize predefined quantities */ - CCTK_REAL const p1o12dx = 0.0833333333333333333333333333333*INV(dx); - CCTK_REAL const p1o12dy = 0.0833333333333333333333333333333*INV(dy); - CCTK_REAL const p1o12dz = 0.0833333333333333333333333333333*INV(dz); - CCTK_REAL const p1o144dxdy = 0.00694444444444444444444444444444*INV(dx)*INV(dy); - CCTK_REAL const p1o144dxdz = 0.00694444444444444444444444444444*INV(dx)*INV(dz); - CCTK_REAL const p1o144dydz = 0.00694444444444444444444444444444*INV(dy)*INV(dz); - CCTK_REAL const p1o2dx = 0.5*INV(dx); - CCTK_REAL const p1o2dy = 0.5*INV(dy); - CCTK_REAL const p1o2dz = 0.5*INV(dz); - CCTK_REAL const p1o4dxdy = 0.25*INV(dx)*INV(dy); - CCTK_REAL const p1o4dxdz = 0.25*INV(dx)*INV(dz); - CCTK_REAL const p1o4dydz = 0.25*INV(dy)*INV(dz); - CCTK_REAL const p1odx2 = INV(SQR(dx)); - CCTK_REAL const p1ody2 = INV(SQR(dy)); - CCTK_REAL const p1odz2 = INV(SQR(dz)); - CCTK_REAL const pm1o12dx2 = -0.0833333333333333333333333333333*INV(SQR(dx)); - CCTK_REAL const pm1o12dy2 = -0.0833333333333333333333333333333*INV(SQR(dy)); - CCTK_REAL const pm1o12dz2 = -0.0833333333333333333333333333333*INV(SQR(dz)); - - /* Assign local copies of arrays functions */ - - - - /* Calculate temporaries and arrays functions */ - - /* Copy local copies back to grid functions */ - - /* Loop over the grid points */ - #pragma omp parallel - CCTK_LOOP3 (EM_initial, - i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2], - cctk_lsh[0],cctk_lsh[1],cctk_lsh[2]) - { - ptrdiff_t const index = di*i + dj*j + dk*k; - - /* Assign local copies of grid functions */ - - CCTK_REAL xL = x[index]; - CCTK_REAL yL = y[index]; - - - /* Include user supplied include files */ - - /* Precompute derivatives */ - - switch(fdOrder) - { - case 2: - break; - - case 4: - break; - } - - /* Calculate temporaries and grid functions */ - CCTK_REAL sigma = 1; - - CCTK_REAL El1L = sigma*Cos(2*(xL + yL)*Pi); - - CCTK_REAL El2L = (-1 + sigma)*Cos(2*xL*Pi) - sigma*Cos(2*(xL + - yL)*Pi); - - CCTK_REAL El3L = 0; - - CCTK_REAL B1L = 0; - - CCTK_REAL B2L = 0; - - CCTK_REAL B3L = -((-1 + sigma)*Cos(2*xL*Pi)) + sigma*Cos(2*(xL + - yL)*Pi); - - /* Copy local copies back to grid functions */ - B1[index] = B1L; - B2[index] = B2L; - B3[index] = B3L; - El1[index] = El1L; - El2[index] = El2L; - El3[index] = El3L; - } - CCTK_ENDLOOP3 (EM_initial); } extern "C" void EM_initial(CCTK_ARGUMENTS) @@ -170,7 +238,7 @@ extern "C" void EM_initial(CCTK_ARGUMENTS) return; } - const char *groups[] = {"EMScript::B_group","EMScript::El_group","grid::coordinates"}; + const char *groups[] = {"My_New_Implementation::B_group","My_New_Implementation::El_group","grid::coordinates"}; GenericFD_AssertGroupStorage(cctkGH, "EM_initial", 3, groups); switch(fdOrder) |