Replace SimpleWaveCaKernel with WaveCaKernel

This script now generates WaveCaKernel and WaveHost which do the same
thing, one using CaKernel and one using the host.  This allows easy
cross-comparison between the two methods.

34 files changed, 1517 insertions, 193 deletions

diff --git a/Examples/SimpleWaveCaKernel/schedule.ccl b/Examples/SimpleWaveCaKernel/schedule.ccl
deleted file mode 100644
index 5c777d8..0000000
--- a/Examples/SimpleWaveCaKernel/schedule.ccl
+++ /dev/null
@@ -1,118 +0,0 @@
-# File produced by Kranc
-
-
-if (other_timelevels == 1)
-{
-  STORAGE: xCopy_g[1]
-}
-
-if (timelevels == 1)
-{
-  STORAGE: phi_g[1]
-}
-if (timelevels == 2)
-{
-  STORAGE: phi_g[2]
-}
-
-if (timelevels == 1)
-{
-  STORAGE: pi_g[1]
-}
-if (timelevels == 2)
-{
-  STORAGE: pi_g[2]
-}
-
-if (rhs_timelevels == 1)
-{
-  STORAGE: phi_grhs[1]
-}
-if (rhs_timelevels == 2)
-{
-  STORAGE: phi_grhs[2]
-}
-
-if (rhs_timelevels == 1)
-{
-  STORAGE: pi_grhs[1]
-}
-if (rhs_timelevels == 2)
-{
-  STORAGE: pi_grhs[2]
-}
-
-schedule SimpleWaveCaKernel_Startup at STARTUP
-{
-  LANG: C
-  OPTIONS: meta
-} "create banner"
-
-schedule SimpleWaveCaKernel_RegisterSymmetries in SymmetryRegister
-{
-  LANG: C
-  OPTIONS: meta
-} "register symmetries"
-
-schedule initial_gaussian AT INITIAL
-{
-  LANG: C
-  READS:    grid::coordinates
-  WRITES:   SimpleWaveCaKernel::phi_g
-  WRITES:   SimpleWaveCaKernel::pi_g
-  WRITES:   SimpleWaveCaKernel::xCopy_g
-} "initial_gaussian"
-
-schedule CAKERNEL_Launch_calc_rhs in MoL_CalcRHS
-{
-  LANG: C
-  TAGS: Device=1
-  READS:    SimpleWaveCaKernel::phi_g
-  READS:    SimpleWaveCaKernel::pi_g
-  WRITES:   SimpleWaveCaKernel::phi_grhs
-  WRITES:   SimpleWaveCaKernel::pi_grhs
-} "calc_rhs"
-
-schedule CAKERNEL_Launch_copy_to_device at INITIAL after initial_gaussian
-{
-  LANG: C
-  TAGS: Device=1
-  READS:    SimpleWaveCaKernel::phi_g
-  READS:    SimpleWaveCaKernel::pi_g
-  WRITES:   SimpleWaveCaKernel::phi_g
-  WRITES:   SimpleWaveCaKernel::pi_g
-} "copy_to_device"
-
-schedule CAKERNEL_Launch_calc_bound_rhs in MoL_RHSBoundaries
-{
-  LANG: C
-  TAGS: Device=1
-  READS:    SimpleWaveCaKernel::xCopy_g
-  WRITES:   SimpleWaveCaKernel::phi_grhs
-  WRITES:   SimpleWaveCaKernel::pi_grhs
-} "calc_bound_rhs"
-
-schedule SimpleWaveCaKernel_SelectBoundConds in MoL_PostStep
-{
-  LANG: C
-  OPTIONS: level
-  SYNC: phi_g
-  SYNC: pi_g
-} "select boundary conditions"
-
-schedule SimpleWaveCaKernel_CheckBoundaries at BASEGRID
-{
-  LANG: C
-  OPTIONS: meta
-} "check boundaries treatment"
-
-schedule SimpleWaveCaKernel_RegisterVars in MoL_Register
-{
-  LANG: C
-  OPTIONS: meta
-} "Register Variables for MoL"
-
-schedule group ApplyBCs as SimpleWaveCaKernel_ApplyBCs in MoL_PostStep after SimpleWaveCaKernel_SelectBoundConds
-{
-  # no language specified
-} "Apply boundary conditions controlled by thorn Boundary"
diff --git a/Examples/SimpleWaveCaKernel.m b/Examples/WaveCaKernel.m
index f4991ea..f4991ea 100644
--- a/Examples/SimpleWaveCaKernel.m
+++ b/Examples/WaveCaKernel.m
diff --git a/Examples/SimpleWaveCaKernel/cakernel.ccl b/Examples/WaveCaKernel/cakernel.ccl
index 75c394c..2b33cf1 100644
--- a/Examples/SimpleWaveCaKernel/cakernel.ccl
+++ b/Examples/WaveCaKernel/cakernel.ccl
@@ -25,21 +25,6 @@ CCTK_CUDA_KERNEL calc_rhs TYPE=gpu_cuda/3dblock TILE="8,8,8" SHARECODE=yes STENC
   "pirhs"
 }
 
-CCTK_CUDA_KERNEL copy_to_device TYPE=gpu_cuda/3dblock TILE="8,8,8" SHARECODE=yes STENCIL="0,0,0,0,0,0"
-{
-  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=inout
-  {
-    phi
-  }
-  "phi"
-  
-  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=inout
-  {
-    pi
-  }
-  "pi"
-}
-
 CCTK_CUDA_KERNEL calc_bound_rhs TYPE=gpu_cuda/boundary TILE="8,8,8" SHARECODE=yes
 {
   CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=out
@@ -61,3 +46,18 @@ CCTK_CUDA_KERNEL calc_bound_rhs TYPE=gpu_cuda/boundary TILE="8,8,8" SHARECODE=ye
   "xCopy"
 }
 
+CCTK_CUDA_KERNEL copy_to_device TYPE=gpu_cuda/3dblock TILE="8,8,8" SHARECODE=yes STENCIL="0,0,0,0,0,0"
+{
+  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=inout
+  {
+    phi
+  }
+  "phi"
+  
+  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=inout
+  {
+    pi
+  }
+  "pi"
+}
+
diff --git a/Examples/SimpleWaveCaKernel/configuration.ccl b/Examples/WaveCaKernel/configuration.ccl
index 46dbde5..46dbde5 100644
--- a/Examples/SimpleWaveCaKernel/configuration.ccl
+++ b/Examples/WaveCaKernel/configuration.ccl
diff --git a/Examples/SimpleWaveCaKernel/interface.ccl b/Examples/WaveCaKernel/interface.ccl
index 710d050..11c1416 100644
--- a/Examples/SimpleWaveCaKernel/interface.ccl
+++ b/Examples/WaveCaKernel/interface.ccl
@@ -1,6 +1,6 @@
 # File produced by Kranc
 
-implements: SimpleWaveCaKernel
+implements: WaveCaKernel
 
 inherits:   Grid GenericFD Boundary
 
diff --git a/Examples/SimpleWaveCaKernel/param.ccl b/Examples/WaveCaKernel/param.ccl
index 845e24e..1fbbee4 100644
--- a/Examples/SimpleWaveCaKernel/param.ccl
+++ b/Examples/WaveCaKernel/param.ccl
@@ -17,13 +17,13 @@ CCTK_INT verbose "verbose" STEERABLE=ALWAYS
 } 0
 
 restricted:
-CCTK_INT SimpleWaveCaKernel_MaxNumEvolvedVars "Number of evolved variables used by this thorn" ACCUMULATOR-BASE=MethodofLines::MoL_Num_Evolved_Vars STEERABLE=RECOVER
+CCTK_INT WaveCaKernel_MaxNumEvolvedVars "Number of evolved variables used by this thorn" ACCUMULATOR-BASE=MethodofLines::MoL_Num_Evolved_Vars STEERABLE=RECOVER
 {
   2:2 :: "Number of evolved variables used by this thorn"
 } 2
 
 restricted:
-CCTK_INT SimpleWaveCaKernel_MaxNumArrayEvolvedVars "Number of Array evolved variables used by this thorn" ACCUMULATOR-BASE=MethodofLines::MoL_Num_ArrayEvolved_Vars STEERABLE=RECOVER
+CCTK_INT WaveCaKernel_MaxNumArrayEvolvedVars "Number of Array evolved variables used by this thorn" ACCUMULATOR-BASE=MethodofLines::MoL_Num_ArrayEvolved_Vars STEERABLE=RECOVER
 {
   0:0 :: "Number of Array evolved variables used by this thorn"
 } 0
@@ -59,13 +59,13 @@ CCTK_INT calc_rhs_calc_every "calc_rhs_calc_every" STEERABLE=ALWAYS
 } 1
 
 restricted:
-CCTK_INT copy_to_device_calc_every "copy_to_device_calc_every" STEERABLE=ALWAYS
+CCTK_INT calc_bound_rhs_calc_every "calc_bound_rhs_calc_every" STEERABLE=ALWAYS
 {
   *:* :: ""
 } 1
 
 restricted:
-CCTK_INT calc_bound_rhs_calc_every "calc_bound_rhs_calc_every" STEERABLE=ALWAYS
+CCTK_INT copy_to_device_calc_every "copy_to_device_calc_every" STEERABLE=ALWAYS
 {
   *:* :: ""
 } 1
@@ -83,13 +83,13 @@ CCTK_INT calc_rhs_calc_offset "calc_rhs_calc_offset" STEERABLE=ALWAYS
 } 0
 
 restricted:
-CCTK_INT copy_to_device_calc_offset "copy_to_device_calc_offset" STEERABLE=ALWAYS
+CCTK_INT calc_bound_rhs_calc_offset "calc_bound_rhs_calc_offset" STEERABLE=ALWAYS
 {
   *:* :: ""
 } 0
 
 restricted:
-CCTK_INT calc_bound_rhs_calc_offset "calc_bound_rhs_calc_offset" STEERABLE=ALWAYS
+CCTK_INT copy_to_device_calc_offset "copy_to_device_calc_offset" STEERABLE=ALWAYS
 {
   *:* :: ""
 } 0
diff --git a/Examples/WaveCaKernel/schedule.ccl b/Examples/WaveCaKernel/schedule.ccl
new file mode 100644
index 0000000..8f54e2a
--- /dev/null
+++ b/Examples/WaveCaKernel/schedule.ccl
@@ -0,0 +1,141 @@
+# File produced by Kranc
+
+
+if (other_timelevels == 1)
+{
+  STORAGE: xCopy_g[1]
+}
+
+if (timelevels == 1)
+{
+  STORAGE: phi_g[1]
+}
+if (timelevels == 2)
+{
+  STORAGE: phi_g[2]
+}
+
+if (timelevels == 1)
+{
+  STORAGE: pi_g[1]
+}
+if (timelevels == 2)
+{
+  STORAGE: pi_g[2]
+}
+
+if (rhs_timelevels == 1)
+{
+  STORAGE: phi_grhs[1]
+}
+if (rhs_timelevels == 2)
+{
+  STORAGE: phi_grhs[2]
+}
+
+if (rhs_timelevels == 1)
+{
+  STORAGE: pi_grhs[1]
+}
+if (rhs_timelevels == 2)
+{
+  STORAGE: pi_grhs[2]
+}
+
+schedule WaveCaKernel_Startup at STARTUP
+{
+  LANG: C
+  OPTIONS: meta
+} "create banner"
+
+schedule WaveCaKernel_RegisterSymmetries in SymmetryRegister
+{
+  LANG: C
+  OPTIONS: meta
+} "register symmetries"
+
+schedule initial_gaussian AT INITIAL
+{
+  LANG: C
+  READS:    grid::coordinates
+  WRITES:   WaveCaKernel::phi_g
+  WRITES:   WaveCaKernel::pi_g
+  WRITES:   WaveCaKernel::xCopy_g
+} "initial_gaussian"
+
+schedule CAKERNEL_Launch_calc_rhs in MoL_CalcRHS
+{
+  LANG: C
+  TAGS: Device=1
+  READS:    WaveCaKernel::phi_g
+  READS:    WaveCaKernel::pi_g
+  WRITES:   WaveCaKernel::phi_grhs
+  WRITES:   WaveCaKernel::pi_grhs
+} "calc_rhs"
+
+schedule CAKERNEL_Launch_calc_rhs at ANALYSIS
+{
+  LANG: C
+  SYNC: phi_grhs
+  SYNC: pi_grhs
+  TAGS: Device=1
+  READS:    WaveCaKernel::phi_g
+  READS:    WaveCaKernel::pi_g
+  WRITES:   WaveCaKernel::phi_grhs
+  WRITES:   WaveCaKernel::pi_grhs
+} "calc_rhs"
+
+schedule CAKERNEL_Launch_calc_bound_rhs in MoL_RHSBoundaries
+{
+  LANG: C
+  TAGS: Device=1
+  READS:    WaveCaKernel::xCopy_g
+  WRITES:   WaveCaKernel::phi_grhs
+  WRITES:   WaveCaKernel::pi_grhs
+} "calc_bound_rhs"
+
+schedule CAKERNEL_Launch_calc_bound_rhs at ANALYSIS
+{
+  LANG: C
+  SYNC: phi_grhs
+  SYNC: pi_grhs
+  TAGS: Device=1
+  READS:    WaveCaKernel::xCopy_g
+  WRITES:   WaveCaKernel::phi_grhs
+  WRITES:   WaveCaKernel::pi_grhs
+} "calc_bound_rhs"
+
+schedule CAKERNEL_Launch_copy_to_device at INITIAL after initial_gaussian
+{
+  LANG: C
+  TAGS: Device=1
+  READS:    WaveCaKernel::phi_g
+  READS:    WaveCaKernel::pi_g
+  WRITES:   WaveCaKernel::phi_g
+  WRITES:   WaveCaKernel::pi_g
+} "copy_to_device"
+
+schedule WaveCaKernel_SelectBoundConds in MoL_PostStep
+{
+  LANG: C
+  OPTIONS: level
+  SYNC: phi_g
+  SYNC: pi_g
+} "select boundary conditions"
+
+schedule WaveCaKernel_CheckBoundaries at BASEGRID
+{
+  LANG: C
+  OPTIONS: meta
+} "check boundaries treatment"
+
+schedule WaveCaKernel_RegisterVars in MoL_Register
+{
+  LANG: C
+  OPTIONS: meta
+} "Register Variables for MoL"
+
+schedule group ApplyBCs as WaveCaKernel_ApplyBCs in MoL_PostStep after WaveCaKernel_SelectBoundConds
+{
+  # no language specified
+} "Apply boundary conditions controlled by thorn Boundary"
diff --git a/Examples/SimpleWaveCaKernel/src/Boundaries.cc b/Examples/WaveCaKernel/src/Boundaries.cc
index 11561bf..ad79d17 100644
--- a/Examples/SimpleWaveCaKernel/src/Boundaries.cc
+++ b/Examples/WaveCaKernel/src/Boundaries.cc
@@ -17,7 +17,7 @@
 /* are currently applied in separate functions      */
 
 
-extern "C" void SimpleWaveCaKernel_CheckBoundaries(CCTK_ARGUMENTS)
+extern "C" void WaveCaKernel_CheckBoundaries(CCTK_ARGUMENTS)
 {
   DECLARE_CCTK_ARGUMENTS;
   DECLARE_CCTK_PARAMETERS;
@@ -25,7 +25,7 @@ extern "C" void SimpleWaveCaKernel_CheckBoundaries(CCTK_ARGUMENTS)
   return;
 }
 
-extern "C" void SimpleWaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS)
+extern "C" void WaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS)
 {
   DECLARE_CCTK_ARGUMENTS;
   DECLARE_CCTK_PARAMETERS;
@@ -38,9 +38,9 @@ extern "C" void SimpleWaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS)
       CCTK_EQUALS(phi_g_bound, "zero"  ) )
   {
     ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, -1,
-                      "SimpleWaveCaKernel::phi_g", phi_g_bound);
+                      "WaveCaKernel::phi_g", phi_g_bound);
     if (ierr < 0)
-       CCTK_WARN(0, "Failed to register phi_g_bound BC for SimpleWaveCaKernel::phi_g!");
+       CCTK_WARN(0, "Failed to register phi_g_bound BC for WaveCaKernel::phi_g!");
   }
   
   if (CCTK_EQUALS(pi_g_bound, "none"  ) ||
@@ -49,9 +49,9 @@ extern "C" void SimpleWaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS)
       CCTK_EQUALS(pi_g_bound, "zero"  ) )
   {
     ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, -1,
-                      "SimpleWaveCaKernel::pi_g", pi_g_bound);
+                      "WaveCaKernel::pi_g", pi_g_bound);
     if (ierr < 0)
-       CCTK_WARN(0, "Failed to register pi_g_bound BC for SimpleWaveCaKernel::pi_g!");
+       CCTK_WARN(0, "Failed to register pi_g_bound BC for WaveCaKernel::pi_g!");
   }
   
   if (CCTK_EQUALS(phi_bound, "none"  ) ||
@@ -60,9 +60,9 @@ extern "C" void SimpleWaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS)
       CCTK_EQUALS(phi_bound, "zero"  ) )
   {
     ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, -1,
-                      "SimpleWaveCaKernel::phi", phi_bound);
+                      "WaveCaKernel::phi", phi_bound);
     if (ierr < 0)
-       CCTK_WARN(0, "Failed to register phi_bound BC for SimpleWaveCaKernel::phi!");
+       CCTK_WARN(0, "Failed to register phi_bound BC for WaveCaKernel::phi!");
   }
   
   if (CCTK_EQUALS(pi_bound, "none"  ) ||
@@ -71,9 +71,9 @@ extern "C" void SimpleWaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS)
       CCTK_EQUALS(pi_bound, "zero"  ) )
   {
     ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, -1,
-                      "SimpleWaveCaKernel::pi", pi_bound);
+                      "WaveCaKernel::pi", pi_bound);
     if (ierr < 0)
-       CCTK_WARN(0, "Failed to register pi_bound BC for SimpleWaveCaKernel::pi!");
+       CCTK_WARN(0, "Failed to register pi_bound BC for WaveCaKernel::pi!");
   }
   
   if (CCTK_EQUALS(phi_g_bound, "radiative"))
@@ -88,10 +88,10 @@ extern "C" void SimpleWaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS)
        CCTK_WARN(0, "could not set SPEED value in table!");
   
     ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_phi_g_bound, 
-                      "SimpleWaveCaKernel::phi_g", "Radiation");
+                      "WaveCaKernel::phi_g", "Radiation");
   
     if (ierr < 0)
-       CCTK_WARN(0, "Failed to register Radiation BC for SimpleWaveCaKernel::phi_g!");
+       CCTK_WARN(0, "Failed to register Radiation BC for WaveCaKernel::phi_g!");
   
   }
   
@@ -107,10 +107,10 @@ extern "C" void SimpleWaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS)
        CCTK_WARN(0, "could not set SPEED value in table!");
   
     ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_pi_g_bound, 
-                      "SimpleWaveCaKernel::pi_g", "Radiation");
+                      "WaveCaKernel::pi_g", "Radiation");
   
     if (ierr < 0)
-       CCTK_WARN(0, "Failed to register Radiation BC for SimpleWaveCaKernel::pi_g!");
+       CCTK_WARN(0, "Failed to register Radiation BC for WaveCaKernel::pi_g!");
   
   }
   
@@ -126,10 +126,10 @@ extern "C" void SimpleWaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS)
         CCTK_WARN(0, "could not set SPEED value in table!");
   
     ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_phi_bound, 
-                      "SimpleWaveCaKernel::phi", "Radiation");
+                      "WaveCaKernel::phi", "Radiation");
   
     if (ierr < 0)
-       CCTK_WARN(0, "Failed to register Radiation BC for SimpleWaveCaKernel::phi!");
+       CCTK_WARN(0, "Failed to register Radiation BC for WaveCaKernel::phi!");
   
   }
   
@@ -145,10 +145,10 @@ extern "C" void SimpleWaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS)
         CCTK_WARN(0, "could not set SPEED value in table!");
   
     ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_pi_bound, 
-                      "SimpleWaveCaKernel::pi", "Radiation");
+                      "WaveCaKernel::pi", "Radiation");
   
     if (ierr < 0)
-       CCTK_WARN(0, "Failed to register Radiation BC for SimpleWaveCaKernel::pi!");
+       CCTK_WARN(0, "Failed to register Radiation BC for WaveCaKernel::pi!");
   
   }
   
@@ -162,10 +162,10 @@ extern "C" void SimpleWaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS)
         CCTK_WARN(0, "could not set SCALAR value in table!");
   
     ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_phi_g_bound, 
-                      "SimpleWaveCaKernel::phi_g", "scalar");
+                      "WaveCaKernel::phi_g", "scalar");
   
     if (ierr < 0)
-       CCTK_WARN(0, "Failed to register Scalar BC for SimpleWaveCaKernel::phi_g!");
+       CCTK_WARN(0, "Failed to register Scalar BC for WaveCaKernel::phi_g!");
   
   }
   
@@ -179,10 +179,10 @@ extern "C" void SimpleWaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS)
         CCTK_WARN(0, "could not set SCALAR value in table!");
   
     ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_pi_g_bound, 
-                      "SimpleWaveCaKernel::pi_g", "scalar");
+                      "WaveCaKernel::pi_g", "scalar");
   
     if (ierr < 0)
-       CCTK_WARN(0, "Failed to register Scalar BC for SimpleWaveCaKernel::pi_g!");
+       CCTK_WARN(0, "Failed to register Scalar BC for WaveCaKernel::pi_g!");
   
   }
   
@@ -196,10 +196,10 @@ extern "C" void SimpleWaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS)
       CCTK_WARN(0, "could not set SCALAR value in table!");
   
     ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_phi_bound, 
-                      "SimpleWaveCaKernel::phi", "scalar");
+                      "WaveCaKernel::phi", "scalar");
   
     if (ierr < 0)
-       CCTK_WARN(0, "Error in registering Scalar BC for SimpleWaveCaKernel::phi!");
+       CCTK_WARN(0, "Error in registering Scalar BC for WaveCaKernel::phi!");
   
   }
   
@@ -213,10 +213,10 @@ extern "C" void SimpleWaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS)
       CCTK_WARN(0, "could not set SCALAR value in table!");
   
     ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_pi_bound, 
-                      "SimpleWaveCaKernel::pi", "scalar");
+                      "WaveCaKernel::pi", "scalar");
   
     if (ierr < 0)
-       CCTK_WARN(0, "Error in registering Scalar BC for SimpleWaveCaKernel::pi!");
+       CCTK_WARN(0, "Error in registering Scalar BC for WaveCaKernel::pi!");
   
   }
   return;
@@ -225,25 +225,25 @@ extern "C" void SimpleWaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS)
 
 
 /* template for entries in parameter file:
-#$bound$#SimpleWaveCaKernel::phi_g_bound       = "skip"
-#$bound$#SimpleWaveCaKernel::phi_g_bound_speed = 1.0
-#$bound$#SimpleWaveCaKernel::phi_g_bound_limit = 0.0
-#$bound$#SimpleWaveCaKernel::phi_g_bound_scalar = 0.0
+#$bound$#WaveCaKernel::phi_g_bound       = "skip"
+#$bound$#WaveCaKernel::phi_g_bound_speed = 1.0
+#$bound$#WaveCaKernel::phi_g_bound_limit = 0.0
+#$bound$#WaveCaKernel::phi_g_bound_scalar = 0.0
 
-#$bound$#SimpleWaveCaKernel::pi_g_bound       = "skip"
-#$bound$#SimpleWaveCaKernel::pi_g_bound_speed = 1.0
-#$bound$#SimpleWaveCaKernel::pi_g_bound_limit = 0.0
-#$bound$#SimpleWaveCaKernel::pi_g_bound_scalar = 0.0
+#$bound$#WaveCaKernel::pi_g_bound       = "skip"
+#$bound$#WaveCaKernel::pi_g_bound_speed = 1.0
+#$bound$#WaveCaKernel::pi_g_bound_limit = 0.0
+#$bound$#WaveCaKernel::pi_g_bound_scalar = 0.0
 
-#$bound$#SimpleWaveCaKernel::phi_bound       = "skip"
-#$bound$#SimpleWaveCaKernel::phi_bound_speed = 1.0
-#$bound$#SimpleWaveCaKernel::phi_bound_limit = 0.0
-#$bound$#SimpleWaveCaKernel::phi_bound_scalar = 0.0
+#$bound$#WaveCaKernel::phi_bound       = "skip"
+#$bound$#WaveCaKernel::phi_bound_speed = 1.0
+#$bound$#WaveCaKernel::phi_bound_limit = 0.0
+#$bound$#WaveCaKernel::phi_bound_scalar = 0.0
 
-#$bound$#SimpleWaveCaKernel::pi_bound       = "skip"
-#$bound$#SimpleWaveCaKernel::pi_bound_speed = 1.0
-#$bound$#SimpleWaveCaKernel::pi_bound_limit = 0.0
-#$bound$#SimpleWaveCaKernel::pi_bound_scalar = 0.0
+#$bound$#WaveCaKernel::pi_bound       = "skip"
+#$bound$#WaveCaKernel::pi_bound_speed = 1.0
+#$bound$#WaveCaKernel::pi_bound_limit = 0.0
+#$bound$#WaveCaKernel::pi_bound_scalar = 0.0
 
 */
 
diff --git a/Examples/SimpleWaveCaKernel/src/CaKernel__calc_bound_rhs.code b/Examples/WaveCaKernel/src/CaKernel__calc_bound_rhs.code
index 863ad4f..863ad4f 100644
--- a/Examples/SimpleWaveCaKernel/src/CaKernel__calc_bound_rhs.code
+++ b/Examples/WaveCaKernel/src/CaKernel__calc_bound_rhs.code
diff --git a/Examples/SimpleWaveCaKernel/src/CaKernel__calc_rhs.code b/Examples/WaveCaKernel/src/CaKernel__calc_rhs.code
index d2d7fdb..d2d7fdb 100644
--- a/Examples/SimpleWaveCaKernel/src/CaKernel__calc_rhs.code
+++ b/Examples/WaveCaKernel/src/CaKernel__calc_rhs.code
diff --git a/Examples/SimpleWaveCaKernel/src/CaKernel__copy_to_device.code b/Examples/WaveCaKernel/src/CaKernel__copy_to_device.code
index 24b4566..24b4566 100644
--- a/Examples/SimpleWaveCaKernel/src/CaKernel__copy_to_device.code
+++ b/Examples/WaveCaKernel/src/CaKernel__copy_to_device.code
diff --git a/Examples/SimpleWaveCaKernel/src/Differencing.h b/Examples/WaveCaKernel/src/Differencing.h
index ef21b0b..ef21b0b 100644
--- a/Examples/SimpleWaveCaKernel/src/Differencing.h
+++ b/Examples/WaveCaKernel/src/Differencing.h
diff --git a/Examples/SimpleWaveCaKernel/src/RegisterMoL.cc b/Examples/WaveCaKernel/src/RegisterMoL.cc
index ac8c57b..938ce66 100644
--- a/Examples/SimpleWaveCaKernel/src/RegisterMoL.cc
+++ b/Examples/WaveCaKernel/src/RegisterMoL.cc
@@ -4,7 +4,7 @@
 #include "cctk_Arguments.h"
 #include "cctk_Parameters.h"
 
-extern "C" void SimpleWaveCaKernel_RegisterVars(CCTK_ARGUMENTS)
+extern "C" void WaveCaKernel_RegisterVars(CCTK_ARGUMENTS)
 {
   DECLARE_CCTK_ARGUMENTS;
   DECLARE_CCTK_PARAMETERS;
@@ -12,8 +12,8 @@ extern "C" void SimpleWaveCaKernel_RegisterVars(CCTK_ARGUMENTS)
   CCTK_INT ierr = 0;
   
   /* Register all the evolved grid functions with MoL */
-  ierr += MoLRegisterEvolved(CCTK_VarIndex("SimpleWaveCaKernel::phi"),  CCTK_VarIndex("SimpleWaveCaKernel::phirhs"));
-  ierr += MoLRegisterEvolved(CCTK_VarIndex("SimpleWaveCaKernel::pi"),  CCTK_VarIndex("SimpleWaveCaKernel::pirhs"));
+  ierr += MoLRegisterEvolved(CCTK_VarIndex("WaveCaKernel::phi"),  CCTK_VarIndex("WaveCaKernel::phirhs"));
+  ierr += MoLRegisterEvolved(CCTK_VarIndex("WaveCaKernel::pi"),  CCTK_VarIndex("WaveCaKernel::pirhs"));
   
   /* Register all the evolved Array functions with MoL */
   return;
diff --git a/Examples/SimpleWaveCaKernel/src/RegisterSymmetries.cc b/Examples/WaveCaKernel/src/RegisterSymmetries.cc
index c099f5b..7fd4a11 100644
--- a/Examples/SimpleWaveCaKernel/src/RegisterSymmetries.cc
+++ b/Examples/WaveCaKernel/src/RegisterSymmetries.cc
@@ -5,7 +5,7 @@
 #include "cctk_Parameters.h"
 #include "Symmetry.h"
 
-extern "C" void SimpleWaveCaKernel_RegisterSymmetries(CCTK_ARGUMENTS)
+extern "C" void WaveCaKernel_RegisterSymmetries(CCTK_ARGUMENTS)
 {
   DECLARE_CCTK_ARGUMENTS;
   DECLARE_CCTK_PARAMETERS;
@@ -19,16 +19,16 @@ extern "C" void SimpleWaveCaKernel_RegisterSymmetries(CCTK_ARGUMENTS)
   sym[0] = 1;
   sym[1] = 1;
   sym[2] = 1;
-  SetCartSymVN(cctkGH, sym, "SimpleWaveCaKernel::phi");
+  SetCartSymVN(cctkGH, sym, "WaveCaKernel::phi");
   
   sym[0] = 1;
   sym[1] = 1;
   sym[2] = 1;
-  SetCartSymVN(cctkGH, sym, "SimpleWaveCaKernel::pi");
+  SetCartSymVN(cctkGH, sym, "WaveCaKernel::pi");
   
   sym[0] = 1;
   sym[1] = 1;
   sym[2] = 1;
-  SetCartSymVN(cctkGH, sym, "SimpleWaveCaKernel::xCopy");
+  SetCartSymVN(cctkGH, sym, "WaveCaKernel::xCopy");
   
 }
diff --git a/Examples/SimpleWaveCaKernel/src/Startup.cc b/Examples/WaveCaKernel/src/Startup.cc
index 2761c8b..17ace69 100644
--- a/Examples/SimpleWaveCaKernel/src/Startup.cc
+++ b/Examples/WaveCaKernel/src/Startup.cc
@@ -2,9 +2,9 @@
 
 #include "cctk.h"
 
-extern "C" int SimpleWaveCaKernel_Startup(void)
+extern "C" int WaveCaKernel_Startup(void)
 {
-  const char * banner = "SimpleWaveCaKernel";
+  const char * banner = "WaveCaKernel";
   CCTK_RegisterBanner(banner);
   return 0;
 }
diff --git a/Examples/WaveCaKernel/src/initial_gaussian.cc b/Examples/WaveCaKernel/src/initial_gaussian.cc
new file mode 100644
index 0000000..dab0d62
--- /dev/null
+++ b/Examples/WaveCaKernel/src/initial_gaussian.cc
@@ -0,0 +1,135 @@
+/*  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"
+
+/* 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_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;
+  
+  
+  /* Include user-supplied include files */
+  
+  /* 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 p1o2dx = 0.5*INV(dx);
+  CCTK_REAL const p1o2dy = 0.5*INV(dy);
+  CCTK_REAL const p1o2dz = 0.5*INV(dz);
+  CCTK_REAL const p1odx2 = INV(SQR(dx));
+  CCTK_REAL const p1ody2 = INV(SQR(dy));
+  CCTK_REAL const p1odz2 = 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(initial_gaussian,
+    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];
+    
+    
+    /* Include user supplied include files */
+    
+    /* Precompute derivatives */
+    
+    /* Calculate temporaries and grid functions */
+    CCTK_REAL phiL = exp(-100.*SQR(xL + t));
+    
+    CCTK_REAL piL = -200.*(xL + t)*exp(-100.*SQR(xL + t));
+    
+    CCTK_REAL xCopyL = xL;
+    
+    /* Copy local copies back to grid functions */
+    phi[index] = phiL;
+    pi[index] = piL;
+    xCopy[index] = xCopyL;
+  }
+  CCTK_ENDLOOP3(initial_gaussian);
+}
+
+extern "C" void initial_gaussian(CCTK_ARGUMENTS)
+{
+  DECLARE_CCTK_ARGUMENTS;
+  DECLARE_CCTK_PARAMETERS;
+  
+  
+  if (verbose > 1)
+  {
+    CCTK_VInfo(CCTK_THORNSTRING,"Entering initial_gaussian_Body");
+  }
+  
+  if (cctk_iteration % initial_gaussian_calc_every != initial_gaussian_calc_offset)
+  {
+    return;
+  }
+  
+  const char *const groups[] = {
+    "grid::coordinates",
+    "WaveCaKernel::phi_g",
+    "WaveCaKernel::pi_g",
+    "WaveCaKernel::xCopy_g"};
+  GenericFD_AssertGroupStorage(cctkGH, "initial_gaussian", 4, groups);
+  
+  
+  GenericFD_LoopOverEverything(cctkGH, initial_gaussian_Body);
+  
+  if (verbose > 1)
+  {
+    CCTK_VInfo(CCTK_THORNSTRING,"Leaving initial_gaussian_Body");
+  }
+}
diff --git a/Examples/SimpleWaveCaKernel/src/make.code.defn b/Examples/WaveCaKernel/src/make.code.defn
index c43661b..c43661b 100644
--- a/Examples/SimpleWaveCaKernel/src/make.code.defn
+++ b/Examples/WaveCaKernel/src/make.code.defn
diff --git a/Examples/WaveCaKernel/test b/Examples/WaveCaKernel/test
new file mode 120000
index 0000000..35022b1
--- /dev/null
+++ b/Examples/WaveCaKernel/test
@@ -0,0 +1 @@
+../tests/WaveCaKernel/test
\ No newline at end of file
diff --git a/Examples/WaveHost/configuration.ccl b/Examples/WaveHost/configuration.ccl
new file mode 100644
index 0000000..0a66ec2
--- /dev/null
+++ b/Examples/WaveHost/configuration.ccl
@@ -0,0 +1,6 @@
+# File produced by Kranc
+
+REQUIRES GenericFD
+OPTIONAL LoopControl
+{
+}
diff --git a/Examples/WaveHost/interface.ccl b/Examples/WaveHost/interface.ccl
new file mode 100644
index 0000000..b08be61
--- /dev/null
+++ b/Examples/WaveHost/interface.ccl
@@ -0,0 +1,58 @@
+# File produced by Kranc
+
+implements: WaveHost
+
+inherits:   Grid GenericFD Boundary
+
+
+
+USES INCLUDE: GenericFD.h
+USES INCLUDE: Symmetry.h
+USES INCLUDE: sbp_calc_coeffs.h
+USES INCLUDE: Boundary.h
+USES INCLUDE: loopcontrol.h
+
+CCTK_INT FUNCTION MoLRegisterEvolved(CCTK_INT IN EvolvedIndex, CCTK_INT IN RHSIndex)
+USES FUNCTION MoLRegisterEvolved
+
+SUBROUTINE Diff_coeff(CCTK_POINTER_TO_CONST IN cctkGH, CCTK_INT IN dir, CCTK_INT IN nsize, CCTK_INT OUT ARRAY imin, CCTK_INT OUT ARRAY imax, CCTK_REAL OUT ARRAY q, CCTK_INT IN table_handle)
+USES FUNCTION Diff_coeff
+
+CCTK_INT FUNCTION MultiPatch_GetMap(CCTK_POINTER_TO_CONST IN cctkGH)
+USES FUNCTION MultiPatch_GetMap
+
+CCTK_INT FUNCTION Boundary_SelectGroupForBC(CCTK_POINTER_TO_CONST IN GH, CCTK_INT IN faces, CCTK_INT IN boundary_width, CCTK_INT IN table_handle, CCTK_STRING IN group_name, CCTK_STRING IN bc_name)
+USES FUNCTION Boundary_SelectGroupForBC
+
+CCTK_INT FUNCTION Boundary_SelectVarForBC(CCTK_POINTER_TO_CONST IN GH, CCTK_INT IN faces, CCTK_INT IN boundary_width, CCTK_INT IN table_handle, CCTK_STRING IN var_name, CCTK_STRING IN bc_name)
+USES FUNCTION Boundary_SelectVarForBC
+
+public:
+CCTK_REAL xCopy_g type=GF timelevels=1 tags=''
+{
+  xCopy
+} "xCopy_g"
+
+public:
+CCTK_REAL phi_g type=GF timelevels=2 tags=''
+{
+  phi
+} "phi_g"
+
+public:
+CCTK_REAL pi_g type=GF timelevels=2 tags=''
+{
+  pi
+} "pi_g"
+
+public:
+CCTK_REAL phi_grhs type=GF timelevels=2 tags=''
+{
+  phirhs
+} "phi_grhs"
+
+public:
+CCTK_REAL pi_grhs type=GF timelevels=2 tags=''
+{
+  pirhs
+} "pi_grhs"
diff --git a/Examples/WaveHost/param.ccl b/Examples/WaveHost/param.ccl
new file mode 100644
index 0000000..566e0c2
--- /dev/null
+++ b/Examples/WaveHost/param.ccl
@@ -0,0 +1,204 @@
+# File produced by Kranc
+
+
+shares: GenericFD
+
+
+
+shares: MethodOfLines
+
+USES CCTK_INT MoL_Num_Evolved_Vars
+USES CCTK_INT MoL_Num_ArrayEvolved_Vars
+
+restricted:
+CCTK_INT verbose "verbose" STEERABLE=ALWAYS
+{
+  *:* :: ""
+} 0
+
+restricted:
+CCTK_INT WaveHost_MaxNumEvolvedVars "Number of evolved variables used by this thorn" ACCUMULATOR-BASE=MethodofLines::MoL_Num_Evolved_Vars STEERABLE=RECOVER
+{
+  2:2 :: "Number of evolved variables used by this thorn"
+} 2
+
+restricted:
+CCTK_INT WaveHost_MaxNumArrayEvolvedVars "Number of Array evolved variables used by this thorn" ACCUMULATOR-BASE=MethodofLines::MoL_Num_ArrayEvolved_Vars STEERABLE=RECOVER
+{
+  0:0 :: "Number of Array evolved variables used by this thorn"
+} 0
+
+restricted:
+CCTK_INT timelevels "Number of active timelevels" STEERABLE=RECOVER
+{
+  0:2 :: ""
+} 2
+
+restricted:
+CCTK_INT rhs_timelevels "Number of active RHS timelevels" STEERABLE=RECOVER
+{
+  0:2 :: ""
+} 1
+
+restricted:
+CCTK_INT other_timelevels "Number of active timelevels for non-evolved grid functions" STEERABLE=RECOVER
+{
+  0:2 :: ""
+} 1
+
+restricted:
+CCTK_INT initial_gaussian_calc_every "initial_gaussian_calc_every" STEERABLE=ALWAYS
+{
+  *:* :: ""
+} 1
+
+restricted:
+CCTK_INT calc_rhs_calc_every "calc_rhs_calc_every" STEERABLE=ALWAYS
+{
+  *:* :: ""
+} 1
+
+restricted:
+CCTK_INT calc_bound_rhs_calc_every "calc_bound_rhs_calc_every" STEERABLE=ALWAYS
+{
+  *:* :: ""
+} 1
+
+restricted:
+CCTK_INT initial_gaussian_calc_offset "initial_gaussian_calc_offset" STEERABLE=ALWAYS
+{
+  *:* :: ""
+} 0
+
+restricted:
+CCTK_INT calc_rhs_calc_offset "calc_rhs_calc_offset" STEERABLE=ALWAYS
+{
+  *:* :: ""
+} 0
+
+restricted:
+CCTK_INT calc_bound_rhs_calc_offset "calc_bound_rhs_calc_offset" STEERABLE=ALWAYS
+{
+  *:* :: ""
+} 0
+
+private:
+KEYWORD phi_bound "Boundary condition to implement" STEERABLE=ALWAYS
+{
+  "flat" :: "Flat boundary condition"
+  "none" :: "No boundary condition"
+  "static" :: "Boundaries held fixed"
+  "radiative" :: "Radiation boundary condition"
+  "scalar" :: "Dirichlet boundary condition"
+  "newrad" :: "Improved radiative boundary condition"
+  "skip" :: "skip boundary condition code"
+} "skip"
+
+private:
+KEYWORD pi_bound "Boundary condition to implement" STEERABLE=ALWAYS
+{
+  "flat" :: "Flat boundary condition"
+  "none" :: "No boundary condition"
+  "static" :: "Boundaries held fixed"
+  "radiative" :: "Radiation boundary condition"
+  "scalar" :: "Dirichlet boundary condition"
+  "newrad" :: "Improved radiative boundary condition"
+  "skip" :: "skip boundary condition code"
+} "skip"
+
+private:
+KEYWORD phi_g_bound "Boundary condition to implement" STEERABLE=ALWAYS
+{
+  "flat" :: "Flat boundary condition"
+  "none" :: "No boundary condition"
+  "static" :: "Boundaries held fixed"
+  "radiative" :: "Radiation boundary condition"
+  "scalar" :: "Dirichlet boundary condition"
+  "newrad" :: "Improved radiative boundary condition"
+  "skip" :: "skip boundary condition code"
+} "none"
+
+private:
+KEYWORD pi_g_bound "Boundary condition to implement" STEERABLE=ALWAYS
+{
+  "flat" :: "Flat boundary condition"
+  "none" :: "No boundary condition"
+  "static" :: "Boundaries held fixed"
+  "radiative" :: "Radiation boundary condition"
+  "scalar" :: "Dirichlet boundary condition"
+  "newrad" :: "Improved radiative boundary condition"
+  "skip" :: "skip boundary condition code"
+} "none"
+
+private:
+CCTK_REAL phi_bound_speed "characteristic speed at boundary" STEERABLE=ALWAYS
+{
+  "0:*" :: "outgoing characteristic speed > 0"
+} 1.
+
+private:
+CCTK_REAL pi_bound_speed "characteristic speed at boundary" STEERABLE=ALWAYS
+{
+  "0:*" :: "outgoing characteristic speed > 0"
+} 1.
+
+private:
+CCTK_REAL phi_g_bound_speed "characteristic speed at boundary" STEERABLE=ALWAYS
+{
+  "0:*" :: "outgoing characteristic speed > 0"
+} 1.
+
+private:
+CCTK_REAL pi_g_bound_speed "characteristic speed at boundary" STEERABLE=ALWAYS
+{
+  "0:*" :: "outgoing characteristic speed > 0"
+} 1.
+
+private:
+CCTK_REAL phi_bound_limit "limit value for r -> infinity" STEERABLE=ALWAYS
+{
+  "*:*" :: "value of limit value is unrestricted"
+} 0.
+
+private:
+CCTK_REAL pi_bound_limit "limit value for r -> infinity" STEERABLE=ALWAYS
+{
+  "*:*" :: "value of limit value is unrestricted"
+} 0.
+
+private:
+CCTK_REAL phi_g_bound_limit "limit value for r -> infinity" STEERABLE=ALWAYS
+{
+  "*:*" :: "value of limit value is unrestricted"
+} 0.
+
+private:
+CCTK_REAL pi_g_bound_limit "limit value for r -> infinity" STEERABLE=ALWAYS
+{
+  "*:*" :: "value of limit value is unrestricted"
+} 0.
+
+private:
+CCTK_REAL phi_bound_scalar "Dirichlet boundary value" STEERABLE=ALWAYS
+{
+  "*:*" :: "unrestricted"
+} 0.
+
+private:
+CCTK_REAL pi_bound_scalar "Dirichlet boundary value" STEERABLE=ALWAYS
+{
+  "*:*" :: "unrestricted"
+} 0.
+
+private:
+CCTK_REAL phi_g_bound_scalar "Dirichlet boundary value" STEERABLE=ALWAYS
+{
+  "*:*" :: "unrestricted"
+} 0.
+
+private:
+CCTK_REAL pi_g_bound_scalar "Dirichlet boundary value" STEERABLE=ALWAYS
+{
+  "*:*" :: "unrestricted"
+} 0.
+
diff --git a/Examples/WaveHost/schedule.ccl b/Examples/WaveHost/schedule.ccl
new file mode 100644
index 0000000..b11880e
--- /dev/null
+++ b/Examples/WaveHost/schedule.ccl
@@ -0,0 +1,127 @@
+# File produced by Kranc
+
+
+if (other_timelevels == 1)
+{
+  STORAGE: xCopy_g[1]
+}
+
+if (timelevels == 1)
+{
+  STORAGE: phi_g[1]
+}
+if (timelevels == 2)
+{
+  STORAGE: phi_g[2]
+}
+
+if (timelevels == 1)
+{
+  STORAGE: pi_g[1]
+}
+if (timelevels == 2)
+{
+  STORAGE: pi_g[2]
+}
+
+if (rhs_timelevels == 1)
+{
+  STORAGE: phi_grhs[1]
+}
+if (rhs_timelevels == 2)
+{
+  STORAGE: phi_grhs[2]
+}
+
+if (rhs_timelevels == 1)
+{
+  STORAGE: pi_grhs[1]
+}
+if (rhs_timelevels == 2)
+{
+  STORAGE: pi_grhs[2]
+}
+
+schedule WaveHost_Startup at STARTUP
+{
+  LANG: C
+  OPTIONS: meta
+} "create banner"
+
+schedule WaveHost_RegisterSymmetries in SymmetryRegister
+{
+  LANG: C
+  OPTIONS: meta
+} "register symmetries"
+
+schedule initial_gaussian AT INITIAL
+{
+  LANG: C
+  READS:    grid::coordinates
+  WRITES:   WaveHost::phi_g
+  WRITES:   WaveHost::pi_g
+  WRITES:   WaveHost::xCopy_g
+} "initial_gaussian"
+
+schedule calc_rhs in MoL_CalcRHS
+{
+  LANG: C
+  READS:    WaveHost::phi_g
+  READS:    WaveHost::pi_g
+  WRITES:   WaveHost::phi_grhs
+  WRITES:   WaveHost::pi_grhs
+} "calc_rhs"
+
+schedule calc_rhs at ANALYSIS
+{
+  LANG: C
+  SYNC: phi_grhs
+  SYNC: pi_grhs
+  READS:    WaveHost::phi_g
+  READS:    WaveHost::pi_g
+  WRITES:   WaveHost::phi_grhs
+  WRITES:   WaveHost::pi_grhs
+} "calc_rhs"
+
+schedule calc_bound_rhs in MoL_RHSBoundaries
+{
+  LANG: C
+  READS:    WaveHost::xCopy_g
+  WRITES:   WaveHost::phi_grhs
+  WRITES:   WaveHost::pi_grhs
+} "calc_bound_rhs"
+
+schedule calc_bound_rhs at ANALYSIS
+{
+  LANG: C
+  SYNC: phi_grhs
+  SYNC: pi_grhs
+  READS:    WaveHost::xCopy_g
+  WRITES:   WaveHost::phi_grhs
+  WRITES:   WaveHost::pi_grhs
+} "calc_bound_rhs"
+
+schedule WaveHost_SelectBoundConds in MoL_PostStep
+{
+  LANG: C
+  OPTIONS: level
+  SYNC: phi_g
+  SYNC: pi_g
+} "select boundary conditions"
+
+schedule WaveHost_CheckBoundaries at BASEGRID
+{
+  LANG: C
+  OPTIONS: meta
+} "check boundaries treatment"
+
+schedule WaveHost_RegisterVars in MoL_Register
+{
+  LANG: C
+  OPTIONS: meta
+} "Register Variables for MoL"
+
+schedule group ApplyBCs as WaveHost_ApplyBCs in MoL_PostStep after WaveHost_SelectBoundConds
+{
+  # no language specified
+} "Apply boundary conditions controlled by thorn Boundary"
diff --git a/Examples/WaveHost/src/Boundaries.cc b/Examples/WaveHost/src/Boundaries.cc
new file mode 100644
index 0000000..8fe5d69
--- /dev/null
+++ b/Examples/WaveHost/src/Boundaries.cc
@@ -0,0 +1,249 @@
+/*  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 WaveHost_CheckBoundaries(CCTK_ARGUMENTS)
+{
+  DECLARE_CCTK_ARGUMENTS;
+  DECLARE_CCTK_PARAMETERS;
+  
+  return;
+}
+
+extern "C" void WaveHost_SelectBoundConds(CCTK_ARGUMENTS)
+{
+  DECLARE_CCTK_ARGUMENTS;
+  DECLARE_CCTK_PARAMETERS;
+  
+  CCTK_INT ierr = 0;
+  
+  if (CCTK_EQUALS(phi_g_bound, "none"  ) ||
+      CCTK_EQUALS(phi_g_bound, "static") ||
+      CCTK_EQUALS(phi_g_bound, "flat"  ) ||
+      CCTK_EQUALS(phi_g_bound, "zero"  ) )
+  {
+    ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, -1,
+                      "WaveHost::phi_g", phi_g_bound);
+    if (ierr < 0)
+       CCTK_WARN(0, "Failed to register phi_g_bound BC for WaveHost::phi_g!");
+  }
+  
+  if (CCTK_EQUALS(pi_g_bound, "none"  ) ||
+      CCTK_EQUALS(pi_g_bound, "static") ||
+      CCTK_EQUALS(pi_g_bound, "flat"  ) ||
+      CCTK_EQUALS(pi_g_bound, "zero"  ) )
+  {
+    ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, -1,
+                      "WaveHost::pi_g", pi_g_bound);
+    if (ierr < 0)
+       CCTK_WARN(0, "Failed to register pi_g_bound BC for WaveHost::pi_g!");
+  }
+  
+  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,
+                      "WaveHost::phi", phi_bound);
+    if (ierr < 0)
+       CCTK_WARN(0, "Failed to register phi_bound BC for WaveHost::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,
+                      "WaveHost::pi", pi_bound);
+    if (ierr < 0)
+       CCTK_WARN(0, "Failed to register pi_bound BC for WaveHost::pi!");
+  }
+  
+  if (CCTK_EQUALS(phi_g_bound, "radiative"))
+  {
+   /* select radiation boundary condition */
+    static CCTK_INT handle_phi_g_bound = -1;
+    if (handle_phi_g_bound < 0) handle_phi_g_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE);
+    if (handle_phi_g_bound < 0) CCTK_WARN(0, "could not create table!");
+    if (Util_TableSetReal(handle_phi_g_bound , phi_g_bound_limit, "LIMIT") < 0)
+       CCTK_WARN(0, "could not set LIMIT value in table!");
+    if (Util_TableSetReal(handle_phi_g_bound ,phi_g_bound_speed, "SPEED") < 0)
+       CCTK_WARN(0, "could not set SPEED value in table!");
+  
+    ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_phi_g_bound, 
+                      "WaveHost::phi_g", "Radiation");
+  
+    if (ierr < 0)
+       CCTK_WARN(0, "Failed to register Radiation BC for WaveHost::phi_g!");
+  
+  }
+  
+  if (CCTK_EQUALS(pi_g_bound, "radiative"))
+  {
+   /* select radiation boundary condition */
+    static CCTK_INT handle_pi_g_bound = -1;
+    if (handle_pi_g_bound < 0) handle_pi_g_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE);
+    if (handle_pi_g_bound < 0) CCTK_WARN(0, "could not create table!");
+    if (Util_TableSetReal(handle_pi_g_bound , pi_g_bound_limit, "LIMIT") < 0)
+       CCTK_WARN(0, "could not set LIMIT value in table!");
+    if (Util_TableSetReal(handle_pi_g_bound ,pi_g_bound_speed, "SPEED") < 0)
+       CCTK_WARN(0, "could not set SPEED value in table!");
+  
+    ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_pi_g_bound, 
+                      "WaveHost::pi_g", "Radiation");
+  
+    if (ierr < 0)
+       CCTK_WARN(0, "Failed to register Radiation BC for WaveHost::pi_g!");
+  
+  }
+  
+  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, 
+                      "WaveHost::phi", "Radiation");
+  
+    if (ierr < 0)
+       CCTK_WARN(0, "Failed to register Radiation BC for WaveHost::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, 
+                      "WaveHost::pi", "Radiation");
+  
+    if (ierr < 0)
+       CCTK_WARN(0, "Failed to register Radiation BC for WaveHost::pi!");
+  
+  }
+  
+  if (CCTK_EQUALS(phi_g_bound, "scalar"))
+  {
+   /* select scalar boundary condition */
+    static CCTK_INT handle_phi_g_bound = -1;
+    if (handle_phi_g_bound < 0) handle_phi_g_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE);
+    if (handle_phi_g_bound < 0) CCTK_WARN(0, "could not create table!");
+    if (Util_TableSetReal(handle_phi_g_bound ,phi_g_bound_scalar, "SCALAR") < 0)
+        CCTK_WARN(0, "could not set SCALAR value in table!");
+  
+    ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_phi_g_bound, 
+                      "WaveHost::phi_g", "scalar");
+  
+    if (ierr < 0)
+       CCTK_WARN(0, "Failed to register Scalar BC for WaveHost::phi_g!");
+  
+  }
+  
+  if (CCTK_EQUALS(pi_g_bound, "scalar"))
+  {
+   /* select scalar boundary condition */
+    static CCTK_INT handle_pi_g_bound = -1;
+    if (handle_pi_g_bound < 0) handle_pi_g_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE);
+    if (handle_pi_g_bound < 0) CCTK_WARN(0, "could not create table!");
+    if (Util_TableSetReal(handle_pi_g_bound ,pi_g_bound_scalar, "SCALAR") < 0)
+        CCTK_WARN(0, "could not set SCALAR value in table!");
+  
+    ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_pi_g_bound, 
+                      "WaveHost::pi_g", "scalar");
+  
+    if (ierr < 0)
+       CCTK_WARN(0, "Failed to register Scalar BC for WaveHost::pi_g!");
+  
+  }
+  
+  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, 
+                      "WaveHost::phi", "scalar");
+  
+    if (ierr < 0)
+       CCTK_WARN(0, "Error in registering Scalar BC for WaveHost::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, 
+                      "WaveHost::pi", "scalar");
+  
+    if (ierr < 0)
+       CCTK_WARN(0, "Error in registering Scalar BC for WaveHost::pi!");
+  
+  }
+  return;
+}
+
+
+
+/* template for entries in parameter file:
+#$bound$#WaveHost::phi_g_bound       = "skip"
+#$bound$#WaveHost::phi_g_bound_speed = 1.0
+#$bound$#WaveHost::phi_g_bound_limit = 0.0
+#$bound$#WaveHost::phi_g_bound_scalar = 0.0
+
+#$bound$#WaveHost::pi_g_bound       = "skip"
+#$bound$#WaveHost::pi_g_bound_speed = 1.0
+#$bound$#WaveHost::pi_g_bound_limit = 0.0
+#$bound$#WaveHost::pi_g_bound_scalar = 0.0
+
+#$bound$#WaveHost::phi_bound       = "skip"
+#$bound$#WaveHost::phi_bound_speed = 1.0
+#$bound$#WaveHost::phi_bound_limit = 0.0
+#$bound$#WaveHost::phi_bound_scalar = 0.0
+
+#$bound$#WaveHost::pi_bound       = "skip"
+#$bound$#WaveHost::pi_bound_speed = 1.0
+#$bound$#WaveHost::pi_bound_limit = 0.0
+#$bound$#WaveHost::pi_bound_scalar = 0.0
+
+*/
+
diff --git a/Examples/WaveHost/src/CaKernel__copy_to_device.code b/Examples/WaveHost/src/CaKernel__copy_to_device.code
new file mode 100644
index 0000000..24b4566
--- /dev/null
+++ b/Examples/WaveHost/src/CaKernel__copy_to_device.code
@@ -0,0 +1,77 @@
+#undef KRANC_DIFF_FUNCTIONS
+#define KRANC_C
+#include "Differencing.h"
+#include "GenericFD.h"
+
+#undef KRANC_GFOFFSET3D
+#define KRANC_GFOFFSET3D(u,i,j,k) I3D(u,i,j,k)
+
+
+/* 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))
+
+CAKERNEL_copy_to_device_Begin
+  
+  /* Include user-supplied include files */
+  
+  /* Initialise finite differencing variables */
+  CCTK_REAL const dx = params.cagh_dx;
+  CCTK_REAL const dy = params.cagh_dy;
+  CCTK_REAL const dz = params.cagh_dz;
+  CCTK_REAL const dt = params.cagh_dt;
+  CCTK_REAL const t = params.cagh_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 p1o2dx = 0.5*INV(dx);
+  CCTK_REAL const p1o2dy = 0.5*INV(dy);
+  CCTK_REAL const p1o2dz = 0.5*INV(dz);
+  CCTK_REAL const p1odx2 = INV(SQR(dx));
+  CCTK_REAL const p1ody2 = INV(SQR(dy));
+  CCTK_REAL const p1odz2 = INV(SQR(dz));
+  
+  /* Assign local copies of arrays functions */
+  
+  
+  
+  /* Calculate temporaries and arrays functions */
+  
+  /* Copy local copies back to grid functions */
+  CAKERNEL_copy_to_device_Computations_Begin
+    
+    /* Assign local copies of grid functions */
+    
+    CCTK_REAL phiL = I3D(phi,0,0,0);
+    CCTK_REAL piL = I3D(pi,0,0,0);
+    
+    
+    /* Include user supplied include files */
+    
+    /* Precompute derivatives */
+    
+    /* Calculate temporaries and grid functions */
+    phiL = phiL;
+    
+    piL = piL;
+    
+    /* Copy local copies back to grid functions */
+    I3D(phi,0,0,0) = phiL;
+    I3D(pi,0,0,0) = piL;
+    
+  CAKERNEL_copy_to_device_Computations_End
+  
+CAKERNEL_copy_to_device_End
diff --git a/Examples/WaveHost/src/Differencing.h b/Examples/WaveHost/src/Differencing.h
new file mode 100644
index 0000000..ef21b0b
--- /dev/null
+++ b/Examples/WaveHost/src/Differencing.h
@@ -0,0 +1,72 @@
+#ifndef KRANC_DIFF_FUNCTIONS
+#  define PDstandard2nd1(u) ((-KRANC_GFOFFSET3D(u,-1,0,0) + KRANC_GFOFFSET3D(u,1,0,0))*p1o2dx)
+#else
+#  define PDstandard2nd1(u) (PDstandard2nd1_impl(u,p1o2dx,cdj,cdk))
+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;
+static CCTK_REAL PDstandard2nd1_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dx, ptrdiff_t const cdj, ptrdiff_t const cdk)
+{
+  ptrdiff_t const cdi=sizeof(CCTK_REAL);
+  return (-KRANC_GFOFFSET3D(u,-1,0,0) + KRANC_GFOFFSET3D(u,1,0,0))*p1o2dx;
+}
+#endif
+
+#ifndef KRANC_DIFF_FUNCTIONS
+#  define PDstandard2nd2(u) ((-KRANC_GFOFFSET3D(u,0,-1,0) + KRANC_GFOFFSET3D(u,0,1,0))*p1o2dy)
+#else
+#  define PDstandard2nd2(u) (PDstandard2nd2_impl(u,p1o2dy,cdj,cdk))
+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;
+static CCTK_REAL PDstandard2nd2_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dy, ptrdiff_t const cdj, ptrdiff_t const cdk)
+{
+  ptrdiff_t const cdi=sizeof(CCTK_REAL);
+  return (-KRANC_GFOFFSET3D(u,0,-1,0) + KRANC_GFOFFSET3D(u,0,1,0))*p1o2dy;
+}
+#endif
+
+#ifndef KRANC_DIFF_FUNCTIONS
+#  define PDstandard2nd3(u) ((-KRANC_GFOFFSET3D(u,0,0,-1) + KRANC_GFOFFSET3D(u,0,0,1))*p1o2dz)
+#else
+#  define PDstandard2nd3(u) (PDstandard2nd3_impl(u,p1o2dz,cdj,cdk))
+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;
+static CCTK_REAL PDstandard2nd3_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dz, ptrdiff_t const cdj, ptrdiff_t const cdk)
+{
+  ptrdiff_t const cdi=sizeof(CCTK_REAL);
+  return (-KRANC_GFOFFSET3D(u,0,0,-1) + KRANC_GFOFFSET3D(u,0,0,1))*p1o2dz;
+}
+#endif
+
+#ifndef KRANC_DIFF_FUNCTIONS
+#  define PDstandard2nd11(u) ((-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,-1,0,0) + KRANC_GFOFFSET3D(u,1,0,0))*p1odx2)
+#else
+#  define PDstandard2nd11(u) (PDstandard2nd11_impl(u,p1odx2,cdj,cdk))
+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;
+static CCTK_REAL PDstandard2nd11_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1odx2, ptrdiff_t const cdj, ptrdiff_t const cdk)
+{
+  ptrdiff_t const cdi=sizeof(CCTK_REAL);
+  return (-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,-1,0,0) + KRANC_GFOFFSET3D(u,1,0,0))*p1odx2;
+}
+#endif
+
+#ifndef KRANC_DIFF_FUNCTIONS
+#  define PDstandard2nd22(u) ((-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,0,-1,0) + KRANC_GFOFFSET3D(u,0,1,0))*p1ody2)
+#else
+#  define PDstandard2nd22(u) (PDstandard2nd22_impl(u,p1ody2,cdj,cdk))
+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;
+static CCTK_REAL PDstandard2nd22_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1ody2, ptrdiff_t const cdj, ptrdiff_t const cdk)
+{
+  ptrdiff_t const cdi=sizeof(CCTK_REAL);
+  return (-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,0,-1,0) + KRANC_GFOFFSET3D(u,0,1,0))*p1ody2;
+}
+#endif
+
+#ifndef KRANC_DIFF_FUNCTIONS
+#  define PDstandard2nd33(u) ((-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,0,0,-1) + KRANC_GFOFFSET3D(u,0,0,1))*p1odz2)
+#else
+#  define PDstandard2nd33(u) (PDstandard2nd33_impl(u,p1odz2,cdj,cdk))
+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;
+static CCTK_REAL PDstandard2nd33_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1odz2, ptrdiff_t const cdj, ptrdiff_t const cdk)
+{
+  ptrdiff_t const cdi=sizeof(CCTK_REAL);
+  return (-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,0,0,-1) + KRANC_GFOFFSET3D(u,0,0,1))*p1odz2;
+}
+#endif
+
diff --git a/Examples/WaveHost/src/RegisterMoL.cc b/Examples/WaveHost/src/RegisterMoL.cc
new file mode 100644
index 0000000..0e263ab
--- /dev/null
+++ b/Examples/WaveHost/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 WaveHost_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("WaveHost::phi"),  CCTK_VarIndex("WaveHost::phirhs"));
+  ierr += MoLRegisterEvolved(CCTK_VarIndex("WaveHost::pi"),  CCTK_VarIndex("WaveHost::pirhs"));
+  
+  /* Register all the evolved Array functions with MoL */
+  return;
+}
diff --git a/Examples/WaveHost/src/RegisterSymmetries.cc b/Examples/WaveHost/src/RegisterSymmetries.cc
new file mode 100644
index 0000000..7e44b14
--- /dev/null
+++ b/Examples/WaveHost/src/RegisterSymmetries.cc
@@ -0,0 +1,34 @@
+/*  File produced by Kranc */
+
+#include "cctk.h"
+#include "cctk_Arguments.h"
+#include "cctk_Parameters.h"
+#include "Symmetry.h"
+
+extern "C" void WaveHost_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, "WaveHost::phi");
+  
+  sym[0] = 1;
+  sym[1] = 1;
+  sym[2] = 1;
+  SetCartSymVN(cctkGH, sym, "WaveHost::pi");
+  
+  sym[0] = 1;
+  sym[1] = 1;
+  sym[2] = 1;
+  SetCartSymVN(cctkGH, sym, "WaveHost::xCopy");
+  
+}
diff --git a/Examples/WaveHost/src/Startup.cc b/Examples/WaveHost/src/Startup.cc
new file mode 100644
index 0000000..a779f90
--- /dev/null
+++ b/Examples/WaveHost/src/Startup.cc
@@ -0,0 +1,10 @@
+/*  File produced by Kranc */
+
+#include "cctk.h"
+
+extern "C" int WaveHost_Startup(void)
+{
+  const char * banner = "WaveHost";
+  CCTK_RegisterBanner(banner);
+  return 0;
+}
diff --git a/Examples/WaveHost/src/calc_bound_rhs.cc b/Examples/WaveHost/src/calc_bound_rhs.cc
new file mode 100644
index 0000000..7098ac5
--- /dev/null
+++ b/Examples/WaveHost/src/calc_bound_rhs.cc
@@ -0,0 +1,147 @@
+/*  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"
+
+/* 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_bound_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 */, "WaveHost::phi_grhs","flat");
+  if (ierr < 0)
+    CCTK_WARN(1, "Failed to register flat BC for WaveHost::phi_grhs.");
+  ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "WaveHost::pi_grhs","flat");
+  if (ierr < 0)
+    CCTK_WARN(1, "Failed to register flat BC for WaveHost::pi_grhs.");
+  return;
+}
+
+static void calc_bound_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;
+  
+  
+  /* Include user-supplied include files */
+  
+  /* 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 p1o2dx = 0.5*INV(dx);
+  CCTK_REAL const p1o2dy = 0.5*INV(dy);
+  CCTK_REAL const p1o2dz = 0.5*INV(dz);
+  CCTK_REAL const p1odx2 = INV(SQR(dx));
+  CCTK_REAL const p1ody2 = INV(SQR(dy));
+  CCTK_REAL const p1odz2 = 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(calc_bound_rhs,
+    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 xCopyL = xCopy[index];
+    
+    
+    /* Include user supplied include files */
+    
+    /* Precompute derivatives */
+    
+    /* Calculate temporaries and grid functions */
+    CCTK_REAL phirhsL = -200.*(xCopyL + t)*exp(-100.*SQR(xCopyL + t));
+    
+    CCTK_REAL pirhsL = exp(-100.*SQR(xCopyL + t))*(-200. + 
+      80000.*xCopyL*t + 40000.*(SQR(xCopyL) + SQR(t)));
+    
+    /* Copy local copies back to grid functions */
+    phirhs[index] = phirhsL;
+    pirhs[index] = pirhsL;
+  }
+  CCTK_ENDLOOP3(calc_bound_rhs);
+}
+
+extern "C" void calc_bound_rhs(CCTK_ARGUMENTS)
+{
+  DECLARE_CCTK_ARGUMENTS;
+  DECLARE_CCTK_PARAMETERS;
+  
+  
+  if (verbose > 1)
+  {
+    CCTK_VInfo(CCTK_THORNSTRING,"Entering calc_bound_rhs_Body");
+  }
+  
+  if (cctk_iteration % calc_bound_rhs_calc_every != calc_bound_rhs_calc_offset)
+  {
+    return;
+  }
+  
+  const char *const groups[] = {
+    "WaveHost::phi_grhs",
+    "WaveHost::pi_grhs",
+    "WaveHost::xCopy_g"};
+  GenericFD_AssertGroupStorage(cctkGH, "calc_bound_rhs", 3, groups);
+  
+  
+  GenericFD_LoopOverBoundary(cctkGH, calc_bound_rhs_Body);
+  
+  if (verbose > 1)
+  {
+    CCTK_VInfo(CCTK_THORNSTRING,"Leaving calc_bound_rhs_Body");
+  }
+}
diff --git a/Examples/WaveHost/src/calc_rhs.cc b/Examples/WaveHost/src/calc_rhs.cc
new file mode 100644
index 0000000..ee114e7
--- /dev/null
+++ b/Examples/WaveHost/src/calc_rhs.cc
@@ -0,0 +1,153 @@
+/*  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"
+
+/* 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 */, "WaveHost::phi_grhs","flat");
+  if (ierr < 0)
+    CCTK_WARN(1, "Failed to register flat BC for WaveHost::phi_grhs.");
+  ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "WaveHost::pi_grhs","flat");
+  if (ierr < 0)
+    CCTK_WARN(1, "Failed to register flat BC for WaveHost::pi_grhs.");
+  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;
+  
+  
+  /* Include user-supplied include files */
+  
+  /* 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 p1o2dx = 0.5*INV(dx);
+  CCTK_REAL const p1o2dy = 0.5*INV(dy);
+  CCTK_REAL const p1o2dz = 0.5*INV(dz);
+  CCTK_REAL const p1odx2 = INV(SQR(dx));
+  CCTK_REAL const p1ody2 = INV(SQR(dy));
+  CCTK_REAL const p1odz2 = 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(calc_rhs,
+    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 phiL = phi[index];
+    CCTK_REAL piL = pi[index];
+    
+    
+    /* Include user supplied include files */
+    
+    /* Precompute derivatives */
+    CCTK_REAL const PDstandard2nd11phi = PDstandard2nd11(&phi[index]);
+    CCTK_REAL const PDstandard2nd22phi = PDstandard2nd22(&phi[index]);
+    CCTK_REAL const PDstandard2nd33phi = PDstandard2nd33(&phi[index]);
+    
+    /* Calculate temporaries and grid functions */
+    CCTK_REAL phirhsL = piL;
+    
+    CCTK_REAL pirhsL = PDstandard2nd11phi + PDstandard2nd22phi + 
+      PDstandard2nd33phi;
+    
+    /* Copy local copies back to grid functions */
+    phirhs[index] = phirhsL;
+    pirhs[index] = pirhsL;
+  }
+  CCTK_ENDLOOP3(calc_rhs);
+}
+
+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 *const groups[] = {
+    "WaveHost::phi_g",
+    "WaveHost::phi_grhs",
+    "WaveHost::pi_g",
+    "WaveHost::pi_grhs"};
+  GenericFD_AssertGroupStorage(cctkGH, "calc_rhs", 4, 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/SimpleWaveCaKernel/src/initial_gaussian.cc b/Examples/WaveHost/src/initial_gaussian.cc
index 55bd228..2831936 100644
--- a/Examples/SimpleWaveCaKernel/src/initial_gaussian.cc
+++ b/Examples/WaveHost/src/initial_gaussian.cc
@@ -118,11 +118,15 @@ extern "C" void initial_gaussian(CCTK_ARGUMENTS)
     return;
   }
   
-  const char *groups[] = {"grid::coordinates","SimpleWaveCaKernel::phi_g","SimpleWaveCaKernel::pi_g","SimpleWaveCaKernel::xCopy_g"};
+  const char *const groups[] = {
+    "grid::coordinates",
+    "WaveHost::phi_g",
+    "WaveHost::pi_g",
+    "WaveHost::xCopy_g"};
   GenericFD_AssertGroupStorage(cctkGH, "initial_gaussian", 4, groups);
   
   
-  GenericFD_LoopOverEverything(cctkGH, &initial_gaussian_Body);
+  GenericFD_LoopOverEverything(cctkGH, initial_gaussian_Body);
   
   if (verbose > 1)
   {
diff --git a/Examples/WaveHost/src/make.code.defn b/Examples/WaveHost/src/make.code.defn
new file mode 100644
index 0000000..f30a323
--- /dev/null
+++ b/Examples/WaveHost/src/make.code.defn
@@ -0,0 +1,3 @@
+# File produced by Kranc
+
+SRCS = Startup.cc RegisterSymmetries.cc RegisterMoL.cc initial_gaussian.cc calc_rhs.cc calc_bound_rhs.cc Boundaries.cc 
diff --git a/Examples/WaveHost/test b/Examples/WaveHost/test
new file mode 120000
index 0000000..bbc45b0
--- /dev/null
+++ b/Examples/WaveHost/test
@@ -0,0 +1 @@
+../tests/WaveHost/test
\ No newline at end of file
diff --git a/Examples/simplewavecakernel_gaussian.par b/Examples/wavecakernel_gaussian.par
index 9ffeeae..9ffeeae 100644
--- a/Examples/simplewavecakernel_gaussian.par
+++ b/Examples/wavecakernel_gaussian.par