diff options
Diffstat (limited to 'Examples/EMScript/src')
| -rw-r--r-- | Examples/EMScript/src/Boundaries.cc | 457 | ||||
| -rw-r--r-- | Examples/EMScript/src/Differencing.h | 291 | ||||
| -rw-r--r-- | Examples/EMScript/src/EM_constraints.cc | 354 | ||||
| -rw-r--r-- | Examples/EMScript/src/EM_energy.cc | 235 | ||||
| -rw-r--r-- | Examples/EMScript/src/EM_evol.cc | 400 | ||||
| -rw-r--r-- | Examples/EMScript/src/EM_initial.cc | 259 | ||||
| -rw-r--r-- | Examples/EMScript/src/RegisterMoL.cc | 24 | ||||
| -rw-r--r-- | Examples/EMScript/src/RegisterSymmetries.cc | 64 | ||||
| -rw-r--r-- | Examples/EMScript/src/Startup.cc | 10 | ||||
| -rw-r--r-- | Examples/EMScript/src/make.code.defn | 3 |
10 files changed, 2097 insertions, 0 deletions
diff --git a/Examples/EMScript/src/Boundaries.cc b/Examples/EMScript/src/Boundaries.cc new file mode 100644 index 0000000..b091436 --- /dev/null +++ b/Examples/EMScript/src/Boundaries.cc @@ -0,0 +1,457 @@ +/* File produced by Kranc */ + +#include "cctk.h" +#include "cctk_Arguments.h" +#include "cctk_Parameters.h" +#include "cctk_Faces.h" +#include "util_Table.h" +#include "Symmetry.h" + + +/* the boundary treatment is split into 3 steps: */ +/* 1. excision */ +/* 2. symmetries */ +/* 3. "other" boundary conditions, e.g. radiative */ + +/* to simplify scheduling and testing, the 3 steps */ +/* are currently applied in separate functions */ + + +extern "C" void EMScript_CheckBoundaries(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + return; +} + +extern "C" void EMScript_SelectBoundConds(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + CCTK_INT ierr = 0; + + if (CCTK_EQUALS(B_group_bound, "none" ) || + CCTK_EQUALS(B_group_bound, "static") || + CCTK_EQUALS(B_group_bound, "flat" ) || + CCTK_EQUALS(B_group_bound, "zero" ) ) + { + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, -1, + "My_New_Implementation::B_group", B_group_bound); + if (ierr < 0) + CCTK_WARN(0, "Failed to register B_group_bound BC for My_New_Implementation::B_group!"); + } + + if (CCTK_EQUALS(El_group_bound, "none" ) || + CCTK_EQUALS(El_group_bound, "static") || + CCTK_EQUALS(El_group_bound, "flat" ) || + CCTK_EQUALS(El_group_bound, "zero" ) ) + { + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, -1, + "My_New_Implementation::El_group", El_group_bound); + if (ierr < 0) + CCTK_WARN(0, "Failed to register El_group_bound BC for My_New_Implementation::El_group!"); + } + + if (CCTK_EQUALS(B1_bound, "none" ) || + CCTK_EQUALS(B1_bound, "static") || + CCTK_EQUALS(B1_bound, "flat" ) || + CCTK_EQUALS(B1_bound, "zero" ) ) + { + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, -1, + "My_New_Implementation::B1", B1_bound); + if (ierr < 0) + CCTK_WARN(0, "Failed to register B1_bound BC for My_New_Implementation::B1!"); + } + + if (CCTK_EQUALS(B2_bound, "none" ) || + CCTK_EQUALS(B2_bound, "static") || + CCTK_EQUALS(B2_bound, "flat" ) || + CCTK_EQUALS(B2_bound, "zero" ) ) + { + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, -1, + "My_New_Implementation::B2", B2_bound); + if (ierr < 0) + CCTK_WARN(0, "Failed to register B2_bound BC for My_New_Implementation::B2!"); + } + + if (CCTK_EQUALS(B3_bound, "none" ) || + CCTK_EQUALS(B3_bound, "static") || + CCTK_EQUALS(B3_bound, "flat" ) || + CCTK_EQUALS(B3_bound, "zero" ) ) + { + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, -1, + "My_New_Implementation::B3", B3_bound); + if (ierr < 0) + CCTK_WARN(0, "Failed to register B3_bound BC for My_New_Implementation::B3!"); + } + + if (CCTK_EQUALS(El1_bound, "none" ) || + CCTK_EQUALS(El1_bound, "static") || + CCTK_EQUALS(El1_bound, "flat" ) || + CCTK_EQUALS(El1_bound, "zero" ) ) + { + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, -1, + "My_New_Implementation::El1", El1_bound); + if (ierr < 0) + CCTK_WARN(0, "Failed to register El1_bound BC for My_New_Implementation::El1!"); + } + + if (CCTK_EQUALS(El2_bound, "none" ) || + CCTK_EQUALS(El2_bound, "static") || + CCTK_EQUALS(El2_bound, "flat" ) || + CCTK_EQUALS(El2_bound, "zero" ) ) + { + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, -1, + "My_New_Implementation::El2", El2_bound); + if (ierr < 0) + CCTK_WARN(0, "Failed to register El2_bound BC for My_New_Implementation::El2!"); + } + + if (CCTK_EQUALS(El3_bound, "none" ) || + CCTK_EQUALS(El3_bound, "static") || + CCTK_EQUALS(El3_bound, "flat" ) || + CCTK_EQUALS(El3_bound, "zero" ) ) + { + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, -1, + "My_New_Implementation::El3", El3_bound); + if (ierr < 0) + CCTK_WARN(0, "Failed to register El3_bound BC for My_New_Implementation::El3!"); + } + + if (CCTK_EQUALS(B_group_bound, "radiative")) + { + /* select radiation boundary condition */ + static CCTK_INT handle_B_group_bound = -1; + if (handle_B_group_bound < 0) handle_B_group_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_B_group_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_B_group_bound , B_group_bound_limit, "LIMIT") < 0) + CCTK_WARN(0, "could not set LIMIT value in table!"); + if (Util_TableSetReal(handle_B_group_bound ,B_group_bound_speed, "SPEED") < 0) + CCTK_WARN(0, "could not set SPEED value in table!"); + + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_B_group_bound, + "My_New_Implementation::B_group", "Radiation"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Radiation BC for My_New_Implementation::B_group!"); + + } + + if (CCTK_EQUALS(El_group_bound, "radiative")) + { + /* select radiation boundary condition */ + static CCTK_INT handle_El_group_bound = -1; + if (handle_El_group_bound < 0) handle_El_group_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_El_group_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_El_group_bound , El_group_bound_limit, "LIMIT") < 0) + CCTK_WARN(0, "could not set LIMIT value in table!"); + if (Util_TableSetReal(handle_El_group_bound ,El_group_bound_speed, "SPEED") < 0) + CCTK_WARN(0, "could not set SPEED value in table!"); + + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_El_group_bound, + "My_New_Implementation::El_group", "Radiation"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Radiation BC for My_New_Implementation::El_group!"); + + } + + if (CCTK_EQUALS(B1_bound, "radiative")) + { + /* select radiation boundary condition */ + static CCTK_INT handle_B1_bound = -1; + if (handle_B1_bound < 0) handle_B1_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_B1_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_B1_bound , B1_bound_limit, "LIMIT") < 0) + CCTK_WARN(0, "could not set LIMIT value in table!"); + if (Util_TableSetReal(handle_B1_bound ,B1_bound_speed, "SPEED") < 0) + CCTK_WARN(0, "could not set SPEED value in table!"); + + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_B1_bound, + "My_New_Implementation::B1", "Radiation"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Radiation BC for My_New_Implementation::B1!"); + + } + + if (CCTK_EQUALS(B2_bound, "radiative")) + { + /* select radiation boundary condition */ + static CCTK_INT handle_B2_bound = -1; + if (handle_B2_bound < 0) handle_B2_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_B2_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_B2_bound , B2_bound_limit, "LIMIT") < 0) + CCTK_WARN(0, "could not set LIMIT value in table!"); + if (Util_TableSetReal(handle_B2_bound ,B2_bound_speed, "SPEED") < 0) + CCTK_WARN(0, "could not set SPEED value in table!"); + + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_B2_bound, + "My_New_Implementation::B2", "Radiation"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Radiation BC for My_New_Implementation::B2!"); + + } + + if (CCTK_EQUALS(B3_bound, "radiative")) + { + /* select radiation boundary condition */ + static CCTK_INT handle_B3_bound = -1; + if (handle_B3_bound < 0) handle_B3_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_B3_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_B3_bound , B3_bound_limit, "LIMIT") < 0) + CCTK_WARN(0, "could not set LIMIT value in table!"); + if (Util_TableSetReal(handle_B3_bound ,B3_bound_speed, "SPEED") < 0) + CCTK_WARN(0, "could not set SPEED value in table!"); + + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_B3_bound, + "My_New_Implementation::B3", "Radiation"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Radiation BC for My_New_Implementation::B3!"); + + } + + if (CCTK_EQUALS(El1_bound, "radiative")) + { + /* select radiation boundary condition */ + static CCTK_INT handle_El1_bound = -1; + if (handle_El1_bound < 0) handle_El1_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_El1_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_El1_bound , El1_bound_limit, "LIMIT") < 0) + CCTK_WARN(0, "could not set LIMIT value in table!"); + if (Util_TableSetReal(handle_El1_bound ,El1_bound_speed, "SPEED") < 0) + CCTK_WARN(0, "could not set SPEED value in table!"); + + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_El1_bound, + "My_New_Implementation::El1", "Radiation"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Radiation BC for My_New_Implementation::El1!"); + + } + + if (CCTK_EQUALS(El2_bound, "radiative")) + { + /* select radiation boundary condition */ + static CCTK_INT handle_El2_bound = -1; + if (handle_El2_bound < 0) handle_El2_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_El2_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_El2_bound , El2_bound_limit, "LIMIT") < 0) + CCTK_WARN(0, "could not set LIMIT value in table!"); + if (Util_TableSetReal(handle_El2_bound ,El2_bound_speed, "SPEED") < 0) + CCTK_WARN(0, "could not set SPEED value in table!"); + + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_El2_bound, + "My_New_Implementation::El2", "Radiation"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Radiation BC for My_New_Implementation::El2!"); + + } + + if (CCTK_EQUALS(El3_bound, "radiative")) + { + /* select radiation boundary condition */ + static CCTK_INT handle_El3_bound = -1; + if (handle_El3_bound < 0) handle_El3_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_El3_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_El3_bound , El3_bound_limit, "LIMIT") < 0) + CCTK_WARN(0, "could not set LIMIT value in table!"); + if (Util_TableSetReal(handle_El3_bound ,El3_bound_speed, "SPEED") < 0) + CCTK_WARN(0, "could not set SPEED value in table!"); + + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_El3_bound, + "My_New_Implementation::El3", "Radiation"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Radiation BC for My_New_Implementation::El3!"); + + } + + if (CCTK_EQUALS(B_group_bound, "scalar")) + { + /* select scalar boundary condition */ + static CCTK_INT handle_B_group_bound = -1; + if (handle_B_group_bound < 0) handle_B_group_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_B_group_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_B_group_bound ,B_group_bound_scalar, "SCALAR") < 0) + CCTK_WARN(0, "could not set SCALAR value in table!"); + + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_B_group_bound, + "My_New_Implementation::B_group", "scalar"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Scalar BC for My_New_Implementation::B_group!"); + + } + + if (CCTK_EQUALS(El_group_bound, "scalar")) + { + /* select scalar boundary condition */ + static CCTK_INT handle_El_group_bound = -1; + if (handle_El_group_bound < 0) handle_El_group_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_El_group_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_El_group_bound ,El_group_bound_scalar, "SCALAR") < 0) + CCTK_WARN(0, "could not set SCALAR value in table!"); + + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_El_group_bound, + "My_New_Implementation::El_group", "scalar"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Scalar BC for My_New_Implementation::El_group!"); + + } + + if (CCTK_EQUALS(B1_bound, "scalar")) + { + /* select scalar boundary condition */ + static CCTK_INT handle_B1_bound = -1; + if (handle_B1_bound < 0) handle_B1_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_B1_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_B1_bound ,B1_bound_scalar, "SCALAR") < 0) + CCTK_WARN(0, "could not set SCALAR value in table!"); + + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_B1_bound, + "My_New_Implementation::B1", "scalar"); + + if (ierr < 0) + CCTK_WARN(0, "Error in registering Scalar BC for My_New_Implementation::B1!"); + + } + + if (CCTK_EQUALS(B2_bound, "scalar")) + { + /* select scalar boundary condition */ + static CCTK_INT handle_B2_bound = -1; + if (handle_B2_bound < 0) handle_B2_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_B2_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_B2_bound ,B2_bound_scalar, "SCALAR") < 0) + CCTK_WARN(0, "could not set SCALAR value in table!"); + + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_B2_bound, + "My_New_Implementation::B2", "scalar"); + + if (ierr < 0) + CCTK_WARN(0, "Error in registering Scalar BC for My_New_Implementation::B2!"); + + } + + if (CCTK_EQUALS(B3_bound, "scalar")) + { + /* select scalar boundary condition */ + static CCTK_INT handle_B3_bound = -1; + if (handle_B3_bound < 0) handle_B3_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_B3_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_B3_bound ,B3_bound_scalar, "SCALAR") < 0) + CCTK_WARN(0, "could not set SCALAR value in table!"); + + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_B3_bound, + "My_New_Implementation::B3", "scalar"); + + if (ierr < 0) + CCTK_WARN(0, "Error in registering Scalar BC for My_New_Implementation::B3!"); + + } + + if (CCTK_EQUALS(El1_bound, "scalar")) + { + /* select scalar boundary condition */ + static CCTK_INT handle_El1_bound = -1; + if (handle_El1_bound < 0) handle_El1_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_El1_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_El1_bound ,El1_bound_scalar, "SCALAR") < 0) + CCTK_WARN(0, "could not set SCALAR value in table!"); + + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_El1_bound, + "My_New_Implementation::El1", "scalar"); + + if (ierr < 0) + CCTK_WARN(0, "Error in registering Scalar BC for My_New_Implementation::El1!"); + + } + + if (CCTK_EQUALS(El2_bound, "scalar")) + { + /* select scalar boundary condition */ + static CCTK_INT handle_El2_bound = -1; + if (handle_El2_bound < 0) handle_El2_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_El2_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_El2_bound ,El2_bound_scalar, "SCALAR") < 0) + CCTK_WARN(0, "could not set SCALAR value in table!"); + + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_El2_bound, + "My_New_Implementation::El2", "scalar"); + + if (ierr < 0) + CCTK_WARN(0, "Error in registering Scalar BC for My_New_Implementation::El2!"); + + } + + if (CCTK_EQUALS(El3_bound, "scalar")) + { + /* select scalar boundary condition */ + static CCTK_INT handle_El3_bound = -1; + if (handle_El3_bound < 0) handle_El3_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_El3_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_El3_bound ,El3_bound_scalar, "SCALAR") < 0) + CCTK_WARN(0, "could not set SCALAR value in table!"); + + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_El3_bound, + "My_New_Implementation::El3", "scalar"); + + if (ierr < 0) + CCTK_WARN(0, "Error in registering Scalar BC for My_New_Implementation::El3!"); + + } + return; +} + + + +/* template for entries in parameter file: +#$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$#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$#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$#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$#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$#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$#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$#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 new file mode 100644 index 0000000..1933a29 --- /dev/null +++ b/Examples/EMScript/src/Differencing.h @@ -0,0 +1,291 @@ +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 new file mode 100644 index 0000000..fc8e97e --- /dev/null +++ b/Examples/EMScript/src/EM_constraints.cc @@ -0,0 +1,354 @@ +/* File produced by Kranc */ + +#define KRANC_C + +#include <assert.h> +#include <math.h> +#include <stdio.h> +#include <stdlib.h> +#include <string.h> +#include "cctk.h" +#include "cctk_Arguments.h" +#include "cctk_Parameters.h" +#include "GenericFD.h" +#include "Differencing.h" +#include "cctk_Loop.h" +#include "loopcontrol.h" +#include "OpenCLRunTime.h" +#include "vectors.h" + +/* Define macros used in calculations */ +#define INITVALUE (42) +#define QAD(x) (SQR(SQR(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) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + CCTK_INT ierr = 0; + 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 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 My_New_Implementation::CEl_group."); + return; +} + +static void EM_constraints_Body(cGH const * restrict const cctkGH, int const dir, int const face, CCTK_REAL const normal[3], CCTK_REAL const tangentA[3], CCTK_REAL const tangentB[3], int const imin[3], int const imax[3], int const n_subblock_gfs, CCTK_REAL * restrict const subblock_gfs[]) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + char const * const source = + "\n" + "/* 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" + "" + ; + + 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}; + + 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); + +} + +extern "C" void EM_constraints(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Entering EM_constraints_Body"); + } + + if (cctk_iteration % EM_constraints_calc_every != EM_constraints_calc_offset) + { + return; + } + + 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) + { + case 2: + GenericFD_EnsureStencilFits(cctkGH, "EM_constraints", 1, 1, 1); + break; + + case 4: + GenericFD_EnsureStencilFits(cctkGH, "EM_constraints", 2, 2, 2); + break; + } + + GenericFD_LoopOverInterior(cctkGH, &EM_constraints_Body); + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Leaving EM_constraints_Body"); + } +} diff --git a/Examples/EMScript/src/EM_energy.cc b/Examples/EMScript/src/EM_energy.cc new file mode 100644 index 0000000..5a7875c --- /dev/null +++ b/Examples/EMScript/src/EM_energy.cc @@ -0,0 +1,235 @@ +/* File produced by Kranc */ + +#define KRANC_C + +#include <assert.h> +#include <math.h> +#include <stdio.h> +#include <stdlib.h> +#include <string.h> +#include "cctk.h" +#include "cctk_Arguments.h" +#include "cctk_Parameters.h" +#include "GenericFD.h" +#include "Differencing.h" +#include "cctk_Loop.h" +#include "loopcontrol.h" +#include "OpenCLRunTime.h" +#include "vectors.h" + +/* Define macros used in calculations */ +#define INITVALUE (42) +#define QAD(x) (SQR(SQR(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) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + CCTK_INT ierr = 0; + 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 My_New_Implementation::rho_group."); + return; +} + +static void EM_energy_Body(cGH const * restrict const cctkGH, int const dir, int const face, CCTK_REAL const normal[3], CCTK_REAL const tangentA[3], CCTK_REAL const tangentB[3], int const imin[3], int const imax[3], int const n_subblock_gfs, CCTK_REAL * restrict const subblock_gfs[]) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + char const * const source = + "\n" + "/* 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); + +} + +extern "C" void EM_energy(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Entering EM_energy_Body"); + } + + if (cctk_iteration % EM_energy_calc_every != EM_energy_calc_offset) + { + return; + } + + 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) + { + case 2: + break; + + case 4: + break; + } + + GenericFD_LoopOverEverything(cctkGH, &EM_energy_Body); + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Leaving EM_energy_Body"); + } +} diff --git a/Examples/EMScript/src/EM_evol.cc b/Examples/EMScript/src/EM_evol.cc new file mode 100644 index 0000000..4e66517 --- /dev/null +++ b/Examples/EMScript/src/EM_evol.cc @@ -0,0 +1,400 @@ +/* File produced by Kranc */ + +#define KRANC_C + +#include <assert.h> +#include <math.h> +#include <stdio.h> +#include <stdlib.h> +#include <string.h> +#include "cctk.h" +#include "cctk_Arguments.h" +#include "cctk_Parameters.h" +#include "GenericFD.h" +#include "Differencing.h" +#include "cctk_Loop.h" +#include "loopcontrol.h" +#include "OpenCLRunTime.h" +#include "vectors.h" + +/* Define macros used in calculations */ +#define INITVALUE (42) +#define QAD(x) (SQR(SQR(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) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + CCTK_INT ierr = 0; + 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 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 My_New_Implementation::El_grouprhs."); + return; +} + +static void EM_evol_Body(cGH const * restrict const cctkGH, int const dir, int const face, CCTK_REAL const normal[3], CCTK_REAL const tangentA[3], CCTK_REAL const tangentB[3], int const imin[3], int const imax[3], int const n_subblock_gfs, CCTK_REAL * restrict const subblock_gfs[]) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + char const * const source = + "\n" + "/* 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" + "" + ; + + 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}; + + 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); + +} + +extern "C" void EM_evol(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Entering EM_evol_Body"); + } + + if (cctk_iteration % EM_evol_calc_every != EM_evol_calc_offset) + { + return; + } + + 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) + { + case 2: + GenericFD_EnsureStencilFits(cctkGH, "EM_evol", 1, 1, 1); + break; + + case 4: + GenericFD_EnsureStencilFits(cctkGH, "EM_evol", 2, 2, 2); + break; + } + + GenericFD_LoopOverInterior(cctkGH, &EM_evol_Body); + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Leaving EM_evol_Body"); + } +} diff --git a/Examples/EMScript/src/EM_initial.cc b/Examples/EMScript/src/EM_initial.cc new file mode 100644 index 0000000..34e72f1 --- /dev/null +++ b/Examples/EMScript/src/EM_initial.cc @@ -0,0 +1,259 @@ +/* File produced by Kranc */ + +#define KRANC_C + +#include <assert.h> +#include <math.h> +#include <stdio.h> +#include <stdlib.h> +#include <string.h> +#include "cctk.h" +#include "cctk_Arguments.h" +#include "cctk_Parameters.h" +#include "GenericFD.h" +#include "Differencing.h" +#include "cctk_Loop.h" +#include "loopcontrol.h" +#include "OpenCLRunTime.h" +#include "vectors.h" + +/* Define macros used in calculations */ +#define INITVALUE (42) +#define QAD(x) (SQR(SQR(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) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + CCTK_INT ierr = 0; + 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 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 My_New_Implementation::El_group."); + return; +} + +static void EM_initial_Body(cGH const * restrict const cctkGH, int const dir, int const face, CCTK_REAL const normal[3], CCTK_REAL const tangentA[3], CCTK_REAL const tangentB[3], int const imin[3], int const imax[3], int const n_subblock_gfs, CCTK_REAL * restrict const subblock_gfs[]) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + char const * const source = + "\n" + "/* 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" + "" + ; + + char const * const groups[] = {"My_New_Implementation::B_group","My_New_Implementation::El_group","grid::coordinates",NULL}; + + 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); + +} + +extern "C" void EM_initial(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Entering EM_initial_Body"); + } + + if (cctk_iteration % EM_initial_calc_every != EM_initial_calc_offset) + { + return; + } + + const char *groups[] = {"My_New_Implementation::B_group","My_New_Implementation::El_group","grid::coordinates"}; + GenericFD_AssertGroupStorage(cctkGH, "EM_initial", 3, groups); + + switch(fdOrder) + { + case 2: + break; + + case 4: + break; + } + + GenericFD_LoopOverEverything(cctkGH, &EM_initial_Body); + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Leaving EM_initial_Body"); + } +} diff --git a/Examples/EMScript/src/RegisterMoL.cc b/Examples/EMScript/src/RegisterMoL.cc new file mode 100644 index 0000000..1f1ea18 --- /dev/null +++ b/Examples/EMScript/src/RegisterMoL.cc @@ -0,0 +1,24 @@ +/* File produced by Kranc */ + +#include "cctk.h" +#include "cctk_Arguments.h" +#include "cctk_Parameters.h" + +extern "C" void EMScript_RegisterVars(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + CCTK_INT ierr = 0; + + /* Register all the evolved grid functions with MoL */ + ierr += MoLRegisterEvolved(CCTK_VarIndex("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 new file mode 100644 index 0000000..29654ac --- /dev/null +++ b/Examples/EMScript/src/RegisterSymmetries.cc @@ -0,0 +1,64 @@ +/* File produced by Kranc */ + +#include "cctk.h" +#include "cctk_Arguments.h" +#include "cctk_Parameters.h" +#include "Symmetry.h" + +extern "C" void EMScript_RegisterSymmetries(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + + /* array holding symmetry definitions */ + CCTK_INT sym[3]; + + + /* Register symmetries of grid functions */ + sym[0] = -1; + sym[1] = 1; + sym[2] = 1; + SetCartSymVN(cctkGH, sym, "My_New_Implementation::B1"); + + sym[0] = 1; + sym[1] = -1; + sym[2] = 1; + SetCartSymVN(cctkGH, sym, "My_New_Implementation::B2"); + + sym[0] = 1; + sym[1] = 1; + sym[2] = -1; + SetCartSymVN(cctkGH, sym, "My_New_Implementation::B3"); + + sym[0] = -1; + sym[1] = 1; + sym[2] = 1; + SetCartSymVN(cctkGH, sym, "My_New_Implementation::El1"); + + sym[0] = 1; + sym[1] = -1; + sym[2] = 1; + SetCartSymVN(cctkGH, sym, "My_New_Implementation::El2"); + + sym[0] = 1; + sym[1] = 1; + sym[2] = -1; + SetCartSymVN(cctkGH, sym, "My_New_Implementation::El3"); + + sym[0] = 1; + sym[1] = 1; + sym[2] = 1; + SetCartSymVN(cctkGH, sym, "My_New_Implementation::CB"); + + sym[0] = 1; + sym[1] = 1; + sym[2] = 1; + SetCartSymVN(cctkGH, sym, "My_New_Implementation::CEl"); + + sym[0] = 1; + sym[1] = 1; + sym[2] = 1; + SetCartSymVN(cctkGH, sym, "My_New_Implementation::rho"); + +} diff --git a/Examples/EMScript/src/Startup.cc b/Examples/EMScript/src/Startup.cc new file mode 100644 index 0000000..00bdff6 --- /dev/null +++ b/Examples/EMScript/src/Startup.cc @@ -0,0 +1,10 @@ +/* File produced by Kranc */ + +#include "cctk.h" + +extern "C" int EMScript_Startup(void) +{ + const char * banner = "EMScript"; + CCTK_RegisterBanner(banner); + return 0; +} diff --git a/Examples/EMScript/src/make.code.defn b/Examples/EMScript/src/make.code.defn new file mode 100644 index 0000000..0d52b41 --- /dev/null +++ b/Examples/EMScript/src/make.code.defn @@ -0,0 +1,3 @@ +# File produced by Kranc + +SRCS = Startup.cc RegisterMoL.cc RegisterSymmetries.cc EM_initial.cc EM_evol.cc EM_constraints.cc EM_energy.cc Boundaries.cc |