From 658fa6fcb2dfea4354d734af7cb3d946e54342fa Mon Sep 17 00:00:00 2001 From: Ian Hinder Date: Tue, 31 Jan 2012 14:22:14 -0600 Subject: Generate SimpleWaveOpenCL --- Examples/SimpleWaveOpenCL/cakernel.ccl | 1 + Examples/SimpleWaveOpenCL/configuration.ccl | 7 + Examples/SimpleWaveOpenCL/interface.ccl | 43 +++++ Examples/SimpleWaveOpenCL/param.ccl | 162 +++++++++++++++++ Examples/SimpleWaveOpenCL/schedule.ccl | 74 ++++++++ Examples/SimpleWaveOpenCL/src/Boundaries.cc | 197 +++++++++++++++++++++ Examples/SimpleWaveOpenCL/src/Differencing.h | 75 ++++++++ Examples/SimpleWaveOpenCL/src/RegisterMoL.cc | 20 +++ .../SimpleWaveOpenCL/src/RegisterSymmetries.cc | 29 +++ Examples/SimpleWaveOpenCL/src/Startup.cc | 10 ++ Examples/SimpleWaveOpenCL/src/calc_rhs.cc | 159 +++++++++++++++++ Examples/SimpleWaveOpenCL/src/initial_sine.cc | 141 +++++++++++++++ Examples/SimpleWaveOpenCL/src/make.code.defn | 3 + 13 files changed, 921 insertions(+) create mode 100644 Examples/SimpleWaveOpenCL/cakernel.ccl create mode 100644 Examples/SimpleWaveOpenCL/configuration.ccl create mode 100644 Examples/SimpleWaveOpenCL/interface.ccl create mode 100644 Examples/SimpleWaveOpenCL/param.ccl create mode 100644 Examples/SimpleWaveOpenCL/schedule.ccl create mode 100644 Examples/SimpleWaveOpenCL/src/Boundaries.cc create mode 100644 Examples/SimpleWaveOpenCL/src/Differencing.h create mode 100644 Examples/SimpleWaveOpenCL/src/RegisterMoL.cc create mode 100644 Examples/SimpleWaveOpenCL/src/RegisterSymmetries.cc create mode 100644 Examples/SimpleWaveOpenCL/src/Startup.cc create mode 100644 Examples/SimpleWaveOpenCL/src/calc_rhs.cc create mode 100644 Examples/SimpleWaveOpenCL/src/initial_sine.cc create mode 100644 Examples/SimpleWaveOpenCL/src/make.code.defn (limited to 'Examples') diff --git a/Examples/SimpleWaveOpenCL/cakernel.ccl b/Examples/SimpleWaveOpenCL/cakernel.ccl new file mode 100644 index 0000000..32dc04a --- /dev/null +++ b/Examples/SimpleWaveOpenCL/cakernel.ccl @@ -0,0 +1 @@ +KrancThorn`Private`cakernel$438 \ No newline at end of file diff --git a/Examples/SimpleWaveOpenCL/configuration.ccl b/Examples/SimpleWaveOpenCL/configuration.ccl new file mode 100644 index 0000000..6274e84 --- /dev/null +++ b/Examples/SimpleWaveOpenCL/configuration.ccl @@ -0,0 +1,7 @@ +# File produced by Kranc + +REQUIRES GenericFD +OPTIONAL LoopControl +{ +} +REQUIRES OpenCL OpenCLRunTime diff --git a/Examples/SimpleWaveOpenCL/interface.ccl b/Examples/SimpleWaveOpenCL/interface.ccl new file mode 100644 index 0000000..2c987ce --- /dev/null +++ b/Examples/SimpleWaveOpenCL/interface.ccl @@ -0,0 +1,43 @@ +# File produced by Kranc + +implements: SimpleWaveOpenCL + +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 + +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 evolved_group type=GF timelevels=3 tags='' +{ + phi, + pi +} "evolved_group" + +public: +CCTK_REAL evolved_grouprhs type=GF timelevels=3 tags='' +{ + phirhs, + pirhs +} "evolved_grouprhs" diff --git a/Examples/SimpleWaveOpenCL/param.ccl b/Examples/SimpleWaveOpenCL/param.ccl new file mode 100644 index 0000000..a1082da --- /dev/null +++ b/Examples/SimpleWaveOpenCL/param.ccl @@ -0,0 +1,162 @@ +# 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 SimpleWaveOpenCL_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 SimpleWaveOpenCL_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:3 :: "" +} 3 + +restricted: +CCTK_INT rhs_timelevels "Number of active RHS timelevels" STEERABLE=RECOVER +{ + 0:3 :: "" +} 1 + +restricted: +CCTK_INT other_timelevels "Number of active timelevels for non-evolved grid functions" STEERABLE=RECOVER +{ + 0:3 :: "" +} 1 + +restricted: +CCTK_INT initial_sine_calc_every "initial_sine_calc_every" STEERABLE=ALWAYS +{ + *:* :: "" +} 1 + +restricted: +CCTK_INT calc_rhs_calc_every "calc_rhs_calc_every" STEERABLE=ALWAYS +{ + *:* :: "" +} 1 + +restricted: +CCTK_INT initial_sine_calc_offset "initial_sine_calc_offset" STEERABLE=ALWAYS +{ + *:* :: "" +} 0 + +restricted: +CCTK_INT calc_rhs_calc_offset "calc_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 evolved_group_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 evolved_group_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 evolved_group_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 evolved_group_bound_scalar "Dirichlet boundary value" STEERABLE=ALWAYS +{ + "*:*" :: "unrestricted" +} 0. + diff --git a/Examples/SimpleWaveOpenCL/schedule.ccl b/Examples/SimpleWaveOpenCL/schedule.ccl new file mode 100644 index 0000000..8395c0a --- /dev/null +++ b/Examples/SimpleWaveOpenCL/schedule.ccl @@ -0,0 +1,74 @@ +# File produced by Kranc + + +if (timelevels == 1) +{ + STORAGE: evolved_group[1] +} +if (timelevels == 2) +{ + STORAGE: evolved_group[2] +} +if (timelevels == 3) +{ + STORAGE: evolved_group[3] +} + +if (rhs_timelevels == 1) +{ + STORAGE: evolved_grouprhs[1] +} +if (rhs_timelevels == 2) +{ + STORAGE: evolved_grouprhs[2] +} +if (rhs_timelevels == 3) +{ + STORAGE: evolved_grouprhs[3] +} + +schedule SimpleWaveOpenCL_Startup at STARTUP +{ + LANG: C + OPTIONS: meta +} "create banner" + +schedule SimpleWaveOpenCL_RegisterVars in MoL_Register +{ + LANG: C + OPTIONS: meta +} "Register Variables for MoL" + +schedule SimpleWaveOpenCL_RegisterSymmetries in SymmetryRegister +{ + LANG: C + OPTIONS: meta +} "register symmetries" + +schedule initial_sine AT INITIAL +{ + LANG: C +} "initial_sine" + +schedule calc_rhs in MoL_CalcRHS +{ + LANG: C +} "calc_rhs" + +schedule SimpleWaveOpenCL_SelectBoundConds in MoL_PostStep +{ + LANG: C + OPTIONS: level + SYNC: evolved_group +} "select boundary conditions" + +schedule SimpleWaveOpenCL_CheckBoundaries at BASEGRID +{ + LANG: C + OPTIONS: meta +} "check boundaries treatment" + +schedule group ApplyBCs as SimpleWaveOpenCL_ApplyBCs in MoL_PostStep after SimpleWaveOpenCL_SelectBoundConds +{ + # no language specified +} "Apply boundary conditions controlled by thorn Boundary" diff --git a/Examples/SimpleWaveOpenCL/src/Boundaries.cc b/Examples/SimpleWaveOpenCL/src/Boundaries.cc new file mode 100644 index 0000000..fc19996 --- /dev/null +++ b/Examples/SimpleWaveOpenCL/src/Boundaries.cc @@ -0,0 +1,197 @@ +/* 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 SimpleWaveOpenCL_CheckBoundaries(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + return; +} + +extern "C" void SimpleWaveOpenCL_SelectBoundConds(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + CCTK_INT ierr = 0; + + if (CCTK_EQUALS(evolved_group_bound, "none" ) || + CCTK_EQUALS(evolved_group_bound, "static") || + CCTK_EQUALS(evolved_group_bound, "flat" ) || + CCTK_EQUALS(evolved_group_bound, "zero" ) ) + { + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, -1, + "SimpleWaveOpenCL::evolved_group", evolved_group_bound); + if (ierr < 0) + CCTK_WARN(0, "Failed to register evolved_group_bound BC for SimpleWaveOpenCL::evolved_group!"); + } + + 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, + "SimpleWaveOpenCL::phi", phi_bound); + if (ierr < 0) + CCTK_WARN(0, "Failed to register phi_bound BC for SimpleWaveOpenCL::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, + "SimpleWaveOpenCL::pi", pi_bound); + if (ierr < 0) + CCTK_WARN(0, "Failed to register pi_bound BC for SimpleWaveOpenCL::pi!"); + } + + if (CCTK_EQUALS(evolved_group_bound, "radiative")) + { + /* select radiation boundary condition */ + static CCTK_INT handle_evolved_group_bound = -1; + if (handle_evolved_group_bound < 0) handle_evolved_group_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_evolved_group_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_evolved_group_bound , evolved_group_bound_limit, "LIMIT") < 0) + CCTK_WARN(0, "could not set LIMIT value in table!"); + if (Util_TableSetReal(handle_evolved_group_bound ,evolved_group_bound_speed, "SPEED") < 0) + CCTK_WARN(0, "could not set SPEED value in table!"); + + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_evolved_group_bound, + "SimpleWaveOpenCL::evolved_group", "Radiation"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Radiation BC for SimpleWaveOpenCL::evolved_group!"); + + } + + 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, + "SimpleWaveOpenCL::phi", "Radiation"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Radiation BC for SimpleWaveOpenCL::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, + "SimpleWaveOpenCL::pi", "Radiation"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Radiation BC for SimpleWaveOpenCL::pi!"); + + } + + if (CCTK_EQUALS(evolved_group_bound, "scalar")) + { + /* select scalar boundary condition */ + static CCTK_INT handle_evolved_group_bound = -1; + if (handle_evolved_group_bound < 0) handle_evolved_group_bound = Util_TableCreate(UTIL_TABLE_FLAGS_CASE_INSENSITIVE); + if (handle_evolved_group_bound < 0) CCTK_WARN(0, "could not create table!"); + if (Util_TableSetReal(handle_evolved_group_bound ,evolved_group_bound_scalar, "SCALAR") < 0) + CCTK_WARN(0, "could not set SCALAR value in table!"); + + ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, 1, handle_evolved_group_bound, + "SimpleWaveOpenCL::evolved_group", "scalar"); + + if (ierr < 0) + CCTK_WARN(0, "Failed to register Scalar BC for SimpleWaveOpenCL::evolved_group!"); + + } + + 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, + "SimpleWaveOpenCL::phi", "scalar"); + + if (ierr < 0) + CCTK_WARN(0, "Error in registering Scalar BC for SimpleWaveOpenCL::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, + "SimpleWaveOpenCL::pi", "scalar"); + + if (ierr < 0) + CCTK_WARN(0, "Error in registering Scalar BC for SimpleWaveOpenCL::pi!"); + + } + return; +} + + + +/* template for entries in parameter file: +#$bound$#SimpleWaveOpenCL::evolved_group_bound = "skip" +#$bound$#SimpleWaveOpenCL::evolved_group_bound_speed = 1.0 +#$bound$#SimpleWaveOpenCL::evolved_group_bound_limit = 0.0 +#$bound$#SimpleWaveOpenCL::evolved_group_bound_scalar = 0.0 + +#$bound$#SimpleWaveOpenCL::phi_bound = "skip" +#$bound$#SimpleWaveOpenCL::phi_bound_speed = 1.0 +#$bound$#SimpleWaveOpenCL::phi_bound_limit = 0.0 +#$bound$#SimpleWaveOpenCL::phi_bound_scalar = 0.0 + +#$bound$#SimpleWaveOpenCL::pi_bound = "skip" +#$bound$#SimpleWaveOpenCL::pi_bound_speed = 1.0 +#$bound$#SimpleWaveOpenCL::pi_bound_limit = 0.0 +#$bound$#SimpleWaveOpenCL::pi_bound_scalar = 0.0 + +*/ + diff --git a/Examples/SimpleWaveOpenCL/src/Differencing.h b/Examples/SimpleWaveOpenCL/src/Differencing.h new file mode 100644 index 0000000..476bbe1 --- /dev/null +++ b/Examples/SimpleWaveOpenCL/src/Differencing.h @@ -0,0 +1,75 @@ +static char const * const differencing = +"#ifndef KRANC_DIFF_FUNCTIONS\n" +"# define PDstandard2nd1(u) ((-KRANC_GFOFFSET3D(u,-1,0,0) + KRANC_GFOFFSET3D(u,1,0,0))*p1o2dx)\n" +"#else\n" +"# define PDstandard2nd1(u) (PDstandard2nd1_impl(u,p1o2dx,cdj,cdk))\n" +"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;\n" +"static CCTK_REAL PDstandard2nd1_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dx, ptrdiff_t const cdj, ptrdiff_t const cdk)\n" +"{\n" +" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n" +" return (-KRANC_GFOFFSET3D(u,-1,0,0) + KRANC_GFOFFSET3D(u,1,0,0))*p1o2dx;\n" +"}\n" +"#endif\n" +"\n" +"#ifndef KRANC_DIFF_FUNCTIONS\n" +"# define PDstandard2nd2(u) ((-KRANC_GFOFFSET3D(u,0,-1,0) + KRANC_GFOFFSET3D(u,0,1,0))*p1o2dy)\n" +"#else\n" +"# define PDstandard2nd2(u) (PDstandard2nd2_impl(u,p1o2dy,cdj,cdk))\n" +"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;\n" +"static CCTK_REAL PDstandard2nd2_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dy, ptrdiff_t const cdj, ptrdiff_t const cdk)\n" +"{\n" +" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n" +" return (-KRANC_GFOFFSET3D(u,0,-1,0) + KRANC_GFOFFSET3D(u,0,1,0))*p1o2dy;\n" +"}\n" +"#endif\n" +"\n" +"#ifndef KRANC_DIFF_FUNCTIONS\n" +"# define PDstandard2nd3(u) ((-KRANC_GFOFFSET3D(u,0,0,-1) + KRANC_GFOFFSET3D(u,0,0,1))*p1o2dz)\n" +"#else\n" +"# define PDstandard2nd3(u) (PDstandard2nd3_impl(u,p1o2dz,cdj,cdk))\n" +"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;\n" +"static CCTK_REAL PDstandard2nd3_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1o2dz, ptrdiff_t const cdj, ptrdiff_t const cdk)\n" +"{\n" +" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n" +" return (-KRANC_GFOFFSET3D(u,0,0,-1) + KRANC_GFOFFSET3D(u,0,0,1))*p1o2dz;\n" +"}\n" +"#endif\n" +"\n" +"#ifndef KRANC_DIFF_FUNCTIONS\n" +"# define PDstandard2nd11(u) ((-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,-1,0,0) + KRANC_GFOFFSET3D(u,1,0,0))*p1odx2)\n" +"#else\n" +"# define PDstandard2nd11(u) (PDstandard2nd11_impl(u,p1odx2,cdj,cdk))\n" +"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;\n" +"static CCTK_REAL PDstandard2nd11_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1odx2, ptrdiff_t const cdj, ptrdiff_t const cdk)\n" +"{\n" +" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n" +" return (-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,-1,0,0) + KRANC_GFOFFSET3D(u,1,0,0))*p1odx2;\n" +"}\n" +"#endif\n" +"\n" +"#ifndef KRANC_DIFF_FUNCTIONS\n" +"# define PDstandard2nd22(u) ((-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,0,-1,0) + KRANC_GFOFFSET3D(u,0,1,0))*p1ody2)\n" +"#else\n" +"# define PDstandard2nd22(u) (PDstandard2nd22_impl(u,p1ody2,cdj,cdk))\n" +"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;\n" +"static CCTK_REAL PDstandard2nd22_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1ody2, ptrdiff_t const cdj, ptrdiff_t const cdk)\n" +"{\n" +" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n" +" return (-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,0,-1,0) + KRANC_GFOFFSET3D(u,0,1,0))*p1ody2;\n" +"}\n" +"#endif\n" +"\n" +"#ifndef KRANC_DIFF_FUNCTIONS\n" +"# define PDstandard2nd33(u) ((-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,0,0,-1) + KRANC_GFOFFSET3D(u,0,0,1))*p1odz2)\n" +"#else\n" +"# define PDstandard2nd33(u) (PDstandard2nd33_impl(u,p1odz2,cdj,cdk))\n" +"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;\n" +"static CCTK_REAL PDstandard2nd33_impl(CCTK_REAL const* restrict const u, CCTK_REAL const p1odz2, ptrdiff_t const cdj, ptrdiff_t const cdk)\n" +"{\n" +" ptrdiff_t const cdi=sizeof(CCTK_REAL);\n" +" return (-2*KRANC_GFOFFSET3D(u,0,0,0) + KRANC_GFOFFSET3D(u,0,0,-1) + KRANC_GFOFFSET3D(u,0,0,1))*p1odz2;\n" +"}\n" +"#endif\n" +"\n" +"" +; diff --git a/Examples/SimpleWaveOpenCL/src/RegisterMoL.cc b/Examples/SimpleWaveOpenCL/src/RegisterMoL.cc new file mode 100644 index 0000000..7ad8024 --- /dev/null +++ b/Examples/SimpleWaveOpenCL/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 SimpleWaveOpenCL_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("SimpleWaveOpenCL::phi"), CCTK_VarIndex("SimpleWaveOpenCL::phirhs")); + ierr += MoLRegisterEvolved(CCTK_VarIndex("SimpleWaveOpenCL::pi"), CCTK_VarIndex("SimpleWaveOpenCL::pirhs")); + + /* Register all the evolved Array functions with MoL */ + return; +} diff --git a/Examples/SimpleWaveOpenCL/src/RegisterSymmetries.cc b/Examples/SimpleWaveOpenCL/src/RegisterSymmetries.cc new file mode 100644 index 0000000..795d652 --- /dev/null +++ b/Examples/SimpleWaveOpenCL/src/RegisterSymmetries.cc @@ -0,0 +1,29 @@ +/* File produced by Kranc */ + +#include "cctk.h" +#include "cctk_Arguments.h" +#include "cctk_Parameters.h" +#include "Symmetry.h" + +extern "C" void SimpleWaveOpenCL_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, "SimpleWaveOpenCL::phi"); + + sym[0] = 1; + sym[1] = 1; + sym[2] = 1; + SetCartSymVN(cctkGH, sym, "SimpleWaveOpenCL::pi"); + +} diff --git a/Examples/SimpleWaveOpenCL/src/Startup.cc b/Examples/SimpleWaveOpenCL/src/Startup.cc new file mode 100644 index 0000000..10e5dbf --- /dev/null +++ b/Examples/SimpleWaveOpenCL/src/Startup.cc @@ -0,0 +1,10 @@ +/* File produced by Kranc */ + +#include "cctk.h" + +extern "C" int SimpleWaveOpenCL_Startup(void) +{ + const char * banner = "SimpleWaveOpenCL"; + CCTK_RegisterBanner(banner); + return 0; +} diff --git a/Examples/SimpleWaveOpenCL/src/calc_rhs.cc b/Examples/SimpleWaveOpenCL/src/calc_rhs.cc new file mode 100644 index 0000000..f463547 --- /dev/null +++ b/Examples/SimpleWaveOpenCL/src/calc_rhs.cc @@ -0,0 +1,159 @@ +/* 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" + +/* 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 */, "SimpleWaveOpenCL::evolved_grouprhs","flat"); + if (ierr < 0) + CCTK_WARN(1, "Failed to register flat BC for SimpleWaveOpenCL::evolved_grouprhs."); + 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; + + 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 const dx = ToReal(CCTK_DELTA_SPACE(0));\n" + "CCTK_REAL const dy = ToReal(CCTK_DELTA_SPACE(1));\n" + "CCTK_REAL const dz = ToReal(CCTK_DELTA_SPACE(2));\n" + "CCTK_REAL const dt = ToReal(CCTK_DELTA_TIME);\n" + "CCTK_REAL const t = ToReal(cctk_time);\n" + "CCTK_REAL const dxi = INV(dx);\n" + "CCTK_REAL const dyi = INV(dy);\n" + "CCTK_REAL const dzi = INV(dz);\n" + "CCTK_REAL const khalf = 0.5;\n" + "CCTK_REAL const kthird = 1/3.0;\n" + "CCTK_REAL const ktwothird = 2.0/3.0;\n" + "CCTK_REAL const kfourthird = 4.0/3.0;\n" + "CCTK_REAL const keightthird = 8.0/3.0;\n" + "CCTK_REAL const hdxi = 0.5 * dxi;\n" + "CCTK_REAL const hdyi = 0.5 * dyi;\n" + "CCTK_REAL const hdzi = 0.5 * dzi;\n" + "\n" + "/* Initialize predefined quantities */\n" + "CCTK_REAL const p1o2dx = 0.5*INV(dx);\n" + "CCTK_REAL const p1o2dy = 0.5*INV(dy);\n" + "CCTK_REAL const p1o2dz = 0.5*INV(dz);\n" + "CCTK_REAL const p1odx2 = INV(SQR(dx));\n" + "CCTK_REAL const p1ody2 = INV(SQR(dy));\n" + "CCTK_REAL const p1odz2 = INV(SQR(dz));\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" + "CCTK_LOOP3 (calc_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" + "{\n" + " ptrdiff_t const index = di*i + dj*j + dk*k;\n" + " \n" + " /* Assign local copies of grid functions */\n" + " \n" + " CCTK_REAL phiL = phi[index];\n" + " CCTK_REAL piL = pi[index];\n" + " \n" + " \n" + " /* Include user supplied include files */\n" + " \n" + " /* Precompute derivatives */\n" + " CCTK_REAL const PDstandard2nd11phi = PDstandard2nd11(&phi[index]);\n" + " CCTK_REAL const PDstandard2nd22phi = PDstandard2nd22(&phi[index]);\n" + " CCTK_REAL const PDstandard2nd33phi = PDstandard2nd33(&phi[index]);\n" + " \n" + " /* Calculate temporaries and grid functions */\n" + " CCTK_REAL phirhsL = piL;\n" + " \n" + " CCTK_REAL pirhsL = PDstandard2nd11phi + PDstandard2nd22phi + \n" + " PDstandard2nd33phi;\n" + " \n" + " /* Copy local copies back to grid functions */\n" + " vec_store_nta_partial(phirhs[index],phirhsL);\n" + " vec_store_nta_partial(pirhs[index],pirhsL);\n" + "}\n" + "CCTK_ENDLOOP3 (calc_rhs);\n" + "" + ; + + char const * const groups[] = {"SimpleWaveOpenCL::evolved_group","SimpleWaveOpenCL::evolved_grouprhs",NULL}; + + static struct OpenCLKernel * kernel = NULL; + char const * const sources[] = {differencing, source, NULL}; + OpenCLRunTime_CallKernel (cctkGH, CCTK_THORNSTRING, "calc_rhs", + sources, groups, NULL, NULL, NULL, -1, + imin, imax, &kernel); + +} + +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 *groups[] = {"SimpleWaveOpenCL::evolved_group","SimpleWaveOpenCL::evolved_grouprhs"}; + GenericFD_AssertGroupStorage(cctkGH, "calc_rhs", 2, 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/SimpleWaveOpenCL/src/initial_sine.cc b/Examples/SimpleWaveOpenCL/src/initial_sine.cc new file mode 100644 index 0000000..fa553dd --- /dev/null +++ b/Examples/SimpleWaveOpenCL/src/initial_sine.cc @@ -0,0 +1,141 @@ +/* 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" + +/* 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_sine_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 const dx = ToReal(CCTK_DELTA_SPACE(0));\n" + "CCTK_REAL const dy = ToReal(CCTK_DELTA_SPACE(1));\n" + "CCTK_REAL const dz = ToReal(CCTK_DELTA_SPACE(2));\n" + "CCTK_REAL const dt = ToReal(CCTK_DELTA_TIME);\n" + "CCTK_REAL const t = ToReal(cctk_time);\n" + "CCTK_REAL const dxi = INV(dx);\n" + "CCTK_REAL const dyi = INV(dy);\n" + "CCTK_REAL const dzi = INV(dz);\n" + "CCTK_REAL const khalf = 0.5;\n" + "CCTK_REAL const kthird = 1/3.0;\n" + "CCTK_REAL const ktwothird = 2.0/3.0;\n" + "CCTK_REAL const kfourthird = 4.0/3.0;\n" + "CCTK_REAL const keightthird = 8.0/3.0;\n" + "CCTK_REAL const hdxi = 0.5 * dxi;\n" + "CCTK_REAL const hdyi = 0.5 * dyi;\n" + "CCTK_REAL const hdzi = 0.5 * dzi;\n" + "\n" + "/* Initialize predefined quantities */\n" + "CCTK_REAL const p1o2dx = 0.5*INV(dx);\n" + "CCTK_REAL const p1o2dy = 0.5*INV(dy);\n" + "CCTK_REAL const p1o2dz = 0.5*INV(dz);\n" + "CCTK_REAL const p1odx2 = INV(SQR(dx));\n" + "CCTK_REAL const p1ody2 = INV(SQR(dy));\n" + "CCTK_REAL const p1odz2 = INV(SQR(dz));\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" + "CCTK_LOOP3 (initial_sine,\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" + "{\n" + " ptrdiff_t const index = di*i + dj*j + dk*k;\n" + " \n" + " /* Assign local copies of grid functions */\n" + " \n" + " CCTK_REAL xL = x[index];\n" + " \n" + " \n" + " /* Include user supplied include files */\n" + " \n" + " /* Precompute derivatives */\n" + " \n" + " /* Calculate temporaries and grid functions */\n" + " CCTK_REAL phiL = Sin(2*Pi*(xL - t));\n" + " \n" + " CCTK_REAL piL = -2*Pi*Cos(2*Pi*(xL - t));\n" + " \n" + " /* Copy local copies back to grid functions */\n" + " vec_store_nta_partial(phi[index],phiL);\n" + " vec_store_nta_partial(pi[index],piL);\n" + "}\n" + "CCTK_ENDLOOP3 (initial_sine);\n" + "" + ; + + char const * const groups[] = {"SimpleWaveOpenCL::evolved_group","grid::coordinates",NULL}; + + static struct OpenCLKernel * kernel = NULL; + char const * const sources[] = {differencing, source, NULL}; + OpenCLRunTime_CallKernel (cctkGH, CCTK_THORNSTRING, "initial_sine", + sources, groups, NULL, NULL, NULL, -1, + imin, imax, &kernel); + +} + +extern "C" void initial_sine(CCTK_ARGUMENTS) +{ + DECLARE_CCTK_ARGUMENTS; + DECLARE_CCTK_PARAMETERS; + + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Entering initial_sine_Body"); + } + + if (cctk_iteration % initial_sine_calc_every != initial_sine_calc_offset) + { + return; + } + + const char *groups[] = {"SimpleWaveOpenCL::evolved_group","grid::coordinates"}; + GenericFD_AssertGroupStorage(cctkGH, "initial_sine", 2, groups); + + + GenericFD_LoopOverEverything(cctkGH, &initial_sine_Body); + + if (verbose > 1) + { + CCTK_VInfo(CCTK_THORNSTRING,"Leaving initial_sine_Body"); + } +} diff --git a/Examples/SimpleWaveOpenCL/src/make.code.defn b/Examples/SimpleWaveOpenCL/src/make.code.defn new file mode 100644 index 0000000..535f5d0 --- /dev/null +++ b/Examples/SimpleWaveOpenCL/src/make.code.defn @@ -0,0 +1,3 @@ +# File produced by Kranc + +SRCS = Startup.cc RegisterMoL.cc RegisterSymmetries.cc initial_sine.cc calc_rhs.cc Boundaries.cc -- cgit v1.2.3