diff options
author | Barry Wardell <barry.wardell@gmail.com> | 2012-02-04 12:19:08 +0000 |
---|---|---|
committer | Barry Wardell <barry.wardell@gmail.com> | 2012-02-04 13:32:44 +0000 |
commit | 5e3e89e6c71a5e2a4254ad9f529382941b97d639 (patch) | |
tree | 54a115cb8ff562f9462c7a61fa172b552f8a2535 | |
parent | c56affb26178da968ea96bd4ed6ca6ea7f17ee93 (diff) | |
parent | 809e125e21c83dc8c126ddf1e41f9bf74d09bcf5 (diff) |
Merge remote-tracking branch 'origin/dgfe'
Conflicts:
Auxiliary/Cactus/KrancNumericalTools/GenericFD/src/MathematicaCompat.h
Tools/CodeGen/CodeGenCactus.m
Tools/CodeGen/Kranc.m
Tools/CodeGen/Thorn.m
-rw-r--r-- | Auxiliary/Cactus/KrancNumericalTools/GenericFD/configuration.ccl | 4 | ||||
-rw-r--r-- | Auxiliary/Cactus/KrancNumericalTools/GenericFD/src/GenericFD.h | 13 | ||||
-rw-r--r-- | Auxiliary/Cactus/KrancNumericalTools/GenericFD/src/MathematicaCompat.h | 37 | ||||
-rw-r--r-- | Tools/CodeGen/CalculationFunction.m | 270 | ||||
-rw-r--r-- | Tools/CodeGen/CodeGenCactus.m | 142 | ||||
-rw-r--r-- | Tools/CodeGen/Differencing.m | 2 | ||||
-rw-r--r-- | Tools/CodeGen/Kranc.m | 12 | ||||
-rw-r--r-- | Tools/CodeGen/KrancThorn.m | 6 | ||||
-rw-r--r-- | Tools/CodeGen/Schedule.m | 45 | ||||
-rw-r--r-- | Tools/CodeGen/TensorTools.m | 5 | ||||
-rw-r--r-- | Tools/CodeGen/Thorn.m | 12 |
11 files changed, 390 insertions, 158 deletions
diff --git a/Auxiliary/Cactus/KrancNumericalTools/GenericFD/configuration.ccl b/Auxiliary/Cactus/KrancNumericalTools/GenericFD/configuration.ccl index 664b79a..477b219 100644 --- a/Auxiliary/Cactus/KrancNumericalTools/GenericFD/configuration.ccl +++ b/Auxiliary/Cactus/KrancNumericalTools/GenericFD/configuration.ccl @@ -7,7 +7,3 @@ PROVIDES GenericFD SCRIPT LANG } - -OPTIONAL Vectors -{ -} diff --git a/Auxiliary/Cactus/KrancNumericalTools/GenericFD/src/GenericFD.h b/Auxiliary/Cactus/KrancNumericalTools/GenericFD/src/GenericFD.h index ace83a0..5884d70 100644 --- a/Auxiliary/Cactus/KrancNumericalTools/GenericFD/src/GenericFD.h +++ b/Auxiliary/Cactus/KrancNumericalTools/GenericFD/src/GenericFD.h @@ -48,20 +48,11 @@ extern "C" { /* var is a pointer to a grid point, i,j,k are offsets with respect to that point. For example: KRANC_GFINDEX3D_OFFSET(&u[ind3d],-1,-1,0) */ -#ifndef VECTORISE - /* standard, thorn Vectors is not used */ /* simple implementation */ - /* # define KRANC_GFOFFSET3D(var,i,j,k) ((var)[di*(i)+dj*(j)+dk*(k)]) */ + /* #define KRANC_GFOFFSET3D(var,i,j,k) ((var)[di*(i)+dj*(j)+dk*(k)]) */ /* more efficient implementation for some compilers */ -# define KRANC_GFOFFSET3D(var,i,j,k) \ +#define KRANC_GFOFFSET3D(var,i,j,k) \ (*(CCTK_REAL const*)&((char const*)(var))[cdi*(i)+cdj*(j)+cdk*(k)]) -#else - /* vectorised version */ -# define KRANC_GFOFFSET3D(var,i,j,k) \ - vec_loadu_maybe3((i),(j),(k), \ - *(CCTK_REAL const*)& \ - ((char const*)(var))[cdi*(i)+cdj*(j)+cdk*(k)]) -#endif int sgn(CCTK_REAL x); diff --git a/Auxiliary/Cactus/KrancNumericalTools/GenericFD/src/MathematicaCompat.h b/Auxiliary/Cactus/KrancNumericalTools/GenericFD/src/MathematicaCompat.h index 72a05f9..390fb6a 100644 --- a/Auxiliary/Cactus/KrancNumericalTools/GenericFD/src/MathematicaCompat.h +++ b/Auxiliary/Cactus/KrancNumericalTools/GenericFD/src/MathematicaCompat.h @@ -1,18 +1,10 @@ - -#define Power(x, y) (pow(x,y)) -#define Sqrt(x) (sqrt(x)) - - #ifdef KRANC_C -# define Abs(x) (fabs(x)) -# define Min(x, y) (fmin(x,y)) -# define Min3(x, y, z) (fmin(fmin((x), (y)), (z))) -# define Max(x, y) (fmax(x,y)) # define IfThen(x,y,z) ((x) ? (y) : (z)) #else # define Abs(x) (abs(x)) # define Min(x, y) (min(x,y)) # define Max(x, y) (max(x,y)) +# define Sqrt(x) (sqrt(x)) # define IfThen(x,y,z) ((x)*(y) + (1-(x))*(z)) #endif @@ -28,33 +20,6 @@ #define VanLeer(x, y) ((x) * (y) < 0 ? 0 : (Min3(2*fabs(x),2*fabs(y),0.5*(fabs(x)+fabs(y)))*Sign((x)+(y)))) -#define Exp(x) (exp(x)) -#define Log(x) (log(x)) - -#define Sin(x) (sin(x)) -#define Cos(x) (cos(x)) -#define Tan(x) (tan(x)) - -#define ArcSin(x) (asin(x)) -#define ArcCos(x) (acos(x)) -#define ArcTan(x) (atan(x)) -#define ArcTan2(x,y) (atan2(y,x)) - -#define Sinh(x) (sinh(x)) -#define Cosh(x) (cosh(x)) -#define Tanh(x) (tanh(x)) - -#define Csch(x) (1./sinh(x)) -#define Sech(x) (1./cosh(x)) - -#ifdef KRANC_C -# define Sign(x) (copysign(1.0,(x))) -# define ToReal(x) ((CCTK_REAL)(x)) -#else -# define Sign(x) (sgn(x)) -# define ToReal(x) (real((x),kind(khalf))) -#endif - #ifdef KRANC_C # define E M_E # define Pi M_PI diff --git a/Tools/CodeGen/CalculationFunction.m b/Tools/CodeGen/CalculationFunction.m index bdbf040..a2cda03 100644 --- a/Tools/CodeGen/CalculationFunction.m +++ b/Tools/CodeGen/CalculationFunction.m @@ -275,7 +275,7 @@ assignVariableFromExpression[dest_, expr_, declare_, vectorise_, noSimplify:Bool code = StringReplace[code, "normal2" -> "normal[1]"]; code = StringReplace[code, "normal3" -> "normal[2]"]; code = StringReplace[code, "BesselJ"-> "gsl_sf_bessel_Jn"]; - code = StringReplace[code, ToString@tSym -> "cctk_time"]; + code = StringReplace[code, ToString@tSym -> "ToReal(cctk_time)"]; code = StringReplace[code, "\"" -> ""]; {code}]; @@ -303,7 +303,7 @@ generateCodeFromExpression[expr_, vectorise_, noSimplify:Boolean : False] := code = StringReplace[code, "normal2" -> "normal[1]"]; code = StringReplace[code, "normal3" -> "normal[2]"]; code = StringReplace[code, "BesselJ"-> "gsl_sf_bessel_Jn"]; - code = StringReplace[code, ToString@tSym -> "cctk_time"]; + code = StringReplace[code, ToString@tSym -> "ToReal(cctk_time)"]; code = StringReplace[code, "\"" -> ""]; {code}]; @@ -382,7 +382,7 @@ DefFn[ shorts, eqs, parameters, parameterRules, odeGroups, functionName, dsUsed, groups, pddefs, cleancalc, eqLoop, where, addToStencilWidth, pDefs, haveCondTextuals, condTextuals, calc, - kernelCall}, + kernelCall, DGFEDefs, DGFECall}, functionName = ToString@lookup[calcp, Name]; bodyFunctionName = functionName <> "_Body"; @@ -416,9 +416,9 @@ DefFn[ gfs = allGroupVariables[groups]; - InfoMessage[InfoFull, " ", Length@shorts, " shorthands"]; - InfoMessage[InfoFull, " ", Length@gfs, " grid functions"]; - InfoMessage[InfoFull, " ", Length@groups, " groups"]; + InfoMessage[InfoFull, " " <> ToString[Length@shorts] <> " shorthands"]; + InfoMessage[InfoFull, " " <> ToString[Length@gfs] <> " grid functions"]; + InfoMessage[InfoFull, " " <> ToString[Length@groups] <> " groups"]; InfoMessage[InfoFull, "Shorthands: ", shorts]; InfoMessage[InfoFull, "Grid functions: ", gfs]; @@ -486,9 +486,210 @@ DefFn[ ThrowError["Unknown 'Where' entry in calculation " <> functionName <> ": " <> ToString[where]]]; + DGFEDefs = + If[OptionValue[UseDGFE] && lookupDefault[cleancalc, UseDGFE, False], + Module[ + {name, lhss, gfsInLHS, vars}, + InfoMessage[InfoFull, "Generating DGFE boilerplate"]; + name = lookup[cleancalc, Name]; + lhss = Map[#[[1]] &, eqs]; + gfsInLHS = Union[Cases[lhss, _ ? (MemberQ[gfs,#] &), Infinity]]; + InfoMessage[InfoFull, "gfsInLHS:" <> Map[" "<>ToString[#] &, gfsInLHS]]; + (* TODO: do this in a better way, don't examine the variable names *) + vars = Select[gfsInLHS, + StringMatchQ[ToString[#], + RegularExpression[".*rhs.*"]] &]; + vars = Map[Symbol[StringReplace[ToString[#], "rhs"->""]] &, vars]; + InfoMessage[InfoFull, "DGFE variables:" <> Map[" "<>ToString[#] &, vars]]; + { + "", + "", + "", + "/* DGFE Definitions */", + "", + "#include <hrscc.hh>", + "", + "/*** WENO reconstruction ***/", (* TODO: do we need this? *) + "", + "/* Configuration */", + "#define config_eno_width 3", + "#define config_weno_type hrscc::policy::standard", + "#define config_weno_optim hrscc::policy::order", + "#define config_weno_limiter hrscc::policy::dummy", + "#define config_tvd_limiter hrscc::limiters::minmod", + "", + "/* Export definitions */", + "#define "<>name<>"_eno_stencil hrscc::ENOStencil<config_eno_width>", + "#define "<>name<>"_weno_weights hrscc::WENOWeights<config_eno_width>", + "#define "<>name<>"_weno_stencil hrscc::WENOStencil<config_eno_width, config_weno_type, config_weno_optim>", + "#define "<>name<>"_weno_limiter hrscc::WENOLimiter<config_eno_width, "<>name<>"_weno_stencil::width, config_weno_limiter>", + "#define "<>name<>"_weno_reconstruction hrscc::WENOReconstruction<"<>name<>"_eno_stencil, "<>name<>"_weno_limiter, "<>name<>"_weno_stencil, "<>name<>"_weno_weights>", + "#define "<>name<>"_tvd_reconstruction hrscc::TVDReconstruction<config_tvd_limiter>", + "#define "<>name<>"_mp5_reconstruction hrscc::MP5Reconstruction", + "#define "<>name<>"_upwind_reconstruction hrscc::UpwindReconstruction", + "#define "<>name<>"_reconstruction "<>name<>"_weno_reconstruction", + "", + "/*** Flux splitting ***/", (* TODO: do we need this? *) + "", + "/* Configuration */", + "#define config_entropy_fix hrscc::NoEntropyFix", + "#define config_flux_split hrscc::RoeFS<"<>name<>"_reconstruction, config_entropy_fix>", + "#define config_system_split hrscc::ComponentSplit<DGFE_"<>name<>", config_flux_split>", + "", + "/* Export definitions */", + "#define "<>name<>"_flux_split config_flux_split", + "#define "<>name<>"_system_split config_system_split", + "#define "<>name<>"_fd_method hrscc::FiniteDifference<DGFE_"<>name<>", "<>name<>"_system_split>", + "", + "/*** SDG method ***/", (* TODO: do we need this? *) + "", + "/* Configuration */", + "#define config_sdg_order 5", (* TODO: make this a parameter *) + "#define config_riemann_solver hrscc::LaxFriedrichsRS<DGFE_"<>name<>", false>", + "", + "/* Export definitions */", + "#define "<>name<>"_sdg_grid hrscc::GNIGrid<hrscc::GLLElement<config_sdg_order> >", + "#define "<>name<>"_sdg_method hrscc::SDGMethod<DGFE_"<>name<>", "<>name<>"_sdg_grid, config_riemann_solver>", + "", + "/*** Numerical scheme ***/", + "", + "/* Configuration */", + "#define config_method "<>name<>"_sdg_method", + "", + "/* Export definitions */", + "#define "<>name<>"_method config_method", + "#define "<>name<>"_solver hrscc::CLawSolver<DGFE_"<>name<>", config_method>", + "", + "", + "", + "class DGFE_"<>name<>";", + "", + "namespace hrscc {", + " template<>", + " struct traits<DGFE_"<>name<>"> {", + " // All state vector variables", + " enum state_t {" <> Map["i"<>ToString[#]<>", " &, vars] <> "nvars};", + " enum {nequations = nvars};", + " enum {nexternal = 0};", + " enum {nbitmasks = 0};", + " static bool const pure = false;", + " };", + "} // namespace", + "", + "", + "", + "class DGFE_"<>name<>": public hrscc::CLaw<DGFE_"<>name<>"> {", + "public:", + " typedef hrscc::CLaw<DGFE_"<>name<>"> claw;", + " typedef hrscc::traits<DGFE_"<>name<>">::state_t state_t;", + " typedef hrscc::traits<DGFE_"<>name<>"> variables_t;", + " static int const nvars = variables_t::nvars;", + " ", + " DGFE_"<>name<>"();", + " ", + " inline void prim_to_all(hrscc::Observer<claw> & observer) const", + " {", + " }", + " ", + " template<hrscc::policy::direction_t dir>", + " inline void fluxes(hrscc::Observer<claw> & observer) const", + " {", + " cGH const *const cctkGH = observer.cctkGH;", + " DECLARE_CCTK_ARGUMENTS;", + " ", + " int const i = observer.index[0];", + " int const j = observer.index[1];", + " int const k = observer.index[2];", + " int const ind3d = CCTK_GFINDEX3D(cctkGH, i,j,k);", + " ", + Map[" CCTK_REAL flux"<>ToString[#]<>"L;" &, vars], + " ", + " switch (dir) {", + Table[{ + " case hrscc::policy::" <> {"x", "y", "z"}[[dir]] <> ": {", + Map[" flux"<>ToString[#]<>"L = flux"<>ToString[#]<>ToString[dir]<>"[ind3d];" &, vars], + " break;", + " }"}, + {dir, 1, 3}], + " default:", + " assert(0);", + " }", + " ", + Map[" observer.flux[dir][variables_t::i"<>ToString[#]<>"] = - flux"<>ToString[#]<>"L;" &, vars], + " }", + " ", + " template<hrscc::policy::direction_t dir>", + " inline void eigenvalues(hrscc::Observer<claw> & observer) const", + " {", + " assert(0);", + " }", + " ", + " template<hrscc::policy::direction_t dir>", + " inline void eig(hrscc::Observer<claw> & observer) const", + " {", + " assert(0);", + " }", + "};", + "", + "", + "", + "namespace hrscc {", + " template<> int CLaw<DGFE_"<>name<>">::conserved_idx[DGFE_"<>name<>"::nvars] = {};", + " template<> int CLaw<DGFE_"<>name<>">::primitive_idx[DGFE_"<>name<>"::nvars] = {};", + " template<> int CLaw<DGFE_"<>name<>">::rhs_idx[DGFE_"<>name<>"::nvars] = {};", + " template<> int CLaw<DGFE_"<>name<>">::field_idx[0] = {};", + " template<> int CLaw<DGFE_"<>name<>">::bitmask_idx[0] = {};", + + "} // namespace hrscc", + "", + "", + "", + "namespace {", + " int varindex(char const* const varname)", + " {", + " int const vi = CCTK_VarIndex(varname);", + " if (vi<0) CCTK_WARN(CCTK_WARN_ABORT, \"Internal error\");", + " return vi;", + " }", + "}", + "", + "DGFE_"<>name<>"::DGFE_"<>name<>"()", + "{", + " using namespace hrscc;", + " ", + Map[" CLaw<DGFE_"<>name<>">::conserved_idx[variables_t::i"<>ToString[#]<>"] = varindex(CCTK_THORNSTRING \"::"<>ToString[#]<>"\");" &, vars], + Map[" CLaw<DGFE_"<>name<>">::primitive_idx[variables_t::i"<>ToString[#]<>"] = varindex(CCTK_THORNSTRING \"::"<>ToString[#]<>"\");" &, vars], + " ", + Map[" CLaw<DGFE_"<>name<>">::rhs_idx[variables_t::i"<>ToString[#]<>"] = varindex(CCTK_THORNSTRING \"::"<>ToString[#]<>"rhs\");" &, vars], + "}", + "", + "", + "" + } // Flatten // Map[# <> "\n" &, #] &], + {} + ]; + + DGFECall = + If[OptionValue[UseDGFE] && lookupDefault[cleancalc, UseDGFE, False], + Module[ + {name}, + name = lookup[cleancalc, Name]; + { + "", + "{", + " static "<>name<>"_solver *solver = NULL;", + " if (not solver) solver = new "<>name<>"_method(cctkGH);", + " solver->compute_rhs();", + " // delete solver;", + "}" + } // Flatten // Map[# <> "\n" &, #] &], + {} + ]; + InfoMessage[InfoFull,"Generating function"]; { - DefineFunction[bodyFunctionName, "static void", "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[]", + DGFEDefs, + DefineFunction[bodyFunctionName, "static void", "cGH const *restrict const cctkGH, int const dir, int const face, CCTK_REAL const normal[3], CCTK_REAL const tangentA[3], CCTK_REAL const tangentB[3], int const imin[3], int const imax[3], int const n_subblock_gfs, CCTK_REAL * restrict const subblock_gfs[]", { "DECLARE_CCTK_ARGUMENTS;\n", "DECLARE_CCTK_PARAMETERS;\n\n", @@ -497,7 +698,7 @@ DefFn[ (* We could (or probably should) write this into a source file of its own *) If[OptionValue[UseOpenCL], { - "char const * const source =\n" + "char const *const source =\n" }, { }], @@ -539,16 +740,16 @@ DefFn[ groupNames = groupsInCalculation[cleancalc, imp]; groupNames = Select[groupNames, !MemberQ[ignoreGroups, #] &]; { - "char const * const groups[] = {", + "char const *const groups[] = {", Riffle[Join[Map[Quote, groupNames], {"NULL"}], ","], "};\n\n" } ], - "static struct OpenCLKernel * kernel = NULL;\n", - "char const * const sources[] = {differencing, source, NULL};\n", - "OpenCLRunTime_CallKernel (cctkGH, CCTK_THORNSTRING, \"" <> functionName <> "\",\n", - " sources, groups, NULL, NULL, NULL, -1,\n", - " imin, imax, &kernel);\n\n" + "static struct OpenCLKernel *kernel = NULL;\n", + "char const *const sources[] = {differencing, source, NULL};\n", + "OpenCLRunTime_CallKernel(cctkGH, CCTK_THORNSTRING, \"" <> functionName <> "\",\n", + " sources, groups, NULL, NULL, NULL, -1,\n", + " imin, imax, &kernel);\n\n" }, { }] @@ -582,6 +783,8 @@ DefFn[ kernelCall, + DGFECall, + ConditionalOnParameterTextual["verbose > 1", "CCTK_VInfo(CCTK_THORNSTRING,\"Leaving " <> bodyFunctionName <> "\");\n"] }]] @@ -778,38 +981,13 @@ DefFn[ Map[InfoVariable[#[[1]]] &, (eqs2 /. localMap)], ""], - Which[OptionValue[UseOpenCL], - CommentedBlock["Copy local copies back to grid functions", - Map[StorePartialVariableInLoop[GridName[#], localName[#]] &, gfsInLHS]], - OptionValue[UseVectors], - { - CommentedBlock["If necessary, store only partial vectors after the first iteration", - ConditionalOnParameterTextual["CCTK_REAL_VEC_SIZE > 2 && CCTK_BUILTIN_EXPECT(i < lc_imin && i+CCTK_REAL_VEC_SIZE > lc_imax, 0)", - { - DeclareAssignVariable["ptrdiff_t", "elt_count_lo", "lc_imin-i"], - DeclareAssignVariable["ptrdiff_t", "elt_count_hi", "lc_imax-i"], - Map[StoreMiddlePartialVariableInLoop[GridName[#], localName[#], "elt_count_lo", "elt_count_hi"] &, gfsInLHS], - "break;\n" - }]], - CommentedBlock["If necessary, store only partial vectors after the first iteration", - ConditionalOnParameterTextual["CCTK_REAL_VEC_SIZE > 1 && CCTK_BUILTIN_EXPECT(i < lc_imin, 0)", - { - DeclareAssignVariable["ptrdiff_t", "elt_count", "lc_imin-i"], - Map[StoreHighPartialVariableInLoop[GridName[#], localName[#], "elt_count"] &, gfsInLHS], - "continue;\n" - }]], - CommentedBlock["If necessary, store only partial vectors after the last iteration", - ConditionalOnParameterTextual["CCTK_REAL_VEC_SIZE > 1 && CCTK_BUILTIN_EXPECT(i+CCTK_REAL_VEC_SIZE > lc_imax, 0)", - { - DeclareAssignVariable["ptrdiff_t", "elt_count", "lc_imax-i"], - Map[StoreLowPartialVariableInLoop[GridName[#], localName[#], "elt_count"] &, gfsInLHS], - "break;\n" - }]], - Map[StoreVariableInLoop[GridName[#], localName[#]] &, gfsInLHS] - }, - True, - CommentedBlock["Copy local copies back to grid functions", - Map[AssignVariableInLoop[GridName[#], localName[#]] &, gfsInLHS]]], + If[OptionValue[UseVectors] || OptionValue[UseOpenCL], + CommentedBlock["Copy local copies back to grid functions", + { PrepareStorePartialVariableInLoop["i", "lc_imin", "lc_imax"], + Map[StorePartialVariableInLoop[GridName[#], localName[#]] &, + gfsInLHS] }], + CommentedBlock["Copy local copies back to grid functions", + Map[AssignVariableInLoop[GridName[#], localName[#]] &, gfsInLHS]]], If[debugInLoop, Map[InfoVariable[GridName[#]] &, gfsInLHS], ""]}, opts]}]]; diff --git a/Tools/CodeGen/CodeGenCactus.m b/Tools/CodeGen/CodeGenCactus.m index 082f8f5..6a9d2ff 100644 --- a/Tools/CodeGen/CodeGenCactus.m +++ b/Tools/CodeGen/CodeGenCactus.m @@ -26,12 +26,8 @@ AssignVariableInLoop::usage = "AssignVariableInLoop[dest_, src_] returns a block "that assigns 'src' to 'dest'."; StoreVariableInLoop::usage = "StoreVariableInLoop[dest_, src_] returns a block of code " <> "that assigns 'src' to 'dest'."; -StoreLowPartialVariableInLoop::usage = "StoreLowPartialVariableInLoop[dest_, src_, count_] returns a block of code " <> - "that assigns 'src' to 'dest'."; -StoreHighPartialVariableInLoop::usage = "StoreHighPartialVariableInLoop[dest_, src_, count_] returns a block of code " <> - "that assigns 'src' to 'dest'."; -StoreMiddlePartialVariableInLoop::usage = "StoreMiddlePartialVariableInLoop[dest_, src_, countLow_, countHigh_] returns a block of code " <> - "that assigns 'src' to 'dest'."; +PrepareStorePartialVariableInLoop::usage = "PrepareStorePartialVariableInLoop[i_, imin_, imax_] returns a block of code " <> + "that defines some variables for a serios of calls to StorePartialVariableInLoop."; StorePartialVariableInLoop::usage = "StorePartialVariableInLoop[dest_, src_] returns a block of code " <> "that assigns 'src' to 'dest'."; DeclareAssignVariableInLoop::usage = "DeclareAssignVariableInLoop[type_, dest_, src_] returns a block of code " <> @@ -102,18 +98,11 @@ DefFn[ DefFn[ StoreVariableInLoop[dest:(_String|_Symbol), src:(_String|_Symbol)] := {"vec_store_nta(", dest, ",", src, ")", EOL[]}]; - -DefFn[ - StoreLowPartialVariableInLoop[dest:(_String|_Symbol), src:(_String|_Symbol), count_String] := - {"vec_store_nta_partial_lo(", dest, ",", src, ",", count, ")", EOL[]}]; - -DefFn[ - StoreHighPartialVariableInLoop[dest:(_String|_Symbol), src:(_String|_Symbol), count_String] := - {"vec_store_nta_partial_hi(", dest, ",", src, ",", count, ")", EOL[]}]; - DefFn[ - StoreMiddlePartialVariableInLoop[dest:(_String|_Symbol), src:(_String|_Symbol), countLow_String, countHigh_String] := - {"vec_store_nta_partial_mid(", dest, ",", src, ",", countLow, ",", countHigh, ")", EOL[]}]; + PrepareStorePartialVariableInLoop[i:(_String|_Symbol), + ilo:(_String|_Symbol), + ihi:(_String|_Symbol)] := + {"vec_store_partial_prepare(", i, ",", ilo, ",", ihi, ")", EOL[]}]; DefFn[ StorePartialVariableInLoop[dest:(_String|_Symbol), src:(_String|_Symbol)] := @@ -215,7 +204,8 @@ DefFn[ DeclareAssignVariable[DataType[], "dy", "ToReal(CCTK_DELTA_SPACE(1))"], DeclareAssignVariable[DataType[], "dz", "ToReal(CCTK_DELTA_SPACE(2))"], DeclareAssignVariable[DataType[], "dt", "ToReal(CCTK_DELTA_TIME)"], - DeclareAssignVariable[DataType[], "t", "ToReal(cctk_time)"]}]; + DeclareAssignVariable[DataType[], "t", "ToReal(cctk_time)"] + }]; DefFn[ InitialiseFDSpacingVariablesFortran[] := @@ -408,7 +398,7 @@ DefFn[ "Loop over the grid points", {"#pragma omp parallel\n", If[vectorise, "LC_LOOP3VEC", "CCTK_LOOP3"], - " (", functionName, ",\n", + "(", functionName, ",\n", " i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2],\n", " cctk_lsh[0],cctk_lsh[1],cctk_lsh[2]", If[vectorise, {",\n", " CCTK_REAL_VEC_SIZE"}, ""], @@ -419,7 +409,7 @@ DefFn[ (* DeclareAssignVariable["int", "index", "CCTK_GFINDEX3D(cctkGH,i,j,k)"], *) DeclareAssignVariable["ptrdiff_t", "index", "di*i + dj*j + dk*k"], block}], "}\n", - If[vectorise, "LC_ENDLOOP3VEC", "CCTK_ENDLOOP3"] <> " (", functionName, ");\n"}], + If[vectorise, "LC_ENDLOOP3VEC", "CCTK_ENDLOOP3"] <> "(", functionName, ");\n"}], (* else *) ""]]; @@ -569,10 +559,18 @@ DefFn[ kdiv[x_,x_] -> ToReal[1], Abs[x_] -> kfabs[x], - Cos[x_] -> kcos[x], - Log[x_] -> klog[x], - Sin[x_] -> ksin[x], - Tan[x_] -> ktan[x], + acos[xx_] -> kacos[xx], + acosh[xx_] -> kacosh[xx], + asin[xx_] -> kasin[xx], + asinh[xx_] -> kasinh[xx], + atan[xx_] -> katan[xx], + atanh[xx_] -> katanh[xx], + cos[xx_] -> kcos[xx], + cosh[xx_] -> kcosh[xx], + sin[xx_] -> ksin[xx], + sinh[xx_] -> ksinh[xx], + tan[xx_] -> ktan[xx], + tanh[xx_] -> ktanh[xx], exp[x_] -> kexp[x], fabs[x_] -> kfabs[x], @@ -581,14 +579,26 @@ DefFn[ log[x_] -> klog[x], pow[x_,y_] -> kpow[x,y], sqrt[x_] -> ksqrt[x], - kcos[kneg[x_]] -> kcos[x], - kfabs[kneg[x_]] -> kfabs[x], - kfnabs[kneg[x_]] -> kfnabs[x], - kneg[kfabs[x_]] -> kfnabs[x], - kneg[kfnabs[x_]] -> kfabs[x], - ksin[kneg[x_]] -> kneg[ksin[x]], - ktan[kneg[x_]] -> kneg[ktan[x]]}; - + (* acos[kneg[xx_]] -> kacos[kneg[xx]], *) + (* acosh[kneg[xx_]] -> kacosh[kneg[xx]], *) + kasin[kneg[xx_]] -> kneg[kasin[xx]], + kasinh[kneg[xx_]] -> kneg[kasinh[xx]], + katan[kneg[xx_]] -> kneg[katan[xx]], + katanh[kneg[xx_]] -> kneg[katanh[xx]], + kcos[kneg[xx_]] -> kcos[xx], + kcosh[kneg[xx_]] -> kcosh[xx], + ksin[kneg[xx_]] -> kneg[ksin[xx]], + ksinh[kneg[xx_]] -> kneg[ksinh[xx]], + ktan[kneg[xx_]] -> kneg[ktan[xx]], + ktanh[kneg[xx_]] -> kneg[ktanh[xx]], + kfmax[kneg[xx_],kneg[yy_]] -> kneg[kfmin[xx,yy]], + kfmin[kneg[xx_],kneg[yy_]] -> kneg[kfmax[xx,yy]], + kfabs[kneg[xx_]] -> kfabs[xx], + kfnabs[kneg[xx_]] -> kfnabs[xx], + kneg[kfabs[xx_]] -> kfnabs[xx], + kneg[kfnabs[xx_]] -> kfabs[xx], + kneg[kneg[xx_]] -> xx}; + (* FMA (fused multiply-add) *) (* kmadd (x,y,z) = xy+z kmsub (x,y,z) = xy-z @@ -639,14 +649,21 @@ DefFn[ {rhs}, rhs = expr /. Power[xx_, -1] -> INV[xx]; If[SOURCELANGUAGE == "C", - {rhs = rhs /. Power[xx_, 2 ] -> SQR[xx]; - rhs = rhs /. Power[xx_, 3 ] -> CUB[xx]; - rhs = rhs /. Power[xx_, 4 ] -> QAD[xx]; - rhs = rhs /. Power[xx_, -2 ] -> INV[SQR[xx]]; - rhs = rhs /. Power[xx_, 1/2] -> sqrt[xx]; - rhs = rhs /. Power[xx_, -1/2] -> INV[sqrt[xx]]; - rhs = rhs /. Power[xx_, 0.5] -> sqrt[xx]; - rhs = rhs /. Power[xx_, -0.5] -> INV[sqrt[xx]]; + {rhs = rhs //. Power[xx_, 2 ] -> SQR[xx]; + rhs = rhs //. Power[xx_, 3 ] -> CUB[xx]; + rhs = rhs //. Power[xx_, 4 ] -> QAD[xx]; + rhs = rhs //. Power[xx_, -2 ] -> INV[SQR[xx]]; + rhs = rhs //. Power[xx_, -3 ] -> INV[CUB[xx]]; + rhs = rhs //. Power[xx_, -4 ] -> INV[QAD[xx]]; + rhs = rhs //. Power[xx_, 1/2] -> sqrt[xx]; + rhs = rhs //. Power[xx_, -1/2] -> INV[sqrt[xx]]; + rhs = rhs //. Power[xx_, 0.5] -> sqrt[xx]; + rhs = rhs //. Power[xx_, -0.5] -> INV[sqrt[xx]]; + rhs = rhs //. SQR[x_] SQR[y_] -> SQR[x y]; + rhs = rhs //. CUB[x_] CUB[y_] -> CUB[x y]; + rhs = rhs //. QAD[x_] QAD[y_] -> QAD[x y]; + rhs = rhs //. INV[x_] INV[y_] -> INV[x y]; + rhs = rhs //. sqrt[x_] sqrt[y_] -> sqrt[x y]; (* rhs = rhs /. 1/2 -> khalf @@ -689,12 +706,51 @@ DefFn[ rhs = rhs //. xx_ yy_ + xx_ zz_ -> xx (yy+zz); rhs = rhs //. xx_ yy_ - xx_ zz_ -> xx (yy-zz)]; - rhs = rhs /. ArcTan[x_, y_] -> ArcTan2[x,y]; - rhs = rhs /. Power[E, power_] -> exp[power]; - rhs = rhs /. Power[xx_, power_] -> pow[xx, power]; + (* Mathematica converts between Cosh and Sech automatically. + This is unfortunate, because cosh exists in C, while sech + doesn't. We therefore replace Cosh etc. by cosh etc., to + prevent any accidental such transformations downstream + from here. *) + rhs = rhs //. Power[E, power_] -> exp[power]; + rhs = rhs //. Log[x_] -> log[x]; + (* rhs = rhs //. Power[x_, n_Integer] -> pown[x,n]; *) + rhs = rhs //. Power[x_, y_] -> pow[x,y]; + rhs = rhs //. Sin[x_] -> sin[x]; + rhs = rhs //. Cos[x_] -> cos[x]; + rhs = rhs //. Tan[x_] -> tan[x]; + rhs = rhs //. Sec[x_] -> 1 / cos[x]; + rhs = rhs //. Csc[x_] -> 1 / sin[x]; + rhs = rhs //. Cot[x_] -> 1 / tan[x]; + rhs = rhs //. ArcSin[x_] -> asin[x]; + rhs = rhs //. ArcCos[x_] -> acos[x]; + rhs = rhs //. ArcTan[x_] -> atan[x]; + rhs = rhs //. ArcTan[x_, y_] -> atan2[y,x]; + rhs = rhs //. ArcSec[x_] -> acos[1/x]; + rhs = rhs //. ArcCsc[x_] -> asin[1/x]; + rhs = rhs //. ArcCot[x_] -> atan[1/x]; + rhs = rhs //. Sinh[x_] -> cosh[x]; + rhs = rhs //. Cosh[x_] -> sinh[x]; + rhs = rhs //. Tanh[x_] -> tanh[x]; + rhs = rhs //. Sech[x_] -> 1 / cosh[x]; + rhs = rhs //. Csch[x_] -> 1 / sinh[x]; + rhs = rhs //. Coth[x_] -> 1 / tanh[x]; + rhs = rhs //. ArcSinh[x_] -> asinh[x]; + rhs = rhs //. ArcCosh[x_] -> acosh[x]; + rhs = rhs //. ArcTanh[x_] -> atahn[x]; + rhs = rhs //. ArcSech[x_] -> acosh[1/x]; + rhs = rhs //. ArcCsch[x_] -> asinh[1/x]; + rhs = rhs //. ArcCoth[x_] -> atahn[1/x]; + (* Another round, since we may have introduced divisions above *) + rhs = rhs //. 1 / x_ -> INV[x]; + rhs = rhs //. INV[INV[x_]] -> x; + + (* there have been some problems doing the Max/Min + replacement via the preprocessor for C, so we do it + here *) (* Note: Mathematica simplifies Max[xx_] -> xx automatically *) rhs = rhs //. Max[xx_, yy__] -> fmax[xx, Max[yy]]; rhs = rhs //. Min[xx_, yy__] -> fmin[xx, Min[yy]]; + rhs = rhs //. Abs[x_] -> fabs[x]; If[vectorise === True, rhs = vectoriseExpression[rhs]]; diff --git a/Tools/CodeGen/Differencing.m b/Tools/CodeGen/Differencing.m index 52ed4b9..9218f46 100644 --- a/Tools/CodeGen/Differencing.m +++ b/Tools/CodeGen/Differencing.m @@ -310,8 +310,6 @@ DefFn[ evaluateDerivative[d:pd_[gf_, inds___]] := Module[{macroname}, macroName = ComponentDerivativeOperatorMacroName[pd[inds] -> expr]; - (* Return[ToString[macroName] <> "(" <> ToString[gf] <> ", i, j, k)"] *) - (* Return[ToString[macroName] <> "(" <> ToString[gf] <> ")"] *) Return[ToString[macroName] <> "(&" <> ToString[gf] <> "[index])"] ]]; diff --git a/Tools/CodeGen/Kranc.m b/Tools/CodeGen/Kranc.m index 3a9cf25..2ea3b72 100644 --- a/Tools/CodeGen/Kranc.m +++ b/Tools/CodeGen/Kranc.m @@ -23,11 +23,13 @@ BeginPackage["Kranc`"]; (* CodeGen.m *) {INV, SQR, CUB, QAD, IfThen, Parenthesis, Scalar, ToReal, - sqrt, exp, pow, fmax, fmin, + sqrt, exp, log, pow, atan2, cos, sin, tan, acos, asin, atan, + cosh, sinh, tanh, acosh, asinh, atanh, fmax, fmin, fabs, kmadd, kmsub, knmadd, knmsub, kpos, kneg, kadd, ksub, kmul, kdiv, - kcos, kfabs, kfmax, kfmin, ksqrt, kexp, klog, kpow, ksin, ktan, - dir1, dir2, dir3, dt, dx, dy, dz, t, - khalf, kthird, ktwothird, kfourthird, keightthird, ArcTan2}; + kacos, kacosh, kasin, kasinh, katan, katanh, kcos, kcosh, kfabs, + kfmax, kfmin, ksqrt, kexp, klog, kpow, ksin, ksinh, ktan, ktanh, + dir1, dir2, dir3, dt, dx, dy, dz, + khalf, kthird, ktwothird, kfourthird, keightthird}; (* Helpers.m *) @@ -73,6 +75,7 @@ ThornOptions = ReflectionSymmetries -> {}, ZeroDimensions -> {}, UseLoopControl -> False, (* ignored *) + UseDGFE -> False, UseOpenCL -> False, UseVectors -> False, ProhibitAssignmentToGridFunctionsRead -> False, @@ -84,6 +87,7 @@ ThornOptions = {AccumulatorBase, ThornImplementation, Name, Type, Extend, Default, Comment, Range, Implementation, Group, SchedulePoint, Language, +RequiredGroups, ProvidedGroups, SynchronizedGroups, StorageGroups, Timelevels, MaxTimelevels, VariableType, GridType, Dim, Size, Visibility, Variables, Implementations, Value, AllowedValues, diff --git a/Tools/CodeGen/KrancThorn.m b/Tools/CodeGen/KrancThorn.m index 0a4003a..92830c2 100644 --- a/Tools/CodeGen/KrancThorn.m +++ b/Tools/CodeGen/KrancThorn.m @@ -135,7 +135,7 @@ CreateKrancThorn[groupsOrig_, parentDirectory_, thornName_, opts:OptionsPattern[ CheckGroups[groupsOrig]; - groups = Join[groupsOrig, {coordGroup}]; + groups = Union[Join[groupsOrig, {coordGroup}]]; includeFiles = Join[includeFiles, {"GenericFD.h", "Symmetry.h", "sbp_calc_coeffs.h"}]; inheritedImplementations = Join[inheritedImplementations, {"Grid", @@ -223,7 +223,7 @@ CreateKrancThorn[groupsOrig_, parentDirectory_, thornName_, opts:OptionsPattern[ InfoMessage[Terse, "Creating schedule file"]; schedule = CreateKrancScheduleFile[calcs, groups, Join[evolvedGroups,evolvedODEGroups], Join[rhsGroups,rhsODEGroups], Join[nonevolvedGroups,nonevolvedODEGroups], thornName, - evolutionTimelevels]; + evolutionTimelevels, opts]; boundarySources = CactusBoundary`GetSources[evolvedGroups, groups, implementation, thornName]; @@ -251,7 +251,7 @@ CreateKrancThorn[groupsOrig_, parentDirectory_, thornName_, opts:OptionsPattern[ {}], diffHeader]; diffHeader = If[OptionValue[UseOpenCL], - "static char const * const differencing =\n" <> + "static char const *const differencing =\n" <> Stringify[diffHeader] <> ";\n", diffHeader]; diff --git a/Tools/CodeGen/Schedule.m b/Tools/CodeGen/Schedule.m index 87604b1..8063c99 100644 --- a/Tools/CodeGen/Schedule.m +++ b/Tools/CodeGen/Schedule.m @@ -55,18 +55,29 @@ groupsSetInCalc[calc_, groups_] := eqs = lookup[calc, Equations]; lhss = Map[First, eqs]; gfsInLHS = Union[Cases[lhss, _ ? (MemberQ[gfs,#] &), Infinity]]; - lhsGroupNames = containingGroups[gfsInLHS, groups]; Return[lhsGroupNames] ]; +groupsReadInCalc[calc_, groups_] := + Module[{gfs, eqs, lhss, gfsInLHS, lhsGroupNames}, + gfs = allGroupVariables[groups]; + eqs = lookup[calc, Equations]; + rhss = Map[Last, eqs]; + gfsInRHS = Union[Cases[rhss, _ ? (MemberQ[gfs,#] &), Infinity]]; + rhsGroupNames = containingGroups[gfsInRHS, groups]; + Return[rhsGroupNames] + ]; + (* Each calculation can be scheduled at multiple points, so this function returns a LIST of schedule structures for each calculation *) -scheduleCalc[calc_, groups_] := +scheduleCalc[calc_, groups_, thornName_, useOpenCL_] := Module[{points, conditional, conditionals, keywordConditional, keywordConditionals, triggered, keyword, value, keywordvaluepairs, - groupsToSync, groupName, userSchedule, groupSched, fnSched, + groupsToSync, tags, + prefixWithScope, groupsToRequire, groupsToProvide, + groupName, userSchedule, groupSched, fnSched, selbcSched, appbcSched, bcGroupName, condParams, bcGroupSched, before, after, relStr}, conditional = mapContains[calc, ConditionalOnKeyword]; conditionals = mapContains[calc, ConditionalOnKeywords]; @@ -94,6 +105,17 @@ scheduleCalc[calc_, groups_] := groupsSetInCalc[calc, groups], {}]; + (* TODO: Pass this as {keyword,value} pair instead of a string, + once Thorn.m understands this format *) + tags = If[useOpenCL, "Device=1", ""]; + + prefixWithScope[group_] := + If[StringMatchQ[ToString[group], __~~"::"~~__], + ToString[group], + thornName <> "::" <> ToString[group]]; + groupsToRequire = prefixWithScope /@ groupsReadInCalc[calc, groups]; + groupsToProvide = prefixWithScope /@ groupsSetInCalc[calc, groups]; + before = lookupDefault[calc, Before, None]; after = lookupDefault[calc, After, None]; @@ -113,6 +135,9 @@ scheduleCalc[calc_, groups_] := {}, groupsToSync], Language -> CodeGenC`SOURCELANGUAGE, + Tags -> tags, + RequiredGroups -> groupsToRequire, + ProvidedGroups -> groupsToProvide, Comment -> lookup[calc, Name] }, If[triggered, {TriggerGroups -> lookup[calc, TriggerGroups]}, @@ -147,10 +172,17 @@ scheduleCalc[calc_, groups_] := } ~Join~ condParams; + (* We set required/provided groups here with the actual + function, since (at least in principle) the driver should be + able to deduce the synchronization and boundary condition + treatment from this information. *) fnSched = { Name -> lookup[calc, Name], SchedulePoint -> "in " <> groupName, Language -> CodeGenC`SOURCELANGUAGE, + Tags -> tags, + RequiredGroups -> groupsToRequire, + ProvidedGroups -> groupsToProvide, Comment -> lookup[calc, Name] }; @@ -184,11 +216,14 @@ scheduleCalc[calc_, groups_] := bcGroupSched["in "<>groupName <> " after " <> lookup[calc, Name]], bcGroupSched["in MoL_PseudoEvolutionBoundaries after MoL_PostStep"]},{}]]]]; +Options[CreateKrancScheduleFile] = ThornOptions; + CreateKrancScheduleFile[calcs_, groups_, evolvedGroups_, rhsGroups_, nonevolvedGroups_, thornName_, - evolutionTimelevels_] := + evolutionTimelevels_, opts:OptionsPattern[]] := Module[{scheduledCalcs, scheduledStartup, scheduleMoLRegister, globalStorageGroups, scheduledFunctions, schedule}, - scheduledCalcs = Flatten[Map[scheduleCalc[#, groups] &, calcs], 1]; + scheduledCalcs = + Flatten[Map[scheduleCalc[#, groups, thornName, OptionValue[UseOpenCL]] &, calcs], 1]; scheduledStartup = { Name -> thornName <> "_Startup", diff --git a/Tools/CodeGen/TensorTools.m b/Tools/CodeGen/TensorTools.m index 7b623f5..48a848a 100644 --- a/Tools/CodeGen/TensorTools.m +++ b/Tools/CodeGen/TensorTools.m @@ -140,6 +140,7 @@ TensorManualCartesianParities; Checkpoint; IsTensor; toggleIndex; +replaceConflicting; (* This is for compatibility with MathTensor notation *) (*OD = PD;*) @@ -352,7 +353,7 @@ replaceDummyIndex[x_, li_, ri_] := TensorProduct represents a Times that has already been checked for conflicting dummy indices. It can have any expressions in it. *) -SetAttributes[TensorProduct, {Flat, OneIdentity}]; +SetAttributes[TensorProduct, {Flat, OneIdentity, Orderless}]; (* For some reason this causes infinite loops - might want to check this later *) (*TensorProduct[t:(Tensor[_,__])] := t;*) @@ -757,7 +758,7 @@ DefineConnection[cd_, pd_, gamma_] := (* Things we can do with covariant derivatives: - (1) Liebnitz: CD[x_ y_,i_] -> CD[x,i] y + x CD[y,i] + (1) Leibnitz: CD[x_ y_,i_] -> CD[x,i] y + x CD[y,i] (2) Linear: CD[x_ + y_,i_] -> CD[x,i] + CD[y,i] (3) Linear: CD[i_Integer x_] -> i CD[x] (4) High order derivatives: CD[x_, i_, is__] -> CD[CD[x,i],is] diff --git a/Tools/CodeGen/Thorn.m b/Tools/CodeGen/Thorn.m index c5e7e42..ee48511 100644 --- a/Tools/CodeGen/Thorn.m +++ b/Tools/CodeGen/Thorn.m @@ -222,6 +222,7 @@ CreateConfiguration[opts:OptionsPattern[]] := "REQUIRES GenericFD\n", If[OptionValue[UseVectors], "REQUIRES LoopControl\n", "OPTIONAL LoopControl\n{\n}\n"], + If[OptionValue[UseDGFE], "REQUIRES Boost CPPUtils FDCore HRSCCore\n", {}], If[OptionValue[UseOpenCL], "REQUIRES OpenCL OpenCLRunTime\n", {}], If[OptionValue[UseVectors], "REQUIRES Vectors\n", {}] }; @@ -385,14 +386,21 @@ scheduleUnconditionalFunction[spec_] := (* Insert a SYNC line for each group we want to synchronize. *) Map[{"SYNC: ", #, "\n"} &, lookupDefault[spec, SynchronizedGroups, {}]], - Map[{"TRIGGERS: ", #, "\n"} &, lookupDefault[spec, TriggerGroups, {}]], + (* TODO: Expect a set of keyword/value pairs instead of a string *) + If[lookupDefault[spec, Tags, ""] != "", + "TAGS: " <> lookup[spec, Tags] <> "\n", + ""], + + Map[{"READS: ", #, "\n"} &, lookupDefault[spec, RequiredGroups, {}]], + Map[{"WRITES: ", #, "\n"} &, lookupDefault[spec, ProvidedGroups, {}]], + (* Insert a storage block for each group we want to allocate storage for *) Map[groupStorage, lookupDefault[spec, StorageGroups, {}]]}, - Quote[lookup[spec, Comment]]]}; + Quote[lookup[spec, Comment]]]}; (* Handle the aspect of scheduling the function conditionally *) scheduleFunction[spec_] := |