Regenerate EMScript

1 files changed, 322 insertions, 151 deletions

diff --git a/Examples/EMScript/src/EM_evol.cc b/Examples/EMScript/src/EM_evol.cc
index ff1c8ce..4e66517 100644
--- a/Examples/EMScript/src/EM_evol.cc
+++ b/Examples/EMScript/src/EM_evol.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_evol_SelectBCs(CCTK_ARGUMENTS)
 {
@@ -28,12 +30,12 @@ extern "C" void EM_evol_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_grouprhs","flat");
+  ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "My_New_Implementation::B_grouprhs","flat");
   if (ierr < 0)
-    CCTK_WARN(1, "Failed to register flat BC for EMScript::B_grouprhs.");
-  ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "EMScript::El_grouprhs","flat");
+    CCTK_WARN(1, "Failed to register flat BC for My_New_Implementation::B_grouprhs.");
+  ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "My_New_Implementation::El_grouprhs","flat");
   if (ierr < 0)
-    CCTK_WARN(1, "Failed to register flat BC for EMScript::El_grouprhs.");
+    CCTK_WARN(1, "Failed to register flat BC for My_New_Implementation::El_grouprhs.");
   return;
 }
 
@@ -42,152 +44,321 @@ static void EM_evol_Body(cGH const * restrict const cctkGH, int const dir, int c
   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_evol,\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 B1L = vec_load(B1[index]);\n"
+  "  CCTK_REAL_VEC B2L = vec_load(B2[index]);\n"
+  "  CCTK_REAL_VEC B3L = vec_load(B3[index]);\n"
+  "  CCTK_REAL_VEC El1L = vec_load(El1[index]);\n"
+  "  CCTK_REAL_VEC El2L = vec_load(El2[index]);\n"
+  "  CCTK_REAL_VEC El3L = vec_load(El3[index]);\n"
+  "  \n"
+  "  \n"
+  "  CCTK_REAL_VEC J11L, J12L, J13L, J21L, J22L, J23L, J31L, J32L, J33L;\n"
+  "  \n"
+  "  if (use_jacobian)\n"
+  "  {\n"
+  "    J11L = vec_load(J11[index]);\n"
+  "    J12L = vec_load(J12[index]);\n"
+  "    J13L = vec_load(J13[index]);\n"
+  "    J21L = vec_load(J21[index]);\n"
+  "    J22L = vec_load(J22[index]);\n"
+  "    J23L = vec_load(J23[index]);\n"
+  "    J31L = vec_load(J31[index]);\n"
+  "    J32L = vec_load(J32[index]);\n"
+  "    J33L = vec_load(J33[index]);\n"
+  "  }\n"
+  "  \n"
+  "  /* Include user supplied include files */\n"
+  "  \n"
+  "  /* Precompute derivatives */\n"
+  "  CCTK_REAL_VEC PDstandard1B1;\n"
+  "  CCTK_REAL_VEC PDstandard2B1;\n"
+  "  CCTK_REAL_VEC PDstandard3B1;\n"
+  "  CCTK_REAL_VEC PDstandard1B2;\n"
+  "  CCTK_REAL_VEC PDstandard2B2;\n"
+  "  CCTK_REAL_VEC PDstandard3B2;\n"
+  "  CCTK_REAL_VEC PDstandard1B3;\n"
+  "  CCTK_REAL_VEC PDstandard2B3;\n"
+  "  CCTK_REAL_VEC PDstandard3B3;\n"
+  "  CCTK_REAL_VEC PDstandard1El1;\n"
+  "  CCTK_REAL_VEC PDstandard2El1;\n"
+  "  CCTK_REAL_VEC PDstandard3El1;\n"
+  "  CCTK_REAL_VEC PDstandard1El2;\n"
+  "  CCTK_REAL_VEC PDstandard2El2;\n"
+  "  CCTK_REAL_VEC PDstandard3El2;\n"
+  "  CCTK_REAL_VEC PDstandard1El3;\n"
+  "  CCTK_REAL_VEC PDstandard2El3;\n"
+  "  CCTK_REAL_VEC PDstandard3El3;\n"
+  "  \n"
+  "  switch(fdOrder)\n"
+  "  {\n"
+  "    case 2:\n"
+  "      PDstandard1B1 = PDstandardfdOrder21(&B1[index]);\n"
+  "      PDstandard2B1 = PDstandardfdOrder22(&B1[index]);\n"
+  "      PDstandard3B1 = PDstandardfdOrder23(&B1[index]);\n"
+  "      PDstandard1B2 = PDstandardfdOrder21(&B2[index]);\n"
+  "      PDstandard2B2 = PDstandardfdOrder22(&B2[index]);\n"
+  "      PDstandard3B2 = PDstandardfdOrder23(&B2[index]);\n"
+  "      PDstandard1B3 = PDstandardfdOrder21(&B3[index]);\n"
+  "      PDstandard2B3 = PDstandardfdOrder22(&B3[index]);\n"
+  "      PDstandard3B3 = PDstandardfdOrder23(&B3[index]);\n"
+  "      PDstandard1El1 = PDstandardfdOrder21(&El1[index]);\n"
+  "      PDstandard2El1 = PDstandardfdOrder22(&El1[index]);\n"
+  "      PDstandard3El1 = PDstandardfdOrder23(&El1[index]);\n"
+  "      PDstandard1El2 = PDstandardfdOrder21(&El2[index]);\n"
+  "      PDstandard2El2 = PDstandardfdOrder22(&El2[index]);\n"
+  "      PDstandard3El2 = PDstandardfdOrder23(&El2[index]);\n"
+  "      PDstandard1El3 = PDstandardfdOrder21(&El3[index]);\n"
+  "      PDstandard2El3 = PDstandardfdOrder22(&El3[index]);\n"
+  "      PDstandard3El3 = PDstandardfdOrder23(&El3[index]);\n"
+  "      break;\n"
+  "    \n"
+  "    case 4:\n"
+  "      PDstandard1B1 = PDstandardfdOrder41(&B1[index]);\n"
+  "      PDstandard2B1 = PDstandardfdOrder42(&B1[index]);\n"
+  "      PDstandard3B1 = PDstandardfdOrder43(&B1[index]);\n"
+  "      PDstandard1B2 = PDstandardfdOrder41(&B2[index]);\n"
+  "      PDstandard2B2 = PDstandardfdOrder42(&B2[index]);\n"
+  "      PDstandard3B2 = PDstandardfdOrder43(&B2[index]);\n"
+  "      PDstandard1B3 = PDstandardfdOrder41(&B3[index]);\n"
+  "      PDstandard2B3 = PDstandardfdOrder42(&B3[index]);\n"
+  "      PDstandard3B3 = PDstandardfdOrder43(&B3[index]);\n"
+  "      PDstandard1El1 = PDstandardfdOrder41(&El1[index]);\n"
+  "      PDstandard2El1 = PDstandardfdOrder42(&El1[index]);\n"
+  "      PDstandard3El1 = PDstandardfdOrder43(&El1[index]);\n"
+  "      PDstandard1El2 = PDstandardfdOrder41(&El2[index]);\n"
+  "      PDstandard2El2 = PDstandardfdOrder42(&El2[index]);\n"
+  "      PDstandard3El2 = PDstandardfdOrder43(&El2[index]);\n"
+  "      PDstandard1El3 = PDstandardfdOrder41(&El3[index]);\n"
+  "      PDstandard2El3 = PDstandardfdOrder42(&El3[index]);\n"
+  "      PDstandard3El3 = PDstandardfdOrder43(&El3[index]);\n"
+  "      break;\n"
+  "  }\n"
+  "  \n"
+  "  /* Calculate temporaries and grid functions */\n"
+  "  CCTK_REAL_VEC JacPDstandard1B2;\n"
+  "  CCTK_REAL_VEC JacPDstandard1B3;\n"
+  "  CCTK_REAL_VEC JacPDstandard1El2;\n"
+  "  CCTK_REAL_VEC JacPDstandard1El3;\n"
+  "  CCTK_REAL_VEC JacPDstandard2B1;\n"
+  "  CCTK_REAL_VEC JacPDstandard2B3;\n"
+  "  CCTK_REAL_VEC JacPDstandard2El1;\n"
+  "  CCTK_REAL_VEC JacPDstandard2El3;\n"
+  "  CCTK_REAL_VEC JacPDstandard3B1;\n"
+  "  CCTK_REAL_VEC JacPDstandard3B2;\n"
+  "  CCTK_REAL_VEC JacPDstandard3El1;\n"
+  "  CCTK_REAL_VEC JacPDstandard3El2;\n"
+  "  \n"
+  "  if (use_jacobian)\n"
+  "  {\n"
+  "    JacPDstandard1B2 = \n"
+  "      kmadd(J11L,PDstandard1B2,kmadd(J21L,PDstandard2B2,kmul(J31L,PDstandard3B2)));\n"
+  "    \n"
+  "    JacPDstandard1B3 = \n"
+  "      kmadd(J11L,PDstandard1B3,kmadd(J21L,PDstandard2B3,kmul(J31L,PDstandard3B3)));\n"
+  "    \n"
+  "    JacPDstandard1El2 = \n"
+  "      kmadd(J11L,PDstandard1El2,kmadd(J21L,PDstandard2El2,kmul(J31L,PDstandard3El2)));\n"
+  "    \n"
+  "    JacPDstandard1El3 = \n"
+  "      kmadd(J11L,PDstandard1El3,kmadd(J21L,PDstandard2El3,kmul(J31L,PDstandard3El3)));\n"
+  "    \n"
+  "    JacPDstandard2B1 = \n"
+  "      kmadd(J12L,PDstandard1B1,kmadd(J22L,PDstandard2B1,kmul(J32L,PDstandard3B1)));\n"
+  "    \n"
+  "    JacPDstandard2B3 = \n"
+  "      kmadd(J12L,PDstandard1B3,kmadd(J22L,PDstandard2B3,kmul(J32L,PDstandard3B3)));\n"
+  "    \n"
+  "    JacPDstandard2El1 = \n"
+  "      kmadd(J12L,PDstandard1El1,kmadd(J22L,PDstandard2El1,kmul(J32L,PDstandard3El1)));\n"
+  "    \n"
+  "    JacPDstandard2El3 = \n"
+  "      kmadd(J12L,PDstandard1El3,kmadd(J22L,PDstandard2El3,kmul(J32L,PDstandard3El3)));\n"
+  "    \n"
+  "    JacPDstandard3B1 = \n"
+  "      kmadd(J13L,PDstandard1B1,kmadd(J23L,PDstandard2B1,kmul(J33L,PDstandard3B1)));\n"
+  "    \n"
+  "    JacPDstandard3B2 = \n"
+  "      kmadd(J13L,PDstandard1B2,kmadd(J23L,PDstandard2B2,kmul(J33L,PDstandard3B2)));\n"
+  "    \n"
+  "    JacPDstandard3El1 = \n"
+  "      kmadd(J13L,PDstandard1El1,kmadd(J23L,PDstandard2El1,kmul(J33L,PDstandard3El1)));\n"
+  "    \n"
+  "    JacPDstandard3El2 = \n"
+  "      kmadd(J13L,PDstandard1El2,kmadd(J23L,PDstandard2El2,kmul(J33L,PDstandard3El2)));\n"
+  "  }\n"
+  "  else\n"
+  "  {\n"
+  "    JacPDstandard1B2 = PDstandard1B2;\n"
+  "    \n"
+  "    JacPDstandard1B3 = PDstandard1B3;\n"
+  "    \n"
+  "    JacPDstandard1El2 = PDstandard1El2;\n"
+  "    \n"
+  "    JacPDstandard1El3 = PDstandard1El3;\n"
+  "    \n"
+  "    JacPDstandard2B1 = PDstandard2B1;\n"
+  "    \n"
+  "    JacPDstandard2B3 = PDstandard2B3;\n"
+  "    \n"
+  "    JacPDstandard2El1 = PDstandard2El1;\n"
+  "    \n"
+  "    JacPDstandard2El3 = PDstandard2El3;\n"
+  "    \n"
+  "    JacPDstandard3B1 = PDstandard3B1;\n"
+  "    \n"
+  "    JacPDstandard3B2 = PDstandard3B2;\n"
+  "    \n"
+  "    JacPDstandard3El1 = PDstandard3El1;\n"
+  "    \n"
+  "    JacPDstandard3El2 = PDstandard3El2;\n"
+  "  }\n"
+  "  \n"
+  "  CCTK_REAL_VEC El1rhsL = ksub(JacPDstandard2B3,JacPDstandard3B2);\n"
+  "  \n"
+  "  CCTK_REAL_VEC El2rhsL = ksub(JacPDstandard3B1,JacPDstandard1B3);\n"
+  "  \n"
+  "  CCTK_REAL_VEC El3rhsL = ksub(JacPDstandard1B2,JacPDstandard2B1);\n"
+  "  \n"
+  "  CCTK_REAL_VEC B1rhsL = ksub(JacPDstandard3El2,JacPDstandard2El3);\n"
+  "  \n"
+  "  CCTK_REAL_VEC B2rhsL = ksub(JacPDstandard1El3,JacPDstandard3El1);\n"
+  "  \n"
+  "  CCTK_REAL_VEC B3rhsL = ksub(JacPDstandard2El1,JacPDstandard1El2);\n"
+  "  \n"
+  "  /* Copy local copies back to grid functions */\n"
+  "  vec_store_nta_partial(B1rhs[index],B1rhsL);\n"
+  "  vec_store_nta_partial(B2rhs[index],B2rhsL);\n"
+  "  vec_store_nta_partial(B3rhs[index],B3rhsL);\n"
+  "  vec_store_nta_partial(El1rhs[index],El1rhsL);\n"
+  "  vec_store_nta_partial(El2rhs[index],El2rhsL);\n"
+  "  vec_store_nta_partial(El3rhs[index],El3rhsL);\n"
+  "}\n"
+  "LC_ENDLOOP3VEC (EM_evol);\n"
+  ""
+  ;
   
-  /* Declare finite differencing variables */
+  char const * const groups[] = {"My_New_Implementation::B_group","My_New_Implementation::B_grouprhs","My_New_Implementation::El_group","My_New_Implementation::El_grouprhs",NULL};
   
-  /* Include user-supplied include files */
+  static struct OpenCLKernel * kernel = NULL;
+  char const * const sources[] = {differencing, source, NULL};
+  OpenCLRunTime_CallKernel (cctkGH, CCTK_THORNSTRING, "EM_evol",
+                            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_evol,
-    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 B1L = B1[index];
-    CCTK_REAL B2L = B2[index];
-    CCTK_REAL B3L = B3[index];
-    CCTK_REAL El1L = El1[index];
-    CCTK_REAL El2L = El2[index];
-    CCTK_REAL El3L = El3[index];
-    
-    
-    /* Include user supplied include files */
-    
-    /* Precompute derivatives */
-    CCTK_REAL PDstandard2B1;
-    CCTK_REAL PDstandard3B1;
-    CCTK_REAL PDstandard1B2;
-    CCTK_REAL PDstandard3B2;
-    CCTK_REAL PDstandard1B3;
-    CCTK_REAL PDstandard2B3;
-    CCTK_REAL PDstandard2El1;
-    CCTK_REAL PDstandard3El1;
-    CCTK_REAL PDstandard1El2;
-    CCTK_REAL PDstandard3El2;
-    CCTK_REAL PDstandard1El3;
-    CCTK_REAL PDstandard2El3;
-    
-    switch(fdOrder)
-    {
-      case 2:
-        PDstandard2B1 = PDstandardfdOrder22(&B1[index]);
-        PDstandard3B1 = PDstandardfdOrder23(&B1[index]);
-        PDstandard1B2 = PDstandardfdOrder21(&B2[index]);
-        PDstandard3B2 = PDstandardfdOrder23(&B2[index]);
-        PDstandard1B3 = PDstandardfdOrder21(&B3[index]);
-        PDstandard2B3 = PDstandardfdOrder22(&B3[index]);
-        PDstandard2El1 = PDstandardfdOrder22(&El1[index]);
-        PDstandard3El1 = PDstandardfdOrder23(&El1[index]);
-        PDstandard1El2 = PDstandardfdOrder21(&El2[index]);
-        PDstandard3El2 = PDstandardfdOrder23(&El2[index]);
-        PDstandard1El3 = PDstandardfdOrder21(&El3[index]);
-        PDstandard2El3 = PDstandardfdOrder22(&El3[index]);
-        break;
-      
-      case 4:
-        PDstandard2B1 = PDstandardfdOrder42(&B1[index]);
-        PDstandard3B1 = PDstandardfdOrder43(&B1[index]);
-        PDstandard1B2 = PDstandardfdOrder41(&B2[index]);
-        PDstandard3B2 = PDstandardfdOrder43(&B2[index]);
-        PDstandard1B3 = PDstandardfdOrder41(&B3[index]);
-        PDstandard2B3 = PDstandardfdOrder42(&B3[index]);
-        PDstandard2El1 = PDstandardfdOrder42(&El1[index]);
-        PDstandard3El1 = PDstandardfdOrder43(&El1[index]);
-        PDstandard1El2 = PDstandardfdOrder41(&El2[index]);
-        PDstandard3El2 = PDstandardfdOrder43(&El2[index]);
-        PDstandard1El3 = PDstandardfdOrder41(&El3[index]);
-        PDstandard2El3 = PDstandardfdOrder42(&El3[index]);
-        break;
-    }
-    
-    /* Calculate temporaries and grid functions */
-    CCTK_REAL El1rhsL = PDstandard2B3 - PDstandard3B2;
-    
-    CCTK_REAL El2rhsL = -PDstandard1B3 + PDstandard3B1;
-    
-    CCTK_REAL El3rhsL = PDstandard1B2 - PDstandard2B1;
-    
-    CCTK_REAL B1rhsL = -PDstandard2El3 + PDstandard3El2;
-    
-    CCTK_REAL B2rhsL = PDstandard1El3 - PDstandard3El1;
-    
-    CCTK_REAL B3rhsL = -PDstandard1El2 + PDstandard2El1;
-    
-    /* Copy local copies back to grid functions */
-    B1rhs[index] = B1rhsL;
-    B2rhs[index] = B2rhsL;
-    B3rhs[index] = B3rhsL;
-    El1rhs[index] = El1rhsL;
-    El2rhs[index] = El2rhsL;
-    El3rhs[index] = El3rhsL;
-  }
-  CCTK_ENDLOOP3 (EM_evol);
 }
 
 extern "C" void EM_evol(CCTK_ARGUMENTS)
@@ -206,7 +377,7 @@ extern "C" void EM_evol(CCTK_ARGUMENTS)
     return;
   }
   
-  const char *groups[] = {"EMScript::B_group","EMScript::B_grouprhs","EMScript::El_group","EMScript::El_grouprhs"};
+  const char *groups[] = {"My_New_Implementation::B_group","My_New_Implementation::B_grouprhs","My_New_Implementation::El_group","My_New_Implementation::El_grouprhs"};
   GenericFD_AssertGroupStorage(cctkGH, "EM_evol", 4, groups);
   
   switch(fdOrder)