From fa52744053f212b79087a40dac8c9df8d6d40158 Mon Sep 17 00:00:00 2001 From: Ian Hinder Date: Thu, 15 Mar 2012 17:23:09 +0100 Subject: 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.m | 72 ------ Examples/SimpleWaveCaKernel/cakernel.ccl | 63 ------ Examples/SimpleWaveCaKernel/configuration.ccl | 7 - Examples/SimpleWaveCaKernel/interface.ccl | 69 ------ Examples/SimpleWaveCaKernel/param.ccl | 216 ------------------ Examples/SimpleWaveCaKernel/schedule.ccl | 118 ---------- Examples/SimpleWaveCaKernel/src/Boundaries.cc | 249 --------------------- .../src/CaKernel__calc_bound_rhs.code | 74 ------ .../SimpleWaveCaKernel/src/CaKernel__calc_rhs.code | 81 ------- .../src/CaKernel__copy_to_device.code | 77 ------- Examples/SimpleWaveCaKernel/src/Differencing.h | 72 ------ Examples/SimpleWaveCaKernel/src/RegisterMoL.cc | 20 -- .../SimpleWaveCaKernel/src/RegisterSymmetries.cc | 34 --- Examples/SimpleWaveCaKernel/src/Startup.cc | 10 - .../SimpleWaveCaKernel/src/initial_gaussian.cc | 131 ----------- Examples/SimpleWaveCaKernel/src/make.code.defn | 3 - Examples/WaveCaKernel.m | 72 ++++++ Examples/WaveCaKernel/cakernel.ccl | 63 ++++++ Examples/WaveCaKernel/configuration.ccl | 7 + Examples/WaveCaKernel/interface.ccl | 69 ++++++ Examples/WaveCaKernel/param.ccl | 216 ++++++++++++++++++ Examples/WaveCaKernel/schedule.ccl | 141 ++++++++++++ Examples/WaveCaKernel/src/Boundaries.cc | 249 +++++++++++++++++++++ .../WaveCaKernel/src/CaKernel__calc_bound_rhs.code | 74 ++++++ Examples/WaveCaKernel/src/CaKernel__calc_rhs.code | 81 +++++++ .../WaveCaKernel/src/CaKernel__copy_to_device.code | 77 +++++++ Examples/WaveCaKernel/src/Differencing.h | 72 ++++++ Examples/WaveCaKernel/src/RegisterMoL.cc | 20 ++ Examples/WaveCaKernel/src/RegisterSymmetries.cc | 34 +++ Examples/WaveCaKernel/src/Startup.cc | 10 + Examples/WaveCaKernel/src/initial_gaussian.cc | 135 +++++++++++ Examples/WaveCaKernel/src/make.code.defn | 3 + 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 +++++++++++++++++++++ .../WaveHost/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 +++++++++++++ Examples/WaveHost/src/initial_gaussian.cc | 135 +++++++++++ Examples/WaveHost/src/make.code.defn | 3 + Examples/WaveHost/test | 1 + Examples/simplewavecakernel_gaussian.par | 140 ------------ Examples/wavecakernel_gaussian.par | 140 ++++++++++++ 50 files changed, 2760 insertions(+), 1436 deletions(-) delete mode 100644 Examples/SimpleWaveCaKernel.m delete mode 100644 Examples/SimpleWaveCaKernel/cakernel.ccl delete mode 100644 Examples/SimpleWaveCaKernel/configuration.ccl delete mode 100644 Examples/SimpleWaveCaKernel/interface.ccl delete mode 100644 Examples/SimpleWaveCaKernel/param.ccl delete mode 100644 Examples/SimpleWaveCaKernel/schedule.ccl delete mode 100644 Examples/SimpleWaveCaKernel/src/Boundaries.cc delete mode 100644 Examples/SimpleWaveCaKernel/src/CaKernel__calc_bound_rhs.code delete mode 100644 Examples/SimpleWaveCaKernel/src/CaKernel__calc_rhs.code delete mode 100644 Examples/SimpleWaveCaKernel/src/CaKernel__copy_to_device.code delete mode 100644 Examples/SimpleWaveCaKernel/src/Differencing.h delete mode 100644 Examples/SimpleWaveCaKernel/src/RegisterMoL.cc delete mode 100644 Examples/SimpleWaveCaKernel/src/RegisterSymmetries.cc delete mode 100644 Examples/SimpleWaveCaKernel/src/Startup.cc delete mode 100644 Examples/SimpleWaveCaKernel/src/initial_gaussian.cc delete mode 100644 Examples/SimpleWaveCaKernel/src/make.code.defn create mode 100644 Examples/WaveCaKernel.m create mode 100644 Examples/WaveCaKernel/cakernel.ccl create mode 100644 Examples/WaveCaKernel/configuration.ccl create mode 100644 Examples/WaveCaKernel/interface.ccl create mode 100644 Examples/WaveCaKernel/param.ccl create mode 100644 Examples/WaveCaKernel/schedule.ccl create mode 100644 Examples/WaveCaKernel/src/Boundaries.cc create mode 100644 Examples/WaveCaKernel/src/CaKernel__calc_bound_rhs.code create mode 100644 Examples/WaveCaKernel/src/CaKernel__calc_rhs.code create mode 100644 Examples/WaveCaKernel/src/CaKernel__copy_to_device.code create mode 100644 Examples/WaveCaKernel/src/Differencing.h create mode 100644 Examples/WaveCaKernel/src/RegisterMoL.cc create mode 100644 Examples/WaveCaKernel/src/RegisterSymmetries.cc create mode 100644 Examples/WaveCaKernel/src/Startup.cc create mode 100644 Examples/WaveCaKernel/src/initial_gaussian.cc create mode 100644 Examples/WaveCaKernel/src/make.code.defn 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 create mode 100644 Examples/WaveHost/src/initial_gaussian.cc create mode 100644 Examples/WaveHost/src/make.code.defn create mode 120000 Examples/WaveHost/test delete mode 100644 Examples/simplewavecakernel_gaussian.par create mode 100644 Examples/wavecakernel_gaussian.par diff --git a/Examples/SimpleWaveCaKernel.m b/Examples/SimpleWaveCaKernel.m deleted file mode 100644 index f4991ea..0000000 --- a/Examples/SimpleWaveCaKernel.m +++ /dev/null @@ -1,72 +0,0 @@ -<< "KrancThorn.m"; - -groups = {{"phi_g", {phi}}, {"pi_g", {pi}}, {"xCopy_g", {xCopy}}}; - -derivatives = -{ - PDstandard2nd[i_] -> StandardCenteredDifferenceOperator[1,1,i], - PDstandard2nd[i_, i_] -> StandardCenteredDifferenceOperator[2,1,i] -}; - -PD = PDstandard2nd; - -f[x_] := Exp[-(x/0.1)^2]; - -initialGaussianCalc = -{ - Name -> "initial_gaussian", - Schedule -> {"AT INITIAL"}, - ExecuteOn -> Host, - Equations -> - { - phi -> f[t+x], - pi -> D[f[t+x],t], - xCopy -> x - } -}; - -evolveCalc = -{ - Name -> "calc_rhs", - Schedule -> {"in MoL_CalcRHS"}, - Where -> Interior, - Equations -> - { - dot[phi] -> pi, - dot[pi] -> Euc[ui,uj] PD[phi,li,lj] - } -}; - -boundCalc = -{ - Name -> "calc_bound_rhs", - Schedule -> {"in MoL_RHSBoundaries"}, - Where -> Boundary, - Equations -> - { - dot[phi] -> D[f[t+xCopy],t], - dot[pi] -> D[f[t+xCopy],t,t] - } -}; - - -copyCalc = -{ - Name -> "copy_to_device", - Schedule -> {"at INITIAL after initial_gaussian"}, - Where -> Everywhere, - ExecuteOn -> Device, - Equations -> - { - phi -> phi, - pi -> pi - } -}; - -CreateKrancThornTT[groups, ".", - "SimpleWaveCaKernel", - Calculations -> {initialGaussianCalc, evolveCalc, copyCalc, boundCalc}, - PartialDerivatives -> derivatives, - UseCaKernel -> True, - EvolutionTimelevels -> 2, - DeclaredGroups -> {"phi_g","pi_g","xCopy_g"}]; diff --git a/Examples/SimpleWaveCaKernel/cakernel.ccl b/Examples/SimpleWaveCaKernel/cakernel.ccl deleted file mode 100644 index 75c394c..0000000 --- a/Examples/SimpleWaveCaKernel/cakernel.ccl +++ /dev/null @@ -1,63 +0,0 @@ -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 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 - { - phirhs - } - "phirhs" - - CCTK_CUDA_KERNEL_VARIABLE cached=no intent=out - { - pirhs - } - "pirhs" - - CCTK_CUDA_KERNEL_VARIABLE cached=no intent=in - { - xCopy - } - "xCopy" -} - diff --git a/Examples/SimpleWaveCaKernel/configuration.ccl b/Examples/SimpleWaveCaKernel/configuration.ccl deleted file mode 100644 index 46dbde5..0000000 --- a/Examples/SimpleWaveCaKernel/configuration.ccl +++ /dev/null @@ -1,7 +0,0 @@ -# File produced by Kranc - -REQUIRES GenericFD -OPTIONAL LoopControl -{ -} -REQUIRES CUDA \ No newline at end of file diff --git a/Examples/SimpleWaveCaKernel/interface.ccl b/Examples/SimpleWaveCaKernel/interface.ccl deleted file mode 100644 index 710d050..0000000 --- a/Examples/SimpleWaveCaKernel/interface.ccl +++ /dev/null @@ -1,69 +0,0 @@ -# File produced by Kranc - -implements: SimpleWaveCaKernel - -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/SimpleWaveCaKernel/param.ccl b/Examples/SimpleWaveCaKernel/param.ccl deleted file mode 100644 index 845e24e..0000000 --- a/Examples/SimpleWaveCaKernel/param.ccl +++ /dev/null @@ -1,216 +0,0 @@ -# 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 SimpleWaveCaKernel_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 -{ - 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 copy_to_device_calc_every "copy_to_device_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 copy_to_device_calc_offset "copy_to_device_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/SimpleWaveCaKernel/schedule.ccl b/Examples/SimpleWaveCaKernel/schedule.ccl deleted file mode 100644 index 5c777d8..0000000 --- a/Examples/SimpleWaveCaKernel/schedule.ccl +++ /dev/null @@ -1,118 +0,0 @@ -# File produced by Kranc - - -if (other_timelevels == 1) -{ - STORAGE: xCopy_g[1] -} - -if (timelevels == 1) -{ - STORAGE: phi_g[1] -} -if (timelevels == 2) -{ - STORAGE: phi_g[2] -} - -if (timelevels == 1) -{ - STORAGE: pi_g[1] -} -if (timelevels == 2) -{ - STORAGE: pi_g[2] -} - -if (rhs_timelevels == 1) -{ - STORAGE: phi_grhs[1] -} -if (rhs_timelevels == 2) -{ - STORAGE: phi_grhs[2] -} - -if (rhs_timelevels == 1) -{ - STORAGE: pi_grhs[1] -} -if (rhs_timelevels == 2) -{ - STORAGE: pi_grhs[2] -} - -schedule SimpleWaveCaKernel_Startup at STARTUP -{ - LANG: C - OPTIONS: meta -} "create banner" - -schedule SimpleWaveCaKernel_RegisterSymmetries in SymmetryRegister -{ - LANG: C - OPTIONS: meta -} "register symmetries" - -schedule initial_gaussian AT INITIAL -{ - LANG: C - READS: grid::coordinates - WRITES: SimpleWaveCaKernel::phi_g - WRITES: SimpleWaveCaKernel::pi_g - WRITES: SimpleWaveCaKernel::xCopy_g -} "initial_gaussian" - -schedule CAKERNEL_Launch_calc_rhs in MoL_CalcRHS -{ - LANG: C - TAGS: Device=1 - READS: SimpleWaveCaKernel::phi_g - READS: SimpleWaveCaKernel::pi_g - WRITES: SimpleWaveCaKernel::phi_grhs - WRITES: SimpleWaveCaKernel::pi_grhs -} "calc_rhs" - -schedule CAKERNEL_Launch_copy_to_device at INITIAL after initial_gaussian -{ - LANG: C - TAGS: Device=1 - READS: SimpleWaveCaKernel::phi_g - READS: SimpleWaveCaKernel::pi_g - WRITES: SimpleWaveCaKernel::phi_g - WRITES: SimpleWaveCaKernel::pi_g -} "copy_to_device" - -schedule CAKERNEL_Launch_calc_bound_rhs in MoL_RHSBoundaries -{ - LANG: C - TAGS: Device=1 - READS: SimpleWaveCaKernel::xCopy_g - WRITES: SimpleWaveCaKernel::phi_grhs - WRITES: SimpleWaveCaKernel::pi_grhs -} "calc_bound_rhs" - -schedule SimpleWaveCaKernel_SelectBoundConds in MoL_PostStep -{ - LANG: C - OPTIONS: level - SYNC: phi_g - SYNC: pi_g -} "select boundary conditions" - -schedule SimpleWaveCaKernel_CheckBoundaries at BASEGRID -{ - LANG: C - OPTIONS: meta -} "check boundaries treatment" - -schedule SimpleWaveCaKernel_RegisterVars in MoL_Register -{ - LANG: C - OPTIONS: meta -} "Register Variables for MoL" - -schedule group ApplyBCs as SimpleWaveCaKernel_ApplyBCs in MoL_PostStep after SimpleWaveCaKernel_SelectBoundConds -{ - # no language specified -} "Apply boundary conditions controlled by thorn Boundary" diff --git a/Examples/SimpleWaveCaKernel/src/Boundaries.cc b/Examples/SimpleWaveCaKernel/src/Boundaries.cc deleted file mode 100644 index 11561bf..0000000 --- a/Examples/SimpleWaveCaKernel/src/Boundaries.cc +++ /dev/null @@ -1,249 +0,0 @@ -/* 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 SimpleWaveCaKernel_CheckBoundaries(CCTK_ARGUMENTS) -{ - DECLARE_CCTK_ARGUMENTS; - DECLARE_CCTK_PARAMETERS; - - return; -} - -extern "C" void SimpleWaveCaKernel_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, - "SimpleWaveCaKernel::phi_g", phi_g_bound); - if (ierr < 0) - CCTK_WARN(0, "Failed to register phi_g_bound BC for SimpleWaveCaKernel::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, - "SimpleWaveCaKernel::pi_g", pi_g_bound); - if (ierr < 0) - CCTK_WARN(0, "Failed to register pi_g_bound BC for SimpleWaveCaKernel::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, - "SimpleWaveCaKernel::phi", phi_bound); - if (ierr < 0) - CCTK_WARN(0, "Failed to register phi_bound BC for SimpleWaveCaKernel::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, - "SimpleWaveCaKernel::pi", pi_bound); - if (ierr < 0) - CCTK_WARN(0, "Failed to register pi_bound BC for SimpleWaveCaKernel::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, - "SimpleWaveCaKernel::phi_g", "Radiation"); - - if (ierr < 0) - CCTK_WARN(0, "Failed to register Radiation BC for SimpleWaveCaKernel::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, - "SimpleWaveCaKernel::pi_g", "Radiation"); - - if (ierr < 0) - CCTK_WARN(0, "Failed to register Radiation BC for SimpleWaveCaKernel::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, - "SimpleWaveCaKernel::phi", "Radiation"); - - if (ierr < 0) - CCTK_WARN(0, "Failed to register Radiation BC for SimpleWaveCaKernel::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, - "SimpleWaveCaKernel::pi", "Radiation"); - - if (ierr < 0) - CCTK_WARN(0, "Failed to register Radiation BC for SimpleWaveCaKernel::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, - "SimpleWaveCaKernel::phi_g", "scalar"); - - if (ierr < 0) - CCTK_WARN(0, "Failed to register Scalar BC for SimpleWaveCaKernel::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, - "SimpleWaveCaKernel::pi_g", "scalar"); - - if (ierr < 0) - CCTK_WARN(0, "Failed to register Scalar BC for SimpleWaveCaKernel::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, - "SimpleWaveCaKernel::phi", "scalar"); - - if (ierr < 0) - CCTK_WARN(0, "Error in registering Scalar BC for SimpleWaveCaKernel::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, - "SimpleWaveCaKernel::pi", "scalar"); - - if (ierr < 0) - CCTK_WARN(0, "Error in registering Scalar BC for SimpleWaveCaKernel::pi!"); - - } - return; -} - - - -/* 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$#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$#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$#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 - -*/ - diff --git a/Examples/SimpleWaveCaKernel/src/CaKernel__calc_bound_rhs.code b/Examples/SimpleWaveCaKernel/src/CaKernel__calc_bound_rhs.code deleted file mode 100644 index 863ad4f..0000000 --- a/Examples/SimpleWaveCaKernel/src/CaKernel__calc_bound_rhs.code +++ /dev/null @@ -1,74 +0,0 @@ -#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/SimpleWaveCaKernel/src/CaKernel__calc_rhs.code b/Examples/SimpleWaveCaKernel/src/CaKernel__calc_rhs.code deleted file mode 100644 index d2d7fdb..0000000 --- a/Examples/SimpleWaveCaKernel/src/CaKernel__calc_rhs.code +++ /dev/null @@ -1,81 +0,0 @@ -#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/SimpleWaveCaKernel/src/CaKernel__copy_to_device.code b/Examples/SimpleWaveCaKernel/src/CaKernel__copy_to_device.code deleted file mode 100644 index 24b4566..0000000 --- a/Examples/SimpleWaveCaKernel/src/CaKernel__copy_to_device.code +++ /dev/null @@ -1,77 +0,0 @@ -#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/SimpleWaveCaKernel/src/Differencing.h b/Examples/SimpleWaveCaKernel/src/Differencing.h deleted file mode 100644 index ef21b0b..0000000 --- a/Examples/SimpleWaveCaKernel/src/Differencing.h +++ /dev/null @@ -1,72 +0,0 @@ -#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/SimpleWaveCaKernel/src/RegisterMoL.cc b/Examples/SimpleWaveCaKernel/src/RegisterMoL.cc deleted file mode 100644 index ac8c57b..0000000 --- a/Examples/SimpleWaveCaKernel/src/RegisterMoL.cc +++ /dev/null @@ -1,20 +0,0 @@ -/* File produced by Kranc */ - -#include "cctk.h" -#include "cctk_Arguments.h" -#include "cctk_Parameters.h" - -extern "C" void SimpleWaveCaKernel_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("SimpleWaveCaKernel::phi"), CCTK_VarIndex("SimpleWaveCaKernel::phirhs")); - ierr += MoLRegisterEvolved(CCTK_VarIndex("SimpleWaveCaKernel::pi"), CCTK_VarIndex("SimpleWaveCaKernel::pirhs")); - - /* Register all the evolved Array functions with MoL */ - return; -} diff --git a/Examples/SimpleWaveCaKernel/src/RegisterSymmetries.cc b/Examples/SimpleWaveCaKernel/src/RegisterSymmetries.cc deleted file mode 100644 index c099f5b..0000000 --- a/Examples/SimpleWaveCaKernel/src/RegisterSymmetries.cc +++ /dev/null @@ -1,34 +0,0 @@ -/* File produced by Kranc */ - -#include "cctk.h" -#include "cctk_Arguments.h" -#include "cctk_Parameters.h" -#include "Symmetry.h" - -extern "C" void SimpleWaveCaKernel_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, "SimpleWaveCaKernel::phi"); - - sym[0] = 1; - sym[1] = 1; - sym[2] = 1; - SetCartSymVN(cctkGH, sym, "SimpleWaveCaKernel::pi"); - - sym[0] = 1; - sym[1] = 1; - sym[2] = 1; - SetCartSymVN(cctkGH, sym, "SimpleWaveCaKernel::xCopy"); - -} diff --git a/Examples/SimpleWaveCaKernel/src/Startup.cc b/Examples/SimpleWaveCaKernel/src/Startup.cc deleted file mode 100644 index 2761c8b..0000000 --- a/Examples/SimpleWaveCaKernel/src/Startup.cc +++ /dev/null @@ -1,10 +0,0 @@ -/* File produced by Kranc */ - -#include "cctk.h" - -extern "C" int SimpleWaveCaKernel_Startup(void) -{ - const char * banner = "SimpleWaveCaKernel"; - CCTK_RegisterBanner(banner); - return 0; -} diff --git a/Examples/SimpleWaveCaKernel/src/initial_gaussian.cc b/Examples/SimpleWaveCaKernel/src/initial_gaussian.cc deleted file mode 100644 index 55bd228..0000000 --- a/Examples/SimpleWaveCaKernel/src/initial_gaussian.cc +++ /dev/null @@ -1,131 +0,0 @@ -/* 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 *groups[] = {"grid::coordinates","SimpleWaveCaKernel::phi_g","SimpleWaveCaKernel::pi_g","SimpleWaveCaKernel::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/SimpleWaveCaKernel/src/make.code.defn deleted file mode 100644 index c43661b..0000000 --- a/Examples/SimpleWaveCaKernel/src/make.code.defn +++ /dev/null @@ -1,3 +0,0 @@ -# File produced by Kranc - -SRCS = Startup.cc RegisterSymmetries.cc RegisterMoL.cc initial_gaussian.cc Boundaries.cc diff --git a/Examples/WaveCaKernel.m b/Examples/WaveCaKernel.m new file mode 100644 index 0000000..f4991ea --- /dev/null +++ b/Examples/WaveCaKernel.m @@ -0,0 +1,72 @@ +<< "KrancThorn.m"; + +groups = {{"phi_g", {phi}}, {"pi_g", {pi}}, {"xCopy_g", {xCopy}}}; + +derivatives = +{ + PDstandard2nd[i_] -> StandardCenteredDifferenceOperator[1,1,i], + PDstandard2nd[i_, i_] -> StandardCenteredDifferenceOperator[2,1,i] +}; + +PD = PDstandard2nd; + +f[x_] := Exp[-(x/0.1)^2]; + +initialGaussianCalc = +{ + Name -> "initial_gaussian", + Schedule -> {"AT INITIAL"}, + ExecuteOn -> Host, + Equations -> + { + phi -> f[t+x], + pi -> D[f[t+x],t], + xCopy -> x + } +}; + +evolveCalc = +{ + Name -> "calc_rhs", + Schedule -> {"in MoL_CalcRHS"}, + Where -> Interior, + Equations -> + { + dot[phi] -> pi, + dot[pi] -> Euc[ui,uj] PD[phi,li,lj] + } +}; + +boundCalc = +{ + Name -> "calc_bound_rhs", + Schedule -> {"in MoL_RHSBoundaries"}, + Where -> Boundary, + Equations -> + { + dot[phi] -> D[f[t+xCopy],t], + dot[pi] -> D[f[t+xCopy],t,t] + } +}; + + +copyCalc = +{ + Name -> "copy_to_device", + Schedule -> {"at INITIAL after initial_gaussian"}, + Where -> Everywhere, + ExecuteOn -> Device, + Equations -> + { + phi -> phi, + pi -> pi + } +}; + +CreateKrancThornTT[groups, ".", + "SimpleWaveCaKernel", + Calculations -> {initialGaussianCalc, evolveCalc, copyCalc, boundCalc}, + PartialDerivatives -> derivatives, + UseCaKernel -> True, + EvolutionTimelevels -> 2, + DeclaredGroups -> {"phi_g","pi_g","xCopy_g"}]; 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 +#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/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 diff --git a/Examples/WaveHost/configuration.ccl b/Examples/WaveHost/configuration.ccl new file mode 100644 index 0000000..0a66ec2 --- /dev/null +++ b/Examples/WaveHost/configuration.ccl @@ -0,0 +1,6 @@ +# File produced by Kranc + +REQUIRES GenericFD +OPTIONAL LoopControl +{ +} diff --git a/Examples/WaveHost/interface.ccl b/Examples/WaveHost/interface.ccl new file mode 100644 index 0000000..b08be61 --- /dev/null +++ b/Examples/WaveHost/interface.ccl @@ -0,0 +1,58 @@ +# File produced by Kranc + +implements: WaveHost + +inherits: Grid GenericFD Boundary + + + +USES INCLUDE: GenericFD.h +USES INCLUDE: Symmetry.h +USES INCLUDE: sbp_calc_coeffs.h +USES INCLUDE: Boundary.h +USES INCLUDE: loopcontrol.h + +CCTK_INT FUNCTION MoLRegisterEvolved(CCTK_INT IN EvolvedIndex, CCTK_INT IN RHSIndex) +USES FUNCTION MoLRegisterEvolved + +SUBROUTINE Diff_coeff(CCTK_POINTER_TO_CONST IN cctkGH, CCTK_INT IN dir, CCTK_INT IN nsize, CCTK_INT OUT ARRAY imin, CCTK_INT OUT ARRAY imax, CCTK_REAL OUT ARRAY q, CCTK_INT IN table_handle) +USES FUNCTION Diff_coeff + +CCTK_INT FUNCTION MultiPatch_GetMap(CCTK_POINTER_TO_CONST IN cctkGH) +USES FUNCTION MultiPatch_GetMap + +CCTK_INT FUNCTION Boundary_SelectGroupForBC(CCTK_POINTER_TO_CONST IN GH, CCTK_INT IN faces, CCTK_INT IN boundary_width, CCTK_INT IN table_handle, CCTK_STRING IN group_name, CCTK_STRING IN bc_name) +USES FUNCTION Boundary_SelectGroupForBC + +CCTK_INT FUNCTION Boundary_SelectVarForBC(CCTK_POINTER_TO_CONST IN GH, CCTK_INT IN faces, CCTK_INT IN boundary_width, CCTK_INT IN table_handle, CCTK_STRING IN var_name, CCTK_STRING IN bc_name) +USES FUNCTION Boundary_SelectVarForBC + +public: +CCTK_REAL xCopy_g type=GF timelevels=1 tags='' +{ + xCopy +} "xCopy_g" + +public: +CCTK_REAL phi_g type=GF timelevels=2 tags='' +{ + phi +} "phi_g" + +public: +CCTK_REAL pi_g type=GF timelevels=2 tags='' +{ + pi +} "pi_g" + +public: +CCTK_REAL phi_grhs type=GF timelevels=2 tags='' +{ + phirhs +} "phi_grhs" + +public: +CCTK_REAL pi_grhs type=GF timelevels=2 tags='' +{ + pirhs +} "pi_grhs" diff --git a/Examples/WaveHost/param.ccl b/Examples/WaveHost/param.ccl new file mode 100644 index 0000000..566e0c2 --- /dev/null +++ b/Examples/WaveHost/param.ccl @@ -0,0 +1,204 @@ +# File produced by Kranc + + +shares: GenericFD + + + +shares: MethodOfLines + +USES CCTK_INT MoL_Num_Evolved_Vars +USES CCTK_INT MoL_Num_ArrayEvolved_Vars + +restricted: +CCTK_INT verbose "verbose" STEERABLE=ALWAYS +{ + *:* :: "" +} 0 + +restricted: +CCTK_INT WaveHost_MaxNumEvolvedVars "Number of evolved variables used by this thorn" ACCUMULATOR-BASE=MethodofLines::MoL_Num_Evolved_Vars STEERABLE=RECOVER +{ + 2:2 :: "Number of evolved variables used by this thorn" +} 2 + +restricted: +CCTK_INT WaveHost_MaxNumArrayEvolvedVars "Number of Array evolved variables used by this thorn" ACCUMULATOR-BASE=MethodofLines::MoL_Num_ArrayEvolved_Vars STEERABLE=RECOVER +{ + 0:0 :: "Number of Array evolved variables used by this thorn" +} 0 + +restricted: +CCTK_INT timelevels "Number of active timelevels" STEERABLE=RECOVER +{ + 0:2 :: "" +} 2 + +restricted: +CCTK_INT rhs_timelevels "Number of active RHS timelevels" STEERABLE=RECOVER +{ + 0:2 :: "" +} 1 + +restricted: +CCTK_INT other_timelevels "Number of active timelevels for non-evolved grid functions" STEERABLE=RECOVER +{ + 0:2 :: "" +} 1 + +restricted: +CCTK_INT initial_gaussian_calc_every "initial_gaussian_calc_every" STEERABLE=ALWAYS +{ + *:* :: "" +} 1 + +restricted: +CCTK_INT calc_rhs_calc_every "calc_rhs_calc_every" STEERABLE=ALWAYS +{ + *:* :: "" +} 1 + +restricted: +CCTK_INT calc_bound_rhs_calc_every "calc_bound_rhs_calc_every" STEERABLE=ALWAYS +{ + *:* :: "" +} 1 + +restricted: +CCTK_INT initial_gaussian_calc_offset "initial_gaussian_calc_offset" STEERABLE=ALWAYS +{ + *:* :: "" +} 0 + +restricted: +CCTK_INT calc_rhs_calc_offset "calc_rhs_calc_offset" STEERABLE=ALWAYS +{ + *:* :: "" +} 0 + +restricted: +CCTK_INT calc_bound_rhs_calc_offset "calc_bound_rhs_calc_offset" STEERABLE=ALWAYS +{ + *:* :: "" +} 0 + +private: +KEYWORD phi_bound "Boundary condition to implement" STEERABLE=ALWAYS +{ + "flat" :: "Flat boundary condition" + "none" :: "No boundary condition" + "static" :: "Boundaries held fixed" + "radiative" :: "Radiation boundary condition" + "scalar" :: "Dirichlet boundary condition" + "newrad" :: "Improved radiative boundary condition" + "skip" :: "skip boundary condition code" +} "skip" + +private: +KEYWORD pi_bound "Boundary condition to implement" STEERABLE=ALWAYS +{ + "flat" :: "Flat boundary condition" + "none" :: "No boundary condition" + "static" :: "Boundaries held fixed" + "radiative" :: "Radiation boundary condition" + "scalar" :: "Dirichlet boundary condition" + "newrad" :: "Improved radiative boundary condition" + "skip" :: "skip boundary condition code" +} "skip" + +private: +KEYWORD phi_g_bound "Boundary condition to implement" STEERABLE=ALWAYS +{ + "flat" :: "Flat boundary condition" + "none" :: "No boundary condition" + "static" :: "Boundaries held fixed" + "radiative" :: "Radiation boundary condition" + "scalar" :: "Dirichlet boundary condition" + "newrad" :: "Improved radiative boundary condition" + "skip" :: "skip boundary condition code" +} "none" + +private: +KEYWORD pi_g_bound "Boundary condition to implement" STEERABLE=ALWAYS +{ + "flat" :: "Flat boundary condition" + "none" :: "No boundary condition" + "static" :: "Boundaries held fixed" + "radiative" :: "Radiation boundary condition" + "scalar" :: "Dirichlet boundary condition" + "newrad" :: "Improved radiative boundary condition" + "skip" :: "skip boundary condition code" +} "none" + +private: +CCTK_REAL phi_bound_speed "characteristic speed at boundary" STEERABLE=ALWAYS +{ + "0:*" :: "outgoing characteristic speed > 0" +} 1. + +private: +CCTK_REAL pi_bound_speed "characteristic speed at boundary" STEERABLE=ALWAYS +{ + "0:*" :: "outgoing characteristic speed > 0" +} 1. + +private: +CCTK_REAL phi_g_bound_speed "characteristic speed at boundary" STEERABLE=ALWAYS +{ + "0:*" :: "outgoing characteristic speed > 0" +} 1. + +private: +CCTK_REAL pi_g_bound_speed "characteristic speed at boundary" STEERABLE=ALWAYS +{ + "0:*" :: "outgoing characteristic speed > 0" +} 1. + +private: +CCTK_REAL phi_bound_limit "limit value for r -> infinity" STEERABLE=ALWAYS +{ + "*:*" :: "value of limit value is unrestricted" +} 0. + +private: +CCTK_REAL pi_bound_limit "limit value for r -> infinity" STEERABLE=ALWAYS +{ + "*:*" :: "value of limit value is unrestricted" +} 0. + +private: +CCTK_REAL phi_g_bound_limit "limit value for r -> infinity" STEERABLE=ALWAYS +{ + "*:*" :: "value of limit value is unrestricted" +} 0. + +private: +CCTK_REAL pi_g_bound_limit "limit value for r -> infinity" STEERABLE=ALWAYS +{ + "*:*" :: "value of limit value is unrestricted" +} 0. + +private: +CCTK_REAL phi_bound_scalar "Dirichlet boundary value" STEERABLE=ALWAYS +{ + "*:*" :: "unrestricted" +} 0. + +private: +CCTK_REAL pi_bound_scalar "Dirichlet boundary value" STEERABLE=ALWAYS +{ + "*:*" :: "unrestricted" +} 0. + +private: +CCTK_REAL phi_g_bound_scalar "Dirichlet boundary value" STEERABLE=ALWAYS +{ + "*:*" :: "unrestricted" +} 0. + +private: +CCTK_REAL pi_g_bound_scalar "Dirichlet boundary value" STEERABLE=ALWAYS +{ + "*:*" :: "unrestricted" +} 0. + diff --git a/Examples/WaveHost/schedule.ccl b/Examples/WaveHost/schedule.ccl new file mode 100644 index 0000000..b11880e --- /dev/null +++ b/Examples/WaveHost/schedule.ccl @@ -0,0 +1,127 @@ +# File produced by Kranc + + +if (other_timelevels == 1) +{ + STORAGE: xCopy_g[1] +} + +if (timelevels == 1) +{ + STORAGE: phi_g[1] +} +if (timelevels == 2) +{ + STORAGE: phi_g[2] +} + +if (timelevels == 1) +{ + STORAGE: pi_g[1] +} +if (timelevels == 2) +{ + STORAGE: pi_g[2] +} + +if (rhs_timelevels == 1) +{ + STORAGE: phi_grhs[1] +} +if (rhs_timelevels == 2) +{ + STORAGE: phi_grhs[2] +} + +if (rhs_timelevels == 1) +{ + STORAGE: pi_grhs[1] +} +if (rhs_timelevels == 2) +{ + STORAGE: pi_grhs[2] +} + +schedule WaveHost_Startup at STARTUP +{ + LANG: C + OPTIONS: meta +} "create banner" + +schedule WaveHost_RegisterSymmetries in SymmetryRegister +{ + LANG: C + OPTIONS: meta +} "register symmetries" + +schedule initial_gaussian AT INITIAL +{ + LANG: C + READS: grid::coordinates + WRITES: WaveHost::phi_g + WRITES: WaveHost::pi_g + WRITES: WaveHost::xCopy_g +} "initial_gaussian" + +schedule calc_rhs in MoL_CalcRHS +{ + LANG: C + READS: WaveHost::phi_g + READS: WaveHost::pi_g + WRITES: WaveHost::phi_grhs + WRITES: WaveHost::pi_grhs +} "calc_rhs" + +schedule calc_rhs at ANALYSIS +{ + LANG: C + SYNC: phi_grhs + SYNC: pi_grhs + READS: WaveHost::phi_g + READS: WaveHost::pi_g + WRITES: WaveHost::phi_grhs + WRITES: WaveHost::pi_grhs +} "calc_rhs" + +schedule calc_bound_rhs in MoL_RHSBoundaries +{ + LANG: C + READS: WaveHost::xCopy_g + WRITES: WaveHost::phi_grhs + WRITES: WaveHost::pi_grhs +} "calc_bound_rhs" + +schedule calc_bound_rhs at ANALYSIS +{ + LANG: C + SYNC: phi_grhs + SYNC: pi_grhs + READS: WaveHost::xCopy_g + WRITES: WaveHost::phi_grhs + WRITES: WaveHost::pi_grhs +} "calc_bound_rhs" + +schedule WaveHost_SelectBoundConds in MoL_PostStep +{ + LANG: C + OPTIONS: level + SYNC: phi_g + SYNC: pi_g +} "select boundary conditions" + +schedule WaveHost_CheckBoundaries at BASEGRID +{ + LANG: C + OPTIONS: meta +} "check boundaries treatment" + +schedule WaveHost_RegisterVars in MoL_Register +{ + LANG: C + OPTIONS: meta +} "Register Variables for MoL" + +schedule group ApplyBCs as WaveHost_ApplyBCs in MoL_PostStep after WaveHost_SelectBoundConds +{ + # no language specified +} "Apply boundary conditions controlled by thorn Boundary" diff --git a/Examples/WaveHost/src/Boundaries.cc b/Examples/WaveHost/src/Boundaries.cc new file mode 100644 index 0000000..8fe5d69 --- /dev/null +++ b/Examples/WaveHost/src/Boundaries.cc @@ -0,0 +1,249 @@ +/* File produced by Kranc */ + +#include "cctk.h" +#include "cctk_Arguments.h" +#include "cctk_Parameters.h" +#include "cctk_Faces.h" +#include "util_Table.h" +#include "Symmetry.h" + + +/* the boundary treatment is split into 3 steps: */ +/* 1. excision */ +/* 2. symmetries */ +/* 3. "other" boundary conditions, e.g. radiative */ + +/* to simplify scheduling and testing, the 3 steps */ +/* are currently applied in separate functions */ + + +extern "C" void WaveHost_CheckBoundaries(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + return; +} + +extern "C" void WaveHost_SelectBoundConds(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + CCTK_INT ierr = 0; + + if (CCTK_EQUALS(phi_g_bound, "none" ) || + CCTK_EQUALS(phi_g_bound, "static") || + CCTK_EQUALS(phi_g_bound, "flat" ) || + CCTK_EQUALS(phi_g_bound, "zero" ) ) + { + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, -1, + "WaveHost::phi_g", phi_g_bound); + if (ierr < 0) + CCTK_WARN(0, "Failed to register phi_g_bound BC for WaveHost::phi_g!"); + } + + if (CCTK_EQUALS(pi_g_bound, "none" ) || + CCTK_EQUALS(pi_g_bound, "static") || + CCTK_EQUALS(pi_g_bound, "flat" ) || + CCTK_EQUALS(pi_g_bound, "zero" ) ) + { + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, -1, + "WaveHost::pi_g", pi_g_bound); + if (ierr < 0) + CCTK_WARN(0, "Failed to register pi_g_bound BC for WaveHost::pi_g!"); + } + + if (CCTK_EQUALS(phi_bound, "none" ) || + CCTK_EQUALS(phi_bound, "static") || + CCTK_EQUALS(phi_bound, "flat" ) || + CCTK_EQUALS(phi_bound, "zero" ) ) + { + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, -1, + "WaveHost::phi", phi_bound); + if (ierr < 0) + CCTK_WARN(0, "Failed to register phi_bound BC for WaveHost::phi!"); + } + + if (CCTK_EQUALS(pi_bound, "none" ) || + CCTK_EQUALS(pi_bound, "static") || + CCTK_EQUALS(pi_bound, "flat" ) || + CCTK_EQUALS(pi_bound, "zero" ) ) + { + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, -1, + "WaveHost::pi", pi_bound); + if (ierr < 0) + CCTK_WARN(0, "Failed to register pi_bound BC for WaveHost::pi!"); + } + + if (CCTK_EQUALS(phi_g_bound, "radiative")) + { + /* select radiation boundary condition */ + static CCTK_INT handle_phi_g_bound = -1; + if (handle_phi_g_bound < 0) handle_phi_g_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_phi_g_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_phi_g_bound , phi_g_bound_limit, "LIMIT") < 0) + CCTK_WARN(0, "could not set LIMIT value in table!"); + if (Util_TableSetReal(handle_phi_g_bound ,phi_g_bound_speed, "SPEED") < 0) + CCTK_WARN(0, "could not set SPEED value in table!"); + + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_phi_g_bound, + "WaveHost::phi_g", "Radiation"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Radiation BC for WaveHost::phi_g!"); + + } + + if (CCTK_EQUALS(pi_g_bound, "radiative")) + { + /* select radiation boundary condition */ + static CCTK_INT handle_pi_g_bound = -1; + if (handle_pi_g_bound < 0) handle_pi_g_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_pi_g_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_pi_g_bound , pi_g_bound_limit, "LIMIT") < 0) + CCTK_WARN(0, "could not set LIMIT value in table!"); + if (Util_TableSetReal(handle_pi_g_bound ,pi_g_bound_speed, "SPEED") < 0) + CCTK_WARN(0, "could not set SPEED value in table!"); + + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_pi_g_bound, + "WaveHost::pi_g", "Radiation"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Radiation BC for WaveHost::pi_g!"); + + } + + if (CCTK_EQUALS(phi_bound, "radiative")) + { + /* select radiation boundary condition */ + static CCTK_INT handle_phi_bound = -1; + if (handle_phi_bound < 0) handle_phi_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_phi_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_phi_bound , phi_bound_limit, "LIMIT") < 0) + CCTK_WARN(0, "could not set LIMIT value in table!"); + if (Util_TableSetReal(handle_phi_bound ,phi_bound_speed, "SPEED") < 0) + CCTK_WARN(0, "could not set SPEED value in table!"); + + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_phi_bound, + "WaveHost::phi", "Radiation"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Radiation BC for WaveHost::phi!"); + + } + + if (CCTK_EQUALS(pi_bound, "radiative")) + { + /* select radiation boundary condition */ + static CCTK_INT handle_pi_bound = -1; + if (handle_pi_bound < 0) handle_pi_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_pi_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_pi_bound , pi_bound_limit, "LIMIT") < 0) + CCTK_WARN(0, "could not set LIMIT value in table!"); + if (Util_TableSetReal(handle_pi_bound ,pi_bound_speed, "SPEED") < 0) + CCTK_WARN(0, "could not set SPEED value in table!"); + + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_pi_bound, + "WaveHost::pi", "Radiation"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Radiation BC for WaveHost::pi!"); + + } + + if (CCTK_EQUALS(phi_g_bound, "scalar")) + { + /* select scalar boundary condition */ + static CCTK_INT handle_phi_g_bound = -1; + if (handle_phi_g_bound < 0) handle_phi_g_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_phi_g_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_phi_g_bound ,phi_g_bound_scalar, "SCALAR") < 0) + CCTK_WARN(0, "could not set SCALAR value in table!"); + + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_phi_g_bound, + "WaveHost::phi_g", "scalar"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Scalar BC for WaveHost::phi_g!"); + + } + + if (CCTK_EQUALS(pi_g_bound, "scalar")) + { + /* select scalar boundary condition */ + static CCTK_INT handle_pi_g_bound = -1; + if (handle_pi_g_bound < 0) handle_pi_g_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_pi_g_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_pi_g_bound ,pi_g_bound_scalar, "SCALAR") < 0) + CCTK_WARN(0, "could not set SCALAR value in table!"); + + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_pi_g_bound, + "WaveHost::pi_g", "scalar"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Scalar BC for WaveHost::pi_g!"); + + } + + if (CCTK_EQUALS(phi_bound, "scalar")) + { + /* select scalar boundary condition */ + static CCTK_INT handle_phi_bound = -1; + if (handle_phi_bound < 0) handle_phi_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_phi_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_phi_bound ,phi_bound_scalar, "SCALAR") < 0) + CCTK_WARN(0, "could not set SCALAR value in table!"); + + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_phi_bound, + "WaveHost::phi", "scalar"); + + if (ierr < 0) + CCTK_WARN(0, "Error in registering Scalar BC for WaveHost::phi!"); + + } + + if (CCTK_EQUALS(pi_bound, "scalar")) + { + /* select scalar boundary condition */ + static CCTK_INT handle_pi_bound = -1; + if (handle_pi_bound < 0) handle_pi_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_pi_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_pi_bound ,pi_bound_scalar, "SCALAR") < 0) + CCTK_WARN(0, "could not set SCALAR value in table!"); + + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_pi_bound, + "WaveHost::pi", "scalar"); + + if (ierr < 0) + CCTK_WARN(0, "Error in registering Scalar BC for WaveHost::pi!"); + + } + return; +} + + + +/* template for entries in parameter file: +#$bound$#WaveHost::phi_g_bound = "skip" +#$bound$#WaveHost::phi_g_bound_speed = 1.0 +#$bound$#WaveHost::phi_g_bound_limit = 0.0 +#$bound$#WaveHost::phi_g_bound_scalar = 0.0 + +#$bound$#WaveHost::pi_g_bound = "skip" +#$bound$#WaveHost::pi_g_bound_speed = 1.0 +#$bound$#WaveHost::pi_g_bound_limit = 0.0 +#$bound$#WaveHost::pi_g_bound_scalar = 0.0 + +#$bound$#WaveHost::phi_bound = "skip" +#$bound$#WaveHost::phi_bound_speed = 1.0 +#$bound$#WaveHost::phi_bound_limit = 0.0 +#$bound$#WaveHost::phi_bound_scalar = 0.0 + +#$bound$#WaveHost::pi_bound = "skip" +#$bound$#WaveHost::pi_bound_speed = 1.0 +#$bound$#WaveHost::pi_bound_limit = 0.0 +#$bound$#WaveHost::pi_bound_scalar = 0.0 + +*/ + diff --git a/Examples/WaveHost/src/CaKernel__copy_to_device.code b/Examples/WaveHost/src/CaKernel__copy_to_device.code new file mode 100644 index 0000000..24b4566 --- /dev/null +++ b/Examples/WaveHost/src/CaKernel__copy_to_device.code @@ -0,0 +1,77 @@ +#undef KRANC_DIFF_FUNCTIONS +#define KRANC_C +#include "Differencing.h" +#include "GenericFD.h" + +#undef KRANC_GFOFFSET3D +#define KRANC_GFOFFSET3D(u,i,j,k) I3D(u,i,j,k) + + +/* Define macros used in calculations */ +#define INITVALUE (42) +#define QAD(x) (SQR(SQR(x))) +#define INV(x) ((1.0) / (x)) +#define SQR(x) ((x) * (x)) +#define CUB(x) ((x) * (x) * (x)) + +CAKERNEL_copy_to_device_Begin + + /* Include user-supplied include files */ + + /* Initialise finite differencing variables */ + CCTK_REAL const dx = params.cagh_dx; + CCTK_REAL const dy = params.cagh_dy; + CCTK_REAL const dz = params.cagh_dz; + CCTK_REAL const dt = params.cagh_dt; + CCTK_REAL const t = params.cagh_time; + CCTK_REAL const dxi = INV(dx); + CCTK_REAL const dyi = INV(dy); + CCTK_REAL const dzi = INV(dz); + CCTK_REAL const khalf = 0.5; + CCTK_REAL const kthird = 1/3.0; + CCTK_REAL const ktwothird = 2.0/3.0; + CCTK_REAL const kfourthird = 4.0/3.0; + CCTK_REAL const keightthird = 8.0/3.0; + CCTK_REAL const hdxi = 0.5 * dxi; + CCTK_REAL const hdyi = 0.5 * dyi; + CCTK_REAL const hdzi = 0.5 * dzi; + + /* Initialize predefined quantities */ + CCTK_REAL const p1o2dx = 0.5*INV(dx); + CCTK_REAL const p1o2dy = 0.5*INV(dy); + CCTK_REAL const p1o2dz = 0.5*INV(dz); + CCTK_REAL const p1odx2 = INV(SQR(dx)); + CCTK_REAL const p1ody2 = INV(SQR(dy)); + CCTK_REAL const p1odz2 = INV(SQR(dz)); + + /* Assign local copies of arrays functions */ + + + + /* Calculate temporaries and arrays functions */ + + /* Copy local copies back to grid functions */ + CAKERNEL_copy_to_device_Computations_Begin + + /* Assign local copies of grid functions */ + + CCTK_REAL phiL = I3D(phi,0,0,0); + CCTK_REAL piL = I3D(pi,0,0,0); + + + /* Include user supplied include files */ + + /* Precompute derivatives */ + + /* Calculate temporaries and grid functions */ + phiL = phiL; + + piL = piL; + + /* Copy local copies back to grid functions */ + I3D(phi,0,0,0) = phiL; + I3D(pi,0,0,0) = piL; + + CAKERNEL_copy_to_device_Computations_End + +CAKERNEL_copy_to_device_End diff --git a/Examples/WaveHost/src/Differencing.h b/Examples/WaveHost/src/Differencing.h new file mode 100644 index 0000000..ef21b0b --- /dev/null +++ b/Examples/WaveHost/src/Differencing.h @@ -0,0 +1,72 @@ +#ifndef KRANC_DIFF_FUNCTIONS +# define PDstandard2nd1(u) ((-KRANC_GFOFFSET3D(u,-1,0,0) + KRANC_GFOFFSET3D(u,1,0,0))*p1o2dx) +#else +# define PDstandard2nd1(u) (PDstandard2nd1_impl(u,p1o2dx,cdj,cdk)) +static CCTK_REAL PDstandard2nd1_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dx, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED; +static CCTK_REAL PDstandard2nd1_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dx, ptrdiff_t const cdj, ptrdiff_t const cdk) +{ + ptrdiff_t const cdi=sizeof(CCTK_REAL); + return (-KRANC_GFOFFSET3D(u,-1,0,0) + KRANC_GFOFFSET3D(u,1,0,0))*p1o2dx; +} +#endif + +#ifndef KRANC_DIFF_FUNCTIONS +# define PDstandard2nd2(u) ((-KRANC_GFOFFSET3D(u,0,-1,0) + KRANC_GFOFFSET3D(u,0,1,0))*p1o2dy) +#else +# define PDstandard2nd2(u) (PDstandard2nd2_impl(u,p1o2dy,cdj,cdk)) +static CCTK_REAL PDstandard2nd2_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dy, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED; +static CCTK_REAL PDstandard2nd2_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dy, ptrdiff_t const cdj, ptrdiff_t const cdk) +{ + ptrdiff_t const cdi=sizeof(CCTK_REAL); + return (-KRANC_GFOFFSET3D(u,0,-1,0) + KRANC_GFOFFSET3D(u,0,1,0))*p1o2dy; +} +#endif + +#ifndef KRANC_DIFF_FUNCTIONS +# define PDstandard2nd3(u) ((-KRANC_GFOFFSET3D(u,0,0,-1) + KRANC_GFOFFSET3D(u,0,0,1))*p1o2dz) +#else +# define PDstandard2nd3(u) (PDstandard2nd3_impl(u,p1o2dz,cdj,cdk)) +static CCTK_REAL PDstandard2nd3_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dz, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED; +static CCTK_REAL PDstandard2nd3_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dz, ptrdiff_t const cdj, ptrdiff_t const cdk) +{ + ptrdiff_t const cdi=sizeof(CCTK_REAL); + return (-KRANC_GFOFFSET3D(u,0,0,-1) + KRANC_GFOFFSET3D(u,0,0,1))*p1o2dz; +} +#endif + +#ifndef KRANC_DIFF_FUNCTIONS +# define PDstandard2nd11(u) ((-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,-1,0,0) + KRANC_GFOFFSET3D(u,1,0,0))*p1odx2) +#else +# define PDstandard2nd11(u) (PDstandard2nd11_impl(u,p1odx2,cdj,cdk)) +static CCTK_REAL PDstandard2nd11_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1odx2, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED; +static CCTK_REAL PDstandard2nd11_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1odx2, ptrdiff_t const cdj, ptrdiff_t const cdk) +{ + ptrdiff_t const cdi=sizeof(CCTK_REAL); + return (-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,-1,0,0) + KRANC_GFOFFSET3D(u,1,0,0))*p1odx2; +} +#endif + +#ifndef KRANC_DIFF_FUNCTIONS +# define PDstandard2nd22(u) ((-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,0,-1,0) + KRANC_GFOFFSET3D(u,0,1,0))*p1ody2) +#else +# define PDstandard2nd22(u) (PDstandard2nd22_impl(u,p1ody2,cdj,cdk)) +static CCTK_REAL PDstandard2nd22_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1ody2, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED; +static CCTK_REAL PDstandard2nd22_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1ody2, ptrdiff_t const cdj, ptrdiff_t const cdk) +{ + ptrdiff_t const cdi=sizeof(CCTK_REAL); + return (-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,0,-1,0) + KRANC_GFOFFSET3D(u,0,1,0))*p1ody2; +} +#endif + +#ifndef KRANC_DIFF_FUNCTIONS +# define PDstandard2nd33(u) ((-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,0,0,-1) + KRANC_GFOFFSET3D(u,0,0,1))*p1odz2) +#else +# define PDstandard2nd33(u) (PDstandard2nd33_impl(u,p1odz2,cdj,cdk)) +static CCTK_REAL PDstandard2nd33_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1odz2, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED; +static CCTK_REAL PDstandard2nd33_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1odz2, ptrdiff_t const cdj, ptrdiff_t const cdk) +{ + ptrdiff_t const cdi=sizeof(CCTK_REAL); + return (-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,0,0,-1) + KRANC_GFOFFSET3D(u,0,0,1))*p1odz2; +} +#endif + diff --git a/Examples/WaveHost/src/RegisterMoL.cc b/Examples/WaveHost/src/RegisterMoL.cc new file mode 100644 index 0000000..0e263ab --- /dev/null +++ b/Examples/WaveHost/src/RegisterMoL.cc @@ -0,0 +1,20 @@ +/* File produced by Kranc */ + +#include "cctk.h" +#include "cctk_Arguments.h" +#include "cctk_Parameters.h" + +extern "C" void WaveHost_RegisterVars(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + CCTK_INT ierr = 0; + + /* Register all the evolved grid functions with MoL */ + ierr += MoLRegisterEvolved(CCTK_VarIndex("WaveHost::phi"), CCTK_VarIndex("WaveHost::phirhs")); + ierr += MoLRegisterEvolved(CCTK_VarIndex("WaveHost::pi"), CCTK_VarIndex("WaveHost::pirhs")); + + /* Register all the evolved Array functions with MoL */ + return; +} diff --git a/Examples/WaveHost/src/RegisterSymmetries.cc b/Examples/WaveHost/src/RegisterSymmetries.cc new file mode 100644 index 0000000..7e44b14 --- /dev/null +++ b/Examples/WaveHost/src/RegisterSymmetries.cc @@ -0,0 +1,34 @@ +/* File produced by Kranc */ + +#include "cctk.h" +#include "cctk_Arguments.h" +#include "cctk_Parameters.h" +#include "Symmetry.h" + +extern "C" void WaveHost_RegisterSymmetries(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + + /* array holding symmetry definitions */ + CCTK_INT sym[3]; + + + /* Register symmetries of grid functions */ + sym[0] = 1; + sym[1] = 1; + sym[2] = 1; + SetCartSymVN(cctkGH, sym, "WaveHost::phi"); + + sym[0] = 1; + sym[1] = 1; + sym[2] = 1; + SetCartSymVN(cctkGH, sym, "WaveHost::pi"); + + sym[0] = 1; + sym[1] = 1; + sym[2] = 1; + SetCartSymVN(cctkGH, sym, "WaveHost::xCopy"); + +} diff --git a/Examples/WaveHost/src/Startup.cc b/Examples/WaveHost/src/Startup.cc new file mode 100644 index 0000000..a779f90 --- /dev/null +++ b/Examples/WaveHost/src/Startup.cc @@ -0,0 +1,10 @@ +/* File produced by Kranc */ + +#include "cctk.h" + +extern "C" int WaveHost_Startup(void) +{ + const char * banner = "WaveHost"; + CCTK_RegisterBanner(banner); + return 0; +} diff --git a/Examples/WaveHost/src/calc_bound_rhs.cc b/Examples/WaveHost/src/calc_bound_rhs.cc new file mode 100644 index 0000000..7098ac5 --- /dev/null +++ b/Examples/WaveHost/src/calc_bound_rhs.cc @@ -0,0 +1,147 @@ +/* File produced by Kranc */ + +#define KRANC_C + +#include +#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 0000000..ee114e7 --- /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/WaveHost/src/initial_gaussian.cc b/Examples/WaveHost/src/initial_gaussian.cc new file mode 100644 index 0000000..2831936 --- /dev/null +++ b/Examples/WaveHost/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", + "WaveHost::phi_g", + "WaveHost::pi_g", + "WaveHost::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/WaveHost/src/make.code.defn b/Examples/WaveHost/src/make.code.defn new file mode 100644 index 0000000..f30a323 --- /dev/null +++ b/Examples/WaveHost/src/make.code.defn @@ -0,0 +1,3 @@ +# File produced by Kranc + +SRCS = Startup.cc RegisterSymmetries.cc RegisterMoL.cc initial_gaussian.cc calc_rhs.cc calc_bound_rhs.cc Boundaries.cc diff --git a/Examples/WaveHost/test b/Examples/WaveHost/test new file mode 120000 index 0000000..bbc45b0 --- /dev/null +++ b/Examples/WaveHost/test @@ -0,0 +1 @@ +../tests/WaveHost/test \ No newline at end of file diff --git a/Examples/simplewavecakernel_gaussian.par b/Examples/simplewavecakernel_gaussian.par deleted file mode 100644 index 9ffeeae..0000000 --- a/Examples/simplewavecakernel_gaussian.par +++ /dev/null @@ -1,140 +0,0 @@ - -ActiveThorns = " -CarpetIOHDF5 -Accelerator -Boundary -CaKernel -Carpet -CarpetIOASCII -CarpetIOBasic -CarpetIOScalar -CarpetLib -CarpetReduce -CartGrid3d -CoordBase -GenericFD -IOUtil -LoopControl -MoL -NanChecker -SimpleWaveCaKernel -SymBase -Time -TimerReport -" - -############################################################# -# Grid -############################################################# - -CoordBase::domainsize = minmax - -CoordBase::boundary_size_x_lower = 1 -CoordBase::boundary_size_y_lower = 1 -CoordBase::boundary_size_z_lower = 1 -CoordBase::boundary_shiftout_x_lower = 1 -CoordBase::boundary_shiftout_y_lower = 1 -CoordBase::boundary_shiftout_z_lower = 1 - -CoordBase::boundary_size_x_upper = 1 -CoordBase::boundary_size_y_upper = 1 -CoordBase::boundary_size_z_upper = 1 -CoordBase::boundary_shiftout_x_upper = 0 -CoordBase::boundary_shiftout_y_upper = 0 -CoordBase::boundary_shiftout_z_upper = 0 - -CartGrid3D::type = "coordbase" -CartGrid3D::domain = "full" -CartGrid3D::avoid_origin = "no" - -CoordBase::xmin = -1 -CoordBase::ymin = -1 -CoordBase::zmin = -1 - -CoordBase::xmax = 1 -CoordBase::ymax = 1 -CoordBase::zmax = 1 - -CoordBase::dx = 0.025 -CoordBase::dy = 0.025 -CoordBase::dz = 0.025 - -# CoordBase::dx = 0.05 -# CoordBase::dy = 0.05 -# CoordBase::dz = 0.05 - -############################################################# -# Carpet -############################################################# - -Carpet::ghost_size = 1 -Carpet::domain_from_coordbase = "yes" -Carpet::init_3_timelevels = "no" -# Carpet::verbose = yes -# Carpet::veryverbose = yes -Carpet::poison_value = 0 # 113 -Carpet::init_fill_timelevels = yes - -############################################################# -# Timers -############################################################# - -# Do not commit these as testsuite output as they will change on -# different machines. They are useful for performance monitoring in -# automated build and test systems. -#TimerReport::output_all_timers_readable = yes -#TimerReport::n_top_timers = 40 -TimerReport::output_schedule_timers = no - -############################################################# -# Time integration -############################################################# - -Cactus::terminate = "iteration" -Cactus::cctk_itlast = 32 - -Time::dtfac = 0.25 - -# MethodOfLines::ode_method = "generic" -# MethodOfLines::generic_type = "RK" -# MethodOfLines::MoL_Intermediate_Steps = 1 -# MethodOfLines::MoL_Num_Scratch_Levels = 0 - -MethodOfLines::ode_method = "RK3" -MethodOfLines::MoL_Intermediate_Steps = 3 -MethodOfLines::MoL_Num_Scratch_Levels = 1 -MethodOfLines::MoL_NaN_Check = "yes" - -############################################################# -# Boundary conditions -############################################################# - -LoopControl::printstats = no - -############################################################# -# Output -############################################################# - -IO::out_dir = $parfile -IO::out_fileinfo = "none" - -IOBasic::outInfo_every = 1 -IOBasic::outInfo_vars = "SimpleWaveCaKernel::phi" - -IOASCII::out1D_every = 1 -IOASCII::out1D_x = "yes" -IOASCII::out1D_y = "no" -IOASCII::out1D_z = "no" -IOASCII::out1D_d = "no" -IOASCII::out1D_vars = "SimpleWaveCaKernel::phi SimpleWaveCaKernel::pi" # SimpleWaveCaKernel::phirhs SimpleWaveCaKernel::pirhs -# IOHDF5::out_every = 1 -# IOHDF5::out_vars = "SimpleWaveCaKernel::phi SimpleWaveCaKernel::pi -# SimpleWaveCaKernel::phirhs SimpleWaveCaKernel::pirhs" - -# CaKernel::verbose = yes -# CaKernel::veryverbose = yes -# Accelerator::verbose = yes -# Accelerator::veryverbose = yes -# SimpleWaveCaKernel::verbose = 1 -# Carpet::veryverbose = yes - diff --git a/Examples/wavecakernel_gaussian.par b/Examples/wavecakernel_gaussian.par new file mode 100644 index 0000000..9ffeeae --- /dev/null +++ b/Examples/wavecakernel_gaussian.par @@ -0,0 +1,140 @@ + +ActiveThorns = " +CarpetIOHDF5 +Accelerator +Boundary +CaKernel +Carpet +CarpetIOASCII +CarpetIOBasic +CarpetIOScalar +CarpetLib +CarpetReduce +CartGrid3d +CoordBase +GenericFD +IOUtil +LoopControl +MoL +NanChecker +SimpleWaveCaKernel +SymBase +Time +TimerReport +" + +############################################################# +# Grid +############################################################# + +CoordBase::domainsize = minmax + +CoordBase::boundary_size_x_lower = 1 +CoordBase::boundary_size_y_lower = 1 +CoordBase::boundary_size_z_lower = 1 +CoordBase::boundary_shiftout_x_lower = 1 +CoordBase::boundary_shiftout_y_lower = 1 +CoordBase::boundary_shiftout_z_lower = 1 + +CoordBase::boundary_size_x_upper = 1 +CoordBase::boundary_size_y_upper = 1 +CoordBase::boundary_size_z_upper = 1 +CoordBase::boundary_shiftout_x_upper = 0 +CoordBase::boundary_shiftout_y_upper = 0 +CoordBase::boundary_shiftout_z_upper = 0 + +CartGrid3D::type = "coordbase" +CartGrid3D::domain = "full" +CartGrid3D::avoid_origin = "no" + +CoordBase::xmin = -1 +CoordBase::ymin = -1 +CoordBase::zmin = -1 + +CoordBase::xmax = 1 +CoordBase::ymax = 1 +CoordBase::zmax = 1 + +CoordBase::dx = 0.025 +CoordBase::dy = 0.025 +CoordBase::dz = 0.025 + +# CoordBase::dx = 0.05 +# CoordBase::dy = 0.05 +# CoordBase::dz = 0.05 + +############################################################# +# Carpet +############################################################# + +Carpet::ghost_size = 1 +Carpet::domain_from_coordbase = "yes" +Carpet::init_3_timelevels = "no" +# Carpet::verbose = yes +# Carpet::veryverbose = yes +Carpet::poison_value = 0 # 113 +Carpet::init_fill_timelevels = yes + +############################################################# +# Timers +############################################################# + +# Do not commit these as testsuite output as they will change on +# different machines. They are useful for performance monitoring in +# automated build and test systems. +#TimerReport::output_all_timers_readable = yes +#TimerReport::n_top_timers = 40 +TimerReport::output_schedule_timers = no + +############################################################# +# Time integration +############################################################# + +Cactus::terminate = "iteration" +Cactus::cctk_itlast = 32 + +Time::dtfac = 0.25 + +# MethodOfLines::ode_method = "generic" +# MethodOfLines::generic_type = "RK" +# MethodOfLines::MoL_Intermediate_Steps = 1 +# MethodOfLines::MoL_Num_Scratch_Levels = 0 + +MethodOfLines::ode_method = "RK3" +MethodOfLines::MoL_Intermediate_Steps = 3 +MethodOfLines::MoL_Num_Scratch_Levels = 1 +MethodOfLines::MoL_NaN_Check = "yes" + +############################################################# +# Boundary conditions +############################################################# + +LoopControl::printstats = no + +############################################################# +# Output +############################################################# + +IO::out_dir = $parfile +IO::out_fileinfo = "none" + +IOBasic::outInfo_every = 1 +IOBasic::outInfo_vars = "SimpleWaveCaKernel::phi" + +IOASCII::out1D_every = 1 +IOASCII::out1D_x = "yes" +IOASCII::out1D_y = "no" +IOASCII::out1D_z = "no" +IOASCII::out1D_d = "no" +IOASCII::out1D_vars = "SimpleWaveCaKernel::phi SimpleWaveCaKernel::pi" # SimpleWaveCaKernel::phirhs SimpleWaveCaKernel::pirhs +# IOHDF5::out_every = 1 +# IOHDF5::out_vars = "SimpleWaveCaKernel::phi SimpleWaveCaKernel::pi +# SimpleWaveCaKernel::phirhs SimpleWaveCaKernel::pirhs" + +# CaKernel::verbose = yes +# CaKernel::veryverbose = yes +# Accelerator::verbose = yes +# Accelerator::veryverbose = yes +# SimpleWaveCaKernel::verbose = 1 +# Carpet::veryverbose = yes + -- cgit v1.2.3