Support FunHPC parallelized loops

ianhinder · Jun 27, 2016 · 349fe2b · 349fe2b
1 parent 06f44ae
commit 349fe2b
Show file tree

Hide file tree

Showing 9 changed files with 104 additions and 20 deletions.
diff --git a/Auxiliary/Cactus/SourceFiles/Kranc.hh b/Auxiliary/Cactus/SourceFiles/Kranc.hh
@@ -98,17 +98,21 @@ void AssertGroupStorage(cGH const *restrict const cctkGH, const char *calc,
 /* simple implementation */
 /* #define KRANC_GFOFFSET3D(var,i,j,k) ((var)[di*(i)+dj*(j)+dk*(k)]) */
 /* more efficient implementation for some compilers */
+#ifndef KRANC_GFOFFSET3D
 #define KRANC_GFOFFSET3D(var, i, j, k)                                         \
   (*(CCTK_REAL const *)&(                                                      \
        (char const *)(var))[cdi * (i) + cdj * (j) + cdk * (k)])
+#endif
 
 #define GFOffset(u, di, dj, dk) KRANC_GFOFFSET3D(&(u)[index], di, dj, dk)
 
 /*********************************************************************
  * Macros used in Kranc expressions
  *********************************************************************/
 
+#ifndef IfThen
 #define IfThen(x, y, z) ((x) ? (y) : (z))
+#endif
 #define MinMod(x, y) ((x) * (y) < 0 ? 0 : (fabs((x)) < fabs((y)) ? (x) : (y)))
 #define VanLeer(x, y)                                                          \
   ((x) * (y) < 0 ? 0 : (Min3(2 * fabs(x), 2 * fabs(y),                         \

diff --git a/Tools/CodeGen/Calculation.m b/Tools/CodeGen/Calculation.m
@@ -48,7 +48,9 @@
 TileCalculationQ;
 DGTileCalculationQ;
 CalculationLoopControlQ;
+CalculationLoopControlQQ;
 SetCalculationLoopControl;
+SetCalculationLoopControlQ;
 
 Begin["`Private`"];
 
@@ -436,7 +438,7 @@
          LocalGroups, NoSimplify, UseDGFE, SimpleCode, UseCaKernel,
          UseJacobian,
          ScheduleGroups, TriggerGroups, ThornName, Tile, DGTile,
-         UseLoopControl};
+         UseLoopControl, UseLoopControlQ};
 
     usedKeys = Map[First, calc];
     unknownKeys = Complement[usedKeys, allowedKeys];
@@ -501,12 +503,21 @@ pathalogical enough (e.g. {s1 -> s2, s2 -> s1} would not be
 DefFn[CalculationLoopControlQ[calc_] :=
   lookup[calc,UseLoopControl]];
 
+DefFn[CalculationLoopControlQQ[calc_] :=
+  lookup[calc,UseLoopControlQ]];
+
 Options[SetCalculationLoopControl] = ThornOptions;
 DefFn[SetCalculationLoopControl[calc_, opts:OptionsPattern[]] :=
   mapReplaceAdd[calc, UseLoopControl,
     OptionValue[UseLoopControl] &&
     lookupDefault[calc, UseLoopControl, True] =!= False]];
 
+Options[SetCalculationLoopControlQ] = ThornOptions;
+DefFn[SetCalculationLoopControlQ[calc_, opts:OptionsPattern[]] :=
+  mapReplaceAdd[calc, UseLoopControlQ,
+    OptionValue[UseLoopControlQ] &&
+    lookupDefault[calc, UseLoopControlQ, True] =!= False]];
+
 End[];
 
 EndPackage[];
diff --git a/Tools/CodeGen/CodeGenCalculation.m b/Tools/CodeGen/CodeGenCalculation.m
@@ -104,6 +104,7 @@ General Utility Functions (could be moved outside this package)
                          If[!OptionValue[DGTile], {"Differencing.h"}, {}],
                          include,
                          If[CalculationLoopControlQ[calc], {"loopcontrol.h"},{}],
+                         If[CalculationLoopControlQQ[calc], {"LoopControlQ.hpp", "RHSFutures.hh"},{}],
                          If[OptionValue[UseDGFE], {"hrscc.hh"}, {}],
                          If[OptionValue[DGTile], {"StencilOps.hh"}, {}],
                          If[OptionValue[UseOpenCL], OpenCLIncludeFiles[], {}],

diff --git a/Tools/CodeGen/CodeGenConfiguration.m b/Tools/CodeGen/CodeGenConfiguration.m
@@ -39,6 +39,7 @@
     If[OptionValue[UseLoopControl],
       If[OptionValue[UseVectors], 
         "REQUIRES LoopControl\n", "OPTIONAL LoopControl\n{\n}\n"],{}],
+   If[OptionValue[UseLoopControlQ], "REQUIRES FunHPC\n", {}],
    If[OptionValue[UseDGFE], DGFEConfigurationCCL[], {}],
    If[OptionValue[UseOpenCL], OpenCLConfigurationCCL[], {}],
    If[OptionValue[UseVectors], VectorisationConfigurationCCL[], {}],

diff --git a/Tools/CodeGen/CodeGenKranc.m b/Tools/CodeGen/CodeGenKranc.m
@@ -145,20 +145,47 @@
        }],
      CommentedBlock[
        "Loop over the grid points",
-       { (* Circumvent a compiler bug on Blue Gene/Q *)
-         "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\n",
-         If[OptionValue[UseVectors], "CCTK_LOOP3STR", "CCTK_LOOP3"],
-         "(", functionName, ",\n",
-         "  i,j,k, imin0,imin1,imin2, imax0,imax1,imax2,\n",
-         "  cctk_ash[0],cctk_ash[1],cctk_ash[2]",
-         If[OptionValue[UseVectors], {",\n", "  vecimin,vecimax, CCTK_REAL_VEC_SIZE"}, ""],
-         ")\n",
+       { If[!OptionValue[UseLoopControlQ],
+         { (* Circumvent a compiler bug on Blue Gene/Q *)
+           "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\n",
+           If[OptionValue[UseVectors], "CCTK_LOOP3STR", "CCTK_LOOP3"],
+           "(", functionName, ",\n",
+           "  i,j,k, imin0,imin1,imin2, imax0,imax1,imax2,\n",
+           "  cctk_ash[0],cctk_ash[1],cctk_ash[2]",
+           If[OptionValue[UseVectors], {",\n", "  vecimin,vecimax, CCTK_REAL_VEC_SIZE"}, ""],
+           ")\n"
+         },
+         { (* LoopControlQ enabled *)
+           "typedef LoopControlQ::iarray<3> iarray3;\n",
+           "iarray3 iamin, iamax, iablock, iaash;\n",
+           "for (int d=0; d<3; ++d) {\n",
+           "  iamin[d] = imin[d];\n",
+           "  iamax[d] = imax[d];\n",
+           "  iaash[d] = cctk_ash[d];\n",
+           "}\n",
+           "iablock[0] = block_size_i;\n",
+           "iablock[1] = block_size_j;\n",
+           "iablock[2] = block_size_k;\n",
+           If[!OptionValue[UseVectors],
+              {
+                "rhs_loop(\n",
+                "  [=](const iarray3& ipos) {\n"
+              }, {
+                "rhs_loop_str<CCTK_REAL_VEC_SIZE>(\n",
+                "  [=](const iarray3& imin, const iarray3& imax, const iarray3& ipos) {\n",
+                "    const ptrdiff_t vecimin = imin[0];\n",
+                "    const ptrdiff_t vecimax = imax[0];\n"
+                 }],
+           "    const ptrdiff_t i = ipos[0];\n",
+           "    const ptrdiff_t j = ipos[1];\n",
+           "    const ptrdiff_t k = ipos[2];\n"
+         }],
          "{\n",
          IndentBlock[
            {DefineConstant["index", "ptrdiff_t", "di*i + dj*j + dk*k"],
@@ -177,8 +204,14 @@
                {}],
             block}],
          "}\n",
-         If[OptionValue[UseVectors], "CCTK_ENDLOOP3STR", "CCTK_ENDLOOP3"] <>
-         "(", functionName, ");\n" 
+         If[!OptionValue[UseLoopControlQ],
+         {
+           If[OptionValue[UseVectors], "CCTK_ENDLOOP3STR", "CCTK_ENDLOOP3"] <>
+           "(", functionName, ");\n"
+         }, {
+           "  },\n",
+           "  cctkGH, iamin, iamax, iablock, iaash);\n"
+         }]
        }]]];
 
 DefFn[
@@ -191,7 +224,7 @@
   InfoVariable[name_String] :=
   onceInGridLoop[
     {"char buffer[255];\n",
-     "sprintf(buffer,\"" , name , " == %f\", " , name , ");\n",
+     "snprintf(buffer, sizeof buffer, \"" , name , " == %f\", " , name , ");\n",
      "CCTK_INFO(buffer);\n"}]];
 
 (* Take an expression x and replace occurrences of Powers with the C

diff --git a/Tools/CodeGen/Interface.m b/Tools/CodeGen/Interface.m
@@ -103,6 +103,7 @@
     interface = Join[CreateInterface[implementation, inheritedImplementations,
       Join[includeFiles, {CactusBoundary`GetIncludeFiles[]},
            If[OptionValue[UseLoopControl], {"loopcontrol.h"}, {}],
+           If[OptionValue[UseLoopControlQ], {"LoopControlQ.hpp", "RHSFutures.hh"}, {}],
            If[OptionValue[UseOpenCL], OpenCLIncludeFiles[], {}],
            If[OptionValue[UseVectors], {"vectors.h"}, {}]],
       declaredGroupStructures,

diff --git a/Tools/CodeGen/Kranc.m b/Tools/CodeGen/Kranc.m
@@ -101,6 +101,7 @@
   ReflectionSymmetries -> {},
   ZeroDimensions -> {},
   UseLoopControl -> True,
+  UseLoopControlQ -> False,
   UseDGFE -> False,
   Tile -> False,
   DGTile -> False,
@@ -160,6 +161,7 @@
 (* KrancScript.m *)
 ScriptFlags =
   {"loopcontrol" -> UseLoopControl,
+   "loopcontrolq" -> UseLoopControlQ,
    "vectors" -> UseVectors,
    "opencl" -> UseOpenCL,
    "jacobian"-> UseJacobian,

diff --git a/Tools/CodeGen/KrancThorn.m b/Tools/CodeGen/KrancThorn.m
@@ -142,10 +142,19 @@
   loopControlProcessCode[cIn_Code, opts:OptionsPattern[]] :=
   Module[
     {},
-    If[OptionValue[UseVectors] && !OptionValue[UseLoopControl],
+    If[OptionValue[UseVectors] &&
+       !(OptionValue[UseLoopControl] || OptionValue[UseLoopControlQ]),
+      (* TODO: Is this still true? *)
       ThrowError["UseVectors -> True requires UseLoopControl -> True"]];
     SetObjectField[cIn, "Calculations", SetCalculationLoopControl[#,opts] & /@ GetObjectField[cIn, "Calculations"]]]];
 
+Options[loopControlQProcessCode] = ThornOptions;
+DefFn[
+  loopControlQProcessCode[cIn_Code, opts:OptionsPattern[]] :=
+  Module[
+    {},
+    SetObjectField[cIn, "Calculations", SetCalculationLoopControlQ[#,opts] & /@ GetObjectField[cIn, "Calculations"]]]];
+
 
 (* --------------------------------------------------------------------------
    Thorn generation (main entry point for non-tensorial thorns)
@@ -290,6 +299,12 @@ Thorn generation (main entry point for non-tensorial thorns)
 
     c = loopControlProcessCode[c, opts];
 
+    (* ------------------------------------------------------------------------ 
+       LoopControlQ
+       ------------------------------------------------------------------------ *)
+
+    c = loopControlQProcessCode[c, opts];
+
     (* ------------------------------------------------------------------------ 
        Declared groups
        ------------------------------------------------------------------------ *)

diff --git a/Tools/CodeGen/Param.m b/Tools/CodeGen/Param.m
@@ -111,6 +111,21 @@
     Default -> -1
    }};
 
+blockingIntParameters[opts___] :=
+  {{
+    Name -> "block_size_i",
+    Description -> "Loop block size",
+    Default -> 8
+   },{
+    Name -> "block_size_j",
+    Description -> "Loop block size",
+    Default -> 8
+   },{
+    Name -> "block_size_k",
+    Description -> "Loop block size",
+    Default -> 8
+   }};
+
 DefFn[ParameterDatabase[opts:OptionsPattern[]] :=
   Module[
     {realParams, intParams, keywordParams,
@@ -123,7 +138,8 @@
        realParams = Join[realParams,ConservationDifferencingRealParameters[]]];
     intParams = OptionValue[IntParameters];
 
-    intParams = Join[intParams, tilingIntParameters[opts]];
+    intParams = Join[intParams,
+                     tilingIntParameters[opts], blockingIntParameters[opts]];
 
     realParamDefs = MakeFullParamDefs[realParams];
     intParamDefs = MakeFullParamDefs[intParams];