Generate SimpleWaveOpenCL

8 files changed, 634 insertions, 0 deletions

diff --git a/Examples/SimpleWaveOpenCL/src/Boundaries.cc b/Examples/SimpleWaveOpenCL/src/Boundaries.cc
new file mode 100644
index 0000000..fc19996
--- /dev/null
+++ b/Examples/SimpleWaveOpenCL/src/Boundaries.cc
@@ -0,0 +1,197 @@
+/*  File produced by Kranc */
+
+#include "cctk.h"
+#include "cctk_Arguments.h"
+#include "cctk_Parameters.h"
+#include "cctk_Faces.h"
+#include "util_Table.h"
+#include "Symmetry.h"
+
+
+/* the boundary treatment is split into 3 steps:    */
+/* 1. excision                                      */
+/* 2. symmetries                                    */
+/* 3. "other" boundary conditions, e.g. radiative */
+
+/* to simplify scheduling and testing, the 3 steps  */
+/* are currently applied in separate functions      */
+
+
+extern "C" void SimpleWaveOpenCL_CheckBoundaries(CCTK_ARGUMENTS)
+{
+  DECLARE_CCTK_ARGUMENTS;
+  DECLARE_CCTK_PARAMETERS;
+  
+  return;
+}
+
+extern "C" void SimpleWaveOpenCL_SelectBoundConds(CCTK_ARGUMENTS)
+{
+  DECLARE_CCTK_ARGUMENTS;
+  DECLARE_CCTK_PARAMETERS;
+  
+  CCTK_INT ierr = 0;
+  
+  if (CCTK_EQUALS(evolved_group_bound, "none"  ) ||
+      CCTK_EQUALS(evolved_group_bound, "static") ||
+      CCTK_EQUALS(evolved_group_bound, "flat"  ) ||
+      CCTK_EQUALS(evolved_group_bound, "zero"  ) )
+  {
+    ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, -1,
+                      "SimpleWaveOpenCL::evolved_group", evolved_group_bound);
+    if (ierr < 0)
+       CCTK_WARN(0, "Failed to register evolved_group_bound BC for SimpleWaveOpenCL::evolved_group!");
+  }
+  
+  if (CCTK_EQUALS(phi_bound, "none"  ) ||
+      CCTK_EQUALS(phi_bound, "static") ||
+      CCTK_EQUALS(phi_bound, "flat"  ) ||
+      CCTK_EQUALS(phi_bound, "zero"  ) )
+  {
+    ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, -1,
+                      "SimpleWaveOpenCL::phi", phi_bound);
+    if (ierr < 0)
+       CCTK_WARN(0, "Failed to register phi_bound BC for SimpleWaveOpenCL::phi!");
+  }
+  
+  if (CCTK_EQUALS(pi_bound, "none"  ) ||
+      CCTK_EQUALS(pi_bound, "static") ||
+      CCTK_EQUALS(pi_bound, "flat"  ) ||
+      CCTK_EQUALS(pi_bound, "zero"  ) )
+  {
+    ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, -1,
+                      "SimpleWaveOpenCL::pi", pi_bound);
+    if (ierr < 0)
+       CCTK_WARN(0, "Failed to register pi_bound BC for SimpleWaveOpenCL::pi!");
+  }
+  
+  if (CCTK_EQUALS(evolved_group_bound, "radiative"))
+  {
+   /* select radiation boundary condition */
+    static CCTK_INT handle_evolved_group_bound = -1;
+    if (handle_evolved_group_bound < 0) handle_evolved_group_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE);
+    if (handle_evolved_group_bound < 0) CCTK_WARN(0, "could not create table!");
+    if (Util_TableSetReal(handle_evolved_group_bound , evolved_group_bound_limit, "LIMIT") < 0)
+       CCTK_WARN(0, "could not set LIMIT value in table!");
+    if (Util_TableSetReal(handle_evolved_group_bound ,evolved_group_bound_speed, "SPEED") < 0)
+       CCTK_WARN(0, "could not set SPEED value in table!");
+  
+    ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_evolved_group_bound, 
+                      "SimpleWaveOpenCL::evolved_group", "Radiation");
+  
+    if (ierr < 0)
+       CCTK_WARN(0, "Failed to register Radiation BC for SimpleWaveOpenCL::evolved_group!");
+  
+  }
+  
+  if (CCTK_EQUALS(phi_bound, "radiative"))
+  {
+   /* select radiation boundary condition */
+    static CCTK_INT handle_phi_bound = -1;
+    if (handle_phi_bound < 0) handle_phi_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE);
+    if (handle_phi_bound < 0) CCTK_WARN(0, "could not create table!");
+    if (Util_TableSetReal(handle_phi_bound , phi_bound_limit, "LIMIT") < 0)
+       CCTK_WARN(0, "could not set LIMIT value in table!");
+    if (Util_TableSetReal(handle_phi_bound ,phi_bound_speed, "SPEED") < 0)
+        CCTK_WARN(0, "could not set SPEED value in table!");
+  
+    ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_phi_bound, 
+                      "SimpleWaveOpenCL::phi", "Radiation");
+  
+    if (ierr < 0)
+       CCTK_WARN(0, "Failed to register Radiation BC for SimpleWaveOpenCL::phi!");
+  
+  }
+  
+  if (CCTK_EQUALS(pi_bound, "radiative"))
+  {
+   /* select radiation boundary condition */
+    static CCTK_INT handle_pi_bound = -1;
+    if (handle_pi_bound < 0) handle_pi_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE);
+    if (handle_pi_bound < 0) CCTK_WARN(0, "could not create table!");
+    if (Util_TableSetReal(handle_pi_bound , pi_bound_limit, "LIMIT") < 0)
+       CCTK_WARN(0, "could not set LIMIT value in table!");
+    if (Util_TableSetReal(handle_pi_bound ,pi_bound_speed, "SPEED") < 0)
+        CCTK_WARN(0, "could not set SPEED value in table!");
+  
+    ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_pi_bound, 
+                      "SimpleWaveOpenCL::pi", "Radiation");
+  
+    if (ierr < 0)
+       CCTK_WARN(0, "Failed to register Radiation BC for SimpleWaveOpenCL::pi!");
+  
+  }
+  
+  if (CCTK_EQUALS(evolved_group_bound, "scalar"))
+  {
+   /* select scalar boundary condition */
+    static CCTK_INT handle_evolved_group_bound = -1;
+    if (handle_evolved_group_bound < 0) handle_evolved_group_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE);
+    if (handle_evolved_group_bound < 0) CCTK_WARN(0, "could not create table!");
+    if (Util_TableSetReal(handle_evolved_group_bound ,evolved_group_bound_scalar, "SCALAR") < 0)
+        CCTK_WARN(0, "could not set SCALAR value in table!");
+  
+    ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_evolved_group_bound, 
+                      "SimpleWaveOpenCL::evolved_group", "scalar");
+  
+    if (ierr < 0)
+       CCTK_WARN(0, "Failed to register Scalar BC for SimpleWaveOpenCL::evolved_group!");
+  
+  }
+  
+  if (CCTK_EQUALS(phi_bound, "scalar"))
+  {
+   /* select scalar boundary condition */
+    static CCTK_INT handle_phi_bound = -1;
+    if (handle_phi_bound < 0) handle_phi_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE);
+    if (handle_phi_bound < 0) CCTK_WARN(0, "could not create table!");
+    if (Util_TableSetReal(handle_phi_bound ,phi_bound_scalar, "SCALAR") < 0)
+      CCTK_WARN(0, "could not set SCALAR value in table!");
+  
+    ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_phi_bound, 
+                      "SimpleWaveOpenCL::phi", "scalar");
+  
+    if (ierr < 0)
+       CCTK_WARN(0, "Error in registering Scalar BC for SimpleWaveOpenCL::phi!");
+  
+  }
+  
+  if (CCTK_EQUALS(pi_bound, "scalar"))
+  {
+   /* select scalar boundary condition */
+    static CCTK_INT handle_pi_bound = -1;
+    if (handle_pi_bound < 0) handle_pi_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE);
+    if (handle_pi_bound < 0) CCTK_WARN(0, "could not create table!");
+    if (Util_TableSetReal(handle_pi_bound ,pi_bound_scalar, "SCALAR") < 0)
+      CCTK_WARN(0, "could not set SCALAR value in table!");
+  
+    ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_pi_bound, 
+                      "SimpleWaveOpenCL::pi", "scalar");
+  
+    if (ierr < 0)
+       CCTK_WARN(0, "Error in registering Scalar BC for SimpleWaveOpenCL::pi!");
+  
+  }
+  return;
+}
+
+
+
+/* template for entries in parameter file:
+#$bound$#SimpleWaveOpenCL::evolved_group_bound       = "skip"
+#$bound$#SimpleWaveOpenCL::evolved_group_bound_speed = 1.0
+#$bound$#SimpleWaveOpenCL::evolved_group_bound_limit = 0.0
+#$bound$#SimpleWaveOpenCL::evolved_group_bound_scalar = 0.0
+
+#$bound$#SimpleWaveOpenCL::phi_bound       = "skip"
+#$bound$#SimpleWaveOpenCL::phi_bound_speed = 1.0
+#$bound$#SimpleWaveOpenCL::phi_bound_limit = 0.0
+#$bound$#SimpleWaveOpenCL::phi_bound_scalar = 0.0
+
+#$bound$#SimpleWaveOpenCL::pi_bound       = "skip"
+#$bound$#SimpleWaveOpenCL::pi_bound_speed = 1.0
+#$bound$#SimpleWaveOpenCL::pi_bound_limit = 0.0
+#$bound$#SimpleWaveOpenCL::pi_bound_scalar = 0.0
+
+*/
+
diff --git a/Examples/SimpleWaveOpenCL/src/Differencing.h b/Examples/SimpleWaveOpenCL/src/Differencing.h
new file mode 100644
index 0000000..476bbe1
--- /dev/null
+++ b/Examples/SimpleWaveOpenCL/src/Differencing.h
@@ -0,0 +1,75 @@
+static char const * const differencing =
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"#  define PDstandard2nd1(u) ((-KRANC_GFOFFSET3D(u,-1,0,0) + KRANC_GFOFFSET3D(u,1,0,0))*p1o2dx)\n"
+"#else\n"
+"#  define PDstandard2nd1(u) (PDstandard2nd1_impl(u,p1o2dx,cdj,cdk))\n"
+"static CCTK_REAL PDstandard2nd1_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dx, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL PDstandard2nd1_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dx, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+"  ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+"  return (-KRANC_GFOFFSET3D(u,-1,0,0) + KRANC_GFOFFSET3D(u,1,0,0))*p1o2dx;\n"
+"}\n"
+"#endif\n"
+"\n"
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"#  define PDstandard2nd2(u) ((-KRANC_GFOFFSET3D(u,0,-1,0) + KRANC_GFOFFSET3D(u,0,1,0))*p1o2dy)\n"
+"#else\n"
+"#  define PDstandard2nd2(u) (PDstandard2nd2_impl(u,p1o2dy,cdj,cdk))\n"
+"static CCTK_REAL PDstandard2nd2_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dy, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL PDstandard2nd2_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dy, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+"  ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+"  return (-KRANC_GFOFFSET3D(u,0,-1,0) + KRANC_GFOFFSET3D(u,0,1,0))*p1o2dy;\n"
+"}\n"
+"#endif\n"
+"\n"
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"#  define PDstandard2nd3(u) ((-KRANC_GFOFFSET3D(u,0,0,-1) + KRANC_GFOFFSET3D(u,0,0,1))*p1o2dz)\n"
+"#else\n"
+"#  define PDstandard2nd3(u) (PDstandard2nd3_impl(u,p1o2dz,cdj,cdk))\n"
+"static CCTK_REAL PDstandard2nd3_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dz, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL PDstandard2nd3_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dz, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+"  ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+"  return (-KRANC_GFOFFSET3D(u,0,0,-1) + KRANC_GFOFFSET3D(u,0,0,1))*p1o2dz;\n"
+"}\n"
+"#endif\n"
+"\n"
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"#  define PDstandard2nd11(u) ((-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,-1,0,0) + KRANC_GFOFFSET3D(u,1,0,0))*p1odx2)\n"
+"#else\n"
+"#  define PDstandard2nd11(u) (PDstandard2nd11_impl(u,p1odx2,cdj,cdk))\n"
+"static CCTK_REAL PDstandard2nd11_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1odx2, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL PDstandard2nd11_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1odx2, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+"  ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+"  return (-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,-1,0,0) + KRANC_GFOFFSET3D(u,1,0,0))*p1odx2;\n"
+"}\n"
+"#endif\n"
+"\n"
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"#  define PDstandard2nd22(u) ((-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,0,-1,0) + KRANC_GFOFFSET3D(u,0,1,0))*p1ody2)\n"
+"#else\n"
+"#  define PDstandard2nd22(u) (PDstandard2nd22_impl(u,p1ody2,cdj,cdk))\n"
+"static CCTK_REAL PDstandard2nd22_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1ody2, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL PDstandard2nd22_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1ody2, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+"  ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+"  return (-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,0,-1,0) + KRANC_GFOFFSET3D(u,0,1,0))*p1ody2;\n"
+"}\n"
+"#endif\n"
+"\n"
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"#  define PDstandard2nd33(u) ((-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,0,0,-1) + KRANC_GFOFFSET3D(u,0,0,1))*p1odz2)\n"
+"#else\n"
+"#  define PDstandard2nd33(u) (PDstandard2nd33_impl(u,p1odz2,cdj,cdk))\n"
+"static CCTK_REAL PDstandard2nd33_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1odz2, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL PDstandard2nd33_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1odz2, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+"  ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+"  return (-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,0,0,-1) + KRANC_GFOFFSET3D(u,0,0,1))*p1odz2;\n"
+"}\n"
+"#endif\n"
+"\n"
+""
+;
diff --git a/Examples/SimpleWaveOpenCL/src/RegisterMoL.cc b/Examples/SimpleWaveOpenCL/src/RegisterMoL.cc
new file mode 100644
index 0000000..7ad8024
--- /dev/null
+++ b/Examples/SimpleWaveOpenCL/src/RegisterMoL.cc
@@ -0,0 +1,20 @@
+/*  File produced by Kranc */
+
+#include "cctk.h"
+#include "cctk_Arguments.h"
+#include "cctk_Parameters.h"
+
+extern "C" void SimpleWaveOpenCL_RegisterVars(CCTK_ARGUMENTS)
+{
+  DECLARE_CCTK_ARGUMENTS;
+  DECLARE_CCTK_PARAMETERS;
+  
+  CCTK_INT ierr = 0;
+  
+  /* Register all the evolved grid functions with MoL */
+  ierr += MoLRegisterEvolved(CCTK_VarIndex("SimpleWaveOpenCL::phi"),  CCTK_VarIndex("SimpleWaveOpenCL::phirhs"));
+  ierr += MoLRegisterEvolved(CCTK_VarIndex("SimpleWaveOpenCL::pi"),  CCTK_VarIndex("SimpleWaveOpenCL::pirhs"));
+  
+  /* Register all the evolved Array functions with MoL */
+  return;
+}
diff --git a/Examples/SimpleWaveOpenCL/src/RegisterSymmetries.cc b/Examples/SimpleWaveOpenCL/src/RegisterSymmetries.cc
new file mode 100644
index 0000000..795d652
--- /dev/null
+++ b/Examples/SimpleWaveOpenCL/src/RegisterSymmetries.cc
@@ -0,0 +1,29 @@
+/*  File produced by Kranc */
+
+#include "cctk.h"
+#include "cctk_Arguments.h"
+#include "cctk_Parameters.h"
+#include "Symmetry.h"
+
+extern "C" void SimpleWaveOpenCL_RegisterSymmetries(CCTK_ARGUMENTS)
+{
+  DECLARE_CCTK_ARGUMENTS;
+  DECLARE_CCTK_PARAMETERS;
+  
+  
+  /* array holding symmetry definitions */
+  CCTK_INT sym[3];
+  
+  
+  /* Register symmetries of grid functions */
+  sym[0] = 1;
+  sym[1] = 1;
+  sym[2] = 1;
+  SetCartSymVN(cctkGH, sym, "SimpleWaveOpenCL::phi");
+  
+  sym[0] = 1;
+  sym[1] = 1;
+  sym[2] = 1;
+  SetCartSymVN(cctkGH, sym, "SimpleWaveOpenCL::pi");
+  
+}
diff --git a/Examples/SimpleWaveOpenCL/src/Startup.cc b/Examples/SimpleWaveOpenCL/src/Startup.cc
new file mode 100644
index 0000000..10e5dbf
--- /dev/null
+++ b/Examples/SimpleWaveOpenCL/src/Startup.cc
@@ -0,0 +1,10 @@
+/*  File produced by Kranc */
+
+#include "cctk.h"
+
+extern "C" int SimpleWaveOpenCL_Startup(void)
+{
+  const char * banner = "SimpleWaveOpenCL";
+  CCTK_RegisterBanner(banner);
+  return 0;
+}
diff --git a/Examples/SimpleWaveOpenCL/src/calc_rhs.cc b/Examples/SimpleWaveOpenCL/src/calc_rhs.cc
new file mode 100644
index 0000000..f463547
--- /dev/null
+++ b/Examples/SimpleWaveOpenCL/src/calc_rhs.cc
@@ -0,0 +1,159 @@
+/*  File produced by Kranc */
+
+#define KRANC_C
+
+#include <assert.h>
+#include <math.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#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"
+
+/* 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))
+
+extern "C" void calc_rhs_SelectBCs(CCTK_ARGUMENTS)
+{
+  DECLARE_CCTK_ARGUMENTS;
+  DECLARE_CCTK_PARAMETERS;
+  
+  CCTK_INT ierr = 0;
+  ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "SimpleWaveOpenCL::evolved_grouprhs","flat");
+  if (ierr < 0)
+    CCTK_WARN(1, "Failed to register flat BC for SimpleWaveOpenCL::evolved_grouprhs.");
+  return;
+}
+
+static void calc_rhs_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 const dx = ToReal(CCTK_DELTA_SPACE(0));\n"
+  "CCTK_REAL const dy = ToReal(CCTK_DELTA_SPACE(1));\n"
+  "CCTK_REAL const dz = ToReal(CCTK_DELTA_SPACE(2));\n"
+  "CCTK_REAL const dt = ToReal(CCTK_DELTA_TIME);\n"
+  "CCTK_REAL const t = ToReal(cctk_time);\n"
+  "CCTK_REAL const dxi = INV(dx);\n"
+  "CCTK_REAL const dyi = INV(dy);\n"
+  "CCTK_REAL const dzi = INV(dz);\n"
+  "CCTK_REAL const khalf = 0.5;\n"
+  "CCTK_REAL const kthird = 1/3.0;\n"
+  "CCTK_REAL const ktwothird = 2.0/3.0;\n"
+  "CCTK_REAL const kfourthird = 4.0/3.0;\n"
+  "CCTK_REAL const keightthird = 8.0/3.0;\n"
+  "CCTK_REAL const hdxi = 0.5 * dxi;\n"
+  "CCTK_REAL const hdyi = 0.5 * dyi;\n"
+  "CCTK_REAL const hdzi = 0.5 * dzi;\n"
+  "\n"
+  "/* Initialize predefined quantities */\n"
+  "CCTK_REAL const p1o2dx = 0.5*INV(dx);\n"
+  "CCTK_REAL const p1o2dy = 0.5*INV(dy);\n"
+  "CCTK_REAL const p1o2dz = 0.5*INV(dz);\n"
+  "CCTK_REAL const p1odx2 = INV(SQR(dx));\n"
+  "CCTK_REAL const p1ody2 = INV(SQR(dy));\n"
+  "CCTK_REAL const p1odz2 = INV(SQR(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"
+  "CCTK_LOOP3 (calc_rhs,\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"
+  "{\n"
+  "  ptrdiff_t const index = di*i + dj*j + dk*k;\n"
+  "  \n"
+  "  /* Assign local copies of grid functions */\n"
+  "  \n"
+  "  CCTK_REAL phiL = phi[index];\n"
+  "  CCTK_REAL piL = pi[index];\n"
+  "  \n"
+  "  \n"
+  "  /* Include user supplied include files */\n"
+  "  \n"
+  "  /* Precompute derivatives */\n"
+  "  CCTK_REAL const PDstandard2nd11phi = PDstandard2nd11(&phi[index]);\n"
+  "  CCTK_REAL const PDstandard2nd22phi = PDstandard2nd22(&phi[index]);\n"
+  "  CCTK_REAL const PDstandard2nd33phi = PDstandard2nd33(&phi[index]);\n"
+  "  \n"
+  "  /* Calculate temporaries and grid functions */\n"
+  "  CCTK_REAL phirhsL = piL;\n"
+  "  \n"
+  "  CCTK_REAL pirhsL = PDstandard2nd11phi + PDstandard2nd22phi + \n"
+  "    PDstandard2nd33phi;\n"
+  "  \n"
+  "  /* Copy local copies back to grid functions */\n"
+  "  vec_store_nta_partial(phirhs[index],phirhsL);\n"
+  "  vec_store_nta_partial(pirhs[index],pirhsL);\n"
+  "}\n"
+  "CCTK_ENDLOOP3 (calc_rhs);\n"
+  ""
+  ;
+  
+  char const * const groups[] = {"SimpleWaveOpenCL::evolved_group","SimpleWaveOpenCL::evolved_grouprhs",NULL};
+  
+  static struct OpenCLKernel * kernel = NULL;
+  char const * const sources[] = {differencing, source, NULL};
+  OpenCLRunTime_CallKernel (cctkGH, CCTK_THORNSTRING, "calc_rhs",
+                            sources, groups, NULL, NULL, NULL, -1,
+                            imin, imax, &kernel);
+  
+}
+
+extern "C" void calc_rhs(CCTK_ARGUMENTS)
+{
+  DECLARE_CCTK_ARGUMENTS;
+  DECLARE_CCTK_PARAMETERS;
+  
+  
+  if (verbose > 1)
+  {
+    CCTK_VInfo(CCTK_THORNSTRING,"Entering calc_rhs_Body");
+  }
+  
+  if (cctk_iteration % calc_rhs_calc_every != calc_rhs_calc_offset)
+  {
+    return;
+  }
+  
+  const char *groups[] = {"SimpleWaveOpenCL::evolved_group","SimpleWaveOpenCL::evolved_grouprhs"};
+  GenericFD_AssertGroupStorage(cctkGH, "calc_rhs", 2, groups);
+  
+  GenericFD_EnsureStencilFits(cctkGH, "calc_rhs", 1, 1, 1);
+  
+  GenericFD_LoopOverInterior(cctkGH, &calc_rhs_Body);
+  
+  if (verbose > 1)
+  {
+    CCTK_VInfo(CCTK_THORNSTRING,"Leaving calc_rhs_Body");
+  }
+}
diff --git a/Examples/SimpleWaveOpenCL/src/initial_sine.cc b/Examples/SimpleWaveOpenCL/src/initial_sine.cc
new file mode 100644
index 0000000..fa553dd
--- /dev/null
+++ b/Examples/SimpleWaveOpenCL/src/initial_sine.cc
@@ -0,0 +1,141 @@
+/*  File produced by Kranc */
+
+#define KRANC_C
+
+#include <assert.h>
+#include <math.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#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"
+
+/* 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))
+
+static void initial_sine_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 const dx = ToReal(CCTK_DELTA_SPACE(0));\n"
+  "CCTK_REAL const dy = ToReal(CCTK_DELTA_SPACE(1));\n"
+  "CCTK_REAL const dz = ToReal(CCTK_DELTA_SPACE(2));\n"
+  "CCTK_REAL const dt = ToReal(CCTK_DELTA_TIME);\n"
+  "CCTK_REAL const t = ToReal(cctk_time);\n"
+  "CCTK_REAL const dxi = INV(dx);\n"
+  "CCTK_REAL const dyi = INV(dy);\n"
+  "CCTK_REAL const dzi = INV(dz);\n"
+  "CCTK_REAL const khalf = 0.5;\n"
+  "CCTK_REAL const kthird = 1/3.0;\n"
+  "CCTK_REAL const ktwothird = 2.0/3.0;\n"
+  "CCTK_REAL const kfourthird = 4.0/3.0;\n"
+  "CCTK_REAL const keightthird = 8.0/3.0;\n"
+  "CCTK_REAL const hdxi = 0.5 * dxi;\n"
+  "CCTK_REAL const hdyi = 0.5 * dyi;\n"
+  "CCTK_REAL const hdzi = 0.5 * dzi;\n"
+  "\n"
+  "/* Initialize predefined quantities */\n"
+  "CCTK_REAL const p1o2dx = 0.5*INV(dx);\n"
+  "CCTK_REAL const p1o2dy = 0.5*INV(dy);\n"
+  "CCTK_REAL const p1o2dz = 0.5*INV(dz);\n"
+  "CCTK_REAL const p1odx2 = INV(SQR(dx));\n"
+  "CCTK_REAL const p1ody2 = INV(SQR(dy));\n"
+  "CCTK_REAL const p1odz2 = INV(SQR(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"
+  "CCTK_LOOP3 (initial_sine,\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"
+  "{\n"
+  "  ptrdiff_t const index = di*i + dj*j + dk*k;\n"
+  "  \n"
+  "  /* Assign local copies of grid functions */\n"
+  "  \n"
+  "  CCTK_REAL xL = x[index];\n"
+  "  \n"
+  "  \n"
+  "  /* Include user supplied include files */\n"
+  "  \n"
+  "  /* Precompute derivatives */\n"
+  "  \n"
+  "  /* Calculate temporaries and grid functions */\n"
+  "  CCTK_REAL phiL = Sin(2*Pi*(xL - t));\n"
+  "  \n"
+  "  CCTK_REAL piL = -2*Pi*Cos(2*Pi*(xL - t));\n"
+  "  \n"
+  "  /* Copy local copies back to grid functions */\n"
+  "  vec_store_nta_partial(phi[index],phiL);\n"
+  "  vec_store_nta_partial(pi[index],piL);\n"
+  "}\n"
+  "CCTK_ENDLOOP3 (initial_sine);\n"
+  ""
+  ;
+  
+  char const * const groups[] = {"SimpleWaveOpenCL::evolved_group","grid::coordinates",NULL};
+  
+  static struct OpenCLKernel * kernel = NULL;
+  char const * const sources[] = {differencing, source, NULL};
+  OpenCLRunTime_CallKernel (cctkGH, CCTK_THORNSTRING, "initial_sine",
+                            sources, groups, NULL, NULL, NULL, -1,
+                            imin, imax, &kernel);
+  
+}
+
+extern "C" void initial_sine(CCTK_ARGUMENTS)
+{
+  DECLARE_CCTK_ARGUMENTS;
+  DECLARE_CCTK_PARAMETERS;
+  
+  
+  if (verbose > 1)
+  {
+    CCTK_VInfo(CCTK_THORNSTRING,"Entering initial_sine_Body");
+  }
+  
+  if (cctk_iteration % initial_sine_calc_every != initial_sine_calc_offset)
+  {
+    return;
+  }
+  
+  const char *groups[] = {"SimpleWaveOpenCL::evolved_group","grid::coordinates"};
+  GenericFD_AssertGroupStorage(cctkGH, "initial_sine", 2, groups);
+  
+  
+  GenericFD_LoopOverEverything(cctkGH, &initial_sine_Body);
+  
+  if (verbose > 1)
+  {
+    CCTK_VInfo(CCTK_THORNSTRING,"Leaving initial_sine_Body");
+  }
+}
diff --git a/Examples/SimpleWaveOpenCL/src/make.code.defn b/Examples/SimpleWaveOpenCL/src/make.code.defn
new file mode 100644
index 0000000..535f5d0
--- /dev/null
+++ b/Examples/SimpleWaveOpenCL/src/make.code.defn
@@ -0,0 +1,3 @@
+# File produced by Kranc
+
+SRCS = Startup.cc RegisterMoL.cc RegisterSymmetries.cc initial_sine.cc calc_rhs.cc Boundaries.cc