From 280545da21ad246f73e815c09f43b2d08804ebca Mon Sep 17 00:00:00 2001 From: Ian Hinder Date: Thu, 5 Sep 2013 22:58:00 +0200 Subject: Regenerate McLachlan (Kranc commit caf12ddfb11ccd9a7bf822cc8ec0d7691cded0aa) --- ML_WaveToy_CL/param.ccl | 1 + ML_WaveToy_CL/schedule.ccl | 81 ++++++-------------- ML_WaveToy_CL/src/Boundaries.cc | 18 ++--- ML_WaveToy_CL/src/Differencing.h | 122 +++++++++++++++--------------- ML_WaveToy_CL/src/RegisterMoL.cc | 2 +- ML_WaveToy_CL/src/Startup.cc | 2 +- ML_WaveToy_CL/src/WT_CL_Dirichlet.cc | 95 ++++++++++++----------- ML_WaveToy_CL/src/WT_CL_Energy.cc | 97 +++++++++++++----------- ML_WaveToy_CL/src/WT_CL_EnergyBoundary.cc | 93 +++++++++++++---------- ML_WaveToy_CL/src/WT_CL_Gaussian.cc | 89 ++++++++++++---------- ML_WaveToy_CL/src/WT_CL_RHS.cc | 97 +++++++++++++----------- ML_WaveToy_CL/src/WT_CL_Standing.cc | 91 ++++++++++++---------- 12 files changed, 404 insertions(+), 384 deletions(-) (limited to 'ML_WaveToy_CL') diff --git a/ML_WaveToy_CL/param.ccl b/ML_WaveToy_CL/param.ccl index 15d28f0..616b192 100644 --- a/ML_WaveToy_CL/param.ccl +++ b/ML_WaveToy_CL/param.ccl @@ -3,6 +3,7 @@ shares: GenericFD +USES CCTK_INT assume_stress_energy_state shares: MethodOfLines diff --git a/ML_WaveToy_CL/schedule.ccl b/ML_WaveToy_CL/schedule.ccl index 4f466bd..03ca58f 100644 --- a/ML_WaveToy_CL/schedule.ccl +++ b/ML_WaveToy_CL/schedule.ccl @@ -1,46 +1,15 @@ # File produced by Kranc -if (other_timelevels == 1) -{ - STORAGE: WT_eps[1] -} +STORAGE: WT_eps[other_timelevels] -if (timelevels == 1) -{ - STORAGE: WT_rho[1] -} -if (timelevels == 2) -{ - STORAGE: WT_rho[2] -} +STORAGE: WT_rho[timelevels] -if (timelevels == 1) -{ - STORAGE: WT_u[1] -} -if (timelevels == 2) -{ - STORAGE: WT_u[2] -} +STORAGE: WT_u[timelevels] -if (rhs_timelevels == 1) -{ - STORAGE: WT_rhorhs[1] -} -if (rhs_timelevels == 2) -{ - STORAGE: WT_rhorhs[2] -} +STORAGE: WT_rhorhs[rhs_timelevels] -if (rhs_timelevels == 1) -{ - STORAGE: WT_urhs[1] -} -if (rhs_timelevels == 2) -{ - STORAGE: WT_urhs[2] -} +STORAGE: WT_urhs[rhs_timelevels] schedule ML_WaveToy_CL_Startup at STARTUP { @@ -61,9 +30,9 @@ if (CCTK_EQUALS(initial_data, "Gaussian")) { LANG: C TAGS: Device=1 - READS: grid::r(Everywhere) - WRITES: ML_WaveToy_CL::rho(Everywhere) - WRITES: ML_WaveToy_CL::u(Everywhere) + READS: grid::r(Everywhere) + WRITES: ML_WaveToy_CL::rho(Everywhere) + WRITES: ML_WaveToy_CL::u(Everywhere) } "WT_CL_Gaussian" } @@ -74,11 +43,11 @@ if (CCTK_EQUALS(initial_data, "Standing")) { LANG: C TAGS: Device=1 - READS: grid::x(Everywhere) - READS: grid::y(Everywhere) - READS: grid::z(Everywhere) - WRITES: ML_WaveToy_CL::rho(Everywhere) - WRITES: ML_WaveToy_CL::u(Everywhere) + READS: grid::x(Everywhere) + READS: grid::y(Everywhere) + READS: grid::z(Everywhere) + WRITES: ML_WaveToy_CL::rho(Everywhere) + WRITES: ML_WaveToy_CL::u(Everywhere) } "WT_CL_Standing" } @@ -86,18 +55,18 @@ schedule WT_CL_RHS IN MoL_CalcRHS { LANG: C TAGS: Device=1 - READS: ML_WaveToy_CL::rho(Everywhere) - READS: ML_WaveToy_CL::u(Everywhere) - WRITES: ML_WaveToy_CL::rhorhs(Interior) - WRITES: ML_WaveToy_CL::urhs(Interior) + READS: ML_WaveToy_CL::rho(Everywhere) + READS: ML_WaveToy_CL::u(Everywhere) + WRITES: ML_WaveToy_CL::rhorhs(Interior) + WRITES: ML_WaveToy_CL::urhs(Interior) } "WT_CL_RHS" schedule WT_CL_Dirichlet IN MoL_CalcRHS { LANG: C TAGS: Device=1 - WRITES: ML_WaveToy_CL::rhorhs(Boundary) - WRITES: ML_WaveToy_CL::urhs(Boundary) + WRITES: ML_WaveToy_CL::rhorhs(Boundary) + WRITES: ML_WaveToy_CL::urhs(Boundary) } "WT_CL_Dirichlet" schedule WT_CL_Dirichlet AT analysis @@ -106,8 +75,8 @@ schedule WT_CL_Dirichlet AT analysis SYNC: WT_rhorhs SYNC: WT_urhs TAGS: Device=1 - WRITES: ML_WaveToy_CL::rhorhs(Boundary) - WRITES: ML_WaveToy_CL::urhs(Boundary) + WRITES: ML_WaveToy_CL::rhorhs(Boundary) + WRITES: ML_WaveToy_CL::urhs(Boundary) } "WT_CL_Dirichlet" schedule WT_CL_Energy AT analysis @@ -115,9 +84,9 @@ schedule WT_CL_Energy AT analysis LANG: C SYNC: WT_eps TAGS: Device=1 - READS: ML_WaveToy_CL::rho(Everywhere) - READS: ML_WaveToy_CL::u(Everywhere) - WRITES: ML_WaveToy_CL::eps(Interior) + READS: ML_WaveToy_CL::rho(Everywhere) + READS: ML_WaveToy_CL::u(Everywhere) + WRITES: ML_WaveToy_CL::eps(Interior) } "WT_CL_Energy" schedule WT_CL_EnergyBoundary AT analysis @@ -125,7 +94,7 @@ schedule WT_CL_EnergyBoundary AT analysis LANG: C SYNC: WT_eps TAGS: Device=1 - WRITES: ML_WaveToy_CL::eps(Boundary) + WRITES: ML_WaveToy_CL::eps(Boundary) } "WT_CL_EnergyBoundary" schedule ML_WaveToy_CL_SelectBoundConds in MoL_PostStep diff --git a/ML_WaveToy_CL/src/Boundaries.cc b/ML_WaveToy_CL/src/Boundaries.cc index 274144c..6f346a7 100644 --- a/ML_WaveToy_CL/src/Boundaries.cc +++ b/ML_WaveToy_CL/src/Boundaries.cc @@ -30,7 +30,7 @@ extern "C" void ML_WaveToy_CL_SelectBoundConds(CCTK_ARGUMENTS) DECLARE_CCTK_ARGUMENTS; DECLARE_CCTK_PARAMETERS; - CCTK_INT ierr CCTK_ATTRIBUTE_UNUSED = 0; + CCTK_INT ierr CCTK_ATTRIBUTE_UNUSED = 0; if (CCTK_EQUALS(WT_rho_bound, "none" ) || CCTK_EQUALS(WT_rho_bound, "static") || @@ -79,7 +79,7 @@ extern "C" void ML_WaveToy_CL_SelectBoundConds(CCTK_ARGUMENTS) if (CCTK_EQUALS(WT_rho_bound, "radiative")) { /* select radiation boundary condition */ - static CCTK_INT handle_WT_rho_bound CCTK_ATTRIBUTE_UNUSED = -1; + static CCTK_INT handle_WT_rho_bound CCTK_ATTRIBUTE_UNUSED = -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) @@ -98,7 +98,7 @@ extern "C" void ML_WaveToy_CL_SelectBoundConds(CCTK_ARGUMENTS) if (CCTK_EQUALS(WT_u_bound, "radiative")) { /* select radiation boundary condition */ - static CCTK_INT handle_WT_u_bound CCTK_ATTRIBUTE_UNUSED = -1; + static CCTK_INT handle_WT_u_bound CCTK_ATTRIBUTE_UNUSED = -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) @@ -117,7 +117,7 @@ extern "C" void ML_WaveToy_CL_SelectBoundConds(CCTK_ARGUMENTS) if (CCTK_EQUALS(rho_bound, "radiative")) { /* select radiation boundary condition */ - static CCTK_INT handle_rho_bound CCTK_ATTRIBUTE_UNUSED = -1; + static CCTK_INT handle_rho_bound CCTK_ATTRIBUTE_UNUSED = -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) @@ -136,7 +136,7 @@ extern "C" void ML_WaveToy_CL_SelectBoundConds(CCTK_ARGUMENTS) if (CCTK_EQUALS(u_bound, "radiative")) { /* select radiation boundary condition */ - static CCTK_INT handle_u_bound CCTK_ATTRIBUTE_UNUSED = -1; + static CCTK_INT handle_u_bound CCTK_ATTRIBUTE_UNUSED = -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) @@ -155,7 +155,7 @@ extern "C" void ML_WaveToy_CL_SelectBoundConds(CCTK_ARGUMENTS) if (CCTK_EQUALS(WT_rho_bound, "scalar")) { /* select scalar boundary condition */ - static CCTK_INT handle_WT_rho_bound CCTK_ATTRIBUTE_UNUSED = -1; + static CCTK_INT handle_WT_rho_bound CCTK_ATTRIBUTE_UNUSED = -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) @@ -172,7 +172,7 @@ extern "C" void ML_WaveToy_CL_SelectBoundConds(CCTK_ARGUMENTS) if (CCTK_EQUALS(WT_u_bound, "scalar")) { /* select scalar boundary condition */ - static CCTK_INT handle_WT_u_bound CCTK_ATTRIBUTE_UNUSED = -1; + static CCTK_INT handle_WT_u_bound CCTK_ATTRIBUTE_UNUSED = -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) @@ -189,7 +189,7 @@ extern "C" void ML_WaveToy_CL_SelectBoundConds(CCTK_ARGUMENTS) if (CCTK_EQUALS(rho_bound, "scalar")) { /* select scalar boundary condition */ - static CCTK_INT handle_rho_bound CCTK_ATTRIBUTE_UNUSED = -1; + static CCTK_INT handle_rho_bound CCTK_ATTRIBUTE_UNUSED = -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) @@ -206,7 +206,7 @@ extern "C" void ML_WaveToy_CL_SelectBoundConds(CCTK_ARGUMENTS) if (CCTK_EQUALS(u_bound, "scalar")) { /* select scalar boundary condition */ - static CCTK_INT handle_u_bound CCTK_ATTRIBUTE_UNUSED = -1; + static CCTK_INT handle_u_bound CCTK_ATTRIBUTE_UNUSED = -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) diff --git a/ML_WaveToy_CL/src/Differencing.h b/ML_WaveToy_CL/src/Differencing.h index 9a5e0c6..e992787 100644 --- a/ML_WaveToy_CL/src/Differencing.h +++ b/ML_WaveToy_CL/src/Differencing.h @@ -1,145 +1,145 @@ -static char const *const differencing = +static const char* 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.),kmadd(KRANC_GFOFFSET3D(u,2,0,0),ToReal(-1.),kmul(KRANC_GFOFFSET3D(u,1,0,0),ToReal(8.)))))))\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" +"static CCTK_REAL_VEC PDstandardNth1_impl(const CCTK_REAL* restrict const u, const CCTK_REAL_VEC p1o12dx, const ptrdiff_t cdj, const ptrdiff_t cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n" +"static CCTK_REAL_VEC PDstandardNth1_impl(const CCTK_REAL* restrict const u, const CCTK_REAL_VEC p1o12dx, const ptrdiff_t cdj, const ptrdiff_t cdk)\n" "{\n" -" ptrdiff_t const cdi CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL);\n" -" return kmul(p1o12dx,kadd(KRANC_GFOFFSET3D(u,-2,0,0),kmadd(KRANC_GFOFFSET3D(u,-1,0,0),ToReal(-8.),kmadd(KRANC_GFOFFSET3D(u,2,0,0),ToReal(-1.),kmul(KRANC_GFOFFSET3D(u,1,0,0),ToReal(8.))))));\n" +" const ptrdiff_t cdi CCTK_ATTRIBUTE_UNUSED = 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.),kmadd(KRANC_GFOFFSET3D(u,0,2,0),ToReal(-1.),kmul(KRANC_GFOFFSET3D(u,0,1,0),ToReal(8.)))))))\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" +"static CCTK_REAL_VEC PDstandardNth2_impl(const CCTK_REAL* restrict const u, const CCTK_REAL_VEC p1o12dy, const ptrdiff_t cdj, const ptrdiff_t cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n" +"static CCTK_REAL_VEC PDstandardNth2_impl(const CCTK_REAL* restrict const u, const CCTK_REAL_VEC p1o12dy, const ptrdiff_t cdj, const ptrdiff_t cdk)\n" "{\n" -" ptrdiff_t const cdi CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL);\n" -" return kmul(p1o12dy,kadd(KRANC_GFOFFSET3D(u,0,-2,0),kmadd(KRANC_GFOFFSET3D(u,0,-1,0),ToReal(-8.),kmadd(KRANC_GFOFFSET3D(u,0,2,0),ToReal(-1.),kmul(KRANC_GFOFFSET3D(u,0,1,0),ToReal(8.))))));\n" +" const ptrdiff_t cdi CCTK_ATTRIBUTE_UNUSED = 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.),kmadd(KRANC_GFOFFSET3D(u,0,0,2),ToReal(-1.),kmul(KRANC_GFOFFSET3D(u,0,0,1),ToReal(8.)))))))\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" +"static CCTK_REAL_VEC PDstandardNth3_impl(const CCTK_REAL* restrict const u, const CCTK_REAL_VEC p1o12dz, const ptrdiff_t cdj, const ptrdiff_t cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n" +"static CCTK_REAL_VEC PDstandardNth3_impl(const CCTK_REAL* restrict const u, const CCTK_REAL_VEC p1o12dz, const ptrdiff_t cdj, const ptrdiff_t cdk)\n" "{\n" -" ptrdiff_t const cdi CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL);\n" -" return kmul(p1o12dz,kadd(KRANC_GFOFFSET3D(u,0,0,-2),kmadd(KRANC_GFOFFSET3D(u,0,0,-1),ToReal(-8.),kmadd(KRANC_GFOFFSET3D(u,0,0,2),ToReal(-1.),kmul(KRANC_GFOFFSET3D(u,0,0,1),ToReal(8.))))));\n" +" const ptrdiff_t cdi CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL);\n" +" return PDstandardNth2_impl(u, p1o12dz, cdk, cdj);\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" +"# 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" +"static CCTK_REAL_VEC PDstandardNth11_impl(const CCTK_REAL* restrict const u, const CCTK_REAL_VEC pm1o12dx2, const ptrdiff_t cdj, const ptrdiff_t cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n" +"static CCTK_REAL_VEC PDstandardNth11_impl(const CCTK_REAL* restrict const u, const CCTK_REAL_VEC pm1o12dx2, const ptrdiff_t cdj, const ptrdiff_t cdk)\n" "{\n" -" ptrdiff_t const cdi CCTK_ATTRIBUTE_UNUSED = 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" +" const ptrdiff_t cdi CCTK_ATTRIBUTE_UNUSED = 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" +"# 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" +"static CCTK_REAL_VEC PDstandardNth22_impl(const CCTK_REAL* restrict const u, const CCTK_REAL_VEC pm1o12dy2, const ptrdiff_t cdj, const ptrdiff_t cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n" +"static CCTK_REAL_VEC PDstandardNth22_impl(const CCTK_REAL* restrict const u, const CCTK_REAL_VEC pm1o12dy2, const ptrdiff_t cdj, const ptrdiff_t cdk)\n" "{\n" -" ptrdiff_t const cdi CCTK_ATTRIBUTE_UNUSED = 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" +" const ptrdiff_t cdi CCTK_ATTRIBUTE_UNUSED = 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" +"# 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" +"static CCTK_REAL_VEC PDstandardNth33_impl(const CCTK_REAL* restrict const u, const CCTK_REAL_VEC pm1o12dz2, const ptrdiff_t cdj, const ptrdiff_t cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n" +"static CCTK_REAL_VEC PDstandardNth33_impl(const CCTK_REAL* restrict const u, const CCTK_REAL_VEC pm1o12dz2, const ptrdiff_t cdj, const ptrdiff_t cdk)\n" "{\n" -" ptrdiff_t const cdi CCTK_ATTRIBUTE_UNUSED = 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" +" const ptrdiff_t cdi CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL);\n" +" return PDstandardNth22_impl(u, pm1o12dz2, cdk, cdj);\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.),kmadd(kadd(KRANC_GFOFFSET3D(u,-2,2,0),KRANC_GFOFFSET3D(u,2,-2,0)),ToReal(-1.),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.))))))))))\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" +"static CCTK_REAL_VEC PDstandardNth12_impl(const CCTK_REAL* restrict const u, const CCTK_REAL_VEC p1o144dxdy, const ptrdiff_t cdj, const ptrdiff_t cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n" +"static CCTK_REAL_VEC PDstandardNth12_impl(const CCTK_REAL* restrict const u, const CCTK_REAL_VEC p1o144dxdy, const ptrdiff_t cdj, const ptrdiff_t cdk)\n" "{\n" -" ptrdiff_t const cdi CCTK_ATTRIBUTE_UNUSED = 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.),kmadd(kadd(KRANC_GFOFFSET3D(u,-2,2,0),KRANC_GFOFFSET3D(u,2,-2,0)),ToReal(-1.),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.)))))))));\n" +" const ptrdiff_t cdi CCTK_ATTRIBUTE_UNUSED = 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.),kmadd(kadd(KRANC_GFOFFSET3D(u,-2,0,2),KRANC_GFOFFSET3D(u,2,0,-2)),ToReal(-1.),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.))))))))))\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" +"static CCTK_REAL_VEC PDstandardNth13_impl(const CCTK_REAL* restrict const u, const CCTK_REAL_VEC p1o144dxdz, const ptrdiff_t cdj, const ptrdiff_t cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n" +"static CCTK_REAL_VEC PDstandardNth13_impl(const CCTK_REAL* restrict const u, const CCTK_REAL_VEC p1o144dxdz, const ptrdiff_t cdj, const ptrdiff_t cdk)\n" "{\n" -" ptrdiff_t const cdi CCTK_ATTRIBUTE_UNUSED = 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.),kmadd(kadd(KRANC_GFOFFSET3D(u,-2,0,2),KRANC_GFOFFSET3D(u,2,0,-2)),ToReal(-1.),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.)))))))));\n" +" const ptrdiff_t cdi CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL);\n" +" return PDstandardNth12_impl(u, p1o144dxdz, cdk, cdj);\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.),kmadd(kadd(KRANC_GFOFFSET3D(u,-2,2,0),KRANC_GFOFFSET3D(u,2,-2,0)),ToReal(-1.),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.))))))))))\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" +"static CCTK_REAL_VEC PDstandardNth21_impl(const CCTK_REAL* restrict const u, const CCTK_REAL_VEC p1o144dxdy, const ptrdiff_t cdj, const ptrdiff_t cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n" +"static CCTK_REAL_VEC PDstandardNth21_impl(const CCTK_REAL* restrict const u, const CCTK_REAL_VEC p1o144dxdy, const ptrdiff_t cdj, const ptrdiff_t cdk)\n" "{\n" -" ptrdiff_t const cdi CCTK_ATTRIBUTE_UNUSED = 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.),kmadd(kadd(KRANC_GFOFFSET3D(u,-2,2,0),KRANC_GFOFFSET3D(u,2,-2,0)),ToReal(-1.),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.)))))))));\n" +" const ptrdiff_t cdi CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL);\n" +" return PDstandardNth12_impl(u, p1o144dxdy, cdj, cdk);\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.),kmadd(kadd(KRANC_GFOFFSET3D(u,0,-2,2),KRANC_GFOFFSET3D(u,0,2,-2)),ToReal(-1.),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.))))))))))\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" +"static CCTK_REAL_VEC PDstandardNth23_impl(const CCTK_REAL* restrict const u, const CCTK_REAL_VEC p1o144dydz, const ptrdiff_t cdj, const ptrdiff_t cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n" +"static CCTK_REAL_VEC PDstandardNth23_impl(const CCTK_REAL* restrict const u, const CCTK_REAL_VEC p1o144dydz, const ptrdiff_t cdj, const ptrdiff_t cdk)\n" "{\n" -" ptrdiff_t const cdi CCTK_ATTRIBUTE_UNUSED = 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.),kmadd(kadd(KRANC_GFOFFSET3D(u,0,-2,2),KRANC_GFOFFSET3D(u,0,2,-2)),ToReal(-1.),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.)))))))));\n" +" const ptrdiff_t cdi CCTK_ATTRIBUTE_UNUSED = 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.),kmadd(kadd(KRANC_GFOFFSET3D(u,-2,0,2),KRANC_GFOFFSET3D(u,2,0,-2)),ToReal(-1.),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.))))))))))\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" +"static CCTK_REAL_VEC PDstandardNth31_impl(const CCTK_REAL* restrict const u, const CCTK_REAL_VEC p1o144dxdz, const ptrdiff_t cdj, const ptrdiff_t cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n" +"static CCTK_REAL_VEC PDstandardNth31_impl(const CCTK_REAL* restrict const u, const CCTK_REAL_VEC p1o144dxdz, const ptrdiff_t cdj, const ptrdiff_t cdk)\n" "{\n" -" ptrdiff_t const cdi CCTK_ATTRIBUTE_UNUSED = 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.),kmadd(kadd(KRANC_GFOFFSET3D(u,-2,0,2),KRANC_GFOFFSET3D(u,2,0,-2)),ToReal(-1.),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.)))))))));\n" +" const ptrdiff_t cdi CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL);\n" +" return PDstandardNth12_impl(u, p1o144dxdz, cdk, cdj);\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.),kmadd(kadd(KRANC_GFOFFSET3D(u,0,-2,2),KRANC_GFOFFSET3D(u,0,2,-2)),ToReal(-1.),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.))))))))))\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" +"static CCTK_REAL_VEC PDstandardNth32_impl(const CCTK_REAL* restrict const u, const CCTK_REAL_VEC p1o144dydz, const ptrdiff_t cdj, const ptrdiff_t cdk) CCTK_ATTRIBUTE_NOINLINE CCTK_ATTRIBUTE_UNUSED;\n" +"static CCTK_REAL_VEC PDstandardNth32_impl(const CCTK_REAL* restrict const u, const CCTK_REAL_VEC p1o144dydz, const ptrdiff_t cdj, const ptrdiff_t cdk)\n" "{\n" -" ptrdiff_t const cdi CCTK_ATTRIBUTE_UNUSED = 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.),kmadd(kadd(KRANC_GFOFFSET3D(u,0,-2,2),KRANC_GFOFFSET3D(u,0,2,-2)),ToReal(-1.),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.)))))))));\n" +" const ptrdiff_t cdi CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL);\n" +" return PDstandardNth23_impl(u, p1o144dydz, cdj, cdk);\n" "}\n" "#endif\n" "\n" diff --git a/ML_WaveToy_CL/src/RegisterMoL.cc b/ML_WaveToy_CL/src/RegisterMoL.cc index 472f9f7..c507ddc 100644 --- a/ML_WaveToy_CL/src/RegisterMoL.cc +++ b/ML_WaveToy_CL/src/RegisterMoL.cc @@ -9,7 +9,7 @@ extern "C" void ML_WaveToy_CL_RegisterVars(CCTK_ARGUMENTS) DECLARE_CCTK_ARGUMENTS; DECLARE_CCTK_PARAMETERS; - CCTK_INT ierr CCTK_ATTRIBUTE_UNUSED = 0; + CCTK_INT ierr CCTK_ATTRIBUTE_UNUSED = 0; /* Register all the evolved grid functions with MoL */ ierr += MoLRegisterEvolved(CCTK_VarIndex("ML_WaveToy_CL::rho"), CCTK_VarIndex("ML_WaveToy_CL::rhorhs")); diff --git a/ML_WaveToy_CL/src/Startup.cc b/ML_WaveToy_CL/src/Startup.cc index 5d1cd87..5a9e1cf 100644 --- a/ML_WaveToy_CL/src/Startup.cc +++ b/ML_WaveToy_CL/src/Startup.cc @@ -4,7 +4,7 @@ extern "C" int ML_WaveToy_CL_Startup(void) { - const char * banner CCTK_ATTRIBUTE_UNUSED = "ML_WaveToy_CL"; + const char* banner CCTK_ATTRIBUTE_UNUSED = "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 index 1974909..d83016a 100644 --- a/ML_WaveToy_CL/src/WT_CL_Dirichlet.cc +++ b/ML_WaveToy_CL/src/WT_CL_Dirichlet.cc @@ -33,7 +33,9 @@ extern "C" void WT_CL_Dirichlet_SelectBCs(CCTK_ARGUMENTS) DECLARE_CCTK_ARGUMENTS; DECLARE_CCTK_PARAMETERS; - CCTK_INT ierr CCTK_ATTRIBUTE_UNUSED = 0; + if (cctk_iteration % WT_CL_Dirichlet_calc_every != WT_CL_Dirichlet_calc_offset) + return; + CCTK_INT ierr CCTK_ATTRIBUTE_UNUSED = 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."); @@ -43,49 +45,49 @@ extern "C" void WT_CL_Dirichlet_SelectBCs(CCTK_ARGUMENTS) 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[]) +static void WT_CL_Dirichlet_Body(const cGH* restrict const cctkGH, const int dir, const int face, const CCTK_REAL normal[3], const CCTK_REAL tangentA[3], const CCTK_REAL tangentB[3], const int imin[3], const int imax[3], const int n_subblock_gfs, CCTK_REAL* restrict const subblock_gfs[]) { DECLARE_CCTK_ARGUMENTS; DECLARE_CCTK_PARAMETERS; - char const *const source = + const char* const source = "\n" "/* Include user-supplied include files */\n" "\n" "/* Initialise finite differencing variables */\n" - "ptrdiff_t /*const*/ di CCTK_ATTRIBUTE_UNUSED = 1;\n" - "ptrdiff_t /*const*/ dj CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" - "ptrdiff_t /*const*/ dk CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" - "ptrdiff_t /*const*/ cdi CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * di;\n" - "ptrdiff_t /*const*/ cdj CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dj;\n" - "ptrdiff_t /*const*/ cdk CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dk;\n" - "CCTK_REAL_VEC /*const*/ dx CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(0));\n" - "CCTK_REAL_VEC /*const*/ dy CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(1));\n" - "CCTK_REAL_VEC /*const*/ dz CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(2));\n" - "CCTK_REAL_VEC /*const*/ dt CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_TIME);\n" - "CCTK_REAL_VEC /*const*/ t CCTK_ATTRIBUTE_UNUSED = ToReal(cctk_time);\n" - "CCTK_REAL_VEC /*const*/ dxi CCTK_ATTRIBUTE_UNUSED = INV(dx);\n" - "CCTK_REAL_VEC /*const*/ dyi CCTK_ATTRIBUTE_UNUSED = INV(dy);\n" - "CCTK_REAL_VEC /*const*/ dzi CCTK_ATTRIBUTE_UNUSED = INV(dz);\n" - "CCTK_REAL_VEC /*const*/ khalf CCTK_ATTRIBUTE_UNUSED = ToReal(0.5);\n" - "CCTK_REAL_VEC /*const*/ kthird CCTK_ATTRIBUTE_UNUSED = ToReal(1.0/3.0);\n" - "CCTK_REAL_VEC /*const*/ ktwothird CCTK_ATTRIBUTE_UNUSED = ToReal(2.0/3.0);\n" - "CCTK_REAL_VEC /*const*/ kfourthird CCTK_ATTRIBUTE_UNUSED = ToReal(4.0/3.0);\n" - "CCTK_REAL_VEC /*const*/ keightthird CCTK_ATTRIBUTE_UNUSED = ToReal(8.0/3.0);\n" - "CCTK_REAL_VEC /*const*/ hdxi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dxi);\n" - "CCTK_REAL_VEC /*const*/ hdyi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dyi);\n" - "CCTK_REAL_VEC /*const*/ hdzi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dzi);\n" + "const ptrdiff_t di CCTK_ATTRIBUTE_UNUSED = 1;\n" + "const ptrdiff_t dj CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" + "const ptrdiff_t dk CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" + "const ptrdiff_t cdi CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * di;\n" + "const ptrdiff_t cdj CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dj;\n" + "const ptrdiff_t cdk CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dk;\n" + "const CCTK_REAL_VEC dx CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(0));\n" + "const CCTK_REAL_VEC dy CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(1));\n" + "const CCTK_REAL_VEC dz CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(2));\n" + "const CCTK_REAL_VEC dt CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_TIME);\n" + "const CCTK_REAL_VEC t CCTK_ATTRIBUTE_UNUSED = ToReal(cctk_time);\n" + "const CCTK_REAL_VEC dxi CCTK_ATTRIBUTE_UNUSED = INV(dx);\n" + "const CCTK_REAL_VEC dyi CCTK_ATTRIBUTE_UNUSED = INV(dy);\n" + "const CCTK_REAL_VEC dzi CCTK_ATTRIBUTE_UNUSED = INV(dz);\n" + "const CCTK_REAL_VEC khalf CCTK_ATTRIBUTE_UNUSED = ToReal(0.5);\n" + "const CCTK_REAL_VEC kthird CCTK_ATTRIBUTE_UNUSED = ToReal(1.0/3.0);\n" + "const CCTK_REAL_VEC ktwothird CCTK_ATTRIBUTE_UNUSED = ToReal(2.0/3.0);\n" + "const CCTK_REAL_VEC kfourthird CCTK_ATTRIBUTE_UNUSED = ToReal(4.0/3.0);\n" + "const CCTK_REAL_VEC keightthird CCTK_ATTRIBUTE_UNUSED = ToReal(8.0/3.0);\n" + "const CCTK_REAL_VEC hdxi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dxi);\n" + "const CCTK_REAL_VEC hdyi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dyi);\n" + "const CCTK_REAL_VEC hdzi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dzi);\n" "\n" "/* Initialize predefined quantities */\n" - "CCTK_REAL_VEC /*const*/ p1o12dx CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dx);\n" - "CCTK_REAL_VEC /*const*/ p1o12dy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dy);\n" - "CCTK_REAL_VEC /*const*/ p1o12dz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dz);\n" - "CCTK_REAL_VEC /*const*/ p1o144dxdy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dy,dx));\n" - "CCTK_REAL_VEC /*const*/ p1o144dxdz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dx));\n" - "CCTK_REAL_VEC /*const*/ p1o144dydz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dy));\n" - "CCTK_REAL_VEC /*const*/ pm1o12dx2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dx,dx));\n" - "CCTK_REAL_VEC /*const*/ pm1o12dy2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dy,dy));\n" - "CCTK_REAL_VEC /*const*/ pm1o12dz2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dz,dz));\n" + "const CCTK_REAL_VEC p1o12dx CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dx);\n" + "const CCTK_REAL_VEC p1o12dy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dy);\n" + "const CCTK_REAL_VEC p1o12dz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dz);\n" + "const CCTK_REAL_VEC p1o144dxdy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dy,dx));\n" + "const CCTK_REAL_VEC p1o144dxdz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dx));\n" + "const CCTK_REAL_VEC p1o144dydz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dy));\n" + "const CCTK_REAL_VEC pm1o12dx2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dx,dx));\n" + "const CCTK_REAL_VEC pm1o12dy2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dy,dy));\n" + "const CCTK_REAL_VEC pm1o12dz2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dz,dz));\n" "\n" "/* Assign local copies of arrays functions */\n" "\n" @@ -96,13 +98,20 @@ static void WT_CL_Dirichlet_Body(cGH const * restrict const cctkGH, int const di "/* Copy local copies back to grid functions */\n" "\n" "/* Loop over the grid points */\n" - "#pragma omp parallel\n" + "const int imin0=imin[0];\n" + "const int imin1=imin[1];\n" + "const int imin2=imin[2];\n" + "const int imax0=imax[0];\n" + "const int imax1=imax[1];\n" + "const int imax2=imax[2];\n" + "#pragma omp parallel // reduction(+: vec_iter_counter, vec_op_counter, vec_mem_counter)\n" "CCTK_LOOP3STR(WT_CL_Dirichlet,\n" - " i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2],\n" + " i,j,k, imin0,imin1,imin2, imax0,imax1,imax2,\n" " cctk_ash[0],cctk_ash[1],cctk_ash[2],\n" - " kimin,kimax, CCTK_REAL_VEC_SIZE)\n" + " vecimin,vecimax, CCTK_REAL_VEC_SIZE)\n" "{\n" - " ptrdiff_t /*const*/ index CCTK_ATTRIBUTE_UNUSED = di*i + dj*j + dk*k;\n" + " const ptrdiff_t index CCTK_ATTRIBUTE_UNUSED = di*i + dj*j + dk*k;\n" + " // vec_iter_counter+=CCTK_REAL_VEC_SIZE;\n" " \n" " /* Assign local copies of grid functions */\n" " \n" @@ -113,9 +122,9 @@ static void WT_CL_Dirichlet_Body(cGH const * restrict const cctkGH, int const di " /* Precompute derivatives */\n" " \n" " /* Calculate temporaries and grid functions */\n" - " CCTK_REAL_VEC CCTK_ATTRIBUTE_UNUSED urhsL = ToReal(0.);\n" + " CCTK_REAL_VEC CCTK_ATTRIBUTE_UNUSED urhsL = ToReal(0);\n" " \n" - " CCTK_REAL_VEC CCTK_ATTRIBUTE_UNUSED rhorhsL = ToReal(0.);\n" + " CCTK_REAL_VEC CCTK_ATTRIBUTE_UNUSED rhorhsL = ToReal(0);\n" " \n" " /* Copy local copies back to grid functions */\n" " vec_store_partial_prepare(i,lc_imin,lc_imax);\n" @@ -126,13 +135,13 @@ static void WT_CL_Dirichlet_Body(cGH const * restrict const cctkGH, int const di "" ; - char const *const groups[] = { + const char* 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}; + const char* const sources[] = {differencing, source, NULL}; OpenCLRunTime_CallKernel(cctkGH, CCTK_THORNSTRING, "WT_CL_Dirichlet", sources, groups, NULL, NULL, NULL, -1, imin, imax, &kernel); @@ -155,7 +164,7 @@ extern "C" void WT_CL_Dirichlet(CCTK_ARGUMENTS) return; } - const char *const groups[] = { + const char* const groups[] = { "ML_WaveToy_CL::WT_rhorhs", "ML_WaveToy_CL::WT_urhs"}; GenericFD_AssertGroupStorage(cctkGH, "WT_CL_Dirichlet", 2, groups); diff --git a/ML_WaveToy_CL/src/WT_CL_Energy.cc b/ML_WaveToy_CL/src/WT_CL_Energy.cc index 89b8524..4ed7d60 100644 --- a/ML_WaveToy_CL/src/WT_CL_Energy.cc +++ b/ML_WaveToy_CL/src/WT_CL_Energy.cc @@ -33,56 +33,58 @@ extern "C" void WT_CL_Energy_SelectBCs(CCTK_ARGUMENTS) DECLARE_CCTK_ARGUMENTS; DECLARE_CCTK_PARAMETERS; - CCTK_INT ierr CCTK_ATTRIBUTE_UNUSED = 0; + if (cctk_iteration % WT_CL_Energy_calc_every != WT_CL_Energy_calc_offset) + return; + CCTK_INT ierr CCTK_ATTRIBUTE_UNUSED = 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[]) +static void WT_CL_Energy_Body(const cGH* restrict const cctkGH, const int dir, const int face, const CCTK_REAL normal[3], const CCTK_REAL tangentA[3], const CCTK_REAL tangentB[3], const int imin[3], const int imax[3], const int n_subblock_gfs, CCTK_REAL* restrict const subblock_gfs[]) { DECLARE_CCTK_ARGUMENTS; DECLARE_CCTK_PARAMETERS; - char const *const source = + const char* const source = "\n" "/* Include user-supplied include files */\n" "\n" "/* Initialise finite differencing variables */\n" - "ptrdiff_t /*const*/ di CCTK_ATTRIBUTE_UNUSED = 1;\n" - "ptrdiff_t /*const*/ dj CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" - "ptrdiff_t /*const*/ dk CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" - "ptrdiff_t /*const*/ cdi CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * di;\n" - "ptrdiff_t /*const*/ cdj CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dj;\n" - "ptrdiff_t /*const*/ cdk CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dk;\n" - "CCTK_REAL_VEC /*const*/ dx CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(0));\n" - "CCTK_REAL_VEC /*const*/ dy CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(1));\n" - "CCTK_REAL_VEC /*const*/ dz CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(2));\n" - "CCTK_REAL_VEC /*const*/ dt CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_TIME);\n" - "CCTK_REAL_VEC /*const*/ t CCTK_ATTRIBUTE_UNUSED = ToReal(cctk_time);\n" - "CCTK_REAL_VEC /*const*/ dxi CCTK_ATTRIBUTE_UNUSED = INV(dx);\n" - "CCTK_REAL_VEC /*const*/ dyi CCTK_ATTRIBUTE_UNUSED = INV(dy);\n" - "CCTK_REAL_VEC /*const*/ dzi CCTK_ATTRIBUTE_UNUSED = INV(dz);\n" - "CCTK_REAL_VEC /*const*/ khalf CCTK_ATTRIBUTE_UNUSED = ToReal(0.5);\n" - "CCTK_REAL_VEC /*const*/ kthird CCTK_ATTRIBUTE_UNUSED = ToReal(1.0/3.0);\n" - "CCTK_REAL_VEC /*const*/ ktwothird CCTK_ATTRIBUTE_UNUSED = ToReal(2.0/3.0);\n" - "CCTK_REAL_VEC /*const*/ kfourthird CCTK_ATTRIBUTE_UNUSED = ToReal(4.0/3.0);\n" - "CCTK_REAL_VEC /*const*/ keightthird CCTK_ATTRIBUTE_UNUSED = ToReal(8.0/3.0);\n" - "CCTK_REAL_VEC /*const*/ hdxi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dxi);\n" - "CCTK_REAL_VEC /*const*/ hdyi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dyi);\n" - "CCTK_REAL_VEC /*const*/ hdzi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dzi);\n" + "const ptrdiff_t di CCTK_ATTRIBUTE_UNUSED = 1;\n" + "const ptrdiff_t dj CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" + "const ptrdiff_t dk CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" + "const ptrdiff_t cdi CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * di;\n" + "const ptrdiff_t cdj CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dj;\n" + "const ptrdiff_t cdk CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dk;\n" + "const CCTK_REAL_VEC dx CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(0));\n" + "const CCTK_REAL_VEC dy CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(1));\n" + "const CCTK_REAL_VEC dz CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(2));\n" + "const CCTK_REAL_VEC dt CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_TIME);\n" + "const CCTK_REAL_VEC t CCTK_ATTRIBUTE_UNUSED = ToReal(cctk_time);\n" + "const CCTK_REAL_VEC dxi CCTK_ATTRIBUTE_UNUSED = INV(dx);\n" + "const CCTK_REAL_VEC dyi CCTK_ATTRIBUTE_UNUSED = INV(dy);\n" + "const CCTK_REAL_VEC dzi CCTK_ATTRIBUTE_UNUSED = INV(dz);\n" + "const CCTK_REAL_VEC khalf CCTK_ATTRIBUTE_UNUSED = ToReal(0.5);\n" + "const CCTK_REAL_VEC kthird CCTK_ATTRIBUTE_UNUSED = ToReal(1.0/3.0);\n" + "const CCTK_REAL_VEC ktwothird CCTK_ATTRIBUTE_UNUSED = ToReal(2.0/3.0);\n" + "const CCTK_REAL_VEC kfourthird CCTK_ATTRIBUTE_UNUSED = ToReal(4.0/3.0);\n" + "const CCTK_REAL_VEC keightthird CCTK_ATTRIBUTE_UNUSED = ToReal(8.0/3.0);\n" + "const CCTK_REAL_VEC hdxi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dxi);\n" + "const CCTK_REAL_VEC hdyi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dyi);\n" + "const CCTK_REAL_VEC hdzi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dzi);\n" "\n" "/* Initialize predefined quantities */\n" - "CCTK_REAL_VEC /*const*/ p1o12dx CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dx);\n" - "CCTK_REAL_VEC /*const*/ p1o12dy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dy);\n" - "CCTK_REAL_VEC /*const*/ p1o12dz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dz);\n" - "CCTK_REAL_VEC /*const*/ p1o144dxdy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dy,dx));\n" - "CCTK_REAL_VEC /*const*/ p1o144dxdz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dx));\n" - "CCTK_REAL_VEC /*const*/ p1o144dydz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dy));\n" - "CCTK_REAL_VEC /*const*/ pm1o12dx2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dx,dx));\n" - "CCTK_REAL_VEC /*const*/ pm1o12dy2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dy,dy));\n" - "CCTK_REAL_VEC /*const*/ pm1o12dz2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dz,dz));\n" + "const CCTK_REAL_VEC p1o12dx CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dx);\n" + "const CCTK_REAL_VEC p1o12dy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dy);\n" + "const CCTK_REAL_VEC p1o12dz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dz);\n" + "const CCTK_REAL_VEC p1o144dxdy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dy,dx));\n" + "const CCTK_REAL_VEC p1o144dxdz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dx));\n" + "const CCTK_REAL_VEC p1o144dydz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dy));\n" + "const CCTK_REAL_VEC pm1o12dx2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dx,dx));\n" + "const CCTK_REAL_VEC pm1o12dy2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dy,dy));\n" + "const CCTK_REAL_VEC pm1o12dz2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dz,dz));\n" "\n" "/* Assign local copies of arrays functions */\n" "\n" @@ -93,13 +95,20 @@ static void WT_CL_Energy_Body(cGH const * restrict const cctkGH, int const dir, "/* Copy local copies back to grid functions */\n" "\n" "/* Loop over the grid points */\n" - "#pragma omp parallel\n" + "const int imin0=imin[0];\n" + "const int imin1=imin[1];\n" + "const int imin2=imin[2];\n" + "const int imax0=imax[0];\n" + "const int imax1=imax[1];\n" + "const int imax2=imax[2];\n" + "#pragma omp parallel // reduction(+: vec_iter_counter, vec_op_counter, vec_mem_counter)\n" "CCTK_LOOP3STR(WT_CL_Energy,\n" - " i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2],\n" + " i,j,k, imin0,imin1,imin2, imax0,imax1,imax2,\n" " cctk_ash[0],cctk_ash[1],cctk_ash[2],\n" - " kimin,kimax, CCTK_REAL_VEC_SIZE)\n" + " vecimin,vecimax, CCTK_REAL_VEC_SIZE)\n" "{\n" - " ptrdiff_t /*const*/ index CCTK_ATTRIBUTE_UNUSED = di*i + dj*j + dk*k;\n" + " const ptrdiff_t index CCTK_ATTRIBUTE_UNUSED = di*i + dj*j + dk*k;\n" + " // vec_iter_counter+=CCTK_REAL_VEC_SIZE;\n" " \n" " /* Assign local copies of grid functions */\n" " \n" @@ -110,9 +119,9 @@ static void WT_CL_Energy_Body(cGH const * restrict const cctkGH, int const dir, " /* Include user supplied include files */\n" " \n" " /* Precompute derivatives */\n" - " CCTK_REAL_VEC /*const*/ PDstandardNth1u CCTK_ATTRIBUTE_UNUSED = PDstandardNth1(&u[index]);\n" - " CCTK_REAL_VEC /*const*/ PDstandardNth2u CCTK_ATTRIBUTE_UNUSED = PDstandardNth2(&u[index]);\n" - " CCTK_REAL_VEC /*const*/ PDstandardNth3u CCTK_ATTRIBUTE_UNUSED = PDstandardNth3(&u[index]);\n" + " const CCTK_REAL_VEC PDstandardNth1u CCTK_ATTRIBUTE_UNUSED = PDstandardNth1(&u[index]);\n" + " const CCTK_REAL_VEC PDstandardNth2u CCTK_ATTRIBUTE_UNUSED = PDstandardNth2(&u[index]);\n" + " const CCTK_REAL_VEC PDstandardNth3u CCTK_ATTRIBUTE_UNUSED = PDstandardNth3(&u[index]);\n" " \n" " /* Calculate temporaries and grid functions */\n" " CCTK_REAL_VEC CCTK_ATTRIBUTE_UNUSED epsL = \n" @@ -126,14 +135,14 @@ static void WT_CL_Energy_Body(cGH const * restrict const cctkGH, int const dir, "" ; - char const *const groups[] = { + const char* 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}; + const char* const sources[] = {differencing, source, NULL}; OpenCLRunTime_CallKernel(cctkGH, CCTK_THORNSTRING, "WT_CL_Energy", sources, groups, NULL, NULL, NULL, -1, imin, imax, &kernel); @@ -156,7 +165,7 @@ extern "C" void WT_CL_Energy(CCTK_ARGUMENTS) return; } - const char *const groups[] = { + const char* const groups[] = { "ML_WaveToy_CL::WT_eps", "ML_WaveToy_CL::WT_rho", "ML_WaveToy_CL::WT_u"}; diff --git a/ML_WaveToy_CL/src/WT_CL_EnergyBoundary.cc b/ML_WaveToy_CL/src/WT_CL_EnergyBoundary.cc index 592d29e..8161612 100644 --- a/ML_WaveToy_CL/src/WT_CL_EnergyBoundary.cc +++ b/ML_WaveToy_CL/src/WT_CL_EnergyBoundary.cc @@ -33,56 +33,58 @@ extern "C" void WT_CL_EnergyBoundary_SelectBCs(CCTK_ARGUMENTS) DECLARE_CCTK_ARGUMENTS; DECLARE_CCTK_PARAMETERS; - CCTK_INT ierr CCTK_ATTRIBUTE_UNUSED = 0; + if (cctk_iteration % WT_CL_EnergyBoundary_calc_every != WT_CL_EnergyBoundary_calc_offset) + return; + CCTK_INT ierr CCTK_ATTRIBUTE_UNUSED = 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[]) +static void WT_CL_EnergyBoundary_Body(const cGH* restrict const cctkGH, const int dir, const int face, const CCTK_REAL normal[3], const CCTK_REAL tangentA[3], const CCTK_REAL tangentB[3], const int imin[3], const int imax[3], const int n_subblock_gfs, CCTK_REAL* restrict const subblock_gfs[]) { DECLARE_CCTK_ARGUMENTS; DECLARE_CCTK_PARAMETERS; - char const *const source = + const char* const source = "\n" "/* Include user-supplied include files */\n" "\n" "/* Initialise finite differencing variables */\n" - "ptrdiff_t /*const*/ di CCTK_ATTRIBUTE_UNUSED = 1;\n" - "ptrdiff_t /*const*/ dj CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" - "ptrdiff_t /*const*/ dk CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" - "ptrdiff_t /*const*/ cdi CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * di;\n" - "ptrdiff_t /*const*/ cdj CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dj;\n" - "ptrdiff_t /*const*/ cdk CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dk;\n" - "CCTK_REAL_VEC /*const*/ dx CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(0));\n" - "CCTK_REAL_VEC /*const*/ dy CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(1));\n" - "CCTK_REAL_VEC /*const*/ dz CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(2));\n" - "CCTK_REAL_VEC /*const*/ dt CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_TIME);\n" - "CCTK_REAL_VEC /*const*/ t CCTK_ATTRIBUTE_UNUSED = ToReal(cctk_time);\n" - "CCTK_REAL_VEC /*const*/ dxi CCTK_ATTRIBUTE_UNUSED = INV(dx);\n" - "CCTK_REAL_VEC /*const*/ dyi CCTK_ATTRIBUTE_UNUSED = INV(dy);\n" - "CCTK_REAL_VEC /*const*/ dzi CCTK_ATTRIBUTE_UNUSED = INV(dz);\n" - "CCTK_REAL_VEC /*const*/ khalf CCTK_ATTRIBUTE_UNUSED = ToReal(0.5);\n" - "CCTK_REAL_VEC /*const*/ kthird CCTK_ATTRIBUTE_UNUSED = ToReal(1.0/3.0);\n" - "CCTK_REAL_VEC /*const*/ ktwothird CCTK_ATTRIBUTE_UNUSED = ToReal(2.0/3.0);\n" - "CCTK_REAL_VEC /*const*/ kfourthird CCTK_ATTRIBUTE_UNUSED = ToReal(4.0/3.0);\n" - "CCTK_REAL_VEC /*const*/ keightthird CCTK_ATTRIBUTE_UNUSED = ToReal(8.0/3.0);\n" - "CCTK_REAL_VEC /*const*/ hdxi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dxi);\n" - "CCTK_REAL_VEC /*const*/ hdyi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dyi);\n" - "CCTK_REAL_VEC /*const*/ hdzi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dzi);\n" + "const ptrdiff_t di CCTK_ATTRIBUTE_UNUSED = 1;\n" + "const ptrdiff_t dj CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" + "const ptrdiff_t dk CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" + "const ptrdiff_t cdi CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * di;\n" + "const ptrdiff_t cdj CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dj;\n" + "const ptrdiff_t cdk CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dk;\n" + "const CCTK_REAL_VEC dx CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(0));\n" + "const CCTK_REAL_VEC dy CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(1));\n" + "const CCTK_REAL_VEC dz CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(2));\n" + "const CCTK_REAL_VEC dt CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_TIME);\n" + "const CCTK_REAL_VEC t CCTK_ATTRIBUTE_UNUSED = ToReal(cctk_time);\n" + "const CCTK_REAL_VEC dxi CCTK_ATTRIBUTE_UNUSED = INV(dx);\n" + "const CCTK_REAL_VEC dyi CCTK_ATTRIBUTE_UNUSED = INV(dy);\n" + "const CCTK_REAL_VEC dzi CCTK_ATTRIBUTE_UNUSED = INV(dz);\n" + "const CCTK_REAL_VEC khalf CCTK_ATTRIBUTE_UNUSED = ToReal(0.5);\n" + "const CCTK_REAL_VEC kthird CCTK_ATTRIBUTE_UNUSED = ToReal(1.0/3.0);\n" + "const CCTK_REAL_VEC ktwothird CCTK_ATTRIBUTE_UNUSED = ToReal(2.0/3.0);\n" + "const CCTK_REAL_VEC kfourthird CCTK_ATTRIBUTE_UNUSED = ToReal(4.0/3.0);\n" + "const CCTK_REAL_VEC keightthird CCTK_ATTRIBUTE_UNUSED = ToReal(8.0/3.0);\n" + "const CCTK_REAL_VEC hdxi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dxi);\n" + "const CCTK_REAL_VEC hdyi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dyi);\n" + "const CCTK_REAL_VEC hdzi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dzi);\n" "\n" "/* Initialize predefined quantities */\n" - "CCTK_REAL_VEC /*const*/ p1o12dx CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dx);\n" - "CCTK_REAL_VEC /*const*/ p1o12dy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dy);\n" - "CCTK_REAL_VEC /*const*/ p1o12dz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dz);\n" - "CCTK_REAL_VEC /*const*/ p1o144dxdy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dy,dx));\n" - "CCTK_REAL_VEC /*const*/ p1o144dxdz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dx));\n" - "CCTK_REAL_VEC /*const*/ p1o144dydz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dy));\n" - "CCTK_REAL_VEC /*const*/ pm1o12dx2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dx,dx));\n" - "CCTK_REAL_VEC /*const*/ pm1o12dy2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dy,dy));\n" - "CCTK_REAL_VEC /*const*/ pm1o12dz2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dz,dz));\n" + "const CCTK_REAL_VEC p1o12dx CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dx);\n" + "const CCTK_REAL_VEC p1o12dy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dy);\n" + "const CCTK_REAL_VEC p1o12dz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dz);\n" + "const CCTK_REAL_VEC p1o144dxdy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dy,dx));\n" + "const CCTK_REAL_VEC p1o144dxdz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dx));\n" + "const CCTK_REAL_VEC p1o144dydz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dy));\n" + "const CCTK_REAL_VEC pm1o12dx2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dx,dx));\n" + "const CCTK_REAL_VEC pm1o12dy2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dy,dy));\n" + "const CCTK_REAL_VEC pm1o12dz2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dz,dz));\n" "\n" "/* Assign local copies of arrays functions */\n" "\n" @@ -93,13 +95,20 @@ static void WT_CL_EnergyBoundary_Body(cGH const * restrict const cctkGH, int con "/* Copy local copies back to grid functions */\n" "\n" "/* Loop over the grid points */\n" - "#pragma omp parallel\n" + "const int imin0=imin[0];\n" + "const int imin1=imin[1];\n" + "const int imin2=imin[2];\n" + "const int imax0=imax[0];\n" + "const int imax1=imax[1];\n" + "const int imax2=imax[2];\n" + "#pragma omp parallel // reduction(+: vec_iter_counter, vec_op_counter, vec_mem_counter)\n" "CCTK_LOOP3STR(WT_CL_EnergyBoundary,\n" - " i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2],\n" + " i,j,k, imin0,imin1,imin2, imax0,imax1,imax2,\n" " cctk_ash[0],cctk_ash[1],cctk_ash[2],\n" - " kimin,kimax, CCTK_REAL_VEC_SIZE)\n" + " vecimin,vecimax, CCTK_REAL_VEC_SIZE)\n" "{\n" - " ptrdiff_t /*const*/ index CCTK_ATTRIBUTE_UNUSED = di*i + dj*j + dk*k;\n" + " const ptrdiff_t index CCTK_ATTRIBUTE_UNUSED = di*i + dj*j + dk*k;\n" + " // vec_iter_counter+=CCTK_REAL_VEC_SIZE;\n" " \n" " /* Assign local copies of grid functions */\n" " \n" @@ -110,7 +119,7 @@ static void WT_CL_EnergyBoundary_Body(cGH const * restrict const cctkGH, int con " /* Precompute derivatives */\n" " \n" " /* Calculate temporaries and grid functions */\n" - " CCTK_REAL_VEC CCTK_ATTRIBUTE_UNUSED epsL = ToReal(0.);\n" + " CCTK_REAL_VEC CCTK_ATTRIBUTE_UNUSED epsL = ToReal(0);\n" " \n" " /* Copy local copies back to grid functions */\n" " vec_store_partial_prepare(i,lc_imin,lc_imax);\n" @@ -120,12 +129,12 @@ static void WT_CL_EnergyBoundary_Body(cGH const * restrict const cctkGH, int con "" ; - char const *const groups[] = { + const char* const groups[] = { "ML_WaveToy_CL::WT_eps", NULL}; static struct OpenCLKernel *kernel = NULL; - char const *const sources[] = {differencing, source, NULL}; + const char* const sources[] = {differencing, source, NULL}; OpenCLRunTime_CallKernel(cctkGH, CCTK_THORNSTRING, "WT_CL_EnergyBoundary", sources, groups, NULL, NULL, NULL, -1, imin, imax, &kernel); @@ -148,7 +157,7 @@ extern "C" void WT_CL_EnergyBoundary(CCTK_ARGUMENTS) return; } - const char *const groups[] = { + const char* const groups[] = { "ML_WaveToy_CL::WT_eps"}; GenericFD_AssertGroupStorage(cctkGH, "WT_CL_EnergyBoundary", 1, groups); diff --git a/ML_WaveToy_CL/src/WT_CL_Gaussian.cc b/ML_WaveToy_CL/src/WT_CL_Gaussian.cc index fde5721..9cfe036 100644 --- a/ML_WaveToy_CL/src/WT_CL_Gaussian.cc +++ b/ML_WaveToy_CL/src/WT_CL_Gaussian.cc @@ -28,49 +28,49 @@ #define CUB(x) (kmul(x,SQR(x))) #define QAD(x) (SQR(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[]) +static void WT_CL_Gaussian_Body(const cGH* restrict const cctkGH, const int dir, const int face, const CCTK_REAL normal[3], const CCTK_REAL tangentA[3], const CCTK_REAL tangentB[3], const int imin[3], const int imax[3], const int n_subblock_gfs, CCTK_REAL* restrict const subblock_gfs[]) { DECLARE_CCTK_ARGUMENTS; DECLARE_CCTK_PARAMETERS; - char const *const source = + const char* const source = "\n" "/* Include user-supplied include files */\n" "\n" "/* Initialise finite differencing variables */\n" - "ptrdiff_t /*const*/ di CCTK_ATTRIBUTE_UNUSED = 1;\n" - "ptrdiff_t /*const*/ dj CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" - "ptrdiff_t /*const*/ dk CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" - "ptrdiff_t /*const*/ cdi CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * di;\n" - "ptrdiff_t /*const*/ cdj CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dj;\n" - "ptrdiff_t /*const*/ cdk CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dk;\n" - "CCTK_REAL_VEC /*const*/ dx CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(0));\n" - "CCTK_REAL_VEC /*const*/ dy CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(1));\n" - "CCTK_REAL_VEC /*const*/ dz CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(2));\n" - "CCTK_REAL_VEC /*const*/ dt CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_TIME);\n" - "CCTK_REAL_VEC /*const*/ t CCTK_ATTRIBUTE_UNUSED = ToReal(cctk_time);\n" - "CCTK_REAL_VEC /*const*/ dxi CCTK_ATTRIBUTE_UNUSED = INV(dx);\n" - "CCTK_REAL_VEC /*const*/ dyi CCTK_ATTRIBUTE_UNUSED = INV(dy);\n" - "CCTK_REAL_VEC /*const*/ dzi CCTK_ATTRIBUTE_UNUSED = INV(dz);\n" - "CCTK_REAL_VEC /*const*/ khalf CCTK_ATTRIBUTE_UNUSED = ToReal(0.5);\n" - "CCTK_REAL_VEC /*const*/ kthird CCTK_ATTRIBUTE_UNUSED = ToReal(1.0/3.0);\n" - "CCTK_REAL_VEC /*const*/ ktwothird CCTK_ATTRIBUTE_UNUSED = ToReal(2.0/3.0);\n" - "CCTK_REAL_VEC /*const*/ kfourthird CCTK_ATTRIBUTE_UNUSED = ToReal(4.0/3.0);\n" - "CCTK_REAL_VEC /*const*/ keightthird CCTK_ATTRIBUTE_UNUSED = ToReal(8.0/3.0);\n" - "CCTK_REAL_VEC /*const*/ hdxi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dxi);\n" - "CCTK_REAL_VEC /*const*/ hdyi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dyi);\n" - "CCTK_REAL_VEC /*const*/ hdzi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dzi);\n" + "const ptrdiff_t di CCTK_ATTRIBUTE_UNUSED = 1;\n" + "const ptrdiff_t dj CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" + "const ptrdiff_t dk CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" + "const ptrdiff_t cdi CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * di;\n" + "const ptrdiff_t cdj CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dj;\n" + "const ptrdiff_t cdk CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dk;\n" + "const CCTK_REAL_VEC dx CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(0));\n" + "const CCTK_REAL_VEC dy CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(1));\n" + "const CCTK_REAL_VEC dz CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(2));\n" + "const CCTK_REAL_VEC dt CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_TIME);\n" + "const CCTK_REAL_VEC t CCTK_ATTRIBUTE_UNUSED = ToReal(cctk_time);\n" + "const CCTK_REAL_VEC dxi CCTK_ATTRIBUTE_UNUSED = INV(dx);\n" + "const CCTK_REAL_VEC dyi CCTK_ATTRIBUTE_UNUSED = INV(dy);\n" + "const CCTK_REAL_VEC dzi CCTK_ATTRIBUTE_UNUSED = INV(dz);\n" + "const CCTK_REAL_VEC khalf CCTK_ATTRIBUTE_UNUSED = ToReal(0.5);\n" + "const CCTK_REAL_VEC kthird CCTK_ATTRIBUTE_UNUSED = ToReal(1.0/3.0);\n" + "const CCTK_REAL_VEC ktwothird CCTK_ATTRIBUTE_UNUSED = ToReal(2.0/3.0);\n" + "const CCTK_REAL_VEC kfourthird CCTK_ATTRIBUTE_UNUSED = ToReal(4.0/3.0);\n" + "const CCTK_REAL_VEC keightthird CCTK_ATTRIBUTE_UNUSED = ToReal(8.0/3.0);\n" + "const CCTK_REAL_VEC hdxi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dxi);\n" + "const CCTK_REAL_VEC hdyi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dyi);\n" + "const CCTK_REAL_VEC hdzi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dzi);\n" "\n" "/* Initialize predefined quantities */\n" - "CCTK_REAL_VEC /*const*/ p1o12dx CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dx);\n" - "CCTK_REAL_VEC /*const*/ p1o12dy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dy);\n" - "CCTK_REAL_VEC /*const*/ p1o12dz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dz);\n" - "CCTK_REAL_VEC /*const*/ p1o144dxdy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dy,dx));\n" - "CCTK_REAL_VEC /*const*/ p1o144dxdz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dx));\n" - "CCTK_REAL_VEC /*const*/ p1o144dydz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dy));\n" - "CCTK_REAL_VEC /*const*/ pm1o12dx2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dx,dx));\n" - "CCTK_REAL_VEC /*const*/ pm1o12dy2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dy,dy));\n" - "CCTK_REAL_VEC /*const*/ pm1o12dz2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dz,dz));\n" + "const CCTK_REAL_VEC p1o12dx CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dx);\n" + "const CCTK_REAL_VEC p1o12dy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dy);\n" + "const CCTK_REAL_VEC p1o12dz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dz);\n" + "const CCTK_REAL_VEC p1o144dxdy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dy,dx));\n" + "const CCTK_REAL_VEC p1o144dxdz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dx));\n" + "const CCTK_REAL_VEC p1o144dydz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dy));\n" + "const CCTK_REAL_VEC pm1o12dx2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dx,dx));\n" + "const CCTK_REAL_VEC pm1o12dy2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dy,dy));\n" + "const CCTK_REAL_VEC pm1o12dz2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dz,dz));\n" "\n" "/* Assign local copies of arrays functions */\n" "\n" @@ -81,13 +81,20 @@ static void WT_CL_Gaussian_Body(cGH const * restrict const cctkGH, int const dir "/* Copy local copies back to grid functions */\n" "\n" "/* Loop over the grid points */\n" - "#pragma omp parallel\n" + "const int imin0=imin[0];\n" + "const int imin1=imin[1];\n" + "const int imin2=imin[2];\n" + "const int imax0=imax[0];\n" + "const int imax1=imax[1];\n" + "const int imax2=imax[2];\n" + "#pragma omp parallel // reduction(+: vec_iter_counter, vec_op_counter, vec_mem_counter)\n" "CCTK_LOOP3STR(WT_CL_Gaussian,\n" - " i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2],\n" + " i,j,k, imin0,imin1,imin2, imax0,imax1,imax2,\n" " cctk_ash[0],cctk_ash[1],cctk_ash[2],\n" - " kimin,kimax, CCTK_REAL_VEC_SIZE)\n" + " vecimin,vecimax, CCTK_REAL_VEC_SIZE)\n" "{\n" - " ptrdiff_t /*const*/ index CCTK_ATTRIBUTE_UNUSED = di*i + dj*j + dk*k;\n" + " const ptrdiff_t index CCTK_ATTRIBUTE_UNUSED = di*i + dj*j + dk*k;\n" + " // vec_iter_counter+=CCTK_REAL_VEC_SIZE;\n" " \n" " /* Assign local copies of grid functions */\n" " \n" @@ -102,7 +109,7 @@ static void WT_CL_Gaussian_Body(cGH const * restrict const cctkGH, int const dir " CCTK_REAL_VEC CCTK_ATTRIBUTE_UNUSED uL = \n" " kmul(kexp(kmul(kmul(kmul(rL,rL),ToReal(-0.5)),ToReal(ScalarINV(ScalarSQR(width))))),ToReal(amplitude));\n" " \n" - " CCTK_REAL_VEC CCTK_ATTRIBUTE_UNUSED rhoL = ToReal(0.);\n" + " CCTK_REAL_VEC CCTK_ATTRIBUTE_UNUSED rhoL = ToReal(0);\n" " \n" " /* Copy local copies back to grid functions */\n" " vec_store_partial_prepare(i,lc_imin,lc_imax);\n" @@ -113,14 +120,14 @@ static void WT_CL_Gaussian_Body(cGH const * restrict const cctkGH, int const dir "" ; - char const *const groups[] = { + const char* 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}; + const char* const sources[] = {differencing, source, NULL}; OpenCLRunTime_CallKernel(cctkGH, CCTK_THORNSTRING, "WT_CL_Gaussian", sources, groups, NULL, NULL, NULL, -1, imin, imax, &kernel); @@ -143,7 +150,7 @@ extern "C" void WT_CL_Gaussian(CCTK_ARGUMENTS) return; } - const char *const groups[] = { + const char* const groups[] = { "grid::coordinates", "ML_WaveToy_CL::WT_rho", "ML_WaveToy_CL::WT_u"}; diff --git a/ML_WaveToy_CL/src/WT_CL_RHS.cc b/ML_WaveToy_CL/src/WT_CL_RHS.cc index a8212f6..128496a 100644 --- a/ML_WaveToy_CL/src/WT_CL_RHS.cc +++ b/ML_WaveToy_CL/src/WT_CL_RHS.cc @@ -33,7 +33,9 @@ extern "C" void WT_CL_RHS_SelectBCs(CCTK_ARGUMENTS) DECLARE_CCTK_ARGUMENTS; DECLARE_CCTK_PARAMETERS; - CCTK_INT ierr CCTK_ATTRIBUTE_UNUSED = 0; + if (cctk_iteration % WT_CL_RHS_calc_every != WT_CL_RHS_calc_offset) + return; + CCTK_INT ierr CCTK_ATTRIBUTE_UNUSED = 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."); @@ -43,49 +45,49 @@ extern "C" void WT_CL_RHS_SelectBCs(CCTK_ARGUMENTS) 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[]) +static void WT_CL_RHS_Body(const cGH* restrict const cctkGH, const int dir, const int face, const CCTK_REAL normal[3], const CCTK_REAL tangentA[3], const CCTK_REAL tangentB[3], const int imin[3], const int imax[3], const int n_subblock_gfs, CCTK_REAL* restrict const subblock_gfs[]) { DECLARE_CCTK_ARGUMENTS; DECLARE_CCTK_PARAMETERS; - char const *const source = + const char* const source = "\n" "/* Include user-supplied include files */\n" "\n" "/* Initialise finite differencing variables */\n" - "ptrdiff_t /*const*/ di CCTK_ATTRIBUTE_UNUSED = 1;\n" - "ptrdiff_t /*const*/ dj CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" - "ptrdiff_t /*const*/ dk CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" - "ptrdiff_t /*const*/ cdi CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * di;\n" - "ptrdiff_t /*const*/ cdj CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dj;\n" - "ptrdiff_t /*const*/ cdk CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dk;\n" - "CCTK_REAL_VEC /*const*/ dx CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(0));\n" - "CCTK_REAL_VEC /*const*/ dy CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(1));\n" - "CCTK_REAL_VEC /*const*/ dz CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(2));\n" - "CCTK_REAL_VEC /*const*/ dt CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_TIME);\n" - "CCTK_REAL_VEC /*const*/ t CCTK_ATTRIBUTE_UNUSED = ToReal(cctk_time);\n" - "CCTK_REAL_VEC /*const*/ dxi CCTK_ATTRIBUTE_UNUSED = INV(dx);\n" - "CCTK_REAL_VEC /*const*/ dyi CCTK_ATTRIBUTE_UNUSED = INV(dy);\n" - "CCTK_REAL_VEC /*const*/ dzi CCTK_ATTRIBUTE_UNUSED = INV(dz);\n" - "CCTK_REAL_VEC /*const*/ khalf CCTK_ATTRIBUTE_UNUSED = ToReal(0.5);\n" - "CCTK_REAL_VEC /*const*/ kthird CCTK_ATTRIBUTE_UNUSED = ToReal(1.0/3.0);\n" - "CCTK_REAL_VEC /*const*/ ktwothird CCTK_ATTRIBUTE_UNUSED = ToReal(2.0/3.0);\n" - "CCTK_REAL_VEC /*const*/ kfourthird CCTK_ATTRIBUTE_UNUSED = ToReal(4.0/3.0);\n" - "CCTK_REAL_VEC /*const*/ keightthird CCTK_ATTRIBUTE_UNUSED = ToReal(8.0/3.0);\n" - "CCTK_REAL_VEC /*const*/ hdxi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dxi);\n" - "CCTK_REAL_VEC /*const*/ hdyi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dyi);\n" - "CCTK_REAL_VEC /*const*/ hdzi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dzi);\n" + "const ptrdiff_t di CCTK_ATTRIBUTE_UNUSED = 1;\n" + "const ptrdiff_t dj CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" + "const ptrdiff_t dk CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" + "const ptrdiff_t cdi CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * di;\n" + "const ptrdiff_t cdj CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dj;\n" + "const ptrdiff_t cdk CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dk;\n" + "const CCTK_REAL_VEC dx CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(0));\n" + "const CCTK_REAL_VEC dy CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(1));\n" + "const CCTK_REAL_VEC dz CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(2));\n" + "const CCTK_REAL_VEC dt CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_TIME);\n" + "const CCTK_REAL_VEC t CCTK_ATTRIBUTE_UNUSED = ToReal(cctk_time);\n" + "const CCTK_REAL_VEC dxi CCTK_ATTRIBUTE_UNUSED = INV(dx);\n" + "const CCTK_REAL_VEC dyi CCTK_ATTRIBUTE_UNUSED = INV(dy);\n" + "const CCTK_REAL_VEC dzi CCTK_ATTRIBUTE_UNUSED = INV(dz);\n" + "const CCTK_REAL_VEC khalf CCTK_ATTRIBUTE_UNUSED = ToReal(0.5);\n" + "const CCTK_REAL_VEC kthird CCTK_ATTRIBUTE_UNUSED = ToReal(1.0/3.0);\n" + "const CCTK_REAL_VEC ktwothird CCTK_ATTRIBUTE_UNUSED = ToReal(2.0/3.0);\n" + "const CCTK_REAL_VEC kfourthird CCTK_ATTRIBUTE_UNUSED = ToReal(4.0/3.0);\n" + "const CCTK_REAL_VEC keightthird CCTK_ATTRIBUTE_UNUSED = ToReal(8.0/3.0);\n" + "const CCTK_REAL_VEC hdxi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dxi);\n" + "const CCTK_REAL_VEC hdyi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dyi);\n" + "const CCTK_REAL_VEC hdzi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dzi);\n" "\n" "/* Initialize predefined quantities */\n" - "CCTK_REAL_VEC /*const*/ p1o12dx CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dx);\n" - "CCTK_REAL_VEC /*const*/ p1o12dy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dy);\n" - "CCTK_REAL_VEC /*const*/ p1o12dz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dz);\n" - "CCTK_REAL_VEC /*const*/ p1o144dxdy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dy,dx));\n" - "CCTK_REAL_VEC /*const*/ p1o144dxdz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dx));\n" - "CCTK_REAL_VEC /*const*/ p1o144dydz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dy));\n" - "CCTK_REAL_VEC /*const*/ pm1o12dx2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dx,dx));\n" - "CCTK_REAL_VEC /*const*/ pm1o12dy2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dy,dy));\n" - "CCTK_REAL_VEC /*const*/ pm1o12dz2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dz,dz));\n" + "const CCTK_REAL_VEC p1o12dx CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dx);\n" + "const CCTK_REAL_VEC p1o12dy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dy);\n" + "const CCTK_REAL_VEC p1o12dz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dz);\n" + "const CCTK_REAL_VEC p1o144dxdy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dy,dx));\n" + "const CCTK_REAL_VEC p1o144dxdz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dx));\n" + "const CCTK_REAL_VEC p1o144dydz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dy));\n" + "const CCTK_REAL_VEC pm1o12dx2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dx,dx));\n" + "const CCTK_REAL_VEC pm1o12dy2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dy,dy));\n" + "const CCTK_REAL_VEC pm1o12dz2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dz,dz));\n" "\n" "/* Assign local copies of arrays functions */\n" "\n" @@ -96,13 +98,20 @@ static void WT_CL_RHS_Body(cGH const * restrict const cctkGH, int const dir, int "/* Copy local copies back to grid functions */\n" "\n" "/* Loop over the grid points */\n" - "#pragma omp parallel\n" + "const int imin0=imin[0];\n" + "const int imin1=imin[1];\n" + "const int imin2=imin[2];\n" + "const int imax0=imax[0];\n" + "const int imax1=imax[1];\n" + "const int imax2=imax[2];\n" + "#pragma omp parallel // reduction(+: vec_iter_counter, vec_op_counter, vec_mem_counter)\n" "CCTK_LOOP3STR(WT_CL_RHS,\n" - " i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2],\n" + " i,j,k, imin0,imin1,imin2, imax0,imax1,imax2,\n" " cctk_ash[0],cctk_ash[1],cctk_ash[2],\n" - " kimin,kimax, CCTK_REAL_VEC_SIZE)\n" + " vecimin,vecimax, CCTK_REAL_VEC_SIZE)\n" "{\n" - " ptrdiff_t /*const*/ index CCTK_ATTRIBUTE_UNUSED = di*i + dj*j + dk*k;\n" + " const ptrdiff_t index CCTK_ATTRIBUTE_UNUSED = di*i + dj*j + dk*k;\n" + " // vec_iter_counter+=CCTK_REAL_VEC_SIZE;\n" " \n" " /* Assign local copies of grid functions */\n" " \n" @@ -113,9 +122,9 @@ static void WT_CL_RHS_Body(cGH const * restrict const cctkGH, int const dir, int " /* Include user supplied include files */\n" " \n" " /* Precompute derivatives */\n" - " CCTK_REAL_VEC /*const*/ PDstandardNth11u CCTK_ATTRIBUTE_UNUSED = PDstandardNth11(&u[index]);\n" - " CCTK_REAL_VEC /*const*/ PDstandardNth22u CCTK_ATTRIBUTE_UNUSED = PDstandardNth22(&u[index]);\n" - " CCTK_REAL_VEC /*const*/ PDstandardNth33u CCTK_ATTRIBUTE_UNUSED = PDstandardNth33(&u[index]);\n" + " const CCTK_REAL_VEC PDstandardNth11u CCTK_ATTRIBUTE_UNUSED = PDstandardNth11(&u[index]);\n" + " const CCTK_REAL_VEC PDstandardNth22u CCTK_ATTRIBUTE_UNUSED = PDstandardNth22(&u[index]);\n" + " const CCTK_REAL_VEC PDstandardNth33u CCTK_ATTRIBUTE_UNUSED = PDstandardNth33(&u[index]);\n" " \n" " /* Calculate temporaries and grid functions */\n" " CCTK_REAL_VEC CCTK_ATTRIBUTE_UNUSED urhsL = rhoL;\n" @@ -132,7 +141,7 @@ static void WT_CL_RHS_Body(cGH const * restrict const cctkGH, int const dir, int "" ; - char const *const groups[] = { + const char* const groups[] = { "ML_WaveToy_CL::WT_rho", "ML_WaveToy_CL::WT_rhorhs", "ML_WaveToy_CL::WT_u", @@ -140,7 +149,7 @@ static void WT_CL_RHS_Body(cGH const * restrict const cctkGH, int const dir, int NULL}; static struct OpenCLKernel *kernel = NULL; - char const *const sources[] = {differencing, source, NULL}; + const char* const sources[] = {differencing, source, NULL}; OpenCLRunTime_CallKernel(cctkGH, CCTK_THORNSTRING, "WT_CL_RHS", sources, groups, NULL, NULL, NULL, -1, imin, imax, &kernel); @@ -163,7 +172,7 @@ extern "C" void WT_CL_RHS(CCTK_ARGUMENTS) return; } - const char *const groups[] = { + const char* const groups[] = { "ML_WaveToy_CL::WT_rho", "ML_WaveToy_CL::WT_rhorhs", "ML_WaveToy_CL::WT_u", diff --git a/ML_WaveToy_CL/src/WT_CL_Standing.cc b/ML_WaveToy_CL/src/WT_CL_Standing.cc index 3f79b33..6d590b8 100644 --- a/ML_WaveToy_CL/src/WT_CL_Standing.cc +++ b/ML_WaveToy_CL/src/WT_CL_Standing.cc @@ -28,49 +28,49 @@ #define CUB(x) (kmul(x,SQR(x))) #define QAD(x) (SQR(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[]) +static void WT_CL_Standing_Body(const cGH* restrict const cctkGH, const int dir, const int face, const CCTK_REAL normal[3], const CCTK_REAL tangentA[3], const CCTK_REAL tangentB[3], const int imin[3], const int imax[3], const int n_subblock_gfs, CCTK_REAL* restrict const subblock_gfs[]) { DECLARE_CCTK_ARGUMENTS; DECLARE_CCTK_PARAMETERS; - char const *const source = + const char* const source = "\n" "/* Include user-supplied include files */\n" "\n" "/* Initialise finite differencing variables */\n" - "ptrdiff_t /*const*/ di CCTK_ATTRIBUTE_UNUSED = 1;\n" - "ptrdiff_t /*const*/ dj CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" - "ptrdiff_t /*const*/ dk CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" - "ptrdiff_t /*const*/ cdi CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * di;\n" - "ptrdiff_t /*const*/ cdj CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dj;\n" - "ptrdiff_t /*const*/ cdk CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dk;\n" - "CCTK_REAL_VEC /*const*/ dx CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(0));\n" - "CCTK_REAL_VEC /*const*/ dy CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(1));\n" - "CCTK_REAL_VEC /*const*/ dz CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(2));\n" - "CCTK_REAL_VEC /*const*/ dt CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_TIME);\n" - "CCTK_REAL_VEC /*const*/ t CCTK_ATTRIBUTE_UNUSED = ToReal(cctk_time);\n" - "CCTK_REAL_VEC /*const*/ dxi CCTK_ATTRIBUTE_UNUSED = INV(dx);\n" - "CCTK_REAL_VEC /*const*/ dyi CCTK_ATTRIBUTE_UNUSED = INV(dy);\n" - "CCTK_REAL_VEC /*const*/ dzi CCTK_ATTRIBUTE_UNUSED = INV(dz);\n" - "CCTK_REAL_VEC /*const*/ khalf CCTK_ATTRIBUTE_UNUSED = ToReal(0.5);\n" - "CCTK_REAL_VEC /*const*/ kthird CCTK_ATTRIBUTE_UNUSED = ToReal(1.0/3.0);\n" - "CCTK_REAL_VEC /*const*/ ktwothird CCTK_ATTRIBUTE_UNUSED = ToReal(2.0/3.0);\n" - "CCTK_REAL_VEC /*const*/ kfourthird CCTK_ATTRIBUTE_UNUSED = ToReal(4.0/3.0);\n" - "CCTK_REAL_VEC /*const*/ keightthird CCTK_ATTRIBUTE_UNUSED = ToReal(8.0/3.0);\n" - "CCTK_REAL_VEC /*const*/ hdxi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dxi);\n" - "CCTK_REAL_VEC /*const*/ hdyi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dyi);\n" - "CCTK_REAL_VEC /*const*/ hdzi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dzi);\n" + "const ptrdiff_t di CCTK_ATTRIBUTE_UNUSED = 1;\n" + "const ptrdiff_t dj CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" + "const ptrdiff_t dk CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n" + "const ptrdiff_t cdi CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * di;\n" + "const ptrdiff_t cdj CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dj;\n" + "const ptrdiff_t cdk CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dk;\n" + "const CCTK_REAL_VEC dx CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(0));\n" + "const CCTK_REAL_VEC dy CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(1));\n" + "const CCTK_REAL_VEC dz CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(2));\n" + "const CCTK_REAL_VEC dt CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_TIME);\n" + "const CCTK_REAL_VEC t CCTK_ATTRIBUTE_UNUSED = ToReal(cctk_time);\n" + "const CCTK_REAL_VEC dxi CCTK_ATTRIBUTE_UNUSED = INV(dx);\n" + "const CCTK_REAL_VEC dyi CCTK_ATTRIBUTE_UNUSED = INV(dy);\n" + "const CCTK_REAL_VEC dzi CCTK_ATTRIBUTE_UNUSED = INV(dz);\n" + "const CCTK_REAL_VEC khalf CCTK_ATTRIBUTE_UNUSED = ToReal(0.5);\n" + "const CCTK_REAL_VEC kthird CCTK_ATTRIBUTE_UNUSED = ToReal(1.0/3.0);\n" + "const CCTK_REAL_VEC ktwothird CCTK_ATTRIBUTE_UNUSED = ToReal(2.0/3.0);\n" + "const CCTK_REAL_VEC kfourthird CCTK_ATTRIBUTE_UNUSED = ToReal(4.0/3.0);\n" + "const CCTK_REAL_VEC keightthird CCTK_ATTRIBUTE_UNUSED = ToReal(8.0/3.0);\n" + "const CCTK_REAL_VEC hdxi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dxi);\n" + "const CCTK_REAL_VEC hdyi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dyi);\n" + "const CCTK_REAL_VEC hdzi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dzi);\n" "\n" "/* Initialize predefined quantities */\n" - "CCTK_REAL_VEC /*const*/ p1o12dx CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dx);\n" - "CCTK_REAL_VEC /*const*/ p1o12dy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dy);\n" - "CCTK_REAL_VEC /*const*/ p1o12dz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dz);\n" - "CCTK_REAL_VEC /*const*/ p1o144dxdy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dy,dx));\n" - "CCTK_REAL_VEC /*const*/ p1o144dxdz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dx));\n" - "CCTK_REAL_VEC /*const*/ p1o144dydz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dy));\n" - "CCTK_REAL_VEC /*const*/ pm1o12dx2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dx,dx));\n" - "CCTK_REAL_VEC /*const*/ pm1o12dy2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dy,dy));\n" - "CCTK_REAL_VEC /*const*/ pm1o12dz2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dz,dz));\n" + "const CCTK_REAL_VEC p1o12dx CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dx);\n" + "const CCTK_REAL_VEC p1o12dy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dy);\n" + "const CCTK_REAL_VEC p1o12dz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dz);\n" + "const CCTK_REAL_VEC p1o144dxdy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dy,dx));\n" + "const CCTK_REAL_VEC p1o144dxdz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dx));\n" + "const CCTK_REAL_VEC p1o144dydz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dy));\n" + "const CCTK_REAL_VEC pm1o12dx2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dx,dx));\n" + "const CCTK_REAL_VEC pm1o12dy2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dy,dy));\n" + "const CCTK_REAL_VEC pm1o12dz2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dz,dz));\n" "\n" "/* Assign local copies of arrays functions */\n" "\n" @@ -81,13 +81,20 @@ static void WT_CL_Standing_Body(cGH const * restrict const cctkGH, int const dir "/* Copy local copies back to grid functions */\n" "\n" "/* Loop over the grid points */\n" - "#pragma omp parallel\n" + "const int imin0=imin[0];\n" + "const int imin1=imin[1];\n" + "const int imin2=imin[2];\n" + "const int imax0=imax[0];\n" + "const int imax1=imax[1];\n" + "const int imax2=imax[2];\n" + "#pragma omp parallel // reduction(+: vec_iter_counter, vec_op_counter, vec_mem_counter)\n" "CCTK_LOOP3STR(WT_CL_Standing,\n" - " i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2],\n" + " i,j,k, imin0,imin1,imin2, imax0,imax1,imax2,\n" " cctk_ash[0],cctk_ash[1],cctk_ash[2],\n" - " kimin,kimax, CCTK_REAL_VEC_SIZE)\n" + " vecimin,vecimax, CCTK_REAL_VEC_SIZE)\n" "{\n" - " ptrdiff_t /*const*/ index CCTK_ATTRIBUTE_UNUSED = di*i + dj*j + dk*k;\n" + " const ptrdiff_t index CCTK_ATTRIBUTE_UNUSED = di*i + dj*j + dk*k;\n" + " // vec_iter_counter+=CCTK_REAL_VEC_SIZE;\n" " \n" " /* Assign local copies of grid functions */\n" " \n" @@ -105,13 +112,13 @@ static void WT_CL_Standing_Body(cGH const * restrict const cctkGH, int const dir " ToReal(3.14159265358979323846264338328*ScalarINV(width));\n" " \n" " CCTK_REAL_VEC CCTK_ATTRIBUTE_UNUSED omega = \n" - " ksqrt(kmul(kmul(kvec,kvec),ToReal(3.)));\n" + " ksqrt(kmul(kmul(kvec,kvec),ToReal(3)));\n" " \n" " CCTK_REAL_VEC CCTK_ATTRIBUTE_UNUSED uL = \n" " kmul(kcos(kmul(xL,kvec)),kmul(kcos(kmul(yL,kvec)),kmul(kcos(kmul(zL,kvec)),kmul(kcos(kmul(omega,t)),ToReal(amplitude)))));\n" " \n" " CCTK_REAL_VEC CCTK_ATTRIBUTE_UNUSED rhoL = \n" - " kmul(omega,kmul(kcos(kmul(xL,kvec)),kmul(kcos(kmul(yL,kvec)),kmul(kcos(kmul(zL,kvec)),kmul(ksin(kmul(omega,t)),ToReal(-1.*amplitude))))));\n" + " kneg(kmul(omega,kmul(kcos(kmul(xL,kvec)),kmul(kcos(kmul(yL,kvec)),kmul(kcos(kmul(zL,kvec)),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" @@ -122,14 +129,14 @@ static void WT_CL_Standing_Body(cGH const * restrict const cctkGH, int const dir "" ; - char const *const groups[] = { + const char* 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}; + const char* const sources[] = {differencing, source, NULL}; OpenCLRunTime_CallKernel(cctkGH, CCTK_THORNSTRING, "WT_CL_Standing", sources, groups, NULL, NULL, NULL, -1, imin, imax, &kernel); @@ -152,7 +159,7 @@ extern "C" void WT_CL_Standing(CCTK_ARGUMENTS) return; } - const char *const groups[] = { + const char* const groups[] = { "grid::coordinates", "ML_WaveToy_CL::WT_rho", "ML_WaveToy_CL::WT_u"}; -- cgit v1.2.3