aboutsummaryrefslogtreecommitdiff
path: root/Examples
diff options
context:
space:
mode:
authorIan Hinder <ian.hinder@aei.mpg.de>2012-01-31 17:50:27 -0600
committerIan Hinder <ian.hinder@aei.mpg.de>2012-01-31 17:50:27 -0600
commitc403011488f0201749428007142445a2e98a7447 (patch)
tree8631ef9dd6c29e98704f600e362a7a80ce40f5b7 /Examples
parent7cf69990ac3891c0e5258a7010952ab1fcf763d5 (diff)
Regenerate EMScript
Diffstat (limited to 'Examples')
-rw-r--r--Examples/EMScript/configuration.ccl6
-rw-r--r--Examples/EMScript/interface.ccl6
-rw-r--r--Examples/EMScript/param.ccl3
-rw-r--r--Examples/EMScript/src/Boundaries.cc160
-rw-r--r--Examples/EMScript/src/Differencing.h579
-rw-r--r--Examples/EMScript/src/EM_constraints.cc397
-rw-r--r--Examples/EMScript/src/EM_energy.cc263
-rw-r--r--Examples/EMScript/src/EM_evol.cc473
-rw-r--r--Examples/EMScript/src/EM_initial.cc298
-rw-r--r--Examples/EMScript/src/RegisterMoL.cc12
-rw-r--r--Examples/EMScript/src/RegisterSymmetries.cc18
11 files changed, 1339 insertions, 876 deletions
diff --git a/Examples/EMScript/configuration.ccl b/Examples/EMScript/configuration.ccl
index 0a66ec2..585220a 100644
--- a/Examples/EMScript/configuration.ccl
+++ b/Examples/EMScript/configuration.ccl
@@ -1,6 +1,6 @@
# File produced by Kranc
REQUIRES GenericFD
-OPTIONAL LoopControl
-{
-}
+REQUIRES LoopControl
+REQUIRES OpenCL OpenCLRunTime
+REQUIRES Vectors
diff --git a/Examples/EMScript/interface.ccl b/Examples/EMScript/interface.ccl
index 8537b0d..569d68b 100644
--- a/Examples/EMScript/interface.ccl
+++ b/Examples/EMScript/interface.ccl
@@ -1,8 +1,8 @@
# File produced by Kranc
-implements: EMScript
+implements: My_New_Implementation
-inherits: Grid GenericFD Boundary
+inherits: ADMBase Grid GenericFD Boundary
@@ -11,6 +11,8 @@ USES INCLUDE: Symmetry.h
USES INCLUDE: sbp_calc_coeffs.h
USES INCLUDE: Boundary.h
USES INCLUDE: loopcontrol.h
+USES INCLUDE: OpenCLRunTime.h
+USES INCLUDE: vectors.h
CCTK_INT FUNCTION MoLRegisterEvolved(CCTK_INT IN EvolvedIndex, CCTK_INT IN RHSIndex)
USES FUNCTION MoLRegisterEvolved
diff --git a/Examples/EMScript/param.ccl b/Examples/EMScript/param.ccl
index a83e18c..5f97309 100644
--- a/Examples/EMScript/param.ccl
+++ b/Examples/EMScript/param.ccl
@@ -3,6 +3,9 @@
shares: GenericFD
+USES CCTK_STRING jacobian_group
+USES CCTK_STRING jacobian_derivative_group
+USES CCTK_INT jacobian_identity_map
shares: MethodOfLines
diff --git a/Examples/EMScript/src/Boundaries.cc b/Examples/EMScript/src/Boundaries.cc
index 2d15538..b091436 100644
--- a/Examples/EMScript/src/Boundaries.cc
+++ b/Examples/EMScript/src/Boundaries.cc
@@ -38,9 +38,9 @@ extern "C" void EMScript_SelectBoundConds(CCTK_ARGUMENTS)
CCTK_EQUALS(B_group_bound, "zero" ) )
{
ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, -1,
- "EMScript::B_group", B_group_bound);
+ "My_New_Implementation::B_group", B_group_bound);
if (ierr < 0)
- CCTK_WARN(0, "Failed to register B_group_bound BC for EMScript::B_group!");
+ CCTK_WARN(0, "Failed to register B_group_bound BC for My_New_Implementation::B_group!");
}
if (CCTK_EQUALS(El_group_bound, "none" ) ||
@@ -49,9 +49,9 @@ extern "C" void EMScript_SelectBoundConds(CCTK_ARGUMENTS)
CCTK_EQUALS(El_group_bound, "zero" ) )
{
ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, -1,
- "EMScript::El_group", El_group_bound);
+ "My_New_Implementation::El_group", El_group_bound);
if (ierr < 0)
- CCTK_WARN(0, "Failed to register El_group_bound BC for EMScript::El_group!");
+ CCTK_WARN(0, "Failed to register El_group_bound BC for My_New_Implementation::El_group!");
}
if (CCTK_EQUALS(B1_bound, "none" ) ||
@@ -60,9 +60,9 @@ extern "C" void EMScript_SelectBoundConds(CCTK_ARGUMENTS)
CCTK_EQUALS(B1_bound, "zero" ) )
{
ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, -1,
- "EMScript::B1", B1_bound);
+ "My_New_Implementation::B1", B1_bound);
if (ierr < 0)
- CCTK_WARN(0, "Failed to register B1_bound BC for EMScript::B1!");
+ CCTK_WARN(0, "Failed to register B1_bound BC for My_New_Implementation::B1!");
}
if (CCTK_EQUALS(B2_bound, "none" ) ||
@@ -71,9 +71,9 @@ extern "C" void EMScript_SelectBoundConds(CCTK_ARGUMENTS)
CCTK_EQUALS(B2_bound, "zero" ) )
{
ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, -1,
- "EMScript::B2", B2_bound);
+ "My_New_Implementation::B2", B2_bound);
if (ierr < 0)
- CCTK_WARN(0, "Failed to register B2_bound BC for EMScript::B2!");
+ CCTK_WARN(0, "Failed to register B2_bound BC for My_New_Implementation::B2!");
}
if (CCTK_EQUALS(B3_bound, "none" ) ||
@@ -82,9 +82,9 @@ extern "C" void EMScript_SelectBoundConds(CCTK_ARGUMENTS)
CCTK_EQUALS(B3_bound, "zero" ) )
{
ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, -1,
- "EMScript::B3", B3_bound);
+ "My_New_Implementation::B3", B3_bound);
if (ierr < 0)
- CCTK_WARN(0, "Failed to register B3_bound BC for EMScript::B3!");
+ CCTK_WARN(0, "Failed to register B3_bound BC for My_New_Implementation::B3!");
}
if (CCTK_EQUALS(El1_bound, "none" ) ||
@@ -93,9 +93,9 @@ extern "C" void EMScript_SelectBoundConds(CCTK_ARGUMENTS)
CCTK_EQUALS(El1_bound, "zero" ) )
{
ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, -1,
- "EMScript::El1", El1_bound);
+ "My_New_Implementation::El1", El1_bound);
if (ierr < 0)
- CCTK_WARN(0, "Failed to register El1_bound BC for EMScript::El1!");
+ CCTK_WARN(0, "Failed to register El1_bound BC for My_New_Implementation::El1!");
}
if (CCTK_EQUALS(El2_bound, "none" ) ||
@@ -104,9 +104,9 @@ extern "C" void EMScript_SelectBoundConds(CCTK_ARGUMENTS)
CCTK_EQUALS(El2_bound, "zero" ) )
{
ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, -1,
- "EMScript::El2", El2_bound);
+ "My_New_Implementation::El2", El2_bound);
if (ierr < 0)
- CCTK_WARN(0, "Failed to register El2_bound BC for EMScript::El2!");
+ CCTK_WARN(0, "Failed to register El2_bound BC for My_New_Implementation::El2!");
}
if (CCTK_EQUALS(El3_bound, "none" ) ||
@@ -115,9 +115,9 @@ extern "C" void EMScript_SelectBoundConds(CCTK_ARGUMENTS)
CCTK_EQUALS(El3_bound, "zero" ) )
{
ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, -1,
- "EMScript::El3", El3_bound);
+ "My_New_Implementation::El3", El3_bound);
if (ierr < 0)
- CCTK_WARN(0, "Failed to register El3_bound BC for EMScript::El3!");
+ CCTK_WARN(0, "Failed to register El3_bound BC for My_New_Implementation::El3!");
}
if (CCTK_EQUALS(B_group_bound, "radiative"))
@@ -132,10 +132,10 @@ extern "C" void EMScript_SelectBoundConds(CCTK_ARGUMENTS)
CCTK_WARN(0, "could not set SPEED value in table!");
ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_B_group_bound,
- "EMScript::B_group", "Radiation");
+ "My_New_Implementation::B_group", "Radiation");
if (ierr < 0)
- CCTK_WARN(0, "Failed to register Radiation BC for EMScript::B_group!");
+ CCTK_WARN(0, "Failed to register Radiation BC for My_New_Implementation::B_group!");
}
@@ -151,10 +151,10 @@ extern "C" void EMScript_SelectBoundConds(CCTK_ARGUMENTS)
CCTK_WARN(0, "could not set SPEED value in table!");
ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_El_group_bound,
- "EMScript::El_group", "Radiation");
+ "My_New_Implementation::El_group", "Radiation");
if (ierr < 0)
- CCTK_WARN(0, "Failed to register Radiation BC for EMScript::El_group!");
+ CCTK_WARN(0, "Failed to register Radiation BC for My_New_Implementation::El_group!");
}
@@ -170,10 +170,10 @@ extern "C" void EMScript_SelectBoundConds(CCTK_ARGUMENTS)
CCTK_WARN(0, "could not set SPEED value in table!");
ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_B1_bound,
- "EMScript::B1", "Radiation");
+ "My_New_Implementation::B1", "Radiation");
if (ierr < 0)
- CCTK_WARN(0, "Failed to register Radiation BC for EMScript::B1!");
+ CCTK_WARN(0, "Failed to register Radiation BC for My_New_Implementation::B1!");
}
@@ -189,10 +189,10 @@ extern "C" void EMScript_SelectBoundConds(CCTK_ARGUMENTS)
CCTK_WARN(0, "could not set SPEED value in table!");
ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_B2_bound,
- "EMScript::B2", "Radiation");
+ "My_New_Implementation::B2", "Radiation");
if (ierr < 0)
- CCTK_WARN(0, "Failed to register Radiation BC for EMScript::B2!");
+ CCTK_WARN(0, "Failed to register Radiation BC for My_New_Implementation::B2!");
}
@@ -208,10 +208,10 @@ extern "C" void EMScript_SelectBoundConds(CCTK_ARGUMENTS)
CCTK_WARN(0, "could not set SPEED value in table!");
ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_B3_bound,
- "EMScript::B3", "Radiation");
+ "My_New_Implementation::B3", "Radiation");
if (ierr < 0)
- CCTK_WARN(0, "Failed to register Radiation BC for EMScript::B3!");
+ CCTK_WARN(0, "Failed to register Radiation BC for My_New_Implementation::B3!");
}
@@ -227,10 +227,10 @@ extern "C" void EMScript_SelectBoundConds(CCTK_ARGUMENTS)
CCTK_WARN(0, "could not set SPEED value in table!");
ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_El1_bound,
- "EMScript::El1", "Radiation");
+ "My_New_Implementation::El1", "Radiation");
if (ierr < 0)
- CCTK_WARN(0, "Failed to register Radiation BC for EMScript::El1!");
+ CCTK_WARN(0, "Failed to register Radiation BC for My_New_Implementation::El1!");
}
@@ -246,10 +246,10 @@ extern "C" void EMScript_SelectBoundConds(CCTK_ARGUMENTS)
CCTK_WARN(0, "could not set SPEED value in table!");
ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_El2_bound,
- "EMScript::El2", "Radiation");
+ "My_New_Implementation::El2", "Radiation");
if (ierr < 0)
- CCTK_WARN(0, "Failed to register Radiation BC for EMScript::El2!");
+ CCTK_WARN(0, "Failed to register Radiation BC for My_New_Implementation::El2!");
}
@@ -265,10 +265,10 @@ extern "C" void EMScript_SelectBoundConds(CCTK_ARGUMENTS)
CCTK_WARN(0, "could not set SPEED value in table!");
ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_El3_bound,
- "EMScript::El3", "Radiation");
+ "My_New_Implementation::El3", "Radiation");
if (ierr < 0)
- CCTK_WARN(0, "Failed to register Radiation BC for EMScript::El3!");
+ CCTK_WARN(0, "Failed to register Radiation BC for My_New_Implementation::El3!");
}
@@ -282,10 +282,10 @@ extern "C" void EMScript_SelectBoundConds(CCTK_ARGUMENTS)
CCTK_WARN(0, "could not set SCALAR value in table!");
ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_B_group_bound,
- "EMScript::B_group", "scalar");
+ "My_New_Implementation::B_group", "scalar");
if (ierr < 0)
- CCTK_WARN(0, "Failed to register Scalar BC for EMScript::B_group!");
+ CCTK_WARN(0, "Failed to register Scalar BC for My_New_Implementation::B_group!");
}
@@ -299,10 +299,10 @@ extern "C" void EMScript_SelectBoundConds(CCTK_ARGUMENTS)
CCTK_WARN(0, "could not set SCALAR value in table!");
ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_El_group_bound,
- "EMScript::El_group", "scalar");
+ "My_New_Implementation::El_group", "scalar");
if (ierr < 0)
- CCTK_WARN(0, "Failed to register Scalar BC for EMScript::El_group!");
+ CCTK_WARN(0, "Failed to register Scalar BC for My_New_Implementation::El_group!");
}
@@ -316,10 +316,10 @@ extern "C" void EMScript_SelectBoundConds(CCTK_ARGUMENTS)
CCTK_WARN(0, "could not set SCALAR value in table!");
ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_B1_bound,
- "EMScript::B1", "scalar");
+ "My_New_Implementation::B1", "scalar");
if (ierr < 0)
- CCTK_WARN(0, "Error in registering Scalar BC for EMScript::B1!");
+ CCTK_WARN(0, "Error in registering Scalar BC for My_New_Implementation::B1!");
}
@@ -333,10 +333,10 @@ extern "C" void EMScript_SelectBoundConds(CCTK_ARGUMENTS)
CCTK_WARN(0, "could not set SCALAR value in table!");
ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_B2_bound,
- "EMScript::B2", "scalar");
+ "My_New_Implementation::B2", "scalar");
if (ierr < 0)
- CCTK_WARN(0, "Error in registering Scalar BC for EMScript::B2!");
+ CCTK_WARN(0, "Error in registering Scalar BC for My_New_Implementation::B2!");
}
@@ -350,10 +350,10 @@ extern "C" void EMScript_SelectBoundConds(CCTK_ARGUMENTS)
CCTK_WARN(0, "could not set SCALAR value in table!");
ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_B3_bound,
- "EMScript::B3", "scalar");
+ "My_New_Implementation::B3", "scalar");
if (ierr < 0)
- CCTK_WARN(0, "Error in registering Scalar BC for EMScript::B3!");
+ CCTK_WARN(0, "Error in registering Scalar BC for My_New_Implementation::B3!");
}
@@ -367,10 +367,10 @@ extern "C" void EMScript_SelectBoundConds(CCTK_ARGUMENTS)
CCTK_WARN(0, "could not set SCALAR value in table!");
ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_El1_bound,
- "EMScript::El1", "scalar");
+ "My_New_Implementation::El1", "scalar");
if (ierr < 0)
- CCTK_WARN(0, "Error in registering Scalar BC for EMScript::El1!");
+ CCTK_WARN(0, "Error in registering Scalar BC for My_New_Implementation::El1!");
}
@@ -384,10 +384,10 @@ extern "C" void EMScript_SelectBoundConds(CCTK_ARGUMENTS)
CCTK_WARN(0, "could not set SCALAR value in table!");
ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_El2_bound,
- "EMScript::El2", "scalar");
+ "My_New_Implementation::El2", "scalar");
if (ierr < 0)
- CCTK_WARN(0, "Error in registering Scalar BC for EMScript::El2!");
+ CCTK_WARN(0, "Error in registering Scalar BC for My_New_Implementation::El2!");
}
@@ -401,10 +401,10 @@ extern "C" void EMScript_SelectBoundConds(CCTK_ARGUMENTS)
CCTK_WARN(0, "could not set SCALAR value in table!");
ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_El3_bound,
- "EMScript::El3", "scalar");
+ "My_New_Implementation::El3", "scalar");
if (ierr < 0)
- CCTK_WARN(0, "Error in registering Scalar BC for EMScript::El3!");
+ CCTK_WARN(0, "Error in registering Scalar BC for My_New_Implementation::El3!");
}
return;
@@ -413,45 +413,45 @@ extern "C" void EMScript_SelectBoundConds(CCTK_ARGUMENTS)
/* template for entries in parameter file:
-#$bound$#EMScript::B_group_bound = "skip"
-#$bound$#EMScript::B_group_bound_speed = 1.0
-#$bound$#EMScript::B_group_bound_limit = 0.0
-#$bound$#EMScript::B_group_bound_scalar = 0.0
+#$bound$#My_New_Implementation::B_group_bound = "skip"
+#$bound$#My_New_Implementation::B_group_bound_speed = 1.0
+#$bound$#My_New_Implementation::B_group_bound_limit = 0.0
+#$bound$#My_New_Implementation::B_group_bound_scalar = 0.0
-#$bound$#EMScript::El_group_bound = "skip"
-#$bound$#EMScript::El_group_bound_speed = 1.0
-#$bound$#EMScript::El_group_bound_limit = 0.0
-#$bound$#EMScript::El_group_bound_scalar = 0.0
+#$bound$#My_New_Implementation::El_group_bound = "skip"
+#$bound$#My_New_Implementation::El_group_bound_speed = 1.0
+#$bound$#My_New_Implementation::El_group_bound_limit = 0.0
+#$bound$#My_New_Implementation::El_group_bound_scalar = 0.0
-#$bound$#EMScript::B1_bound = "skip"
-#$bound$#EMScript::B1_bound_speed = 1.0
-#$bound$#EMScript::B1_bound_limit = 0.0
-#$bound$#EMScript::B1_bound_scalar = 0.0
+#$bound$#My_New_Implementation::B1_bound = "skip"
+#$bound$#My_New_Implementation::B1_bound_speed = 1.0
+#$bound$#My_New_Implementation::B1_bound_limit = 0.0
+#$bound$#My_New_Implementation::B1_bound_scalar = 0.0
-#$bound$#EMScript::B2_bound = "skip"
-#$bound$#EMScript::B2_bound_speed = 1.0
-#$bound$#EMScript::B2_bound_limit = 0.0
-#$bound$#EMScript::B2_bound_scalar = 0.0
+#$bound$#My_New_Implementation::B2_bound = "skip"
+#$bound$#My_New_Implementation::B2_bound_speed = 1.0
+#$bound$#My_New_Implementation::B2_bound_limit = 0.0
+#$bound$#My_New_Implementation::B2_bound_scalar = 0.0
-#$bound$#EMScript::B3_bound = "skip"
-#$bound$#EMScript::B3_bound_speed = 1.0
-#$bound$#EMScript::B3_bound_limit = 0.0
-#$bound$#EMScript::B3_bound_scalar = 0.0
+#$bound$#My_New_Implementation::B3_bound = "skip"
+#$bound$#My_New_Implementation::B3_bound_speed = 1.0
+#$bound$#My_New_Implementation::B3_bound_limit = 0.0
+#$bound$#My_New_Implementation::B3_bound_scalar = 0.0
-#$bound$#EMScript::El1_bound = "skip"
-#$bound$#EMScript::El1_bound_speed = 1.0
-#$bound$#EMScript::El1_bound_limit = 0.0
-#$bound$#EMScript::El1_bound_scalar = 0.0
+#$bound$#My_New_Implementation::El1_bound = "skip"
+#$bound$#My_New_Implementation::El1_bound_speed = 1.0
+#$bound$#My_New_Implementation::El1_bound_limit = 0.0
+#$bound$#My_New_Implementation::El1_bound_scalar = 0.0
-#$bound$#EMScript::El2_bound = "skip"
-#$bound$#EMScript::El2_bound_speed = 1.0
-#$bound$#EMScript::El2_bound_limit = 0.0
-#$bound$#EMScript::El2_bound_scalar = 0.0
+#$bound$#My_New_Implementation::El2_bound = "skip"
+#$bound$#My_New_Implementation::El2_bound_speed = 1.0
+#$bound$#My_New_Implementation::El2_bound_limit = 0.0
+#$bound$#My_New_Implementation::El2_bound_scalar = 0.0
-#$bound$#EMScript::El3_bound = "skip"
-#$bound$#EMScript::El3_bound_speed = 1.0
-#$bound$#EMScript::El3_bound_limit = 0.0
-#$bound$#EMScript::El3_bound_scalar = 0.0
+#$bound$#My_New_Implementation::El3_bound = "skip"
+#$bound$#My_New_Implementation::El3_bound_speed = 1.0
+#$bound$#My_New_Implementation::El3_bound_limit = 0.0
+#$bound$#My_New_Implementation::El3_bound_scalar = 0.0
*/
diff --git a/Examples/EMScript/src/Differencing.h b/Examples/EMScript/src/Differencing.h
index 98de1a1..1933a29 100644
--- a/Examples/EMScript/src/Differencing.h
+++ b/Examples/EMScript/src/Differencing.h
@@ -1,288 +1,291 @@
-#ifndef KRANC_DIFF_FUNCTIONS
-# define PDstandardfdOrder21(u) ((-KRANC_GFOFFSET3D(u,-1,0,0) + KRANC_GFOFFSET3D(u,1,0,0))*p1o2dx)
-#else
-# define PDstandardfdOrder21(u) (PDstandardfdOrder21_impl(u,p1o2dx,cdj,cdk))
-static CCTK_REAL PDstandardfdOrder21_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 PDstandardfdOrder21_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 PDstandardfdOrder22(u) ((-KRANC_GFOFFSET3D(u,0,-1,0) + KRANC_GFOFFSET3D(u,0,1,0))*p1o2dy)
-#else
-# define PDstandardfdOrder22(u) (PDstandardfdOrder22_impl(u,p1o2dy,cdj,cdk))
-static CCTK_REAL PDstandardfdOrder22_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 PDstandardfdOrder22_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 PDstandardfdOrder23(u) ((-KRANC_GFOFFSET3D(u,0,0,-1) + KRANC_GFOFFSET3D(u,0,0,1))*p1o2dz)
-#else
-# define PDstandardfdOrder23(u) (PDstandardfdOrder23_impl(u,p1o2dz,cdj,cdk))
-static CCTK_REAL PDstandardfdOrder23_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 PDstandardfdOrder23_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 PDstandardfdOrder41(u) ((-8*KRANC_GFOFFSET3D(u,-1,0,0) + 8*KRANC_GFOFFSET3D(u,1,0,0) + KRANC_GFOFFSET3D(u,-2,0,0) - KRANC_GFOFFSET3D(u,2,0,0))*p1o12dx)
-#else
-# define PDstandardfdOrder41(u) (PDstandardfdOrder41_impl(u,p1o12dx,cdj,cdk))
-static CCTK_REAL PDstandardfdOrder41_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o12dx, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;
-static CCTK_REAL PDstandardfdOrder41_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o12dx, ptrdiff_t const cdj, ptrdiff_t const cdk)
-{
- ptrdiff_t const cdi=sizeof(CCTK_REAL);
- return (-8*KRANC_GFOFFSET3D(u,-1,0,0) + 8*KRANC_GFOFFSET3D(u,1,0,0) + KRANC_GFOFFSET3D(u,-2,0,0) - KRANC_GFOFFSET3D(u,2,0,0))*p1o12dx;
-}
-#endif
-
-#ifndef KRANC_DIFF_FUNCTIONS
-# define PDstandardfdOrder42(u) ((-8*KRANC_GFOFFSET3D(u,0,-1,0) + 8*KRANC_GFOFFSET3D(u,0,1,0) + KRANC_GFOFFSET3D(u,0,-2,0) - KRANC_GFOFFSET3D(u,0,2,0))*p1o12dy)
-#else
-# define PDstandardfdOrder42(u) (PDstandardfdOrder42_impl(u,p1o12dy,cdj,cdk))
-static CCTK_REAL PDstandardfdOrder42_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o12dy, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;
-static CCTK_REAL PDstandardfdOrder42_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o12dy, ptrdiff_t const cdj, ptrdiff_t const cdk)
-{
- ptrdiff_t const cdi=sizeof(CCTK_REAL);
- return (-8*KRANC_GFOFFSET3D(u,0,-1,0) + 8*KRANC_GFOFFSET3D(u,0,1,0) + KRANC_GFOFFSET3D(u,0,-2,0) - KRANC_GFOFFSET3D(u,0,2,0))*p1o12dy;
-}
-#endif
-
-#ifndef KRANC_DIFF_FUNCTIONS
-# define PDstandardfdOrder43(u) ((-8*KRANC_GFOFFSET3D(u,0,0,-1) + 8*KRANC_GFOFFSET3D(u,0,0,1) + KRANC_GFOFFSET3D(u,0,0,-2) - KRANC_GFOFFSET3D(u,0,0,2))*p1o12dz)
-#else
-# define PDstandardfdOrder43(u) (PDstandardfdOrder43_impl(u,p1o12dz,cdj,cdk))
-static CCTK_REAL PDstandardfdOrder43_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o12dz, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;
-static CCTK_REAL PDstandardfdOrder43_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o12dz, ptrdiff_t const cdj, ptrdiff_t const cdk)
-{
- ptrdiff_t const cdi=sizeof(CCTK_REAL);
- return (-8*KRANC_GFOFFSET3D(u,0,0,-1) + 8*KRANC_GFOFFSET3D(u,0,0,1) + KRANC_GFOFFSET3D(u,0,0,-2) - KRANC_GFOFFSET3D(u,0,0,2))*p1o12dz;
-}
-#endif
-
-#ifndef KRANC_DIFF_FUNCTIONS
-# define PDstandardfdOrder211(u) ((-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,-1,0,0) + KRANC_GFOFFSET3D(u,1,0,0))*p1odx2)
-#else
-# define PDstandardfdOrder211(u) (PDstandardfdOrder211_impl(u,p1odx2,cdj,cdk))
-static CCTK_REAL PDstandardfdOrder211_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 PDstandardfdOrder211_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 PDstandardfdOrder222(u) ((-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,0,-1,0) + KRANC_GFOFFSET3D(u,0,1,0))*p1ody2)
-#else
-# define PDstandardfdOrder222(u) (PDstandardfdOrder222_impl(u,p1ody2,cdj,cdk))
-static CCTK_REAL PDstandardfdOrder222_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 PDstandardfdOrder222_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 PDstandardfdOrder233(u) ((-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,0,0,-1) + KRANC_GFOFFSET3D(u,0,0,1))*p1odz2)
-#else
-# define PDstandardfdOrder233(u) (PDstandardfdOrder233_impl(u,p1odz2,cdj,cdk))
-static CCTK_REAL PDstandardfdOrder233_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 PDstandardfdOrder233_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
-
-#ifndef KRANC_DIFF_FUNCTIONS
-# define PDstandardfdOrder411(u) ((30*KRANC_GFOFFSET3D(u,0,0,0) - 16*(KRANC_GFOFFSET3D(u,-1,0,0) + KRANC_GFOFFSET3D(u,1,0,0)) + KRANC_GFOFFSET3D(u,-2,0,0) + KRANC_GFOFFSET3D(u,2,0,0))*pm1o12dx2)
-#else
-# define PDstandardfdOrder411(u) (PDstandardfdOrder411_impl(u,pm1o12dx2,cdj,cdk))
-static CCTK_REAL PDstandardfdOrder411_impl(CCTK_REAL const* restrict const u, CCTK_REAL const pm1o12dx2, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;
-static CCTK_REAL PDstandardfdOrder411_impl(CCTK_REAL const* restrict const u, CCTK_REAL const pm1o12dx2, ptrdiff_t const cdj, ptrdiff_t const cdk)
-{
- ptrdiff_t const cdi=sizeof(CCTK_REAL);
- return (30*KRANC_GFOFFSET3D(u,0,0,0) - 16*(KRANC_GFOFFSET3D(u,-1,0,0) + KRANC_GFOFFSET3D(u,1,0,0)) + KRANC_GFOFFSET3D(u,-2,0,0) + KRANC_GFOFFSET3D(u,2,0,0))*pm1o12dx2;
-}
-#endif
-
-#ifndef KRANC_DIFF_FUNCTIONS
-# define PDstandardfdOrder422(u) ((30*KRANC_GFOFFSET3D(u,0,0,0) - 16*(KRANC_GFOFFSET3D(u,0,-1,0) + KRANC_GFOFFSET3D(u,0,1,0)) + KRANC_GFOFFSET3D(u,0,-2,0) + KRANC_GFOFFSET3D(u,0,2,0))*pm1o12dy2)
-#else
-# define PDstandardfdOrder422(u) (PDstandardfdOrder422_impl(u,pm1o12dy2,cdj,cdk))
-static CCTK_REAL PDstandardfdOrder422_impl(CCTK_REAL const* restrict const u, CCTK_REAL const pm1o12dy2, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;
-static CCTK_REAL PDstandardfdOrder422_impl(CCTK_REAL const* restrict const u, CCTK_REAL const pm1o12dy2, ptrdiff_t const cdj, ptrdiff_t const cdk)
-{
- ptrdiff_t const cdi=sizeof(CCTK_REAL);
- return (30*KRANC_GFOFFSET3D(u,0,0,0) - 16*(KRANC_GFOFFSET3D(u,0,-1,0) + KRANC_GFOFFSET3D(u,0,1,0)) + KRANC_GFOFFSET3D(u,0,-2,0) + KRANC_GFOFFSET3D(u,0,2,0))*pm1o12dy2;
-}
-#endif
-
-#ifndef KRANC_DIFF_FUNCTIONS
-# define PDstandardfdOrder433(u) ((30*KRANC_GFOFFSET3D(u,0,0,0) - 16*(KRANC_GFOFFSET3D(u,0,0,-1) + KRANC_GFOFFSET3D(u,0,0,1)) + KRANC_GFOFFSET3D(u,0,0,-2) + KRANC_GFOFFSET3D(u,0,0,2))*pm1o12dz2)
-#else
-# define PDstandardfdOrder433(u) (PDstandardfdOrder433_impl(u,pm1o12dz2,cdj,cdk))
-static CCTK_REAL PDstandardfdOrder433_impl(CCTK_REAL const* restrict const u, CCTK_REAL const pm1o12dz2, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;
-static CCTK_REAL PDstandardfdOrder433_impl(CCTK_REAL const* restrict const u, CCTK_REAL const pm1o12dz2, ptrdiff_t const cdj, ptrdiff_t const cdk)
-{
- ptrdiff_t const cdi=sizeof(CCTK_REAL);
- return (30*KRANC_GFOFFSET3D(u,0,0,0) - 16*(KRANC_GFOFFSET3D(u,0,0,-1) + KRANC_GFOFFSET3D(u,0,0,1)) + KRANC_GFOFFSET3D(u,0,0,-2) + KRANC_GFOFFSET3D(u,0,0,2))*pm1o12dz2;
-}
-#endif
-
-#ifndef KRANC_DIFF_FUNCTIONS
-# define PDstandardfdOrder212(u) ((KRANC_GFOFFSET3D(u,-1,-1,0) - KRANC_GFOFFSET3D(u,-1,1,0) - KRANC_GFOFFSET3D(u,1,-1,0) + KRANC_GFOFFSET3D(u,1,1,0))*p1o4dxdy)
-#else
-# define PDstandardfdOrder212(u) (PDstandardfdOrder212_impl(u,p1o4dxdy,cdj,cdk))
-static CCTK_REAL PDstandardfdOrder212_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o4dxdy, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;
-static CCTK_REAL PDstandardfdOrder212_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o4dxdy, ptrdiff_t const cdj, ptrdiff_t const cdk)
-{
- ptrdiff_t const cdi=sizeof(CCTK_REAL);
- return (KRANC_GFOFFSET3D(u,-1,-1,0) - KRANC_GFOFFSET3D(u,-1,1,0) - KRANC_GFOFFSET3D(u,1,-1,0) + KRANC_GFOFFSET3D(u,1,1,0))*p1o4dxdy;
-}
-#endif
-
-#ifndef KRANC_DIFF_FUNCTIONS
-# define PDstandardfdOrder213(u) ((KRANC_GFOFFSET3D(u,-1,0,-1) - KRANC_GFOFFSET3D(u,-1,0,1) - KRANC_GFOFFSET3D(u,1,0,-1) + KRANC_GFOFFSET3D(u,1,0,1))*p1o4dxdz)
-#else
-# define PDstandardfdOrder213(u) (PDstandardfdOrder213_impl(u,p1o4dxdz,cdj,cdk))
-static CCTK_REAL PDstandardfdOrder213_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o4dxdz, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;
-static CCTK_REAL PDstandardfdOrder213_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o4dxdz, ptrdiff_t const cdj, ptrdiff_t const cdk)
-{
- ptrdiff_t const cdi=sizeof(CCTK_REAL);
- return (KRANC_GFOFFSET3D(u,-1,0,-1) - KRANC_GFOFFSET3D(u,-1,0,1) - KRANC_GFOFFSET3D(u,1,0,-1) + KRANC_GFOFFSET3D(u,1,0,1))*p1o4dxdz;
-}
-#endif
-
-#ifndef KRANC_DIFF_FUNCTIONS
-# define PDstandardfdOrder221(u) ((KRANC_GFOFFSET3D(u,-1,-1,0) - KRANC_GFOFFSET3D(u,-1,1,0) - KRANC_GFOFFSET3D(u,1,-1,0) + KRANC_GFOFFSET3D(u,1,1,0))*p1o4dxdy)
-#else
-# define PDstandardfdOrder221(u) (PDstandardfdOrder221_impl(u,p1o4dxdy,cdj,cdk))
-static CCTK_REAL PDstandardfdOrder221_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o4dxdy, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;
-static CCTK_REAL PDstandardfdOrder221_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o4dxdy, ptrdiff_t const cdj, ptrdiff_t const cdk)
-{
- ptrdiff_t const cdi=sizeof(CCTK_REAL);
- return (KRANC_GFOFFSET3D(u,-1,-1,0) - KRANC_GFOFFSET3D(u,-1,1,0) - KRANC_GFOFFSET3D(u,1,-1,0) + KRANC_GFOFFSET3D(u,1,1,0))*p1o4dxdy;
-}
-#endif
-
-#ifndef KRANC_DIFF_FUNCTIONS
-# define PDstandardfdOrder223(u) ((KRANC_GFOFFSET3D(u,0,-1,-1) - KRANC_GFOFFSET3D(u,0,-1,1) - KRANC_GFOFFSET3D(u,0,1,-1) + KRANC_GFOFFSET3D(u,0,1,1))*p1o4dydz)
-#else
-# define PDstandardfdOrder223(u) (PDstandardfdOrder223_impl(u,p1o4dydz,cdj,cdk))
-static CCTK_REAL PDstandardfdOrder223_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o4dydz, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;
-static CCTK_REAL PDstandardfdOrder223_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o4dydz, ptrdiff_t const cdj, ptrdiff_t const cdk)
-{
- ptrdiff_t const cdi=sizeof(CCTK_REAL);
- return (KRANC_GFOFFSET3D(u,0,-1,-1) - KRANC_GFOFFSET3D(u,0,-1,1) - KRANC_GFOFFSET3D(u,0,1,-1) + KRANC_GFOFFSET3D(u,0,1,1))*p1o4dydz;
-}
-#endif
-
-#ifndef KRANC_DIFF_FUNCTIONS
-# define PDstandardfdOrder231(u) ((KRANC_GFOFFSET3D(u,-1,0,-1) - KRANC_GFOFFSET3D(u,-1,0,1) - KRANC_GFOFFSET3D(u,1,0,-1) + KRANC_GFOFFSET3D(u,1,0,1))*p1o4dxdz)
-#else
-# define PDstandardfdOrder231(u) (PDstandardfdOrder231_impl(u,p1o4dxdz,cdj,cdk))
-static CCTK_REAL PDstandardfdOrder231_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o4dxdz, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;
-static CCTK_REAL PDstandardfdOrder231_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o4dxdz, ptrdiff_t const cdj, ptrdiff_t const cdk)
-{
- ptrdiff_t const cdi=sizeof(CCTK_REAL);
- return (KRANC_GFOFFSET3D(u,-1,0,-1) - KRANC_GFOFFSET3D(u,-1,0,1) - KRANC_GFOFFSET3D(u,1,0,-1) + KRANC_GFOFFSET3D(u,1,0,1))*p1o4dxdz;
-}
-#endif
-
-#ifndef KRANC_DIFF_FUNCTIONS
-# define PDstandardfdOrder232(u) ((KRANC_GFOFFSET3D(u,0,-1,-1) - KRANC_GFOFFSET3D(u,0,-1,1) - KRANC_GFOFFSET3D(u,0,1,-1) + KRANC_GFOFFSET3D(u,0,1,1))*p1o4dydz)
-#else
-# define PDstandardfdOrder232(u) (PDstandardfdOrder232_impl(u,p1o4dydz,cdj,cdk))
-static CCTK_REAL PDstandardfdOrder232_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o4dydz, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;
-static CCTK_REAL PDstandardfdOrder232_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o4dydz, ptrdiff_t const cdj, ptrdiff_t const cdk)
-{
- ptrdiff_t const cdi=sizeof(CCTK_REAL);
- return (KRANC_GFOFFSET3D(u,0,-1,-1) - KRANC_GFOFFSET3D(u,0,-1,1) - KRANC_GFOFFSET3D(u,0,1,-1) + KRANC_GFOFFSET3D(u,0,1,1))*p1o4dydz;
-}
-#endif
-
-#ifndef KRANC_DIFF_FUNCTIONS
-# define PDstandardfdOrder412(u) ((-64*(KRANC_GFOFFSET3D(u,-1,1,0) + KRANC_GFOFFSET3D(u,1,-1,0)) + 64*(KRANC_GFOFFSET3D(u,-1,-1,0) + KRANC_GFOFFSET3D(u,1,1,0)) + 8*(KRANC_GFOFFSET3D(u,-1,2,0) + KRANC_GFOFFSET3D(u,1,-2,0) + KRANC_GFOFFSET3D(u,-2,1,0) + KRANC_GFOFFSET3D(u,2,-1,0)) - 8*(KRANC_GFOFFSET3D(u,-1,-2,0) + KRANC_GFOFFSET3D(u,1,2,0) + KRANC_GFOFFSET3D(u,-2,-1,0) + KRANC_GFOFFSET3D(u,2,1,0)) + KRANC_GFOFFSET3D(u,-2,-2,0) - KRANC_GFOFFSET3D(u,-2,2,0) - KRANC_GFOFFSET3D(u,2,-2,0) + KRANC_GFOFFSET3D(u,2,2,0))*p1o144dxdy)
-#else
-# define PDstandardfdOrder412(u) (PDstandardfdOrder412_impl(u,p1o144dxdy,cdj,cdk))
-static CCTK_REAL PDstandardfdOrder412_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o144dxdy, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;
-static CCTK_REAL PDstandardfdOrder412_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o144dxdy, ptrdiff_t const cdj, ptrdiff_t const cdk)
-{
- ptrdiff_t const cdi=sizeof(CCTK_REAL);
- return (-64*(KRANC_GFOFFSET3D(u,-1,1,0) + KRANC_GFOFFSET3D(u,1,-1,0)) + 64*(KRANC_GFOFFSET3D(u,-1,-1,0) + KRANC_GFOFFSET3D(u,1,1,0)) + 8*(KRANC_GFOFFSET3D(u,-1,2,0) + KRANC_GFOFFSET3D(u,1,-2,0) + KRANC_GFOFFSET3D(u,-2,1,0) + KRANC_GFOFFSET3D(u,2,-1,0)) - 8*(KRANC_GFOFFSET3D(u,-1,-2,0) + KRANC_GFOFFSET3D(u,1,2,0) + KRANC_GFOFFSET3D(u,-2,-1,0) + KRANC_GFOFFSET3D(u,2,1,0)) + KRANC_GFOFFSET3D(u,-2,-2,0) - KRANC_GFOFFSET3D(u,-2,2,0) - KRANC_GFOFFSET3D(u,2,-2,0) + KRANC_GFOFFSET3D(u,2,2,0))*p1o144dxdy;
-}
-#endif
-
-#ifndef KRANC_DIFF_FUNCTIONS
-# define PDstandardfdOrder413(u) ((-64*(KRANC_GFOFFSET3D(u,-1,0,1) + KRANC_GFOFFSET3D(u,1,0,-1)) + 64*(KRANC_GFOFFSET3D(u,-1,0,-1) + KRANC_GFOFFSET3D(u,1,0,1)) + 8*(KRANC_GFOFFSET3D(u,-1,0,2) + KRANC_GFOFFSET3D(u,1,0,-2) + KRANC_GFOFFSET3D(u,-2,0,1) + KRANC_GFOFFSET3D(u,2,0,-1)) - 8*(KRANC_GFOFFSET3D(u,-1,0,-2) + KRANC_GFOFFSET3D(u,1,0,2) + KRANC_GFOFFSET3D(u,-2,0,-1) + KRANC_GFOFFSET3D(u,2,0,1)) + KRANC_GFOFFSET3D(u,-2,0,-2) - KRANC_GFOFFSET3D(u,-2,0,2) - KRANC_GFOFFSET3D(u,2,0,-2) + KRANC_GFOFFSET3D(u,2,0,2))*p1o144dxdz)
-#else
-# define PDstandardfdOrder413(u) (PDstandardfdOrder413_impl(u,p1o144dxdz,cdj,cdk))
-static CCTK_REAL PDstandardfdOrder413_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o144dxdz, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;
-static CCTK_REAL PDstandardfdOrder413_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o144dxdz, ptrdiff_t const cdj, ptrdiff_t const cdk)
-{
- ptrdiff_t const cdi=sizeof(CCTK_REAL);
- return (-64*(KRANC_GFOFFSET3D(u,-1,0,1) + KRANC_GFOFFSET3D(u,1,0,-1)) + 64*(KRANC_GFOFFSET3D(u,-1,0,-1) + KRANC_GFOFFSET3D(u,1,0,1)) + 8*(KRANC_GFOFFSET3D(u,-1,0,2) + KRANC_GFOFFSET3D(u,1,0,-2) + KRANC_GFOFFSET3D(u,-2,0,1) + KRANC_GFOFFSET3D(u,2,0,-1)) - 8*(KRANC_GFOFFSET3D(u,-1,0,-2) + KRANC_GFOFFSET3D(u,1,0,2) + KRANC_GFOFFSET3D(u,-2,0,-1) + KRANC_GFOFFSET3D(u,2,0,1)) + KRANC_GFOFFSET3D(u,-2,0,-2) - KRANC_GFOFFSET3D(u,-2,0,2) - KRANC_GFOFFSET3D(u,2,0,-2) + KRANC_GFOFFSET3D(u,2,0,2))*p1o144dxdz;
-}
-#endif
-
-#ifndef KRANC_DIFF_FUNCTIONS
-# define PDstandardfdOrder421(u) ((-64*(KRANC_GFOFFSET3D(u,-1,1,0) + KRANC_GFOFFSET3D(u,1,-1,0)) + 64*(KRANC_GFOFFSET3D(u,-1,-1,0) + KRANC_GFOFFSET3D(u,1,1,0)) + 8*(KRANC_GFOFFSET3D(u,-1,2,0) + KRANC_GFOFFSET3D(u,1,-2,0) + KRANC_GFOFFSET3D(u,-2,1,0) + KRANC_GFOFFSET3D(u,2,-1,0)) - 8*(KRANC_GFOFFSET3D(u,-1,-2,0) + KRANC_GFOFFSET3D(u,1,2,0) + KRANC_GFOFFSET3D(u,-2,-1,0) + KRANC_GFOFFSET3D(u,2,1,0)) + KRANC_GFOFFSET3D(u,-2,-2,0) - KRANC_GFOFFSET3D(u,-2,2,0) - KRANC_GFOFFSET3D(u,2,-2,0) + KRANC_GFOFFSET3D(u,2,2,0))*p1o144dxdy)
-#else
-# define PDstandardfdOrder421(u) (PDstandardfdOrder421_impl(u,p1o144dxdy,cdj,cdk))
-static CCTK_REAL PDstandardfdOrder421_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o144dxdy, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;
-static CCTK_REAL PDstandardfdOrder421_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o144dxdy, ptrdiff_t const cdj, ptrdiff_t const cdk)
-{
- ptrdiff_t const cdi=sizeof(CCTK_REAL);
- return (-64*(KRANC_GFOFFSET3D(u,-1,1,0) + KRANC_GFOFFSET3D(u,1,-1,0)) + 64*(KRANC_GFOFFSET3D(u,-1,-1,0) + KRANC_GFOFFSET3D(u,1,1,0)) + 8*(KRANC_GFOFFSET3D(u,-1,2,0) + KRANC_GFOFFSET3D(u,1,-2,0) + KRANC_GFOFFSET3D(u,-2,1,0) + KRANC_GFOFFSET3D(u,2,-1,0)) - 8*(KRANC_GFOFFSET3D(u,-1,-2,0) + KRANC_GFOFFSET3D(u,1,2,0) + KRANC_GFOFFSET3D(u,-2,-1,0) + KRANC_GFOFFSET3D(u,2,1,0)) + KRANC_GFOFFSET3D(u,-2,-2,0) - KRANC_GFOFFSET3D(u,-2,2,0) - KRANC_GFOFFSET3D(u,2,-2,0) + KRANC_GFOFFSET3D(u,2,2,0))*p1o144dxdy;
-}
-#endif
-
-#ifndef KRANC_DIFF_FUNCTIONS
-# define PDstandardfdOrder423(u) ((-64*(KRANC_GFOFFSET3D(u,0,-1,1) + KRANC_GFOFFSET3D(u,0,1,-1)) + 64*(KRANC_GFOFFSET3D(u,0,-1,-1) + KRANC_GFOFFSET3D(u,0,1,1)) + 8*(KRANC_GFOFFSET3D(u,0,-1,2) + KRANC_GFOFFSET3D(u,0,1,-2) + KRANC_GFOFFSET3D(u,0,-2,1) + KRANC_GFOFFSET3D(u,0,2,-1)) - 8*(KRANC_GFOFFSET3D(u,0,-1,-2) + KRANC_GFOFFSET3D(u,0,1,2) + KRANC_GFOFFSET3D(u,0,-2,-1) + KRANC_GFOFFSET3D(u,0,2,1)) + KRANC_GFOFFSET3D(u,0,-2,-2) - KRANC_GFOFFSET3D(u,0,-2,2) - KRANC_GFOFFSET3D(u,0,2,-2) + KRANC_GFOFFSET3D(u,0,2,2))*p1o144dydz)
-#else
-# define PDstandardfdOrder423(u) (PDstandardfdOrder423_impl(u,p1o144dydz,cdj,cdk))
-static CCTK_REAL PDstandardfdOrder423_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o144dydz, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;
-static CCTK_REAL PDstandardfdOrder423_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o144dydz, ptrdiff_t const cdj, ptrdiff_t const cdk)
-{
- ptrdiff_t const cdi=sizeof(CCTK_REAL);
- return (-64*(KRANC_GFOFFSET3D(u,0,-1,1) + KRANC_GFOFFSET3D(u,0,1,-1)) + 64*(KRANC_GFOFFSET3D(u,0,-1,-1) + KRANC_GFOFFSET3D(u,0,1,1)) + 8*(KRANC_GFOFFSET3D(u,0,-1,2) + KRANC_GFOFFSET3D(u,0,1,-2) + KRANC_GFOFFSET3D(u,0,-2,1) + KRANC_GFOFFSET3D(u,0,2,-1)) - 8*(KRANC_GFOFFSET3D(u,0,-1,-2) + KRANC_GFOFFSET3D(u,0,1,2) + KRANC_GFOFFSET3D(u,0,-2,-1) + KRANC_GFOFFSET3D(u,0,2,1)) + KRANC_GFOFFSET3D(u,0,-2,-2) - KRANC_GFOFFSET3D(u,0,-2,2) - KRANC_GFOFFSET3D(u,0,2,-2) + KRANC_GFOFFSET3D(u,0,2,2))*p1o144dydz;
-}
-#endif
-
-#ifndef KRANC_DIFF_FUNCTIONS
-# define PDstandardfdOrder431(u) ((-64*(KRANC_GFOFFSET3D(u,-1,0,1) + KRANC_GFOFFSET3D(u,1,0,-1)) + 64*(KRANC_GFOFFSET3D(u,-1,0,-1) + KRANC_GFOFFSET3D(u,1,0,1)) + 8*(KRANC_GFOFFSET3D(u,-1,0,2) + KRANC_GFOFFSET3D(u,1,0,-2) + KRANC_GFOFFSET3D(u,-2,0,1) + KRANC_GFOFFSET3D(u,2,0,-1)) - 8*(KRANC_GFOFFSET3D(u,-1,0,-2) + KRANC_GFOFFSET3D(u,1,0,2) + KRANC_GFOFFSET3D(u,-2,0,-1) + KRANC_GFOFFSET3D(u,2,0,1)) + KRANC_GFOFFSET3D(u,-2,0,-2) - KRANC_GFOFFSET3D(u,-2,0,2) - KRANC_GFOFFSET3D(u,2,0,-2) + KRANC_GFOFFSET3D(u,2,0,2))*p1o144dxdz)
-#else
-# define PDstandardfdOrder431(u) (PDstandardfdOrder431_impl(u,p1o144dxdz,cdj,cdk))
-static CCTK_REAL PDstandardfdOrder431_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o144dxdz, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;
-static CCTK_REAL PDstandardfdOrder431_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o144dxdz, ptrdiff_t const cdj, ptrdiff_t const cdk)
-{
- ptrdiff_t const cdi=sizeof(CCTK_REAL);
- return (-64*(KRANC_GFOFFSET3D(u,-1,0,1) + KRANC_GFOFFSET3D(u,1,0,-1)) + 64*(KRANC_GFOFFSET3D(u,-1,0,-1) + KRANC_GFOFFSET3D(u,1,0,1)) + 8*(KRANC_GFOFFSET3D(u,-1,0,2) + KRANC_GFOFFSET3D(u,1,0,-2) + KRANC_GFOFFSET3D(u,-2,0,1) + KRANC_GFOFFSET3D(u,2,0,-1)) - 8*(KRANC_GFOFFSET3D(u,-1,0,-2) + KRANC_GFOFFSET3D(u,1,0,2) + KRANC_GFOFFSET3D(u,-2,0,-1) + KRANC_GFOFFSET3D(u,2,0,1)) + KRANC_GFOFFSET3D(u,-2,0,-2) - KRANC_GFOFFSET3D(u,-2,0,2) - KRANC_GFOFFSET3D(u,2,0,-2) + KRANC_GFOFFSET3D(u,2,0,2))*p1o144dxdz;
-}
-#endif
-
-#ifndef KRANC_DIFF_FUNCTIONS
-# define PDstandardfdOrder432(u) ((-64*(KRANC_GFOFFSET3D(u,0,-1,1) + KRANC_GFOFFSET3D(u,0,1,-1)) + 64*(KRANC_GFOFFSET3D(u,0,-1,-1) + KRANC_GFOFFSET3D(u,0,1,1)) + 8*(KRANC_GFOFFSET3D(u,0,-1,2) + KRANC_GFOFFSET3D(u,0,1,-2) + KRANC_GFOFFSET3D(u,0,-2,1) + KRANC_GFOFFSET3D(u,0,2,-1)) - 8*(KRANC_GFOFFSET3D(u,0,-1,-2) + KRANC_GFOFFSET3D(u,0,1,2) + KRANC_GFOFFSET3D(u,0,-2,-1) + KRANC_GFOFFSET3D(u,0,2,1)) + KRANC_GFOFFSET3D(u,0,-2,-2) - KRANC_GFOFFSET3D(u,0,-2,2) - KRANC_GFOFFSET3D(u,0,2,-2) + KRANC_GFOFFSET3D(u,0,2,2))*p1o144dydz)
-#else
-# define PDstandardfdOrder432(u) (PDstandardfdOrder432_impl(u,p1o144dydz,cdj,cdk))
-static CCTK_REAL PDstandardfdOrder432_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o144dydz, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;
-static CCTK_REAL PDstandardfdOrder432_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o144dydz, ptrdiff_t const cdj, ptrdiff_t const cdk)
-{
- ptrdiff_t const cdi=sizeof(CCTK_REAL);
- return (-64*(KRANC_GFOFFSET3D(u,0,-1,1) + KRANC_GFOFFSET3D(u,0,1,-1)) + 64*(KRANC_GFOFFSET3D(u,0,-1,-1) + KRANC_GFOFFSET3D(u,0,1,1)) + 8*(KRANC_GFOFFSET3D(u,0,-1,2) + KRANC_GFOFFSET3D(u,0,1,-2) + KRANC_GFOFFSET3D(u,0,-2,1) + KRANC_GFOFFSET3D(u,0,2,-1)) - 8*(KRANC_GFOFFSET3D(u,0,-1,-2) + KRANC_GFOFFSET3D(u,0,1,2) + KRANC_GFOFFSET3D(u,0,-2,-1) + KRANC_GFOFFSET3D(u,0,2,1)) + KRANC_GFOFFSET3D(u,0,-2,-2) - KRANC_GFOFFSET3D(u,0,-2,2) - KRANC_GFOFFSET3D(u,0,2,-2) + KRANC_GFOFFSET3D(u,0,2,2))*p1o144dydz;
-}
-#endif
-
+static char const * const differencing =
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"# define PDstandardfdOrder21(u) (kmul(p1o2dx,ksub(KRANC_GFOFFSET3D(u,1,0,0),KRANC_GFOFFSET3D(u,-1,0,0))))\n"
+"#else\n"
+"# define PDstandardfdOrder21(u) (PDstandardfdOrder21_impl(u,p1o2dx,cdj,cdk))\n"
+"static CCTK_REAL_VEC PDstandardfdOrder21_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o2dx, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL_VEC PDstandardfdOrder21_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o2dx, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+" return kmul(p1o2dx,ksub(KRANC_GFOFFSET3D(u,1,0,0),KRANC_GFOFFSET3D(u,-1,0,0)));\n"
+"}\n"
+"#endif\n"
+"\n"
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"# define PDstandardfdOrder22(u) (kmul(p1o2dy,ksub(KRANC_GFOFFSET3D(u,0,1,0),KRANC_GFOFFSET3D(u,0,-1,0))))\n"
+"#else\n"
+"# define PDstandardfdOrder22(u) (PDstandardfdOrder22_impl(u,p1o2dy,cdj,cdk))\n"
+"static CCTK_REAL_VEC PDstandardfdOrder22_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o2dy, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL_VEC PDstandardfdOrder22_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o2dy, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+" return kmul(p1o2dy,ksub(KRANC_GFOFFSET3D(u,0,1,0),KRANC_GFOFFSET3D(u,0,-1,0)));\n"
+"}\n"
+"#endif\n"
+"\n"
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"# define PDstandardfdOrder23(u) (kmul(p1o2dz,ksub(KRANC_GFOFFSET3D(u,0,0,1),KRANC_GFOFFSET3D(u,0,0,-1))))\n"
+"#else\n"
+"# define PDstandardfdOrder23(u) (PDstandardfdOrder23_impl(u,p1o2dz,cdj,cdk))\n"
+"static CCTK_REAL_VEC PDstandardfdOrder23_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o2dz, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL_VEC PDstandardfdOrder23_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o2dz, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+" return kmul(p1o2dz,ksub(KRANC_GFOFFSET3D(u,0,0,1),KRANC_GFOFFSET3D(u,0,0,-1)));\n"
+"}\n"
+"#endif\n"
+"\n"
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"# define PDstandardfdOrder41(u) (kmul(p1o12dx,kadd(KRANC_GFOFFSET3D(u,-2,0,0),kmadd(KRANC_GFOFFSET3D(u,-1,0,0),ToReal(-8),kmsub(KRANC_GFOFFSET3D(u,1,0,0),ToReal(8),KRANC_GFOFFSET3D(u,2,0,0))))))\n"
+"#else\n"
+"# define PDstandardfdOrder41(u) (PDstandardfdOrder41_impl(u,p1o12dx,cdj,cdk))\n"
+"static CCTK_REAL_VEC PDstandardfdOrder41_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o12dx, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL_VEC PDstandardfdOrder41_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o12dx, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+" return kmul(p1o12dx,kadd(KRANC_GFOFFSET3D(u,-2,0,0),kmadd(KRANC_GFOFFSET3D(u,-1,0,0),ToReal(-8),kmsub(KRANC_GFOFFSET3D(u,1,0,0),ToReal(8),KRANC_GFOFFSET3D(u,2,0,0)))));\n"
+"}\n"
+"#endif\n"
+"\n"
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"# define PDstandardfdOrder42(u) (kmul(p1o12dy,kadd(KRANC_GFOFFSET3D(u,0,-2,0),kmadd(KRANC_GFOFFSET3D(u,0,-1,0),ToReal(-8),kmsub(KRANC_GFOFFSET3D(u,0,1,0),ToReal(8),KRANC_GFOFFSET3D(u,0,2,0))))))\n"
+"#else\n"
+"# define PDstandardfdOrder42(u) (PDstandardfdOrder42_impl(u,p1o12dy,cdj,cdk))\n"
+"static CCTK_REAL_VEC PDstandardfdOrder42_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o12dy, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL_VEC PDstandardfdOrder42_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o12dy, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+" return kmul(p1o12dy,kadd(KRANC_GFOFFSET3D(u,0,-2,0),kmadd(KRANC_GFOFFSET3D(u,0,-1,0),ToReal(-8),kmsub(KRANC_GFOFFSET3D(u,0,1,0),ToReal(8),KRANC_GFOFFSET3D(u,0,2,0)))));\n"
+"}\n"
+"#endif\n"
+"\n"
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"# define PDstandardfdOrder43(u) (kmul(p1o12dz,kadd(KRANC_GFOFFSET3D(u,0,0,-2),kmadd(KRANC_GFOFFSET3D(u,0,0,-1),ToReal(-8),kmsub(KRANC_GFOFFSET3D(u,0,0,1),ToReal(8),KRANC_GFOFFSET3D(u,0,0,2))))))\n"
+"#else\n"
+"# define PDstandardfdOrder43(u) (PDstandardfdOrder43_impl(u,p1o12dz,cdj,cdk))\n"
+"static CCTK_REAL_VEC PDstandardfdOrder43_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o12dz, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL_VEC PDstandardfdOrder43_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o12dz, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+" return kmul(p1o12dz,kadd(KRANC_GFOFFSET3D(u,0,0,-2),kmadd(KRANC_GFOFFSET3D(u,0,0,-1),ToReal(-8),kmsub(KRANC_GFOFFSET3D(u,0,0,1),ToReal(8),KRANC_GFOFFSET3D(u,0,0,2)))));\n"
+"}\n"
+"#endif\n"
+"\n"
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"# define PDstandardfdOrder211(u) (kmul(p1odx2,kadd(KRANC_GFOFFSET3D(u,-1,0,0),kmadd(KRANC_GFOFFSET3D(u,0,0,0),ToReal(-2),KRANC_GFOFFSET3D(u,1,0,0)))))\n"
+"#else\n"
+"# define PDstandardfdOrder211(u) (PDstandardfdOrder211_impl(u,p1odx2,cdj,cdk))\n"
+"static CCTK_REAL_VEC PDstandardfdOrder211_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1odx2, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL_VEC PDstandardfdOrder211_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1odx2, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+" return kmul(p1odx2,kadd(KRANC_GFOFFSET3D(u,-1,0,0),kmadd(KRANC_GFOFFSET3D(u,0,0,0),ToReal(-2),KRANC_GFOFFSET3D(u,1,0,0))));\n"
+"}\n"
+"#endif\n"
+"\n"
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"# define PDstandardfdOrder222(u) (kmul(p1ody2,kadd(KRANC_GFOFFSET3D(u,0,-1,0),kmadd(KRANC_GFOFFSET3D(u,0,0,0),ToReal(-2),KRANC_GFOFFSET3D(u,0,1,0)))))\n"
+"#else\n"
+"# define PDstandardfdOrder222(u) (PDstandardfdOrder222_impl(u,p1ody2,cdj,cdk))\n"
+"static CCTK_REAL_VEC PDstandardfdOrder222_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1ody2, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL_VEC PDstandardfdOrder222_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1ody2, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+" return kmul(p1ody2,kadd(KRANC_GFOFFSET3D(u,0,-1,0),kmadd(KRANC_GFOFFSET3D(u,0,0,0),ToReal(-2),KRANC_GFOFFSET3D(u,0,1,0))));\n"
+"}\n"
+"#endif\n"
+"\n"
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"# define PDstandardfdOrder233(u) (kmul(p1odz2,kadd(KRANC_GFOFFSET3D(u,0,0,-1),kmadd(KRANC_GFOFFSET3D(u,0,0,0),ToReal(-2),KRANC_GFOFFSET3D(u,0,0,1)))))\n"
+"#else\n"
+"# define PDstandardfdOrder233(u) (PDstandardfdOrder233_impl(u,p1odz2,cdj,cdk))\n"
+"static CCTK_REAL_VEC PDstandardfdOrder233_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1odz2, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL_VEC PDstandardfdOrder233_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1odz2, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+" return kmul(p1odz2,kadd(KRANC_GFOFFSET3D(u,0,0,-1),kmadd(KRANC_GFOFFSET3D(u,0,0,0),ToReal(-2),KRANC_GFOFFSET3D(u,0,0,1))));\n"
+"}\n"
+"#endif\n"
+"\n"
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"# define PDstandardfdOrder411(u) (kmul(pm1o12dx2,kadd(KRANC_GFOFFSET3D(u,-2,0,0),kadd(KRANC_GFOFFSET3D(u,2,0,0),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,0,0),KRANC_GFOFFSET3D(u,1,0,0)),ToReal(-16),kmul(KRANC_GFOFFSET3D(u,0,0,0),ToReal(30)))))))\n"
+"#else\n"
+"# define PDstandardfdOrder411(u) (PDstandardfdOrder411_impl(u,pm1o12dx2,cdj,cdk))\n"
+"static CCTK_REAL_VEC PDstandardfdOrder411_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const pm1o12dx2, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL_VEC PDstandardfdOrder411_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const pm1o12dx2, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+" return kmul(pm1o12dx2,kadd(KRANC_GFOFFSET3D(u,-2,0,0),kadd(KRANC_GFOFFSET3D(u,2,0,0),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,0,0),KRANC_GFOFFSET3D(u,1,0,0)),ToReal(-16),kmul(KRANC_GFOFFSET3D(u,0,0,0),ToReal(30))))));\n"
+"}\n"
+"#endif\n"
+"\n"
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"# define PDstandardfdOrder422(u) (kmul(pm1o12dy2,kadd(KRANC_GFOFFSET3D(u,0,-2,0),kadd(KRANC_GFOFFSET3D(u,0,2,0),kmadd(kadd(KRANC_GFOFFSET3D(u,0,-1,0),KRANC_GFOFFSET3D(u,0,1,0)),ToReal(-16),kmul(KRANC_GFOFFSET3D(u,0,0,0),ToReal(30)))))))\n"
+"#else\n"
+"# define PDstandardfdOrder422(u) (PDstandardfdOrder422_impl(u,pm1o12dy2,cdj,cdk))\n"
+"static CCTK_REAL_VEC PDstandardfdOrder422_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const pm1o12dy2, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL_VEC PDstandardfdOrder422_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const pm1o12dy2, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+" return kmul(pm1o12dy2,kadd(KRANC_GFOFFSET3D(u,0,-2,0),kadd(KRANC_GFOFFSET3D(u,0,2,0),kmadd(kadd(KRANC_GFOFFSET3D(u,0,-1,0),KRANC_GFOFFSET3D(u,0,1,0)),ToReal(-16),kmul(KRANC_GFOFFSET3D(u,0,0,0),ToReal(30))))));\n"
+"}\n"
+"#endif\n"
+"\n"
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"# define PDstandardfdOrder433(u) (kmul(pm1o12dz2,kadd(KRANC_GFOFFSET3D(u,0,0,-2),kadd(KRANC_GFOFFSET3D(u,0,0,2),kmadd(kadd(KRANC_GFOFFSET3D(u,0,0,-1),KRANC_GFOFFSET3D(u,0,0,1)),ToReal(-16),kmul(KRANC_GFOFFSET3D(u,0,0,0),ToReal(30)))))))\n"
+"#else\n"
+"# define PDstandardfdOrder433(u) (PDstandardfdOrder433_impl(u,pm1o12dz2,cdj,cdk))\n"
+"static CCTK_REAL_VEC PDstandardfdOrder433_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const pm1o12dz2, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL_VEC PDstandardfdOrder433_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const pm1o12dz2, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+" return kmul(pm1o12dz2,kadd(KRANC_GFOFFSET3D(u,0,0,-2),kadd(KRANC_GFOFFSET3D(u,0,0,2),kmadd(kadd(KRANC_GFOFFSET3D(u,0,0,-1),KRANC_GFOFFSET3D(u,0,0,1)),ToReal(-16),kmul(KRANC_GFOFFSET3D(u,0,0,0),ToReal(30))))));\n"
+"}\n"
+"#endif\n"
+"\n"
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"# define PDstandardfdOrder212(u) (kmul(p1o4dxdy,kadd(KRANC_GFOFFSET3D(u,-1,-1,0),ksub(KRANC_GFOFFSET3D(u,1,1,0),kadd(KRANC_GFOFFSET3D(u,1,-1,0),KRANC_GFOFFSET3D(u,-1,1,0))))))\n"
+"#else\n"
+"# define PDstandardfdOrder212(u) (PDstandardfdOrder212_impl(u,p1o4dxdy,cdj,cdk))\n"
+"static CCTK_REAL_VEC PDstandardfdOrder212_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o4dxdy, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL_VEC PDstandardfdOrder212_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o4dxdy, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+" return kmul(p1o4dxdy,kadd(KRANC_GFOFFSET3D(u,-1,-1,0),ksub(KRANC_GFOFFSET3D(u,1,1,0),kadd(KRANC_GFOFFSET3D(u,1,-1,0),KRANC_GFOFFSET3D(u,-1,1,0)))));\n"
+"}\n"
+"#endif\n"
+"\n"
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"# define PDstandardfdOrder213(u) (kmul(p1o4dxdz,kadd(KRANC_GFOFFSET3D(u,-1,0,-1),ksub(KRANC_GFOFFSET3D(u,1,0,1),kadd(KRANC_GFOFFSET3D(u,1,0,-1),KRANC_GFOFFSET3D(u,-1,0,1))))))\n"
+"#else\n"
+"# define PDstandardfdOrder213(u) (PDstandardfdOrder213_impl(u,p1o4dxdz,cdj,cdk))\n"
+"static CCTK_REAL_VEC PDstandardfdOrder213_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o4dxdz, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL_VEC PDstandardfdOrder213_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o4dxdz, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+" return kmul(p1o4dxdz,kadd(KRANC_GFOFFSET3D(u,-1,0,-1),ksub(KRANC_GFOFFSET3D(u,1,0,1),kadd(KRANC_GFOFFSET3D(u,1,0,-1),KRANC_GFOFFSET3D(u,-1,0,1)))));\n"
+"}\n"
+"#endif\n"
+"\n"
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"# define PDstandardfdOrder221(u) (kmul(p1o4dxdy,kadd(KRANC_GFOFFSET3D(u,-1,-1,0),ksub(KRANC_GFOFFSET3D(u,1,1,0),kadd(KRANC_GFOFFSET3D(u,1,-1,0),KRANC_GFOFFSET3D(u,-1,1,0))))))\n"
+"#else\n"
+"# define PDstandardfdOrder221(u) (PDstandardfdOrder221_impl(u,p1o4dxdy,cdj,cdk))\n"
+"static CCTK_REAL_VEC PDstandardfdOrder221_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o4dxdy, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL_VEC PDstandardfdOrder221_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o4dxdy, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+" return kmul(p1o4dxdy,kadd(KRANC_GFOFFSET3D(u,-1,-1,0),ksub(KRANC_GFOFFSET3D(u,1,1,0),kadd(KRANC_GFOFFSET3D(u,1,-1,0),KRANC_GFOFFSET3D(u,-1,1,0)))));\n"
+"}\n"
+"#endif\n"
+"\n"
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"# define PDstandardfdOrder223(u) (kmul(p1o4dydz,kadd(KRANC_GFOFFSET3D(u,0,-1,-1),ksub(KRANC_GFOFFSET3D(u,0,1,1),kadd(KRANC_GFOFFSET3D(u,0,1,-1),KRANC_GFOFFSET3D(u,0,-1,1))))))\n"
+"#else\n"
+"# define PDstandardfdOrder223(u) (PDstandardfdOrder223_impl(u,p1o4dydz,cdj,cdk))\n"
+"static CCTK_REAL_VEC PDstandardfdOrder223_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o4dydz, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL_VEC PDstandardfdOrder223_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o4dydz, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+" return kmul(p1o4dydz,kadd(KRANC_GFOFFSET3D(u,0,-1,-1),ksub(KRANC_GFOFFSET3D(u,0,1,1),kadd(KRANC_GFOFFSET3D(u,0,1,-1),KRANC_GFOFFSET3D(u,0,-1,1)))));\n"
+"}\n"
+"#endif\n"
+"\n"
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"# define PDstandardfdOrder231(u) (kmul(p1o4dxdz,kadd(KRANC_GFOFFSET3D(u,-1,0,-1),ksub(KRANC_GFOFFSET3D(u,1,0,1),kadd(KRANC_GFOFFSET3D(u,1,0,-1),KRANC_GFOFFSET3D(u,-1,0,1))))))\n"
+"#else\n"
+"# define PDstandardfdOrder231(u) (PDstandardfdOrder231_impl(u,p1o4dxdz,cdj,cdk))\n"
+"static CCTK_REAL_VEC PDstandardfdOrder231_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o4dxdz, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL_VEC PDstandardfdOrder231_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o4dxdz, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+" return kmul(p1o4dxdz,kadd(KRANC_GFOFFSET3D(u,-1,0,-1),ksub(KRANC_GFOFFSET3D(u,1,0,1),kadd(KRANC_GFOFFSET3D(u,1,0,-1),KRANC_GFOFFSET3D(u,-1,0,1)))));\n"
+"}\n"
+"#endif\n"
+"\n"
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"# define PDstandardfdOrder232(u) (kmul(p1o4dydz,kadd(KRANC_GFOFFSET3D(u,0,-1,-1),ksub(KRANC_GFOFFSET3D(u,0,1,1),kadd(KRANC_GFOFFSET3D(u,0,1,-1),KRANC_GFOFFSET3D(u,0,-1,1))))))\n"
+"#else\n"
+"# define PDstandardfdOrder232(u) (PDstandardfdOrder232_impl(u,p1o4dydz,cdj,cdk))\n"
+"static CCTK_REAL_VEC PDstandardfdOrder232_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o4dydz, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL_VEC PDstandardfdOrder232_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o4dydz, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+" return kmul(p1o4dydz,kadd(KRANC_GFOFFSET3D(u,0,-1,-1),ksub(KRANC_GFOFFSET3D(u,0,1,1),kadd(KRANC_GFOFFSET3D(u,0,1,-1),KRANC_GFOFFSET3D(u,0,-1,1)))));\n"
+"}\n"
+"#endif\n"
+"\n"
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"# define PDstandardfdOrder412(u) (kmul(p1o144dxdy,kadd(KRANC_GFOFFSET3D(u,-2,-2,0),kadd(KRANC_GFOFFSET3D(u,2,2,0),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,1,0),KRANC_GFOFFSET3D(u,1,-1,0)),ToReal(-64),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,-2,0),kadd(KRANC_GFOFFSET3D(u,1,2,0),kadd(KRANC_GFOFFSET3D(u,-2,-1,0),KRANC_GFOFFSET3D(u,2,1,0)))),ToReal(-8),ksub(ksub(kmadd(kadd(KRANC_GFOFFSET3D(u,-1,2,0),kadd(KRANC_GFOFFSET3D(u,1,-2,0),kadd(KRANC_GFOFFSET3D(u,-2,1,0),KRANC_GFOFFSET3D(u,2,-1,0)))),ToReal(8),kmul(kadd(KRANC_GFOFFSET3D(u,-1,-1,0),KRANC_GFOFFSET3D(u,1,1,0)),ToReal(64))),KRANC_GFOFFSET3D(u,2,-2,0)),KRANC_GFOFFSET3D(u,-2,2,0))))))))\n"
+"#else\n"
+"# define PDstandardfdOrder412(u) (PDstandardfdOrder412_impl(u,p1o144dxdy,cdj,cdk))\n"
+"static CCTK_REAL_VEC PDstandardfdOrder412_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o144dxdy, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL_VEC PDstandardfdOrder412_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o144dxdy, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+" return kmul(p1o144dxdy,kadd(KRANC_GFOFFSET3D(u,-2,-2,0),kadd(KRANC_GFOFFSET3D(u,2,2,0),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,1,0),KRANC_GFOFFSET3D(u,1,-1,0)),ToReal(-64),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,-2,0),kadd(KRANC_GFOFFSET3D(u,1,2,0),kadd(KRANC_GFOFFSET3D(u,-2,-1,0),KRANC_GFOFFSET3D(u,2,1,0)))),ToReal(-8),ksub(ksub(kmadd(kadd(KRANC_GFOFFSET3D(u,-1,2,0),kadd(KRANC_GFOFFSET3D(u,1,-2,0),kadd(KRANC_GFOFFSET3D(u,-2,1,0),KRANC_GFOFFSET3D(u,2,-1,0)))),ToReal(8),kmul(kadd(KRANC_GFOFFSET3D(u,-1,-1,0),KRANC_GFOFFSET3D(u,1,1,0)),ToReal(64))),KRANC_GFOFFSET3D(u,2,-2,0)),KRANC_GFOFFSET3D(u,-2,2,0)))))));\n"
+"}\n"
+"#endif\n"
+"\n"
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"# define PDstandardfdOrder413(u) (kmul(p1o144dxdz,kadd(KRANC_GFOFFSET3D(u,-2,0,-2),kadd(KRANC_GFOFFSET3D(u,2,0,2),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,0,1),KRANC_GFOFFSET3D(u,1,0,-1)),ToReal(-64),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,0,-2),kadd(KRANC_GFOFFSET3D(u,1,0,2),kadd(KRANC_GFOFFSET3D(u,-2,0,-1),KRANC_GFOFFSET3D(u,2,0,1)))),ToReal(-8),ksub(ksub(kmadd(kadd(KRANC_GFOFFSET3D(u,-1,0,2),kadd(KRANC_GFOFFSET3D(u,1,0,-2),kadd(KRANC_GFOFFSET3D(u,-2,0,1),KRANC_GFOFFSET3D(u,2,0,-1)))),ToReal(8),kmul(kadd(KRANC_GFOFFSET3D(u,-1,0,-1),KRANC_GFOFFSET3D(u,1,0,1)),ToReal(64))),KRANC_GFOFFSET3D(u,2,0,-2)),KRANC_GFOFFSET3D(u,-2,0,2))))))))\n"
+"#else\n"
+"# define PDstandardfdOrder413(u) (PDstandardfdOrder413_impl(u,p1o144dxdz,cdj,cdk))\n"
+"static CCTK_REAL_VEC PDstandardfdOrder413_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o144dxdz, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL_VEC PDstandardfdOrder413_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o144dxdz, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+" return kmul(p1o144dxdz,kadd(KRANC_GFOFFSET3D(u,-2,0,-2),kadd(KRANC_GFOFFSET3D(u,2,0,2),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,0,1),KRANC_GFOFFSET3D(u,1,0,-1)),ToReal(-64),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,0,-2),kadd(KRANC_GFOFFSET3D(u,1,0,2),kadd(KRANC_GFOFFSET3D(u,-2,0,-1),KRANC_GFOFFSET3D(u,2,0,1)))),ToReal(-8),ksub(ksub(kmadd(kadd(KRANC_GFOFFSET3D(u,-1,0,2),kadd(KRANC_GFOFFSET3D(u,1,0,-2),kadd(KRANC_GFOFFSET3D(u,-2,0,1),KRANC_GFOFFSET3D(u,2,0,-1)))),ToReal(8),kmul(kadd(KRANC_GFOFFSET3D(u,-1,0,-1),KRANC_GFOFFSET3D(u,1,0,1)),ToReal(64))),KRANC_GFOFFSET3D(u,2,0,-2)),KRANC_GFOFFSET3D(u,-2,0,2)))))));\n"
+"}\n"
+"#endif\n"
+"\n"
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"# define PDstandardfdOrder421(u) (kmul(p1o144dxdy,kadd(KRANC_GFOFFSET3D(u,-2,-2,0),kadd(KRANC_GFOFFSET3D(u,2,2,0),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,1,0),KRANC_GFOFFSET3D(u,1,-1,0)),ToReal(-64),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,-2,0),kadd(KRANC_GFOFFSET3D(u,1,2,0),kadd(KRANC_GFOFFSET3D(u,-2,-1,0),KRANC_GFOFFSET3D(u,2,1,0)))),ToReal(-8),ksub(ksub(kmadd(kadd(KRANC_GFOFFSET3D(u,-1,2,0),kadd(KRANC_GFOFFSET3D(u,1,-2,0),kadd(KRANC_GFOFFSET3D(u,-2,1,0),KRANC_GFOFFSET3D(u,2,-1,0)))),ToReal(8),kmul(kadd(KRANC_GFOFFSET3D(u,-1,-1,0),KRANC_GFOFFSET3D(u,1,1,0)),ToReal(64))),KRANC_GFOFFSET3D(u,2,-2,0)),KRANC_GFOFFSET3D(u,-2,2,0))))))))\n"
+"#else\n"
+"# define PDstandardfdOrder421(u) (PDstandardfdOrder421_impl(u,p1o144dxdy,cdj,cdk))\n"
+"static CCTK_REAL_VEC PDstandardfdOrder421_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o144dxdy, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL_VEC PDstandardfdOrder421_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o144dxdy, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+" return kmul(p1o144dxdy,kadd(KRANC_GFOFFSET3D(u,-2,-2,0),kadd(KRANC_GFOFFSET3D(u,2,2,0),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,1,0),KRANC_GFOFFSET3D(u,1,-1,0)),ToReal(-64),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,-2,0),kadd(KRANC_GFOFFSET3D(u,1,2,0),kadd(KRANC_GFOFFSET3D(u,-2,-1,0),KRANC_GFOFFSET3D(u,2,1,0)))),ToReal(-8),ksub(ksub(kmadd(kadd(KRANC_GFOFFSET3D(u,-1,2,0),kadd(KRANC_GFOFFSET3D(u,1,-2,0),kadd(KRANC_GFOFFSET3D(u,-2,1,0),KRANC_GFOFFSET3D(u,2,-1,0)))),ToReal(8),kmul(kadd(KRANC_GFOFFSET3D(u,-1,-1,0),KRANC_GFOFFSET3D(u,1,1,0)),ToReal(64))),KRANC_GFOFFSET3D(u,2,-2,0)),KRANC_GFOFFSET3D(u,-2,2,0)))))));\n"
+"}\n"
+"#endif\n"
+"\n"
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"# define PDstandardfdOrder423(u) (kmul(p1o144dydz,kadd(KRANC_GFOFFSET3D(u,0,-2,-2),kadd(KRANC_GFOFFSET3D(u,0,2,2),kmadd(kadd(KRANC_GFOFFSET3D(u,0,-1,1),KRANC_GFOFFSET3D(u,0,1,-1)),ToReal(-64),kmadd(kadd(KRANC_GFOFFSET3D(u,0,-1,-2),kadd(KRANC_GFOFFSET3D(u,0,1,2),kadd(KRANC_GFOFFSET3D(u,0,-2,-1),KRANC_GFOFFSET3D(u,0,2,1)))),ToReal(-8),ksub(ksub(kmadd(kadd(KRANC_GFOFFSET3D(u,0,-1,2),kadd(KRANC_GFOFFSET3D(u,0,1,-2),kadd(KRANC_GFOFFSET3D(u,0,-2,1),KRANC_GFOFFSET3D(u,0,2,-1)))),ToReal(8),kmul(kadd(KRANC_GFOFFSET3D(u,0,-1,-1),KRANC_GFOFFSET3D(u,0,1,1)),ToReal(64))),KRANC_GFOFFSET3D(u,0,2,-2)),KRANC_GFOFFSET3D(u,0,-2,2))))))))\n"
+"#else\n"
+"# define PDstandardfdOrder423(u) (PDstandardfdOrder423_impl(u,p1o144dydz,cdj,cdk))\n"
+"static CCTK_REAL_VEC PDstandardfdOrder423_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o144dydz, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL_VEC PDstandardfdOrder423_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o144dydz, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+" return kmul(p1o144dydz,kadd(KRANC_GFOFFSET3D(u,0,-2,-2),kadd(KRANC_GFOFFSET3D(u,0,2,2),kmadd(kadd(KRANC_GFOFFSET3D(u,0,-1,1),KRANC_GFOFFSET3D(u,0,1,-1)),ToReal(-64),kmadd(kadd(KRANC_GFOFFSET3D(u,0,-1,-2),kadd(KRANC_GFOFFSET3D(u,0,1,2),kadd(KRANC_GFOFFSET3D(u,0,-2,-1),KRANC_GFOFFSET3D(u,0,2,1)))),ToReal(-8),ksub(ksub(kmadd(kadd(KRANC_GFOFFSET3D(u,0,-1,2),kadd(KRANC_GFOFFSET3D(u,0,1,-2),kadd(KRANC_GFOFFSET3D(u,0,-2,1),KRANC_GFOFFSET3D(u,0,2,-1)))),ToReal(8),kmul(kadd(KRANC_GFOFFSET3D(u,0,-1,-1),KRANC_GFOFFSET3D(u,0,1,1)),ToReal(64))),KRANC_GFOFFSET3D(u,0,2,-2)),KRANC_GFOFFSET3D(u,0,-2,2)))))));\n"
+"}\n"
+"#endif\n"
+"\n"
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"# define PDstandardfdOrder431(u) (kmul(p1o144dxdz,kadd(KRANC_GFOFFSET3D(u,-2,0,-2),kadd(KRANC_GFOFFSET3D(u,2,0,2),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,0,1),KRANC_GFOFFSET3D(u,1,0,-1)),ToReal(-64),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,0,-2),kadd(KRANC_GFOFFSET3D(u,1,0,2),kadd(KRANC_GFOFFSET3D(u,-2,0,-1),KRANC_GFOFFSET3D(u,2,0,1)))),ToReal(-8),ksub(ksub(kmadd(kadd(KRANC_GFOFFSET3D(u,-1,0,2),kadd(KRANC_GFOFFSET3D(u,1,0,-2),kadd(KRANC_GFOFFSET3D(u,-2,0,1),KRANC_GFOFFSET3D(u,2,0,-1)))),ToReal(8),kmul(kadd(KRANC_GFOFFSET3D(u,-1,0,-1),KRANC_GFOFFSET3D(u,1,0,1)),ToReal(64))),KRANC_GFOFFSET3D(u,2,0,-2)),KRANC_GFOFFSET3D(u,-2,0,2))))))))\n"
+"#else\n"
+"# define PDstandardfdOrder431(u) (PDstandardfdOrder431_impl(u,p1o144dxdz,cdj,cdk))\n"
+"static CCTK_REAL_VEC PDstandardfdOrder431_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o144dxdz, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL_VEC PDstandardfdOrder431_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o144dxdz, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+" return kmul(p1o144dxdz,kadd(KRANC_GFOFFSET3D(u,-2,0,-2),kadd(KRANC_GFOFFSET3D(u,2,0,2),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,0,1),KRANC_GFOFFSET3D(u,1,0,-1)),ToReal(-64),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,0,-2),kadd(KRANC_GFOFFSET3D(u,1,0,2),kadd(KRANC_GFOFFSET3D(u,-2,0,-1),KRANC_GFOFFSET3D(u,2,0,1)))),ToReal(-8),ksub(ksub(kmadd(kadd(KRANC_GFOFFSET3D(u,-1,0,2),kadd(KRANC_GFOFFSET3D(u,1,0,-2),kadd(KRANC_GFOFFSET3D(u,-2,0,1),KRANC_GFOFFSET3D(u,2,0,-1)))),ToReal(8),kmul(kadd(KRANC_GFOFFSET3D(u,-1,0,-1),KRANC_GFOFFSET3D(u,1,0,1)),ToReal(64))),KRANC_GFOFFSET3D(u,2,0,-2)),KRANC_GFOFFSET3D(u,-2,0,2)))))));\n"
+"}\n"
+"#endif\n"
+"\n"
+"#ifndef KRANC_DIFF_FUNCTIONS\n"
+"# define PDstandardfdOrder432(u) (kmul(p1o144dydz,kadd(KRANC_GFOFFSET3D(u,0,-2,-2),kadd(KRANC_GFOFFSET3D(u,0,2,2),kmadd(kadd(KRANC_GFOFFSET3D(u,0,-1,1),KRANC_GFOFFSET3D(u,0,1,-1)),ToReal(-64),kmadd(kadd(KRANC_GFOFFSET3D(u,0,-1,-2),kadd(KRANC_GFOFFSET3D(u,0,1,2),kadd(KRANC_GFOFFSET3D(u,0,-2,-1),KRANC_GFOFFSET3D(u,0,2,1)))),ToReal(-8),ksub(ksub(kmadd(kadd(KRANC_GFOFFSET3D(u,0,-1,2),kadd(KRANC_GFOFFSET3D(u,0,1,-2),kadd(KRANC_GFOFFSET3D(u,0,-2,1),KRANC_GFOFFSET3D(u,0,2,-1)))),ToReal(8),kmul(kadd(KRANC_GFOFFSET3D(u,0,-1,-1),KRANC_GFOFFSET3D(u,0,1,1)),ToReal(64))),KRANC_GFOFFSET3D(u,0,2,-2)),KRANC_GFOFFSET3D(u,0,-2,2))))))))\n"
+"#else\n"
+"# define PDstandardfdOrder432(u) (PDstandardfdOrder432_impl(u,p1o144dydz,cdj,cdk))\n"
+"static CCTK_REAL_VEC PDstandardfdOrder432_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o144dydz, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n"
+"static CCTK_REAL_VEC PDstandardfdOrder432_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o144dydz, ptrdiff_t const cdj, ptrdiff_t const cdk)\n"
+"{\n"
+" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n"
+" return kmul(p1o144dydz,kadd(KRANC_GFOFFSET3D(u,0,-2,-2),kadd(KRANC_GFOFFSET3D(u,0,2,2),kmadd(kadd(KRANC_GFOFFSET3D(u,0,-1,1),KRANC_GFOFFSET3D(u,0,1,-1)),ToReal(-64),kmadd(kadd(KRANC_GFOFFSET3D(u,0,-1,-2),kadd(KRANC_GFOFFSET3D(u,0,1,2),kadd(KRANC_GFOFFSET3D(u,0,-2,-1),KRANC_GFOFFSET3D(u,0,2,1)))),ToReal(-8),ksub(ksub(kmadd(kadd(KRANC_GFOFFSET3D(u,0,-1,2),kadd(KRANC_GFOFFSET3D(u,0,1,-2),kadd(KRANC_GFOFFSET3D(u,0,-2,1),KRANC_GFOFFSET3D(u,0,2,-1)))),ToReal(8),kmul(kadd(KRANC_GFOFFSET3D(u,0,-1,-1),KRANC_GFOFFSET3D(u,0,1,1)),ToReal(64))),KRANC_GFOFFSET3D(u,0,2,-2)),KRANC_GFOFFSET3D(u,0,-2,2)))))));\n"
+"}\n"
+"#endif\n"
+"\n"
+""
+;
diff --git a/Examples/EMScript/src/EM_constraints.cc b/Examples/EMScript/src/EM_constraints.cc
index 2387960..fc8e97e 100644
--- a/Examples/EMScript/src/EM_constraints.cc
+++ b/Examples/EMScript/src/EM_constraints.cc
@@ -14,13 +14,15 @@
#include "Differencing.h"
#include "cctk_Loop.h"
#include "loopcontrol.h"
+#include "OpenCLRunTime.h"
+#include "vectors.h"
/* Define macros used in calculations */
#define INITVALUE (42)
#define QAD(x) (SQR(SQR(x)))
-#define INV(x) ((1.0) / (x))
-#define SQR(x) ((x) * (x))
-#define CUB(x) ((x) * (x) * (x))
+#define INV(x) (kdiv(ToReal(1.0),x))
+#define SQR(x) (kmul(x,x))
+#define CUB(x) (kmul(x,SQR(x)))
extern "C" void EM_constraints_SelectBCs(CCTK_ARGUMENTS)
{
@@ -28,12 +30,12 @@ extern "C" void EM_constraints_SelectBCs(CCTK_ARGUMENTS)
DECLARE_CCTK_PARAMETERS;
CCTK_INT ierr = 0;
- ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "EMScript::CB_group","flat");
+ ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "My_New_Implementation::CB_group","flat");
if (ierr < 0)
- CCTK_WARN(1, "Failed to register flat BC for EMScript::CB_group.");
- ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "EMScript::CEl_group","flat");
+ CCTK_WARN(1, "Failed to register flat BC for My_New_Implementation::CB_group.");
+ ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "My_New_Implementation::CEl_group","flat");
if (ierr < 0)
- CCTK_WARN(1, "Failed to register flat BC for EMScript::CEl_group.");
+ CCTK_WARN(1, "Failed to register flat BC for My_New_Implementation::CEl_group.");
return;
}
@@ -42,122 +44,275 @@ static void EM_constraints_Body(cGH const * restrict const cctkGH, int const dir
DECLARE_CCTK_ARGUMENTS;
DECLARE_CCTK_PARAMETERS;
+ char const * const source =
+ "\n"
+ "/* Declare finite differencing variables */\n"
+ "\n"
+ "/* Include user-supplied include files */\n"
+ "\n"
+ "/* Initialise finite differencing variables */\n"
+ "ptrdiff_t const di = 1;\n"
+ "ptrdiff_t const dj = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n"
+ "ptrdiff_t const dk = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n"
+ "ptrdiff_t const cdi = sizeof(CCTK_REAL) * di;\n"
+ "ptrdiff_t const cdj = sizeof(CCTK_REAL) * dj;\n"
+ "ptrdiff_t const cdk = sizeof(CCTK_REAL) * dk;\n"
+ "CCTK_REAL_VEC const dx = ToReal(CCTK_DELTA_SPACE(0));\n"
+ "CCTK_REAL_VEC const dy = ToReal(CCTK_DELTA_SPACE(1));\n"
+ "CCTK_REAL_VEC const dz = ToReal(CCTK_DELTA_SPACE(2));\n"
+ "CCTK_REAL_VEC const dt = ToReal(CCTK_DELTA_TIME);\n"
+ "CCTK_REAL_VEC const t = ToReal(cctk_time);\n"
+ "CCTK_REAL_VEC const dxi = INV(dx);\n"
+ "CCTK_REAL_VEC const dyi = INV(dy);\n"
+ "CCTK_REAL_VEC const dzi = INV(dz);\n"
+ "CCTK_REAL_VEC const khalf = ToReal(0.5);\n"
+ "CCTK_REAL_VEC const kthird = ToReal(1.0/3.0);\n"
+ "CCTK_REAL_VEC const ktwothird = ToReal(2.0/3.0);\n"
+ "CCTK_REAL_VEC const kfourthird = ToReal(4.0/3.0);\n"
+ "CCTK_REAL_VEC const keightthird = ToReal(8.0/3.0);\n"
+ "CCTK_REAL_VEC const hdxi = kmul(ToReal(0.5), dxi);\n"
+ "CCTK_REAL_VEC const hdyi = kmul(ToReal(0.5), dyi);\n"
+ "CCTK_REAL_VEC const hdzi = kmul(ToReal(0.5), dzi);\n"
+ "\n"
+ "/* Initialize predefined quantities */\n"
+ "CCTK_REAL_VEC const p1o12dx = kmul(INV(dx),ToReal(0.0833333333333333333333333333333));\n"
+ "CCTK_REAL_VEC const p1o12dy = kmul(INV(dy),ToReal(0.0833333333333333333333333333333));\n"
+ "CCTK_REAL_VEC const p1o12dz = kmul(INV(dz),ToReal(0.0833333333333333333333333333333));\n"
+ "CCTK_REAL_VEC const p1o144dxdy = kmul(INV(dx),kmul(INV(dy),ToReal(0.00694444444444444444444444444444)));\n"
+ "CCTK_REAL_VEC const p1o144dxdz = kmul(INV(dx),kmul(INV(dz),ToReal(0.00694444444444444444444444444444)));\n"
+ "CCTK_REAL_VEC const p1o144dydz = kmul(INV(dy),kmul(INV(dz),ToReal(0.00694444444444444444444444444444)));\n"
+ "CCTK_REAL_VEC const p1o2dx = kmul(INV(dx),ToReal(0.5));\n"
+ "CCTK_REAL_VEC const p1o2dy = kmul(INV(dy),ToReal(0.5));\n"
+ "CCTK_REAL_VEC const p1o2dz = kmul(INV(dz),ToReal(0.5));\n"
+ "CCTK_REAL_VEC const p1o4dxdy = kmul(INV(dx),kmul(INV(dy),ToReal(0.25)));\n"
+ "CCTK_REAL_VEC const p1o4dxdz = kmul(INV(dx),kmul(INV(dz),ToReal(0.25)));\n"
+ "CCTK_REAL_VEC const p1o4dydz = kmul(INV(dy),kmul(INV(dz),ToReal(0.25)));\n"
+ "CCTK_REAL_VEC const p1odx2 = INV(SQR(dx));\n"
+ "CCTK_REAL_VEC const p1ody2 = INV(SQR(dy));\n"
+ "CCTK_REAL_VEC const p1odz2 = INV(SQR(dz));\n"
+ "CCTK_REAL_VEC const pm1o12dx2 = kmul(INV(SQR(dx)),ToReal(-0.0833333333333333333333333333333));\n"
+ "CCTK_REAL_VEC const pm1o12dy2 = kmul(INV(SQR(dy)),ToReal(-0.0833333333333333333333333333333));\n"
+ "CCTK_REAL_VEC const pm1o12dz2 = kmul(INV(SQR(dz)),ToReal(-0.0833333333333333333333333333333));\n"
+ "\n"
+ "/* Jacobian variable pointers */\n"
+ "bool const use_jacobian = (!CCTK_IsFunctionAliased(\"MultiPatch_GetMap\") || MultiPatch_GetMap(cctkGH) != jacobian_identity_map)\n"
+ " && strlen(jacobian_group) > 0;\n"
+ "if (use_jacobian && strlen(jacobian_derivative_group) == 0)\n"
+ "{\n"
+ " CCTK_WARN (1, \"GenericFD::jacobian_group and GenericFD::jacobian_derivative_group must both be set to valid group names\");\n"
+ "}\n"
+ "\n"
+ "CCTK_REAL const *restrict jacobian_ptrs[9];\n"
+ "if (use_jacobian) GenericFD_GroupDataPointers(cctkGH, jacobian_group,\n"
+ " 9, jacobian_ptrs);\n"
+ "\n"
+ "CCTK_REAL const *restrict const J11 = use_jacobian ? jacobian_ptrs[0] : 0;\n"
+ "CCTK_REAL const *restrict const J12 = use_jacobian ? jacobian_ptrs[1] : 0;\n"
+ "CCTK_REAL const *restrict const J13 = use_jacobian ? jacobian_ptrs[2] : 0;\n"
+ "CCTK_REAL const *restrict const J21 = use_jacobian ? jacobian_ptrs[3] : 0;\n"
+ "CCTK_REAL const *restrict const J22 = use_jacobian ? jacobian_ptrs[4] : 0;\n"
+ "CCTK_REAL const *restrict const J23 = use_jacobian ? jacobian_ptrs[5] : 0;\n"
+ "CCTK_REAL const *restrict const J31 = use_jacobian ? jacobian_ptrs[6] : 0;\n"
+ "CCTK_REAL const *restrict const J32 = use_jacobian ? jacobian_ptrs[7] : 0;\n"
+ "CCTK_REAL const *restrict const J33 = use_jacobian ? jacobian_ptrs[8] : 0;\n"
+ "\n"
+ "CCTK_REAL const *restrict jacobian_derivative_ptrs[18];\n"
+ "if (use_jacobian) GenericFD_GroupDataPointers(cctkGH, jacobian_derivative_group,\n"
+ " 18, jacobian_derivative_ptrs);\n"
+ "\n"
+ "CCTK_REAL const *restrict const dJ111 = use_jacobian ? jacobian_derivative_ptrs[0] : 0;\n"
+ "CCTK_REAL const *restrict const dJ112 = use_jacobian ? jacobian_derivative_ptrs[1] : 0;\n"
+ "CCTK_REAL const *restrict const dJ113 = use_jacobian ? jacobian_derivative_ptrs[2] : 0;\n"
+ "CCTK_REAL const *restrict const dJ122 = use_jacobian ? jacobian_derivative_ptrs[3] : 0;\n"
+ "CCTK_REAL const *restrict const dJ123 = use_jacobian ? jacobian_derivative_ptrs[4] : 0;\n"
+ "CCTK_REAL const *restrict const dJ133 = use_jacobian ? jacobian_derivative_ptrs[5] : 0;\n"
+ "CCTK_REAL const *restrict const dJ211 = use_jacobian ? jacobian_derivative_ptrs[6] : 0;\n"
+ "CCTK_REAL const *restrict const dJ212 = use_jacobian ? jacobian_derivative_ptrs[7] : 0;\n"
+ "CCTK_REAL const *restrict const dJ213 = use_jacobian ? jacobian_derivative_ptrs[8] : 0;\n"
+ "CCTK_REAL const *restrict const dJ222 = use_jacobian ? jacobian_derivative_ptrs[9] : 0;\n"
+ "CCTK_REAL const *restrict const dJ223 = use_jacobian ? jacobian_derivative_ptrs[10] : 0;\n"
+ "CCTK_REAL const *restrict const dJ233 = use_jacobian ? jacobian_derivative_ptrs[11] : 0;\n"
+ "CCTK_REAL const *restrict const dJ311 = use_jacobian ? jacobian_derivative_ptrs[12] : 0;\n"
+ "CCTK_REAL const *restrict const dJ312 = use_jacobian ? jacobian_derivative_ptrs[13] : 0;\n"
+ "CCTK_REAL const *restrict const dJ313 = use_jacobian ? jacobian_derivative_ptrs[14] : 0;\n"
+ "CCTK_REAL const *restrict const dJ322 = use_jacobian ? jacobian_derivative_ptrs[15] : 0;\n"
+ "CCTK_REAL const *restrict const dJ323 = use_jacobian ? jacobian_derivative_ptrs[16] : 0;\n"
+ "CCTK_REAL const *restrict const dJ333 = use_jacobian ? jacobian_derivative_ptrs[17] : 0;\n"
+ "\n"
+ "/* Assign local copies of arrays functions */\n"
+ "\n"
+ "\n"
+ "\n"
+ "/* Calculate temporaries and arrays functions */\n"
+ "\n"
+ "/* Copy local copies back to grid functions */\n"
+ "\n"
+ "/* Loop over the grid points */\n"
+ "#pragma omp parallel\n"
+ "LC_LOOP3VEC (EM_constraints,\n"
+ " i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2],\n"
+ " cctk_lsh[0],cctk_lsh[1],cctk_lsh[2],\n"
+ " CCTK_REAL_VEC_SIZE)\n"
+ "{\n"
+ " ptrdiff_t const index = di*i + dj*j + dk*k;\n"
+ " \n"
+ " /* Assign local copies of grid functions */\n"
+ " \n"
+ " CCTK_REAL_VEC B1L = vec_load(B1[index]);\n"
+ " CCTK_REAL_VEC B2L = vec_load(B2[index]);\n"
+ " CCTK_REAL_VEC B3L = vec_load(B3[index]);\n"
+ " CCTK_REAL_VEC El1L = vec_load(El1[index]);\n"
+ " CCTK_REAL_VEC El2L = vec_load(El2[index]);\n"
+ " CCTK_REAL_VEC El3L = vec_load(El3[index]);\n"
+ " \n"
+ " \n"
+ " CCTK_REAL_VEC J11L, J12L, J13L, J21L, J22L, J23L, J31L, J32L, J33L;\n"
+ " \n"
+ " if (use_jacobian)\n"
+ " {\n"
+ " J11L = vec_load(J11[index]);\n"
+ " J12L = vec_load(J12[index]);\n"
+ " J13L = vec_load(J13[index]);\n"
+ " J21L = vec_load(J21[index]);\n"
+ " J22L = vec_load(J22[index]);\n"
+ " J23L = vec_load(J23[index]);\n"
+ " J31L = vec_load(J31[index]);\n"
+ " J32L = vec_load(J32[index]);\n"
+ " J33L = vec_load(J33[index]);\n"
+ " }\n"
+ " \n"
+ " /* Include user supplied include files */\n"
+ " \n"
+ " /* Precompute derivatives */\n"
+ " CCTK_REAL_VEC PDstandard1B1;\n"
+ " CCTK_REAL_VEC PDstandard2B1;\n"
+ " CCTK_REAL_VEC PDstandard3B1;\n"
+ " CCTK_REAL_VEC PDstandard1B2;\n"
+ " CCTK_REAL_VEC PDstandard2B2;\n"
+ " CCTK_REAL_VEC PDstandard3B2;\n"
+ " CCTK_REAL_VEC PDstandard1B3;\n"
+ " CCTK_REAL_VEC PDstandard2B3;\n"
+ " CCTK_REAL_VEC PDstandard3B3;\n"
+ " CCTK_REAL_VEC PDstandard1El1;\n"
+ " CCTK_REAL_VEC PDstandard2El1;\n"
+ " CCTK_REAL_VEC PDstandard3El1;\n"
+ " CCTK_REAL_VEC PDstandard1El2;\n"
+ " CCTK_REAL_VEC PDstandard2El2;\n"
+ " CCTK_REAL_VEC PDstandard3El2;\n"
+ " CCTK_REAL_VEC PDstandard1El3;\n"
+ " CCTK_REAL_VEC PDstandard2El3;\n"
+ " CCTK_REAL_VEC PDstandard3El3;\n"
+ " \n"
+ " switch(fdOrder)\n"
+ " {\n"
+ " case 2:\n"
+ " PDstandard1B1 = PDstandardfdOrder21(&B1[index]);\n"
+ " PDstandard2B1 = PDstandardfdOrder22(&B1[index]);\n"
+ " PDstandard3B1 = PDstandardfdOrder23(&B1[index]);\n"
+ " PDstandard1B2 = PDstandardfdOrder21(&B2[index]);\n"
+ " PDstandard2B2 = PDstandardfdOrder22(&B2[index]);\n"
+ " PDstandard3B2 = PDstandardfdOrder23(&B2[index]);\n"
+ " PDstandard1B3 = PDstandardfdOrder21(&B3[index]);\n"
+ " PDstandard2B3 = PDstandardfdOrder22(&B3[index]);\n"
+ " PDstandard3B3 = PDstandardfdOrder23(&B3[index]);\n"
+ " PDstandard1El1 = PDstandardfdOrder21(&El1[index]);\n"
+ " PDstandard2El1 = PDstandardfdOrder22(&El1[index]);\n"
+ " PDstandard3El1 = PDstandardfdOrder23(&El1[index]);\n"
+ " PDstandard1El2 = PDstandardfdOrder21(&El2[index]);\n"
+ " PDstandard2El2 = PDstandardfdOrder22(&El2[index]);\n"
+ " PDstandard3El2 = PDstandardfdOrder23(&El2[index]);\n"
+ " PDstandard1El3 = PDstandardfdOrder21(&El3[index]);\n"
+ " PDstandard2El3 = PDstandardfdOrder22(&El3[index]);\n"
+ " PDstandard3El3 = PDstandardfdOrder23(&El3[index]);\n"
+ " break;\n"
+ " \n"
+ " case 4:\n"
+ " PDstandard1B1 = PDstandardfdOrder41(&B1[index]);\n"
+ " PDstandard2B1 = PDstandardfdOrder42(&B1[index]);\n"
+ " PDstandard3B1 = PDstandardfdOrder43(&B1[index]);\n"
+ " PDstandard1B2 = PDstandardfdOrder41(&B2[index]);\n"
+ " PDstandard2B2 = PDstandardfdOrder42(&B2[index]);\n"
+ " PDstandard3B2 = PDstandardfdOrder43(&B2[index]);\n"
+ " PDstandard1B3 = PDstandardfdOrder41(&B3[index]);\n"
+ " PDstandard2B3 = PDstandardfdOrder42(&B3[index]);\n"
+ " PDstandard3B3 = PDstandardfdOrder43(&B3[index]);\n"
+ " PDstandard1El1 = PDstandardfdOrder41(&El1[index]);\n"
+ " PDstandard2El1 = PDstandardfdOrder42(&El1[index]);\n"
+ " PDstandard3El1 = PDstandardfdOrder43(&El1[index]);\n"
+ " PDstandard1El2 = PDstandardfdOrder41(&El2[index]);\n"
+ " PDstandard2El2 = PDstandardfdOrder42(&El2[index]);\n"
+ " PDstandard3El2 = PDstandardfdOrder43(&El2[index]);\n"
+ " PDstandard1El3 = PDstandardfdOrder41(&El3[index]);\n"
+ " PDstandard2El3 = PDstandardfdOrder42(&El3[index]);\n"
+ " PDstandard3El3 = PDstandardfdOrder43(&El3[index]);\n"
+ " break;\n"
+ " }\n"
+ " \n"
+ " /* Calculate temporaries and grid functions */\n"
+ " CCTK_REAL_VEC JacPDstandard1B1;\n"
+ " CCTK_REAL_VEC JacPDstandard1El1;\n"
+ " CCTK_REAL_VEC JacPDstandard2B2;\n"
+ " CCTK_REAL_VEC JacPDstandard2El2;\n"
+ " CCTK_REAL_VEC JacPDstandard3B3;\n"
+ " CCTK_REAL_VEC JacPDstandard3El3;\n"
+ " \n"
+ " if (use_jacobian)\n"
+ " {\n"
+ " JacPDstandard1B1 = \n"
+ " kmadd(J11L,PDstandard1B1,kmadd(J21L,PDstandard2B1,kmul(J31L,PDstandard3B1)));\n"
+ " \n"
+ " JacPDstandard1El1 = \n"
+ " kmadd(J11L,PDstandard1El1,kmadd(J21L,PDstandard2El1,kmul(J31L,PDstandard3El1)));\n"
+ " \n"
+ " JacPDstandard2B2 = \n"
+ " kmadd(J12L,PDstandard1B2,kmadd(J22L,PDstandard2B2,kmul(J32L,PDstandard3B2)));\n"
+ " \n"
+ " JacPDstandard2El2 = \n"
+ " kmadd(J12L,PDstandard1El2,kmadd(J22L,PDstandard2El2,kmul(J32L,PDstandard3El2)));\n"
+ " \n"
+ " JacPDstandard3B3 = \n"
+ " kmadd(J13L,PDstandard1B3,kmadd(J23L,PDstandard2B3,kmul(J33L,PDstandard3B3)));\n"
+ " \n"
+ " JacPDstandard3El3 = \n"
+ " kmadd(J13L,PDstandard1El3,kmadd(J23L,PDstandard2El3,kmul(J33L,PDstandard3El3)));\n"
+ " }\n"
+ " else\n"
+ " {\n"
+ " JacPDstandard1B1 = PDstandard1B1;\n"
+ " \n"
+ " JacPDstandard1El1 = PDstandard1El1;\n"
+ " \n"
+ " JacPDstandard2B2 = PDstandard2B2;\n"
+ " \n"
+ " JacPDstandard2El2 = PDstandard2El2;\n"
+ " \n"
+ " JacPDstandard3B3 = PDstandard3B3;\n"
+ " \n"
+ " JacPDstandard3El3 = PDstandard3El3;\n"
+ " }\n"
+ " \n"
+ " CCTK_REAL_VEC CElL = \n"
+ " kadd(JacPDstandard1El1,kadd(JacPDstandard2El2,JacPDstandard3El3));\n"
+ " \n"
+ " CCTK_REAL_VEC CBL = \n"
+ " kadd(JacPDstandard1B1,kadd(JacPDstandard2B2,JacPDstandard3B3));\n"
+ " \n"
+ " /* Copy local copies back to grid functions */\n"
+ " vec_store_nta_partial(CB[index],CBL);\n"
+ " vec_store_nta_partial(CEl[index],CElL);\n"
+ "}\n"
+ "LC_ENDLOOP3VEC (EM_constraints);\n"
+ ""
+ ;
- /* Declare finite differencing variables */
+ char const * const groups[] = {"My_New_Implementation::B_group","My_New_Implementation::CB_group","My_New_Implementation::CEl_group","My_New_Implementation::El_group",NULL};
- /* Include user-supplied include files */
+ static struct OpenCLKernel * kernel = NULL;
+ char const * const sources[] = {differencing, source, NULL};
+ OpenCLRunTime_CallKernel (cctkGH, CCTK_THORNSTRING, "EM_constraints",
+ sources, groups, NULL, NULL, NULL, -1,
+ imin, imax, &kernel);
- /* Initialise finite differencing variables */
- ptrdiff_t const di = 1;
- ptrdiff_t const dj = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);
- ptrdiff_t const dk = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);
- ptrdiff_t const cdi = sizeof(CCTK_REAL) * di;
- ptrdiff_t const cdj = sizeof(CCTK_REAL) * dj;
- ptrdiff_t const cdk = sizeof(CCTK_REAL) * dk;
- CCTK_REAL const dx = ToReal(CCTK_DELTA_SPACE(0));
- CCTK_REAL const dy = ToReal(CCTK_DELTA_SPACE(1));
- CCTK_REAL const dz = ToReal(CCTK_DELTA_SPACE(2));
- CCTK_REAL const dt = ToReal(CCTK_DELTA_TIME);
- CCTK_REAL const t = ToReal(cctk_time);
- CCTK_REAL const dxi = INV(dx);
- CCTK_REAL const dyi = INV(dy);
- CCTK_REAL const dzi = INV(dz);
- CCTK_REAL const khalf = 0.5;
- CCTK_REAL const kthird = 1/3.0;
- CCTK_REAL const ktwothird = 2.0/3.0;
- CCTK_REAL const kfourthird = 4.0/3.0;
- CCTK_REAL const keightthird = 8.0/3.0;
- CCTK_REAL const hdxi = 0.5 * dxi;
- CCTK_REAL const hdyi = 0.5 * dyi;
- CCTK_REAL const hdzi = 0.5 * dzi;
-
- /* Initialize predefined quantities */
- CCTK_REAL const p1o12dx = 0.0833333333333333333333333333333*INV(dx);
- CCTK_REAL const p1o12dy = 0.0833333333333333333333333333333*INV(dy);
- CCTK_REAL const p1o12dz = 0.0833333333333333333333333333333*INV(dz);
- CCTK_REAL const p1o144dxdy = 0.00694444444444444444444444444444*INV(dx)*INV(dy);
- CCTK_REAL const p1o144dxdz = 0.00694444444444444444444444444444*INV(dx)*INV(dz);
- CCTK_REAL const p1o144dydz = 0.00694444444444444444444444444444*INV(dy)*INV(dz);
- CCTK_REAL const p1o2dx = 0.5*INV(dx);
- CCTK_REAL const p1o2dy = 0.5*INV(dy);
- CCTK_REAL const p1o2dz = 0.5*INV(dz);
- CCTK_REAL const p1o4dxdy = 0.25*INV(dx)*INV(dy);
- CCTK_REAL const p1o4dxdz = 0.25*INV(dx)*INV(dz);
- CCTK_REAL const p1o4dydz = 0.25*INV(dy)*INV(dz);
- CCTK_REAL const p1odx2 = INV(SQR(dx));
- CCTK_REAL const p1ody2 = INV(SQR(dy));
- CCTK_REAL const p1odz2 = INV(SQR(dz));
- CCTK_REAL const pm1o12dx2 = -0.0833333333333333333333333333333*INV(SQR(dx));
- CCTK_REAL const pm1o12dy2 = -0.0833333333333333333333333333333*INV(SQR(dy));
- CCTK_REAL const pm1o12dz2 = -0.0833333333333333333333333333333*INV(SQR(dz));
-
- /* Assign local copies of arrays functions */
-
-
-
- /* Calculate temporaries and arrays functions */
-
- /* Copy local copies back to grid functions */
-
- /* Loop over the grid points */
- #pragma omp parallel
- CCTK_LOOP3 (EM_constraints,
- i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2],
- cctk_lsh[0],cctk_lsh[1],cctk_lsh[2])
- {
- ptrdiff_t const index = di*i + dj*j + dk*k;
-
- /* Assign local copies of grid functions */
-
- CCTK_REAL B1L = B1[index];
- CCTK_REAL B2L = B2[index];
- CCTK_REAL B3L = B3[index];
- CCTK_REAL El1L = El1[index];
- CCTK_REAL El2L = El2[index];
- CCTK_REAL El3L = El3[index];
-
-
- /* Include user supplied include files */
-
- /* Precompute derivatives */
- CCTK_REAL PDstandard1B1;
- CCTK_REAL PDstandard2B2;
- CCTK_REAL PDstandard3B3;
- CCTK_REAL PDstandard1El1;
- CCTK_REAL PDstandard2El2;
- CCTK_REAL PDstandard3El3;
-
- switch(fdOrder)
- {
- case 2:
- PDstandard1B1 = PDstandardfdOrder21(&B1[index]);
- PDstandard2B2 = PDstandardfdOrder22(&B2[index]);
- PDstandard3B3 = PDstandardfdOrder23(&B3[index]);
- PDstandard1El1 = PDstandardfdOrder21(&El1[index]);
- PDstandard2El2 = PDstandardfdOrder22(&El2[index]);
- PDstandard3El3 = PDstandardfdOrder23(&El3[index]);
- break;
-
- case 4:
- PDstandard1B1 = PDstandardfdOrder41(&B1[index]);
- PDstandard2B2 = PDstandardfdOrder42(&B2[index]);
- PDstandard3B3 = PDstandardfdOrder43(&B3[index]);
- PDstandard1El1 = PDstandardfdOrder41(&El1[index]);
- PDstandard2El2 = PDstandardfdOrder42(&El2[index]);
- PDstandard3El3 = PDstandardfdOrder43(&El3[index]);
- break;
- }
-
- /* Calculate temporaries and grid functions */
- CCTK_REAL CElL = PDstandard1El1 + PDstandard2El2 + PDstandard3El3;
-
- CCTK_REAL CBL = PDstandard1B1 + PDstandard2B2 + PDstandard3B3;
-
- /* Copy local copies back to grid functions */
- CB[index] = CBL;
- CEl[index] = CElL;
- }
- CCTK_ENDLOOP3 (EM_constraints);
}
extern "C" void EM_constraints(CCTK_ARGUMENTS)
@@ -176,7 +331,7 @@ extern "C" void EM_constraints(CCTK_ARGUMENTS)
return;
}
- const char *groups[] = {"EMScript::B_group","EMScript::CB_group","EMScript::CEl_group","EMScript::El_group"};
+ const char *groups[] = {"My_New_Implementation::B_group","My_New_Implementation::CB_group","My_New_Implementation::CEl_group","My_New_Implementation::El_group"};
GenericFD_AssertGroupStorage(cctkGH, "EM_constraints", 4, groups);
switch(fdOrder)
diff --git a/Examples/EMScript/src/EM_energy.cc b/Examples/EMScript/src/EM_energy.cc
index 50e9f94..5a7875c 100644
--- a/Examples/EMScript/src/EM_energy.cc
+++ b/Examples/EMScript/src/EM_energy.cc
@@ -14,13 +14,15 @@
#include "Differencing.h"
#include "cctk_Loop.h"
#include "loopcontrol.h"
+#include "OpenCLRunTime.h"
+#include "vectors.h"
/* Define macros used in calculations */
#define INITVALUE (42)
#define QAD(x) (SQR(SQR(x)))
-#define INV(x) ((1.0) / (x))
-#define SQR(x) ((x) * (x))
-#define CUB(x) ((x) * (x) * (x))
+#define INV(x) (kdiv(ToReal(1.0),x))
+#define SQR(x) (kmul(x,x))
+#define CUB(x) (kmul(x,SQR(x)))
extern "C" void EM_energy_SelectBCs(CCTK_ARGUMENTS)
{
@@ -28,9 +30,9 @@ extern "C" void EM_energy_SelectBCs(CCTK_ARGUMENTS)
DECLARE_CCTK_PARAMETERS;
CCTK_INT ierr = 0;
- ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "EMScript::rho_group","flat");
+ ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "My_New_Implementation::rho_group","flat");
if (ierr < 0)
- CCTK_WARN(1, "Failed to register flat BC for EMScript::rho_group.");
+ CCTK_WARN(1, "Failed to register flat BC for My_New_Implementation::rho_group.");
return;
}
@@ -39,102 +41,161 @@ static void EM_energy_Body(cGH const * restrict const cctkGH, int const dir, int
DECLARE_CCTK_ARGUMENTS;
DECLARE_CCTK_PARAMETERS;
+ char const * const source =
+ "\n"
+ "/* Declare finite differencing variables */\n"
+ "\n"
+ "/* Include user-supplied include files */\n"
+ "\n"
+ "/* Initialise finite differencing variables */\n"
+ "ptrdiff_t const di = 1;\n"
+ "ptrdiff_t const dj = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n"
+ "ptrdiff_t const dk = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n"
+ "ptrdiff_t const cdi = sizeof(CCTK_REAL) * di;\n"
+ "ptrdiff_t const cdj = sizeof(CCTK_REAL) * dj;\n"
+ "ptrdiff_t const cdk = sizeof(CCTK_REAL) * dk;\n"
+ "CCTK_REAL_VEC const dx = ToReal(CCTK_DELTA_SPACE(0));\n"
+ "CCTK_REAL_VEC const dy = ToReal(CCTK_DELTA_SPACE(1));\n"
+ "CCTK_REAL_VEC const dz = ToReal(CCTK_DELTA_SPACE(2));\n"
+ "CCTK_REAL_VEC const dt = ToReal(CCTK_DELTA_TIME);\n"
+ "CCTK_REAL_VEC const t = ToReal(cctk_time);\n"
+ "CCTK_REAL_VEC const dxi = INV(dx);\n"
+ "CCTK_REAL_VEC const dyi = INV(dy);\n"
+ "CCTK_REAL_VEC const dzi = INV(dz);\n"
+ "CCTK_REAL_VEC const khalf = ToReal(0.5);\n"
+ "CCTK_REAL_VEC const kthird = ToReal(1.0/3.0);\n"
+ "CCTK_REAL_VEC const ktwothird = ToReal(2.0/3.0);\n"
+ "CCTK_REAL_VEC const kfourthird = ToReal(4.0/3.0);\n"
+ "CCTK_REAL_VEC const keightthird = ToReal(8.0/3.0);\n"
+ "CCTK_REAL_VEC const hdxi = kmul(ToReal(0.5), dxi);\n"
+ "CCTK_REAL_VEC const hdyi = kmul(ToReal(0.5), dyi);\n"
+ "CCTK_REAL_VEC const hdzi = kmul(ToReal(0.5), dzi);\n"
+ "\n"
+ "/* Initialize predefined quantities */\n"
+ "CCTK_REAL_VEC const p1o12dx = kmul(INV(dx),ToReal(0.0833333333333333333333333333333));\n"
+ "CCTK_REAL_VEC const p1o12dy = kmul(INV(dy),ToReal(0.0833333333333333333333333333333));\n"
+ "CCTK_REAL_VEC const p1o12dz = kmul(INV(dz),ToReal(0.0833333333333333333333333333333));\n"
+ "CCTK_REAL_VEC const p1o144dxdy = kmul(INV(dx),kmul(INV(dy),ToReal(0.00694444444444444444444444444444)));\n"
+ "CCTK_REAL_VEC const p1o144dxdz = kmul(INV(dx),kmul(INV(dz),ToReal(0.00694444444444444444444444444444)));\n"
+ "CCTK_REAL_VEC const p1o144dydz = kmul(INV(dy),kmul(INV(dz),ToReal(0.00694444444444444444444444444444)));\n"
+ "CCTK_REAL_VEC const p1o2dx = kmul(INV(dx),ToReal(0.5));\n"
+ "CCTK_REAL_VEC const p1o2dy = kmul(INV(dy),ToReal(0.5));\n"
+ "CCTK_REAL_VEC const p1o2dz = kmul(INV(dz),ToReal(0.5));\n"
+ "CCTK_REAL_VEC const p1o4dxdy = kmul(INV(dx),kmul(INV(dy),ToReal(0.25)));\n"
+ "CCTK_REAL_VEC const p1o4dxdz = kmul(INV(dx),kmul(INV(dz),ToReal(0.25)));\n"
+ "CCTK_REAL_VEC const p1o4dydz = kmul(INV(dy),kmul(INV(dz),ToReal(0.25)));\n"
+ "CCTK_REAL_VEC const p1odx2 = INV(SQR(dx));\n"
+ "CCTK_REAL_VEC const p1ody2 = INV(SQR(dy));\n"
+ "CCTK_REAL_VEC const p1odz2 = INV(SQR(dz));\n"
+ "CCTK_REAL_VEC const pm1o12dx2 = kmul(INV(SQR(dx)),ToReal(-0.0833333333333333333333333333333));\n"
+ "CCTK_REAL_VEC const pm1o12dy2 = kmul(INV(SQR(dy)),ToReal(-0.0833333333333333333333333333333));\n"
+ "CCTK_REAL_VEC const pm1o12dz2 = kmul(INV(SQR(dz)),ToReal(-0.0833333333333333333333333333333));\n"
+ "\n"
+ "/* Jacobian variable pointers */\n"
+ "bool const use_jacobian = (!CCTK_IsFunctionAliased(\"MultiPatch_GetMap\") || MultiPatch_GetMap(cctkGH) != jacobian_identity_map)\n"
+ " && strlen(jacobian_group) > 0;\n"
+ "if (use_jacobian && strlen(jacobian_derivative_group) == 0)\n"
+ "{\n"
+ " CCTK_WARN (1, \"GenericFD::jacobian_group and GenericFD::jacobian_derivative_group must both be set to valid group names\");\n"
+ "}\n"
+ "\n"
+ "CCTK_REAL const *restrict jacobian_ptrs[9];\n"
+ "if (use_jacobian) GenericFD_GroupDataPointers(cctkGH, jacobian_group,\n"
+ " 9, jacobian_ptrs);\n"
+ "\n"
+ "CCTK_REAL const *restrict const J11 = use_jacobian ? jacobian_ptrs[0] : 0;\n"
+ "CCTK_REAL const *restrict const J12 = use_jacobian ? jacobian_ptrs[1] : 0;\n"
+ "CCTK_REAL const *restrict const J13 = use_jacobian ? jacobian_ptrs[2] : 0;\n"
+ "CCTK_REAL const *restrict const J21 = use_jacobian ? jacobian_ptrs[3] : 0;\n"
+ "CCTK_REAL const *restrict const J22 = use_jacobian ? jacobian_ptrs[4] : 0;\n"
+ "CCTK_REAL const *restrict const J23 = use_jacobian ? jacobian_ptrs[5] : 0;\n"
+ "CCTK_REAL const *restrict const J31 = use_jacobian ? jacobian_ptrs[6] : 0;\n"
+ "CCTK_REAL const *restrict const J32 = use_jacobian ? jacobian_ptrs[7] : 0;\n"
+ "CCTK_REAL const *restrict const J33 = use_jacobian ? jacobian_ptrs[8] : 0;\n"
+ "\n"
+ "CCTK_REAL const *restrict jacobian_derivative_ptrs[18];\n"
+ "if (use_jacobian) GenericFD_GroupDataPointers(cctkGH, jacobian_derivative_group,\n"
+ " 18, jacobian_derivative_ptrs);\n"
+ "\n"
+ "CCTK_REAL const *restrict const dJ111 = use_jacobian ? jacobian_derivative_ptrs[0] : 0;\n"
+ "CCTK_REAL const *restrict const dJ112 = use_jacobian ? jacobian_derivative_ptrs[1] : 0;\n"
+ "CCTK_REAL const *restrict const dJ113 = use_jacobian ? jacobian_derivative_ptrs[2] : 0;\n"
+ "CCTK_REAL const *restrict const dJ122 = use_jacobian ? jacobian_derivative_ptrs[3] : 0;\n"
+ "CCTK_REAL const *restrict const dJ123 = use_jacobian ? jacobian_derivative_ptrs[4] : 0;\n"
+ "CCTK_REAL const *restrict const dJ133 = use_jacobian ? jacobian_derivative_ptrs[5] : 0;\n"
+ "CCTK_REAL const *restrict const dJ211 = use_jacobian ? jacobian_derivative_ptrs[6] : 0;\n"
+ "CCTK_REAL const *restrict const dJ212 = use_jacobian ? jacobian_derivative_ptrs[7] : 0;\n"
+ "CCTK_REAL const *restrict const dJ213 = use_jacobian ? jacobian_derivative_ptrs[8] : 0;\n"
+ "CCTK_REAL const *restrict const dJ222 = use_jacobian ? jacobian_derivative_ptrs[9] : 0;\n"
+ "CCTK_REAL const *restrict const dJ223 = use_jacobian ? jacobian_derivative_ptrs[10] : 0;\n"
+ "CCTK_REAL const *restrict const dJ233 = use_jacobian ? jacobian_derivative_ptrs[11] : 0;\n"
+ "CCTK_REAL const *restrict const dJ311 = use_jacobian ? jacobian_derivative_ptrs[12] : 0;\n"
+ "CCTK_REAL const *restrict const dJ312 = use_jacobian ? jacobian_derivative_ptrs[13] : 0;\n"
+ "CCTK_REAL const *restrict const dJ313 = use_jacobian ? jacobian_derivative_ptrs[14] : 0;\n"
+ "CCTK_REAL const *restrict const dJ322 = use_jacobian ? jacobian_derivative_ptrs[15] : 0;\n"
+ "CCTK_REAL const *restrict const dJ323 = use_jacobian ? jacobian_derivative_ptrs[16] : 0;\n"
+ "CCTK_REAL const *restrict const dJ333 = use_jacobian ? jacobian_derivative_ptrs[17] : 0;\n"
+ "\n"
+ "/* Assign local copies of arrays functions */\n"
+ "\n"
+ "\n"
+ "\n"
+ "/* Calculate temporaries and arrays functions */\n"
+ "\n"
+ "/* Copy local copies back to grid functions */\n"
+ "\n"
+ "/* Loop over the grid points */\n"
+ "#pragma omp parallel\n"
+ "LC_LOOP3VEC (EM_energy,\n"
+ " i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2],\n"
+ " cctk_lsh[0],cctk_lsh[1],cctk_lsh[2],\n"
+ " CCTK_REAL_VEC_SIZE)\n"
+ "{\n"
+ " ptrdiff_t const index = di*i + dj*j + dk*k;\n"
+ " \n"
+ " /* Assign local copies of grid functions */\n"
+ " \n"
+ " CCTK_REAL_VEC B1L = vec_load(B1[index]);\n"
+ " CCTK_REAL_VEC B2L = vec_load(B2[index]);\n"
+ " CCTK_REAL_VEC B3L = vec_load(B3[index]);\n"
+ " CCTK_REAL_VEC El1L = vec_load(El1[index]);\n"
+ " CCTK_REAL_VEC El2L = vec_load(El2[index]);\n"
+ " CCTK_REAL_VEC El3L = vec_load(El3[index]);\n"
+ " \n"
+ " \n"
+ " \n"
+ " /* Include user supplied include files */\n"
+ " \n"
+ " /* Precompute derivatives */\n"
+ " \n"
+ " switch(fdOrder)\n"
+ " {\n"
+ " case 2:\n"
+ " break;\n"
+ " \n"
+ " case 4:\n"
+ " break;\n"
+ " }\n"
+ " \n"
+ " /* Calculate temporaries and grid functions */\n"
+ " CCTK_REAL_VEC rhoL = \n"
+ " kmul(kadd(SQR(B1L),kadd(SQR(B2L),kadd(SQR(B3L),kadd(SQR(El1L),kadd(SQR(El2L),SQR(El3L)))))),ToReal(0.5));\n"
+ " \n"
+ " /* Copy local copies back to grid functions */\n"
+ " vec_store_nta_partial(rho[index],rhoL);\n"
+ "}\n"
+ "LC_ENDLOOP3VEC (EM_energy);\n"
+ ""
+ ;
+
+ char const * const groups[] = {"My_New_Implementation::B_group","My_New_Implementation::El_group","My_New_Implementation::rho_group",NULL};
+
+ static struct OpenCLKernel * kernel = NULL;
+ char const * const sources[] = {differencing, source, NULL};
+ OpenCLRunTime_CallKernel (cctkGH, CCTK_THORNSTRING, "EM_energy",
+ sources, groups, NULL, NULL, NULL, -1,
+ imin, imax, &kernel);
- /* Declare finite differencing variables */
-
- /* 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 p1o12dx = 0.0833333333333333333333333333333*INV(dx);
- CCTK_REAL const p1o12dy = 0.0833333333333333333333333333333*INV(dy);
- CCTK_REAL const p1o12dz = 0.0833333333333333333333333333333*INV(dz);
- CCTK_REAL const p1o144dxdy = 0.00694444444444444444444444444444*INV(dx)*INV(dy);
- CCTK_REAL const p1o144dxdz = 0.00694444444444444444444444444444*INV(dx)*INV(dz);
- CCTK_REAL const p1o144dydz = 0.00694444444444444444444444444444*INV(dy)*INV(dz);
- CCTK_REAL const p1o2dx = 0.5*INV(dx);
- CCTK_REAL const p1o2dy = 0.5*INV(dy);
- CCTK_REAL const p1o2dz = 0.5*INV(dz);
- CCTK_REAL const p1o4dxdy = 0.25*INV(dx)*INV(dy);
- CCTK_REAL const p1o4dxdz = 0.25*INV(dx)*INV(dz);
- CCTK_REAL const p1o4dydz = 0.25*INV(dy)*INV(dz);
- CCTK_REAL const p1odx2 = INV(SQR(dx));
- CCTK_REAL const p1ody2 = INV(SQR(dy));
- CCTK_REAL const p1odz2 = INV(SQR(dz));
- CCTK_REAL const pm1o12dx2 = -0.0833333333333333333333333333333*INV(SQR(dx));
- CCTK_REAL const pm1o12dy2 = -0.0833333333333333333333333333333*INV(SQR(dy));
- CCTK_REAL const pm1o12dz2 = -0.0833333333333333333333333333333*INV(SQR(dz));
-
- /* Assign local copies of arrays functions */
-
-
-
- /* Calculate temporaries and arrays functions */
-
- /* Copy local copies back to grid functions */
-
- /* Loop over the grid points */
- #pragma omp parallel
- CCTK_LOOP3 (EM_energy,
- i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2],
- cctk_lsh[0],cctk_lsh[1],cctk_lsh[2])
- {
- ptrdiff_t const index = di*i + dj*j + dk*k;
-
- /* Assign local copies of grid functions */
-
- CCTK_REAL B1L = B1[index];
- CCTK_REAL B2L = B2[index];
- CCTK_REAL B3L = B3[index];
- CCTK_REAL El1L = El1[index];
- CCTK_REAL El2L = El2[index];
- CCTK_REAL El3L = El3[index];
-
-
- /* Include user supplied include files */
-
- /* Precompute derivatives */
-
- switch(fdOrder)
- {
- case 2:
- break;
-
- case 4:
- break;
- }
-
- /* Calculate temporaries and grid functions */
- CCTK_REAL rhoL = 0.5*(SQR(B1L) + SQR(B2L) + SQR(B3L) +
- SQR(El1L) + SQR(El2L) + SQR(El3L));
-
- /* Copy local copies back to grid functions */
- rho[index] = rhoL;
- }
- CCTK_ENDLOOP3 (EM_energy);
}
extern "C" void EM_energy(CCTK_ARGUMENTS)
@@ -153,7 +214,7 @@ extern "C" void EM_energy(CCTK_ARGUMENTS)
return;
}
- const char *groups[] = {"EMScript::B_group","EMScript::El_group","EMScript::rho_group"};
+ const char *groups[] = {"My_New_Implementation::B_group","My_New_Implementation::El_group","My_New_Implementation::rho_group"};
GenericFD_AssertGroupStorage(cctkGH, "EM_energy", 3, groups);
switch(fdOrder)
diff --git a/Examples/EMScript/src/EM_evol.cc b/Examples/EMScript/src/EM_evol.cc
index ff1c8ce..4e66517 100644
--- a/Examples/EMScript/src/EM_evol.cc
+++ b/Examples/EMScript/src/EM_evol.cc
@@ -14,13 +14,15 @@
#include "Differencing.h"
#include "cctk_Loop.h"
#include "loopcontrol.h"
+#include "OpenCLRunTime.h"
+#include "vectors.h"
/* Define macros used in calculations */
#define INITVALUE (42)
#define QAD(x) (SQR(SQR(x)))
-#define INV(x) ((1.0) / (x))
-#define SQR(x) ((x) * (x))
-#define CUB(x) ((x) * (x) * (x))
+#define INV(x) (kdiv(ToReal(1.0),x))
+#define SQR(x) (kmul(x,x))
+#define CUB(x) (kmul(x,SQR(x)))
extern "C" void EM_evol_SelectBCs(CCTK_ARGUMENTS)
{
@@ -28,12 +30,12 @@ extern "C" void EM_evol_SelectBCs(CCTK_ARGUMENTS)
DECLARE_CCTK_PARAMETERS;
CCTK_INT ierr = 0;
- ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "EMScript::B_grouprhs","flat");
+ ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "My_New_Implementation::B_grouprhs","flat");
if (ierr < 0)
- CCTK_WARN(1, "Failed to register flat BC for EMScript::B_grouprhs.");
- ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "EMScript::El_grouprhs","flat");
+ CCTK_WARN(1, "Failed to register flat BC for My_New_Implementation::B_grouprhs.");
+ ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "My_New_Implementation::El_grouprhs","flat");
if (ierr < 0)
- CCTK_WARN(1, "Failed to register flat BC for EMScript::El_grouprhs.");
+ CCTK_WARN(1, "Failed to register flat BC for My_New_Implementation::El_grouprhs.");
return;
}
@@ -42,152 +44,321 @@ static void EM_evol_Body(cGH const * restrict const cctkGH, int const dir, int c
DECLARE_CCTK_ARGUMENTS;
DECLARE_CCTK_PARAMETERS;
+ char const * const source =
+ "\n"
+ "/* Declare finite differencing variables */\n"
+ "\n"
+ "/* Include user-supplied include files */\n"
+ "\n"
+ "/* Initialise finite differencing variables */\n"
+ "ptrdiff_t const di = 1;\n"
+ "ptrdiff_t const dj = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n"
+ "ptrdiff_t const dk = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n"
+ "ptrdiff_t const cdi = sizeof(CCTK_REAL) * di;\n"
+ "ptrdiff_t const cdj = sizeof(CCTK_REAL) * dj;\n"
+ "ptrdiff_t const cdk = sizeof(CCTK_REAL) * dk;\n"
+ "CCTK_REAL_VEC const dx = ToReal(CCTK_DELTA_SPACE(0));\n"
+ "CCTK_REAL_VEC const dy = ToReal(CCTK_DELTA_SPACE(1));\n"
+ "CCTK_REAL_VEC const dz = ToReal(CCTK_DELTA_SPACE(2));\n"
+ "CCTK_REAL_VEC const dt = ToReal(CCTK_DELTA_TIME);\n"
+ "CCTK_REAL_VEC const t = ToReal(cctk_time);\n"
+ "CCTK_REAL_VEC const dxi = INV(dx);\n"
+ "CCTK_REAL_VEC const dyi = INV(dy);\n"
+ "CCTK_REAL_VEC const dzi = INV(dz);\n"
+ "CCTK_REAL_VEC const khalf = ToReal(0.5);\n"
+ "CCTK_REAL_VEC const kthird = ToReal(1.0/3.0);\n"
+ "CCTK_REAL_VEC const ktwothird = ToReal(2.0/3.0);\n"
+ "CCTK_REAL_VEC const kfourthird = ToReal(4.0/3.0);\n"
+ "CCTK_REAL_VEC const keightthird = ToReal(8.0/3.0);\n"
+ "CCTK_REAL_VEC const hdxi = kmul(ToReal(0.5), dxi);\n"
+ "CCTK_REAL_VEC const hdyi = kmul(ToReal(0.5), dyi);\n"
+ "CCTK_REAL_VEC const hdzi = kmul(ToReal(0.5), dzi);\n"
+ "\n"
+ "/* Initialize predefined quantities */\n"
+ "CCTK_REAL_VEC const p1o12dx = kmul(INV(dx),ToReal(0.0833333333333333333333333333333));\n"
+ "CCTK_REAL_VEC const p1o12dy = kmul(INV(dy),ToReal(0.0833333333333333333333333333333));\n"
+ "CCTK_REAL_VEC const p1o12dz = kmul(INV(dz),ToReal(0.0833333333333333333333333333333));\n"
+ "CCTK_REAL_VEC const p1o144dxdy = kmul(INV(dx),kmul(INV(dy),ToReal(0.00694444444444444444444444444444)));\n"
+ "CCTK_REAL_VEC const p1o144dxdz = kmul(INV(dx),kmul(INV(dz),ToReal(0.00694444444444444444444444444444)));\n"
+ "CCTK_REAL_VEC const p1o144dydz = kmul(INV(dy),kmul(INV(dz),ToReal(0.00694444444444444444444444444444)));\n"
+ "CCTK_REAL_VEC const p1o2dx = kmul(INV(dx),ToReal(0.5));\n"
+ "CCTK_REAL_VEC const p1o2dy = kmul(INV(dy),ToReal(0.5));\n"
+ "CCTK_REAL_VEC const p1o2dz = kmul(INV(dz),ToReal(0.5));\n"
+ "CCTK_REAL_VEC const p1o4dxdy = kmul(INV(dx),kmul(INV(dy),ToReal(0.25)));\n"
+ "CCTK_REAL_VEC const p1o4dxdz = kmul(INV(dx),kmul(INV(dz),ToReal(0.25)));\n"
+ "CCTK_REAL_VEC const p1o4dydz = kmul(INV(dy),kmul(INV(dz),ToReal(0.25)));\n"
+ "CCTK_REAL_VEC const p1odx2 = INV(SQR(dx));\n"
+ "CCTK_REAL_VEC const p1ody2 = INV(SQR(dy));\n"
+ "CCTK_REAL_VEC const p1odz2 = INV(SQR(dz));\n"
+ "CCTK_REAL_VEC const pm1o12dx2 = kmul(INV(SQR(dx)),ToReal(-0.0833333333333333333333333333333));\n"
+ "CCTK_REAL_VEC const pm1o12dy2 = kmul(INV(SQR(dy)),ToReal(-0.0833333333333333333333333333333));\n"
+ "CCTK_REAL_VEC const pm1o12dz2 = kmul(INV(SQR(dz)),ToReal(-0.0833333333333333333333333333333));\n"
+ "\n"
+ "/* Jacobian variable pointers */\n"
+ "bool const use_jacobian = (!CCTK_IsFunctionAliased(\"MultiPatch_GetMap\") || MultiPatch_GetMap(cctkGH) != jacobian_identity_map)\n"
+ " && strlen(jacobian_group) > 0;\n"
+ "if (use_jacobian && strlen(jacobian_derivative_group) == 0)\n"
+ "{\n"
+ " CCTK_WARN (1, \"GenericFD::jacobian_group and GenericFD::jacobian_derivative_group must both be set to valid group names\");\n"
+ "}\n"
+ "\n"
+ "CCTK_REAL const *restrict jacobian_ptrs[9];\n"
+ "if (use_jacobian) GenericFD_GroupDataPointers(cctkGH, jacobian_group,\n"
+ " 9, jacobian_ptrs);\n"
+ "\n"
+ "CCTK_REAL const *restrict const J11 = use_jacobian ? jacobian_ptrs[0] : 0;\n"
+ "CCTK_REAL const *restrict const J12 = use_jacobian ? jacobian_ptrs[1] : 0;\n"
+ "CCTK_REAL const *restrict const J13 = use_jacobian ? jacobian_ptrs[2] : 0;\n"
+ "CCTK_REAL const *restrict const J21 = use_jacobian ? jacobian_ptrs[3] : 0;\n"
+ "CCTK_REAL const *restrict const J22 = use_jacobian ? jacobian_ptrs[4] : 0;\n"
+ "CCTK_REAL const *restrict const J23 = use_jacobian ? jacobian_ptrs[5] : 0;\n"
+ "CCTK_REAL const *restrict const J31 = use_jacobian ? jacobian_ptrs[6] : 0;\n"
+ "CCTK_REAL const *restrict const J32 = use_jacobian ? jacobian_ptrs[7] : 0;\n"
+ "CCTK_REAL const *restrict const J33 = use_jacobian ? jacobian_ptrs[8] : 0;\n"
+ "\n"
+ "CCTK_REAL const *restrict jacobian_derivative_ptrs[18];\n"
+ "if (use_jacobian) GenericFD_GroupDataPointers(cctkGH, jacobian_derivative_group,\n"
+ " 18, jacobian_derivative_ptrs);\n"
+ "\n"
+ "CCTK_REAL const *restrict const dJ111 = use_jacobian ? jacobian_derivative_ptrs[0] : 0;\n"
+ "CCTK_REAL const *restrict const dJ112 = use_jacobian ? jacobian_derivative_ptrs[1] : 0;\n"
+ "CCTK_REAL const *restrict const dJ113 = use_jacobian ? jacobian_derivative_ptrs[2] : 0;\n"
+ "CCTK_REAL const *restrict const dJ122 = use_jacobian ? jacobian_derivative_ptrs[3] : 0;\n"
+ "CCTK_REAL const *restrict const dJ123 = use_jacobian ? jacobian_derivative_ptrs[4] : 0;\n"
+ "CCTK_REAL const *restrict const dJ133 = use_jacobian ? jacobian_derivative_ptrs[5] : 0;\n"
+ "CCTK_REAL const *restrict const dJ211 = use_jacobian ? jacobian_derivative_ptrs[6] : 0;\n"
+ "CCTK_REAL const *restrict const dJ212 = use_jacobian ? jacobian_derivative_ptrs[7] : 0;\n"
+ "CCTK_REAL const *restrict const dJ213 = use_jacobian ? jacobian_derivative_ptrs[8] : 0;\n"
+ "CCTK_REAL const *restrict const dJ222 = use_jacobian ? jacobian_derivative_ptrs[9] : 0;\n"
+ "CCTK_REAL const *restrict const dJ223 = use_jacobian ? jacobian_derivative_ptrs[10] : 0;\n"
+ "CCTK_REAL const *restrict const dJ233 = use_jacobian ? jacobian_derivative_ptrs[11] : 0;\n"
+ "CCTK_REAL const *restrict const dJ311 = use_jacobian ? jacobian_derivative_ptrs[12] : 0;\n"
+ "CCTK_REAL const *restrict const dJ312 = use_jacobian ? jacobian_derivative_ptrs[13] : 0;\n"
+ "CCTK_REAL const *restrict const dJ313 = use_jacobian ? jacobian_derivative_ptrs[14] : 0;\n"
+ "CCTK_REAL const *restrict const dJ322 = use_jacobian ? jacobian_derivative_ptrs[15] : 0;\n"
+ "CCTK_REAL const *restrict const dJ323 = use_jacobian ? jacobian_derivative_ptrs[16] : 0;\n"
+ "CCTK_REAL const *restrict const dJ333 = use_jacobian ? jacobian_derivative_ptrs[17] : 0;\n"
+ "\n"
+ "/* Assign local copies of arrays functions */\n"
+ "\n"
+ "\n"
+ "\n"
+ "/* Calculate temporaries and arrays functions */\n"
+ "\n"
+ "/* Copy local copies back to grid functions */\n"
+ "\n"
+ "/* Loop over the grid points */\n"
+ "#pragma omp parallel\n"
+ "LC_LOOP3VEC (EM_evol,\n"
+ " i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2],\n"
+ " cctk_lsh[0],cctk_lsh[1],cctk_lsh[2],\n"
+ " CCTK_REAL_VEC_SIZE)\n"
+ "{\n"
+ " ptrdiff_t const index = di*i + dj*j + dk*k;\n"
+ " \n"
+ " /* Assign local copies of grid functions */\n"
+ " \n"
+ " CCTK_REAL_VEC B1L = vec_load(B1[index]);\n"
+ " CCTK_REAL_VEC B2L = vec_load(B2[index]);\n"
+ " CCTK_REAL_VEC B3L = vec_load(B3[index]);\n"
+ " CCTK_REAL_VEC El1L = vec_load(El1[index]);\n"
+ " CCTK_REAL_VEC El2L = vec_load(El2[index]);\n"
+ " CCTK_REAL_VEC El3L = vec_load(El3[index]);\n"
+ " \n"
+ " \n"
+ " CCTK_REAL_VEC J11L, J12L, J13L, J21L, J22L, J23L, J31L, J32L, J33L;\n"
+ " \n"
+ " if (use_jacobian)\n"
+ " {\n"
+ " J11L = vec_load(J11[index]);\n"
+ " J12L = vec_load(J12[index]);\n"
+ " J13L = vec_load(J13[index]);\n"
+ " J21L = vec_load(J21[index]);\n"
+ " J22L = vec_load(J22[index]);\n"
+ " J23L = vec_load(J23[index]);\n"
+ " J31L = vec_load(J31[index]);\n"
+ " J32L = vec_load(J32[index]);\n"
+ " J33L = vec_load(J33[index]);\n"
+ " }\n"
+ " \n"
+ " /* Include user supplied include files */\n"
+ " \n"
+ " /* Precompute derivatives */\n"
+ " CCTK_REAL_VEC PDstandard1B1;\n"
+ " CCTK_REAL_VEC PDstandard2B1;\n"
+ " CCTK_REAL_VEC PDstandard3B1;\n"
+ " CCTK_REAL_VEC PDstandard1B2;\n"
+ " CCTK_REAL_VEC PDstandard2B2;\n"
+ " CCTK_REAL_VEC PDstandard3B2;\n"
+ " CCTK_REAL_VEC PDstandard1B3;\n"
+ " CCTK_REAL_VEC PDstandard2B3;\n"
+ " CCTK_REAL_VEC PDstandard3B3;\n"
+ " CCTK_REAL_VEC PDstandard1El1;\n"
+ " CCTK_REAL_VEC PDstandard2El1;\n"
+ " CCTK_REAL_VEC PDstandard3El1;\n"
+ " CCTK_REAL_VEC PDstandard1El2;\n"
+ " CCTK_REAL_VEC PDstandard2El2;\n"
+ " CCTK_REAL_VEC PDstandard3El2;\n"
+ " CCTK_REAL_VEC PDstandard1El3;\n"
+ " CCTK_REAL_VEC PDstandard2El3;\n"
+ " CCTK_REAL_VEC PDstandard3El3;\n"
+ " \n"
+ " switch(fdOrder)\n"
+ " {\n"
+ " case 2:\n"
+ " PDstandard1B1 = PDstandardfdOrder21(&B1[index]);\n"
+ " PDstandard2B1 = PDstandardfdOrder22(&B1[index]);\n"
+ " PDstandard3B1 = PDstandardfdOrder23(&B1[index]);\n"
+ " PDstandard1B2 = PDstandardfdOrder21(&B2[index]);\n"
+ " PDstandard2B2 = PDstandardfdOrder22(&B2[index]);\n"
+ " PDstandard3B2 = PDstandardfdOrder23(&B2[index]);\n"
+ " PDstandard1B3 = PDstandardfdOrder21(&B3[index]);\n"
+ " PDstandard2B3 = PDstandardfdOrder22(&B3[index]);\n"
+ " PDstandard3B3 = PDstandardfdOrder23(&B3[index]);\n"
+ " PDstandard1El1 = PDstandardfdOrder21(&El1[index]);\n"
+ " PDstandard2El1 = PDstandardfdOrder22(&El1[index]);\n"
+ " PDstandard3El1 = PDstandardfdOrder23(&El1[index]);\n"
+ " PDstandard1El2 = PDstandardfdOrder21(&El2[index]);\n"
+ " PDstandard2El2 = PDstandardfdOrder22(&El2[index]);\n"
+ " PDstandard3El2 = PDstandardfdOrder23(&El2[index]);\n"
+ " PDstandard1El3 = PDstandardfdOrder21(&El3[index]);\n"
+ " PDstandard2El3 = PDstandardfdOrder22(&El3[index]);\n"
+ " PDstandard3El3 = PDstandardfdOrder23(&El3[index]);\n"
+ " break;\n"
+ " \n"
+ " case 4:\n"
+ " PDstandard1B1 = PDstandardfdOrder41(&B1[index]);\n"
+ " PDstandard2B1 = PDstandardfdOrder42(&B1[index]);\n"
+ " PDstandard3B1 = PDstandardfdOrder43(&B1[index]);\n"
+ " PDstandard1B2 = PDstandardfdOrder41(&B2[index]);\n"
+ " PDstandard2B2 = PDstandardfdOrder42(&B2[index]);\n"
+ " PDstandard3B2 = PDstandardfdOrder43(&B2[index]);\n"
+ " PDstandard1B3 = PDstandardfdOrder41(&B3[index]);\n"
+ " PDstandard2B3 = PDstandardfdOrder42(&B3[index]);\n"
+ " PDstandard3B3 = PDstandardfdOrder43(&B3[index]);\n"
+ " PDstandard1El1 = PDstandardfdOrder41(&El1[index]);\n"
+ " PDstandard2El1 = PDstandardfdOrder42(&El1[index]);\n"
+ " PDstandard3El1 = PDstandardfdOrder43(&El1[index]);\n"
+ " PDstandard1El2 = PDstandardfdOrder41(&El2[index]);\n"
+ " PDstandard2El2 = PDstandardfdOrder42(&El2[index]);\n"
+ " PDstandard3El2 = PDstandardfdOrder43(&El2[index]);\n"
+ " PDstandard1El3 = PDstandardfdOrder41(&El3[index]);\n"
+ " PDstandard2El3 = PDstandardfdOrder42(&El3[index]);\n"
+ " PDstandard3El3 = PDstandardfdOrder43(&El3[index]);\n"
+ " break;\n"
+ " }\n"
+ " \n"
+ " /* Calculate temporaries and grid functions */\n"
+ " CCTK_REAL_VEC JacPDstandard1B2;\n"
+ " CCTK_REAL_VEC JacPDstandard1B3;\n"
+ " CCTK_REAL_VEC JacPDstandard1El2;\n"
+ " CCTK_REAL_VEC JacPDstandard1El3;\n"
+ " CCTK_REAL_VEC JacPDstandard2B1;\n"
+ " CCTK_REAL_VEC JacPDstandard2B3;\n"
+ " CCTK_REAL_VEC JacPDstandard2El1;\n"
+ " CCTK_REAL_VEC JacPDstandard2El3;\n"
+ " CCTK_REAL_VEC JacPDstandard3B1;\n"
+ " CCTK_REAL_VEC JacPDstandard3B2;\n"
+ " CCTK_REAL_VEC JacPDstandard3El1;\n"
+ " CCTK_REAL_VEC JacPDstandard3El2;\n"
+ " \n"
+ " if (use_jacobian)\n"
+ " {\n"
+ " JacPDstandard1B2 = \n"
+ " kmadd(J11L,PDstandard1B2,kmadd(J21L,PDstandard2B2,kmul(J31L,PDstandard3B2)));\n"
+ " \n"
+ " JacPDstandard1B3 = \n"
+ " kmadd(J11L,PDstandard1B3,kmadd(J21L,PDstandard2B3,kmul(J31L,PDstandard3B3)));\n"
+ " \n"
+ " JacPDstandard1El2 = \n"
+ " kmadd(J11L,PDstandard1El2,kmadd(J21L,PDstandard2El2,kmul(J31L,PDstandard3El2)));\n"
+ " \n"
+ " JacPDstandard1El3 = \n"
+ " kmadd(J11L,PDstandard1El3,kmadd(J21L,PDstandard2El3,kmul(J31L,PDstandard3El3)));\n"
+ " \n"
+ " JacPDstandard2B1 = \n"
+ " kmadd(J12L,PDstandard1B1,kmadd(J22L,PDstandard2B1,kmul(J32L,PDstandard3B1)));\n"
+ " \n"
+ " JacPDstandard2B3 = \n"
+ " kmadd(J12L,PDstandard1B3,kmadd(J22L,PDstandard2B3,kmul(J32L,PDstandard3B3)));\n"
+ " \n"
+ " JacPDstandard2El1 = \n"
+ " kmadd(J12L,PDstandard1El1,kmadd(J22L,PDstandard2El1,kmul(J32L,PDstandard3El1)));\n"
+ " \n"
+ " JacPDstandard2El3 = \n"
+ " kmadd(J12L,PDstandard1El3,kmadd(J22L,PDstandard2El3,kmul(J32L,PDstandard3El3)));\n"
+ " \n"
+ " JacPDstandard3B1 = \n"
+ " kmadd(J13L,PDstandard1B1,kmadd(J23L,PDstandard2B1,kmul(J33L,PDstandard3B1)));\n"
+ " \n"
+ " JacPDstandard3B2 = \n"
+ " kmadd(J13L,PDstandard1B2,kmadd(J23L,PDstandard2B2,kmul(J33L,PDstandard3B2)));\n"
+ " \n"
+ " JacPDstandard3El1 = \n"
+ " kmadd(J13L,PDstandard1El1,kmadd(J23L,PDstandard2El1,kmul(J33L,PDstandard3El1)));\n"
+ " \n"
+ " JacPDstandard3El2 = \n"
+ " kmadd(J13L,PDstandard1El2,kmadd(J23L,PDstandard2El2,kmul(J33L,PDstandard3El2)));\n"
+ " }\n"
+ " else\n"
+ " {\n"
+ " JacPDstandard1B2 = PDstandard1B2;\n"
+ " \n"
+ " JacPDstandard1B3 = PDstandard1B3;\n"
+ " \n"
+ " JacPDstandard1El2 = PDstandard1El2;\n"
+ " \n"
+ " JacPDstandard1El3 = PDstandard1El3;\n"
+ " \n"
+ " JacPDstandard2B1 = PDstandard2B1;\n"
+ " \n"
+ " JacPDstandard2B3 = PDstandard2B3;\n"
+ " \n"
+ " JacPDstandard2El1 = PDstandard2El1;\n"
+ " \n"
+ " JacPDstandard2El3 = PDstandard2El3;\n"
+ " \n"
+ " JacPDstandard3B1 = PDstandard3B1;\n"
+ " \n"
+ " JacPDstandard3B2 = PDstandard3B2;\n"
+ " \n"
+ " JacPDstandard3El1 = PDstandard3El1;\n"
+ " \n"
+ " JacPDstandard3El2 = PDstandard3El2;\n"
+ " }\n"
+ " \n"
+ " CCTK_REAL_VEC El1rhsL = ksub(JacPDstandard2B3,JacPDstandard3B2);\n"
+ " \n"
+ " CCTK_REAL_VEC El2rhsL = ksub(JacPDstandard3B1,JacPDstandard1B3);\n"
+ " \n"
+ " CCTK_REAL_VEC El3rhsL = ksub(JacPDstandard1B2,JacPDstandard2B1);\n"
+ " \n"
+ " CCTK_REAL_VEC B1rhsL = ksub(JacPDstandard3El2,JacPDstandard2El3);\n"
+ " \n"
+ " CCTK_REAL_VEC B2rhsL = ksub(JacPDstandard1El3,JacPDstandard3El1);\n"
+ " \n"
+ " CCTK_REAL_VEC B3rhsL = ksub(JacPDstandard2El1,JacPDstandard1El2);\n"
+ " \n"
+ " /* Copy local copies back to grid functions */\n"
+ " vec_store_nta_partial(B1rhs[index],B1rhsL);\n"
+ " vec_store_nta_partial(B2rhs[index],B2rhsL);\n"
+ " vec_store_nta_partial(B3rhs[index],B3rhsL);\n"
+ " vec_store_nta_partial(El1rhs[index],El1rhsL);\n"
+ " vec_store_nta_partial(El2rhs[index],El2rhsL);\n"
+ " vec_store_nta_partial(El3rhs[index],El3rhsL);\n"
+ "}\n"
+ "LC_ENDLOOP3VEC (EM_evol);\n"
+ ""
+ ;
- /* Declare finite differencing variables */
+ char const * const groups[] = {"My_New_Implementation::B_group","My_New_Implementation::B_grouprhs","My_New_Implementation::El_group","My_New_Implementation::El_grouprhs",NULL};
- /* Include user-supplied include files */
+ static struct OpenCLKernel * kernel = NULL;
+ char const * const sources[] = {differencing, source, NULL};
+ OpenCLRunTime_CallKernel (cctkGH, CCTK_THORNSTRING, "EM_evol",
+ sources, groups, NULL, NULL, NULL, -1,
+ imin, imax, &kernel);
- /* Initialise finite differencing variables */
- ptrdiff_t const di = 1;
- ptrdiff_t const dj = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);
- ptrdiff_t const dk = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);
- ptrdiff_t const cdi = sizeof(CCTK_REAL) * di;
- ptrdiff_t const cdj = sizeof(CCTK_REAL) * dj;
- ptrdiff_t const cdk = sizeof(CCTK_REAL) * dk;
- CCTK_REAL const dx = ToReal(CCTK_DELTA_SPACE(0));
- CCTK_REAL const dy = ToReal(CCTK_DELTA_SPACE(1));
- CCTK_REAL const dz = ToReal(CCTK_DELTA_SPACE(2));
- CCTK_REAL const dt = ToReal(CCTK_DELTA_TIME);
- CCTK_REAL const t = ToReal(cctk_time);
- CCTK_REAL const dxi = INV(dx);
- CCTK_REAL const dyi = INV(dy);
- CCTK_REAL const dzi = INV(dz);
- CCTK_REAL const khalf = 0.5;
- CCTK_REAL const kthird = 1/3.0;
- CCTK_REAL const ktwothird = 2.0/3.0;
- CCTK_REAL const kfourthird = 4.0/3.0;
- CCTK_REAL const keightthird = 8.0/3.0;
- CCTK_REAL const hdxi = 0.5 * dxi;
- CCTK_REAL const hdyi = 0.5 * dyi;
- CCTK_REAL const hdzi = 0.5 * dzi;
-
- /* Initialize predefined quantities */
- CCTK_REAL const p1o12dx = 0.0833333333333333333333333333333*INV(dx);
- CCTK_REAL const p1o12dy = 0.0833333333333333333333333333333*INV(dy);
- CCTK_REAL const p1o12dz = 0.0833333333333333333333333333333*INV(dz);
- CCTK_REAL const p1o144dxdy = 0.00694444444444444444444444444444*INV(dx)*INV(dy);
- CCTK_REAL const p1o144dxdz = 0.00694444444444444444444444444444*INV(dx)*INV(dz);
- CCTK_REAL const p1o144dydz = 0.00694444444444444444444444444444*INV(dy)*INV(dz);
- CCTK_REAL const p1o2dx = 0.5*INV(dx);
- CCTK_REAL const p1o2dy = 0.5*INV(dy);
- CCTK_REAL const p1o2dz = 0.5*INV(dz);
- CCTK_REAL const p1o4dxdy = 0.25*INV(dx)*INV(dy);
- CCTK_REAL const p1o4dxdz = 0.25*INV(dx)*INV(dz);
- CCTK_REAL const p1o4dydz = 0.25*INV(dy)*INV(dz);
- CCTK_REAL const p1odx2 = INV(SQR(dx));
- CCTK_REAL const p1ody2 = INV(SQR(dy));
- CCTK_REAL const p1odz2 = INV(SQR(dz));
- CCTK_REAL const pm1o12dx2 = -0.0833333333333333333333333333333*INV(SQR(dx));
- CCTK_REAL const pm1o12dy2 = -0.0833333333333333333333333333333*INV(SQR(dy));
- CCTK_REAL const pm1o12dz2 = -0.0833333333333333333333333333333*INV(SQR(dz));
-
- /* Assign local copies of arrays functions */
-
-
-
- /* Calculate temporaries and arrays functions */
-
- /* Copy local copies back to grid functions */
-
- /* Loop over the grid points */
- #pragma omp parallel
- CCTK_LOOP3 (EM_evol,
- i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2],
- cctk_lsh[0],cctk_lsh[1],cctk_lsh[2])
- {
- ptrdiff_t const index = di*i + dj*j + dk*k;
-
- /* Assign local copies of grid functions */
-
- CCTK_REAL B1L = B1[index];
- CCTK_REAL B2L = B2[index];
- CCTK_REAL B3L = B3[index];
- CCTK_REAL El1L = El1[index];
- CCTK_REAL El2L = El2[index];
- CCTK_REAL El3L = El3[index];
-
-
- /* Include user supplied include files */
-
- /* Precompute derivatives */
- CCTK_REAL PDstandard2B1;
- CCTK_REAL PDstandard3B1;
- CCTK_REAL PDstandard1B2;
- CCTK_REAL PDstandard3B2;
- CCTK_REAL PDstandard1B3;
- CCTK_REAL PDstandard2B3;
- CCTK_REAL PDstandard2El1;
- CCTK_REAL PDstandard3El1;
- CCTK_REAL PDstandard1El2;
- CCTK_REAL PDstandard3El2;
- CCTK_REAL PDstandard1El3;
- CCTK_REAL PDstandard2El3;
-
- switch(fdOrder)
- {
- case 2:
- PDstandard2B1 = PDstandardfdOrder22(&B1[index]);
- PDstandard3B1 = PDstandardfdOrder23(&B1[index]);
- PDstandard1B2 = PDstandardfdOrder21(&B2[index]);
- PDstandard3B2 = PDstandardfdOrder23(&B2[index]);
- PDstandard1B3 = PDstandardfdOrder21(&B3[index]);
- PDstandard2B3 = PDstandardfdOrder22(&B3[index]);
- PDstandard2El1 = PDstandardfdOrder22(&El1[index]);
- PDstandard3El1 = PDstandardfdOrder23(&El1[index]);
- PDstandard1El2 = PDstandardfdOrder21(&El2[index]);
- PDstandard3El2 = PDstandardfdOrder23(&El2[index]);
- PDstandard1El3 = PDstandardfdOrder21(&El3[index]);
- PDstandard2El3 = PDstandardfdOrder22(&El3[index]);
- break;
-
- case 4:
- PDstandard2B1 = PDstandardfdOrder42(&B1[index]);
- PDstandard3B1 = PDstandardfdOrder43(&B1[index]);
- PDstandard1B2 = PDstandardfdOrder41(&B2[index]);
- PDstandard3B2 = PDstandardfdOrder43(&B2[index]);
- PDstandard1B3 = PDstandardfdOrder41(&B3[index]);
- PDstandard2B3 = PDstandardfdOrder42(&B3[index]);
- PDstandard2El1 = PDstandardfdOrder42(&El1[index]);
- PDstandard3El1 = PDstandardfdOrder43(&El1[index]);
- PDstandard1El2 = PDstandardfdOrder41(&El2[index]);
- PDstandard3El2 = PDstandardfdOrder43(&El2[index]);
- PDstandard1El3 = PDstandardfdOrder41(&El3[index]);
- PDstandard2El3 = PDstandardfdOrder42(&El3[index]);
- break;
- }
-
- /* Calculate temporaries and grid functions */
- CCTK_REAL El1rhsL = PDstandard2B3 - PDstandard3B2;
-
- CCTK_REAL El2rhsL = -PDstandard1B3 + PDstandard3B1;
-
- CCTK_REAL El3rhsL = PDstandard1B2 - PDstandard2B1;
-
- CCTK_REAL B1rhsL = -PDstandard2El3 + PDstandard3El2;
-
- CCTK_REAL B2rhsL = PDstandard1El3 - PDstandard3El1;
-
- CCTK_REAL B3rhsL = -PDstandard1El2 + PDstandard2El1;
-
- /* Copy local copies back to grid functions */
- B1rhs[index] = B1rhsL;
- B2rhs[index] = B2rhsL;
- B3rhs[index] = B3rhsL;
- El1rhs[index] = El1rhsL;
- El2rhs[index] = El2rhsL;
- El3rhs[index] = El3rhsL;
- }
- CCTK_ENDLOOP3 (EM_evol);
}
extern "C" void EM_evol(CCTK_ARGUMENTS)
@@ -206,7 +377,7 @@ extern "C" void EM_evol(CCTK_ARGUMENTS)
return;
}
- const char *groups[] = {"EMScript::B_group","EMScript::B_grouprhs","EMScript::El_group","EMScript::El_grouprhs"};
+ const char *groups[] = {"My_New_Implementation::B_group","My_New_Implementation::B_grouprhs","My_New_Implementation::El_group","My_New_Implementation::El_grouprhs"};
GenericFD_AssertGroupStorage(cctkGH, "EM_evol", 4, groups);
switch(fdOrder)
diff --git a/Examples/EMScript/src/EM_initial.cc b/Examples/EMScript/src/EM_initial.cc
index ddbd59e..34e72f1 100644
--- a/Examples/EMScript/src/EM_initial.cc
+++ b/Examples/EMScript/src/EM_initial.cc
@@ -14,13 +14,15 @@
#include "Differencing.h"
#include "cctk_Loop.h"
#include "loopcontrol.h"
+#include "OpenCLRunTime.h"
+#include "vectors.h"
/* Define macros used in calculations */
#define INITVALUE (42)
#define QAD(x) (SQR(SQR(x)))
-#define INV(x) ((1.0) / (x))
-#define SQR(x) ((x) * (x))
-#define CUB(x) ((x) * (x) * (x))
+#define INV(x) (kdiv(ToReal(1.0),x))
+#define SQR(x) (kmul(x,x))
+#define CUB(x) (kmul(x,SQR(x)))
extern "C" void EM_initial_SelectBCs(CCTK_ARGUMENTS)
{
@@ -28,12 +30,12 @@ extern "C" void EM_initial_SelectBCs(CCTK_ARGUMENTS)
DECLARE_CCTK_PARAMETERS;
CCTK_INT ierr = 0;
- ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "EMScript::B_group","flat");
+ ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "My_New_Implementation::B_group","flat");
if (ierr < 0)
- CCTK_WARN(1, "Failed to register flat BC for EMScript::B_group.");
- ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "EMScript::El_group","flat");
+ CCTK_WARN(1, "Failed to register flat BC for My_New_Implementation::B_group.");
+ ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "My_New_Implementation::El_group","flat");
if (ierr < 0)
- CCTK_WARN(1, "Failed to register flat BC for EMScript::El_group.");
+ CCTK_WARN(1, "Failed to register flat BC for My_New_Implementation::El_group.");
return;
}
@@ -42,116 +44,182 @@ static void EM_initial_Body(cGH const * restrict const cctkGH, int const dir, in
DECLARE_CCTK_ARGUMENTS;
DECLARE_CCTK_PARAMETERS;
+ char const * const source =
+ "\n"
+ "/* Declare finite differencing variables */\n"
+ "\n"
+ "/* Include user-supplied include files */\n"
+ "\n"
+ "/* Initialise finite differencing variables */\n"
+ "ptrdiff_t const di = 1;\n"
+ "ptrdiff_t const dj = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n"
+ "ptrdiff_t const dk = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n"
+ "ptrdiff_t const cdi = sizeof(CCTK_REAL) * di;\n"
+ "ptrdiff_t const cdj = sizeof(CCTK_REAL) * dj;\n"
+ "ptrdiff_t const cdk = sizeof(CCTK_REAL) * dk;\n"
+ "CCTK_REAL_VEC const dx = ToReal(CCTK_DELTA_SPACE(0));\n"
+ "CCTK_REAL_VEC const dy = ToReal(CCTK_DELTA_SPACE(1));\n"
+ "CCTK_REAL_VEC const dz = ToReal(CCTK_DELTA_SPACE(2));\n"
+ "CCTK_REAL_VEC const dt = ToReal(CCTK_DELTA_TIME);\n"
+ "CCTK_REAL_VEC const t = ToReal(cctk_time);\n"
+ "CCTK_REAL_VEC const dxi = INV(dx);\n"
+ "CCTK_REAL_VEC const dyi = INV(dy);\n"
+ "CCTK_REAL_VEC const dzi = INV(dz);\n"
+ "CCTK_REAL_VEC const khalf = ToReal(0.5);\n"
+ "CCTK_REAL_VEC const kthird = ToReal(1.0/3.0);\n"
+ "CCTK_REAL_VEC const ktwothird = ToReal(2.0/3.0);\n"
+ "CCTK_REAL_VEC const kfourthird = ToReal(4.0/3.0);\n"
+ "CCTK_REAL_VEC const keightthird = ToReal(8.0/3.0);\n"
+ "CCTK_REAL_VEC const hdxi = kmul(ToReal(0.5), dxi);\n"
+ "CCTK_REAL_VEC const hdyi = kmul(ToReal(0.5), dyi);\n"
+ "CCTK_REAL_VEC const hdzi = kmul(ToReal(0.5), dzi);\n"
+ "\n"
+ "/* Initialize predefined quantities */\n"
+ "CCTK_REAL_VEC const p1o12dx = kmul(INV(dx),ToReal(0.0833333333333333333333333333333));\n"
+ "CCTK_REAL_VEC const p1o12dy = kmul(INV(dy),ToReal(0.0833333333333333333333333333333));\n"
+ "CCTK_REAL_VEC const p1o12dz = kmul(INV(dz),ToReal(0.0833333333333333333333333333333));\n"
+ "CCTK_REAL_VEC const p1o144dxdy = kmul(INV(dx),kmul(INV(dy),ToReal(0.00694444444444444444444444444444)));\n"
+ "CCTK_REAL_VEC const p1o144dxdz = kmul(INV(dx),kmul(INV(dz),ToReal(0.00694444444444444444444444444444)));\n"
+ "CCTK_REAL_VEC const p1o144dydz = kmul(INV(dy),kmul(INV(dz),ToReal(0.00694444444444444444444444444444)));\n"
+ "CCTK_REAL_VEC const p1o2dx = kmul(INV(dx),ToReal(0.5));\n"
+ "CCTK_REAL_VEC const p1o2dy = kmul(INV(dy),ToReal(0.5));\n"
+ "CCTK_REAL_VEC const p1o2dz = kmul(INV(dz),ToReal(0.5));\n"
+ "CCTK_REAL_VEC const p1o4dxdy = kmul(INV(dx),kmul(INV(dy),ToReal(0.25)));\n"
+ "CCTK_REAL_VEC const p1o4dxdz = kmul(INV(dx),kmul(INV(dz),ToReal(0.25)));\n"
+ "CCTK_REAL_VEC const p1o4dydz = kmul(INV(dy),kmul(INV(dz),ToReal(0.25)));\n"
+ "CCTK_REAL_VEC const p1odx2 = INV(SQR(dx));\n"
+ "CCTK_REAL_VEC const p1ody2 = INV(SQR(dy));\n"
+ "CCTK_REAL_VEC const p1odz2 = INV(SQR(dz));\n"
+ "CCTK_REAL_VEC const pm1o12dx2 = kmul(INV(SQR(dx)),ToReal(-0.0833333333333333333333333333333));\n"
+ "CCTK_REAL_VEC const pm1o12dy2 = kmul(INV(SQR(dy)),ToReal(-0.0833333333333333333333333333333));\n"
+ "CCTK_REAL_VEC const pm1o12dz2 = kmul(INV(SQR(dz)),ToReal(-0.0833333333333333333333333333333));\n"
+ "\n"
+ "/* Jacobian variable pointers */\n"
+ "bool const use_jacobian = (!CCTK_IsFunctionAliased(\"MultiPatch_GetMap\") || MultiPatch_GetMap(cctkGH) != jacobian_identity_map)\n"
+ " && strlen(jacobian_group) > 0;\n"
+ "if (use_jacobian && strlen(jacobian_derivative_group) == 0)\n"
+ "{\n"
+ " CCTK_WARN (1, \"GenericFD::jacobian_group and GenericFD::jacobian_derivative_group must both be set to valid group names\");\n"
+ "}\n"
+ "\n"
+ "CCTK_REAL const *restrict jacobian_ptrs[9];\n"
+ "if (use_jacobian) GenericFD_GroupDataPointers(cctkGH, jacobian_group,\n"
+ " 9, jacobian_ptrs);\n"
+ "\n"
+ "CCTK_REAL const *restrict const J11 = use_jacobian ? jacobian_ptrs[0] : 0;\n"
+ "CCTK_REAL const *restrict const J12 = use_jacobian ? jacobian_ptrs[1] : 0;\n"
+ "CCTK_REAL const *restrict const J13 = use_jacobian ? jacobian_ptrs[2] : 0;\n"
+ "CCTK_REAL const *restrict const J21 = use_jacobian ? jacobian_ptrs[3] : 0;\n"
+ "CCTK_REAL const *restrict const J22 = use_jacobian ? jacobian_ptrs[4] : 0;\n"
+ "CCTK_REAL const *restrict const J23 = use_jacobian ? jacobian_ptrs[5] : 0;\n"
+ "CCTK_REAL const *restrict const J31 = use_jacobian ? jacobian_ptrs[6] : 0;\n"
+ "CCTK_REAL const *restrict const J32 = use_jacobian ? jacobian_ptrs[7] : 0;\n"
+ "CCTK_REAL const *restrict const J33 = use_jacobian ? jacobian_ptrs[8] : 0;\n"
+ "\n"
+ "CCTK_REAL const *restrict jacobian_derivative_ptrs[18];\n"
+ "if (use_jacobian) GenericFD_GroupDataPointers(cctkGH, jacobian_derivative_group,\n"
+ " 18, jacobian_derivative_ptrs);\n"
+ "\n"
+ "CCTK_REAL const *restrict const dJ111 = use_jacobian ? jacobian_derivative_ptrs[0] : 0;\n"
+ "CCTK_REAL const *restrict const dJ112 = use_jacobian ? jacobian_derivative_ptrs[1] : 0;\n"
+ "CCTK_REAL const *restrict const dJ113 = use_jacobian ? jacobian_derivative_ptrs[2] : 0;\n"
+ "CCTK_REAL const *restrict const dJ122 = use_jacobian ? jacobian_derivative_ptrs[3] : 0;\n"
+ "CCTK_REAL const *restrict const dJ123 = use_jacobian ? jacobian_derivative_ptrs[4] : 0;\n"
+ "CCTK_REAL const *restrict const dJ133 = use_jacobian ? jacobian_derivative_ptrs[5] : 0;\n"
+ "CCTK_REAL const *restrict const dJ211 = use_jacobian ? jacobian_derivative_ptrs[6] : 0;\n"
+ "CCTK_REAL const *restrict const dJ212 = use_jacobian ? jacobian_derivative_ptrs[7] : 0;\n"
+ "CCTK_REAL const *restrict const dJ213 = use_jacobian ? jacobian_derivative_ptrs[8] : 0;\n"
+ "CCTK_REAL const *restrict const dJ222 = use_jacobian ? jacobian_derivative_ptrs[9] : 0;\n"
+ "CCTK_REAL const *restrict const dJ223 = use_jacobian ? jacobian_derivative_ptrs[10] : 0;\n"
+ "CCTK_REAL const *restrict const dJ233 = use_jacobian ? jacobian_derivative_ptrs[11] : 0;\n"
+ "CCTK_REAL const *restrict const dJ311 = use_jacobian ? jacobian_derivative_ptrs[12] : 0;\n"
+ "CCTK_REAL const *restrict const dJ312 = use_jacobian ? jacobian_derivative_ptrs[13] : 0;\n"
+ "CCTK_REAL const *restrict const dJ313 = use_jacobian ? jacobian_derivative_ptrs[14] : 0;\n"
+ "CCTK_REAL const *restrict const dJ322 = use_jacobian ? jacobian_derivative_ptrs[15] : 0;\n"
+ "CCTK_REAL const *restrict const dJ323 = use_jacobian ? jacobian_derivative_ptrs[16] : 0;\n"
+ "CCTK_REAL const *restrict const dJ333 = use_jacobian ? jacobian_derivative_ptrs[17] : 0;\n"
+ "\n"
+ "/* Assign local copies of arrays functions */\n"
+ "\n"
+ "\n"
+ "\n"
+ "/* Calculate temporaries and arrays functions */\n"
+ "\n"
+ "/* Copy local copies back to grid functions */\n"
+ "\n"
+ "/* Loop over the grid points */\n"
+ "#pragma omp parallel\n"
+ "LC_LOOP3VEC (EM_initial,\n"
+ " i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2],\n"
+ " cctk_lsh[0],cctk_lsh[1],cctk_lsh[2],\n"
+ " CCTK_REAL_VEC_SIZE)\n"
+ "{\n"
+ " ptrdiff_t const index = di*i + dj*j + dk*k;\n"
+ " \n"
+ " /* Assign local copies of grid functions */\n"
+ " \n"
+ " CCTK_REAL_VEC xL = vec_load(x[index]);\n"
+ " CCTK_REAL_VEC yL = vec_load(y[index]);\n"
+ " \n"
+ " \n"
+ " \n"
+ " /* Include user supplied include files */\n"
+ " \n"
+ " /* Precompute derivatives */\n"
+ " \n"
+ " switch(fdOrder)\n"
+ " {\n"
+ " case 2:\n"
+ " break;\n"
+ " \n"
+ " case 4:\n"
+ " break;\n"
+ " }\n"
+ " \n"
+ " /* Calculate temporaries and grid functions */\n"
+ " CCTK_REAL_VEC sigma = ToReal(1);\n"
+ " \n"
+ " CCTK_REAL_VEC csetemp0 = kadd(xL,yL);\n"
+ " \n"
+ " CCTK_REAL_VEC El1L = \n"
+ " kmul(sigma,kcos(kmul(csetemp0,kmul(Pi,ToReal(2)))));\n"
+ " \n"
+ " CCTK_REAL_VEC csetemp1 = ksub(ToReal(1),sigma);\n"
+ " \n"
+ " CCTK_REAL_VEC csetemp2 = kmul(xL,kmul(Pi,ToReal(2)));\n"
+ " \n"
+ " CCTK_REAL_VEC El2L = \n"
+ " knmadd(sigma,kcos(kmul(csetemp0,kmul(Pi,ToReal(2)))),kmul(csetemp1,kcos(csetemp2)));\n"
+ " \n"
+ " CCTK_REAL_VEC El3L = ToReal(0);\n"
+ " \n"
+ " CCTK_REAL_VEC B1L = ToReal(0);\n"
+ " \n"
+ " CCTK_REAL_VEC B2L = ToReal(0);\n"
+ " \n"
+ " CCTK_REAL_VEC B3L = \n"
+ " kmadd(csetemp1,kcos(csetemp2),kmul(sigma,kcos(kmul(csetemp0,kmul(Pi,ToReal(2))))));\n"
+ " \n"
+ " /* Copy local copies back to grid functions */\n"
+ " vec_store_nta_partial(B1[index],B1L);\n"
+ " vec_store_nta_partial(B2[index],B2L);\n"
+ " vec_store_nta_partial(B3[index],B3L);\n"
+ " vec_store_nta_partial(El1[index],El1L);\n"
+ " vec_store_nta_partial(El2[index],El2L);\n"
+ " vec_store_nta_partial(El3[index],El3L);\n"
+ "}\n"
+ "LC_ENDLOOP3VEC (EM_initial);\n"
+ ""
+ ;
- /* Declare finite differencing variables */
+ char const * const groups[] = {"My_New_Implementation::B_group","My_New_Implementation::El_group","grid::coordinates",NULL};
- /* Include user-supplied include files */
+ static struct OpenCLKernel * kernel = NULL;
+ char const * const sources[] = {differencing, source, NULL};
+ OpenCLRunTime_CallKernel (cctkGH, CCTK_THORNSTRING, "EM_initial",
+ sources, groups, NULL, NULL, NULL, -1,
+ imin, imax, &kernel);
- /* Initialise finite differencing variables */
- ptrdiff_t const di = 1;
- ptrdiff_t const dj = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);
- ptrdiff_t const dk = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);
- ptrdiff_t const cdi = sizeof(CCTK_REAL) * di;
- ptrdiff_t const cdj = sizeof(CCTK_REAL) * dj;
- ptrdiff_t const cdk = sizeof(CCTK_REAL) * dk;
- CCTK_REAL const dx = ToReal(CCTK_DELTA_SPACE(0));
- CCTK_REAL const dy = ToReal(CCTK_DELTA_SPACE(1));
- CCTK_REAL const dz = ToReal(CCTK_DELTA_SPACE(2));
- CCTK_REAL const dt = ToReal(CCTK_DELTA_TIME);
- CCTK_REAL const t = ToReal(cctk_time);
- CCTK_REAL const dxi = INV(dx);
- CCTK_REAL const dyi = INV(dy);
- CCTK_REAL const dzi = INV(dz);
- CCTK_REAL const khalf = 0.5;
- CCTK_REAL const kthird = 1/3.0;
- CCTK_REAL const ktwothird = 2.0/3.0;
- CCTK_REAL const kfourthird = 4.0/3.0;
- CCTK_REAL const keightthird = 8.0/3.0;
- CCTK_REAL const hdxi = 0.5 * dxi;
- CCTK_REAL const hdyi = 0.5 * dyi;
- CCTK_REAL const hdzi = 0.5 * dzi;
-
- /* Initialize predefined quantities */
- CCTK_REAL const p1o12dx = 0.0833333333333333333333333333333*INV(dx);
- CCTK_REAL const p1o12dy = 0.0833333333333333333333333333333*INV(dy);
- CCTK_REAL const p1o12dz = 0.0833333333333333333333333333333*INV(dz);
- CCTK_REAL const p1o144dxdy = 0.00694444444444444444444444444444*INV(dx)*INV(dy);
- CCTK_REAL const p1o144dxdz = 0.00694444444444444444444444444444*INV(dx)*INV(dz);
- CCTK_REAL const p1o144dydz = 0.00694444444444444444444444444444*INV(dy)*INV(dz);
- CCTK_REAL const p1o2dx = 0.5*INV(dx);
- CCTK_REAL const p1o2dy = 0.5*INV(dy);
- CCTK_REAL const p1o2dz = 0.5*INV(dz);
- CCTK_REAL const p1o4dxdy = 0.25*INV(dx)*INV(dy);
- CCTK_REAL const p1o4dxdz = 0.25*INV(dx)*INV(dz);
- CCTK_REAL const p1o4dydz = 0.25*INV(dy)*INV(dz);
- CCTK_REAL const p1odx2 = INV(SQR(dx));
- CCTK_REAL const p1ody2 = INV(SQR(dy));
- CCTK_REAL const p1odz2 = INV(SQR(dz));
- CCTK_REAL const pm1o12dx2 = -0.0833333333333333333333333333333*INV(SQR(dx));
- CCTK_REAL const pm1o12dy2 = -0.0833333333333333333333333333333*INV(SQR(dy));
- CCTK_REAL const pm1o12dz2 = -0.0833333333333333333333333333333*INV(SQR(dz));
-
- /* Assign local copies of arrays functions */
-
-
-
- /* Calculate temporaries and arrays functions */
-
- /* Copy local copies back to grid functions */
-
- /* Loop over the grid points */
- #pragma omp parallel
- CCTK_LOOP3 (EM_initial,
- i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2],
- cctk_lsh[0],cctk_lsh[1],cctk_lsh[2])
- {
- ptrdiff_t const index = di*i + dj*j + dk*k;
-
- /* Assign local copies of grid functions */
-
- CCTK_REAL xL = x[index];
- CCTK_REAL yL = y[index];
-
-
- /* Include user supplied include files */
-
- /* Precompute derivatives */
-
- switch(fdOrder)
- {
- case 2:
- break;
-
- case 4:
- break;
- }
-
- /* Calculate temporaries and grid functions */
- CCTK_REAL sigma = 1;
-
- CCTK_REAL El1L = sigma*Cos(2*(xL + yL)*Pi);
-
- CCTK_REAL El2L = (-1 + sigma)*Cos(2*xL*Pi) - sigma*Cos(2*(xL +
- yL)*Pi);
-
- CCTK_REAL El3L = 0;
-
- CCTK_REAL B1L = 0;
-
- CCTK_REAL B2L = 0;
-
- CCTK_REAL B3L = -((-1 + sigma)*Cos(2*xL*Pi)) + sigma*Cos(2*(xL +
- yL)*Pi);
-
- /* Copy local copies back to grid functions */
- B1[index] = B1L;
- B2[index] = B2L;
- B3[index] = B3L;
- El1[index] = El1L;
- El2[index] = El2L;
- El3[index] = El3L;
- }
- CCTK_ENDLOOP3 (EM_initial);
}
extern "C" void EM_initial(CCTK_ARGUMENTS)
@@ -170,7 +238,7 @@ extern "C" void EM_initial(CCTK_ARGUMENTS)
return;
}
- const char *groups[] = {"EMScript::B_group","EMScript::El_group","grid::coordinates"};
+ const char *groups[] = {"My_New_Implementation::B_group","My_New_Implementation::El_group","grid::coordinates"};
GenericFD_AssertGroupStorage(cctkGH, "EM_initial", 3, groups);
switch(fdOrder)
diff --git a/Examples/EMScript/src/RegisterMoL.cc b/Examples/EMScript/src/RegisterMoL.cc
index 749736a..1f1ea18 100644
--- a/Examples/EMScript/src/RegisterMoL.cc
+++ b/Examples/EMScript/src/RegisterMoL.cc
@@ -12,12 +12,12 @@ extern "C" void EMScript_RegisterVars(CCTK_ARGUMENTS)
CCTK_INT ierr = 0;
/* Register all the evolved grid functions with MoL */
- ierr += MoLRegisterEvolved(CCTK_VarIndex("EMScript::B1"), CCTK_VarIndex("EMScript::B1rhs"));
- ierr += MoLRegisterEvolved(CCTK_VarIndex("EMScript::B2"), CCTK_VarIndex("EMScript::B2rhs"));
- ierr += MoLRegisterEvolved(CCTK_VarIndex("EMScript::B3"), CCTK_VarIndex("EMScript::B3rhs"));
- ierr += MoLRegisterEvolved(CCTK_VarIndex("EMScript::El1"), CCTK_VarIndex("EMScript::El1rhs"));
- ierr += MoLRegisterEvolved(CCTK_VarIndex("EMScript::El2"), CCTK_VarIndex("EMScript::El2rhs"));
- ierr += MoLRegisterEvolved(CCTK_VarIndex("EMScript::El3"), CCTK_VarIndex("EMScript::El3rhs"));
+ ierr += MoLRegisterEvolved(CCTK_VarIndex("My_New_Implementation::B1"), CCTK_VarIndex("My_New_Implementation::B1rhs"));
+ ierr += MoLRegisterEvolved(CCTK_VarIndex("My_New_Implementation::B2"), CCTK_VarIndex("My_New_Implementation::B2rhs"));
+ ierr += MoLRegisterEvolved(CCTK_VarIndex("My_New_Implementation::B3"), CCTK_VarIndex("My_New_Implementation::B3rhs"));
+ ierr += MoLRegisterEvolved(CCTK_VarIndex("My_New_Implementation::El1"), CCTK_VarIndex("My_New_Implementation::El1rhs"));
+ ierr += MoLRegisterEvolved(CCTK_VarIndex("My_New_Implementation::El2"), CCTK_VarIndex("My_New_Implementation::El2rhs"));
+ ierr += MoLRegisterEvolved(CCTK_VarIndex("My_New_Implementation::El3"), CCTK_VarIndex("My_New_Implementation::El3rhs"));
/* Register all the evolved Array functions with MoL */
return;
diff --git a/Examples/EMScript/src/RegisterSymmetries.cc b/Examples/EMScript/src/RegisterSymmetries.cc
index a7224c9..29654ac 100644
--- a/Examples/EMScript/src/RegisterSymmetries.cc
+++ b/Examples/EMScript/src/RegisterSymmetries.cc
@@ -19,46 +19,46 @@ extern "C" void EMScript_RegisterSymmetries(CCTK_ARGUMENTS)
sym[0] = -1;
sym[1] = 1;
sym[2] = 1;
- SetCartSymVN(cctkGH, sym, "EMScript::B1");
+ SetCartSymVN(cctkGH, sym, "My_New_Implementation::B1");
sym[0] = 1;
sym[1] = -1;
sym[2] = 1;
- SetCartSymVN(cctkGH, sym, "EMScript::B2");
+ SetCartSymVN(cctkGH, sym, "My_New_Implementation::B2");
sym[0] = 1;
sym[1] = 1;
sym[2] = -1;
- SetCartSymVN(cctkGH, sym, "EMScript::B3");
+ SetCartSymVN(cctkGH, sym, "My_New_Implementation::B3");
sym[0] = -1;
sym[1] = 1;
sym[2] = 1;
- SetCartSymVN(cctkGH, sym, "EMScript::El1");
+ SetCartSymVN(cctkGH, sym, "My_New_Implementation::El1");
sym[0] = 1;
sym[1] = -1;
sym[2] = 1;
- SetCartSymVN(cctkGH, sym, "EMScript::El2");
+ SetCartSymVN(cctkGH, sym, "My_New_Implementation::El2");
sym[0] = 1;
sym[1] = 1;
sym[2] = -1;
- SetCartSymVN(cctkGH, sym, "EMScript::El3");
+ SetCartSymVN(cctkGH, sym, "My_New_Implementation::El3");
sym[0] = 1;
sym[1] = 1;
sym[2] = 1;
- SetCartSymVN(cctkGH, sym, "EMScript::CB");
+ SetCartSymVN(cctkGH, sym, "My_New_Implementation::CB");
sym[0] = 1;
sym[1] = 1;
sym[2] = 1;
- SetCartSymVN(cctkGH, sym, "EMScript::CEl");
+ SetCartSymVN(cctkGH, sym, "My_New_Implementation::CEl");
sym[0] = 1;
sym[1] = 1;
sym[2] = 1;
- SetCartSymVN(cctkGH, sym, "EMScript::rho");
+ SetCartSymVN(cctkGH, sym, "My_New_Implementation::rho");
}