diff options
Diffstat (limited to 'Examples/WaveCaKernel')
| -rw-r--r-- | Examples/WaveCaKernel/cakernel.ccl | 63 | ||||
| -rw-r--r-- | Examples/WaveCaKernel/configuration.ccl | 7 | ||||
| -rw-r--r-- | Examples/WaveCaKernel/interface.ccl | 69 | ||||
| -rw-r--r-- | Examples/WaveCaKernel/param.ccl | 216 | ||||
| -rw-r--r-- | Examples/WaveCaKernel/schedule.ccl | 141 | ||||
| -rw-r--r-- | Examples/WaveCaKernel/src/Boundaries.cc | 249 | ||||
| -rw-r--r-- | Examples/WaveCaKernel/src/CaKernel__calc_bound_rhs.code | 74 | ||||
| -rw-r--r-- | Examples/WaveCaKernel/src/CaKernel__calc_rhs.code | 81 | ||||
| -rw-r--r-- | Examples/WaveCaKernel/src/CaKernel__copy_to_device.code | 77 | ||||
| -rw-r--r-- | Examples/WaveCaKernel/src/Differencing.h | 72 | ||||
| -rw-r--r-- | Examples/WaveCaKernel/src/RegisterMoL.cc | 20 | ||||
| -rw-r--r-- | Examples/WaveCaKernel/src/RegisterSymmetries.cc | 34 | ||||
| -rw-r--r-- | Examples/WaveCaKernel/src/Startup.cc | 10 | ||||
| -rw-r--r-- | Examples/WaveCaKernel/src/initial_gaussian.cc | 135 | ||||
| -rw-r--r-- | Examples/WaveCaKernel/src/make.code.defn | 3 | ||||
| l--------- | Examples/WaveCaKernel/test | 1 |
16 files changed, 1252 insertions, 0 deletions
diff --git a/Examples/WaveCaKernel/cakernel.ccl b/Examples/WaveCaKernel/cakernel.ccl new file mode 100644 index 0000000..2b33cf1 --- /dev/null +++ b/Examples/WaveCaKernel/cakernel.ccl @@ -0,0 +1,63 @@ +CCTK_CUDA_KERNEL calc_rhs TYPE=gpu_cuda/3dblock TILE="8,8,8" SHARECODE=yes STENCIL="1,1,1,1,1,1" +{ + CCTK_CUDA_KERNEL_VARIABLE cached=no intent=in + { + phi + } + "phi" + + CCTK_CUDA_KERNEL_VARIABLE cached=no intent=out + { + phirhs + } + "phirhs" + + CCTK_CUDA_KERNEL_VARIABLE cached=no intent=in + { + pi + } + "pi" + + CCTK_CUDA_KERNEL_VARIABLE cached=no intent=out + { + pirhs + } + "pirhs" +} + +CCTK_CUDA_KERNEL calc_bound_rhs TYPE=gpu_cuda/boundary TILE="8,8,8" SHARECODE=yes +{ + CCTK_CUDA_KERNEL_VARIABLE cached=no intent=out + { + phirhs + } + "phirhs" + + CCTK_CUDA_KERNEL_VARIABLE cached=no intent=out + { + pirhs + } + "pirhs" + + CCTK_CUDA_KERNEL_VARIABLE cached=no intent=in + { + xCopy + } + "xCopy" +} + +CCTK_CUDA_KERNEL copy_to_device TYPE=gpu_cuda/3dblock TILE="8,8,8" SHARECODE=yes STENCIL="0,0,0,0,0,0" +{ + CCTK_CUDA_KERNEL_VARIABLE cached=no intent=inout + { + phi + } + "phi" + + CCTK_CUDA_KERNEL_VARIABLE cached=no intent=inout + { + pi + } + "pi" +} + diff --git a/Examples/WaveCaKernel/configuration.ccl b/Examples/WaveCaKernel/configuration.ccl new file mode 100644 index 0000000..46dbde5 --- /dev/null +++ b/Examples/WaveCaKernel/configuration.ccl @@ -0,0 +1,7 @@ +# File produced by Kranc + +REQUIRES GenericFD +OPTIONAL LoopControl +{ +} +REQUIRES CUDA
\ No newline at end of file diff --git a/Examples/WaveCaKernel/interface.ccl b/Examples/WaveCaKernel/interface.ccl new file mode 100644 index 0000000..11c1416 --- /dev/null +++ b/Examples/WaveCaKernel/interface.ccl @@ -0,0 +1,69 @@ +# File produced by Kranc + +implements: WaveCaKernel + +inherits: Grid GenericFD Boundary + + + +USES INCLUDE: GenericFD.h +USES INCLUDE: Symmetry.h +USES INCLUDE: sbp_calc_coeffs.h +USES INCLUDE: Boundary.h +USES INCLUDE: loopcontrol.h + +CCTK_INT FUNCTION MoLRegisterEvolved(CCTK_INT IN EvolvedIndex, CCTK_INT IN RHSIndex) +USES FUNCTION MoLRegisterEvolved + +SUBROUTINE Diff_coeff(CCTK_POINTER_TO_CONST IN cctkGH, CCTK_INT IN dir, CCTK_INT IN nsize, CCTK_INT OUT ARRAY imin, CCTK_INT OUT ARRAY imax, CCTK_REAL OUT ARRAY q, CCTK_INT IN table_handle) +USES FUNCTION Diff_coeff + +CCTK_INT FUNCTION MultiPatch_GetMap(CCTK_POINTER_TO_CONST IN cctkGH) +USES FUNCTION MultiPatch_GetMap + +CCTK_INT FUNCTION Boundary_SelectGroupForBC(CCTK_POINTER_TO_CONST IN GH, CCTK_INT IN faces, CCTK_INT IN boundary_width, CCTK_INT IN table_handle, CCTK_STRING IN group_name, CCTK_STRING IN bc_name) +USES FUNCTION Boundary_SelectGroupForBC + +CCTK_INT FUNCTION Boundary_SelectVarForBC(CCTK_POINTER_TO_CONST IN GH, CCTK_INT IN faces, CCTK_INT IN boundary_width, CCTK_INT IN table_handle, CCTK_STRING IN var_name, CCTK_STRING IN bc_name) +USES FUNCTION Boundary_SelectVarForBC + +public: +CCTK_REAL xCopy_g type=GF timelevels=1 tags='' +{ + xCopy +} "xCopy_g" + +public: +CCTK_REAL phi_g type=GF timelevels=2 tags='' +{ + phi +} "phi_g" + +public: +CCTK_REAL pi_g type=GF timelevels=2 tags='' +{ + pi +} "pi_g" + +public: +CCTK_REAL phi_grhs type=GF timelevels=2 tags='' +{ + phirhs +} "phi_grhs" + +public: +CCTK_REAL pi_grhs type=GF timelevels=2 tags='' +{ + pirhs +} "pi_grhs" + +# These functions are provided by the CaKernel thorn + +CCTK_INT FUNCTION Device_RegisterMem(CCTK_POINTER IN cctkGH, CCTK_INT IN vi, CCTK_INT IN num_tls) +REQUIRES FUNCTION Device_RegisterMem + +CCTK_INT FUNCTION Device_UnRegisterMem(CCTK_POINTER IN cctkGH, CCTK_INT IN vi) +REQUIRES FUNCTION Device_UnRegisterMem + +CCTK_POINTER FUNCTION Device_GetVarI (CCTK_POINTER IN cctkGH, CCTK_INT IN vi, CCTK_INT IN num_tls) +REQUIRES FUNCTION Device_GetVarI diff --git a/Examples/WaveCaKernel/param.ccl b/Examples/WaveCaKernel/param.ccl new file mode 100644 index 0000000..1fbbee4 --- /dev/null +++ b/Examples/WaveCaKernel/param.ccl @@ -0,0 +1,216 @@ +# File produced by Kranc + + +shares: GenericFD + + + +shares: MethodOfLines + +USES CCTK_INT MoL_Num_Evolved_Vars +USES CCTK_INT MoL_Num_ArrayEvolved_Vars + +restricted: +CCTK_INT verbose "verbose" STEERABLE=ALWAYS +{ + *:* :: "" +} 0 + +restricted: +CCTK_INT WaveCaKernel_MaxNumEvolvedVars "Number of evolved variables used by this thorn" ACCUMULATOR-BASE=MethodofLines::MoL_Num_Evolved_Vars STEERABLE=RECOVER +{ + 2:2 :: "Number of evolved variables used by this thorn" +} 2 + +restricted: +CCTK_INT WaveCaKernel_MaxNumArrayEvolvedVars "Number of Array evolved variables used by this thorn" ACCUMULATOR-BASE=MethodofLines::MoL_Num_ArrayEvolved_Vars STEERABLE=RECOVER +{ + 0:0 :: "Number of Array evolved variables used by this thorn" +} 0 + +restricted: +CCTK_INT timelevels "Number of active timelevels" STEERABLE=RECOVER +{ + 0:2 :: "" +} 2 + +restricted: +CCTK_INT rhs_timelevels "Number of active RHS timelevels" STEERABLE=RECOVER +{ + 0:2 :: "" +} 1 + +restricted: +CCTK_INT other_timelevels "Number of active timelevels for non-evolved grid functions" STEERABLE=RECOVER +{ + 0:2 :: "" +} 1 + +restricted: +CCTK_INT initial_gaussian_calc_every "initial_gaussian_calc_every" STEERABLE=ALWAYS +{ + *:* :: "" +} 1 + +restricted: +CCTK_INT calc_rhs_calc_every "calc_rhs_calc_every" STEERABLE=ALWAYS +{ + *:* :: "" +} 1 + +restricted: +CCTK_INT calc_bound_rhs_calc_every "calc_bound_rhs_calc_every" STEERABLE=ALWAYS +{ + *:* :: "" +} 1 + +restricted: +CCTK_INT copy_to_device_calc_every "copy_to_device_calc_every" STEERABLE=ALWAYS +{ + *:* :: "" +} 1 + +restricted: +CCTK_INT initial_gaussian_calc_offset "initial_gaussian_calc_offset" STEERABLE=ALWAYS +{ + *:* :: "" +} 0 + +restricted: +CCTK_INT calc_rhs_calc_offset "calc_rhs_calc_offset" STEERABLE=ALWAYS +{ + *:* :: "" +} 0 + +restricted: +CCTK_INT calc_bound_rhs_calc_offset "calc_bound_rhs_calc_offset" STEERABLE=ALWAYS +{ + *:* :: "" +} 0 + +restricted: +CCTK_INT copy_to_device_calc_offset "copy_to_device_calc_offset" STEERABLE=ALWAYS +{ + *:* :: "" +} 0 + +private: +KEYWORD phi_bound "Boundary condition to implement" STEERABLE=ALWAYS +{ + "flat" :: "Flat boundary condition" + "none" :: "No boundary condition" + "static" :: "Boundaries held fixed" + "radiative" :: "Radiation boundary condition" + "scalar" :: "Dirichlet boundary condition" + "newrad" :: "Improved radiative boundary condition" + "skip" :: "skip boundary condition code" +} "skip" + +private: +KEYWORD pi_bound "Boundary condition to implement" STEERABLE=ALWAYS +{ + "flat" :: "Flat boundary condition" + "none" :: "No boundary condition" + "static" :: "Boundaries held fixed" + "radiative" :: "Radiation boundary condition" + "scalar" :: "Dirichlet boundary condition" + "newrad" :: "Improved radiative boundary condition" + "skip" :: "skip boundary condition code" +} "skip" + +private: +KEYWORD phi_g_bound "Boundary condition to implement" STEERABLE=ALWAYS +{ + "flat" :: "Flat boundary condition" + "none" :: "No boundary condition" + "static" :: "Boundaries held fixed" + "radiative" :: "Radiation boundary condition" + "scalar" :: "Dirichlet boundary condition" + "newrad" :: "Improved radiative boundary condition" + "skip" :: "skip boundary condition code" +} "none" + +private: +KEYWORD pi_g_bound "Boundary condition to implement" STEERABLE=ALWAYS +{ + "flat" :: "Flat boundary condition" + "none" :: "No boundary condition" + "static" :: "Boundaries held fixed" + "radiative" :: "Radiation boundary condition" + "scalar" :: "Dirichlet boundary condition" + "newrad" :: "Improved radiative boundary condition" + "skip" :: "skip boundary condition code" +} "none" + +private: +CCTK_REAL phi_bound_speed "characteristic speed at boundary" STEERABLE=ALWAYS +{ + "0:*" :: "outgoing characteristic speed > 0" +} 1. + +private: +CCTK_REAL pi_bound_speed "characteristic speed at boundary" STEERABLE=ALWAYS +{ + "0:*" :: "outgoing characteristic speed > 0" +} 1. + +private: +CCTK_REAL phi_g_bound_speed "characteristic speed at boundary" STEERABLE=ALWAYS +{ + "0:*" :: "outgoing characteristic speed > 0" +} 1. + +private: +CCTK_REAL pi_g_bound_speed "characteristic speed at boundary" STEERABLE=ALWAYS +{ + "0:*" :: "outgoing characteristic speed > 0" +} 1. + +private: +CCTK_REAL phi_bound_limit "limit value for r -> infinity" STEERABLE=ALWAYS +{ + "*:*" :: "value of limit value is unrestricted" +} 0. + +private: +CCTK_REAL pi_bound_limit "limit value for r -> infinity" STEERABLE=ALWAYS +{ + "*:*" :: "value of limit value is unrestricted" +} 0. + +private: +CCTK_REAL phi_g_bound_limit "limit value for r -> infinity" STEERABLE=ALWAYS +{ + "*:*" :: "value of limit value is unrestricted" +} 0. + +private: +CCTK_REAL pi_g_bound_limit "limit value for r -> infinity" STEERABLE=ALWAYS +{ + "*:*" :: "value of limit value is unrestricted" +} 0. + +private: +CCTK_REAL phi_bound_scalar "Dirichlet boundary value" STEERABLE=ALWAYS +{ + "*:*" :: "unrestricted" +} 0. + +private: +CCTK_REAL pi_bound_scalar "Dirichlet boundary value" STEERABLE=ALWAYS +{ + "*:*" :: "unrestricted" +} 0. + +private: +CCTK_REAL phi_g_bound_scalar "Dirichlet boundary value" STEERABLE=ALWAYS +{ + "*:*" :: "unrestricted" +} 0. + +private: +CCTK_REAL pi_g_bound_scalar "Dirichlet boundary value" STEERABLE=ALWAYS +{ + "*:*" :: "unrestricted" +} 0. + diff --git a/Examples/WaveCaKernel/schedule.ccl b/Examples/WaveCaKernel/schedule.ccl new file mode 100644 index 0000000..8f54e2a --- /dev/null +++ b/Examples/WaveCaKernel/schedule.ccl @@ -0,0 +1,141 @@ +# File produced by Kranc + + +if (other_timelevels == 1) +{ + STORAGE: xCopy_g[1] +} + +if (timelevels == 1) +{ + STORAGE: phi_g[1] +} +if (timelevels == 2) +{ + STORAGE: phi_g[2] +} + +if (timelevels == 1) +{ + STORAGE: pi_g[1] +} +if (timelevels == 2) +{ + STORAGE: pi_g[2] +} + +if (rhs_timelevels == 1) +{ + STORAGE: phi_grhs[1] +} +if (rhs_timelevels == 2) +{ + STORAGE: phi_grhs[2] +} + +if (rhs_timelevels == 1) +{ + STORAGE: pi_grhs[1] +} +if (rhs_timelevels == 2) +{ + STORAGE: pi_grhs[2] +} + +schedule WaveCaKernel_Startup at STARTUP +{ + LANG: C + OPTIONS: meta +} "create banner" + +schedule WaveCaKernel_RegisterSymmetries in SymmetryRegister +{ + LANG: C + OPTIONS: meta +} "register symmetries" + +schedule initial_gaussian AT INITIAL +{ + LANG: C + READS: grid::coordinates + WRITES: WaveCaKernel::phi_g + WRITES: WaveCaKernel::pi_g + WRITES: WaveCaKernel::xCopy_g +} "initial_gaussian" + +schedule CAKERNEL_Launch_calc_rhs in MoL_CalcRHS +{ + LANG: C + TAGS: Device=1 + READS: WaveCaKernel::phi_g + READS: WaveCaKernel::pi_g + WRITES: WaveCaKernel::phi_grhs + WRITES: WaveCaKernel::pi_grhs +} "calc_rhs" + +schedule CAKERNEL_Launch_calc_rhs at ANALYSIS +{ + LANG: C + SYNC: phi_grhs + SYNC: pi_grhs + TAGS: Device=1 + READS: WaveCaKernel::phi_g + READS: WaveCaKernel::pi_g + WRITES: WaveCaKernel::phi_grhs + WRITES: WaveCaKernel::pi_grhs +} "calc_rhs" + +schedule CAKERNEL_Launch_calc_bound_rhs in MoL_RHSBoundaries +{ + LANG: C + TAGS: Device=1 + READS: WaveCaKernel::xCopy_g + WRITES: WaveCaKernel::phi_grhs + WRITES: WaveCaKernel::pi_grhs +} "calc_bound_rhs" + +schedule CAKERNEL_Launch_calc_bound_rhs at ANALYSIS +{ + LANG: C + SYNC: phi_grhs + SYNC: pi_grhs + TAGS: Device=1 + READS: WaveCaKernel::xCopy_g + WRITES: WaveCaKernel::phi_grhs + WRITES: WaveCaKernel::pi_grhs +} "calc_bound_rhs" + +schedule CAKERNEL_Launch_copy_to_device at INITIAL after initial_gaussian +{ + LANG: C + TAGS: Device=1 + READS: WaveCaKernel::phi_g + READS: WaveCaKernel::pi_g + WRITES: WaveCaKernel::phi_g + WRITES: WaveCaKernel::pi_g +} "copy_to_device" + +schedule WaveCaKernel_SelectBoundConds in MoL_PostStep +{ + LANG: C + OPTIONS: level + SYNC: phi_g + SYNC: pi_g +} "select boundary conditions" + +schedule WaveCaKernel_CheckBoundaries at BASEGRID +{ + LANG: C + OPTIONS: meta +} "check boundaries treatment" + +schedule WaveCaKernel_RegisterVars in MoL_Register +{ + LANG: C + OPTIONS: meta +} "Register Variables for MoL" + +schedule group ApplyBCs as WaveCaKernel_ApplyBCs in MoL_PostStep after WaveCaKernel_SelectBoundConds +{ + # no language specified +} "Apply boundary conditions controlled by thorn Boundary" diff --git a/Examples/WaveCaKernel/src/Boundaries.cc b/Examples/WaveCaKernel/src/Boundaries.cc new file mode 100644 index 0000000..ad79d17 --- /dev/null +++ b/Examples/WaveCaKernel/src/Boundaries.cc @@ -0,0 +1,249 @@ +/* File produced by Kranc */ + +#include "cctk.h" +#include "cctk_Arguments.h" +#include "cctk_Parameters.h" +#include "cctk_Faces.h" +#include "util_Table.h" +#include "Symmetry.h" + + +/* the boundary treatment is split into 3 steps: */ +/* 1. excision */ +/* 2. symmetries */ +/* 3. "other" boundary conditions, e.g. radiative */ + +/* to simplify scheduling and testing, the 3 steps */ +/* are currently applied in separate functions */ + + +extern "C" void WaveCaKernel_CheckBoundaries(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + return; +} + +extern "C" void WaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + CCTK_INT ierr = 0; + + if (CCTK_EQUALS(phi_g_bound, "none" ) || + CCTK_EQUALS(phi_g_bound, "static") || + CCTK_EQUALS(phi_g_bound, "flat" ) || + CCTK_EQUALS(phi_g_bound, "zero" ) ) + { + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, -1, + "WaveCaKernel::phi_g", phi_g_bound); + if (ierr < 0) + CCTK_WARN(0, "Failed to register phi_g_bound BC for WaveCaKernel::phi_g!"); + } + + if (CCTK_EQUALS(pi_g_bound, "none" ) || + CCTK_EQUALS(pi_g_bound, "static") || + CCTK_EQUALS(pi_g_bound, "flat" ) || + CCTK_EQUALS(pi_g_bound, "zero" ) ) + { + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, -1, + "WaveCaKernel::pi_g", pi_g_bound); + if (ierr < 0) + CCTK_WARN(0, "Failed to register pi_g_bound BC for WaveCaKernel::pi_g!"); + } + + if (CCTK_EQUALS(phi_bound, "none" ) || + CCTK_EQUALS(phi_bound, "static") || + CCTK_EQUALS(phi_bound, "flat" ) || + CCTK_EQUALS(phi_bound, "zero" ) ) + { + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, -1, + "WaveCaKernel::phi", phi_bound); + if (ierr < 0) + CCTK_WARN(0, "Failed to register phi_bound BC for WaveCaKernel::phi!"); + } + + if (CCTK_EQUALS(pi_bound, "none" ) || + CCTK_EQUALS(pi_bound, "static") || + CCTK_EQUALS(pi_bound, "flat" ) || + CCTK_EQUALS(pi_bound, "zero" ) ) + { + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, -1, + "WaveCaKernel::pi", pi_bound); + if (ierr < 0) + CCTK_WARN(0, "Failed to register pi_bound BC for WaveCaKernel::pi!"); + } + + if (CCTK_EQUALS(phi_g_bound, "radiative")) + { + /* select radiation boundary condition */ + static CCTK_INT handle_phi_g_bound = -1; + if (handle_phi_g_bound < 0) handle_phi_g_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_phi_g_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_phi_g_bound , phi_g_bound_limit, "LIMIT") < 0) + CCTK_WARN(0, "could not set LIMIT value in table!"); + if (Util_TableSetReal(handle_phi_g_bound ,phi_g_bound_speed, "SPEED") < 0) + CCTK_WARN(0, "could not set SPEED value in table!"); + + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_phi_g_bound, + "WaveCaKernel::phi_g", "Radiation"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Radiation BC for WaveCaKernel::phi_g!"); + + } + + if (CCTK_EQUALS(pi_g_bound, "radiative")) + { + /* select radiation boundary condition */ + static CCTK_INT handle_pi_g_bound = -1; + if (handle_pi_g_bound < 0) handle_pi_g_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_pi_g_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_pi_g_bound , pi_g_bound_limit, "LIMIT") < 0) + CCTK_WARN(0, "could not set LIMIT value in table!"); + if (Util_TableSetReal(handle_pi_g_bound ,pi_g_bound_speed, "SPEED") < 0) + CCTK_WARN(0, "could not set SPEED value in table!"); + + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_pi_g_bound, + "WaveCaKernel::pi_g", "Radiation"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Radiation BC for WaveCaKernel::pi_g!"); + + } + + if (CCTK_EQUALS(phi_bound, "radiative")) + { + /* select radiation boundary condition */ + static CCTK_INT handle_phi_bound = -1; + if (handle_phi_bound < 0) handle_phi_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_phi_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_phi_bound , phi_bound_limit, "LIMIT") < 0) + CCTK_WARN(0, "could not set LIMIT value in table!"); + if (Util_TableSetReal(handle_phi_bound ,phi_bound_speed, "SPEED") < 0) + CCTK_WARN(0, "could not set SPEED value in table!"); + + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_phi_bound, + "WaveCaKernel::phi", "Radiation"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Radiation BC for WaveCaKernel::phi!"); + + } + + if (CCTK_EQUALS(pi_bound, "radiative")) + { + /* select radiation boundary condition */ + static CCTK_INT handle_pi_bound = -1; + if (handle_pi_bound < 0) handle_pi_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_pi_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_pi_bound , pi_bound_limit, "LIMIT") < 0) + CCTK_WARN(0, "could not set LIMIT value in table!"); + if (Util_TableSetReal(handle_pi_bound ,pi_bound_speed, "SPEED") < 0) + CCTK_WARN(0, "could not set SPEED value in table!"); + + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_pi_bound, + "WaveCaKernel::pi", "Radiation"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Radiation BC for WaveCaKernel::pi!"); + + } + + if (CCTK_EQUALS(phi_g_bound, "scalar")) + { + /* select scalar boundary condition */ + static CCTK_INT handle_phi_g_bound = -1; + if (handle_phi_g_bound < 0) handle_phi_g_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_phi_g_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_phi_g_bound ,phi_g_bound_scalar, "SCALAR") < 0) + CCTK_WARN(0, "could not set SCALAR value in table!"); + + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_phi_g_bound, + "WaveCaKernel::phi_g", "scalar"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Scalar BC for WaveCaKernel::phi_g!"); + + } + + if (CCTK_EQUALS(pi_g_bound, "scalar")) + { + /* select scalar boundary condition */ + static CCTK_INT handle_pi_g_bound = -1; + if (handle_pi_g_bound < 0) handle_pi_g_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_pi_g_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_pi_g_bound ,pi_g_bound_scalar, "SCALAR") < 0) + CCTK_WARN(0, "could not set SCALAR value in table!"); + + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_pi_g_bound, + "WaveCaKernel::pi_g", "scalar"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Scalar BC for WaveCaKernel::pi_g!"); + + } + + if (CCTK_EQUALS(phi_bound, "scalar")) + { + /* select scalar boundary condition */ + static CCTK_INT handle_phi_bound = -1; + if (handle_phi_bound < 0) handle_phi_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_phi_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_phi_bound ,phi_bound_scalar, "SCALAR") < 0) + CCTK_WARN(0, "could not set SCALAR value in table!"); + + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_phi_bound, + "WaveCaKernel::phi", "scalar"); + + if (ierr < 0) + CCTK_WARN(0, "Error in registering Scalar BC for WaveCaKernel::phi!"); + + } + + if (CCTK_EQUALS(pi_bound, "scalar")) + { + /* select scalar boundary condition */ + static CCTK_INT handle_pi_bound = -1; + if (handle_pi_bound < 0) handle_pi_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_pi_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_pi_bound ,pi_bound_scalar, "SCALAR") < 0) + CCTK_WARN(0, "could not set SCALAR value in table!"); + + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_pi_bound, + "WaveCaKernel::pi", "scalar"); + + if (ierr < 0) + CCTK_WARN(0, "Error in registering Scalar BC for WaveCaKernel::pi!"); + + } + return; +} + + + +/* template for entries in parameter file: +#$bound$#WaveCaKernel::phi_g_bound = "skip" +#$bound$#WaveCaKernel::phi_g_bound_speed = 1.0 +#$bound$#WaveCaKernel::phi_g_bound_limit = 0.0 +#$bound$#WaveCaKernel::phi_g_bound_scalar = 0.0 + +#$bound$#WaveCaKernel::pi_g_bound = "skip" +#$bound$#WaveCaKernel::pi_g_bound_speed = 1.0 +#$bound$#WaveCaKernel::pi_g_bound_limit = 0.0 +#$bound$#WaveCaKernel::pi_g_bound_scalar = 0.0 + +#$bound$#WaveCaKernel::phi_bound = "skip" +#$bound$#WaveCaKernel::phi_bound_speed = 1.0 +#$bound$#WaveCaKernel::phi_bound_limit = 0.0 +#$bound$#WaveCaKernel::phi_bound_scalar = 0.0 + +#$bound$#WaveCaKernel::pi_bound = "skip" +#$bound$#WaveCaKernel::pi_bound_speed = 1.0 +#$bound$#WaveCaKernel::pi_bound_limit = 0.0 +#$bound$#WaveCaKernel::pi_bound_scalar = 0.0 + +*/ + diff --git a/Examples/WaveCaKernel/src/CaKernel__calc_bound_rhs.code b/Examples/WaveCaKernel/src/CaKernel__calc_bound_rhs.code new file mode 100644 index 0000000..863ad4f --- /dev/null +++ b/Examples/WaveCaKernel/src/CaKernel__calc_bound_rhs.code @@ -0,0 +1,74 @@ +#undef KRANC_DIFF_FUNCTIONS +#define KRANC_C +#include "Differencing.h" +#include "GenericFD.h" + +#undef KRANC_GFOFFSET3D +#define KRANC_GFOFFSET3D(u,i,j,k) I3D(u,i,j,k) + + +/* Define macros used in calculations */ +#define INITVALUE (42) +#define QAD(x) (SQR(SQR(x))) +#define INV(x) ((1.0) / (x)) +#define SQR(x) ((x) * (x)) +#define CUB(x) ((x) * (x) * (x)) + +CAKERNEL_calc_bound_rhs_Begin + + /* Include user-supplied include files */ + + /* Initialise finite differencing variables */ + CCTK_REAL const dx = params.cagh_dx; + CCTK_REAL const dy = params.cagh_dy; + CCTK_REAL const dz = params.cagh_dz; + CCTK_REAL const dt = params.cagh_dt; + CCTK_REAL const t = params.cagh_time; + CCTK_REAL const dxi = INV(dx); + CCTK_REAL const dyi = INV(dy); + CCTK_REAL const dzi = INV(dz); + CCTK_REAL const khalf = 0.5; + CCTK_REAL const kthird = 1/3.0; + CCTK_REAL const ktwothird = 2.0/3.0; + CCTK_REAL const kfourthird = 4.0/3.0; + CCTK_REAL const keightthird = 8.0/3.0; + CCTK_REAL const hdxi = 0.5 * dxi; + CCTK_REAL const hdyi = 0.5 * dyi; + CCTK_REAL const hdzi = 0.5 * dzi; + + /* Initialize predefined quantities */ + CCTK_REAL const p1o2dx = 0.5*INV(dx); + CCTK_REAL const p1o2dy = 0.5*INV(dy); + CCTK_REAL const p1o2dz = 0.5*INV(dz); + CCTK_REAL const p1odx2 = INV(SQR(dx)); + CCTK_REAL const p1ody2 = INV(SQR(dy)); + CCTK_REAL const p1odz2 = INV(SQR(dz)); + + /* Assign local copies of arrays functions */ + + + + /* Calculate temporaries and arrays functions */ + + /* Copy local copies back to grid functions */ + + /* Assign local copies of grid functions */ + + CCTK_REAL xCopyL = I3D(xCopy,0,0,0); + + + /* Include user supplied include files */ + + /* Precompute derivatives */ + + /* Calculate temporaries and grid functions */ + CCTK_REAL phirhsL = -200.*(xCopyL + t)*exp(-100.*SQR(xCopyL + t)); + + CCTK_REAL pirhsL = exp(-100.*SQR(xCopyL + t))*(-200. + + 80000.*xCopyL*t + 40000.*(SQR(xCopyL) + SQR(t))); + + /* Copy local copies back to grid functions */ + I3D(phirhs,0,0,0) = phirhsL; + I3D(pirhs,0,0,0) = pirhsL; + +CAKERNEL_calc_bound_rhs_End diff --git a/Examples/WaveCaKernel/src/CaKernel__calc_rhs.code b/Examples/WaveCaKernel/src/CaKernel__calc_rhs.code new file mode 100644 index 0000000..d2d7fdb --- /dev/null +++ b/Examples/WaveCaKernel/src/CaKernel__calc_rhs.code @@ -0,0 +1,81 @@ +#undef KRANC_DIFF_FUNCTIONS +#define KRANC_C +#include "Differencing.h" +#include "GenericFD.h" + +#undef KRANC_GFOFFSET3D +#define KRANC_GFOFFSET3D(u,i,j,k) I3D(u,i,j,k) + + +/* Define macros used in calculations */ +#define INITVALUE (42) +#define QAD(x) (SQR(SQR(x))) +#define INV(x) ((1.0) / (x)) +#define SQR(x) ((x) * (x)) +#define CUB(x) ((x) * (x) * (x)) + +CAKERNEL_calc_rhs_Begin + + /* Include user-supplied include files */ + + /* Initialise finite differencing variables */ + CCTK_REAL const dx = params.cagh_dx; + CCTK_REAL const dy = params.cagh_dy; + CCTK_REAL const dz = params.cagh_dz; + CCTK_REAL const dt = params.cagh_dt; + CCTK_REAL const t = params.cagh_time; + CCTK_REAL const dxi = INV(dx); + CCTK_REAL const dyi = INV(dy); + CCTK_REAL const dzi = INV(dz); + CCTK_REAL const khalf = 0.5; + CCTK_REAL const kthird = 1/3.0; + CCTK_REAL const ktwothird = 2.0/3.0; + CCTK_REAL const kfourthird = 4.0/3.0; + CCTK_REAL const keightthird = 8.0/3.0; + CCTK_REAL const hdxi = 0.5 * dxi; + CCTK_REAL const hdyi = 0.5 * dyi; + CCTK_REAL const hdzi = 0.5 * dzi; + + /* Initialize predefined quantities */ + CCTK_REAL const p1o2dx = 0.5*INV(dx); + CCTK_REAL const p1o2dy = 0.5*INV(dy); + CCTK_REAL const p1o2dz = 0.5*INV(dz); + CCTK_REAL const p1odx2 = INV(SQR(dx)); + CCTK_REAL const p1ody2 = INV(SQR(dy)); + CCTK_REAL const p1odz2 = INV(SQR(dz)); + + /* Assign local copies of arrays functions */ + + + + /* Calculate temporaries and arrays functions */ + + /* Copy local copies back to grid functions */ + CAKERNEL_calc_rhs_Computations_Begin + + /* Assign local copies of grid functions */ + + CCTK_REAL phiL = I3D(phi,0,0,0); + CCTK_REAL piL = I3D(pi,0,0,0); + + + /* Include user supplied include files */ + + /* Precompute derivatives */ + CCTK_REAL const PDstandard2nd11phi = PDstandard2nd11(phi); + CCTK_REAL const PDstandard2nd22phi = PDstandard2nd22(phi); + CCTK_REAL const PDstandard2nd33phi = PDstandard2nd33(phi); + + /* Calculate temporaries and grid functions */ + CCTK_REAL phirhsL = piL; + + CCTK_REAL pirhsL = PDstandard2nd11phi + PDstandard2nd22phi + + PDstandard2nd33phi; + + /* Copy local copies back to grid functions */ + I3D(phirhs,0,0,0) = phirhsL; + I3D(pirhs,0,0,0) = pirhsL; + + CAKERNEL_calc_rhs_Computations_End + +CAKERNEL_calc_rhs_End diff --git a/Examples/WaveCaKernel/src/CaKernel__copy_to_device.code b/Examples/WaveCaKernel/src/CaKernel__copy_to_device.code new file mode 100644 index 0000000..24b4566 --- /dev/null +++ b/Examples/WaveCaKernel/src/CaKernel__copy_to_device.code @@ -0,0 +1,77 @@ +#undef KRANC_DIFF_FUNCTIONS +#define KRANC_C +#include "Differencing.h" +#include "GenericFD.h" + +#undef KRANC_GFOFFSET3D +#define KRANC_GFOFFSET3D(u,i,j,k) I3D(u,i,j,k) + + +/* Define macros used in calculations */ +#define INITVALUE (42) +#define QAD(x) (SQR(SQR(x))) +#define INV(x) ((1.0) / (x)) +#define SQR(x) ((x) * (x)) +#define CUB(x) ((x) * (x) * (x)) + +CAKERNEL_copy_to_device_Begin + + /* Include user-supplied include files */ + + /* Initialise finite differencing variables */ + CCTK_REAL const dx = params.cagh_dx; + CCTK_REAL const dy = params.cagh_dy; + CCTK_REAL const dz = params.cagh_dz; + CCTK_REAL const dt = params.cagh_dt; + CCTK_REAL const t = params.cagh_time; + CCTK_REAL const dxi = INV(dx); + CCTK_REAL const dyi = INV(dy); + CCTK_REAL const dzi = INV(dz); + CCTK_REAL const khalf = 0.5; + CCTK_REAL const kthird = 1/3.0; + CCTK_REAL const ktwothird = 2.0/3.0; + CCTK_REAL const kfourthird = 4.0/3.0; + CCTK_REAL const keightthird = 8.0/3.0; + CCTK_REAL const hdxi = 0.5 * dxi; + CCTK_REAL const hdyi = 0.5 * dyi; + CCTK_REAL const hdzi = 0.5 * dzi; + + /* Initialize predefined quantities */ + CCTK_REAL const p1o2dx = 0.5*INV(dx); + CCTK_REAL const p1o2dy = 0.5*INV(dy); + CCTK_REAL const p1o2dz = 0.5*INV(dz); + CCTK_REAL const p1odx2 = INV(SQR(dx)); + CCTK_REAL const p1ody2 = INV(SQR(dy)); + CCTK_REAL const p1odz2 = INV(SQR(dz)); + + /* Assign local copies of arrays functions */ + + + + /* Calculate temporaries and arrays functions */ + + /* Copy local copies back to grid functions */ + CAKERNEL_copy_to_device_Computations_Begin + + /* Assign local copies of grid functions */ + + CCTK_REAL phiL = I3D(phi,0,0,0); + CCTK_REAL piL = I3D(pi,0,0,0); + + + /* Include user supplied include files */ + + /* Precompute derivatives */ + + /* Calculate temporaries and grid functions */ + phiL = phiL; + + piL = piL; + + /* Copy local copies back to grid functions */ + I3D(phi,0,0,0) = phiL; + I3D(pi,0,0,0) = piL; + + CAKERNEL_copy_to_device_Computations_End + +CAKERNEL_copy_to_device_End diff --git a/Examples/WaveCaKernel/src/Differencing.h b/Examples/WaveCaKernel/src/Differencing.h new file mode 100644 index 0000000..ef21b0b --- /dev/null +++ b/Examples/WaveCaKernel/src/Differencing.h @@ -0,0 +1,72 @@ +#ifndef KRANC_DIFF_FUNCTIONS +# define PDstandard2nd1(u) ((-KRANC_GFOFFSET3D(u,-1,0,0) + KRANC_GFOFFSET3D(u,1,0,0))*p1o2dx) +#else +# define PDstandard2nd1(u) (PDstandard2nd1_impl(u,p1o2dx,cdj,cdk)) +static CCTK_REAL PDstandard2nd1_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dx, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED; +static CCTK_REAL PDstandard2nd1_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dx, ptrdiff_t const cdj, ptrdiff_t const cdk) +{ + ptrdiff_t const cdi=sizeof(CCTK_REAL); + return (-KRANC_GFOFFSET3D(u,-1,0,0) + KRANC_GFOFFSET3D(u,1,0,0))*p1o2dx; +} +#endif + +#ifndef KRANC_DIFF_FUNCTIONS +# define PDstandard2nd2(u) ((-KRANC_GFOFFSET3D(u,0,-1,0) + KRANC_GFOFFSET3D(u,0,1,0))*p1o2dy) +#else +# define PDstandard2nd2(u) (PDstandard2nd2_impl(u,p1o2dy,cdj,cdk)) +static CCTK_REAL PDstandard2nd2_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dy, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED; +static CCTK_REAL PDstandard2nd2_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dy, ptrdiff_t const cdj, ptrdiff_t const cdk) +{ + ptrdiff_t const cdi=sizeof(CCTK_REAL); + return (-KRANC_GFOFFSET3D(u,0,-1,0) + KRANC_GFOFFSET3D(u,0,1,0))*p1o2dy; +} +#endif + +#ifndef KRANC_DIFF_FUNCTIONS +# define PDstandard2nd3(u) ((-KRANC_GFOFFSET3D(u,0,0,-1) + KRANC_GFOFFSET3D(u,0,0,1))*p1o2dz) +#else +# define PDstandard2nd3(u) (PDstandard2nd3_impl(u,p1o2dz,cdj,cdk)) +static CCTK_REAL PDstandard2nd3_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dz, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED; +static CCTK_REAL PDstandard2nd3_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dz, ptrdiff_t const cdj, ptrdiff_t const cdk) +{ + ptrdiff_t const cdi=sizeof(CCTK_REAL); + return (-KRANC_GFOFFSET3D(u,0,0,-1) + KRANC_GFOFFSET3D(u,0,0,1))*p1o2dz; +} +#endif + +#ifndef KRANC_DIFF_FUNCTIONS +# define PDstandard2nd11(u) ((-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,-1,0,0) + KRANC_GFOFFSET3D(u,1,0,0))*p1odx2) +#else +# define PDstandard2nd11(u) (PDstandard2nd11_impl(u,p1odx2,cdj,cdk)) +static CCTK_REAL PDstandard2nd11_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1odx2, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED; +static CCTK_REAL PDstandard2nd11_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1odx2, ptrdiff_t const cdj, ptrdiff_t const cdk) +{ + ptrdiff_t const cdi=sizeof(CCTK_REAL); + return (-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,-1,0,0) + KRANC_GFOFFSET3D(u,1,0,0))*p1odx2; +} +#endif + +#ifndef KRANC_DIFF_FUNCTIONS +# define PDstandard2nd22(u) ((-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,0,-1,0) + KRANC_GFOFFSET3D(u,0,1,0))*p1ody2) +#else +# define PDstandard2nd22(u) (PDstandard2nd22_impl(u,p1ody2,cdj,cdk)) +static CCTK_REAL PDstandard2nd22_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1ody2, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED; +static CCTK_REAL PDstandard2nd22_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1ody2, ptrdiff_t const cdj, ptrdiff_t const cdk) +{ + ptrdiff_t const cdi=sizeof(CCTK_REAL); + return (-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,0,-1,0) + KRANC_GFOFFSET3D(u,0,1,0))*p1ody2; +} +#endif + +#ifndef KRANC_DIFF_FUNCTIONS +# define PDstandard2nd33(u) ((-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,0,0,-1) + KRANC_GFOFFSET3D(u,0,0,1))*p1odz2) +#else +# define PDstandard2nd33(u) (PDstandard2nd33_impl(u,p1odz2,cdj,cdk)) +static CCTK_REAL PDstandard2nd33_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1odz2, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED; +static CCTK_REAL PDstandard2nd33_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1odz2, ptrdiff_t const cdj, ptrdiff_t const cdk) +{ + ptrdiff_t const cdi=sizeof(CCTK_REAL); + return (-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,0,0,-1) + KRANC_GFOFFSET3D(u,0,0,1))*p1odz2; +} +#endif + diff --git a/Examples/WaveCaKernel/src/RegisterMoL.cc b/Examples/WaveCaKernel/src/RegisterMoL.cc new file mode 100644 index 0000000..938ce66 --- /dev/null +++ b/Examples/WaveCaKernel/src/RegisterMoL.cc @@ -0,0 +1,20 @@ +/* File produced by Kranc */ + +#include "cctk.h" +#include "cctk_Arguments.h" +#include "cctk_Parameters.h" + +extern "C" void WaveCaKernel_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("WaveCaKernel::phi"), CCTK_VarIndex("WaveCaKernel::phirhs")); + ierr += MoLRegisterEvolved(CCTK_VarIndex("WaveCaKernel::pi"), CCTK_VarIndex("WaveCaKernel::pirhs")); + + /* Register all the evolved Array functions with MoL */ + return; +} diff --git a/Examples/WaveCaKernel/src/RegisterSymmetries.cc b/Examples/WaveCaKernel/src/RegisterSymmetries.cc new file mode 100644 index 0000000..7fd4a11 --- /dev/null +++ b/Examples/WaveCaKernel/src/RegisterSymmetries.cc @@ -0,0 +1,34 @@ +/* File produced by Kranc */ + +#include "cctk.h" +#include "cctk_Arguments.h" +#include "cctk_Parameters.h" +#include "Symmetry.h" + +extern "C" void WaveCaKernel_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, "WaveCaKernel::phi"); + + sym[0] = 1; + sym[1] = 1; + sym[2] = 1; + SetCartSymVN(cctkGH, sym, "WaveCaKernel::pi"); + + sym[0] = 1; + sym[1] = 1; + sym[2] = 1; + SetCartSymVN(cctkGH, sym, "WaveCaKernel::xCopy"); + +} diff --git a/Examples/WaveCaKernel/src/Startup.cc b/Examples/WaveCaKernel/src/Startup.cc new file mode 100644 index 0000000..17ace69 --- /dev/null +++ b/Examples/WaveCaKernel/src/Startup.cc @@ -0,0 +1,10 @@ +/* File produced by Kranc */ + +#include "cctk.h" + +extern "C" int WaveCaKernel_Startup(void) +{ + const char * banner = "WaveCaKernel"; + CCTK_RegisterBanner(banner); + return 0; +} diff --git a/Examples/WaveCaKernel/src/initial_gaussian.cc b/Examples/WaveCaKernel/src/initial_gaussian.cc new file mode 100644 index 0000000..dab0d62 --- /dev/null +++ b/Examples/WaveCaKernel/src/initial_gaussian.cc @@ -0,0 +1,135 @@ +/* File produced by Kranc */ + +#define KRANC_C + +#include <assert.h> +#include <math.h> +#include <stdio.h> +#include <stdlib.h> +#include <string.h> +#include "cctk.h" +#include "cctk_Arguments.h" +#include "cctk_Parameters.h" +#include "GenericFD.h" +#include "Differencing.h" +#include "cctk_Loop.h" +#include "loopcontrol.h" + +/* Define macros used in calculations */ +#define INITVALUE (42) +#define QAD(x) (SQR(SQR(x))) +#define INV(x) ((1.0) / (x)) +#define SQR(x) ((x) * (x)) +#define CUB(x) ((x) * (x) * (x)) + +static void initial_gaussian_Body(cGH const * restrict const cctkGH, int const dir, int const face, CCTK_REAL const normal[3], CCTK_REAL const tangentA[3], CCTK_REAL const tangentB[3], int const imin[3], int const imax[3], int const n_subblock_gfs, CCTK_REAL * restrict const subblock_gfs[]) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + + /* Include user-supplied include files */ + + /* Initialise finite differencing variables */ + ptrdiff_t const di = 1; + ptrdiff_t const dj = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0); + ptrdiff_t const dk = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0); + ptrdiff_t const cdi = sizeof(CCTK_REAL) * di; + ptrdiff_t const cdj = sizeof(CCTK_REAL) * dj; + ptrdiff_t const cdk = sizeof(CCTK_REAL) * dk; + CCTK_REAL const dx = ToReal(CCTK_DELTA_SPACE(0)); + CCTK_REAL const dy = ToReal(CCTK_DELTA_SPACE(1)); + CCTK_REAL const dz = ToReal(CCTK_DELTA_SPACE(2)); + CCTK_REAL const dt = ToReal(CCTK_DELTA_TIME); + CCTK_REAL const t = ToReal(cctk_time); + CCTK_REAL const dxi = INV(dx); + CCTK_REAL const dyi = INV(dy); + CCTK_REAL const dzi = INV(dz); + CCTK_REAL const khalf = 0.5; + CCTK_REAL const kthird = 1/3.0; + CCTK_REAL const ktwothird = 2.0/3.0; + CCTK_REAL const kfourthird = 4.0/3.0; + CCTK_REAL const keightthird = 8.0/3.0; + CCTK_REAL const hdxi = 0.5 * dxi; + CCTK_REAL const hdyi = 0.5 * dyi; + CCTK_REAL const hdzi = 0.5 * dzi; + + /* Initialize predefined quantities */ + CCTK_REAL const p1o2dx = 0.5*INV(dx); + CCTK_REAL const p1o2dy = 0.5*INV(dy); + CCTK_REAL const p1o2dz = 0.5*INV(dz); + CCTK_REAL const p1odx2 = INV(SQR(dx)); + CCTK_REAL const p1ody2 = INV(SQR(dy)); + CCTK_REAL const p1odz2 = INV(SQR(dz)); + + /* Assign local copies of arrays functions */ + + + + /* Calculate temporaries and arrays functions */ + + /* Copy local copies back to grid functions */ + + /* Loop over the grid points */ + #pragma omp parallel + CCTK_LOOP3(initial_gaussian, + i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2], + cctk_lsh[0],cctk_lsh[1],cctk_lsh[2]) + { + ptrdiff_t const index = di*i + dj*j + dk*k; + + /* Assign local copies of grid functions */ + + CCTK_REAL xL = x[index]; + + + /* Include user supplied include files */ + + /* Precompute derivatives */ + + /* Calculate temporaries and grid functions */ + CCTK_REAL phiL = exp(-100.*SQR(xL + t)); + + CCTK_REAL piL = -200.*(xL + t)*exp(-100.*SQR(xL + t)); + + CCTK_REAL xCopyL = xL; + + /* Copy local copies back to grid functions */ + phi[index] = phiL; + pi[index] = piL; + xCopy[index] = xCopyL; + } + CCTK_ENDLOOP3(initial_gaussian); +} + +extern "C" void initial_gaussian(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Entering initial_gaussian_Body"); + } + + if (cctk_iteration % initial_gaussian_calc_every != initial_gaussian_calc_offset) + { + return; + } + + const char *const groups[] = { + "grid::coordinates", + "WaveCaKernel::phi_g", + "WaveCaKernel::pi_g", + "WaveCaKernel::xCopy_g"}; + GenericFD_AssertGroupStorage(cctkGH, "initial_gaussian", 4, groups); + + + GenericFD_LoopOverEverything(cctkGH, initial_gaussian_Body); + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Leaving initial_gaussian_Body"); + } +} diff --git a/Examples/WaveCaKernel/src/make.code.defn b/Examples/WaveCaKernel/src/make.code.defn new file mode 100644 index 0000000..c43661b --- /dev/null +++ b/Examples/WaveCaKernel/src/make.code.defn @@ -0,0 +1,3 @@ +# File produced by Kranc + +SRCS = Startup.cc RegisterSymmetries.cc RegisterMoL.cc initial_gaussian.cc Boundaries.cc diff --git a/Examples/WaveCaKernel/test b/Examples/WaveCaKernel/test new file mode 120000 index 0000000..35022b1 --- /dev/null +++ b/Examples/WaveCaKernel/test @@ -0,0 +1 @@ +../tests/WaveCaKernel/test
\ No newline at end of file |