From f9c7beb75e0d2bfa91e8de341356c2cff0596d87 Mon Sep 17 00:00:00 2001 From: Erik Schnetter Date: Mon, 21 May 2012 12:59:50 -0400 Subject: Regenerate thorn ML_WaveToy and ML_WaveToy_CL --- ML_WaveToy_CL/configuration.ccl | 6 + ML_WaveToy_CL/interface.ccl | 60 +++++++ ML_WaveToy_CL/param.ccl | 259 ++++++++++++++++++++++++++++++ ML_WaveToy_CL/schedule.ccl | 154 ++++++++++++++++++ ML_WaveToy_CL/src/Boundaries.cc | 249 ++++++++++++++++++++++++++++ ML_WaveToy_CL/src/Differencing.h | 147 +++++++++++++++++ ML_WaveToy_CL/src/RegisterMoL.cc | 20 +++ ML_WaveToy_CL/src/RegisterSymmetries.cc | 34 ++++ ML_WaveToy_CL/src/Startup.cc | 10 ++ ML_WaveToy_CL/src/WT_CL_Dirichlet.cc | 166 +++++++++++++++++++ ML_WaveToy_CL/src/WT_CL_Energy.cc | 169 +++++++++++++++++++ ML_WaveToy_CL/src/WT_CL_EnergyBoundary.cc | 158 ++++++++++++++++++ ML_WaveToy_CL/src/WT_CL_Gaussian.cc | 155 ++++++++++++++++++ ML_WaveToy_CL/src/WT_CL_RHS.cc | 177 ++++++++++++++++++++ ML_WaveToy_CL/src/WT_CL_Standing.cc | 162 +++++++++++++++++++ ML_WaveToy_CL/src/make.code.defn | 3 + 16 files changed, 1929 insertions(+) create mode 100644 ML_WaveToy_CL/configuration.ccl create mode 100644 ML_WaveToy_CL/interface.ccl create mode 100644 ML_WaveToy_CL/param.ccl create mode 100644 ML_WaveToy_CL/schedule.ccl create mode 100644 ML_WaveToy_CL/src/Boundaries.cc create mode 100644 ML_WaveToy_CL/src/Differencing.h create mode 100644 ML_WaveToy_CL/src/RegisterMoL.cc create mode 100644 ML_WaveToy_CL/src/RegisterSymmetries.cc create mode 100644 ML_WaveToy_CL/src/Startup.cc create mode 100644 ML_WaveToy_CL/src/WT_CL_Dirichlet.cc create mode 100644 ML_WaveToy_CL/src/WT_CL_Energy.cc create mode 100644 ML_WaveToy_CL/src/WT_CL_EnergyBoundary.cc create mode 100644 ML_WaveToy_CL/src/WT_CL_Gaussian.cc create mode 100644 ML_WaveToy_CL/src/WT_CL_RHS.cc create mode 100644 ML_WaveToy_CL/src/WT_CL_Standing.cc create mode 100644 ML_WaveToy_CL/src/make.code.defn (limited to 'ML_WaveToy_CL') diff --git a/ML_WaveToy_CL/configuration.ccl b/ML_WaveToy_CL/configuration.ccl new file mode 100644 index 0000000..585220a --- /dev/null +++ b/ML_WaveToy_CL/configuration.ccl @@ -0,0 +1,6 @@ +# File produced by Kranc + +REQUIRES GenericFD +REQUIRES LoopControl +REQUIRES OpenCL OpenCLRunTime +REQUIRES Vectors diff --git a/ML_WaveToy_CL/interface.ccl b/ML_WaveToy_CL/interface.ccl new file mode 100644 index 0000000..2f3dfbd --- /dev/null +++ b/ML_WaveToy_CL/interface.ccl @@ -0,0 +1,60 @@ +# File produced by Kranc + +implements: ML_WaveToy_CL + +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 +USES INCLUDE: OpenCLRunTime.h +USES INCLUDE: vectors.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 WT_eps type=GF timelevels=1 tags='tensortypealias="Scalar" tensorweight=0' +{ + eps +} "WT_eps" + +public: +CCTK_REAL WT_rho type=GF timelevels=2 tags='tensortypealias="Scalar" tensorweight=0' +{ + rho +} "WT_rho" + +public: +CCTK_REAL WT_u type=GF timelevels=2 tags='tensortypealias="Scalar" tensorweight=0' +{ + u +} "WT_u" + +public: +CCTK_REAL WT_rhorhs type=GF timelevels=2 tags='tensortypealias="Scalar" tensorweight=0' +{ + rhorhs +} "WT_rhorhs" + +public: +CCTK_REAL WT_urhs type=GF timelevels=2 tags='tensortypealias="Scalar" tensorweight=0' +{ + urhs +} "WT_urhs" diff --git a/ML_WaveToy_CL/param.ccl b/ML_WaveToy_CL/param.ccl new file mode 100644 index 0000000..15d28f0 --- /dev/null +++ b/ML_WaveToy_CL/param.ccl @@ -0,0 +1,259 @@ +# 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_REAL amplitude "Amplitude of initial Gaussian" +{ + "*:*" :: "" +} 1 + +restricted: +CCTK_REAL width "Width of initial Gaussian" +{ + "*:*" :: "" +} 1 + +private: +KEYWORD initial_data "initial_data" +{ + "Gaussian" :: "Gaussian" + "Standing" :: "Standing" +} "Gaussian" + +restricted: +CCTK_INT ML_WaveToy_CL_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 ML_WaveToy_CL_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 WT_CL_Gaussian_calc_every "WT_CL_Gaussian_calc_every" STEERABLE=ALWAYS +{ + *:* :: "" +} 1 + +restricted: +CCTK_INT WT_CL_Standing_calc_every "WT_CL_Standing_calc_every" STEERABLE=ALWAYS +{ + *:* :: "" +} 1 + +restricted: +CCTK_INT WT_CL_RHS_calc_every "WT_CL_RHS_calc_every" STEERABLE=ALWAYS +{ + *:* :: "" +} 1 + +restricted: +CCTK_INT WT_CL_Dirichlet_calc_every "WT_CL_Dirichlet_calc_every" STEERABLE=ALWAYS +{ + *:* :: "" +} 1 + +restricted: +CCTK_INT WT_CL_Energy_calc_every "WT_CL_Energy_calc_every" STEERABLE=ALWAYS +{ + *:* :: "" +} 1 + +restricted: +CCTK_INT WT_CL_EnergyBoundary_calc_every "WT_CL_EnergyBoundary_calc_every" STEERABLE=ALWAYS +{ + *:* :: "" +} 1 + +restricted: +CCTK_INT WT_CL_Gaussian_calc_offset "WT_CL_Gaussian_calc_offset" STEERABLE=ALWAYS +{ + *:* :: "" +} 0 + +restricted: +CCTK_INT WT_CL_Standing_calc_offset "WT_CL_Standing_calc_offset" STEERABLE=ALWAYS +{ + *:* :: "" +} 0 + +restricted: +CCTK_INT WT_CL_RHS_calc_offset "WT_CL_RHS_calc_offset" STEERABLE=ALWAYS +{ + *:* :: "" +} 0 + +restricted: +CCTK_INT WT_CL_Dirichlet_calc_offset "WT_CL_Dirichlet_calc_offset" STEERABLE=ALWAYS +{ + *:* :: "" +} 0 + +restricted: +CCTK_INT WT_CL_Energy_calc_offset "WT_CL_Energy_calc_offset" STEERABLE=ALWAYS +{ + *:* :: "" +} 0 + +restricted: +CCTK_INT WT_CL_EnergyBoundary_calc_offset "WT_CL_EnergyBoundary_calc_offset" STEERABLE=ALWAYS +{ + *:* :: "" +} 0 + +private: +KEYWORD rho_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 u_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 WT_rho_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 WT_u_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 rho_bound_speed "characteristic speed at boundary" STEERABLE=ALWAYS +{ + "0:*" :: "outgoing characteristic speed > 0" +} 1. + +private: +CCTK_REAL u_bound_speed "characteristic speed at boundary" STEERABLE=ALWAYS +{ + "0:*" :: "outgoing characteristic speed > 0" +} 1. + +private: +CCTK_REAL WT_rho_bound_speed "characteristic speed at boundary" STEERABLE=ALWAYS +{ + "0:*" :: "outgoing characteristic speed > 0" +} 1. + +private: +CCTK_REAL WT_u_bound_speed "characteristic speed at boundary" STEERABLE=ALWAYS +{ + "0:*" :: "outgoing characteristic speed > 0" +} 1. + +private: +CCTK_REAL rho_bound_limit "limit value for r -> infinity" STEERABLE=ALWAYS +{ + "*:*" :: "value of limit value is unrestricted" +} 0. + +private: +CCTK_REAL u_bound_limit "limit value for r -> infinity" STEERABLE=ALWAYS +{ + "*:*" :: "value of limit value is unrestricted" +} 0. + +private: +CCTK_REAL WT_rho_bound_limit "limit value for r -> infinity" STEERABLE=ALWAYS +{ + "*:*" :: "value of limit value is unrestricted" +} 0. + +private: +CCTK_REAL WT_u_bound_limit "limit value for r -> infinity" STEERABLE=ALWAYS +{ + "*:*" :: "value of limit value is unrestricted" +} 0. + +private: +CCTK_REAL rho_bound_scalar "Dirichlet boundary value" STEERABLE=ALWAYS +{ + "*:*" :: "unrestricted" +} 0. + +private: +CCTK_REAL u_bound_scalar "Dirichlet boundary value" STEERABLE=ALWAYS +{ + "*:*" :: "unrestricted" +} 0. + +private: +CCTK_REAL WT_rho_bound_scalar "Dirichlet boundary value" STEERABLE=ALWAYS +{ + "*:*" :: "unrestricted" +} 0. + +private: +CCTK_REAL WT_u_bound_scalar "Dirichlet boundary value" STEERABLE=ALWAYS +{ + "*:*" :: "unrestricted" +} 0. + diff --git a/ML_WaveToy_CL/schedule.ccl b/ML_WaveToy_CL/schedule.ccl new file mode 100644 index 0000000..a6a6441 --- /dev/null +++ b/ML_WaveToy_CL/schedule.ccl @@ -0,0 +1,154 @@ +# File produced by Kranc + + +if (other_timelevels == 1) +{ + STORAGE: WT_eps[1] +} + +if (timelevels == 1) +{ + STORAGE: WT_rho[1] +} +if (timelevels == 2) +{ + STORAGE: WT_rho[2] +} + +if (timelevels == 1) +{ + STORAGE: WT_u[1] +} +if (timelevels == 2) +{ + STORAGE: WT_u[2] +} + +if (rhs_timelevels == 1) +{ + STORAGE: WT_rhorhs[1] +} +if (rhs_timelevels == 2) +{ + STORAGE: WT_rhorhs[2] +} + +if (rhs_timelevels == 1) +{ + STORAGE: WT_urhs[1] +} +if (rhs_timelevels == 2) +{ + STORAGE: WT_urhs[2] +} + +schedule ML_WaveToy_CL_Startup at STARTUP +{ + LANG: C + OPTIONS: meta +} "create banner" + +schedule ML_WaveToy_CL_RegisterSymmetries in SymmetryRegister +{ + LANG: C + OPTIONS: meta +} "register symmetries" + + +if (CCTK_EQUALS(initial_data, "Gaussian")) +{ + schedule WT_CL_Gaussian AT initial + { + LANG: C + TAGS: Device=1 + READS: grid::r + WRITES: ML_WaveToy_CL::rho + WRITES: ML_WaveToy_CL::u + } "WT_CL_Gaussian" +} + + +if (CCTK_EQUALS(initial_data, "Standing")) +{ + schedule WT_CL_Standing AT initial + { + LANG: C + TAGS: Device=1 + READS: grid::x + READS: grid::y + READS: grid::z + WRITES: ML_WaveToy_CL::rho + WRITES: ML_WaveToy_CL::u + } "WT_CL_Standing" +} + +schedule WT_CL_RHS IN MoL_CalcRHS +{ + LANG: C + TAGS: Device=1 + READS: ML_WaveToy_CL::rho + READS: ML_WaveToy_CL::u + WRITES: ML_WaveToy_CL::rhorhs + WRITES: ML_WaveToy_CL::urhs +} "WT_CL_RHS" + +schedule WT_CL_Dirichlet IN MoL_CalcRHS +{ + LANG: C + TAGS: Device=1 + WRITES: ML_WaveToy_CL::rhorhs + WRITES: ML_WaveToy_CL::urhs +} "WT_CL_Dirichlet" + +schedule WT_CL_Dirichlet AT analysis +{ + LANG: C + SYNC: WT_rhorhs + SYNC: WT_urhs + TAGS: Device=1 + WRITES: ML_WaveToy_CL::rhorhs + WRITES: ML_WaveToy_CL::urhs +} "WT_CL_Dirichlet" + +schedule WT_CL_Energy AT analysis +{ + LANG: C + SYNC: WT_eps + TAGS: Device=1 + READS: ML_WaveToy_CL::rho + READS: ML_WaveToy_CL::u + WRITES: ML_WaveToy_CL::eps +} "WT_CL_Energy" + +schedule WT_CL_EnergyBoundary AT analysis +{ + LANG: C + SYNC: WT_eps + TAGS: Device=1 + WRITES: ML_WaveToy_CL::eps +} "WT_CL_EnergyBoundary" + +schedule ML_WaveToy_CL_SelectBoundConds in MoL_PostStep +{ + LANG: C + OPTIONS: level + SYNC: WT_rho + SYNC: WT_u +} "select boundary conditions" + +schedule ML_WaveToy_CL_CheckBoundaries at BASEGRID +{ + LANG: C + OPTIONS: meta +} "check boundaries treatment" + +schedule ML_WaveToy_CL_RegisterVars in MoL_Register +{ + LANG: C + OPTIONS: meta +} "Register Variables for MoL" + +schedule group ApplyBCs as ML_WaveToy_CL_ApplyBCs in MoL_PostStep after ML_WaveToy_CL_SelectBoundConds +{ + # no language specified +} "Apply boundary conditions controlled by thorn Boundary" diff --git a/ML_WaveToy_CL/src/Boundaries.cc b/ML_WaveToy_CL/src/Boundaries.cc new file mode 100644 index 0000000..d9beb11 --- /dev/null +++ b/ML_WaveToy_CL/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 ML_WaveToy_CL_CheckBoundaries(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + return; +} + +extern "C" void ML_WaveToy_CL_SelectBoundConds(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + CCTK_INT ierr = 0; + + if (CCTK_EQUALS(WT_rho_bound, "none" ) || + CCTK_EQUALS(WT_rho_bound, "static") || + CCTK_EQUALS(WT_rho_bound, "flat" ) || + CCTK_EQUALS(WT_rho_bound, "zero" ) ) + { + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, -1, + "ML_WaveToy_CL::WT_rho", WT_rho_bound); + if (ierr < 0) + CCTK_WARN(0, "Failed to register WT_rho_bound BC for ML_WaveToy_CL::WT_rho!"); + } + + if (CCTK_EQUALS(WT_u_bound, "none" ) || + CCTK_EQUALS(WT_u_bound, "static") || + CCTK_EQUALS(WT_u_bound, "flat" ) || + CCTK_EQUALS(WT_u_bound, "zero" ) ) + { + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, -1, + "ML_WaveToy_CL::WT_u", WT_u_bound); + if (ierr < 0) + CCTK_WARN(0, "Failed to register WT_u_bound BC for ML_WaveToy_CL::WT_u!"); + } + + if (CCTK_EQUALS(rho_bound, "none" ) || + CCTK_EQUALS(rho_bound, "static") || + CCTK_EQUALS(rho_bound, "flat" ) || + CCTK_EQUALS(rho_bound, "zero" ) ) + { + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, -1, + "ML_WaveToy_CL::rho", rho_bound); + if (ierr < 0) + CCTK_WARN(0, "Failed to register rho_bound BC for ML_WaveToy_CL::rho!"); + } + + if (CCTK_EQUALS(u_bound, "none" ) || + CCTK_EQUALS(u_bound, "static") || + CCTK_EQUALS(u_bound, "flat" ) || + CCTK_EQUALS(u_bound, "zero" ) ) + { + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, -1, + "ML_WaveToy_CL::u", u_bound); + if (ierr < 0) + CCTK_WARN(0, "Failed to register u_bound BC for ML_WaveToy_CL::u!"); + } + + if (CCTK_EQUALS(WT_rho_bound, "radiative")) + { + /* select radiation boundary condition */ + static CCTK_INT handle_WT_rho_bound = -1; + if (handle_WT_rho_bound < 0) handle_WT_rho_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_WT_rho_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_WT_rho_bound , WT_rho_bound_limit, "LIMIT") < 0) + CCTK_WARN(0, "could not set LIMIT value in table!"); + if (Util_TableSetReal(handle_WT_rho_bound ,WT_rho_bound_speed, "SPEED") < 0) + CCTK_WARN(0, "could not set SPEED value in table!"); + + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_WT_rho_bound, + "ML_WaveToy_CL::WT_rho", "Radiation"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Radiation BC for ML_WaveToy_CL::WT_rho!"); + + } + + if (CCTK_EQUALS(WT_u_bound, "radiative")) + { + /* select radiation boundary condition */ + static CCTK_INT handle_WT_u_bound = -1; + if (handle_WT_u_bound < 0) handle_WT_u_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_WT_u_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_WT_u_bound , WT_u_bound_limit, "LIMIT") < 0) + CCTK_WARN(0, "could not set LIMIT value in table!"); + if (Util_TableSetReal(handle_WT_u_bound ,WT_u_bound_speed, "SPEED") < 0) + CCTK_WARN(0, "could not set SPEED value in table!"); + + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_WT_u_bound, + "ML_WaveToy_CL::WT_u", "Radiation"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Radiation BC for ML_WaveToy_CL::WT_u!"); + + } + + if (CCTK_EQUALS(rho_bound, "radiative")) + { + /* select radiation boundary condition */ + static CCTK_INT handle_rho_bound = -1; + if (handle_rho_bound < 0) handle_rho_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_rho_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_rho_bound , rho_bound_limit, "LIMIT") < 0) + CCTK_WARN(0, "could not set LIMIT value in table!"); + if (Util_TableSetReal(handle_rho_bound ,rho_bound_speed, "SPEED") < 0) + CCTK_WARN(0, "could not set SPEED value in table!"); + + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_rho_bound, + "ML_WaveToy_CL::rho", "Radiation"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Radiation BC for ML_WaveToy_CL::rho!"); + + } + + if (CCTK_EQUALS(u_bound, "radiative")) + { + /* select radiation boundary condition */ + static CCTK_INT handle_u_bound = -1; + if (handle_u_bound < 0) handle_u_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_u_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_u_bound , u_bound_limit, "LIMIT") < 0) + CCTK_WARN(0, "could not set LIMIT value in table!"); + if (Util_TableSetReal(handle_u_bound ,u_bound_speed, "SPEED") < 0) + CCTK_WARN(0, "could not set SPEED value in table!"); + + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_u_bound, + "ML_WaveToy_CL::u", "Radiation"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Radiation BC for ML_WaveToy_CL::u!"); + + } + + if (CCTK_EQUALS(WT_rho_bound, "scalar")) + { + /* select scalar boundary condition */ + static CCTK_INT handle_WT_rho_bound = -1; + if (handle_WT_rho_bound < 0) handle_WT_rho_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_WT_rho_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_WT_rho_bound ,WT_rho_bound_scalar, "SCALAR") < 0) + CCTK_WARN(0, "could not set SCALAR value in table!"); + + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_WT_rho_bound, + "ML_WaveToy_CL::WT_rho", "scalar"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Scalar BC for ML_WaveToy_CL::WT_rho!"); + + } + + if (CCTK_EQUALS(WT_u_bound, "scalar")) + { + /* select scalar boundary condition */ + static CCTK_INT handle_WT_u_bound = -1; + if (handle_WT_u_bound < 0) handle_WT_u_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_WT_u_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_WT_u_bound ,WT_u_bound_scalar, "SCALAR") < 0) + CCTK_WARN(0, "could not set SCALAR value in table!"); + + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_WT_u_bound, + "ML_WaveToy_CL::WT_u", "scalar"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Scalar BC for ML_WaveToy_CL::WT_u!"); + + } + + if (CCTK_EQUALS(rho_bound, "scalar")) + { + /* select scalar boundary condition */ + static CCTK_INT handle_rho_bound = -1; + if (handle_rho_bound < 0) handle_rho_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_rho_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_rho_bound ,rho_bound_scalar, "SCALAR") < 0) + CCTK_WARN(0, "could not set SCALAR value in table!"); + + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_rho_bound, + "ML_WaveToy_CL::rho", "scalar"); + + if (ierr < 0) + CCTK_WARN(0, "Error in registering Scalar BC for ML_WaveToy_CL::rho!"); + + } + + if (CCTK_EQUALS(u_bound, "scalar")) + { + /* select scalar boundary condition */ + static CCTK_INT handle_u_bound = -1; + if (handle_u_bound < 0) handle_u_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_u_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_u_bound ,u_bound_scalar, "SCALAR") < 0) + CCTK_WARN(0, "could not set SCALAR value in table!"); + + ierr = Boundary_SelectVarForBC(cctkGH, CCTK_ALL_FACES, 1, handle_u_bound, + "ML_WaveToy_CL::u", "scalar"); + + if (ierr < 0) + CCTK_WARN(0, "Error in registering Scalar BC for ML_WaveToy_CL::u!"); + + } + return; +} + + + +/* template for entries in parameter file: +#$bound$#ML_WaveToy_CL::WT_rho_bound = "skip" +#$bound$#ML_WaveToy_CL::WT_rho_bound_speed = 1.0 +#$bound$#ML_WaveToy_CL::WT_rho_bound_limit = 0.0 +#$bound$#ML_WaveToy_CL::WT_rho_bound_scalar = 0.0 + +#$bound$#ML_WaveToy_CL::WT_u_bound = "skip" +#$bound$#ML_WaveToy_CL::WT_u_bound_speed = 1.0 +#$bound$#ML_WaveToy_CL::WT_u_bound_limit = 0.0 +#$bound$#ML_WaveToy_CL::WT_u_bound_scalar = 0.0 + +#$bound$#ML_WaveToy_CL::rho_bound = "skip" +#$bound$#ML_WaveToy_CL::rho_bound_speed = 1.0 +#$bound$#ML_WaveToy_CL::rho_bound_limit = 0.0 +#$bound$#ML_WaveToy_CL::rho_bound_scalar = 0.0 + +#$bound$#ML_WaveToy_CL::u_bound = "skip" +#$bound$#ML_WaveToy_CL::u_bound_speed = 1.0 +#$bound$#ML_WaveToy_CL::u_bound_limit = 0.0 +#$bound$#ML_WaveToy_CL::u_bound_scalar = 0.0 + +*/ + diff --git a/ML_WaveToy_CL/src/Differencing.h b/ML_WaveToy_CL/src/Differencing.h new file mode 100644 index 0000000..85002c6 --- /dev/null +++ b/ML_WaveToy_CL/src/Differencing.h @@ -0,0 +1,147 @@ +static char const *const differencing = +"#ifndef KRANC_DIFF_FUNCTIONS\n" +"# define PDstandardNth1(u) (kmul(p1o12dx,kadd(KRANC_GFOFFSET3D(u,-2,0,0),kmadd(KRANC_GFOFFSET3D(u,-1,0,0),ToReal(-8),kmsub(KRANC_GFOFFSET3D(u,1,0,0),ToReal(8),KRANC_GFOFFSET3D(u,2,0,0))))))\n" +"#else\n" +"# define PDstandardNth1(u) (PDstandardNth1_impl(u,p1o12dx,cdj,cdk))\n" +"static CCTK_REAL_VEC PDstandardNth1_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o12dx, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n" +"static CCTK_REAL_VEC PDstandardNth1_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o12dx, ptrdiff_t const cdj, ptrdiff_t const cdk)\n" +"{\n" +" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n" +" return kmul(p1o12dx,kadd(KRANC_GFOFFSET3D(u,-2,0,0),kmadd(KRANC_GFOFFSET3D(u,-1,0,0),ToReal(-8),kmsub(KRANC_GFOFFSET3D(u,1,0,0),ToReal(8),KRANC_GFOFFSET3D(u,2,0,0)))));\n" +"}\n" +"#endif\n" +"\n" +"#ifndef KRANC_DIFF_FUNCTIONS\n" +"# define PDstandardNth2(u) (kmul(p1o12dy,kadd(KRANC_GFOFFSET3D(u,0,-2,0),kmadd(KRANC_GFOFFSET3D(u,0,-1,0),ToReal(-8),kmsub(KRANC_GFOFFSET3D(u,0,1,0),ToReal(8),KRANC_GFOFFSET3D(u,0,2,0))))))\n" +"#else\n" +"# define PDstandardNth2(u) (PDstandardNth2_impl(u,p1o12dy,cdj,cdk))\n" +"static CCTK_REAL_VEC PDstandardNth2_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o12dy, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n" +"static CCTK_REAL_VEC PDstandardNth2_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o12dy, ptrdiff_t const cdj, ptrdiff_t const cdk)\n" +"{\n" +" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n" +" return kmul(p1o12dy,kadd(KRANC_GFOFFSET3D(u,0,-2,0),kmadd(KRANC_GFOFFSET3D(u,0,-1,0),ToReal(-8),kmsub(KRANC_GFOFFSET3D(u,0,1,0),ToReal(8),KRANC_GFOFFSET3D(u,0,2,0)))));\n" +"}\n" +"#endif\n" +"\n" +"#ifndef KRANC_DIFF_FUNCTIONS\n" +"# define PDstandardNth3(u) (kmul(p1o12dz,kadd(KRANC_GFOFFSET3D(u,0,0,-2),kmadd(KRANC_GFOFFSET3D(u,0,0,-1),ToReal(-8),kmsub(KRANC_GFOFFSET3D(u,0,0,1),ToReal(8),KRANC_GFOFFSET3D(u,0,0,2))))))\n" +"#else\n" +"# define PDstandardNth3(u) (PDstandardNth3_impl(u,p1o12dz,cdj,cdk))\n" +"static CCTK_REAL_VEC PDstandardNth3_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o12dz, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n" +"static CCTK_REAL_VEC PDstandardNth3_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o12dz, ptrdiff_t const cdj, ptrdiff_t const cdk)\n" +"{\n" +" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n" +" return kmul(p1o12dz,kadd(KRANC_GFOFFSET3D(u,0,0,-2),kmadd(KRANC_GFOFFSET3D(u,0,0,-1),ToReal(-8),kmsub(KRANC_GFOFFSET3D(u,0,0,1),ToReal(8),KRANC_GFOFFSET3D(u,0,0,2)))));\n" +"}\n" +"#endif\n" +"\n" +"#ifndef KRANC_DIFF_FUNCTIONS\n" +"# define PDstandardNth11(u) (kmul(pm1o12dx2,kadd(KRANC_GFOFFSET3D(u,-2,0,0),kadd(KRANC_GFOFFSET3D(u,2,0,0),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,0,0),KRANC_GFOFFSET3D(u,1,0,0)),ToReal(-16),kmul(KRANC_GFOFFSET3D(u,0,0,0),ToReal(30)))))))\n" +"#else\n" +"# define PDstandardNth11(u) (PDstandardNth11_impl(u,pm1o12dx2,cdj,cdk))\n" +"static CCTK_REAL_VEC PDstandardNth11_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const pm1o12dx2, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n" +"static CCTK_REAL_VEC PDstandardNth11_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const pm1o12dx2, ptrdiff_t const cdj, ptrdiff_t const cdk)\n" +"{\n" +" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n" +" return kmul(pm1o12dx2,kadd(KRANC_GFOFFSET3D(u,-2,0,0),kadd(KRANC_GFOFFSET3D(u,2,0,0),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,0,0),KRANC_GFOFFSET3D(u,1,0,0)),ToReal(-16),kmul(KRANC_GFOFFSET3D(u,0,0,0),ToReal(30))))));\n" +"}\n" +"#endif\n" +"\n" +"#ifndef KRANC_DIFF_FUNCTIONS\n" +"# define PDstandardNth22(u) (kmul(pm1o12dy2,kadd(KRANC_GFOFFSET3D(u,0,-2,0),kadd(KRANC_GFOFFSET3D(u,0,2,0),kmadd(kadd(KRANC_GFOFFSET3D(u,0,-1,0),KRANC_GFOFFSET3D(u,0,1,0)),ToReal(-16),kmul(KRANC_GFOFFSET3D(u,0,0,0),ToReal(30)))))))\n" +"#else\n" +"# define PDstandardNth22(u) (PDstandardNth22_impl(u,pm1o12dy2,cdj,cdk))\n" +"static CCTK_REAL_VEC PDstandardNth22_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const pm1o12dy2, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n" +"static CCTK_REAL_VEC PDstandardNth22_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const pm1o12dy2, ptrdiff_t const cdj, ptrdiff_t const cdk)\n" +"{\n" +" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n" +" return kmul(pm1o12dy2,kadd(KRANC_GFOFFSET3D(u,0,-2,0),kadd(KRANC_GFOFFSET3D(u,0,2,0),kmadd(kadd(KRANC_GFOFFSET3D(u,0,-1,0),KRANC_GFOFFSET3D(u,0,1,0)),ToReal(-16),kmul(KRANC_GFOFFSET3D(u,0,0,0),ToReal(30))))));\n" +"}\n" +"#endif\n" +"\n" +"#ifndef KRANC_DIFF_FUNCTIONS\n" +"# define PDstandardNth33(u) (kmul(pm1o12dz2,kadd(KRANC_GFOFFSET3D(u,0,0,-2),kadd(KRANC_GFOFFSET3D(u,0,0,2),kmadd(kadd(KRANC_GFOFFSET3D(u,0,0,-1),KRANC_GFOFFSET3D(u,0,0,1)),ToReal(-16),kmul(KRANC_GFOFFSET3D(u,0,0,0),ToReal(30)))))))\n" +"#else\n" +"# define PDstandardNth33(u) (PDstandardNth33_impl(u,pm1o12dz2,cdj,cdk))\n" +"static CCTK_REAL_VEC PDstandardNth33_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const pm1o12dz2, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n" +"static CCTK_REAL_VEC PDstandardNth33_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const pm1o12dz2, ptrdiff_t const cdj, ptrdiff_t const cdk)\n" +"{\n" +" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n" +" return kmul(pm1o12dz2,kadd(KRANC_GFOFFSET3D(u,0,0,-2),kadd(KRANC_GFOFFSET3D(u,0,0,2),kmadd(kadd(KRANC_GFOFFSET3D(u,0,0,-1),KRANC_GFOFFSET3D(u,0,0,1)),ToReal(-16),kmul(KRANC_GFOFFSET3D(u,0,0,0),ToReal(30))))));\n" +"}\n" +"#endif\n" +"\n" +"#ifndef KRANC_DIFF_FUNCTIONS\n" +"# define PDstandardNth12(u) (kmul(p1o144dxdy,kadd(KRANC_GFOFFSET3D(u,-2,-2,0),kadd(KRANC_GFOFFSET3D(u,2,2,0),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,1,0),KRANC_GFOFFSET3D(u,1,-1,0)),ToReal(-64),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,-2,0),kadd(KRANC_GFOFFSET3D(u,1,2,0),kadd(KRANC_GFOFFSET3D(u,-2,-1,0),KRANC_GFOFFSET3D(u,2,1,0)))),ToReal(-8),ksub(ksub(kmadd(kadd(KRANC_GFOFFSET3D(u,-1,2,0),kadd(KRANC_GFOFFSET3D(u,1,-2,0),kadd(KRANC_GFOFFSET3D(u,-2,1,0),KRANC_GFOFFSET3D(u,2,-1,0)))),ToReal(8),kmul(kadd(KRANC_GFOFFSET3D(u,-1,-1,0),KRANC_GFOFFSET3D(u,1,1,0)),ToReal(64))),KRANC_GFOFFSET3D(u,2,-2,0)),KRANC_GFOFFSET3D(u,-2,2,0))))))))\n" +"#else\n" +"# define PDstandardNth12(u) (PDstandardNth12_impl(u,p1o144dxdy,cdj,cdk))\n" +"static CCTK_REAL_VEC PDstandardNth12_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o144dxdy, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n" +"static CCTK_REAL_VEC PDstandardNth12_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o144dxdy, ptrdiff_t const cdj, ptrdiff_t const cdk)\n" +"{\n" +" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n" +" return kmul(p1o144dxdy,kadd(KRANC_GFOFFSET3D(u,-2,-2,0),kadd(KRANC_GFOFFSET3D(u,2,2,0),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,1,0),KRANC_GFOFFSET3D(u,1,-1,0)),ToReal(-64),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,-2,0),kadd(KRANC_GFOFFSET3D(u,1,2,0),kadd(KRANC_GFOFFSET3D(u,-2,-1,0),KRANC_GFOFFSET3D(u,2,1,0)))),ToReal(-8),ksub(ksub(kmadd(kadd(KRANC_GFOFFSET3D(u,-1,2,0),kadd(KRANC_GFOFFSET3D(u,1,-2,0),kadd(KRANC_GFOFFSET3D(u,-2,1,0),KRANC_GFOFFSET3D(u,2,-1,0)))),ToReal(8),kmul(kadd(KRANC_GFOFFSET3D(u,-1,-1,0),KRANC_GFOFFSET3D(u,1,1,0)),ToReal(64))),KRANC_GFOFFSET3D(u,2,-2,0)),KRANC_GFOFFSET3D(u,-2,2,0)))))));\n" +"}\n" +"#endif\n" +"\n" +"#ifndef KRANC_DIFF_FUNCTIONS\n" +"# define PDstandardNth13(u) (kmul(p1o144dxdz,kadd(KRANC_GFOFFSET3D(u,-2,0,-2),kadd(KRANC_GFOFFSET3D(u,2,0,2),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,0,1),KRANC_GFOFFSET3D(u,1,0,-1)),ToReal(-64),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,0,-2),kadd(KRANC_GFOFFSET3D(u,1,0,2),kadd(KRANC_GFOFFSET3D(u,-2,0,-1),KRANC_GFOFFSET3D(u,2,0,1)))),ToReal(-8),ksub(ksub(kmadd(kadd(KRANC_GFOFFSET3D(u,-1,0,2),kadd(KRANC_GFOFFSET3D(u,1,0,-2),kadd(KRANC_GFOFFSET3D(u,-2,0,1),KRANC_GFOFFSET3D(u,2,0,-1)))),ToReal(8),kmul(kadd(KRANC_GFOFFSET3D(u,-1,0,-1),KRANC_GFOFFSET3D(u,1,0,1)),ToReal(64))),KRANC_GFOFFSET3D(u,2,0,-2)),KRANC_GFOFFSET3D(u,-2,0,2))))))))\n" +"#else\n" +"# define PDstandardNth13(u) (PDstandardNth13_impl(u,p1o144dxdz,cdj,cdk))\n" +"static CCTK_REAL_VEC PDstandardNth13_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o144dxdz, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n" +"static CCTK_REAL_VEC PDstandardNth13_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o144dxdz, ptrdiff_t const cdj, ptrdiff_t const cdk)\n" +"{\n" +" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n" +" return kmul(p1o144dxdz,kadd(KRANC_GFOFFSET3D(u,-2,0,-2),kadd(KRANC_GFOFFSET3D(u,2,0,2),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,0,1),KRANC_GFOFFSET3D(u,1,0,-1)),ToReal(-64),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,0,-2),kadd(KRANC_GFOFFSET3D(u,1,0,2),kadd(KRANC_GFOFFSET3D(u,-2,0,-1),KRANC_GFOFFSET3D(u,2,0,1)))),ToReal(-8),ksub(ksub(kmadd(kadd(KRANC_GFOFFSET3D(u,-1,0,2),kadd(KRANC_GFOFFSET3D(u,1,0,-2),kadd(KRANC_GFOFFSET3D(u,-2,0,1),KRANC_GFOFFSET3D(u,2,0,-1)))),ToReal(8),kmul(kadd(KRANC_GFOFFSET3D(u,-1,0,-1),KRANC_GFOFFSET3D(u,1,0,1)),ToReal(64))),KRANC_GFOFFSET3D(u,2,0,-2)),KRANC_GFOFFSET3D(u,-2,0,2)))))));\n" +"}\n" +"#endif\n" +"\n" +"#ifndef KRANC_DIFF_FUNCTIONS\n" +"# define PDstandardNth21(u) (kmul(p1o144dxdy,kadd(KRANC_GFOFFSET3D(u,-2,-2,0),kadd(KRANC_GFOFFSET3D(u,2,2,0),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,1,0),KRANC_GFOFFSET3D(u,1,-1,0)),ToReal(-64),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,-2,0),kadd(KRANC_GFOFFSET3D(u,1,2,0),kadd(KRANC_GFOFFSET3D(u,-2,-1,0),KRANC_GFOFFSET3D(u,2,1,0)))),ToReal(-8),ksub(ksub(kmadd(kadd(KRANC_GFOFFSET3D(u,-1,2,0),kadd(KRANC_GFOFFSET3D(u,1,-2,0),kadd(KRANC_GFOFFSET3D(u,-2,1,0),KRANC_GFOFFSET3D(u,2,-1,0)))),ToReal(8),kmul(kadd(KRANC_GFOFFSET3D(u,-1,-1,0),KRANC_GFOFFSET3D(u,1,1,0)),ToReal(64))),KRANC_GFOFFSET3D(u,2,-2,0)),KRANC_GFOFFSET3D(u,-2,2,0))))))))\n" +"#else\n" +"# define PDstandardNth21(u) (PDstandardNth21_impl(u,p1o144dxdy,cdj,cdk))\n" +"static CCTK_REAL_VEC PDstandardNth21_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o144dxdy, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n" +"static CCTK_REAL_VEC PDstandardNth21_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o144dxdy, ptrdiff_t const cdj, ptrdiff_t const cdk)\n" +"{\n" +" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n" +" return kmul(p1o144dxdy,kadd(KRANC_GFOFFSET3D(u,-2,-2,0),kadd(KRANC_GFOFFSET3D(u,2,2,0),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,1,0),KRANC_GFOFFSET3D(u,1,-1,0)),ToReal(-64),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,-2,0),kadd(KRANC_GFOFFSET3D(u,1,2,0),kadd(KRANC_GFOFFSET3D(u,-2,-1,0),KRANC_GFOFFSET3D(u,2,1,0)))),ToReal(-8),ksub(ksub(kmadd(kadd(KRANC_GFOFFSET3D(u,-1,2,0),kadd(KRANC_GFOFFSET3D(u,1,-2,0),kadd(KRANC_GFOFFSET3D(u,-2,1,0),KRANC_GFOFFSET3D(u,2,-1,0)))),ToReal(8),kmul(kadd(KRANC_GFOFFSET3D(u,-1,-1,0),KRANC_GFOFFSET3D(u,1,1,0)),ToReal(64))),KRANC_GFOFFSET3D(u,2,-2,0)),KRANC_GFOFFSET3D(u,-2,2,0)))))));\n" +"}\n" +"#endif\n" +"\n" +"#ifndef KRANC_DIFF_FUNCTIONS\n" +"# define PDstandardNth23(u) (kmul(p1o144dydz,kadd(KRANC_GFOFFSET3D(u,0,-2,-2),kadd(KRANC_GFOFFSET3D(u,0,2,2),kmadd(kadd(KRANC_GFOFFSET3D(u,0,-1,1),KRANC_GFOFFSET3D(u,0,1,-1)),ToReal(-64),kmadd(kadd(KRANC_GFOFFSET3D(u,0,-1,-2),kadd(KRANC_GFOFFSET3D(u,0,1,2),kadd(KRANC_GFOFFSET3D(u,0,-2,-1),KRANC_GFOFFSET3D(u,0,2,1)))),ToReal(-8),ksub(ksub(kmadd(kadd(KRANC_GFOFFSET3D(u,0,-1,2),kadd(KRANC_GFOFFSET3D(u,0,1,-2),kadd(KRANC_GFOFFSET3D(u,0,-2,1),KRANC_GFOFFSET3D(u,0,2,-1)))),ToReal(8),kmul(kadd(KRANC_GFOFFSET3D(u,0,-1,-1),KRANC_GFOFFSET3D(u,0,1,1)),ToReal(64))),KRANC_GFOFFSET3D(u,0,2,-2)),KRANC_GFOFFSET3D(u,0,-2,2))))))))\n" +"#else\n" +"# define PDstandardNth23(u) (PDstandardNth23_impl(u,p1o144dydz,cdj,cdk))\n" +"static CCTK_REAL_VEC PDstandardNth23_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o144dydz, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n" +"static CCTK_REAL_VEC PDstandardNth23_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o144dydz, ptrdiff_t const cdj, ptrdiff_t const cdk)\n" +"{\n" +" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n" +" return kmul(p1o144dydz,kadd(KRANC_GFOFFSET3D(u,0,-2,-2),kadd(KRANC_GFOFFSET3D(u,0,2,2),kmadd(kadd(KRANC_GFOFFSET3D(u,0,-1,1),KRANC_GFOFFSET3D(u,0,1,-1)),ToReal(-64),kmadd(kadd(KRANC_GFOFFSET3D(u,0,-1,-2),kadd(KRANC_GFOFFSET3D(u,0,1,2),kadd(KRANC_GFOFFSET3D(u,0,-2,-1),KRANC_GFOFFSET3D(u,0,2,1)))),ToReal(-8),ksub(ksub(kmadd(kadd(KRANC_GFOFFSET3D(u,0,-1,2),kadd(KRANC_GFOFFSET3D(u,0,1,-2),kadd(KRANC_GFOFFSET3D(u,0,-2,1),KRANC_GFOFFSET3D(u,0,2,-1)))),ToReal(8),kmul(kadd(KRANC_GFOFFSET3D(u,0,-1,-1),KRANC_GFOFFSET3D(u,0,1,1)),ToReal(64))),KRANC_GFOFFSET3D(u,0,2,-2)),KRANC_GFOFFSET3D(u,0,-2,2)))))));\n" +"}\n" +"#endif\n" +"\n" +"#ifndef KRANC_DIFF_FUNCTIONS\n" +"# define PDstandardNth31(u) (kmul(p1o144dxdz,kadd(KRANC_GFOFFSET3D(u,-2,0,-2),kadd(KRANC_GFOFFSET3D(u,2,0,2),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,0,1),KRANC_GFOFFSET3D(u,1,0,-1)),ToReal(-64),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,0,-2),kadd(KRANC_GFOFFSET3D(u,1,0,2),kadd(KRANC_GFOFFSET3D(u,-2,0,-1),KRANC_GFOFFSET3D(u,2,0,1)))),ToReal(-8),ksub(ksub(kmadd(kadd(KRANC_GFOFFSET3D(u,-1,0,2),kadd(KRANC_GFOFFSET3D(u,1,0,-2),kadd(KRANC_GFOFFSET3D(u,-2,0,1),KRANC_GFOFFSET3D(u,2,0,-1)))),ToReal(8),kmul(kadd(KRANC_GFOFFSET3D(u,-1,0,-1),KRANC_GFOFFSET3D(u,1,0,1)),ToReal(64))),KRANC_GFOFFSET3D(u,2,0,-2)),KRANC_GFOFFSET3D(u,-2,0,2))))))))\n" +"#else\n" +"# define PDstandardNth31(u) (PDstandardNth31_impl(u,p1o144dxdz,cdj,cdk))\n" +"static CCTK_REAL_VEC PDstandardNth31_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o144dxdz, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n" +"static CCTK_REAL_VEC PDstandardNth31_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o144dxdz, ptrdiff_t const cdj, ptrdiff_t const cdk)\n" +"{\n" +" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n" +" return kmul(p1o144dxdz,kadd(KRANC_GFOFFSET3D(u,-2,0,-2),kadd(KRANC_GFOFFSET3D(u,2,0,2),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,0,1),KRANC_GFOFFSET3D(u,1,0,-1)),ToReal(-64),kmadd(kadd(KRANC_GFOFFSET3D(u,-1,0,-2),kadd(KRANC_GFOFFSET3D(u,1,0,2),kadd(KRANC_GFOFFSET3D(u,-2,0,-1),KRANC_GFOFFSET3D(u,2,0,1)))),ToReal(-8),ksub(ksub(kmadd(kadd(KRANC_GFOFFSET3D(u,-1,0,2),kadd(KRANC_GFOFFSET3D(u,1,0,-2),kadd(KRANC_GFOFFSET3D(u,-2,0,1),KRANC_GFOFFSET3D(u,2,0,-1)))),ToReal(8),kmul(kadd(KRANC_GFOFFSET3D(u,-1,0,-1),KRANC_GFOFFSET3D(u,1,0,1)),ToReal(64))),KRANC_GFOFFSET3D(u,2,0,-2)),KRANC_GFOFFSET3D(u,-2,0,2)))))));\n" +"}\n" +"#endif\n" +"\n" +"#ifndef KRANC_DIFF_FUNCTIONS\n" +"# define PDstandardNth32(u) (kmul(p1o144dydz,kadd(KRANC_GFOFFSET3D(u,0,-2,-2),kadd(KRANC_GFOFFSET3D(u,0,2,2),kmadd(kadd(KRANC_GFOFFSET3D(u,0,-1,1),KRANC_GFOFFSET3D(u,0,1,-1)),ToReal(-64),kmadd(kadd(KRANC_GFOFFSET3D(u,0,-1,-2),kadd(KRANC_GFOFFSET3D(u,0,1,2),kadd(KRANC_GFOFFSET3D(u,0,-2,-1),KRANC_GFOFFSET3D(u,0,2,1)))),ToReal(-8),ksub(ksub(kmadd(kadd(KRANC_GFOFFSET3D(u,0,-1,2),kadd(KRANC_GFOFFSET3D(u,0,1,-2),kadd(KRANC_GFOFFSET3D(u,0,-2,1),KRANC_GFOFFSET3D(u,0,2,-1)))),ToReal(8),kmul(kadd(KRANC_GFOFFSET3D(u,0,-1,-1),KRANC_GFOFFSET3D(u,0,1,1)),ToReal(64))),KRANC_GFOFFSET3D(u,0,2,-2)),KRANC_GFOFFSET3D(u,0,-2,2))))))))\n" +"#else\n" +"# define PDstandardNth32(u) (PDstandardNth32_impl(u,p1o144dydz,cdj,cdk))\n" +"static CCTK_REAL_VEC PDstandardNth32_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o144dydz, ptrdiff_t const cdj, ptrdiff_t const cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n" +"static CCTK_REAL_VEC PDstandardNth32_impl(CCTK_REAL const* restrict const u, CCTK_REAL_VEC const p1o144dydz, ptrdiff_t const cdj, ptrdiff_t const cdk)\n" +"{\n" +" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n" +" return kmul(p1o144dydz,kadd(KRANC_GFOFFSET3D(u,0,-2,-2),kadd(KRANC_GFOFFSET3D(u,0,2,2),kmadd(kadd(KRANC_GFOFFSET3D(u,0,-1,1),KRANC_GFOFFSET3D(u,0,1,-1)),ToReal(-64),kmadd(kadd(KRANC_GFOFFSET3D(u,0,-1,-2),kadd(KRANC_GFOFFSET3D(u,0,1,2),kadd(KRANC_GFOFFSET3D(u,0,-2,-1),KRANC_GFOFFSET3D(u,0,2,1)))),ToReal(-8),ksub(ksub(kmadd(kadd(KRANC_GFOFFSET3D(u,0,-1,2),kadd(KRANC_GFOFFSET3D(u,0,1,-2),kadd(KRANC_GFOFFSET3D(u,0,-2,1),KRANC_GFOFFSET3D(u,0,2,-1)))),ToReal(8),kmul(kadd(KRANC_GFOFFSET3D(u,0,-1,-1),KRANC_GFOFFSET3D(u,0,1,1)),ToReal(64))),KRANC_GFOFFSET3D(u,0,2,-2)),KRANC_GFOFFSET3D(u,0,-2,2)))))));\n" +"}\n" +"#endif\n" +"\n" +"" +; diff --git a/ML_WaveToy_CL/src/RegisterMoL.cc b/ML_WaveToy_CL/src/RegisterMoL.cc new file mode 100644 index 0000000..b6409b5 --- /dev/null +++ b/ML_WaveToy_CL/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 ML_WaveToy_CL_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("ML_WaveToy_CL::rho"), CCTK_VarIndex("ML_WaveToy_CL::rhorhs")); + ierr += MoLRegisterEvolved(CCTK_VarIndex("ML_WaveToy_CL::u"), CCTK_VarIndex("ML_WaveToy_CL::urhs")); + + /* Register all the evolved Array functions with MoL */ + return; +} diff --git a/ML_WaveToy_CL/src/RegisterSymmetries.cc b/ML_WaveToy_CL/src/RegisterSymmetries.cc new file mode 100644 index 0000000..8b918af --- /dev/null +++ b/ML_WaveToy_CL/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 ML_WaveToy_CL_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, "ML_WaveToy_CL::rho"); + + sym[0] = 1; + sym[1] = 1; + sym[2] = 1; + SetCartSymVN(cctkGH, sym, "ML_WaveToy_CL::u"); + + sym[0] = 1; + sym[1] = 1; + sym[2] = 1; + SetCartSymVN(cctkGH, sym, "ML_WaveToy_CL::eps"); + +} diff --git a/ML_WaveToy_CL/src/Startup.cc b/ML_WaveToy_CL/src/Startup.cc new file mode 100644 index 0000000..257d808 --- /dev/null +++ b/ML_WaveToy_CL/src/Startup.cc @@ -0,0 +1,10 @@ +/* File produced by Kranc */ + +#include "cctk.h" + +extern "C" int ML_WaveToy_CL_Startup(void) +{ + const char * banner = "ML_WaveToy_CL"; + CCTK_RegisterBanner(banner); + return 0; +} diff --git a/ML_WaveToy_CL/src/WT_CL_Dirichlet.cc b/ML_WaveToy_CL/src/WT_CL_Dirichlet.cc new file mode 100644 index 0000000..91ff705 --- /dev/null +++ b/ML_WaveToy_CL/src/WT_CL_Dirichlet.cc @@ -0,0 +1,166 @@ +/* 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" +#include "OpenCLRunTime.h" +#include "vectors.h" + +/* Define macros used in calculations */ +#define INITVALUE (42) +#define QAD(x) (SQR(SQR(x))) +#define INV(x) (kdiv(ToReal(1.0),x)) +#define SQR(x) (kmul(x,x)) +#define CUB(x) (kmul(x,SQR(x))) + +extern "C" void WT_CL_Dirichlet_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 */, "ML_WaveToy_CL::WT_rhorhs","flat"); + if (ierr < 0) + CCTK_WARN(1, "Failed to register flat BC for ML_WaveToy_CL::WT_rhorhs."); + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "ML_WaveToy_CL::WT_urhs","flat"); + if (ierr < 0) + CCTK_WARN(1, "Failed to register flat BC for ML_WaveToy_CL::WT_urhs."); + return; +} + +static void WT_CL_Dirichlet_Body(cGH const * restrict const cctkGH, int const dir, int const face, CCTK_REAL const normal[3], CCTK_REAL const tangentA[3], CCTK_REAL const tangentB[3], int const imin[3], int const imax[3], int const n_subblock_gfs, CCTK_REAL * restrict const subblock_gfs[]) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + char const *const source = + "\n" + "/* Include user-supplied include files */\n" + "\n" + "/* Initialise finite differencing variables */\n" + "ptrdiff_t const di = 1;\n" + "ptrdiff_t const dj = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" + "ptrdiff_t const dk = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" + "ptrdiff_t const cdi = sizeof(CCTK_REAL) * di;\n" + "ptrdiff_t const cdj = sizeof(CCTK_REAL) * dj;\n" + "ptrdiff_t const cdk = sizeof(CCTK_REAL) * dk;\n" + "CCTK_REAL_VEC const dx = ToReal(CCTK_DELTA_SPACE(0));\n" + "CCTK_REAL_VEC const dy = ToReal(CCTK_DELTA_SPACE(1));\n" + "CCTK_REAL_VEC const dz = ToReal(CCTK_DELTA_SPACE(2));\n" + "CCTK_REAL_VEC const dt = ToReal(CCTK_DELTA_TIME);\n" + "CCTK_REAL_VEC const t = ToReal(cctk_time);\n" + "CCTK_REAL_VEC const dxi = INV(dx);\n" + "CCTK_REAL_VEC const dyi = INV(dy);\n" + "CCTK_REAL_VEC const dzi = INV(dz);\n" + "CCTK_REAL_VEC const khalf = ToReal(0.5);\n" + "CCTK_REAL_VEC const kthird = ToReal(1.0/3.0);\n" + "CCTK_REAL_VEC const ktwothird = ToReal(2.0/3.0);\n" + "CCTK_REAL_VEC const kfourthird = ToReal(4.0/3.0);\n" + "CCTK_REAL_VEC const keightthird = ToReal(8.0/3.0);\n" + "CCTK_REAL_VEC const hdxi = kmul(ToReal(0.5), dxi);\n" + "CCTK_REAL_VEC const hdyi = kmul(ToReal(0.5), dyi);\n" + "CCTK_REAL_VEC const hdzi = kmul(ToReal(0.5), dzi);\n" + "\n" + "/* Initialize predefined quantities */\n" + "CCTK_REAL_VEC const p1o12dx = kmul(INV(dx),ToReal(0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const p1o12dy = kmul(INV(dy),ToReal(0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const p1o12dz = kmul(INV(dz),ToReal(0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const p1o144dxdy = kmul(INV(kmul(dx,dy)),ToReal(0.00694444444444444444444444444444));\n" + "CCTK_REAL_VEC const p1o144dxdz = kmul(INV(kmul(dx,dz)),ToReal(0.00694444444444444444444444444444));\n" + "CCTK_REAL_VEC const p1o144dydz = kmul(INV(kmul(dy,dz)),ToReal(0.00694444444444444444444444444444));\n" + "CCTK_REAL_VEC const pm1o12dx2 = kmul(INV(SQR(dx)),ToReal(-0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const pm1o12dy2 = kmul(INV(SQR(dy)),ToReal(-0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const pm1o12dz2 = kmul(INV(SQR(dz)),ToReal(-0.0833333333333333333333333333333));\n" + "\n" + "/* Assign local copies of arrays functions */\n" + "\n" + "\n" + "\n" + "/* Calculate temporaries and arrays functions */\n" + "\n" + "/* Copy local copies back to grid functions */\n" + "\n" + "/* Loop over the grid points */\n" + "#pragma omp parallel\n" + "LC_LOOP3VEC(WT_CL_Dirichlet,\n" + " i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2],\n" + " cctk_lsh[0],cctk_lsh[1],cctk_lsh[2],\n" + " CCTK_REAL_VEC_SIZE)\n" + "{\n" + " ptrdiff_t const index = di*i + dj*j + dk*k;\n" + " \n" + " /* Assign local copies of grid functions */\n" + " \n" + " \n" + " \n" + " /* Include user supplied include files */\n" + " \n" + " /* Precompute derivatives */\n" + " \n" + " /* Calculate temporaries and grid functions */\n" + " CCTK_REAL_VEC urhsL = ToReal(0);\n" + " \n" + " CCTK_REAL_VEC rhorhsL = ToReal(0);\n" + " \n" + " /* Copy local copies back to grid functions */\n" + " vec_store_partial_prepare(i,lc_imin,lc_imax);\n" + " vec_store_nta_partial(rhorhs[index],rhorhsL);\n" + " vec_store_nta_partial(urhs[index],urhsL);\n" + "}\n" + "LC_ENDLOOP3VEC(WT_CL_Dirichlet);\n" + "" + ; + + char const *const groups[] = { + "ML_WaveToy_CL::WT_rhorhs", + "ML_WaveToy_CL::WT_urhs", + NULL}; + + static struct OpenCLKernel *kernel = NULL; + char const *const sources[] = {differencing, source, NULL}; + OpenCLRunTime_CallKernel(cctkGH, CCTK_THORNSTRING, "WT_CL_Dirichlet", + sources, groups, NULL, NULL, NULL, -1, + imin, imax, &kernel); + +} + +extern "C" void WT_CL_Dirichlet(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Entering WT_CL_Dirichlet_Body"); + } + + if (cctk_iteration % WT_CL_Dirichlet_calc_every != WT_CL_Dirichlet_calc_offset) + { + return; + } + + const char *const groups[] = { + "ML_WaveToy_CL::WT_rhorhs", + "ML_WaveToy_CL::WT_urhs"}; + GenericFD_AssertGroupStorage(cctkGH, "WT_CL_Dirichlet", 2, groups); + + + GenericFD_LoopOverBoundary(cctkGH, WT_CL_Dirichlet_Body); + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Leaving WT_CL_Dirichlet_Body"); + } +} diff --git a/ML_WaveToy_CL/src/WT_CL_Energy.cc b/ML_WaveToy_CL/src/WT_CL_Energy.cc new file mode 100644 index 0000000..56a07ea --- /dev/null +++ b/ML_WaveToy_CL/src/WT_CL_Energy.cc @@ -0,0 +1,169 @@ +/* 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" +#include "OpenCLRunTime.h" +#include "vectors.h" + +/* Define macros used in calculations */ +#define INITVALUE (42) +#define QAD(x) (SQR(SQR(x))) +#define INV(x) (kdiv(ToReal(1.0),x)) +#define SQR(x) (kmul(x,x)) +#define CUB(x) (kmul(x,SQR(x))) + +extern "C" void WT_CL_Energy_SelectBCs(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + CCTK_INT ierr = 0; + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "ML_WaveToy_CL::WT_eps","flat"); + if (ierr < 0) + CCTK_WARN(1, "Failed to register flat BC for ML_WaveToy_CL::WT_eps."); + return; +} + +static void WT_CL_Energy_Body(cGH const * restrict const cctkGH, int const dir, int const face, CCTK_REAL const normal[3], CCTK_REAL const tangentA[3], CCTK_REAL const tangentB[3], int const imin[3], int const imax[3], int const n_subblock_gfs, CCTK_REAL * restrict const subblock_gfs[]) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + char const *const source = + "\n" + "/* Include user-supplied include files */\n" + "\n" + "/* Initialise finite differencing variables */\n" + "ptrdiff_t const di = 1;\n" + "ptrdiff_t const dj = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" + "ptrdiff_t const dk = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" + "ptrdiff_t const cdi = sizeof(CCTK_REAL) * di;\n" + "ptrdiff_t const cdj = sizeof(CCTK_REAL) * dj;\n" + "ptrdiff_t const cdk = sizeof(CCTK_REAL) * dk;\n" + "CCTK_REAL_VEC const dx = ToReal(CCTK_DELTA_SPACE(0));\n" + "CCTK_REAL_VEC const dy = ToReal(CCTK_DELTA_SPACE(1));\n" + "CCTK_REAL_VEC const dz = ToReal(CCTK_DELTA_SPACE(2));\n" + "CCTK_REAL_VEC const dt = ToReal(CCTK_DELTA_TIME);\n" + "CCTK_REAL_VEC const t = ToReal(cctk_time);\n" + "CCTK_REAL_VEC const dxi = INV(dx);\n" + "CCTK_REAL_VEC const dyi = INV(dy);\n" + "CCTK_REAL_VEC const dzi = INV(dz);\n" + "CCTK_REAL_VEC const khalf = ToReal(0.5);\n" + "CCTK_REAL_VEC const kthird = ToReal(1.0/3.0);\n" + "CCTK_REAL_VEC const ktwothird = ToReal(2.0/3.0);\n" + "CCTK_REAL_VEC const kfourthird = ToReal(4.0/3.0);\n" + "CCTK_REAL_VEC const keightthird = ToReal(8.0/3.0);\n" + "CCTK_REAL_VEC const hdxi = kmul(ToReal(0.5), dxi);\n" + "CCTK_REAL_VEC const hdyi = kmul(ToReal(0.5), dyi);\n" + "CCTK_REAL_VEC const hdzi = kmul(ToReal(0.5), dzi);\n" + "\n" + "/* Initialize predefined quantities */\n" + "CCTK_REAL_VEC const p1o12dx = kmul(INV(dx),ToReal(0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const p1o12dy = kmul(INV(dy),ToReal(0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const p1o12dz = kmul(INV(dz),ToReal(0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const p1o144dxdy = kmul(INV(kmul(dx,dy)),ToReal(0.00694444444444444444444444444444));\n" + "CCTK_REAL_VEC const p1o144dxdz = kmul(INV(kmul(dx,dz)),ToReal(0.00694444444444444444444444444444));\n" + "CCTK_REAL_VEC const p1o144dydz = kmul(INV(kmul(dy,dz)),ToReal(0.00694444444444444444444444444444));\n" + "CCTK_REAL_VEC const pm1o12dx2 = kmul(INV(SQR(dx)),ToReal(-0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const pm1o12dy2 = kmul(INV(SQR(dy)),ToReal(-0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const pm1o12dz2 = kmul(INV(SQR(dz)),ToReal(-0.0833333333333333333333333333333));\n" + "\n" + "/* Assign local copies of arrays functions */\n" + "\n" + "\n" + "\n" + "/* Calculate temporaries and arrays functions */\n" + "\n" + "/* Copy local copies back to grid functions */\n" + "\n" + "/* Loop over the grid points */\n" + "#pragma omp parallel\n" + "LC_LOOP3VEC(WT_CL_Energy,\n" + " i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2],\n" + " cctk_lsh[0],cctk_lsh[1],cctk_lsh[2],\n" + " CCTK_REAL_VEC_SIZE)\n" + "{\n" + " ptrdiff_t const index = di*i + dj*j + dk*k;\n" + " \n" + " /* Assign local copies of grid functions */\n" + " \n" + " CCTK_REAL_VEC rhoL = vec_load(rho[index]);\n" + " CCTK_REAL_VEC uL = vec_load(u[index]);\n" + " \n" + " \n" + " /* Include user supplied include files */\n" + " \n" + " /* Precompute derivatives */\n" + " CCTK_REAL_VEC const PDstandardNth1u = PDstandardNth1(&u[index]);\n" + " CCTK_REAL_VEC const PDstandardNth2u = PDstandardNth2(&u[index]);\n" + " CCTK_REAL_VEC const PDstandardNth3u = PDstandardNth3(&u[index]);\n" + " \n" + " /* Calculate temporaries and grid functions */\n" + " CCTK_REAL_VEC epsL = \n" + " kmul(kadd(SQR(rhoL),kadd(SQR(PDstandardNth1u),kadd(SQR(PDstandardNth2u),SQR(PDstandardNth3u)))),ToReal(0.5));\n" + " \n" + " /* Copy local copies back to grid functions */\n" + " vec_store_partial_prepare(i,lc_imin,lc_imax);\n" + " vec_store_nta_partial(eps[index],epsL);\n" + "}\n" + "LC_ENDLOOP3VEC(WT_CL_Energy);\n" + "" + ; + + char const *const groups[] = { + "ML_WaveToy_CL::WT_eps", + "ML_WaveToy_CL::WT_rho", + "ML_WaveToy_CL::WT_u", + NULL}; + + static struct OpenCLKernel *kernel = NULL; + char const *const sources[] = {differencing, source, NULL}; + OpenCLRunTime_CallKernel(cctkGH, CCTK_THORNSTRING, "WT_CL_Energy", + sources, groups, NULL, NULL, NULL, -1, + imin, imax, &kernel); + +} + +extern "C" void WT_CL_Energy(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Entering WT_CL_Energy_Body"); + } + + if (cctk_iteration % WT_CL_Energy_calc_every != WT_CL_Energy_calc_offset) + { + return; + } + + const char *const groups[] = { + "ML_WaveToy_CL::WT_eps", + "ML_WaveToy_CL::WT_rho", + "ML_WaveToy_CL::WT_u"}; + GenericFD_AssertGroupStorage(cctkGH, "WT_CL_Energy", 3, groups); + + GenericFD_EnsureStencilFits(cctkGH, "WT_CL_Energy", 2, 2, 2); + + GenericFD_LoopOverInterior(cctkGH, WT_CL_Energy_Body); + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Leaving WT_CL_Energy_Body"); + } +} diff --git a/ML_WaveToy_CL/src/WT_CL_EnergyBoundary.cc b/ML_WaveToy_CL/src/WT_CL_EnergyBoundary.cc new file mode 100644 index 0000000..9425b0d --- /dev/null +++ b/ML_WaveToy_CL/src/WT_CL_EnergyBoundary.cc @@ -0,0 +1,158 @@ +/* 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" +#include "OpenCLRunTime.h" +#include "vectors.h" + +/* Define macros used in calculations */ +#define INITVALUE (42) +#define QAD(x) (SQR(SQR(x))) +#define INV(x) (kdiv(ToReal(1.0),x)) +#define SQR(x) (kmul(x,x)) +#define CUB(x) (kmul(x,SQR(x))) + +extern "C" void WT_CL_EnergyBoundary_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 */, "ML_WaveToy_CL::WT_eps","flat"); + if (ierr < 0) + CCTK_WARN(1, "Failed to register flat BC for ML_WaveToy_CL::WT_eps."); + return; +} + +static void WT_CL_EnergyBoundary_Body(cGH const * restrict const cctkGH, int const dir, int const face, CCTK_REAL const normal[3], CCTK_REAL const tangentA[3], CCTK_REAL const tangentB[3], int const imin[3], int const imax[3], int const n_subblock_gfs, CCTK_REAL * restrict const subblock_gfs[]) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + char const *const source = + "\n" + "/* Include user-supplied include files */\n" + "\n" + "/* Initialise finite differencing variables */\n" + "ptrdiff_t const di = 1;\n" + "ptrdiff_t const dj = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" + "ptrdiff_t const dk = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" + "ptrdiff_t const cdi = sizeof(CCTK_REAL) * di;\n" + "ptrdiff_t const cdj = sizeof(CCTK_REAL) * dj;\n" + "ptrdiff_t const cdk = sizeof(CCTK_REAL) * dk;\n" + "CCTK_REAL_VEC const dx = ToReal(CCTK_DELTA_SPACE(0));\n" + "CCTK_REAL_VEC const dy = ToReal(CCTK_DELTA_SPACE(1));\n" + "CCTK_REAL_VEC const dz = ToReal(CCTK_DELTA_SPACE(2));\n" + "CCTK_REAL_VEC const dt = ToReal(CCTK_DELTA_TIME);\n" + "CCTK_REAL_VEC const t = ToReal(cctk_time);\n" + "CCTK_REAL_VEC const dxi = INV(dx);\n" + "CCTK_REAL_VEC const dyi = INV(dy);\n" + "CCTK_REAL_VEC const dzi = INV(dz);\n" + "CCTK_REAL_VEC const khalf = ToReal(0.5);\n" + "CCTK_REAL_VEC const kthird = ToReal(1.0/3.0);\n" + "CCTK_REAL_VEC const ktwothird = ToReal(2.0/3.0);\n" + "CCTK_REAL_VEC const kfourthird = ToReal(4.0/3.0);\n" + "CCTK_REAL_VEC const keightthird = ToReal(8.0/3.0);\n" + "CCTK_REAL_VEC const hdxi = kmul(ToReal(0.5), dxi);\n" + "CCTK_REAL_VEC const hdyi = kmul(ToReal(0.5), dyi);\n" + "CCTK_REAL_VEC const hdzi = kmul(ToReal(0.5), dzi);\n" + "\n" + "/* Initialize predefined quantities */\n" + "CCTK_REAL_VEC const p1o12dx = kmul(INV(dx),ToReal(0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const p1o12dy = kmul(INV(dy),ToReal(0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const p1o12dz = kmul(INV(dz),ToReal(0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const p1o144dxdy = kmul(INV(kmul(dx,dy)),ToReal(0.00694444444444444444444444444444));\n" + "CCTK_REAL_VEC const p1o144dxdz = kmul(INV(kmul(dx,dz)),ToReal(0.00694444444444444444444444444444));\n" + "CCTK_REAL_VEC const p1o144dydz = kmul(INV(kmul(dy,dz)),ToReal(0.00694444444444444444444444444444));\n" + "CCTK_REAL_VEC const pm1o12dx2 = kmul(INV(SQR(dx)),ToReal(-0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const pm1o12dy2 = kmul(INV(SQR(dy)),ToReal(-0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const pm1o12dz2 = kmul(INV(SQR(dz)),ToReal(-0.0833333333333333333333333333333));\n" + "\n" + "/* Assign local copies of arrays functions */\n" + "\n" + "\n" + "\n" + "/* Calculate temporaries and arrays functions */\n" + "\n" + "/* Copy local copies back to grid functions */\n" + "\n" + "/* Loop over the grid points */\n" + "#pragma omp parallel\n" + "LC_LOOP3VEC(WT_CL_EnergyBoundary,\n" + " i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2],\n" + " cctk_lsh[0],cctk_lsh[1],cctk_lsh[2],\n" + " CCTK_REAL_VEC_SIZE)\n" + "{\n" + " ptrdiff_t const index = di*i + dj*j + dk*k;\n" + " \n" + " /* Assign local copies of grid functions */\n" + " \n" + " \n" + " \n" + " /* Include user supplied include files */\n" + " \n" + " /* Precompute derivatives */\n" + " \n" + " /* Calculate temporaries and grid functions */\n" + " CCTK_REAL_VEC epsL = ToReal(0);\n" + " \n" + " /* Copy local copies back to grid functions */\n" + " vec_store_partial_prepare(i,lc_imin,lc_imax);\n" + " vec_store_nta_partial(eps[index],epsL);\n" + "}\n" + "LC_ENDLOOP3VEC(WT_CL_EnergyBoundary);\n" + "" + ; + + char const *const groups[] = { + "ML_WaveToy_CL::WT_eps", + NULL}; + + static struct OpenCLKernel *kernel = NULL; + char const *const sources[] = {differencing, source, NULL}; + OpenCLRunTime_CallKernel(cctkGH, CCTK_THORNSTRING, "WT_CL_EnergyBoundary", + sources, groups, NULL, NULL, NULL, -1, + imin, imax, &kernel); + +} + +extern "C" void WT_CL_EnergyBoundary(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Entering WT_CL_EnergyBoundary_Body"); + } + + if (cctk_iteration % WT_CL_EnergyBoundary_calc_every != WT_CL_EnergyBoundary_calc_offset) + { + return; + } + + const char *const groups[] = { + "ML_WaveToy_CL::WT_eps"}; + GenericFD_AssertGroupStorage(cctkGH, "WT_CL_EnergyBoundary", 1, groups); + + + GenericFD_LoopOverBoundary(cctkGH, WT_CL_EnergyBoundary_Body); + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Leaving WT_CL_EnergyBoundary_Body"); + } +} diff --git a/ML_WaveToy_CL/src/WT_CL_Gaussian.cc b/ML_WaveToy_CL/src/WT_CL_Gaussian.cc new file mode 100644 index 0000000..e121995 --- /dev/null +++ b/ML_WaveToy_CL/src/WT_CL_Gaussian.cc @@ -0,0 +1,155 @@ +/* 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" +#include "OpenCLRunTime.h" +#include "vectors.h" + +/* Define macros used in calculations */ +#define INITVALUE (42) +#define QAD(x) (SQR(SQR(x))) +#define INV(x) (kdiv(ToReal(1.0),x)) +#define SQR(x) (kmul(x,x)) +#define CUB(x) (kmul(x,SQR(x))) + +static void WT_CL_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; + + char const *const source = + "\n" + "/* Include user-supplied include files */\n" + "\n" + "/* Initialise finite differencing variables */\n" + "ptrdiff_t const di = 1;\n" + "ptrdiff_t const dj = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" + "ptrdiff_t const dk = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" + "ptrdiff_t const cdi = sizeof(CCTK_REAL) * di;\n" + "ptrdiff_t const cdj = sizeof(CCTK_REAL) * dj;\n" + "ptrdiff_t const cdk = sizeof(CCTK_REAL) * dk;\n" + "CCTK_REAL_VEC const dx = ToReal(CCTK_DELTA_SPACE(0));\n" + "CCTK_REAL_VEC const dy = ToReal(CCTK_DELTA_SPACE(1));\n" + "CCTK_REAL_VEC const dz = ToReal(CCTK_DELTA_SPACE(2));\n" + "CCTK_REAL_VEC const dt = ToReal(CCTK_DELTA_TIME);\n" + "CCTK_REAL_VEC const t = ToReal(cctk_time);\n" + "CCTK_REAL_VEC const dxi = INV(dx);\n" + "CCTK_REAL_VEC const dyi = INV(dy);\n" + "CCTK_REAL_VEC const dzi = INV(dz);\n" + "CCTK_REAL_VEC const khalf = ToReal(0.5);\n" + "CCTK_REAL_VEC const kthird = ToReal(1.0/3.0);\n" + "CCTK_REAL_VEC const ktwothird = ToReal(2.0/3.0);\n" + "CCTK_REAL_VEC const kfourthird = ToReal(4.0/3.0);\n" + "CCTK_REAL_VEC const keightthird = ToReal(8.0/3.0);\n" + "CCTK_REAL_VEC const hdxi = kmul(ToReal(0.5), dxi);\n" + "CCTK_REAL_VEC const hdyi = kmul(ToReal(0.5), dyi);\n" + "CCTK_REAL_VEC const hdzi = kmul(ToReal(0.5), dzi);\n" + "\n" + "/* Initialize predefined quantities */\n" + "CCTK_REAL_VEC const p1o12dx = kmul(INV(dx),ToReal(0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const p1o12dy = kmul(INV(dy),ToReal(0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const p1o12dz = kmul(INV(dz),ToReal(0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const p1o144dxdy = kmul(INV(kmul(dx,dy)),ToReal(0.00694444444444444444444444444444));\n" + "CCTK_REAL_VEC const p1o144dxdz = kmul(INV(kmul(dx,dz)),ToReal(0.00694444444444444444444444444444));\n" + "CCTK_REAL_VEC const p1o144dydz = kmul(INV(kmul(dy,dz)),ToReal(0.00694444444444444444444444444444));\n" + "CCTK_REAL_VEC const pm1o12dx2 = kmul(INV(SQR(dx)),ToReal(-0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const pm1o12dy2 = kmul(INV(SQR(dy)),ToReal(-0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const pm1o12dz2 = kmul(INV(SQR(dz)),ToReal(-0.0833333333333333333333333333333));\n" + "\n" + "/* Assign local copies of arrays functions */\n" + "\n" + "\n" + "\n" + "/* Calculate temporaries and arrays functions */\n" + "\n" + "/* Copy local copies back to grid functions */\n" + "\n" + "/* Loop over the grid points */\n" + "#pragma omp parallel\n" + "LC_LOOP3VEC(WT_CL_Gaussian,\n" + " i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2],\n" + " cctk_lsh[0],cctk_lsh[1],cctk_lsh[2],\n" + " CCTK_REAL_VEC_SIZE)\n" + "{\n" + " ptrdiff_t const index = di*i + dj*j + dk*k;\n" + " \n" + " /* Assign local copies of grid functions */\n" + " \n" + " CCTK_REAL_VEC rL = vec_load(r[index]);\n" + " \n" + " \n" + " /* Include user supplied include files */\n" + " \n" + " /* Precompute derivatives */\n" + " \n" + " /* Calculate temporaries and grid functions */\n" + " CCTK_REAL_VEC uL = \n" + " kmul(kexp(kmul(INV(SQR(ToReal(width))),kmul(SQR(rL),ToReal(-0.5)))),ToReal(amplitude));\n" + " \n" + " CCTK_REAL_VEC rhoL = ToReal(0);\n" + " \n" + " /* Copy local copies back to grid functions */\n" + " vec_store_partial_prepare(i,lc_imin,lc_imax);\n" + " vec_store_nta_partial(rho[index],rhoL);\n" + " vec_store_nta_partial(u[index],uL);\n" + "}\n" + "LC_ENDLOOP3VEC(WT_CL_Gaussian);\n" + "" + ; + + char const *const groups[] = { + "grid::coordinates", + "ML_WaveToy_CL::WT_rho", + "ML_WaveToy_CL::WT_u", + NULL}; + + static struct OpenCLKernel *kernel = NULL; + char const *const sources[] = {differencing, source, NULL}; + OpenCLRunTime_CallKernel(cctkGH, CCTK_THORNSTRING, "WT_CL_Gaussian", + sources, groups, NULL, NULL, NULL, -1, + imin, imax, &kernel); + +} + +extern "C" void WT_CL_Gaussian(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Entering WT_CL_Gaussian_Body"); + } + + if (cctk_iteration % WT_CL_Gaussian_calc_every != WT_CL_Gaussian_calc_offset) + { + return; + } + + const char *const groups[] = { + "grid::coordinates", + "ML_WaveToy_CL::WT_rho", + "ML_WaveToy_CL::WT_u"}; + GenericFD_AssertGroupStorage(cctkGH, "WT_CL_Gaussian", 3, groups); + + + GenericFD_LoopOverEverything(cctkGH, WT_CL_Gaussian_Body); + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Leaving WT_CL_Gaussian_Body"); + } +} diff --git a/ML_WaveToy_CL/src/WT_CL_RHS.cc b/ML_WaveToy_CL/src/WT_CL_RHS.cc new file mode 100644 index 0000000..2314408 --- /dev/null +++ b/ML_WaveToy_CL/src/WT_CL_RHS.cc @@ -0,0 +1,177 @@ +/* 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" +#include "OpenCLRunTime.h" +#include "vectors.h" + +/* Define macros used in calculations */ +#define INITVALUE (42) +#define QAD(x) (SQR(SQR(x))) +#define INV(x) (kdiv(ToReal(1.0),x)) +#define SQR(x) (kmul(x,x)) +#define CUB(x) (kmul(x,SQR(x))) + +extern "C" void WT_CL_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 */, "ML_WaveToy_CL::WT_rhorhs","flat"); + if (ierr < 0) + CCTK_WARN(1, "Failed to register flat BC for ML_WaveToy_CL::WT_rhorhs."); + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "ML_WaveToy_CL::WT_urhs","flat"); + if (ierr < 0) + CCTK_WARN(1, "Failed to register flat BC for ML_WaveToy_CL::WT_urhs."); + return; +} + +static void WT_CL_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; + + char const *const source = + "\n" + "/* Include user-supplied include files */\n" + "\n" + "/* Initialise finite differencing variables */\n" + "ptrdiff_t const di = 1;\n" + "ptrdiff_t const dj = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" + "ptrdiff_t const dk = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" + "ptrdiff_t const cdi = sizeof(CCTK_REAL) * di;\n" + "ptrdiff_t const cdj = sizeof(CCTK_REAL) * dj;\n" + "ptrdiff_t const cdk = sizeof(CCTK_REAL) * dk;\n" + "CCTK_REAL_VEC const dx = ToReal(CCTK_DELTA_SPACE(0));\n" + "CCTK_REAL_VEC const dy = ToReal(CCTK_DELTA_SPACE(1));\n" + "CCTK_REAL_VEC const dz = ToReal(CCTK_DELTA_SPACE(2));\n" + "CCTK_REAL_VEC const dt = ToReal(CCTK_DELTA_TIME);\n" + "CCTK_REAL_VEC const t = ToReal(cctk_time);\n" + "CCTK_REAL_VEC const dxi = INV(dx);\n" + "CCTK_REAL_VEC const dyi = INV(dy);\n" + "CCTK_REAL_VEC const dzi = INV(dz);\n" + "CCTK_REAL_VEC const khalf = ToReal(0.5);\n" + "CCTK_REAL_VEC const kthird = ToReal(1.0/3.0);\n" + "CCTK_REAL_VEC const ktwothird = ToReal(2.0/3.0);\n" + "CCTK_REAL_VEC const kfourthird = ToReal(4.0/3.0);\n" + "CCTK_REAL_VEC const keightthird = ToReal(8.0/3.0);\n" + "CCTK_REAL_VEC const hdxi = kmul(ToReal(0.5), dxi);\n" + "CCTK_REAL_VEC const hdyi = kmul(ToReal(0.5), dyi);\n" + "CCTK_REAL_VEC const hdzi = kmul(ToReal(0.5), dzi);\n" + "\n" + "/* Initialize predefined quantities */\n" + "CCTK_REAL_VEC const p1o12dx = kmul(INV(dx),ToReal(0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const p1o12dy = kmul(INV(dy),ToReal(0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const p1o12dz = kmul(INV(dz),ToReal(0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const p1o144dxdy = kmul(INV(kmul(dx,dy)),ToReal(0.00694444444444444444444444444444));\n" + "CCTK_REAL_VEC const p1o144dxdz = kmul(INV(kmul(dx,dz)),ToReal(0.00694444444444444444444444444444));\n" + "CCTK_REAL_VEC const p1o144dydz = kmul(INV(kmul(dy,dz)),ToReal(0.00694444444444444444444444444444));\n" + "CCTK_REAL_VEC const pm1o12dx2 = kmul(INV(SQR(dx)),ToReal(-0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const pm1o12dy2 = kmul(INV(SQR(dy)),ToReal(-0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const pm1o12dz2 = kmul(INV(SQR(dz)),ToReal(-0.0833333333333333333333333333333));\n" + "\n" + "/* Assign local copies of arrays functions */\n" + "\n" + "\n" + "\n" + "/* Calculate temporaries and arrays functions */\n" + "\n" + "/* Copy local copies back to grid functions */\n" + "\n" + "/* Loop over the grid points */\n" + "#pragma omp parallel\n" + "LC_LOOP3VEC(WT_CL_RHS,\n" + " i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2],\n" + " cctk_lsh[0],cctk_lsh[1],cctk_lsh[2],\n" + " CCTK_REAL_VEC_SIZE)\n" + "{\n" + " ptrdiff_t const index = di*i + dj*j + dk*k;\n" + " \n" + " /* Assign local copies of grid functions */\n" + " \n" + " CCTK_REAL_VEC rhoL = vec_load(rho[index]);\n" + " CCTK_REAL_VEC uL = vec_load(u[index]);\n" + " \n" + " \n" + " /* Include user supplied include files */\n" + " \n" + " /* Precompute derivatives */\n" + " CCTK_REAL_VEC const PDstandardNth11u = PDstandardNth11(&u[index]);\n" + " CCTK_REAL_VEC const PDstandardNth22u = PDstandardNth22(&u[index]);\n" + " CCTK_REAL_VEC const PDstandardNth33u = PDstandardNth33(&u[index]);\n" + " \n" + " /* Calculate temporaries and grid functions */\n" + " CCTK_REAL_VEC urhsL = rhoL;\n" + " \n" + " CCTK_REAL_VEC rhorhsL = \n" + " kadd(PDstandardNth11u,kadd(PDstandardNth22u,PDstandardNth33u));\n" + " \n" + " /* Copy local copies back to grid functions */\n" + " vec_store_partial_prepare(i,lc_imin,lc_imax);\n" + " vec_store_nta_partial(rhorhs[index],rhorhsL);\n" + " vec_store_nta_partial(urhs[index],urhsL);\n" + "}\n" + "LC_ENDLOOP3VEC(WT_CL_RHS);\n" + "" + ; + + char const *const groups[] = { + "ML_WaveToy_CL::WT_rho", + "ML_WaveToy_CL::WT_rhorhs", + "ML_WaveToy_CL::WT_u", + "ML_WaveToy_CL::WT_urhs", + NULL}; + + static struct OpenCLKernel *kernel = NULL; + char const *const sources[] = {differencing, source, NULL}; + OpenCLRunTime_CallKernel(cctkGH, CCTK_THORNSTRING, "WT_CL_RHS", + sources, groups, NULL, NULL, NULL, -1, + imin, imax, &kernel); + +} + +extern "C" void WT_CL_RHS(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Entering WT_CL_RHS_Body"); + } + + if (cctk_iteration % WT_CL_RHS_calc_every != WT_CL_RHS_calc_offset) + { + return; + } + + const char *const groups[] = { + "ML_WaveToy_CL::WT_rho", + "ML_WaveToy_CL::WT_rhorhs", + "ML_WaveToy_CL::WT_u", + "ML_WaveToy_CL::WT_urhs"}; + GenericFD_AssertGroupStorage(cctkGH, "WT_CL_RHS", 4, groups); + + GenericFD_EnsureStencilFits(cctkGH, "WT_CL_RHS", 2, 2, 2); + + GenericFD_LoopOverInterior(cctkGH, WT_CL_RHS_Body); + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Leaving WT_CL_RHS_Body"); + } +} diff --git a/ML_WaveToy_CL/src/WT_CL_Standing.cc b/ML_WaveToy_CL/src/WT_CL_Standing.cc new file mode 100644 index 0000000..5c477e1 --- /dev/null +++ b/ML_WaveToy_CL/src/WT_CL_Standing.cc @@ -0,0 +1,162 @@ +/* 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" +#include "OpenCLRunTime.h" +#include "vectors.h" + +/* Define macros used in calculations */ +#define INITVALUE (42) +#define QAD(x) (SQR(SQR(x))) +#define INV(x) (kdiv(ToReal(1.0),x)) +#define SQR(x) (kmul(x,x)) +#define CUB(x) (kmul(x,SQR(x))) + +static void WT_CL_Standing_Body(cGH const * restrict const cctkGH, int const dir, int const face, CCTK_REAL const normal[3], CCTK_REAL const tangentA[3], CCTK_REAL const tangentB[3], int const imin[3], int const imax[3], int const n_subblock_gfs, CCTK_REAL * restrict const subblock_gfs[]) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + char const *const source = + "\n" + "/* Include user-supplied include files */\n" + "\n" + "/* Initialise finite differencing variables */\n" + "ptrdiff_t const di = 1;\n" + "ptrdiff_t const dj = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" + "ptrdiff_t const dk = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" + "ptrdiff_t const cdi = sizeof(CCTK_REAL) * di;\n" + "ptrdiff_t const cdj = sizeof(CCTK_REAL) * dj;\n" + "ptrdiff_t const cdk = sizeof(CCTK_REAL) * dk;\n" + "CCTK_REAL_VEC const dx = ToReal(CCTK_DELTA_SPACE(0));\n" + "CCTK_REAL_VEC const dy = ToReal(CCTK_DELTA_SPACE(1));\n" + "CCTK_REAL_VEC const dz = ToReal(CCTK_DELTA_SPACE(2));\n" + "CCTK_REAL_VEC const dt = ToReal(CCTK_DELTA_TIME);\n" + "CCTK_REAL_VEC const t = ToReal(cctk_time);\n" + "CCTK_REAL_VEC const dxi = INV(dx);\n" + "CCTK_REAL_VEC const dyi = INV(dy);\n" + "CCTK_REAL_VEC const dzi = INV(dz);\n" + "CCTK_REAL_VEC const khalf = ToReal(0.5);\n" + "CCTK_REAL_VEC const kthird = ToReal(1.0/3.0);\n" + "CCTK_REAL_VEC const ktwothird = ToReal(2.0/3.0);\n" + "CCTK_REAL_VEC const kfourthird = ToReal(4.0/3.0);\n" + "CCTK_REAL_VEC const keightthird = ToReal(8.0/3.0);\n" + "CCTK_REAL_VEC const hdxi = kmul(ToReal(0.5), dxi);\n" + "CCTK_REAL_VEC const hdyi = kmul(ToReal(0.5), dyi);\n" + "CCTK_REAL_VEC const hdzi = kmul(ToReal(0.5), dzi);\n" + "\n" + "/* Initialize predefined quantities */\n" + "CCTK_REAL_VEC const p1o12dx = kmul(INV(dx),ToReal(0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const p1o12dy = kmul(INV(dy),ToReal(0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const p1o12dz = kmul(INV(dz),ToReal(0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const p1o144dxdy = kmul(INV(kmul(dx,dy)),ToReal(0.00694444444444444444444444444444));\n" + "CCTK_REAL_VEC const p1o144dxdz = kmul(INV(kmul(dx,dz)),ToReal(0.00694444444444444444444444444444));\n" + "CCTK_REAL_VEC const p1o144dydz = kmul(INV(kmul(dy,dz)),ToReal(0.00694444444444444444444444444444));\n" + "CCTK_REAL_VEC const pm1o12dx2 = kmul(INV(SQR(dx)),ToReal(-0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const pm1o12dy2 = kmul(INV(SQR(dy)),ToReal(-0.0833333333333333333333333333333));\n" + "CCTK_REAL_VEC const pm1o12dz2 = kmul(INV(SQR(dz)),ToReal(-0.0833333333333333333333333333333));\n" + "\n" + "/* Assign local copies of arrays functions */\n" + "\n" + "\n" + "\n" + "/* Calculate temporaries and arrays functions */\n" + "\n" + "/* Copy local copies back to grid functions */\n" + "\n" + "/* Loop over the grid points */\n" + "#pragma omp parallel\n" + "LC_LOOP3VEC(WT_CL_Standing,\n" + " i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2],\n" + " cctk_lsh[0],cctk_lsh[1],cctk_lsh[2],\n" + " CCTK_REAL_VEC_SIZE)\n" + "{\n" + " ptrdiff_t const index = di*i + dj*j + dk*k;\n" + " \n" + " /* Assign local copies of grid functions */\n" + " \n" + " CCTK_REAL_VEC xL = vec_load(x[index]);\n" + " CCTK_REAL_VEC yL = vec_load(y[index]);\n" + " CCTK_REAL_VEC zL = vec_load(z[index]);\n" + " \n" + " \n" + " /* Include user supplied include files */\n" + " \n" + " /* Precompute derivatives */\n" + " \n" + " /* Calculate temporaries and grid functions */\n" + " CCTK_REAL_VEC k = kmul(INV(ToReal(width)),ToReal(Pi));\n" + " \n" + " CCTK_REAL_VEC omega = ksqrt(kmul(SQR(k),ToReal(3)));\n" + " \n" + " CCTK_REAL_VEC uL = \n" + " kmul(kcos(kmul(xL,k)),kmul(kcos(kmul(yL,k)),kmul(kcos(kmul(zL,k)),kmul(kcos(kmul(omega,t)),ToReal(amplitude)))));\n" + " \n" + " CCTK_REAL_VEC rhoL = \n" + " kneg(kmul(omega,kmul(kcos(kmul(xL,k)),kmul(kcos(kmul(yL,k)),kmul(kcos(kmul(zL,k)),kmul(ksin(kmul(omega,t)),ToReal(amplitude)))))));\n" + " \n" + " /* Copy local copies back to grid functions */\n" + " vec_store_partial_prepare(i,lc_imin,lc_imax);\n" + " vec_store_nta_partial(rho[index],rhoL);\n" + " vec_store_nta_partial(u[index],uL);\n" + "}\n" + "LC_ENDLOOP3VEC(WT_CL_Standing);\n" + "" + ; + + char const *const groups[] = { + "grid::coordinates", + "ML_WaveToy_CL::WT_rho", + "ML_WaveToy_CL::WT_u", + NULL}; + + static struct OpenCLKernel *kernel = NULL; + char const *const sources[] = {differencing, source, NULL}; + OpenCLRunTime_CallKernel(cctkGH, CCTK_THORNSTRING, "WT_CL_Standing", + sources, groups, NULL, NULL, NULL, -1, + imin, imax, &kernel); + +} + +extern "C" void WT_CL_Standing(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Entering WT_CL_Standing_Body"); + } + + if (cctk_iteration % WT_CL_Standing_calc_every != WT_CL_Standing_calc_offset) + { + return; + } + + const char *const groups[] = { + "grid::coordinates", + "ML_WaveToy_CL::WT_rho", + "ML_WaveToy_CL::WT_u"}; + GenericFD_AssertGroupStorage(cctkGH, "WT_CL_Standing", 3, groups); + + + GenericFD_LoopOverEverything(cctkGH, WT_CL_Standing_Body); + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Leaving WT_CL_Standing_Body"); + } +} diff --git a/ML_WaveToy_CL/src/make.code.defn b/ML_WaveToy_CL/src/make.code.defn new file mode 100644 index 0000000..0069b28 --- /dev/null +++ b/ML_WaveToy_CL/src/make.code.defn @@ -0,0 +1,3 @@ +# File produced by Kranc + +SRCS = Startup.cc RegisterSymmetries.cc RegisterMoL.cc WT_CL_Gaussian.cc WT_CL_Standing.cc WT_CL_RHS.cc WT_CL_Dirichlet.cc WT_CL_Energy.cc WT_CL_EnergyBoundary.cc Boundaries.cc -- cgit v1.2.3