From fa52744053f212b79087a40dac8c9df8d6d40158 Mon Sep 17 00:00:00 2001 From: Ian Hinder Date: Thu, 15 Mar 2012 17:23:09 +0100 Subject: [PATCH] Replace SimpleWaveCaKernel with WaveCaKernel This script now generates WaveCaKernel and WaveHost which do the same thing, one using CaKernel and one using the host. This allows easy cross-comparison between the two methods. --- Examples/SimpleWaveCaKernel/schedule.ccl | 118 --------- .../{SimpleWaveCaKernel.m => WaveCaKernel.m} | 0 .../cakernel.ccl | 30 +-- .../configuration.ccl | 0 .../interface.ccl | 2 +- .../param.ccl | 12 +- Examples/WaveCaKernel/schedule.ccl | 141 ++++++++++ .../src/Boundaries.cc | 84 +++--- .../src/CaKernel__calc_bound_rhs.code | 0 .../src/CaKernel__calc_rhs.code | 0 .../src/CaKernel__copy_to_device.code | 0 .../src/Differencing.h | 0 .../src/RegisterMoL.cc | 6 +- .../src/RegisterSymmetries.cc | 8 +- .../src/Startup.cc | 4 +- Examples/WaveCaKernel/src/initial_gaussian.cc | 135 ++++++++++ .../src/make.code.defn | 0 Examples/WaveCaKernel/test | 1 + Examples/WaveHost/configuration.ccl | 6 + Examples/WaveHost/interface.ccl | 58 ++++ Examples/WaveHost/param.ccl | 204 ++++++++++++++ Examples/WaveHost/schedule.ccl | 127 +++++++++ Examples/WaveHost/src/Boundaries.cc | 249 ++++++++++++++++++ .../src/CaKernel__copy_to_device.code | 77 ++++++ Examples/WaveHost/src/Differencing.h | 72 +++++ Examples/WaveHost/src/RegisterMoL.cc | 20 ++ Examples/WaveHost/src/RegisterSymmetries.cc | 34 +++ Examples/WaveHost/src/Startup.cc | 10 + Examples/WaveHost/src/calc_bound_rhs.cc | 147 +++++++++++ Examples/WaveHost/src/calc_rhs.cc | 153 +++++++++++ .../src/initial_gaussian.cc | 8 +- Examples/WaveHost/src/make.code.defn | 3 + Examples/WaveHost/test | 1 + ...gaussian.par => wavecakernel_gaussian.par} | 0 34 files changed, 1517 insertions(+), 193 deletions(-) delete mode 100644 Examples/SimpleWaveCaKernel/schedule.ccl rename Examples/{SimpleWaveCaKernel.m => WaveCaKernel.m} (100%) rename Examples/{SimpleWaveCaKernel => WaveCaKernel}/cakernel.ccl (100%) rename Examples/{SimpleWaveCaKernel => WaveCaKernel}/configuration.ccl (100%) rename Examples/{SimpleWaveCaKernel => WaveCaKernel}/interface.ccl (98%) rename Examples/{SimpleWaveCaKernel => WaveCaKernel}/param.ccl (93%) create mode 100644 Examples/WaveCaKernel/schedule.ccl rename Examples/{SimpleWaveCaKernel => WaveCaKernel}/src/Boundaries.cc (71%) rename Examples/{SimpleWaveCaKernel => WaveCaKernel}/src/CaKernel__calc_bound_rhs.code (100%) rename Examples/{SimpleWaveCaKernel => WaveCaKernel}/src/CaKernel__calc_rhs.code (100%) rename Examples/{SimpleWaveCaKernel => WaveCaKernel}/src/CaKernel__copy_to_device.code (100%) rename Examples/{SimpleWaveCaKernel => WaveCaKernel}/src/Differencing.h (100%) rename Examples/{SimpleWaveCaKernel => WaveCaKernel}/src/RegisterMoL.cc (51%) rename Examples/{SimpleWaveCaKernel => WaveCaKernel}/src/RegisterSymmetries.cc (64%) rename Examples/{SimpleWaveCaKernel => WaveCaKernel}/src/Startup.cc (50%) create mode 100644 Examples/WaveCaKernel/src/initial_gaussian.cc rename Examples/{SimpleWaveCaKernel => WaveCaKernel}/src/make.code.defn (100%) create mode 120000 Examples/WaveCaKernel/test create mode 100644 Examples/WaveHost/configuration.ccl create mode 100644 Examples/WaveHost/interface.ccl create mode 100644 Examples/WaveHost/param.ccl create mode 100644 Examples/WaveHost/schedule.ccl create mode 100644 Examples/WaveHost/src/Boundaries.cc create mode 100644 Examples/WaveHost/src/CaKernel__copy_to_device.code create mode 100644 Examples/WaveHost/src/Differencing.h create mode 100644 Examples/WaveHost/src/RegisterMoL.cc create mode 100644 Examples/WaveHost/src/RegisterSymmetries.cc create mode 100644 Examples/WaveHost/src/Startup.cc create mode 100644 Examples/WaveHost/src/calc_bound_rhs.cc create mode 100644 Examples/WaveHost/src/calc_rhs.cc rename Examples/{SimpleWaveCaKernel => WaveHost}/src/initial_gaussian.cc (94%) create mode 100644 Examples/WaveHost/src/make.code.defn create mode 120000 Examples/WaveHost/test rename Examples/{simplewavecakernel_gaussian.par => wavecakernel_gaussian.par} (100%) diff --git a/Examples/SimpleWaveCaKernel/schedule.ccl b/Examples/SimpleWaveCaKernel/schedule.ccl deleted file mode 100644 index 5c777d81..00000000 --- a/Examples/SimpleWaveCaKernel/schedule.ccl +++ /dev/null @@ -1,118 +0,0 @@ -# File produced by Kranc - - -if (other_timelevels == 1) -{ - STORAGE: xCopy_g[1] -} - -if (timelevels == 1) -{ - STORAGE: phi_g[1] -} -if (timelevels == 2) -{ - STORAGE: phi_g[2] -} - -if (timelevels == 1) -{ - STORAGE: pi_g[1] -} -if (timelevels == 2) -{ - STORAGE: pi_g[2] -} - -if (rhs_timelevels == 1) -{ - STORAGE: phi_grhs[1] -} -if (rhs_timelevels == 2) -{ - STORAGE: phi_grhs[2] -} - -if (rhs_timelevels == 1) -{ - STORAGE: pi_grhs[1] -} -if (rhs_timelevels == 2) -{ - STORAGE: pi_grhs[2] -} - -schedule SimpleWaveCaKernel_Startup at STARTUP -{ - LANG: C - OPTIONS: meta -} "create banner" - -schedule SimpleWaveCaKernel_RegisterSymmetries in SymmetryRegister -{ - LANG: C - OPTIONS: meta -} "register symmetries" - -schedule initial_gaussian AT INITIAL -{ - LANG: C - READS: grid::coordinates - WRITES: SimpleWaveCaKernel::phi_g - WRITES: SimpleWaveCaKernel::pi_g - WRITES: SimpleWaveCaKernel::xCopy_g -} "initial_gaussian" - -schedule CAKERNEL_Launch_calc_rhs in MoL_CalcRHS -{ - LANG: C - TAGS: Device=1 - READS: SimpleWaveCaKernel::phi_g - READS: SimpleWaveCaKernel::pi_g - WRITES: SimpleWaveCaKernel::phi_grhs - WRITES: SimpleWaveCaKernel::pi_grhs -} "calc_rhs" - -schedule CAKERNEL_Launch_copy_to_device at INITIAL after initial_gaussian -{ - LANG: C - TAGS: Device=1 - READS: SimpleWaveCaKernel::phi_g - READS: SimpleWaveCaKernel::pi_g - WRITES: SimpleWaveCaKernel::phi_g - WRITES: SimpleWaveCaKernel::pi_g -} "copy_to_device" - -schedule CAKERNEL_Launch_calc_bound_rhs in MoL_RHSBoundaries -{ - LANG: C - TAGS: Device=1 - READS: SimpleWaveCaKernel::xCopy_g - WRITES: SimpleWaveCaKernel::phi_grhs - WRITES: SimpleWaveCaKernel::pi_grhs -} "calc_bound_rhs" - -schedule SimpleWaveCaKernel_SelectBoundConds in MoL_PostStep -{ - LANG: C - OPTIONS: level - SYNC: phi_g - SYNC: pi_g -} "select boundary conditions" - -schedule SimpleWaveCaKernel_CheckBoundaries at BASEGRID -{ - LANG: C - OPTIONS: meta -} "check boundaries treatment" - -schedule SimpleWaveCaKernel_RegisterVars in MoL_Register -{ - LANG: C - OPTIONS: meta -} "Register Variables for MoL" - -schedule group ApplyBCs as SimpleWaveCaKernel_ApplyBCs in MoL_PostStep after SimpleWaveCaKernel_SelectBoundConds -{ - # no language specified -} "Apply boundary conditions controlled by thorn Boundary" diff --git a/Examples/SimpleWaveCaKernel.m b/Examples/WaveCaKernel.m similarity index 100% rename from Examples/SimpleWaveCaKernel.m rename to Examples/WaveCaKernel.m diff --git a/Examples/SimpleWaveCaKernel/cakernel.ccl b/Examples/WaveCaKernel/cakernel.ccl similarity index 100% rename from Examples/SimpleWaveCaKernel/cakernel.ccl rename to Examples/WaveCaKernel/cakernel.ccl index 75c394c5..2b33cf12 100644 --- a/Examples/SimpleWaveCaKernel/cakernel.ccl +++ b/Examples/WaveCaKernel/cakernel.ccl @@ -25,21 +25,6 @@ CCTK_CUDA_KERNEL calc_rhs TYPE=gpu_cuda/3dblock TILE="8,8,8" SHARECODE=yes STENC "pirhs" } -CCTK_CUDA_KERNEL copy_to_device TYPE=gpu_cuda/3dblock TILE="8,8,8" SHARECODE=yes STENCIL="0,0,0,0,0,0" -{ - CCTK_CUDA_KERNEL_VARIABLE cached=no intent=inout - { - phi - } - "phi" - - CCTK_CUDA_KERNEL_VARIABLE cached=no intent=inout - { - pi - } - "pi" -} - CCTK_CUDA_KERNEL calc_bound_rhs TYPE=gpu_cuda/boundary TILE="8,8,8" SHARECODE=yes { CCTK_CUDA_KERNEL_VARIABLE cached=no intent=out @@ -61,3 +46,18 @@ CCTK_CUDA_KERNEL calc_bound_rhs TYPE=gpu_cuda/boundary TILE="8,8,8" SHARECODE=ye "xCopy" } +CCTK_CUDA_KERNEL copy_to_device TYPE=gpu_cuda/3dblock TILE="8,8,8" SHARECODE=yes STENCIL="0,0,0,0,0,0" +{ + CCTK_CUDA_KERNEL_VARIABLE cached=no intent=inout + { + phi + } + "phi" + + CCTK_CUDA_KERNEL_VARIABLE cached=no intent=inout + { + pi + } + "pi" +} + diff --git a/Examples/SimpleWaveCaKernel/configuration.ccl b/Examples/WaveCaKernel/configuration.ccl similarity index 100% rename from Examples/SimpleWaveCaKernel/configuration.ccl rename to Examples/WaveCaKernel/configuration.ccl diff --git a/Examples/SimpleWaveCaKernel/interface.ccl b/Examples/WaveCaKernel/interface.ccl similarity index 98% rename from Examples/SimpleWaveCaKernel/interface.ccl rename to Examples/WaveCaKernel/interface.ccl index 710d050c..11c14166 100644 --- a/Examples/SimpleWaveCaKernel/interface.ccl +++ b/Examples/WaveCaKernel/interface.ccl @@ -1,6 +1,6 @@ # File produced by Kranc -implements: SimpleWaveCaKernel +implements: WaveCaKernel inherits: Grid GenericFD Boundary diff --git a/Examples/SimpleWaveCaKernel/param.ccl b/Examples/WaveCaKernel/param.ccl similarity index 93% rename from Examples/SimpleWaveCaKernel/param.ccl rename to Examples/WaveCaKernel/param.ccl index 845e24e9..1fbbee4d 100644 --- a/Examples/SimpleWaveCaKernel/param.ccl +++ b/Examples/WaveCaKernel/param.ccl @@ -17,13 +17,13 @@ CCTK_INT verbose "verbose" STEERABLE=ALWAYS } 0 restricted: -CCTK_INT SimpleWaveCaKernel_MaxNumEvolvedVars "Number of evolved variables used by this thorn" ACCUMULATOR-BASE=MethodofLines::MoL_Num_Evolved_Vars STEERABLE=RECOVER +CCTK_INT WaveCaKernel_MaxNumEvolvedVars "Number of evolved variables used by this thorn" ACCUMULATOR-BASE=MethodofLines::MoL_Num_Evolved_Vars STEERABLE=RECOVER { 2:2 :: "Number of evolved variables used by this thorn" } 2 restricted: -CCTK_INT SimpleWaveCaKernel_MaxNumArrayEvolvedVars "Number of Array evolved variables used by this thorn" ACCUMULATOR-BASE=MethodofLines::MoL_Num_ArrayEvolved_Vars STEERABLE=RECOVER +CCTK_INT WaveCaKernel_MaxNumArrayEvolvedVars "Number of Array evolved variables used by this thorn" ACCUMULATOR-BASE=MethodofLines::MoL_Num_ArrayEvolved_Vars STEERABLE=RECOVER { 0:0 :: "Number of Array evolved variables used by this thorn" } 0 @@ -59,13 +59,13 @@ CCTK_INT calc_rhs_calc_every "calc_rhs_calc_every" STEERABLE=ALWAYS } 1 restricted: -CCTK_INT copy_to_device_calc_every "copy_to_device_calc_every" STEERABLE=ALWAYS +CCTK_INT calc_bound_rhs_calc_every "calc_bound_rhs_calc_every" STEERABLE=ALWAYS { *:* :: "" } 1 restricted: -CCTK_INT calc_bound_rhs_calc_every "calc_bound_rhs_calc_every" STEERABLE=ALWAYS +CCTK_INT copy_to_device_calc_every "copy_to_device_calc_every" STEERABLE=ALWAYS { *:* :: "" } 1 @@ -83,13 +83,13 @@ CCTK_INT calc_rhs_calc_offset "calc_rhs_calc_offset" STEERABLE=ALWAYS } 0 restricted: -CCTK_INT copy_to_device_calc_offset "copy_to_device_calc_offset" STEERABLE=ALWAYS +CCTK_INT calc_bound_rhs_calc_offset "calc_bound_rhs_calc_offset" STEERABLE=ALWAYS { *:* :: "" } 0 restricted: -CCTK_INT calc_bound_rhs_calc_offset "calc_bound_rhs_calc_offset" STEERABLE=ALWAYS +CCTK_INT copy_to_device_calc_offset "copy_to_device_calc_offset" STEERABLE=ALWAYS { *:* :: "" } 0 diff --git a/Examples/WaveCaKernel/schedule.ccl b/Examples/WaveCaKernel/schedule.ccl new file mode 100644 index 00000000..8f54e2ae --- /dev/null +++ b/Examples/WaveCaKernel/schedule.ccl @@ -0,0 +1,141 @@ +# File produced by Kranc + + +if (other_timelevels == 1) +{ + STORAGE: xCopy_g[1] +} + +if (timelevels == 1) +{ + STORAGE: phi_g[1] +} +if (timelevels == 2) +{ + STORAGE: phi_g[2] +} + +if (timelevels == 1) +{ + STORAGE: pi_g[1] +} +if (timelevels == 2) +{ + STORAGE: pi_g[2] +} + +if (rhs_timelevels == 1) +{ + STORAGE: phi_grhs[1] +} +if (rhs_timelevels == 2) +{ + STORAGE: phi_grhs[2] +} + +if (rhs_timelevels == 1) +{ + STORAGE: pi_grhs[1] +} +if (rhs_timelevels == 2) +{ + STORAGE: pi_grhs[2] +} + +schedule WaveCaKernel_Startup at STARTUP +{ + LANG: C + OPTIONS: meta +} "create banner" + +schedule WaveCaKernel_RegisterSymmetries in SymmetryRegister +{ + LANG: C + OPTIONS: meta +} "register symmetries" + +schedule initial_gaussian AT INITIAL +{ + LANG: C + READS: grid::coordinates + WRITES: WaveCaKernel::phi_g + WRITES: WaveCaKernel::pi_g + WRITES: WaveCaKernel::xCopy_g +} "initial_gaussian" + +schedule CAKERNEL_Launch_calc_rhs in MoL_CalcRHS +{ + LANG: C + TAGS: Device=1 + READS: WaveCaKernel::phi_g + READS: WaveCaKernel::pi_g + WRITES: WaveCaKernel::phi_grhs + WRITES: WaveCaKernel::pi_grhs +} "calc_rhs" + +schedule CAKERNEL_Launch_calc_rhs at ANALYSIS +{ + LANG: C + SYNC: phi_grhs + SYNC: pi_grhs + TAGS: Device=1 + READS: WaveCaKernel::phi_g + READS: WaveCaKernel::pi_g + WRITES: WaveCaKernel::phi_grhs + WRITES: WaveCaKernel::pi_grhs +} "calc_rhs" + +schedule CAKERNEL_Launch_calc_bound_rhs in MoL_RHSBoundaries +{ + LANG: C + TAGS: Device=1 + READS: WaveCaKernel::xCopy_g + WRITES: WaveCaKernel::phi_grhs + WRITES: WaveCaKernel::pi_grhs +} "calc_bound_rhs" + +schedule CAKERNEL_Launch_calc_bound_rhs at ANALYSIS +{ + LANG: C + SYNC: phi_grhs + SYNC: pi_grhs + TAGS: Device=1 + READS: WaveCaKernel::xCopy_g + WRITES: WaveCaKernel::phi_grhs + WRITES: WaveCaKernel::pi_grhs +} "calc_bound_rhs" + +schedule CAKERNEL_Launch_copy_to_device at INITIAL after initial_gaussian +{ + LANG: C + TAGS: Device=1 + READS: WaveCaKernel::phi_g + READS: WaveCaKernel::pi_g + WRITES: WaveCaKernel::phi_g + WRITES: WaveCaKernel::pi_g +} "copy_to_device" + +schedule WaveCaKernel_SelectBoundConds in MoL_PostStep +{ + LANG: C + OPTIONS: level + SYNC: phi_g + SYNC: pi_g +} "select boundary conditions" + +schedule WaveCaKernel_CheckBoundaries at BASEGRID +{ + LANG: C + OPTIONS: meta +} "check boundaries treatment" + +schedule WaveCaKernel_RegisterVars in MoL_Register +{ + LANG: C + OPTIONS: meta +} "Register Variables for MoL" + +schedule group ApplyBCs as WaveCaKernel_ApplyBCs in MoL_PostStep after WaveCaKernel_SelectBoundConds +{ + # no language specified +} "Apply boundary conditions controlled by thorn Boundary" diff --git a/Examples/SimpleWaveCaKernel/src/Boundaries.cc b/Examples/WaveCaKernel/src/Boundaries.cc similarity index 71% rename from Examples/SimpleWaveCaKernel/src/Boundaries.cc rename to Examples/WaveCaKernel/src/Boundaries.cc index 11561bfe..ad79d177 100644 --- a/Examples/SimpleWaveCaKernel/src/Boundaries.cc +++ b/Examples/WaveCaKernel/src/Boundaries.cc @@ -17,7 +17,7 @@ /* are currently applied in separate functions */ -extern "C" void SimpleWaveCaKernel_CheckBoundaries(CCTK_ARGUMENTS) +extern "C" void WaveCaKernel_CheckBoundaries(CCTK_ARGUMENTS) { DECLARE_CCTK_ARGUMENTS; DECLARE_CCTK_PARAMETERS; @@ -25,7 +25,7 @@ extern "C" void SimpleWaveCaKernel_CheckBoundaries(CCTK_ARGUMENTS) return; } -extern "C" void SimpleWaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS) +extern "C" void WaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS) { DECLARE_CCTK_ARGUMENTS; DECLARE_CCTK_PARAMETERS; @@ -38,9 +38,9 @@ extern "C" void SimpleWaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS) CCTK_EQUALS(phi_g_bound, "zero" ) ) { ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, -1, - "SimpleWaveCaKernel::phi_g", phi_g_bound); + "WaveCaKernel::phi_g", phi_g_bound); if (ierr < 0) - CCTK_WARN(0, "Failed to register phi_g_bound BC for SimpleWaveCaKernel::phi_g!"); + CCTK_WARN(0, "Failed to register phi_g_bound BC for WaveCaKernel::phi_g!"); } if (CCTK_EQUALS(pi_g_bound, "none" ) || @@ -49,9 +49,9 @@ extern "C" void SimpleWaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS) CCTK_EQUALS(pi_g_bound, "zero" ) ) { ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, -1, - "SimpleWaveCaKernel::pi_g", pi_g_bound); + "WaveCaKernel::pi_g", pi_g_bound); if (ierr < 0) - CCTK_WARN(0, "Failed to register pi_g_bound BC for SimpleWaveCaKernel::pi_g!"); + CCTK_WARN(0, "Failed to register pi_g_bound BC for WaveCaKernel::pi_g!"); } if (CCTK_EQUALS(phi_bound, "none" ) || @@ -60,9 +60,9 @@ extern "C" void SimpleWaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS) CCTK_EQUALS(phi_bound, "zero" ) ) { ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, -1, - "SimpleWaveCaKernel::phi", phi_bound); + "WaveCaKernel::phi", phi_bound); if (ierr < 0) - CCTK_WARN(0, "Failed to register phi_bound BC for SimpleWaveCaKernel::phi!"); + CCTK_WARN(0, "Failed to register phi_bound BC for WaveCaKernel::phi!"); } if (CCTK_EQUALS(pi_bound, "none" ) || @@ -71,9 +71,9 @@ extern "C" void SimpleWaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS) CCTK_EQUALS(pi_bound, "zero" ) ) { ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, -1, - "SimpleWaveCaKernel::pi", pi_bound); + "WaveCaKernel::pi", pi_bound); if (ierr < 0) - CCTK_WARN(0, "Failed to register pi_bound BC for SimpleWaveCaKernel::pi!"); + CCTK_WARN(0, "Failed to register pi_bound BC for WaveCaKernel::pi!"); } if (CCTK_EQUALS(phi_g_bound, "radiative")) @@ -88,10 +88,10 @@ extern "C" void SimpleWaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS) CCTK_WARN(0, "could not set SPEED value in table!"); ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_phi_g_bound, - "SimpleWaveCaKernel::phi_g", "Radiation"); + "WaveCaKernel::phi_g", "Radiation"); if (ierr < 0) - CCTK_WARN(0, "Failed to register Radiation BC for SimpleWaveCaKernel::phi_g!"); + CCTK_WARN(0, "Failed to register Radiation BC for WaveCaKernel::phi_g!"); } @@ -107,10 +107,10 @@ extern "C" void SimpleWaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS) CCTK_WARN(0, "could not set SPEED value in table!"); ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_pi_g_bound, - "SimpleWaveCaKernel::pi_g", "Radiation"); + "WaveCaKernel::pi_g", "Radiation"); if (ierr < 0) - CCTK_WARN(0, "Failed to register Radiation BC for SimpleWaveCaKernel::pi_g!"); + CCTK_WARN(0, "Failed to register Radiation BC for WaveCaKernel::pi_g!"); } @@ -126,10 +126,10 @@ extern "C" void SimpleWaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS) CCTK_WARN(0, "could not set SPEED value in table!"); ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_phi_bound, - "SimpleWaveCaKernel::phi", "Radiation"); + "WaveCaKernel::phi", "Radiation"); if (ierr < 0) - CCTK_WARN(0, "Failed to register Radiation BC for SimpleWaveCaKernel::phi!"); + CCTK_WARN(0, "Failed to register Radiation BC for WaveCaKernel::phi!"); } @@ -145,10 +145,10 @@ extern "C" void SimpleWaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS) CCTK_WARN(0, "could not set SPEED value in table!"); ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_pi_bound, - "SimpleWaveCaKernel::pi", "Radiation"); + "WaveCaKernel::pi", "Radiation"); if (ierr < 0) - CCTK_WARN(0, "Failed to register Radiation BC for SimpleWaveCaKernel::pi!"); + CCTK_WARN(0, "Failed to register Radiation BC for WaveCaKernel::pi!"); } @@ -162,10 +162,10 @@ extern "C" void SimpleWaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS) CCTK_WARN(0, "could not set SCALAR value in table!"); ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_phi_g_bound, - "SimpleWaveCaKernel::phi_g", "scalar"); + "WaveCaKernel::phi_g", "scalar"); if (ierr < 0) - CCTK_WARN(0, "Failed to register Scalar BC for SimpleWaveCaKernel::phi_g!"); + CCTK_WARN(0, "Failed to register Scalar BC for WaveCaKernel::phi_g!"); } @@ -179,10 +179,10 @@ extern "C" void SimpleWaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS) CCTK_WARN(0, "could not set SCALAR value in table!"); ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_pi_g_bound, - "SimpleWaveCaKernel::pi_g", "scalar"); + "WaveCaKernel::pi_g", "scalar"); if (ierr < 0) - CCTK_WARN(0, "Failed to register Scalar BC for SimpleWaveCaKernel::pi_g!"); + CCTK_WARN(0, "Failed to register Scalar BC for WaveCaKernel::pi_g!"); } @@ -196,10 +196,10 @@ extern "C" void SimpleWaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS) CCTK_WARN(0, "could not set SCALAR value in table!"); ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_phi_bound, - "SimpleWaveCaKernel::phi", "scalar"); + "WaveCaKernel::phi", "scalar"); if (ierr < 0) - CCTK_WARN(0, "Error in registering Scalar BC for SimpleWaveCaKernel::phi!"); + CCTK_WARN(0, "Error in registering Scalar BC for WaveCaKernel::phi!"); } @@ -213,10 +213,10 @@ extern "C" void SimpleWaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS) CCTK_WARN(0, "could not set SCALAR value in table!"); ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_pi_bound, - "SimpleWaveCaKernel::pi", "scalar"); + "WaveCaKernel::pi", "scalar"); if (ierr < 0) - CCTK_WARN(0, "Error in registering Scalar BC for SimpleWaveCaKernel::pi!"); + CCTK_WARN(0, "Error in registering Scalar BC for WaveCaKernel::pi!"); } return; @@ -225,25 +225,25 @@ extern "C" void SimpleWaveCaKernel_SelectBoundConds(CCTK_ARGUMENTS) /* template for entries in parameter file: -#$bound$#SimpleWaveCaKernel::phi_g_bound = "skip" -#$bound$#SimpleWaveCaKernel::phi_g_bound_speed = 1.0 -#$bound$#SimpleWaveCaKernel::phi_g_bound_limit = 0.0 -#$bound$#SimpleWaveCaKernel::phi_g_bound_scalar = 0.0 +#$bound$#WaveCaKernel::phi_g_bound = "skip" +#$bound$#WaveCaKernel::phi_g_bound_speed = 1.0 +#$bound$#WaveCaKernel::phi_g_bound_limit = 0.0 +#$bound$#WaveCaKernel::phi_g_bound_scalar = 0.0 -#$bound$#SimpleWaveCaKernel::pi_g_bound = "skip" -#$bound$#SimpleWaveCaKernel::pi_g_bound_speed = 1.0 -#$bound$#SimpleWaveCaKernel::pi_g_bound_limit = 0.0 -#$bound$#SimpleWaveCaKernel::pi_g_bound_scalar = 0.0 +#$bound$#WaveCaKernel::pi_g_bound = "skip" +#$bound$#WaveCaKernel::pi_g_bound_speed = 1.0 +#$bound$#WaveCaKernel::pi_g_bound_limit = 0.0 +#$bound$#WaveCaKernel::pi_g_bound_scalar = 0.0 -#$bound$#SimpleWaveCaKernel::phi_bound = "skip" -#$bound$#SimpleWaveCaKernel::phi_bound_speed = 1.0 -#$bound$#SimpleWaveCaKernel::phi_bound_limit = 0.0 -#$bound$#SimpleWaveCaKernel::phi_bound_scalar = 0.0 +#$bound$#WaveCaKernel::phi_bound = "skip" +#$bound$#WaveCaKernel::phi_bound_speed = 1.0 +#$bound$#WaveCaKernel::phi_bound_limit = 0.0 +#$bound$#WaveCaKernel::phi_bound_scalar = 0.0 -#$bound$#SimpleWaveCaKernel::pi_bound = "skip" -#$bound$#SimpleWaveCaKernel::pi_bound_speed = 1.0 -#$bound$#SimpleWaveCaKernel::pi_bound_limit = 0.0 -#$bound$#SimpleWaveCaKernel::pi_bound_scalar = 0.0 +#$bound$#WaveCaKernel::pi_bound = "skip" +#$bound$#WaveCaKernel::pi_bound_speed = 1.0 +#$bound$#WaveCaKernel::pi_bound_limit = 0.0 +#$bound$#WaveCaKernel::pi_bound_scalar = 0.0 */ diff --git a/Examples/SimpleWaveCaKernel/src/CaKernel__calc_bound_rhs.code b/Examples/WaveCaKernel/src/CaKernel__calc_bound_rhs.code similarity index 100% rename from Examples/SimpleWaveCaKernel/src/CaKernel__calc_bound_rhs.code rename to Examples/WaveCaKernel/src/CaKernel__calc_bound_rhs.code diff --git a/Examples/SimpleWaveCaKernel/src/CaKernel__calc_rhs.code b/Examples/WaveCaKernel/src/CaKernel__calc_rhs.code similarity index 100% rename from Examples/SimpleWaveCaKernel/src/CaKernel__calc_rhs.code rename to Examples/WaveCaKernel/src/CaKernel__calc_rhs.code diff --git a/Examples/SimpleWaveCaKernel/src/CaKernel__copy_to_device.code b/Examples/WaveCaKernel/src/CaKernel__copy_to_device.code similarity index 100% rename from Examples/SimpleWaveCaKernel/src/CaKernel__copy_to_device.code rename to Examples/WaveCaKernel/src/CaKernel__copy_to_device.code diff --git a/Examples/SimpleWaveCaKernel/src/Differencing.h b/Examples/WaveCaKernel/src/Differencing.h similarity index 100% rename from Examples/SimpleWaveCaKernel/src/Differencing.h rename to Examples/WaveCaKernel/src/Differencing.h diff --git a/Examples/SimpleWaveCaKernel/src/RegisterMoL.cc b/Examples/WaveCaKernel/src/RegisterMoL.cc similarity index 51% rename from Examples/SimpleWaveCaKernel/src/RegisterMoL.cc rename to Examples/WaveCaKernel/src/RegisterMoL.cc index ac8c57b4..938ce664 100644 --- a/Examples/SimpleWaveCaKernel/src/RegisterMoL.cc +++ b/Examples/WaveCaKernel/src/RegisterMoL.cc @@ -4,7 +4,7 @@ #include "cctk_Arguments.h" #include "cctk_Parameters.h" -extern "C" void SimpleWaveCaKernel_RegisterVars(CCTK_ARGUMENTS) +extern "C" void WaveCaKernel_RegisterVars(CCTK_ARGUMENTS) { DECLARE_CCTK_ARGUMENTS; DECLARE_CCTK_PARAMETERS; @@ -12,8 +12,8 @@ extern "C" void SimpleWaveCaKernel_RegisterVars(CCTK_ARGUMENTS) CCTK_INT ierr = 0; /* Register all the evolved grid functions with MoL */ - ierr += MoLRegisterEvolved(CCTK_VarIndex("SimpleWaveCaKernel::phi"), CCTK_VarIndex("SimpleWaveCaKernel::phirhs")); - ierr += MoLRegisterEvolved(CCTK_VarIndex("SimpleWaveCaKernel::pi"), CCTK_VarIndex("SimpleWaveCaKernel::pirhs")); + ierr += MoLRegisterEvolved(CCTK_VarIndex("WaveCaKernel::phi"), CCTK_VarIndex("WaveCaKernel::phirhs")); + ierr += MoLRegisterEvolved(CCTK_VarIndex("WaveCaKernel::pi"), CCTK_VarIndex("WaveCaKernel::pirhs")); /* Register all the evolved Array functions with MoL */ return; diff --git a/Examples/SimpleWaveCaKernel/src/RegisterSymmetries.cc b/Examples/WaveCaKernel/src/RegisterSymmetries.cc similarity index 64% rename from Examples/SimpleWaveCaKernel/src/RegisterSymmetries.cc rename to Examples/WaveCaKernel/src/RegisterSymmetries.cc index c099f5b3..7fd4a118 100644 --- a/Examples/SimpleWaveCaKernel/src/RegisterSymmetries.cc +++ b/Examples/WaveCaKernel/src/RegisterSymmetries.cc @@ -5,7 +5,7 @@ #include "cctk_Parameters.h" #include "Symmetry.h" -extern "C" void SimpleWaveCaKernel_RegisterSymmetries(CCTK_ARGUMENTS) +extern "C" void WaveCaKernel_RegisterSymmetries(CCTK_ARGUMENTS) { DECLARE_CCTK_ARGUMENTS; DECLARE_CCTK_PARAMETERS; @@ -19,16 +19,16 @@ extern "C" void SimpleWaveCaKernel_RegisterSymmetries(CCTK_ARGUMENTS) sym[0] = 1; sym[1] = 1; sym[2] = 1; - SetCartSymVN(cctkGH, sym, "SimpleWaveCaKernel::phi"); + SetCartSymVN(cctkGH, sym, "WaveCaKernel::phi"); sym[0] = 1; sym[1] = 1; sym[2] = 1; - SetCartSymVN(cctkGH, sym, "SimpleWaveCaKernel::pi"); + SetCartSymVN(cctkGH, sym, "WaveCaKernel::pi"); sym[0] = 1; sym[1] = 1; sym[2] = 1; - SetCartSymVN(cctkGH, sym, "SimpleWaveCaKernel::xCopy"); + SetCartSymVN(cctkGH, sym, "WaveCaKernel::xCopy"); } diff --git a/Examples/SimpleWaveCaKernel/src/Startup.cc b/Examples/WaveCaKernel/src/Startup.cc similarity index 50% rename from Examples/SimpleWaveCaKernel/src/Startup.cc rename to Examples/WaveCaKernel/src/Startup.cc index 2761c8b6..17ace695 100644 --- a/Examples/SimpleWaveCaKernel/src/Startup.cc +++ b/Examples/WaveCaKernel/src/Startup.cc @@ -2,9 +2,9 @@ #include "cctk.h" -extern "C" int SimpleWaveCaKernel_Startup(void) +extern "C" int WaveCaKernel_Startup(void) { - const char * banner = "SimpleWaveCaKernel"; + const char * banner = "WaveCaKernel"; CCTK_RegisterBanner(banner); return 0; } diff --git a/Examples/WaveCaKernel/src/initial_gaussian.cc b/Examples/WaveCaKernel/src/initial_gaussian.cc new file mode 100644 index 00000000..dab0d628 --- /dev/null +++ b/Examples/WaveCaKernel/src/initial_gaussian.cc @@ -0,0 +1,135 @@ +/* File produced by Kranc */ + +#define KRANC_C + +#include +#include +#include +#include +#include +#include "cctk.h" +#include "cctk_Arguments.h" +#include "cctk_Parameters.h" +#include "GenericFD.h" +#include "Differencing.h" +#include "cctk_Loop.h" +#include "loopcontrol.h" + +/* Define macros used in calculations */ +#define INITVALUE (42) +#define QAD(x) (SQR(SQR(x))) +#define INV(x) ((1.0) / (x)) +#define SQR(x) ((x) * (x)) +#define CUB(x) ((x) * (x) * (x)) + +static void initial_gaussian_Body(cGH const * restrict const cctkGH, int const dir, int const face, CCTK_REAL const normal[3], CCTK_REAL const tangentA[3], CCTK_REAL const tangentB[3], int const imin[3], int const imax[3], int const n_subblock_gfs, CCTK_REAL * restrict const subblock_gfs[]) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + + /* Include user-supplied include files */ + + /* Initialise finite differencing variables */ + ptrdiff_t const di = 1; + ptrdiff_t const dj = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0); + ptrdiff_t const dk = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0); + ptrdiff_t const cdi = sizeof(CCTK_REAL) * di; + ptrdiff_t const cdj = sizeof(CCTK_REAL) * dj; + ptrdiff_t const cdk = sizeof(CCTK_REAL) * dk; + CCTK_REAL const dx = ToReal(CCTK_DELTA_SPACE(0)); + CCTK_REAL const dy = ToReal(CCTK_DELTA_SPACE(1)); + CCTK_REAL const dz = ToReal(CCTK_DELTA_SPACE(2)); + CCTK_REAL const dt = ToReal(CCTK_DELTA_TIME); + CCTK_REAL const t = ToReal(cctk_time); + CCTK_REAL const dxi = INV(dx); + CCTK_REAL const dyi = INV(dy); + CCTK_REAL const dzi = INV(dz); + CCTK_REAL const khalf = 0.5; + CCTK_REAL const kthird = 1/3.0; + CCTK_REAL const ktwothird = 2.0/3.0; + CCTK_REAL const kfourthird = 4.0/3.0; + CCTK_REAL const keightthird = 8.0/3.0; + CCTK_REAL const hdxi = 0.5 * dxi; + CCTK_REAL const hdyi = 0.5 * dyi; + CCTK_REAL const hdzi = 0.5 * dzi; + + /* Initialize predefined quantities */ + CCTK_REAL const p1o2dx = 0.5*INV(dx); + CCTK_REAL const p1o2dy = 0.5*INV(dy); + CCTK_REAL const p1o2dz = 0.5*INV(dz); + CCTK_REAL const p1odx2 = INV(SQR(dx)); + CCTK_REAL const p1ody2 = INV(SQR(dy)); + CCTK_REAL const p1odz2 = INV(SQR(dz)); + + /* Assign local copies of arrays functions */ + + + + /* Calculate temporaries and arrays functions */ + + /* Copy local copies back to grid functions */ + + /* Loop over the grid points */ + #pragma omp parallel + CCTK_LOOP3(initial_gaussian, + i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2], + cctk_lsh[0],cctk_lsh[1],cctk_lsh[2]) + { + ptrdiff_t const index = di*i + dj*j + dk*k; + + /* Assign local copies of grid functions */ + + CCTK_REAL xL = x[index]; + + + /* Include user supplied include files */ + + /* Precompute derivatives */ + + /* Calculate temporaries and grid functions */ + CCTK_REAL phiL = exp(-100.*SQR(xL + t)); + + CCTK_REAL piL = -200.*(xL + t)*exp(-100.*SQR(xL + t)); + + CCTK_REAL xCopyL = xL; + + /* Copy local copies back to grid functions */ + phi[index] = phiL; + pi[index] = piL; + xCopy[index] = xCopyL; + } + CCTK_ENDLOOP3(initial_gaussian); +} + +extern "C" void initial_gaussian(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Entering initial_gaussian_Body"); + } + + if (cctk_iteration % initial_gaussian_calc_every != initial_gaussian_calc_offset) + { + return; + } + + const char *const groups[] = { + "grid::coordinates", + "WaveCaKernel::phi_g", + "WaveCaKernel::pi_g", + "WaveCaKernel::xCopy_g"}; + GenericFD_AssertGroupStorage(cctkGH, "initial_gaussian", 4, groups); + + + GenericFD_LoopOverEverything(cctkGH, initial_gaussian_Body); + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Leaving initial_gaussian_Body"); + } +} diff --git a/Examples/SimpleWaveCaKernel/src/make.code.defn b/Examples/WaveCaKernel/src/make.code.defn similarity index 100% rename from Examples/SimpleWaveCaKernel/src/make.code.defn rename to Examples/WaveCaKernel/src/make.code.defn diff --git a/Examples/WaveCaKernel/test b/Examples/WaveCaKernel/test new file mode 120000 index 00000000..35022b1b --- /dev/null +++ b/Examples/WaveCaKernel/test @@ -0,0 +1 @@ +../tests/WaveCaKernel/test \ No newline at end of file diff --git a/Examples/WaveHost/configuration.ccl b/Examples/WaveHost/configuration.ccl new file mode 100644 index 00000000..0a66ec25 --- /dev/null +++ b/Examples/WaveHost/configuration.ccl @@ -0,0 +1,6 @@ +# File produced by Kranc + +REQUIRES GenericFD +OPTIONAL LoopControl +{ +} diff --git a/Examples/WaveHost/interface.ccl b/Examples/WaveHost/interface.ccl new file mode 100644 index 00000000..b08be613 --- /dev/null +++ b/Examples/WaveHost/interface.ccl @@ -0,0 +1,58 @@ +# File produced by Kranc + +implements: WaveHost + +inherits: Grid GenericFD Boundary + + + +USES INCLUDE: GenericFD.h +USES INCLUDE: Symmetry.h +USES INCLUDE: sbp_calc_coeffs.h +USES INCLUDE: Boundary.h +USES INCLUDE: loopcontrol.h + +CCTK_INT FUNCTION MoLRegisterEvolved(CCTK_INT IN EvolvedIndex, CCTK_INT IN RHSIndex) +USES FUNCTION MoLRegisterEvolved + +SUBROUTINE Diff_coeff(CCTK_POINTER_TO_CONST IN cctkGH, CCTK_INT IN dir, CCTK_INT IN nsize, CCTK_INT OUT ARRAY imin, CCTK_INT OUT ARRAY imax, CCTK_REAL OUT ARRAY q, CCTK_INT IN table_handle) +USES FUNCTION Diff_coeff + +CCTK_INT FUNCTION MultiPatch_GetMap(CCTK_POINTER_TO_CONST IN cctkGH) +USES FUNCTION MultiPatch_GetMap + +CCTK_INT FUNCTION Boundary_SelectGroupForBC(CCTK_POINTER_TO_CONST IN GH, CCTK_INT IN faces, CCTK_INT IN boundary_width, CCTK_INT IN table_handle, CCTK_STRING IN group_name, CCTK_STRING IN bc_name) +USES FUNCTION Boundary_SelectGroupForBC + +CCTK_INT FUNCTION Boundary_SelectVarForBC(CCTK_POINTER_TO_CONST IN GH, CCTK_INT IN faces, CCTK_INT IN boundary_width, CCTK_INT IN table_handle, CCTK_STRING IN var_name, CCTK_STRING IN bc_name) +USES FUNCTION Boundary_SelectVarForBC + +public: +CCTK_REAL xCopy_g type=GF timelevels=1 tags='' +{ + xCopy +} "xCopy_g" + +public: +CCTK_REAL phi_g type=GF timelevels=2 tags='' +{ + phi +} "phi_g" + +public: +CCTK_REAL pi_g type=GF timelevels=2 tags='' +{ + pi +} "pi_g" + +public: +CCTK_REAL phi_grhs type=GF timelevels=2 tags='' +{ + phirhs +} "phi_grhs" + +public: +CCTK_REAL pi_grhs type=GF timelevels=2 tags='' +{ + pirhs +} "pi_grhs" diff --git a/Examples/WaveHost/param.ccl b/Examples/WaveHost/param.ccl new file mode 100644 index 00000000..566e0c2f --- /dev/null +++ b/Examples/WaveHost/param.ccl @@ -0,0 +1,204 @@ +# File produced by Kranc + + +shares: GenericFD + + + +shares: MethodOfLines + +USES CCTK_INT MoL_Num_Evolved_Vars +USES CCTK_INT MoL_Num_ArrayEvolved_Vars + +restricted: +CCTK_INT verbose "verbose" STEERABLE=ALWAYS +{ + *:* :: "" +} 0 + +restricted: +CCTK_INT WaveHost_MaxNumEvolvedVars "Number of evolved variables used by this thorn" ACCUMULATOR-BASE=MethodofLines::MoL_Num_Evolved_Vars STEERABLE=RECOVER +{ + 2:2 :: "Number of evolved variables used by this thorn" +} 2 + +restricted: +CCTK_INT WaveHost_MaxNumArrayEvolvedVars "Number of Array evolved variables used by this thorn" ACCUMULATOR-BASE=MethodofLines::MoL_Num_ArrayEvolved_Vars STEERABLE=RECOVER +{ + 0:0 :: "Number of Array evolved variables used by this thorn" +} 0 + +restricted: +CCTK_INT timelevels "Number of active timelevels" STEERABLE=RECOVER +{ + 0:2 :: "" +} 2 + +restricted: +CCTK_INT rhs_timelevels "Number of active RHS timelevels" STEERABLE=RECOVER +{ + 0:2 :: "" +} 1 + +restricted: +CCTK_INT other_timelevels "Number of active timelevels for non-evolved grid functions" STEERABLE=RECOVER +{ + 0:2 :: "" +} 1 + +restricted: +CCTK_INT initial_gaussian_calc_every "initial_gaussian_calc_every" STEERABLE=ALWAYS +{ + *:* :: "" +} 1 + +restricted: +CCTK_INT calc_rhs_calc_every "calc_rhs_calc_every" STEERABLE=ALWAYS +{ + *:* :: "" +} 1 + +restricted: +CCTK_INT calc_bound_rhs_calc_every "calc_bound_rhs_calc_every" STEERABLE=ALWAYS +{ + *:* :: "" +} 1 + +restricted: +CCTK_INT initial_gaussian_calc_offset "initial_gaussian_calc_offset" STEERABLE=ALWAYS +{ + *:* :: "" +} 0 + +restricted: +CCTK_INT calc_rhs_calc_offset "calc_rhs_calc_offset" STEERABLE=ALWAYS +{ + *:* :: "" +} 0 + +restricted: +CCTK_INT calc_bound_rhs_calc_offset "calc_bound_rhs_calc_offset" STEERABLE=ALWAYS +{ + *:* :: "" +} 0 + +private: +KEYWORD phi_bound "Boundary condition to implement" STEERABLE=ALWAYS +{ + "flat" :: "Flat boundary condition" + "none" :: "No boundary condition" + "static" :: "Boundaries held fixed" + "radiative" :: "Radiation boundary condition" + "scalar" :: "Dirichlet boundary condition" + "newrad" :: "Improved radiative boundary condition" + "skip" :: "skip boundary condition code" +} "skip" + +private: +KEYWORD pi_bound "Boundary condition to implement" STEERABLE=ALWAYS +{ + "flat" :: "Flat boundary condition" + "none" :: "No boundary condition" + "static" :: "Boundaries held fixed" + "radiative" :: "Radiation boundary condition" + "scalar" :: "Dirichlet boundary condition" + "newrad" :: "Improved radiative boundary condition" + "skip" :: "skip boundary condition code" +} "skip" + +private: +KEYWORD phi_g_bound "Boundary condition to implement" STEERABLE=ALWAYS +{ + "flat" :: "Flat boundary condition" + "none" :: "No boundary condition" + "static" :: "Boundaries held fixed" + "radiative" :: "Radiation boundary condition" + "scalar" :: "Dirichlet boundary condition" + "newrad" :: "Improved radiative boundary condition" + "skip" :: "skip boundary condition code" +} "none" + +private: +KEYWORD pi_g_bound "Boundary condition to implement" STEERABLE=ALWAYS +{ + "flat" :: "Flat boundary condition" + "none" :: "No boundary condition" + "static" :: "Boundaries held fixed" + "radiative" :: "Radiation boundary condition" + "scalar" :: "Dirichlet boundary condition" + "newrad" :: "Improved radiative boundary condition" + "skip" :: "skip boundary condition code" +} "none" + +private: +CCTK_REAL phi_bound_speed "characteristic speed at boundary" STEERABLE=ALWAYS +{ + "0:*" :: "outgoing characteristic speed > 0" +} 1. + +private: +CCTK_REAL pi_bound_speed "characteristic speed at boundary" STEERABLE=ALWAYS +{ + "0:*" :: "outgoing characteristic speed > 0" +} 1. + +private: +CCTK_REAL phi_g_bound_speed "characteristic speed at boundary" STEERABLE=ALWAYS +{ + "0:*" :: "outgoing characteristic speed > 0" +} 1. + +private: +CCTK_REAL pi_g_bound_speed "characteristic speed at boundary" STEERABLE=ALWAYS +{ + "0:*" :: "outgoing characteristic speed > 0" +} 1. + +private: +CCTK_REAL phi_bound_limit "limit value for r -> infinity" STEERABLE=ALWAYS +{ + "*:*" :: "value of limit value is unrestricted" +} 0. + +private: +CCTK_REAL pi_bound_limit "limit value for r -> infinity" STEERABLE=ALWAYS +{ + "*:*" :: "value of limit value is unrestricted" +} 0. + +private: +CCTK_REAL phi_g_bound_limit "limit value for r -> infinity" STEERABLE=ALWAYS +{ + "*:*" :: "value of limit value is unrestricted" +} 0. + +private: +CCTK_REAL pi_g_bound_limit "limit value for r -> infinity" STEERABLE=ALWAYS +{ + "*:*" :: "value of limit value is unrestricted" +} 0. + +private: +CCTK_REAL phi_bound_scalar "Dirichlet boundary value" STEERABLE=ALWAYS +{ + "*:*" :: "unrestricted" +} 0. + +private: +CCTK_REAL pi_bound_scalar "Dirichlet boundary value" STEERABLE=ALWAYS +{ + "*:*" :: "unrestricted" +} 0. + +private: +CCTK_REAL phi_g_bound_scalar "Dirichlet boundary value" STEERABLE=ALWAYS +{ + "*:*" :: "unrestricted" +} 0. + +private: +CCTK_REAL pi_g_bound_scalar "Dirichlet boundary value" STEERABLE=ALWAYS +{ + "*:*" :: "unrestricted" +} 0. + diff --git a/Examples/WaveHost/schedule.ccl b/Examples/WaveHost/schedule.ccl new file mode 100644 index 00000000..b11880ef --- /dev/null +++ b/Examples/WaveHost/schedule.ccl @@ -0,0 +1,127 @@ +# File produced by Kranc + + +if (other_timelevels == 1) +{ + STORAGE: xCopy_g[1] +} + +if (timelevels == 1) +{ + STORAGE: phi_g[1] +} +if (timelevels == 2) +{ + STORAGE: phi_g[2] +} + +if (timelevels == 1) +{ + STORAGE: pi_g[1] +} +if (timelevels == 2) +{ + STORAGE: pi_g[2] +} + +if (rhs_timelevels == 1) +{ + STORAGE: phi_grhs[1] +} +if (rhs_timelevels == 2) +{ + STORAGE: phi_grhs[2] +} + +if (rhs_timelevels == 1) +{ + STORAGE: pi_grhs[1] +} +if (rhs_timelevels == 2) +{ + STORAGE: pi_grhs[2] +} + +schedule WaveHost_Startup at STARTUP +{ + LANG: C + OPTIONS: meta +} "create banner" + +schedule WaveHost_RegisterSymmetries in SymmetryRegister +{ + LANG: C + OPTIONS: meta +} "register symmetries" + +schedule initial_gaussian AT INITIAL +{ + LANG: C + READS: grid::coordinates + WRITES: WaveHost::phi_g + WRITES: WaveHost::pi_g + WRITES: WaveHost::xCopy_g +} "initial_gaussian" + +schedule calc_rhs in MoL_CalcRHS +{ + LANG: C + READS: WaveHost::phi_g + READS: WaveHost::pi_g + WRITES: WaveHost::phi_grhs + WRITES: WaveHost::pi_grhs +} "calc_rhs" + +schedule calc_rhs at ANALYSIS +{ + LANG: C + SYNC: phi_grhs + SYNC: pi_grhs + READS: WaveHost::phi_g + READS: WaveHost::pi_g + WRITES: WaveHost::phi_grhs + WRITES: WaveHost::pi_grhs +} "calc_rhs" + +schedule calc_bound_rhs in MoL_RHSBoundaries +{ + LANG: C + READS: WaveHost::xCopy_g + WRITES: WaveHost::phi_grhs + WRITES: WaveHost::pi_grhs +} "calc_bound_rhs" + +schedule calc_bound_rhs at ANALYSIS +{ + LANG: C + SYNC: phi_grhs + SYNC: pi_grhs + READS: WaveHost::xCopy_g + WRITES: WaveHost::phi_grhs + WRITES: WaveHost::pi_grhs +} "calc_bound_rhs" + +schedule WaveHost_SelectBoundConds in MoL_PostStep +{ + LANG: C + OPTIONS: level + SYNC: phi_g + SYNC: pi_g +} "select boundary conditions" + +schedule WaveHost_CheckBoundaries at BASEGRID +{ + LANG: C + OPTIONS: meta +} "check boundaries treatment" + +schedule WaveHost_RegisterVars in MoL_Register +{ + LANG: C + OPTIONS: meta +} "Register Variables for MoL" + +schedule group ApplyBCs as WaveHost_ApplyBCs in MoL_PostStep after WaveHost_SelectBoundConds +{ + # no language specified +} "Apply boundary conditions controlled by thorn Boundary" diff --git a/Examples/WaveHost/src/Boundaries.cc b/Examples/WaveHost/src/Boundaries.cc new file mode 100644 index 00000000..8fe5d69f --- /dev/null +++ b/Examples/WaveHost/src/Boundaries.cc @@ -0,0 +1,249 @@ +/* File produced by Kranc */ + +#include "cctk.h" +#include "cctk_Arguments.h" +#include "cctk_Parameters.h" +#include "cctk_Faces.h" +#include "util_Table.h" +#include "Symmetry.h" + + +/* the boundary treatment is split into 3 steps: */ +/* 1. excision */ +/* 2. symmetries */ +/* 3. "other" boundary conditions, e.g. radiative */ + +/* to simplify scheduling and testing, the 3 steps */ +/* are currently applied in separate functions */ + + +extern "C" void WaveHost_CheckBoundaries(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + return; +} + +extern "C" void WaveHost_SelectBoundConds(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + CCTK_INT ierr = 0; + + if (CCTK_EQUALS(phi_g_bound, "none" ) || + CCTK_EQUALS(phi_g_bound, "static") || + CCTK_EQUALS(phi_g_bound, "flat" ) || + CCTK_EQUALS(phi_g_bound, "zero" ) ) + { + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, -1, + "WaveHost::phi_g", phi_g_bound); + if (ierr < 0) + CCTK_WARN(0, "Failed to register phi_g_bound BC for WaveHost::phi_g!"); + } + + if (CCTK_EQUALS(pi_g_bound, "none" ) || + CCTK_EQUALS(pi_g_bound, "static") || + CCTK_EQUALS(pi_g_bound, "flat" ) || + CCTK_EQUALS(pi_g_bound, "zero" ) ) + { + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, -1, + "WaveHost::pi_g", pi_g_bound); + if (ierr < 0) + CCTK_WARN(0, "Failed to register pi_g_bound BC for WaveHost::pi_g!"); + } + + if (CCTK_EQUALS(phi_bound, "none" ) || + CCTK_EQUALS(phi_bound, "static") || + CCTK_EQUALS(phi_bound, "flat" ) || + CCTK_EQUALS(phi_bound, "zero" ) ) + { + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, -1, + "WaveHost::phi", phi_bound); + if (ierr < 0) + CCTK_WARN(0, "Failed to register phi_bound BC for WaveHost::phi!"); + } + + if (CCTK_EQUALS(pi_bound, "none" ) || + CCTK_EQUALS(pi_bound, "static") || + CCTK_EQUALS(pi_bound, "flat" ) || + CCTK_EQUALS(pi_bound, "zero" ) ) + { + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, -1, + "WaveHost::pi", pi_bound); + if (ierr < 0) + CCTK_WARN(0, "Failed to register pi_bound BC for WaveHost::pi!"); + } + + if (CCTK_EQUALS(phi_g_bound, "radiative")) + { + /* select radiation boundary condition */ + static CCTK_INT handle_phi_g_bound = -1; + if (handle_phi_g_bound < 0) handle_phi_g_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_phi_g_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_phi_g_bound , phi_g_bound_limit, "LIMIT") < 0) + CCTK_WARN(0, "could not set LIMIT value in table!"); + if (Util_TableSetReal(handle_phi_g_bound ,phi_g_bound_speed, "SPEED") < 0) + CCTK_WARN(0, "could not set SPEED value in table!"); + + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_phi_g_bound, + "WaveHost::phi_g", "Radiation"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Radiation BC for WaveHost::phi_g!"); + + } + + if (CCTK_EQUALS(pi_g_bound, "radiative")) + { + /* select radiation boundary condition */ + static CCTK_INT handle_pi_g_bound = -1; + if (handle_pi_g_bound < 0) handle_pi_g_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_pi_g_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_pi_g_bound , pi_g_bound_limit, "LIMIT") < 0) + CCTK_WARN(0, "could not set LIMIT value in table!"); + if (Util_TableSetReal(handle_pi_g_bound ,pi_g_bound_speed, "SPEED") < 0) + CCTK_WARN(0, "could not set SPEED value in table!"); + + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_pi_g_bound, + "WaveHost::pi_g", "Radiation"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Radiation BC for WaveHost::pi_g!"); + + } + + if (CCTK_EQUALS(phi_bound, "radiative")) + { + /* select radiation boundary condition */ + static CCTK_INT handle_phi_bound = -1; + if (handle_phi_bound < 0) handle_phi_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_phi_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_phi_bound , phi_bound_limit, "LIMIT") < 0) + CCTK_WARN(0, "could not set LIMIT value in table!"); + if (Util_TableSetReal(handle_phi_bound ,phi_bound_speed, "SPEED") < 0) + CCTK_WARN(0, "could not set SPEED value in table!"); + + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_phi_bound, + "WaveHost::phi", "Radiation"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Radiation BC for WaveHost::phi!"); + + } + + if (CCTK_EQUALS(pi_bound, "radiative")) + { + /* select radiation boundary condition */ + static CCTK_INT handle_pi_bound = -1; + if (handle_pi_bound < 0) handle_pi_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_pi_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_pi_bound , pi_bound_limit, "LIMIT") < 0) + CCTK_WARN(0, "could not set LIMIT value in table!"); + if (Util_TableSetReal(handle_pi_bound ,pi_bound_speed, "SPEED") < 0) + CCTK_WARN(0, "could not set SPEED value in table!"); + + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_pi_bound, + "WaveHost::pi", "Radiation"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Radiation BC for WaveHost::pi!"); + + } + + if (CCTK_EQUALS(phi_g_bound, "scalar")) + { + /* select scalar boundary condition */ + static CCTK_INT handle_phi_g_bound = -1; + if (handle_phi_g_bound < 0) handle_phi_g_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_phi_g_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_phi_g_bound ,phi_g_bound_scalar, "SCALAR") < 0) + CCTK_WARN(0, "could not set SCALAR value in table!"); + + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_phi_g_bound, + "WaveHost::phi_g", "scalar"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Scalar BC for WaveHost::phi_g!"); + + } + + if (CCTK_EQUALS(pi_g_bound, "scalar")) + { + /* select scalar boundary condition */ + static CCTK_INT handle_pi_g_bound = -1; + if (handle_pi_g_bound < 0) handle_pi_g_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_pi_g_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_pi_g_bound ,pi_g_bound_scalar, "SCALAR") < 0) + CCTK_WARN(0, "could not set SCALAR value in table!"); + + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_pi_g_bound, + "WaveHost::pi_g", "scalar"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Scalar BC for WaveHost::pi_g!"); + + } + + if (CCTK_EQUALS(phi_bound, "scalar")) + { + /* select scalar boundary condition */ + static CCTK_INT handle_phi_bound = -1; + if (handle_phi_bound < 0) handle_phi_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_phi_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_phi_bound ,phi_bound_scalar, "SCALAR") < 0) + CCTK_WARN(0, "could not set SCALAR value in table!"); + + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_phi_bound, + "WaveHost::phi", "scalar"); + + if (ierr < 0) + CCTK_WARN(0, "Error in registering Scalar BC for WaveHost::phi!"); + + } + + if (CCTK_EQUALS(pi_bound, "scalar")) + { + /* select scalar boundary condition */ + static CCTK_INT handle_pi_bound = -1; + if (handle_pi_bound < 0) handle_pi_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_pi_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_pi_bound ,pi_bound_scalar, "SCALAR") < 0) + CCTK_WARN(0, "could not set SCALAR value in table!"); + + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_pi_bound, + "WaveHost::pi", "scalar"); + + if (ierr < 0) + CCTK_WARN(0, "Error in registering Scalar BC for WaveHost::pi!"); + + } + return; +} + + + +/* template for entries in parameter file: +#$bound$#WaveHost::phi_g_bound = "skip" +#$bound$#WaveHost::phi_g_bound_speed = 1.0 +#$bound$#WaveHost::phi_g_bound_limit = 0.0 +#$bound$#WaveHost::phi_g_bound_scalar = 0.0 + +#$bound$#WaveHost::pi_g_bound = "skip" +#$bound$#WaveHost::pi_g_bound_speed = 1.0 +#$bound$#WaveHost::pi_g_bound_limit = 0.0 +#$bound$#WaveHost::pi_g_bound_scalar = 0.0 + +#$bound$#WaveHost::phi_bound = "skip" +#$bound$#WaveHost::phi_bound_speed = 1.0 +#$bound$#WaveHost::phi_bound_limit = 0.0 +#$bound$#WaveHost::phi_bound_scalar = 0.0 + +#$bound$#WaveHost::pi_bound = "skip" +#$bound$#WaveHost::pi_bound_speed = 1.0 +#$bound$#WaveHost::pi_bound_limit = 0.0 +#$bound$#WaveHost::pi_bound_scalar = 0.0 + +*/ + diff --git a/Examples/WaveHost/src/CaKernel__copy_to_device.code b/Examples/WaveHost/src/CaKernel__copy_to_device.code new file mode 100644 index 00000000..24b4566c --- /dev/null +++ b/Examples/WaveHost/src/CaKernel__copy_to_device.code @@ -0,0 +1,77 @@ +#undef KRANC_DIFF_FUNCTIONS +#define KRANC_C +#include "Differencing.h" +#include "GenericFD.h" + +#undef KRANC_GFOFFSET3D +#define KRANC_GFOFFSET3D(u,i,j,k) I3D(u,i,j,k) + + +/* Define macros used in calculations */ +#define INITVALUE (42) +#define QAD(x) (SQR(SQR(x))) +#define INV(x) ((1.0) / (x)) +#define SQR(x) ((x) * (x)) +#define CUB(x) ((x) * (x) * (x)) + +CAKERNEL_copy_to_device_Begin + + /* Include user-supplied include files */ + + /* Initialise finite differencing variables */ + CCTK_REAL const dx = params.cagh_dx; + CCTK_REAL const dy = params.cagh_dy; + CCTK_REAL const dz = params.cagh_dz; + CCTK_REAL const dt = params.cagh_dt; + CCTK_REAL const t = params.cagh_time; + CCTK_REAL const dxi = INV(dx); + CCTK_REAL const dyi = INV(dy); + CCTK_REAL const dzi = INV(dz); + CCTK_REAL const khalf = 0.5; + CCTK_REAL const kthird = 1/3.0; + CCTK_REAL const ktwothird = 2.0/3.0; + CCTK_REAL const kfourthird = 4.0/3.0; + CCTK_REAL const keightthird = 8.0/3.0; + CCTK_REAL const hdxi = 0.5 * dxi; + CCTK_REAL const hdyi = 0.5 * dyi; + CCTK_REAL const hdzi = 0.5 * dzi; + + /* Initialize predefined quantities */ + CCTK_REAL const p1o2dx = 0.5*INV(dx); + CCTK_REAL const p1o2dy = 0.5*INV(dy); + CCTK_REAL const p1o2dz = 0.5*INV(dz); + CCTK_REAL const p1odx2 = INV(SQR(dx)); + CCTK_REAL const p1ody2 = INV(SQR(dy)); + CCTK_REAL const p1odz2 = INV(SQR(dz)); + + /* Assign local copies of arrays functions */ + + + + /* Calculate temporaries and arrays functions */ + + /* Copy local copies back to grid functions */ + CAKERNEL_copy_to_device_Computations_Begin + + /* Assign local copies of grid functions */ + + CCTK_REAL phiL = I3D(phi,0,0,0); + CCTK_REAL piL = I3D(pi,0,0,0); + + + /* Include user supplied include files */ + + /* Precompute derivatives */ + + /* Calculate temporaries and grid functions */ + phiL = phiL; + + piL = piL; + + /* Copy local copies back to grid functions */ + I3D(phi,0,0,0) = phiL; + I3D(pi,0,0,0) = piL; + + CAKERNEL_copy_to_device_Computations_End + +CAKERNEL_copy_to_device_End diff --git a/Examples/WaveHost/src/Differencing.h b/Examples/WaveHost/src/Differencing.h new file mode 100644 index 00000000..ef21b0bb --- /dev/null +++ b/Examples/WaveHost/src/Differencing.h @@ -0,0 +1,72 @@ +#ifndef KRANC_DIFF_FUNCTIONS +# define PDstandard2nd1(u) ((-KRANC_GFOFFSET3D(u,-1,0,0) + KRANC_GFOFFSET3D(u,1,0,0))*p1o2dx) +#else +# define PDstandard2nd1(u) (PDstandard2nd1_impl(u,p1o2dx,cdj,cdk)) +static CCTK_REAL PDstandard2nd1_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dx, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED; +static CCTK_REAL PDstandard2nd1_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dx, ptrdiff_t const cdj, ptrdiff_t const cdk) +{ + ptrdiff_t const cdi=sizeof(CCTK_REAL); + return (-KRANC_GFOFFSET3D(u,-1,0,0) + KRANC_GFOFFSET3D(u,1,0,0))*p1o2dx; +} +#endif + +#ifndef KRANC_DIFF_FUNCTIONS +# define PDstandard2nd2(u) ((-KRANC_GFOFFSET3D(u,0,-1,0) + KRANC_GFOFFSET3D(u,0,1,0))*p1o2dy) +#else +# define PDstandard2nd2(u) (PDstandard2nd2_impl(u,p1o2dy,cdj,cdk)) +static CCTK_REAL PDstandard2nd2_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dy, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED; +static CCTK_REAL PDstandard2nd2_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dy, ptrdiff_t const cdj, ptrdiff_t const cdk) +{ + ptrdiff_t const cdi=sizeof(CCTK_REAL); + return (-KRANC_GFOFFSET3D(u,0,-1,0) + KRANC_GFOFFSET3D(u,0,1,0))*p1o2dy; +} +#endif + +#ifndef KRANC_DIFF_FUNCTIONS +# define PDstandard2nd3(u) ((-KRANC_GFOFFSET3D(u,0,0,-1) + KRANC_GFOFFSET3D(u,0,0,1))*p1o2dz) +#else +# define PDstandard2nd3(u) (PDstandard2nd3_impl(u,p1o2dz,cdj,cdk)) +static CCTK_REAL PDstandard2nd3_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dz, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED; +static CCTK_REAL PDstandard2nd3_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dz, ptrdiff_t const cdj, ptrdiff_t const cdk) +{ + ptrdiff_t const cdi=sizeof(CCTK_REAL); + return (-KRANC_GFOFFSET3D(u,0,0,-1) + KRANC_GFOFFSET3D(u,0,0,1))*p1o2dz; +} +#endif + +#ifndef KRANC_DIFF_FUNCTIONS +# define PDstandard2nd11(u) ((-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,-1,0,0) + KRANC_GFOFFSET3D(u,1,0,0))*p1odx2) +#else +# define PDstandard2nd11(u) (PDstandard2nd11_impl(u,p1odx2,cdj,cdk)) +static CCTK_REAL PDstandard2nd11_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1odx2, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED; +static CCTK_REAL PDstandard2nd11_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1odx2, ptrdiff_t const cdj, ptrdiff_t const cdk) +{ + ptrdiff_t const cdi=sizeof(CCTK_REAL); + return (-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,-1,0,0) + KRANC_GFOFFSET3D(u,1,0,0))*p1odx2; +} +#endif + +#ifndef KRANC_DIFF_FUNCTIONS +# define PDstandard2nd22(u) ((-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,0,-1,0) + KRANC_GFOFFSET3D(u,0,1,0))*p1ody2) +#else +# define PDstandard2nd22(u) (PDstandard2nd22_impl(u,p1ody2,cdj,cdk)) +static CCTK_REAL PDstandard2nd22_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1ody2, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED; +static CCTK_REAL PDstandard2nd22_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1ody2, ptrdiff_t const cdj, ptrdiff_t const cdk) +{ + ptrdiff_t const cdi=sizeof(CCTK_REAL); + return (-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,0,-1,0) + KRANC_GFOFFSET3D(u,0,1,0))*p1ody2; +} +#endif + +#ifndef KRANC_DIFF_FUNCTIONS +# define PDstandard2nd33(u) ((-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,0,0,-1) + KRANC_GFOFFSET3D(u,0,0,1))*p1odz2) +#else +# define PDstandard2nd33(u) (PDstandard2nd33_impl(u,p1odz2,cdj,cdk)) +static CCTK_REAL PDstandard2nd33_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1odz2, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED; +static CCTK_REAL PDstandard2nd33_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1odz2, ptrdiff_t const cdj, ptrdiff_t const cdk) +{ + ptrdiff_t const cdi=sizeof(CCTK_REAL); + return (-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,0,0,-1) + KRANC_GFOFFSET3D(u,0,0,1))*p1odz2; +} +#endif + diff --git a/Examples/WaveHost/src/RegisterMoL.cc b/Examples/WaveHost/src/RegisterMoL.cc new file mode 100644 index 00000000..0e263abd --- /dev/null +++ b/Examples/WaveHost/src/RegisterMoL.cc @@ -0,0 +1,20 @@ +/* File produced by Kranc */ + +#include "cctk.h" +#include "cctk_Arguments.h" +#include "cctk_Parameters.h" + +extern "C" void WaveHost_RegisterVars(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + CCTK_INT ierr = 0; + + /* Register all the evolved grid functions with MoL */ + ierr += MoLRegisterEvolved(CCTK_VarIndex("WaveHost::phi"), CCTK_VarIndex("WaveHost::phirhs")); + ierr += MoLRegisterEvolved(CCTK_VarIndex("WaveHost::pi"), CCTK_VarIndex("WaveHost::pirhs")); + + /* Register all the evolved Array functions with MoL */ + return; +} diff --git a/Examples/WaveHost/src/RegisterSymmetries.cc b/Examples/WaveHost/src/RegisterSymmetries.cc new file mode 100644 index 00000000..7e44b144 --- /dev/null +++ b/Examples/WaveHost/src/RegisterSymmetries.cc @@ -0,0 +1,34 @@ +/* File produced by Kranc */ + +#include "cctk.h" +#include "cctk_Arguments.h" +#include "cctk_Parameters.h" +#include "Symmetry.h" + +extern "C" void WaveHost_RegisterSymmetries(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + + /* array holding symmetry definitions */ + CCTK_INT sym[3]; + + + /* Register symmetries of grid functions */ + sym[0] = 1; + sym[1] = 1; + sym[2] = 1; + SetCartSymVN(cctkGH, sym, "WaveHost::phi"); + + sym[0] = 1; + sym[1] = 1; + sym[2] = 1; + SetCartSymVN(cctkGH, sym, "WaveHost::pi"); + + sym[0] = 1; + sym[1] = 1; + sym[2] = 1; + SetCartSymVN(cctkGH, sym, "WaveHost::xCopy"); + +} diff --git a/Examples/WaveHost/src/Startup.cc b/Examples/WaveHost/src/Startup.cc new file mode 100644 index 00000000..a779f904 --- /dev/null +++ b/Examples/WaveHost/src/Startup.cc @@ -0,0 +1,10 @@ +/* File produced by Kranc */ + +#include "cctk.h" + +extern "C" int WaveHost_Startup(void) +{ + const char * banner = "WaveHost"; + CCTK_RegisterBanner(banner); + return 0; +} diff --git a/Examples/WaveHost/src/calc_bound_rhs.cc b/Examples/WaveHost/src/calc_bound_rhs.cc new file mode 100644 index 00000000..7098ac51 --- /dev/null +++ b/Examples/WaveHost/src/calc_bound_rhs.cc @@ -0,0 +1,147 @@ +/* File produced by Kranc */ + +#define KRANC_C + +#include +#include +#include +#include +#include +#include "cctk.h" +#include "cctk_Arguments.h" +#include "cctk_Parameters.h" +#include "GenericFD.h" +#include "Differencing.h" +#include "cctk_Loop.h" +#include "loopcontrol.h" + +/* Define macros used in calculations */ +#define INITVALUE (42) +#define QAD(x) (SQR(SQR(x))) +#define INV(x) ((1.0) / (x)) +#define SQR(x) ((x) * (x)) +#define CUB(x) ((x) * (x) * (x)) + +extern "C" void calc_bound_rhs_SelectBCs(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + CCTK_INT ierr = 0; + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "WaveHost::phi_grhs","flat"); + if (ierr < 0) + CCTK_WARN(1, "Failed to register flat BC for WaveHost::phi_grhs."); + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "WaveHost::pi_grhs","flat"); + if (ierr < 0) + CCTK_WARN(1, "Failed to register flat BC for WaveHost::pi_grhs."); + return; +} + +static void calc_bound_rhs_Body(cGH const * restrict const cctkGH, int const dir, int const face, CCTK_REAL const normal[3], CCTK_REAL const tangentA[3], CCTK_REAL const tangentB[3], int const imin[3], int const imax[3], int const n_subblock_gfs, CCTK_REAL * restrict const subblock_gfs[]) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + + /* Include user-supplied include files */ + + /* Initialise finite differencing variables */ + ptrdiff_t const di = 1; + ptrdiff_t const dj = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0); + ptrdiff_t const dk = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0); + ptrdiff_t const cdi = sizeof(CCTK_REAL) * di; + ptrdiff_t const cdj = sizeof(CCTK_REAL) * dj; + ptrdiff_t const cdk = sizeof(CCTK_REAL) * dk; + CCTK_REAL const dx = ToReal(CCTK_DELTA_SPACE(0)); + CCTK_REAL const dy = ToReal(CCTK_DELTA_SPACE(1)); + CCTK_REAL const dz = ToReal(CCTK_DELTA_SPACE(2)); + CCTK_REAL const dt = ToReal(CCTK_DELTA_TIME); + CCTK_REAL const t = ToReal(cctk_time); + CCTK_REAL const dxi = INV(dx); + CCTK_REAL const dyi = INV(dy); + CCTK_REAL const dzi = INV(dz); + CCTK_REAL const khalf = 0.5; + CCTK_REAL const kthird = 1/3.0; + CCTK_REAL const ktwothird = 2.0/3.0; + CCTK_REAL const kfourthird = 4.0/3.0; + CCTK_REAL const keightthird = 8.0/3.0; + CCTK_REAL const hdxi = 0.5 * dxi; + CCTK_REAL const hdyi = 0.5 * dyi; + CCTK_REAL const hdzi = 0.5 * dzi; + + /* Initialize predefined quantities */ + CCTK_REAL const p1o2dx = 0.5*INV(dx); + CCTK_REAL const p1o2dy = 0.5*INV(dy); + CCTK_REAL const p1o2dz = 0.5*INV(dz); + CCTK_REAL const p1odx2 = INV(SQR(dx)); + CCTK_REAL const p1ody2 = INV(SQR(dy)); + CCTK_REAL const p1odz2 = INV(SQR(dz)); + + /* Assign local copies of arrays functions */ + + + + /* Calculate temporaries and arrays functions */ + + /* Copy local copies back to grid functions */ + + /* Loop over the grid points */ + #pragma omp parallel + CCTK_LOOP3(calc_bound_rhs, + i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2], + cctk_lsh[0],cctk_lsh[1],cctk_lsh[2]) + { + ptrdiff_t const index = di*i + dj*j + dk*k; + + /* Assign local copies of grid functions */ + + CCTK_REAL xCopyL = xCopy[index]; + + + /* Include user supplied include files */ + + /* Precompute derivatives */ + + /* Calculate temporaries and grid functions */ + CCTK_REAL phirhsL = -200.*(xCopyL + t)*exp(-100.*SQR(xCopyL + t)); + + CCTK_REAL pirhsL = exp(-100.*SQR(xCopyL + t))*(-200. + + 80000.*xCopyL*t + 40000.*(SQR(xCopyL) + SQR(t))); + + /* Copy local copies back to grid functions */ + phirhs[index] = phirhsL; + pirhs[index] = pirhsL; + } + CCTK_ENDLOOP3(calc_bound_rhs); +} + +extern "C" void calc_bound_rhs(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Entering calc_bound_rhs_Body"); + } + + if (cctk_iteration % calc_bound_rhs_calc_every != calc_bound_rhs_calc_offset) + { + return; + } + + const char *const groups[] = { + "WaveHost::phi_grhs", + "WaveHost::pi_grhs", + "WaveHost::xCopy_g"}; + GenericFD_AssertGroupStorage(cctkGH, "calc_bound_rhs", 3, groups); + + + GenericFD_LoopOverBoundary(cctkGH, calc_bound_rhs_Body); + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Leaving calc_bound_rhs_Body"); + } +} diff --git a/Examples/WaveHost/src/calc_rhs.cc b/Examples/WaveHost/src/calc_rhs.cc new file mode 100644 index 00000000..ee114e7d --- /dev/null +++ b/Examples/WaveHost/src/calc_rhs.cc @@ -0,0 +1,153 @@ +/* File produced by Kranc */ + +#define KRANC_C + +#include +#include +#include +#include +#include +#include "cctk.h" +#include "cctk_Arguments.h" +#include "cctk_Parameters.h" +#include "GenericFD.h" +#include "Differencing.h" +#include "cctk_Loop.h" +#include "loopcontrol.h" + +/* Define macros used in calculations */ +#define INITVALUE (42) +#define QAD(x) (SQR(SQR(x))) +#define INV(x) ((1.0) / (x)) +#define SQR(x) ((x) * (x)) +#define CUB(x) ((x) * (x) * (x)) + +extern "C" void calc_rhs_SelectBCs(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + CCTK_INT ierr = 0; + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "WaveHost::phi_grhs","flat"); + if (ierr < 0) + CCTK_WARN(1, "Failed to register flat BC for WaveHost::phi_grhs."); + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "WaveHost::pi_grhs","flat"); + if (ierr < 0) + CCTK_WARN(1, "Failed to register flat BC for WaveHost::pi_grhs."); + return; +} + +static void calc_rhs_Body(cGH const * restrict const cctkGH, int const dir, int const face, CCTK_REAL const normal[3], CCTK_REAL const tangentA[3], CCTK_REAL const tangentB[3], int const imin[3], int const imax[3], int const n_subblock_gfs, CCTK_REAL * restrict const subblock_gfs[]) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + + /* Include user-supplied include files */ + + /* Initialise finite differencing variables */ + ptrdiff_t const di = 1; + ptrdiff_t const dj = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0); + ptrdiff_t const dk = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0); + ptrdiff_t const cdi = sizeof(CCTK_REAL) * di; + ptrdiff_t const cdj = sizeof(CCTK_REAL) * dj; + ptrdiff_t const cdk = sizeof(CCTK_REAL) * dk; + CCTK_REAL const dx = ToReal(CCTK_DELTA_SPACE(0)); + CCTK_REAL const dy = ToReal(CCTK_DELTA_SPACE(1)); + CCTK_REAL const dz = ToReal(CCTK_DELTA_SPACE(2)); + CCTK_REAL const dt = ToReal(CCTK_DELTA_TIME); + CCTK_REAL const t = ToReal(cctk_time); + CCTK_REAL const dxi = INV(dx); + CCTK_REAL const dyi = INV(dy); + CCTK_REAL const dzi = INV(dz); + CCTK_REAL const khalf = 0.5; + CCTK_REAL const kthird = 1/3.0; + CCTK_REAL const ktwothird = 2.0/3.0; + CCTK_REAL const kfourthird = 4.0/3.0; + CCTK_REAL const keightthird = 8.0/3.0; + CCTK_REAL const hdxi = 0.5 * dxi; + CCTK_REAL const hdyi = 0.5 * dyi; + CCTK_REAL const hdzi = 0.5 * dzi; + + /* Initialize predefined quantities */ + CCTK_REAL const p1o2dx = 0.5*INV(dx); + CCTK_REAL const p1o2dy = 0.5*INV(dy); + CCTK_REAL const p1o2dz = 0.5*INV(dz); + CCTK_REAL const p1odx2 = INV(SQR(dx)); + CCTK_REAL const p1ody2 = INV(SQR(dy)); + CCTK_REAL const p1odz2 = INV(SQR(dz)); + + /* Assign local copies of arrays functions */ + + + + /* Calculate temporaries and arrays functions */ + + /* Copy local copies back to grid functions */ + + /* Loop over the grid points */ + #pragma omp parallel + CCTK_LOOP3(calc_rhs, + i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2], + cctk_lsh[0],cctk_lsh[1],cctk_lsh[2]) + { + ptrdiff_t const index = di*i + dj*j + dk*k; + + /* Assign local copies of grid functions */ + + CCTK_REAL phiL = phi[index]; + CCTK_REAL piL = pi[index]; + + + /* Include user supplied include files */ + + /* Precompute derivatives */ + CCTK_REAL const PDstandard2nd11phi = PDstandard2nd11(&phi[index]); + CCTK_REAL const PDstandard2nd22phi = PDstandard2nd22(&phi[index]); + CCTK_REAL const PDstandard2nd33phi = PDstandard2nd33(&phi[index]); + + /* Calculate temporaries and grid functions */ + CCTK_REAL phirhsL = piL; + + CCTK_REAL pirhsL = PDstandard2nd11phi + PDstandard2nd22phi + + PDstandard2nd33phi; + + /* Copy local copies back to grid functions */ + phirhs[index] = phirhsL; + pirhs[index] = pirhsL; + } + CCTK_ENDLOOP3(calc_rhs); +} + +extern "C" void calc_rhs(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Entering calc_rhs_Body"); + } + + if (cctk_iteration % calc_rhs_calc_every != calc_rhs_calc_offset) + { + return; + } + + const char *const groups[] = { + "WaveHost::phi_g", + "WaveHost::phi_grhs", + "WaveHost::pi_g", + "WaveHost::pi_grhs"}; + GenericFD_AssertGroupStorage(cctkGH, "calc_rhs", 4, groups); + + GenericFD_EnsureStencilFits(cctkGH, "calc_rhs", 1, 1, 1); + + GenericFD_LoopOverInterior(cctkGH, calc_rhs_Body); + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Leaving calc_rhs_Body"); + } +} diff --git a/Examples/SimpleWaveCaKernel/src/initial_gaussian.cc b/Examples/WaveHost/src/initial_gaussian.cc similarity index 94% rename from Examples/SimpleWaveCaKernel/src/initial_gaussian.cc rename to Examples/WaveHost/src/initial_gaussian.cc index 55bd2283..28319362 100644 --- a/Examples/SimpleWaveCaKernel/src/initial_gaussian.cc +++ b/Examples/WaveHost/src/initial_gaussian.cc @@ -118,11 +118,15 @@ extern "C" void initial_gaussian(CCTK_ARGUMENTS) return; } - const char *groups[] = {"grid::coordinates","SimpleWaveCaKernel::phi_g","SimpleWaveCaKernel::pi_g","SimpleWaveCaKernel::xCopy_g"}; + const char *const groups[] = { + "grid::coordinates", + "WaveHost::phi_g", + "WaveHost::pi_g", + "WaveHost::xCopy_g"}; GenericFD_AssertGroupStorage(cctkGH, "initial_gaussian", 4, groups); - GenericFD_LoopOverEverything(cctkGH, &initial_gaussian_Body); + GenericFD_LoopOverEverything(cctkGH, initial_gaussian_Body); if (verbose > 1) { diff --git a/Examples/WaveHost/src/make.code.defn b/Examples/WaveHost/src/make.code.defn new file mode 100644 index 00000000..f30a3230 --- /dev/null +++ b/Examples/WaveHost/src/make.code.defn @@ -0,0 +1,3 @@ +# File produced by Kranc + +SRCS = Startup.cc RegisterSymmetries.cc RegisterMoL.cc initial_gaussian.cc calc_rhs.cc calc_bound_rhs.cc Boundaries.cc diff --git a/Examples/WaveHost/test b/Examples/WaveHost/test new file mode 120000 index 00000000..bbc45b0d --- /dev/null +++ b/Examples/WaveHost/test @@ -0,0 +1 @@ +../tests/WaveHost/test \ No newline at end of file diff --git a/Examples/simplewavecakernel_gaussian.par b/Examples/wavecakernel_gaussian.par similarity index 100% rename from Examples/simplewavecakernel_gaussian.par rename to Examples/wavecakernel_gaussian.par