futhark 0.20.2 → 0.20.3

raw patch · 46 files changed

+1591/−983 lines, 46 filesdep +randomdep −pcg-randomPVP: major bump suggested

API removals or changes: PVP suggests a major version bump

Dependencies added: random

Dependencies removed: pcg-random

API changes (from Hackage documentation)

Files

• docs/c-api.rst +43/−0
• docs/error-index.rst +171/−0
• docs/index.rst +1/−0
• docs/language-reference.rst +2/−0
• docs/man/futhark-c.rst +4/−0
• docs/man/futhark-cuda.rst +4/−0
• docs/man/futhark-literate.rst +1/−1
• docs/man/futhark-multicore.rst +4/−0
• docs/man/futhark-opencl.rst +4/−0
• docs/usage.rst +6/−3
• futhark.cabal +6/−5
• rts/c/cuda.h +50/−22
• src/Futhark/Actions.hs +14/−6
• src/Futhark/CLI/Dataset.hs +13/−14
• src/Futhark/CLI/Dev.hs +2/−2
• src/Futhark/CLI/Literate.hs +1/−1
• src/Futhark/CodeGen/Backends/CCUDA/Boilerplate.hs +5/−1
• src/Futhark/CodeGen/Backends/GenericC.hs +115/−37
• src/Futhark/CodeGen/Backends/GenericC/CLI.hs +150/−172
• src/Futhark/CodeGen/Backends/GenericC/Manifest.hs +161/−0
• src/Futhark/CodeGen/Backends/GenericC/Server.hs +89/−103
• src/Futhark/CodeGen/Backends/GenericPython.hs +7/−8
• src/Futhark/CodeGen/Backends/MulticoreWASM.hs +1/−1
• src/Futhark/CodeGen/Backends/SequentialWASM.hs +1/−1
• src/Futhark/CodeGen/Backends/SimpleRep.hs +22/−10
• src/Futhark/CodeGen/ImpCode.hs +2/−2
• src/Futhark/CodeGen/ImpGen.hs +28/−18
• src/Futhark/CodeGen/ImpGen/GPU/SegScan/SinglePass.hs +164/−99
• src/Futhark/CodeGen/ImpGen/GPU/SegScan/TwoPass.hs +7/−3
• src/Futhark/CodeGen/SetDefaultSpace.hs +1/−1
• src/Futhark/Compiler.hs +4/−3
• src/Futhark/IR/Parse.hs +7/−2
• src/Futhark/IR/Pretty.hs +3/−0
• src/Futhark/IR/Syntax.hs +12/−5
• src/Futhark/Internalise/Exps.hs +26/−20
• src/Futhark/Optimise/MemoryBlockMerging.hs +248/−0
• src/Futhark/Optimise/MemoryBlockMerging/GreedyColoring.hs +59/−0
• src/Futhark/Optimise/ReuseAllocations.hs +0/−248
• src/Futhark/Optimise/ReuseAllocations/GreedyColoring.hs +0/−59
• src/Futhark/Passes.hs +2/−2
• src/Language/Futhark/Syntax.hs +9/−1
• src/Language/Futhark/TypeChecker.hs +18/−9
• src/Language/Futhark/TypeChecker/Terms.hs +52/−52
• unittests/Futhark/Optimise/MemoryBlockMerging/GreedyColoringTests.hs +70/−0
• unittests/Futhark/Optimise/ReuseAllocations/GreedyColoringTests.hs +0/−70
• unittests/futhark_tests.hs +2/−2

@@ -474,3 +474,46 @@ can result in undefined behaviour.  This does not matter for programs whose entry points do not have unique parameter types (:ref:`in-place-updates`).++.. _manifest:++Manifest+--------++The C backends generate a machine-readable *manifest* in JSON format+that describes the API of the compiled Futhark program.  Specifically,+the manifest contains:++* A mapping from the name of each entry point to:++  * The C function name of the entry point.++  * A list of all *inputs*, including their type and whether they are+    *unique* (consuming).++  * A list of all *outputs*, including their type and whether they are+    *unique*.++* A mapping from the name of each non-scalar types to:++  * The C type of used to represent type type (which is practice+    always a pointer of some kind).++  * For arrays, the element type and rank.++  * A mapping from names of *operations* to the name of the C function+    that implements that operation for the type.  The type of the C+    functions are as documented above.  The following operations are+    listed:++    * For arrays: ``free``, ``shape``, ``values``, ``new``.++    * For opaques: ``free``, ``store``, ``restore``.++Manifests are defined by the following JSON Schema:++.. include:: manifest.schema.json+   :code: json++It is likely that we will add more fields in the future, but it is+unlikely that we will remove any.

@@ -0,0 +1,171 @@+.. _error-index:++Compiler Error Index+====================++Elaboration on type errors produced by the compiler.  Many error+messages contain links to the sections below.++Uniqueness errors+-----------------++.. _use-after-consume:++"Using *x*, but this was consumed at *y*."+~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~++A core principle of uniqueness typing (see :ref:`in-place-updates`) is+that after a variable is "consumed", it must not be used again.  For+example, this is invalid, and will result in the error above::++  let y = x with [0] = 0+  in x++Several operations can *consume* a variable: array update expressions,+calling a function with unique-typed parameters, or passing it as the+initial value of a unique-typed loop parameter.  When a variable is+consumed, its *aliases* are also considered consumed.  Aliasing is the+possibility of two variables occupying the same memory at run-time.+For example, this will fail as above, because ``y`` and ``x`` are+aliased::++  let y = x+  let z = y with [0] = 0+  in x++We can always break aliasing by using a ``copy`` expression::++  let y = copy x+  let z = y with [0] = 0+  in x++.. _not-consumable:++"Would consume *x*, which is not consumable"+~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~++This error message occurs for programs that try to perform a+consumption (such as an in-place update) on variables that are not+consumable.  For example, it would occur for the following program::++  let f (a: []i32) =+    let a[0] = a[0]+1+    in a++Only arrays with a a *unique array type* can be consumed.  Such a type+is written by prefixing the array type with an asterisk.  The program+could be fixed by writing it like this::++  let f (a: *[]i32) =+    let a[0] = a[0]+1+    in a++Note that this places extra obligations on the caller of the ``f``+function, since it now *consumes* its argument.  See+:ref:`in-place-updates` for the full details.++You can always obtain a unique copy of an array by using+``copy``::++  let f (a: []i32) =+    let a = copy a+    let a[0] = a[0]+1+    in a++But note that in most cases (although not all), this subverts the+purpose of using in-place updates in the first place.++.. _return-aliased:++"Unique-typed return value of *x* is aliased to *y*, which is not consumable"+~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~++This can be caused by a function like this::++  let f (xs: []i32) : *[]i32 = xs++We are saying that ``f`` returns a *unique* array - meaning it has no+aliases - but at the same time, it aliases the parameter *xs*, which+is not marked as being unique (see :ref:`in-place-updates`).  This+violates one of the core guarantees provided by uniqueness types,+namely that a unique return value does not alias any value that might+be used in the future.  Imagine if this was permitted, and we had a+program that used ``f``::++  let b = f a+  let b[0] = x+  ...++The update of ``b`` is fine, but if ``b`` was allowed to alias ``a``+(hence occupying the same memory), then we would be modifying ``a`` as+well, which is a violation of referential transparency.++As with most uniqueness errors, it can be fixed by using ``copy xs``+to break the aliasing.  We can also change the type of ``f`` to take a+unique array as input::++  let f (xs: *[]i32) : *[]i32 = xs++This makes ``xs`` "consumable", in the sense used by the error message.++.. _unique-return-aliased:++"A unique-typed component of the return value of *x* is aliased to some other component"+~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~++Caused by programs like the following::++  let main (xs: *[]i32) : (*[]i32, *[]i32) = (xs, xs)++While we are allowed to "consume" ``xs``, as it is a unique parameter,+this function is trying to return two unique values that alias each+other.  This violates one of the core guarantees provided by+uniqueness types, namely that a unique return value does not alias any+value that might be used in the future (see :ref:`in-place-updates`) -+and in this case, the two values alias each other.  We can fix this by+inserting copies to break the aliasing::++  let main (xs: *[]i32) : (*[]i32, *[]i32) = (xs, copy xs)++Size errors+-----------++.. _unused-size:++"Size *x* unused in pattern."+~~~~~~~~~~~~~~~~~~~~~~~~~~~~~++Caused by expressions like this::++  let [n] (y: i32) = x++And functions like this::++  let f [n] (x: i32) = x++Since ``n`` is not the size of anything, it cannot be assigned a value+at runtime.  Hence this program is rejected.++.. _causality-check:++"Causality check"+~~~~~~~~~~~~~~~~~++Causality check errors occur when the program is written in such a way+that a size is needed before it is actually computed.  See+:ref:`causality` for the full rules.  Contrived example::++  let f (b: bool) (xs: []i32) =+    let a = [] : [][]i32+    let b = [filter (>0) xs]+    in a[0] == b[0]++Here the inner size of the array ``a`` must be the same as the inner+size of ``b``, but the inner size of ``b`` depends on a ``filter``+operation that is executed after ``a`` is constructed.++There are various ways to fix causality errors.  In the above case, we+could merely change the order of statements, such that ``b`` is bound+first, meaning that the size is available by the time ``a`` is bound.+In many other cases, we can lift out the "size-producing" expressions+into a separate ``let``-binding preceding the problematic expressions.

@@ -33,6 +33,7 @@    c-api.rst    js-api.rst    package-management.rst+   error-index.rst    server-protocol.rst    c-porting-guide.rst    versus-other-languages.rst

@@ -1183,6 +1183,8 @@ Despite their similar syntax, parameter and return type annotations must be valid at compile-time, or type checking will fail. +.. _causality:+ Causality restriction ~~~~~~~~~~~~~~~~~~~~~ 

@@ -104,6 +104,10 @@   Print various low-overhead logging information to stderr while   running. +-n, --no-print-result++  Do not print the program result.+ -r, --runs=NUM    Perform NUM runs of the program.  With ``-t``, the runtime for each

@@ -134,6 +134,10 @@    Load PTX code from the indicated file. +-n, --no-print-result++  Do not print the program result.+ --nvrtc-option=OPT    Add an additional build option to the string passed to NVRTC.  Refer

@@ -149,7 +149,7 @@   * ``[][]f32`` and ``[][]f64``      Interpreted as greyscale. Values should be between 0 and 1, with 0-    being black and 0 being white.+    being black and 1 being white.    * ``[][]u8`` 

@@ -106,6 +106,10 @@   Print various low-overhead logging information to stderr while   running. +-n, --no-print-result++  Do not print the program result.+ -r, --runs=NUM    Perform NUM runs of the program.  With ``-t``, the runtime for each

@@ -150,6 +150,10 @@    Load an OpenCL binary from the indicated file. +-n, --no-print-result++  Do not print the program result.+ -p, --platform=NAME    Use the first OpenCL platform whose name contains the given string.

@@ -384,9 +384,12 @@    $ futhark opencl --library futlib.fut -This produces two files in the current directory: ``futlib.c`` and-``futlib.h``.  If we wish (and are on a Unix system), we can then-compile ``futlib.c`` to an object file like this::+This produces three files in the current directory: ``futlib.c``,+``futlib.h``, and ``futlib.json`` ( see :ref:`manifest` for more on+the latter).++If we wish (and are on a Unix system), we can then compile+``futlib.c`` to an object file like this::    $ gcc futlib.c -c 

@@ -1,6 +1,6 @@ cabal-version: 2.4 name:           futhark-version:        0.20.2+version:        0.20.3 synopsis:       An optimising compiler for a functional, array-oriented language.  description:    Futhark is a small programming language designed to be compiled to@@ -101,6 +101,7 @@       Futhark.CodeGen.Backends.COpenCL.Boilerplate       Futhark.CodeGen.Backends.GenericC       Futhark.CodeGen.Backends.GenericC.CLI+      Futhark.CodeGen.Backends.GenericC.Manifest       Futhark.CodeGen.Backends.GenericC.Options       Futhark.CodeGen.Backends.GenericC.Server       Futhark.CodeGen.Backends.GenericPython@@ -216,8 +217,8 @@       Futhark.Optimise.InPlaceLowering.LowerIntoStm       Futhark.Optimise.InPlaceLowering.SubstituteIndices       Futhark.Optimise.InliningDeadFun-      Futhark.Optimise.ReuseAllocations-      Futhark.Optimise.ReuseAllocations.GreedyColoring+      Futhark.Optimise.MemoryBlockMerging+      Futhark.Optimise.MemoryBlockMerging.GreedyColoring       Futhark.Optimise.Simplify       Futhark.Optimise.Simplify.Engine       Futhark.Optimise.Simplify.Rep@@ -337,7 +338,7 @@     , mtl >=2.2.1     , neat-interpolation >=0.3     , parallel >=3.2.1.0-    , pcg-random >= 0.1+    , random >= 1.2.0     , process >=1.4.3.0     , process-extras >=0.7.2     , regex-tdfa >=1.2@@ -384,7 +385,7 @@       Futhark.IR.PrimitiveTests       Language.Futhark.CoreTests       Language.Futhark.SyntaxTests-      Futhark.Optimise.ReuseAllocations.GreedyColoringTests+      Futhark.Optimise.MemoryBlockMerging.GreedyColoringTests       Paths_futhark   hs-source-dirs:       unittests

@@ -2,10 +2,10 @@  #define CUDA_SUCCEED_FATAL(x) cuda_api_succeed_fatal(x, #x, __FILE__, __LINE__) #define CUDA_SUCCEED_NONFATAL(x) cuda_api_succeed_nonfatal(x, #x, __FILE__, __LINE__)-#define NVRTC_SUCCEED(x) nvrtc_api_succeed(x, #x, __FILE__, __LINE__)+#define NVRTC_SUCCEED_FATAL(x) nvrtc_api_succeed_fatal(x, #x, __FILE__, __LINE__)+#define NVRTC_SUCCEED_NONFATAL(x) nvrtc_api_succeed_nonfatal(x, #x, __FILE__, __LINE__) -#define CUDA_SUCCEED_OR_RETURN(e) {             \-    char *serror = CUDA_SUCCEED_NONFATAL(e);    \+#define SUCCEED_OR_RETURN(serror) {               \     if (serror) {                                 \       if (!ctx->error) {                          \         ctx->error = serror;                      \@@ -16,6 +16,8 @@     }                                             \   } +#define CUDA_SUCCEED_OR_RETURN(e) SUCCEED_OR_RETURN(CUDA_SUCCEED_NONFATAL(e))+ // CUDA_SUCCEED_OR_RETURN returns the value of the variable 'bad' in // scope.  By default, it will be this one.  Create a local variable // of some other type if needed.  This is a bit of a hack, but it@@ -46,8 +48,8 @@   } } -static inline void nvrtc_api_succeed(nvrtcResult res, const char *call,-                                     const char *file, int line) {+static inline void nvrtc_api_succeed_fatal(nvrtcResult res, const char *call,+                                           const char *file, int line) {   if (res != NVRTC_SUCCESS) {     const char *err_str = nvrtcGetErrorString(res);     futhark_panic(-1, "%s:%d: NVRTC call\n  %s\nfailed with error code %d (%s)\n",@@ -55,6 +57,17 @@   } } +static char* nvrtc_api_succeed_nonfatal(nvrtcResult res, const char *call,+                                        const char *file, int line) {+  if (res != NVRTC_SUCCESS) {+    const char *err_str = nvrtcGetErrorString(res);+    return msgprintf("%s:%d: NVRTC call\n  %s\nfailed with error code %d (%s)\n",+                     file, line, call, res, err_str);+  } else {+    return NULL;+  }+}+ struct cuda_config {   int debugging;   int logging;@@ -306,10 +319,17 @@   return x[chosen].arch_str; } -static char *cuda_nvrtc_build(struct cuda_context *ctx, const char *src,-                              const char *extra_opts[]) {+static char* cuda_nvrtc_build(struct cuda_context *ctx, const char *src,+                              const char *extra_opts[], char **ptx) {   nvrtcProgram prog;-  NVRTC_SUCCEED(nvrtcCreateProgram(&prog, src, "futhark-cuda", 0, NULL, NULL));+  char *problem = NULL;++  problem = NVRTC_SUCCEED_NONFATAL(nvrtcCreateProgram(&prog, src, "futhark-cuda", 0, NULL, NULL));++  if (problem) {+    return problem;+  }+   int arch_set = 0, num_extra_opts;    // nvrtc cannot handle multiple -arch options.  Hence, if one of the@@ -356,6 +376,7 @@     opts[i++] = msgprintf("-I%s/include", getenv("CUDA_PATH"));   }   opts[i++] = msgprintf("-I/usr/local/cuda/include");+  opts[i++] = msgprintf("-I/usr/include");    // It is crucial that the extra_opts are last, so that the free()   // logic below does not cause problems.@@ -379,25 +400,26 @@     if (nvrtcGetProgramLogSize(prog, &log_size) == NVRTC_SUCCESS) {       char *log = (char*) malloc(log_size);       if (nvrtcGetProgramLog(prog, log) == NVRTC_SUCCESS) {-        fprintf(stderr,"Compilation log:\n%s\n", log);+        problem = msgprintf("NVRTC compilation failed.\n\n%s\n", log);+      } else {+        problem = msgprintf("Could not retrieve compilation log\n");       }       free(log);     }-    NVRTC_SUCCEED(res);+    return problem;   }    for (i = i_dyn; i < n_opts-num_extra_opts; i++) { free((char *)opts[i]); }   free(opts); -  char *ptx;   size_t ptx_size;-  NVRTC_SUCCEED(nvrtcGetPTXSize(prog, &ptx_size));-  ptx = (char*) malloc(ptx_size);-  NVRTC_SUCCEED(nvrtcGetPTX(prog, ptx));+  NVRTC_SUCCEED_FATAL(nvrtcGetPTXSize(prog, &ptx_size));+  *ptx = (char*) malloc(ptx_size);+  NVRTC_SUCCEED_FATAL(nvrtcGetPTX(prog, *ptx)); -  NVRTC_SUCCEED(nvrtcDestroyProgram(&prog));+  NVRTC_SUCCEED_FATAL(nvrtcDestroyProgram(&prog)); -  return ptx;+  return NULL; }  static void cuda_size_setup(struct cuda_context *ctx)@@ -475,9 +497,9 @@   } } -static void cuda_module_setup(struct cuda_context *ctx,-                              const char *src_fragments[],-                              const char *extra_opts[]) {+static char* cuda_module_setup(struct cuda_context *ctx,+                               const char *src_fragments[],+                               const char *extra_opts[]) {   char *ptx = NULL, *src = NULL;    if (ctx->cfg.load_program_from == NULL) {@@ -500,7 +522,11 @@   }    if (ptx == NULL) {-    ptx = cuda_nvrtc_build(ctx, src, extra_opts);+    char* problem = cuda_nvrtc_build(ctx, src, extra_opts, &ptx);+    if (problem != NULL) {+      free(src);+      return problem;+    }   }    if (ctx->cfg.dump_ptx_to != NULL) {@@ -513,9 +539,11 @@   if (src != NULL) {     free(src);   }++  return NULL; } -static void cuda_setup(struct cuda_context *ctx, const char *src_fragments[], const char *extra_opts[]) {+static char* cuda_setup(struct cuda_context *ctx, const char *src_fragments[], const char *extra_opts[]) {   CUDA_SUCCEED_FATAL(cuInit(0));    if (cuda_device_setup(ctx) != 0) {@@ -534,7 +562,7 @@   ctx->lockstep_width = device_query(ctx->dev, WARP_SIZE);    cuda_size_setup(ctx);-  cuda_module_setup(ctx, src_fragments, extra_opts);+  return cuda_module_setup(ctx, src_fragments, extra_opts); }  // Count up the runtime all the profiling_records that occured during execution.

@@ -179,12 +179,14 @@       cprog <- handleWarnings fcfg $ SequentialC.compileProg prog       let cpath = outpath `addExtension` "c"           hpath = outpath `addExtension` "h"+          jsonpath = outpath `addExtension` "json"        case mode of         ToLibrary -> do-          let (header, impl) = SequentialC.asLibrary cprog+          let (header, impl, manifest) = SequentialC.asLibrary cprog           liftIO $ T.writeFile hpath $ cPrependHeader header           liftIO $ T.writeFile cpath $ cPrependHeader impl+          liftIO $ T.writeFile jsonpath manifest         ToExecutable -> do           liftIO $ T.writeFile cpath $ SequentialC.asExecutable cprog           runCC cpath outpath ["-O3", "-std=c99"] ["-lm"]@@ -205,6 +207,7 @@       cprog <- handleWarnings fcfg $ COpenCL.compileProg prog       let cpath = outpath `addExtension` "c"           hpath = outpath `addExtension` "h"+          jsonpath = outpath `addExtension` "json"           extra_options             | System.Info.os == "darwin" =               ["-framework", "OpenCL"]@@ -215,9 +218,10 @@        case mode of         ToLibrary -> do-          let (header, impl) = COpenCL.asLibrary cprog+          let (header, impl, manifest) = COpenCL.asLibrary cprog           liftIO $ T.writeFile hpath $ cPrependHeader header           liftIO $ T.writeFile cpath $ cPrependHeader impl+          liftIO $ T.writeFile jsonpath manifest         ToExecutable -> do           liftIO $ T.writeFile cpath $ cPrependHeader $ COpenCL.asExecutable cprog           runCC cpath outpath ["-O", "-std=c99"] ("-lm" : extra_options)@@ -238,6 +242,7 @@       cprog <- handleWarnings fcfg $ CCUDA.compileProg prog       let cpath = outpath `addExtension` "c"           hpath = outpath `addExtension` "h"+          jsonpath = outpath `addExtension` "json"           extra_options =             [ "-lcuda",               "-lcudart",@@ -245,9 +250,10 @@             ]       case mode of         ToLibrary -> do-          let (header, impl) = CCUDA.asLibrary cprog+          let (header, impl, manifest) = CCUDA.asLibrary cprog           liftIO $ T.writeFile hpath $ cPrependHeader header           liftIO $ T.writeFile cpath $ cPrependHeader impl+          liftIO $ T.writeFile jsonpath manifest         ToExecutable -> do           liftIO $ T.writeFile cpath $ cPrependHeader $ CCUDA.asExecutable cprog           runCC cpath outpath ["-O", "-std=c99"] ("-lm" : extra_options)@@ -268,12 +274,14 @@       cprog <- handleWarnings fcfg $ MulticoreC.compileProg prog       let cpath = outpath `addExtension` "c"           hpath = outpath `addExtension` "h"+          jsonpath = outpath `addExtension` "json"        case mode of         ToLibrary -> do-          let (header, impl) = MulticoreC.asLibrary cprog+          let (header, impl, manifest) = MulticoreC.asLibrary cprog           liftIO $ T.writeFile hpath $ cPrependHeader header           liftIO $ T.writeFile cpath $ cPrependHeader impl+          liftIO $ T.writeFile jsonpath manifest         ToExecutable -> do           liftIO $ T.writeFile cpath $ cPrependHeader $ MulticoreC.asExecutable cprog           runCC cpath outpath ["-O3", "-std=c99"] ["-lm", "-pthread"]@@ -383,7 +391,7 @@           liftIO $ T.appendFile classpath SequentialWASM.runServer           runEMCC cpath outpath classpath ["-O3", "-msimd128"] ["-lm"] exps False     writeLibs cprog jsprog = do-      let (h, imp) = SequentialC.asLibrary cprog+      let (h, imp, _) = SequentialC.asLibrary cprog       liftIO $ T.writeFile hpath h       liftIO $ T.writeFile cpath imp       liftIO $ T.writeFile classpath jsprog@@ -417,7 +425,7 @@           runEMCC cpath outpath classpath ["-O3", "-msimd128"] ["-lm", "-pthread"] exps False      writeLibs cprog jsprog = do-      let (h, imp) = MulticoreC.asLibrary cprog+      let (h, imp, _) = MulticoreC.asLibrary cprog       liftIO $ T.writeFile hpath h       liftIO $ T.writeFile cpath imp       liftIO $ T.writeFile classpath jsprog

@@ -31,7 +31,8 @@   ) import System.Exit import System.IO-import System.Random.PCG (Variate (..), initialize)+import System.Random (mkStdGen, uniformR)+import System.Random.Stateful (UniformRange (..))  -- | Run @futhark dataset@. main :: String -> [String] -> IO ()@@ -294,34 +295,32 @@     gen range final = randomVector (range conf) final ds seed  randomVector ::-  (SVec.Storable v, Variate v) =>+  (SVec.Storable v, UniformRange v) =>   Range v ->   (SVec.Vector Int -> SVec.Vector v -> V.Value) ->   [Int] ->   Word64 ->   V.Value randomVector range final ds seed = runST $ do-  -- USe some nice impure computation where we can preallocate a+  -- Use some nice impure computation where we can preallocate a   -- vector of the desired size, populate it via the random number   -- generator, and then finally reutrn a frozen binary vector.   arr <- USVec.new n-  g <- initialize 6364136223846793006 seed-  let fill i+  let fill g i         | i < n = do-          v <- uniformR range g+          let (v, g') = uniformR range g           USVec.write arr i v-          fill $! i + 1+          g' `seq` fill g' $! i + 1         | otherwise =-          final (SVec.fromList ds) . SVec.convert <$> freeze arr-  fill 0+          pure ()+  fill (mkStdGen $ fromIntegral seed) 0+  final (SVec.fromList ds) . SVec.convert <$> freeze arr   where     n = product ds  -- XXX: The following instance is an orphan.  Maybe it could be -- avoided with some newtype trickery or refactoring, but it's so -- convenient this way.-instance Variate Half where-  uniformR (a, b) g = do-    (convFloat :: Float -> Half) <$> uniformR (convFloat a, convFloat b) g-  uniform = uniformR (0, 1)-  uniformB b = uniformR (0, b)+instance UniformRange Half where+  uniformRM (a, b) g =+    (convFloat :: Float -> Half) <$> uniformRM (convFloat a, convFloat b) g

@@ -34,7 +34,7 @@ import Futhark.Optimise.Fusion import Futhark.Optimise.InPlaceLowering import Futhark.Optimise.InliningDeadFun-import qualified Futhark.Optimise.ReuseAllocations as ReuseAllocations+import qualified Futhark.Optimise.MemoryBlockMerging as MemoryBlockMerging import Futhark.Optimise.Sink import Futhark.Optimise.TileLoops import Futhark.Optimise.Unstream@@ -535,7 +535,7 @@     allocateOption "a",     kernelsMemPassOption doubleBufferGPU [],     kernelsMemPassOption expandAllocations [],-    kernelsMemPassOption ReuseAllocations.optimise [],+    kernelsMemPassOption MemoryBlockMerging.optimise [],     cseOption [],     simplifyOption "e",     soacsPipelineOption

@@ -851,7 +851,7 @@       "Stop and do not produce output file if any directive fails."   ] --- | Run @futhark script@.+-- | Run @futhark literate@. main :: String -> [String] -> IO () main = mainWithOptions initialOptions commandLineOptions "program" $ \args opts ->   case args of

@@ -381,7 +381,11 @@                  ctx->total_runtime = 0;                  $stms:init_fields -                 cuda_setup(&ctx->cuda, cuda_program, cfg->nvrtc_opts);+                 ctx->error = cuda_setup(&ctx->cuda, cuda_program, cfg->nvrtc_opts);++                 if (ctx->error != NULL) {+                   return NULL;+                 }                   typename int32_t no_error = -1;                  CUDA_SUCCEED_FATAL(cuMemAlloc(&ctx->global_failure, sizeof(no_error)));

@@ -89,6 +89,7 @@ import Data.Maybe import qualified Data.Text as T import Futhark.CodeGen.Backends.GenericC.CLI (cliDefs)+import qualified Futhark.CodeGen.Backends.GenericC.Manifest as Manifest import Futhark.CodeGen.Backends.GenericC.Options import Futhark.CodeGen.Backends.GenericC.Server (serverDefs) import Futhark.CodeGen.Backends.SimpleRep@@ -832,7 +833,7 @@   PrimType ->   Signedness ->   Int ->-  CompilerM op s [C.Definition]+  CompilerM op s Manifest.ArrayOps arrayLibraryFunctions pub space pt signed rank = do   let pt' = primAPIType signed pt       name = arrayName pt signed rank@@ -846,10 +847,10 @@   values_raw_array <- publicName $ "values_raw_" ++ name   shape_array <- publicName $ "shape_" ++ name -  let shape_names = ["dim" ++ show i | i <- [0 .. rank -1]]+  let shape_names = ["dim" ++ show i | i <- [0 .. rank - 1]]       shape_params = [[C.cparam|typename int64_t $id:k|] | k <- shape_names]       arr_size = cproduct [[C.cexp|$id:k|] | k <- shape_names]-      arr_size_array = cproduct [[C.cexp|arr->shape[$int:i]|] | i <- [0 .. rank -1]]+      arr_size_array = cproduct [[C.cexp|arr->shape[$int:i]|] | i <- [0 .. rank - 1]]   copy <- asks envCopy    memty <- rawMemCType space@@ -861,7 +862,7 @@           [C.cexp|$exp:arr_size * $int:(primByteSize pt::Int)|]           space           [C.cstm|return NULL;|]-        forM_ [0 .. rank -1] $ \i ->+        forM_ [0 .. rank - 1] $ \i ->           let dim_s = "dim" ++ show i            in stm [C.cstm|arr->shape[$int:i] = $id:dim_s;|] @@ -922,7 +923,8 @@   proto     [C.cedecl|const typename int64_t* $id:shape_array($ty:ctx_ty *ctx, $ty:array_type *arr);|] -  return+  mapM_+    libDecl     [C.cunit|           $ty:array_type* $id:new_array($ty:ctx_ty *ctx, const $ty:pt' *data, $params:shape_params) {             $ty:array_type* bad = NULL;@@ -967,10 +969,18 @@           }           |] +  pure $+    Manifest.ArrayOps+      { Manifest.arrayFree = T.pack free_array,+        Manifest.arrayShape = T.pack shape_array,+        Manifest.arrayValues = T.pack values_array,+        Manifest.arrayNew = T.pack new_array+      }+ opaqueLibraryFunctions ::   String ->   [ValueDesc] ->-  CompilerM op s [C.Definition]+  CompilerM op s Manifest.OpaqueOps opaqueLibraryFunctions desc vds = do   name <- publicName $ opaqueName desc vds   free_opaque <- publicName $ "free_" ++ opaqueName desc vds@@ -1006,7 +1016,7 @@             shape_array = "futhark_shape_" ++ arr_name             values_array = "futhark_values_" ++ arr_name             shape' = [C.cexp|$id:shape_array(ctx, obj->$id:field)|]-            num_elems = cproduct [[C.cexp|$exp:shape'[$int:j]|] | j <- [0 .. rank -1]]+            num_elems = cproduct [[C.cexp|$exp:shape'[$int:j]|] | j <- [0 .. rank - 1]]          in ( storageSize pt rank shape',               storeValueHeader sign pt rank shape' [C.cexp|out|]                 ++ [C.cstms|ret |= $id:values_array(ctx, obj->$id:field, (void*)out);@@ -1019,7 +1029,7 @@    store_body <- collect $ do     let (sizes, stores) = unzip $ zipWith storeComponent [0 ..] vds-        size_vars = map (("size_" ++) . show) [0 .. length sizes -1]+        size_vars = map (("size_" ++) . show) [0 .. length sizes - 1]         size_sum = csum [[C.cexp|$id:size|] | size <- size_vars]     forM_ (zip size_vars sizes) $ \(v, e) ->       item [C.citem|typename int64_t $id:v = $exp:e;|]@@ -1041,7 +1051,7 @@             new_array = "futhark_new_" ++ arr_name             dataptr = "data_" ++ show i             shapearr = "shape_" ++ show i-            dims = [[C.cexp|$id:shapearr[$int:j]|] | j <- [0 .. rank -1]]+            dims = [[C.cexp|$id:shapearr[$int:j]|] | j <- [0 .. rank - 1]]             num_elems = cproduct dims         item [C.citem|typename int64_t $id:shapearr[$int:rank];|]         stms $ loadValueHeader sign pt rank [C.cexp|$id:shapearr|] [C.cexp|src|]@@ -1076,7 +1086,8 @@   -- We do not need to enclose the body in a critical section, because   -- when we operate on the components of the opaque, we are calling   -- public API functions that do their own locking.-  return+  mapM_+    libDecl     [C.cunit|           int $id:free_opaque($ty:ctx_ty *ctx, $ty:opaque_type *obj) {             int ret = 0, tmp;@@ -1108,6 +1119,13 @@           }     |] +  pure $+    Manifest.OpaqueOps+      { Manifest.opaqueFree = T.pack free_opaque,+        Manifest.opaqueStore = T.pack store_opaque,+        Manifest.opaqueRestore = T.pack restore_opaque+      }+ valueDescToCType :: Publicness -> ValueDesc -> CompilerM op s C.Type valueDescToCType _ (ScalarValue pt signed _) =   return $ primAPIType signed pt@@ -1127,22 +1145,37 @@   mapM_ (valueDescToCType Private) vds   pure [C.cty|struct $id:name|] -generateAPITypes :: CompilerM op s ()+generateAPITypes :: CompilerM op s (M.Map T.Text Manifest.Type) generateAPITypes = do-  mapM_ generateArray . M.toList =<< gets compArrayTypes-  mapM_ generateOpaque . M.toList =<< gets compOpaqueTypes+  array_ts <- mapM generateArray . M.toList =<< gets compArrayTypes+  opaque_ts <- mapM generateOpaque . M.toList =<< gets compOpaqueTypes+  pure $ M.fromList $ catMaybes array_ts <> opaque_ts   where     generateArray ((space, signed, pt, rank), pub) = do       name <- publicName $ arrayName pt signed rank       let memty = fatMemType space       libDecl [C.cedecl|struct $id:name { $ty:memty mem; typename int64_t shape[$int:rank]; };|]-      mapM libDecl =<< arrayLibraryFunctions pub space pt signed rank+      ops <- arrayLibraryFunctions pub space pt signed rank+      let pt_name = T.pack $ prettySigned (signed == TypeUnsigned) pt+          pretty_name = mconcat (replicate rank "[]") <> pt_name+          arr_type = [C.cty|struct $id:name*|]+      case pub of+        Public ->+          pure $+            Just+              ( pretty_name,+                Manifest.TypeArray (prettyText arr_type) pt_name rank ops+              )+        Private ->+          pure Nothing      generateOpaque (desc, vds) = do       name <- publicName $ opaqueName desc vds       members <- zipWithM field vds [(0 :: Int) ..]       libDecl [C.cedecl|struct $id:name { $sdecls:members };|]-      mapM libDecl =<< opaqueLibraryFunctions desc vds+      ops <- opaqueLibraryFunctions desc vds+      let opaque_type = [C.cty|struct $id:name*|]+      pure (T.pack desc, Manifest.TypeOpaque (prettyText opaque_type) ops)      field vd@ScalarValue {} i = do       ct <- valueDescToCType Private vd@@ -1182,7 +1215,7 @@       checks <- map snd <$> zipWithM (prepareValue Private) (zipWith field [0 ..] vds) vds       return         ( [C.cparam|const $ty:ty *$id:pname|],-          if null $ concat checks+          if all null checks             then Nothing             else Just $ allTrue $ concat checks         )@@ -1209,7 +1242,7 @@             )        let (sets, checks) =-            unzip $ zipWith maybeCopyDim shape [0 .. rank -1]+            unzip $ zipWith maybeCopyDim shape [0 .. rank - 1]       stms $ catMaybes sets        return ([C.cty|$ty:ty*|], checks)@@ -1256,13 +1289,13 @@             [C.cstm|$exp:dest->shape[$int:i] = $exp:x;|]           maybeCopyDim (Var d) i =             [C.cstm|$exp:dest->shape[$int:i] = $id:d;|]-      stms $ zipWith maybeCopyDim shape [0 .. rank -1]+      stms $ zipWith maybeCopyDim shape [0 .. rank - 1]  onEntryPoint ::   [C.BlockItem] ->   Name ->   Function op ->-  CompilerM op s (Maybe C.Definition)+  CompilerM op s (Maybe (C.Definition, (T.Text, Manifest.EntryPoint))) onEntryPoint _ _ (Function Nothing _ _ _ _ _) = pure Nothing onEntryPoint get_consts fname (Function (Just ename) outputs inputs _ results args) = do   let out_args = map (\p -> [C.cexp|&$id:(paramName p)|]) outputs@@ -1273,7 +1306,7 @@    entry_point_function_name <- publicName $ "entry_" ++ nameToString ename -  (inputs', unpack_entry_inputs) <- prepareEntryInputs args+  (inputs', unpack_entry_inputs) <- prepareEntryInputs $ map snd args   let (entry_point_input_params, entry_point_input_checks) = unzip inputs'    (entry_point_output_params, pack_entry_outputs) <-@@ -1317,8 +1350,8 @@    ops <- asks envOperations -  pure . Just $-    [C.cedecl|+  let cdef =+        [C.cedecl|        int $id:entry_point_function_name            ($ty:ctx_ty *ctx,             $params:entry_point_output_params,@@ -1332,6 +1365,18 @@           return ret;        }|]++      manifest =+        Manifest.EntryPoint+          { Manifest.entryPointCFun = T.pack entry_point_function_name,+            -- Note that our convention about what is "input/output"+            -- and what is "results/args" is different between the+            -- manifest and ImpCode.+            Manifest.entryPointOutputs = map outputManifest results,+            Manifest.entryPointInputs = map inputManifest args+          }++  pure $ Just (cdef, (nameToText ename, manifest))   where     stubParam (MemParam name space) =       declMem name space@@ -1339,6 +1384,33 @@       let ty' = primTypeToCType ty       decl [C.cdecl|$ty:ty' $id:name;|] +    vdTypeAndUnique (TransparentValue _ (ScalarValue pt signed _)) =+      ( T.pack $ prettySigned (signed == TypeUnsigned) pt,+        False+      )+    vdTypeAndUnique (TransparentValue u (ArrayValue _ _ pt signed shape)) =+      ( T.pack $+          mconcat (replicate (length shape) "[]")+            <> prettySigned (signed == TypeUnsigned) pt,+        u == Unique+      )+    vdTypeAndUnique (OpaqueValue u name _) =+      (T.pack name, u == Unique)++    outputManifest vd =+      let (t, u) = vdTypeAndUnique vd+       in Manifest.Output+            { Manifest.outputType = t,+              Manifest.outputUnique = u+            }+    inputManifest (v, vd) =+      let (t, u) = vdTypeAndUnique vd+       in Manifest.Input+            { Manifest.inputName = nameToText v,+              Manifest.inputType = t,+              Manifest.inputUnique = u+            }+ -- | The result of compilation to C is multiple parts, which can be -- put together in various ways.  The obvious way is to concatenate -- all of them, which yields a CLI program.  Another is to compile the@@ -1350,7 +1422,9 @@     cUtils :: T.Text,     cCLI :: T.Text,     cServer :: T.Text,-    cLib :: T.Text+    cLib :: T.Text,+    -- | The manifest, in JSON format.+    cJsonManifest :: T.Text   }  gnuSource :: T.Text@@ -1387,11 +1461,12 @@ #endif |] --- | Produce header and implementation files.-asLibrary :: CParts -> (T.Text, T.Text)+-- | Produce header, implementation, and manifest files.+asLibrary :: CParts -> (T.Text, T.Text, T.Text) asLibrary parts =   ( "#pragma once\n\n" <> cHeader parts,-    gnuSource <> disableWarnings <> cHeader parts <> cUtils parts <> cLib parts+    gnuSource <> disableWarnings <> cHeader parts <> cUtils parts <> cLib parts,+    cJsonManifest parts   )  -- | As executable with command-line interface.@@ -1418,7 +1493,7 @@   m CParts compileProg backend ops extra header_extra spaces options prog = do   src <- getNameSource-  let ((prototypes, definitions, entry_point_decls), endstate) =+  let ((prototypes, definitions, entry_point_decls, manifest), endstate) =         runCompilerM ops src () compileProg'       initdecls = initDecls endstate       entrydecls = entryDecls endstate@@ -1480,8 +1555,8 @@    let early_decls = T.unlines $ map prettyText $ DL.toList $ compEarlyDecls endstate       lib_decls = T.unlines $ map prettyText $ DL.toList $ compLibDecls endstate-      clidefs = cliDefs options $ Functions entry_funs-      serverdefs = serverDefs options $ Functions entry_funs+      clidefs = cliDefs options manifest+      serverdefs = serverDefs options manifest       libdefs =         [untrimming| #ifdef _MSC_VER@@ -1518,11 +1593,11 @@         cUtils = utildefs,         cCLI = clidefs,         cServer = serverdefs,-        cLib = libdefs+        cLib = libdefs,+        cJsonManifest = Manifest.manifestToJSON manifest       }   where     Definitions consts (Functions funs) = prog-    entry_funs = filter (isJust . functionEntry . snd) funs      compileProg' = do       (memstructs, memfuns, memreport) <- unzip3 <$> mapM defineMemorySpace spaces@@ -1535,19 +1610,20 @@         unzip <$> mapM (compileFun get_consts [[C.cparam|$ty:ctx_ty *ctx|]]) funs        mapM_ earlyDecl memstructs-      entry_points <--        catMaybes <$> mapM (uncurry (onEntryPoint get_consts)) funs+      (entry_points, entry_points_manifest) <-+        unzip . catMaybes <$> mapM (uncurry (onEntryPoint get_consts)) funs        extra        mapM_ earlyDecl $ concat memfuns -      commonLibFuns memreport+      types <- commonLibFuns memreport        return         ( T.unlines $ map prettyText prototypes,           T.unlines $ map (prettyText . funcToDef) functions,-          T.unlines $ map prettyText entry_points+          T.unlines $ map prettyText entry_points,+          Manifest.Manifest (M.fromList entry_points_manifest) types backend         )      funcToDef func = C.FuncDef func loc@@ -1556,9 +1632,9 @@           C.OldFunc _ _ _ _ _ _ l -> l           C.Func _ _ _ _ _ l -> l -commonLibFuns :: [C.BlockItem] -> CompilerM op s ()+commonLibFuns :: [C.BlockItem] -> CompilerM op s (M.Map T.Text Manifest.Type) commonLibFuns memreport = do-  generateAPITypes+  types <- generateAPITypes   ctx <- contextType   ops <- asks envOperations   profilereport <- gets $ DL.toList . compProfileItems@@ -1642,6 +1718,8 @@                          return ctx->error != NULL;                        }|]     )++  pure types  compileConstants :: Constants op -> CompilerM op s [C.BlockItem] compileConstants (Constants ps init_consts) = do

@@ -13,13 +13,19 @@ where  import Data.List (unzip5)-import Data.Maybe+import qualified Data.Map as M import qualified Data.Text as T+import Futhark.CodeGen.Backends.GenericC.Manifest import Futhark.CodeGen.Backends.GenericC.Options import Futhark.CodeGen.Backends.SimpleRep-import Futhark.CodeGen.ImpCode+  ( cproduct,+    fromStorage,+    primAPIType,+    primStorageType,+    scalarToPrim,+  ) import Futhark.CodeGen.RTS.C (tuningH, valuesH)-import Futhark.Util.Pretty (prettyText)+import Futhark.Util.Pretty (pretty, prettyText) import qualified Language.C.Quote.OpenCL as C import qualified Language.C.Syntax as C @@ -68,6 +74,13 @@         optionAction = [C.cstm|binary_output = 1;|]       },     Option+      { optionLongName = "no-print-result",+        optionShortName = Just 'n',+        optionArgument = NoArgument,+        optionDescription = "Do not print the program result.",+        optionAction = [C.cstm|print_result = 0;|]+      },+    Option       { optionLongName = "help",         optionShortName = Just 'h',         optionArgument = NoArgument,@@ -145,188 +158,149 @@           }         }|] -valueDescToCType :: ValueDesc -> C.Type-valueDescToCType (ScalarValue pt signed _) =-  primAPIType signed pt-valueDescToCType (ArrayValue _ _ pt signed shape) =-  let name = "futhark_" ++ arrayName pt signed (length shape)-   in [C.cty|struct $id:name|]--opaqueToCType :: String -> [ValueDesc] -> C.Type-opaqueToCType desc vds =-  let name = "futhark_" ++ opaqueName desc vds-   in [C.cty|struct $id:name|]--externalValueToCType :: ExternalValue -> C.Type-externalValueToCType (TransparentValue _ vd) = valueDescToCType vd-externalValueToCType (OpaqueValue _ desc vds) = opaqueToCType desc vds--primTypeInfo :: PrimType -> Signedness -> C.Exp-primTypeInfo (IntType it) t = case (it, t) of-  (Int8, TypeUnsigned) -> [C.cexp|u8_info|]-  (Int16, TypeUnsigned) -> [C.cexp|u16_info|]-  (Int32, TypeUnsigned) -> [C.cexp|u32_info|]-  (Int64, TypeUnsigned) -> [C.cexp|u64_info|]-  (Int8, _) -> [C.cexp|i8_info|]-  (Int16, _) -> [C.cexp|i16_info|]-  (Int32, _) -> [C.cexp|i32_info|]-  (Int64, _) -> [C.cexp|i64_info|]-primTypeInfo (FloatType Float16) _ = [C.cexp|f16_info|]-primTypeInfo (FloatType Float32) _ = [C.cexp|f32_info|]-primTypeInfo (FloatType Float64) _ = [C.cexp|f64_info|]-primTypeInfo Bool _ = [C.cexp|bool_info|]-primTypeInfo Unit _ = [C.cexp|bool_info|]--readPrimStm :: C.ToIdent a => a -> Int -> PrimType -> Signedness -> C.Stm-readPrimStm place i t ept =-  [C.cstm|if (read_scalar(stdin, &$exp:(primTypeInfo t ept), &$id:place) != 0) {-            futhark_panic(1, "Error when reading input #%d of type %s (errno: %s).\n",-                          $int:i,-                          $exp:(primTypeInfo t ept).type_name,-                          strerror(errno));-          }|]--readInput :: Int -> ExternalValue -> ([C.BlockItem], C.Stm, C.Stm, C.Stm, C.Exp)-readInput i (OpaqueValue _ desc _) =-  ( [C.citems|futhark_panic(1, "Cannot read input #%d of type %s\n", $int:i, $string:desc);|],-    [C.cstm|;|],-    [C.cstm|;|],-    [C.cstm|;|],-    [C.cexp|NULL|]-  )-readInput i (TransparentValue _ (ScalarValue t ept _)) =-  let dest = "read_value_" ++ show i-   in ( [C.citems|$ty:(primStorageType t) $id:dest;-                  $stm:(readPrimStm dest i t ept);|],+readInput :: Manifest -> Int -> T.Text -> ([C.BlockItem], C.Stm, C.Stm, C.Stm, C.Exp)+readInput manifest i tname =+  case M.lookup tname $ manifestTypes manifest of+    Nothing ->+      let (_, t) = scalarToPrim tname+          dest = "read_value_" ++ show i+          info = T.unpack tname <> "_info"+       in ( [C.citems|+             $ty:(primStorageType t) $id:dest;+             if (read_scalar(stdin, &$id:info, &$id:dest) != 0) {+             futhark_panic(1, "Error when reading input #%d of type %s (errno: %s).\n",+                           $int:i,+                           $string:(T.unpack tname),+                           strerror(errno));+                           };|],+            [C.cstm|;|],+            [C.cstm|;|],+            [C.cstm|;|],+            fromStorage t [C.cexp|$id:dest|]+          )+    Just (TypeOpaque desc _) ->+      ( [C.citems|futhark_panic(1, "Cannot read input #%d of type %s\n", $int:i, $string:(T.unpack desc));|],         [C.cstm|;|],         [C.cstm|;|],         [C.cstm|;|],-        fromStorage t [C.cexp|$id:dest|]+        [C.cexp|NULL|]       )-readInput i (TransparentValue _ (ArrayValue _ _ t ept dims)) =-  let dest = "read_value_" ++ show i-      shape = "read_shape_" ++ show i-      arr = "read_arr_" ++ show i+    Just (TypeArray t et rank ops) ->+      let dest = "read_value_" ++ show i+          shape = "read_shape_" ++ show i+          arr = "read_arr_" ++ show i -      name = arrayName t ept rank-      arr_ty_name = "futhark_" ++ name-      ty = [C.cty|struct $id:arr_ty_name|]-      rank = length dims-      dims_exps = [[C.cexp|$id:shape[$int:j]|] | j <- [0 .. rank -1]]-      dims_s = concat $ replicate rank "[]"-      t' = primAPIType ept t+          ty = [C.cty|typename $id:t|]+          dims_exps = [[C.cexp|$id:shape[$int:j]|] | j <- [0 .. rank -1]]+          t' = uncurry primAPIType $ scalarToPrim et -      new_array = "futhark_new_" ++ name-      free_array = "futhark_free_" ++ name+          new_array = arrayNew ops+          free_array = arrayFree ops+          info = T.unpack et <> "_info" -      items =-        [C.citems|-           $ty:ty *$id:dest;-           typename int64_t $id:shape[$int:rank];-           $ty:t' *$id:arr = NULL;-           errno = 0;-           if (read_array(stdin,-                          &$exp:(primTypeInfo t ept),-                          (void**) &$id:arr,-                          $id:shape,-                          $int:(length dims))-               != 0) {-             futhark_panic(1, "Cannot read input #%d of type %s%s (errno: %s).\n",-                           $int:i,-                           $string:dims_s,-                           $exp:(primTypeInfo t ept).type_name,-                           strerror(errno));-           }|]-   in ( items,-        [C.cstm|assert(($id:dest = $id:new_array(ctx, $id:arr, $args:dims_exps)) != NULL);|],-        [C.cstm|assert($id:free_array(ctx, $id:dest) == 0);|],-        [C.cstm|free($id:arr);|],-        [C.cexp|$id:dest|]-      )+          items =+            [C.citems|+               $ty:ty $id:dest;+               typename int64_t $id:shape[$int:rank];+               $ty:t' *$id:arr = NULL;+               errno = 0;+               if (read_array(stdin,+                              &$id:info,+                              (void**) &$id:arr,+                              $id:shape,+                              $int:rank)+                   != 0) {+                 futhark_panic(1, "Cannot read input #%d of type %s%s (errno: %s).\n",+                               $int:i,+                               $string:(T.unpack tname),+                               $id:info.type_name,+                               strerror(errno));+               }|]+       in ( items,+            [C.cstm|assert(($id:dest = $id:new_array(ctx, $id:arr, $args:dims_exps)) != NULL);|],+            [C.cstm|assert($id:free_array(ctx, $id:dest) == 0);|],+            [C.cstm|free($id:arr);|],+            [C.cexp|$id:dest|]+          ) -readInputs :: [ExternalValue] -> [([C.BlockItem], C.Stm, C.Stm, C.Stm, C.Exp)]-readInputs = zipWith readInput [0 ..]+readInputs :: Manifest -> [T.Text] -> [([C.BlockItem], C.Stm, C.Stm, C.Stm, C.Exp)]+readInputs manifest = zipWith (readInput manifest) [0 ..] -prepareOutputs :: [ExternalValue] -> [(C.BlockItem, C.Exp, C.Stm)]-prepareOutputs = zipWith prepareResult [(0 :: Int) ..]+prepareOutputs :: Manifest -> [T.Text] -> [(C.BlockItem, C.Exp, C.Stm)]+prepareOutputs manifest = zipWith prepareResult [(0 :: Int) ..]   where-    prepareResult i ev = do-      let ty = externalValueToCType ev-          result = "result_" ++ show i+    prepareResult i tname = do+      let result = "result_" ++ show i -      case ev of-        TransparentValue _ ScalarValue {} ->-          ( [C.citem|$ty:ty $id:result;|],-            [C.cexp|$id:result|],-            [C.cstm|;|]-          )-        TransparentValue _ (ArrayValue _ _ t ept dims) ->-          let name = arrayName t ept $ length dims-              free_array = "futhark_free_" ++ name-           in ( [C.citem|$ty:ty *$id:result;|],-                [C.cexp|$id:result|],-                [C.cstm|assert($id:free_array(ctx, $id:result) == 0);|]-              )-        OpaqueValue _ desc vds ->-          let free_opaque = "futhark_free_" ++ opaqueName desc vds-           in ( [C.citem|$ty:ty *$id:result;|],+      case M.lookup tname $ manifestTypes manifest of+        Nothing ->+          let (s, pt) = scalarToPrim tname+              ty = primAPIType s pt+           in ( [C.citem|$ty:ty $id:result;|],                 [C.cexp|$id:result|],-                [C.cstm|assert($id:free_opaque(ctx, $id:result) == 0);|]+                [C.cstm|;|]               )--printPrimStm :: (C.ToExp a, C.ToExp b) => a -> b -> PrimType -> Signedness -> C.Stm-printPrimStm dest val bt ept =-  [C.cstm|write_scalar($exp:dest, binary_output, &$exp:(primTypeInfo bt ept), &$exp:val);|]+        Just (TypeArray t _ _ ops) ->+          ( [C.citem|typename $id:t $id:result;|],+            [C.cexp|$id:result|],+            [C.cstm|assert($id:(arrayFree ops)(ctx, $id:result) == 0);|]+          )+        Just (TypeOpaque t ops) ->+          ( [C.citem|typename $id:t $id:result;|],+            [C.cexp|$id:result|],+            [C.cstm|assert($id:(opaqueFree ops)(ctx, $id:result) == 0);|]+          )  -- | Return a statement printing the given external value.-printStm :: ExternalValue -> C.Exp -> C.Stm-printStm (OpaqueValue _ desc _) _ =-  [C.cstm|printf("#<opaque %s>", $string:desc);|]-printStm (TransparentValue _ (ScalarValue bt ept _)) e =-  printPrimStm [C.cexp|stdout|] e bt ept-printStm (TransparentValue _ (ArrayValue _ _ bt ept shape)) e =-  let values_array = "futhark_values_" ++ name-      shape_array = "futhark_shape_" ++ name-      num_elems = cproduct [[C.cexp|$id:shape_array(ctx, $exp:e)[$int:i]|] | i <- [0 .. rank -1]]-   in [C.cstm|{-        $ty:bt' *arr = calloc(sizeof($ty:bt'), $exp:num_elems);-        assert(arr != NULL);-        assert($id:values_array(ctx, $exp:e, arr) == 0);-        write_array(stdout, binary_output, &$exp:(primTypeInfo bt ept), arr,-                    $id:shape_array(ctx, $exp:e), $int:rank);-        free(arr);-      }|]-  where-    rank = length shape-    bt' = primStorageType bt-    name = arrayName bt ept rank+printStm :: Manifest -> T.Text -> C.Exp -> C.Stm+printStm manifest tname e =+  case M.lookup tname $ manifestTypes manifest of+    Nothing ->+      let info = tname <> "_info"+       in [C.cstm|write_scalar(stdout, binary_output, &$id:info, &$exp:e);|]+    Just (TypeOpaque desc _) ->+      [C.cstm|printf("#<opaque %s>", $string:(T.unpack desc));|]+    Just (TypeArray _ et rank ops) ->+      let et' = uncurry primAPIType $ scalarToPrim et+          values_array = arrayValues ops+          shape_array = arrayShape ops+          num_elems =+            cproduct [[C.cexp|$id:shape_array(ctx, $exp:e)[$int:i]|] | i <- [0 .. rank -1]]+          info = et <> "_info"+       in [C.cstm|{+                 $ty:et' *arr = calloc($exp:num_elems, $id:info.size);+                 assert(arr != NULL);+                 assert($id:values_array(ctx, $exp:e, arr) == 0);+                 write_array(stdout, binary_output, &$id:info, arr,+                             $id:shape_array(ctx, $exp:e), $int:rank);+                 free(arr);+                 }|] -printResult :: [(ExternalValue, C.Exp)] -> [C.Stm]-printResult = concatMap f+printResult :: Manifest -> [(T.Text, C.Exp)] -> [C.Stm]+printResult manifest = concatMap f   where-    f (v, e) = [printStm v e, [C.cstm|printf("\n");|]]+    f (v, e) = [printStm manifest v e, [C.cstm|printf("\n");|]]  cliEntryPoint ::-  FunctionT a -> Maybe (C.Definition, C.Initializer)-cliEntryPoint fun@(Function _ _ _ _ results args) = do-  entry_point_name <- nameToString <$> functionEntry fun+  Manifest -> T.Text -> EntryPoint -> (C.Definition, C.Initializer)+cliEntryPoint manifest entry_point_name (EntryPoint cfun outputs inputs) =   let (input_items, pack_input, free_input, free_parsed, input_args) =-        unzip5 $ readInputs args+        unzip5 $ readInputs manifest $ map inputType inputs        (output_decls, output_vals, free_outputs) =-        unzip3 $ prepareOutputs results+        unzip3 $ prepareOutputs manifest $ map outputType outputs -      printstms = printResult $ zip results output_vals+      printstms =+        printResult manifest $ zip (map outputType outputs) output_vals        ctx_ty = [C.cty|struct futhark_context|]-      sync_ctx = "futhark_context_sync" :: Name-      error_ctx = "futhark_context_get_error" :: Name+      sync_ctx = "futhark_context_sync" :: T.Text+      error_ctx = "futhark_context_get_error" :: T.Text -      cli_entry_point_function_name = "futrts_cli_entry_" ++ entry_point_name-      entry_point_function_name = "futhark_entry_" ++ entry_point_name+      cli_entry_point_function_name = "futrts_cli_entry_" ++ T.unpack entry_point_name -      pause_profiling = "futhark_context_pause_profiling" :: Name-      unpause_profiling = "futhark_context_unpause_profiling" :: Name+      pause_profiling = "futhark_context_pause_profiling" :: T.Text+      unpause_profiling = "futhark_context_unpause_profiling" :: T.Text        addrOf e = [C.cexp|&$exp:e|] @@ -343,9 +317,9 @@                   $id:unpause_profiling(ctx);                 }                 t_start = get_wall_time();-                r = $id:entry_point_function_name(ctx,-                                                  $args:(map addrOf output_vals),-                                                  $args:input_args);+                r = $id:cfun(ctx,+                             $args:(map addrOf output_vals),+                             $args:input_args);                 if (r != 0) {                   futhark_panic(1, "%s", $id:error_ctx(ctx));                 }@@ -363,8 +337,7 @@                 }                 $stms:free_input               |]-  Just-    ( [C.cedecl|+   in ( [C.cedecl|    static void $id:cli_entry_point_function_name($ty:ctx_ty *ctx) {      typename int64_t t_start, t_end;      int time_runs = 0, profile_run = 0;@@ -401,27 +374,31 @@      // Free the parsed input.      $stms:free_parsed -     // Print the final result.-     if (binary_output) {-       set_binary_mode(stdout);+     if (print_result) {+       // Print the final result.+       if (binary_output) {+         set_binary_mode(stdout);+       }+       $stms:printstms      }-     $stms:printstms       $stms:free_outputs    }|],-      [C.cinit|{ .name = $string:entry_point_name,+        [C.cinit|{ .name = $string:(T.unpack entry_point_name),                        .fun = $id:cli_entry_point_function_name }|]-    )+      )  {-# NOINLINE cliDefs #-}  -- | Generate Futhark standalone executable code.-cliDefs :: [Option] -> Functions a -> T.Text-cliDefs options (Functions funs) =+cliDefs :: [Option] -> Manifest -> T.Text+cliDefs options manifest =   let option_parser =         generateOptionParser "parse_options" $ genericOptions ++ options       (cli_entry_point_decls, entry_point_inits) =-        unzip $ mapMaybe (cliEntryPoint . snd) funs+        unzip $+          map (uncurry (cliEntryPoint manifest)) $+            M.toList $ manifestEntryPoints manifest    in prettyText         [C.cunit| $esc:("#include <getopt.h>")@@ -431,6 +408,7 @@ $esc:(T.unpack valuesH)  static int binary_output = 0;+static int print_result = 1; static typename FILE *runtime_file; static int perform_warmup = 0; static int num_runs = 1;

@@ -0,0 +1,161 @@+{-# LANGUAGE OverloadedStrings #-}++-- | C manifest data structure and serialisation to JSON.+--+-- A manifest contains machine-readable information about the API of+-- the compiled Futhark program.  Specifically which entry points are+-- available, which types are exposed, and what their C names are.+module Futhark.CodeGen.Backends.GenericC.Manifest+  ( Manifest (..),+    Input (..),+    Output (..),+    EntryPoint (..),+    Type (..),+    ArrayOps (..),+    OpaqueOps (..),+    manifestToJSON,+  )+where++import Data.Aeson (ToJSON (..), object)+import qualified Data.Aeson as JSON+import Data.Aeson.Text (encodeToLazyText)+import qualified Data.Map as M+import qualified Data.Text as T+import Data.Text.Lazy (toStrict)++-- | Manifest info for an entry point parameter.+data Input = Input+  { inputName :: T.Text,+    inputType :: T.Text,+    inputUnique :: Bool+  }+  deriving (Eq, Ord, Show)++-- | Manifest info for an entry point return value.+data Output = Output+  { outputType :: T.Text,+    outputUnique :: Bool+  }+  deriving (Eq, Ord, Show)++-- | Manifest info for an entry point.+data EntryPoint = EntryPoint+  { entryPointCFun :: T.Text,+    entryPointOutputs :: [Output],+    entryPointInputs :: [Input]+  }+  deriving (Eq, Ord, Show)++-- | The names of the C functions implementing the operations on some+-- array type.+data ArrayOps = ArrayOps+  { arrayFree :: T.Text,+    arrayShape :: T.Text,+    arrayValues :: T.Text,+    arrayNew :: T.Text+  }+  deriving (Eq, Ord, Show)++-- | The names of the C functions implementing the operations on some+-- opaque type.+data OpaqueOps = OpaqueOps+  { opaqueFree :: T.Text,+    opaqueStore :: T.Text,+    opaqueRestore :: T.Text+  }+  deriving (Eq, Ord, Show)++-- | Manifest info for a non-scalar type.  Scalar types are not part+-- of the manifest for a program.+data Type+  = -- | ctype, Futhark elemtype, rank.+    TypeArray T.Text T.Text Int ArrayOps+  | TypeOpaque T.Text OpaqueOps+  deriving (Eq, Ord, Show)++-- | A manifest for a compiled program.+data Manifest = Manifest+  { -- | A mapping from Futhark entry points to how they are+    -- represented in C.+    manifestEntryPoints :: M.Map T.Text EntryPoint,+    -- | A mapping from Futhark type name to how they are represented+    -- at the C level.  Should not contain any of the primitive scalar+    -- types.  For array types, these have empty dimensions,+    -- e.g. @[]i32@.+    manifestTypes :: M.Map T.Text Type,+    -- | The compiler backend used to+    -- compile the program, e.g. @c@.+    manifestBackend :: T.Text+  }+  deriving (Eq, Ord, Show)++instance JSON.ToJSON ArrayOps where+  toJSON (ArrayOps free shape values new) =+    object+      [ ("free", toJSON free),+        ("shape", toJSON shape),+        ("values", toJSON values),+        ("new", toJSON new)+      ]++instance JSON.ToJSON OpaqueOps where+  toJSON (OpaqueOps free store restore) =+    object+      [ ("free", toJSON free),+        ("store", toJSON store),+        ("restore", toJSON restore)+      ]++instance JSON.ToJSON Manifest where+  toJSON (Manifest entry_points types backend) =+    object+      [ ("backend", toJSON backend),+        ( "entry_points",+          object $ M.toList $ fmap onEntryPoint entry_points+        ),+        ( "types",+          object $ M.toList $ fmap onType types+        )+      ]+    where+      onEntryPoint (EntryPoint cfun outputs inputs) =+        object+          [ ("cfun", toJSON cfun),+            ("outputs", toJSON $ map onOutput outputs),+            ("inputs", toJSON $ map onInput inputs)+          ]++      onOutput (Output t u) =+        object+          [ ("type", toJSON t),+            ("unique", toJSON u)+          ]++      onInput (Input p t u) =+        object+          [ ("name", toJSON p),+            ("type", toJSON t),+            ("unique", toJSON u)+          ]++      onType (TypeArray t et rank ops) =+        object+          [ ("kind", "array"),+            ("ctype", toJSON t),+            ("rank", toJSON rank),+            ("elemtype", toJSON et),+            ("ops", toJSON ops)+          ]+      onType (TypeOpaque t ops) =+        object+          [ ("kind", "opaque"),+            ("ctype", toJSON t),+            ("ops", toJSON ops)+          ]++-- | Serialise a manifest in JSON format, so it can be read from other+-- tools.  The schema supposed to be at+-- https://futhark-lang.org/manifest.schema.json.+manifestToJSON :: Manifest -> T.Text+manifestToJSON = toStrict . encodeToLazyText

@@ -14,12 +14,12 @@  import Data.Bifunctor (first, second) import qualified Data.Map as M-import Data.Maybe import qualified Data.Text as T+import Futhark.CodeGen.Backends.GenericC.Manifest import Futhark.CodeGen.Backends.GenericC.Options import Futhark.CodeGen.Backends.SimpleRep-import Futhark.CodeGen.ImpCode import Futhark.CodeGen.RTS.C (serverH, tuningH, valuesH)+import Futhark.Util (zEncodeString) import Futhark.Util.Pretty (prettyText) import qualified Language.C.Quote.OpenCL as C import qualified Language.C.Syntax as C@@ -102,187 +102,173 @@       }   ] -typeStructName :: ExternalValue -> String-typeStructName (TransparentValue _ (ScalarValue pt signed _)) =-  let name = prettySigned (signed == TypeUnsigned) pt-   in "type_" ++ name-typeStructName (TransparentValue _ (ArrayValue _ _ pt signed shape)) =-  let rank = length shape-      name = arrayName pt signed rank-   in "type_" ++ name-typeStructName (OpaqueValue _ name vds) =-  "type_" ++ opaqueName name vds+typeStructName :: T.Text -> String+typeStructName tname = "type_" <> zEncodeString (T.unpack tname) -valueDescBoilerplate :: ExternalValue -> (String, (C.Initializer, [C.Definition]))-valueDescBoilerplate ev@(TransparentValue _ (ScalarValue pt signed _)) =-  let name = prettySigned (signed == TypeUnsigned) pt-      type_name = typeStructName ev-   in (name, ([C.cinit|&$id:type_name|], mempty))-valueDescBoilerplate ev@(TransparentValue _ (ArrayValue _ _ pt signed shape)) =-  let rank = length shape-      name = arrayName pt signed rank-      pt_name = prettySigned (signed == TypeUnsigned) pt-      pretty_name = concat (replicate rank "[]") ++ pt_name-      type_name = typeStructName ev+typeBoilerplate :: (T.Text, Type) -> (C.Initializer, [C.Definition])+typeBoilerplate (tname, TypeArray _ et rank ops) =+  let type_name = typeStructName tname       aux_name = type_name ++ "_aux"-      info_name = pt_name ++ "_info"-      array_new = "futhark_new_" ++ name-      array_new_wrap = "futhark_new_" ++ name ++ "_wrap"-      array_free = "futhark_free_" ++ name-      array_shape = "futhark_shape_" ++ name-      array_values = "futhark_values_" ++ name+      info_name = T.unpack et ++ "_info"       shape_args = [[C.cexp|shape[$int:i]|] | i <- [0 .. rank -1]]-   in ( name,-        ( [C.cinit|&$id:type_name|],-          [C.cunit|+      array_new_wrap = arrayNew ops <> "_wrap"+   in ( [C.cinit|&$id:type_name|],+        [C.cunit|               void* $id:array_new_wrap(struct futhark_context *ctx,                                        const void* p,                                        const typename int64_t* shape) {-                return $id:array_new(ctx, p, $args:shape_args);+                return $id:(arrayNew ops)(ctx, p, $args:shape_args);               }               struct array_aux $id:aux_name = {-                .name = $string:pretty_name,+                .name = $string:(T.unpack tname),                 .rank = $int:rank,                 .info = &$id:info_name,                 .new = (typename array_new_fn)$id:array_new_wrap,-                .free = (typename array_free_fn)$id:array_free,-                .shape = (typename array_shape_fn)$id:array_shape,-                .values = (typename array_values_fn)$id:array_values+                .free = (typename array_free_fn)$id:(arrayFree ops),+                .shape = (typename array_shape_fn)$id:(arrayShape ops),+                .values = (typename array_values_fn)$id:(arrayValues ops)               };               struct type $id:type_name = {-                .name = $string:pretty_name,+                .name = $string:(T.unpack tname),                 .restore = (typename restore_fn)restore_array,                 .store = (typename store_fn)store_array,                 .free = (typename free_fn)free_array,                 .aux = &$id:aux_name               };|]-        )       )-valueDescBoilerplate ev@(OpaqueValue _ name vds) =-  let type_name = typeStructName ev+typeBoilerplate (tname, TypeOpaque _ ops) =+  let type_name = typeStructName tname       aux_name = type_name ++ "_aux"-      opaque_free = "futhark_free_" ++ opaqueName name vds-      opaque_store = "futhark_store_" ++ opaqueName name vds-      opaque_restore = "futhark_restore_" ++ opaqueName name vds-   in ( name,-        ( [C.cinit|&$id:type_name|],-          [C.cunit|+   in ( [C.cinit|&$id:type_name|],+        [C.cunit|               struct opaque_aux $id:aux_name = {-                .store = (typename opaque_store_fn)$id:opaque_store,-                .restore = (typename opaque_restore_fn)$id:opaque_restore,-                .free = (typename opaque_free_fn)$id:opaque_free+                .store = (typename opaque_store_fn)$id:(opaqueStore ops),+                .restore = (typename opaque_restore_fn)$id:(opaqueRestore ops),+                .free = (typename opaque_free_fn)$id:(opaqueFree ops)               };               struct type $id:type_name = {-                .name = $string:name,+                .name = $string:(T.unpack tname),                 .restore = (typename restore_fn)restore_opaque,                 .store = (typename store_fn)store_opaque,                 .free = (typename free_fn)free_opaque,                 .aux = &$id:aux_name               };|]-        )       ) -functionExternalValues :: Function a -> [ExternalValue]-functionExternalValues fun = functionResult fun ++ functionArgs fun--entryTypeBoilerplate :: Functions a -> ([C.Initializer], [C.Definition])-entryTypeBoilerplate (Functions funs) =-  second concat . unzip . M.elems . M.fromList . map valueDescBoilerplate-    . concatMap (functionExternalValues . snd)-    . filter (isJust . functionEntry . snd)-    $ funs+entryTypeBoilerplate :: Manifest -> ([C.Initializer], [C.Definition])+entryTypeBoilerplate =+  second concat . unzip . map typeBoilerplate . M.toList . manifestTypes -oneEntryBoilerplate :: (Name, Function a) -> Maybe ([C.Definition], C.Initializer)-oneEntryBoilerplate (name, fun) = do-  ename <- functionEntry fun-  let entry_f = "futhark_entry_" ++ pretty ename-      call_f = "call_" ++ pretty name-      out_types = functionResult fun-      in_types = functionArgs fun-      out_types_name = pretty name ++ "_out_types"-      in_types_name = pretty name ++ "_in_types"-      out_unique_name = pretty name ++ "_out_unique"-      in_unique_name = pretty name ++ "_in_unique"+oneEntryBoilerplate :: Manifest -> (T.Text, EntryPoint) -> ([C.Definition], C.Initializer)+oneEntryBoilerplate manifest (name, EntryPoint cfun outputs inputs) =+  let call_f = "call_" ++ T.unpack name+      out_types = map outputType outputs+      in_types = map inputType inputs+      out_types_name = T.unpack name ++ "_out_types"+      in_types_name = T.unpack name ++ "_in_types"+      out_unique_name = T.unpack name ++ "_out_unique"+      in_unique_name = T.unpack name ++ "_in_unique"       (out_items, out_args)         | null out_types = ([C.citems|(void)outs;|], mempty)         | otherwise = unzip $ zipWith loadOut [0 ..] out_types       (in_items, in_args)         | null in_types = ([C.citems|(void)ins;|], mempty)         | otherwise = unzip $ zipWith loadIn [0 ..] in_types-  pure-    ( [C.cunit|+   in ( [C.cunit|                 struct type* $id:out_types_name[] = {                   $inits:(map typeStructInit out_types),                   NULL                 };                 bool $id:out_unique_name[] = {-                  $inits:(map typeUniqueInit out_types)+                  $inits:(map outputUniqueInit outputs)                 };                 struct type* $id:in_types_name[] = {                   $inits:(map typeStructInit in_types),                   NULL                 };                 bool $id:in_unique_name[] = {-                  $inits:(map typeUniqueInit in_types)+                  $inits:(map inputUniqueInit inputs)                 };                 int $id:call_f(struct futhark_context *ctx, void **outs, void **ins) {                   $items:out_items                   $items:in_items-                  return $id:entry_f(ctx, $args:out_args, $args:in_args);+                  return $id:cfun(ctx, $args:out_args, $args:in_args);                 }                 |],-      [C.cinit|{-            .name = $string:(pretty ename),+        [C.cinit|{+            .name = $string:(T.unpack name),             .f = $id:call_f,             .in_types = $id:in_types_name,             .out_types = $id:out_types_name,             .in_unique = $id:in_unique_name,             .out_unique = $id:out_unique_name             }|]-    )+      )   where-    typeStructInit t = [C.cinit|&$id:(typeStructName t)|]-    typeUniqueInit t =-      case typeUnique t of-        Unique -> [C.cinit|true|]-        Nonunique -> [C.cinit|false|]+    typeStructInit tname = [C.cinit|&$id:(typeStructName tname)|]+    inputUniqueInit = uniqueInit . inputUnique+    outputUniqueInit = uniqueInit . outputUnique+    uniqueInit True = [C.cinit|true|]+    uniqueInit False = [C.cinit|false|] -    typeUnique (TransparentValue u _) = u-    typeUnique (OpaqueValue u _ _) = u+    cType tname =+      case M.lookup tname $ manifestTypes manifest of+        Just (TypeArray ctype _ _ _) -> [C.cty|typename $id:(T.unpack ctype)|]+        Just (TypeOpaque ctype _) -> [C.cty|typename $id:(T.unpack ctype)|]+        Nothing -> uncurry primAPIType $ scalarToPrim tname -    loadOut i ev =+    loadOut i tname =       let v = "out" ++ show (i :: Int)-       in ( [C.citem|$ty:(externalValueType ev) *$id:v = outs[$int:i];|],+       in ( [C.citem|$ty:(cType tname) *$id:v = outs[$int:i];|],             [C.cexp|$id:v|]           )-    loadIn i ev =+    loadIn i tname =       let v = "in" ++ show (i :: Int)-          evt = externalValueType ev-       in ( [C.citem|$ty:evt $id:v = *($ty:evt*)ins[$int:i];|],+       in ( [C.citem|$ty:(cType tname) $id:v = *($ty:(cType tname)*)ins[$int:i];|],             [C.cexp|$id:v|]           ) -entryBoilerplate :: Functions a -> ([C.Definition], [C.Initializer])-entryBoilerplate (Functions funs) =-  first concat $ unzip $ mapMaybe oneEntryBoilerplate funs+entryBoilerplate :: Manifest -> ([C.Definition], [C.Initializer])+entryBoilerplate manifest =+  first concat $+    unzip $+      map (oneEntryBoilerplate manifest) $+        M.toList $ manifestEntryPoints manifest  mkBoilerplate ::-  Functions a ->+  Manifest ->   ([C.Definition], [C.Initializer], [C.Initializer])-mkBoilerplate funs =-  let (type_inits, type_defs) = entryTypeBoilerplate funs-      (entry_defs, entry_inits) = entryBoilerplate funs-   in (type_defs ++ entry_defs, type_inits, entry_inits)+mkBoilerplate manifest =+  let (type_inits, type_defs) = entryTypeBoilerplate manifest+      (entry_defs, entry_inits) = entryBoilerplate manifest+      scalar_type_inits = map scalarTypeInit scalar_types+   in (type_defs ++ entry_defs, scalar_type_inits ++ type_inits, entry_inits)+  where+    scalarTypeInit tname = [C.cinit|&$id:(typeStructName tname)|]+    scalar_types =+      [ "i8",+        "i16",+        "i32",+        "i64",+        "u8",+        "u16",+        "u32",+        "u64",+        "f16",+        "f32",+        "f64",+        "bool"+      ]  {-# NOINLINE serverDefs #-}  -- | Generate Futhark server executable code.-serverDefs :: [Option] -> Functions a -> T.Text-serverDefs options funs =+serverDefs :: [Option] -> Manifest -> T.Text+serverDefs options manifest =   let option_parser =         generateOptionParser "parse_options" $ genericOptions ++ options       (boilerplate_defs, type_inits, entry_point_inits) =-        mkBoilerplate funs+        mkBoilerplate manifest    in prettyText         [C.cunit| $esc:("#include <getopt.h>")

@@ -326,7 +326,7 @@        ]  functionExternalValues :: Imp.Function a -> [Imp.ExternalValue]-functionExternalValues fun = Imp.functionResult fun ++ Imp.functionArgs fun+functionExternalValues fun = Imp.functionResult fun ++ map snd (Imp.functionArgs fun)  opaqueDefs :: Imp.Functions a -> M.Map String [PyExp] opaqueDefs (Imp.Functions funs) =@@ -826,10 +826,9 @@       _ -> return Nothing    prepareIn <- collect $ do-    declEntryPointInputSizes args-    mapM_ entryPointInput $-      zip3 [0 ..] args $-        map (Var . extValueDescName) args+    declEntryPointInputSizes $ map snd args+    mapM_ entryPointInput . zip3 [0 ..] (map snd args) $+      map (Var . extValueDescName . snd) args   (res, prepareOut) <- collect' $ mapM entryPointOutput results    let argexps_lib = map (compileName . Imp.paramName) inputs@@ -848,7 +847,7 @@         ]    return-    ( map extValueDescName args,+    ( map (extValueDescName . snd) args,       prepareIn,       call argexps_lib,       call argexps_bin,@@ -921,7 +920,7 @@  entryTypes :: Imp.Function op -> ([String], [String]) entryTypes func =-  ( map desc $ Imp.functionArgs func,+  ( map (desc . snd) $ Imp.functionArgs func,     map desc $ Imp.functionResult func   )   where@@ -938,7 +937,7 @@ callEntryFun pre_timing entry@(fname, Imp.Function (Just ename) _ _ _ _ decl_args) = do   (_, prepare_in, _, body_bin, _, res, prepare_run) <- prepareEntry entry -  let str_input = map readInput decl_args+  let str_input = map (readInput . snd) decl_args       end_of_input = [Exp $ simpleCall "end_of_input" [String $ pretty fname]]        exitcall = [Exp $ simpleCall "sys.exit" [Field (String "Assertion.{} failed") "format(e)"]]

@@ -55,7 +55,7 @@         Just $           JSEntryPoint             { name = nameToString n,-              parameters = map extToString args,+              parameters = map (extToString . snd) args,               ret = map extToString res             }    in mapMaybe (function . snd) fs

@@ -55,7 +55,7 @@         Just $           JSEntryPoint             { name = nameToString n,-              parameters = map extToString args,+              parameters = map (extToString . snd) args,               ret = map extToString res             }    in mapMaybe (function . snd) fs

@@ -1,3 +1,4 @@+{-# LANGUAGE OverloadedStrings #-} {-# LANGUAGE QuasiQuotes #-} {-# LANGUAGE TemplateHaskell #-} {-# LANGUAGE Trustworthy #-}@@ -20,11 +21,11 @@     primAPIType,     arrayName,     opaqueName,-    externalValueType,     toStorage,     fromStorage,     cproduct,     csum,+    scalarToPrim,      -- * Primitive value operations     cScalarDefs,@@ -139,15 +140,22 @@           )     iter x = ((x :: Word32) `shiftR` 16) `xor` x --- | The type used to expose a Futhark value in the C API.  A pointer--- in the case of arrays and opaques.-externalValueType :: ExternalValue -> C.Type-externalValueType (OpaqueValue _ desc vds) =-  [C.cty|struct $id:("futhark_" ++ opaqueName desc vds)*|]-externalValueType (TransparentValue _ (ArrayValue _ _ pt signed shape)) =-  [C.cty|struct $id:("futhark_" ++ arrayName pt signed (length shape))*|]-externalValueType (TransparentValue _ (ScalarValue pt signed _)) =-  primAPIType signed pt+-- | The 'PrimType' (and sign) correspond to a human-readable scalar+-- type name (e.g. @f64@).  Beware: partial!+scalarToPrim :: T.Text -> (Signedness, PrimType)+scalarToPrim "bool" = (TypeDirect, Bool)+scalarToPrim "i8" = (TypeDirect, IntType Int8)+scalarToPrim "i16" = (TypeDirect, IntType Int16)+scalarToPrim "i32" = (TypeDirect, IntType Int32)+scalarToPrim "i64" = (TypeDirect, IntType Int64)+scalarToPrim "u8" = (TypeUnsigned, IntType Int8)+scalarToPrim "u16" = (TypeUnsigned, IntType Int16)+scalarToPrim "u32" = (TypeUnsigned, IntType Int32)+scalarToPrim "u64" = (TypeUnsigned, IntType Int64)+scalarToPrim "f16" = (TypeDirect, FloatType Float16)+scalarToPrim "f32" = (TypeDirect, FloatType Float32)+scalarToPrim "f64" = (TypeDirect, FloatType Float64)+scalarToPrim tname = error $ "scalarToPrim: " <> T.unpack tname  -- | Return an expression multiplying together the given expressions. -- If an empty list is given, the expression @1@ is returned.@@ -167,6 +175,10 @@  instance C.ToIdent Name where   toIdent = C.toIdent . zEncodeString . nameToString++-- Orphan!+instance C.ToIdent T.Text where+  toIdent = C.toIdent . T.unpack  instance C.ToIdent VName where   toIdent = C.toIdent . zEncodeString . pretty

@@ -171,7 +171,7 @@     functionInput :: [Param],     functionBody :: Code a,     functionResult :: [ExternalValue],-    functionArgs :: [ExternalValue]+    functionArgs :: [(Name, ExternalValue)]   }   deriving (Show) @@ -631,7 +631,7 @@     where       onFun f =         fvBind pnames $-          freeIn' (functionBody f) <> freeIn' (functionResult f <> functionArgs f)+          freeIn' (functionBody f) <> freeIn' (functionResult f <> map snd (functionArgs f))         where           pnames =             namesFromList $ map paramName $ functionInput f <> functionOutput f

@@ -7,6 +7,7 @@ {-# LANGUAGE OverloadedStrings #-} {-# LANGUAGE Strict #-} {-# LANGUAGE Trustworthy #-}+{-# LANGUAGE TupleSections #-} {-# LANGUAGE TypeFamilies #-}  module Futhark.CodeGen.ImpGen@@ -516,11 +517,11 @@ compileInParams ::   Mem rep inner =>   [FParam rep] ->-  [EntryPointType] ->-  ImpM rep r op ([Imp.Param], [ArrayDecl], [Imp.ExternalValue])-compileInParams params orig_epts = do+  [EntryParam] ->+  ImpM rep r op ([Imp.Param], [ArrayDecl], [(Name, Imp.ExternalValue)])+compileInParams params eparams = do   let (ctx_params, val_params) =-        splitAt (length params - sum (map entryPointSize orig_epts)) params+        splitAt (length params - sum (map (entryPointSize . entryParamType) eparams)) params   (inparams, arrayds) <- partitionEithers <$> mapM compileInParam (ctx_params ++ val_params)   let findArray x = find (isArrayDecl x) arrayds @@ -545,22 +546,24 @@           _ ->             Nothing -      mkExts (TypeOpaque u desc n : epts) fparams =+      mkExts (EntryParam v (TypeOpaque u desc n) : epts) fparams =         let (fparams', rest) = splitAt n fparams-         in Imp.OpaqueValue-              u-              desc-              (mapMaybe (`mkValueDesc` Imp.TypeDirect) fparams') :+         in ( v,+              Imp.OpaqueValue+                u+                desc+                (mapMaybe (`mkValueDesc` Imp.TypeDirect) fparams')+            ) :             mkExts epts rest-      mkExts (TypeUnsigned u : epts) (fparam : fparams) =-        maybeToList (Imp.TransparentValue u <$> mkValueDesc fparam Imp.TypeUnsigned)+      mkExts (EntryParam v (TypeUnsigned u) : epts) (fparam : fparams) =+        maybeToList ((v,) . Imp.TransparentValue u <$> mkValueDesc fparam Imp.TypeUnsigned)           ++ mkExts epts fparams-      mkExts (TypeDirect u : epts) (fparam : fparams) =-        maybeToList (Imp.TransparentValue u <$> mkValueDesc fparam Imp.TypeDirect)+      mkExts (EntryParam v (TypeDirect u) : epts) (fparam : fparams) =+        maybeToList ((v,) . Imp.TransparentValue u <$> mkValueDesc fparam Imp.TypeDirect)           ++ mkExts epts fparams       mkExts _ _ = [] -  return (inparams, arrayds, mkExts orig_epts val_params)+  return (inparams, arrayds, mkExts eparams val_params)   where     isArrayDecl x (ArrayDecl y _ _) = x == y @@ -649,7 +652,7 @@     (name_entry, params_entry, ret_entry) = case entry of       Nothing ->         ( Nothing,-          replicate (length params) (TypeDirect mempty),+          replicate (length params) (EntryParam "" $ TypeDirect mempty),           Nothing         )       Just (x, y, z) -> (Just x, y, Just z)@@ -1702,7 +1705,7 @@   let condInBounds (DimFix i) d =         0 .<=. i .&&. i .<. d       condInBounds (DimSlice i n s) d =-        0 .<=. i .&&. i + (n -1) * s .<. d+        0 .<=. i .&&. i + (n - 1) * s .<. d    in foldl1 (.&&.) $ zipWith condInBounds slice dims  --- Building blocks for constructing code.@@ -1736,7 +1739,14 @@ sIf cond tbranch fbranch = do   tbranch' <- collect tbranch   fbranch' <- collect fbranch-  emit $ Imp.If cond tbranch' fbranch'+  -- Avoid generating branch if the condition is known statically.+  emit $+    if cond == true+      then tbranch'+      else+        if cond == false+          then fbranch'+          else Imp.If cond tbranch' fbranch'  sWhen :: Imp.TExp Bool -> ImpM rep r op () -> ImpM rep r op () sWhen cond tbranch = sIf cond tbranch (return ())@@ -1792,7 +1802,7 @@ -- | Uses linear/iota index function. sAllocArray :: String -> PrimType -> ShapeBase SubExp -> Space -> ImpM rep r op VName sAllocArray name pt shape space =-  sAllocArrayPerm name pt shape space [0 .. shapeRank shape -1]+  sAllocArrayPerm name pt shape space [0 .. shapeRank shape - 1]  -- | Uses linear/iota index function. sStaticArray :: String -> Space -> PrimType -> Imp.ArrayContents -> ImpM rep r op VName

@@ -32,7 +32,7 @@   Count GroupSize SubExp ->   SubExp ->   [PrimType] ->-  InKernelGen (VName, [VName], [VName], VName, VName, [VName])+  InKernelGen (VName, [VName], [VName], VName, [VName]) createLocalArrays (Count groupSize) m types = do   let groupSizeE = toInt64Exp groupSize       workSize = toInt64Exp m * groupSizeE@@ -69,7 +69,6 @@   transposeArrayLength <- dPrimV "trans_arr_len" workSize    sharedId <- sArrayInMem "shared_id" int32 (Shape [constant (1 :: Int32)]) localMem-  sharedReadOffset <- sArrayInMem "shared_read_offset" int32 (Shape [constant (1 :: Int32)]) localMem    transposedArrays <-     forM types $ \ty ->@@ -98,8 +97,113 @@         (Shape [constant (warpSize :: Int64)])         $ ArrayIn localMem $ IxFun.iotaOffset off' [warpSize] -  return (sharedId, transposedArrays, prefixArrays, sharedReadOffset, warpscan, warpExchanges)+  return (sharedId, transposedArrays, prefixArrays, warpscan, warpExchanges) +inBlockScanLookback ::+  KernelConstants ->+  Imp.TExp Int64 ->+  VName ->+  [VName] ->+  Lambda GPUMem ->+  InKernelGen ()+inBlockScanLookback constants arrs_full_size flag_arr arrs scan_lam = everythingVolatile $ do+  flg_x <- dPrim "flg_x" p_int8+  flg_y <- dPrim "flg_y" p_int8+  let flg_param_x = Param (tvVar flg_x) (MemPrim p_int8)+      flg_param_y = Param (tvVar flg_y) (MemPrim p_int8)+      flg_y_exp = tvExp flg_y+      statusP = (2 :: Imp.TExp Int8)+      statusX = (0 :: Imp.TExp Int8)++  dLParams (lambdaParams scan_lam)++  skip_threads <- dPrim "skip_threads" int32+  let in_block_thread_active =+        tvExp skip_threads .<=. in_block_id+      actual_params = lambdaParams scan_lam+      (x_params, y_params) =+        splitAt (length actual_params `div` 2) actual_params+      y_to_x =+        forM_ (zip x_params y_params) $ \(x, y) ->+          when (primType (paramType x)) $+            copyDWIM (paramName x) [] (Var (paramName y)) []+      y_to_x_flg =+        copyDWIM (tvVar flg_x) [] (Var (tvVar flg_y)) []++  -- Set initial y values+  sComment "read input for in-block scan" $ do+    zipWithM_ readInitial (flg_param_y : y_params) (flag_arr : arrs)+    -- Since the final result is expected to be in x_params, we may+    -- need to copy it there for the first thread in the block.+    sWhen (in_block_id .==. 0) $ do+      y_to_x+      y_to_x_flg++  when array_scan barrier++  let op_to_x = do+        sIf+          (flg_y_exp .==. statusP .||. flg_y_exp .==. statusX)+          ( do+              y_to_x_flg+              y_to_x+          )+          (compileBody' x_params $ lambdaBody scan_lam)++  sComment "in-block scan (hopefully no barriers needed)" $ do+    skip_threads <-- 1++    sWhile (tvExp skip_threads .<. block_size) $ do+      sWhen in_block_thread_active $ do+        sComment "read operands" $+          zipWithM_+            (readParam (sExt64 $ tvExp skip_threads))+            (flg_param_x : x_params)+            (flag_arr : arrs)+        sComment "perform operation" op_to_x++        sComment "write result" $+          sequence_ $+            zipWith3+              writeResult+              (flg_param_x : x_params)+              (flg_param_y : y_params)+              (flag_arr : arrs)++      skip_threads <-- tvExp skip_threads * 2+  where+    p_int8 = IntType Int8+    block_size = 32+    block_id = ltid32 `quot` block_size+    in_block_id = ltid32 - block_id * block_size+    ltid32 = kernelLocalThreadId constants+    ltid = sExt64 ltid32+    gtid = sExt64 $ kernelGlobalThreadId constants+    array_scan = not $ all primType $ lambdaReturnType scan_lam+    barrier+      | array_scan =+        sOp $ Imp.Barrier Imp.FenceGlobal+      | otherwise =+        sOp $ Imp.Barrier Imp.FenceLocal++    readInitial p arr+      | primType $ paramType p =+        copyDWIM (paramName p) [] (Var arr) [DimFix ltid]+      | otherwise =+        copyDWIM (paramName p) [] (Var arr) [DimFix gtid]+    readParam behind p arr+      | primType $ paramType p =+        copyDWIM (paramName p) [] (Var arr) [DimFix $ ltid - behind]+      | otherwise =+        copyDWIM (paramName p) [] (Var arr) [DimFix $ gtid - behind + arrs_full_size]++    writeResult x y arr+      | primType $ paramType x = do+        copyDWIM arr [DimFix ltid] (Var $ paramName x) []+        copyDWIM (paramName y) [] (Var $ paramName x) []+      | otherwise =+        copyDWIM (paramName y) [] (Var $ paramName x) []+ -- | Compile 'SegScan' instance to host-level code with calls to a -- single-pass kernel. compileSegScan ::@@ -127,12 +231,6 @@       statusX = 0       statusA = 1       statusP = 2-      makeStatusUsed flag used = tvExp flag .|. (tvExp used .<<. 2)-      unmakeStatusUsed :: TV Int8 -> TV Int8 -> TV Int8 -> InKernelGen ()-      unmakeStatusUsed flagUsed flag used = do-        used <-- tvExp flagUsed .>>. 2-        flag <-- tvExp flagUsed .&. 3-       sumT :: Integer       maxT :: Integer       sumT = foldl (\bytes typ -> bytes + primByteSize typ) 0 tys@@ -140,14 +238,18 @@       sumT' = foldl (\bytes typ -> bytes + primByteSize' typ) 0 tys `div` 4       maxT = maximum (map primByteSize tys)       -- TODO: Make these constants dynamic by querying device-      -- RTX 2080 Ti constants (CC 7.5)       k_reg = 64-      k_mem = 48 --12*4+      k_mem = 95       mem_constraint = max k_mem sumT `div` maxT-      reg_constraint = (k_reg -1 - sumT') `div` (2 * sumT' + 3)+      reg_constraint = (k_reg - 1 - sumT') `div` (2 * sumT')       m :: Num a => a       m = fromIntegral $ max 1 $ min mem_constraint reg_constraint +  emit $ Imp.DebugPrint "SegScan: number of elements processed sequentially per thread is m:" $ Just $ untyped (m :: TPrimExp Int32 Imp.ExpLeaf)+  emit $ Imp.DebugPrint "SegScan: memory constraints is: " $ Just $ untyped (fromIntegral mem_constraint :: TPrimExp Int32 Imp.ExpLeaf)+  emit $ Imp.DebugPrint "SegScan: register constraints is: " $ Just $ untyped (fromIntegral reg_constraint :: TPrimExp Int32 Imp.ExpLeaf)+  emit $ Imp.DebugPrint "SegScan: sumT' is: " $ Just $ untyped (fromIntegral sumT' :: TPrimExp Int32 Imp.ExpLeaf)+   -- Allocate the shared memory for output component   numThreads <- dPrimV "numThreads" num_threads   numGroups <- dPrimV "numGroups" $ unCount num_groups@@ -165,7 +267,7 @@   sKernelThread "segscan" num_groups group_size (segFlat space) $ do     constants <- kernelConstants <$> askEnv -    (sharedId, transposedArrays, prefixArrays, sharedReadOffset, warpscan, exchanges) <-+    (sharedId, transposedArrays, prefixArrays, warpscan, exchanges) <-       createLocalArrays (segGroupSize lvl) (intConst Int64 m) tys      dynamicId <- dPrim "dynamic_id" int32@@ -317,7 +419,7 @@     sComment "Perform lookback" $ do       sWhen (blockNewSgm .&&. kernelLocalThreadId constants .==. 0) $ do         everythingVolatile $-          forM_ (zip incprefixArrays accs) $ \(incprefixArray, acc) ->+          forM_ (zip accs incprefixArrays) $ \(acc, incprefixArray) ->             copyDWIMFix incprefixArray [tvExp dynamicId] (tvSize acc) []         sOp globalFence         everythingVolatile $@@ -347,7 +449,8 @@                 everythingVolatile $                   copyDWIMFix statusFlags [tvExp dynamicId] (intConst Int8 statusP) []             )-          copyDWIMFix warpscan [0] (Var statusFlags) [tvExp dynamicId - 1]+          everythingVolatile $+            copyDWIMFix warpscan [0] (Var statusFlags) [tvExp dynamicId - 1]         -- sWhen         sOp localFence @@ -375,8 +478,7 @@                 readI <- dPrimV "read_i" $ tvExp readOffset + kernelLocalThreadId constants                 aggrs <- forM (zip scanOpNe tys) $ \(ne, ty) ->                   dPrimV "aggr" $ TPrimExp $ toExp' ty ne-                flag <- dPrimV "flag" statusX-                used <- dPrimV "used" 0+                flag <- dPrimV "flag" (statusX :: Imp.TExp Int8)                 everythingVolatile . sWhen (tvExp readI .>=. 0) $ do                   sIf                     (sameSegment readI)@@ -390,81 +492,49 @@                           ( sWhen (tvExp flag .==. statusA) $ do                               forM_ (zip aggrs aggregateArrays) $ \(aggr, aggregate) ->                                 copyDWIMFix (tvVar aggr) [] (Var aggregate) [sExt64 $ tvExp readI]-                              used <-- 1                           )                     )                     (copyDWIMFix (tvVar flag) [] (intConst Int8 statusP) [])                 -- end sIf                 -- end sWhen+                 forM_ (zip exchanges aggrs) $ \(exchange, aggr) ->                   copyDWIMFix exchange [sExt64 $ kernelLocalThreadId constants] (tvSize aggr) []-                tmp <- dPrimV "tmp" $ makeStatusUsed flag used-                copyDWIMFix warpscan [sExt64 $ kernelLocalThreadId constants] (tvSize tmp) []-                sOp localFence+                copyDWIMFix warpscan [sExt64 $ kernelLocalThreadId constants] (tvSize flag) [] -                (warpscanMem, warpscanSpace, warpscanOff) <--                  fullyIndexArray warpscan [sExt64 warpSize - 1]-                flag <-- TPrimExp (Imp.index warpscanMem warpscanOff int8 warpscanSpace Imp.Volatile)-                sWhen (kernelLocalThreadId constants .==. 0) $ do-                  -- TODO: This is a single-threaded reduce-                  sIf-                    (bNot $ tvExp flag .==. statusP)-                    ( do-                        scanOp'' <- renameLambda scanOp'-                        let (agg1s, agg2s) = splitAt (length tys) $ map paramName $ lambdaParams scanOp''+                -- execute warp-parallel reduction but only if the last read flag in not STATUS_P+                copyDWIMFix (tvVar flag) [] (Var warpscan) [sExt64 warpSize - 1]+                sWhen (tvExp flag .<. (2 :: Imp.TExp Int8)) $ do+                  lam' <- renameLambda scanOp'+                  inBlockScanLookback+                    constants+                    (tvExp numThreads)+                    warpscan+                    exchanges+                    lam' -                        forM_ (zip3 agg1s scanOpNe tys) $ \(agg1, ne, ty) ->-                          dPrimV_ agg1 $ TPrimExp $ toExp' ty ne-                        zipWithM_ dPrim_ agg2s tys+                -- all threads of the warp read the result of reduction+                copyDWIMFix (tvVar flag) [] (Var warpscan) [sExt64 warpSize - 1]+                forM_ (zip aggrs exchanges) $ \(aggr, exchange) ->+                  copyDWIMFix (tvVar aggr) [] (Var exchange) [sExt64 warpSize - 1]+                -- update read offset+                sIf+                  (tvExp flag .==. statusP)+                  (readOffset <-- loopStop)+                  ( sWhen (tvExp flag .==. statusA) $ do+                      readOffset <-- tvExp readOffset - zExt32 warpSize+                  ) -                        flag1 <- dPrimV "flag1" statusX-                        flag2 <- dPrim "flag2" int8-                        used1 <- dPrimV "used1" 0-                        used2 <- dPrim "used2" int8-                        sFor "i" warpSize $ \i -> do-                          copyDWIMFix (tvVar flag2) [] (Var warpscan) [sExt64 i]-                          unmakeStatusUsed flag2 flag2 used2-                          forM_ (zip agg2s exchanges) $ \(agg2, exchange) ->-                            copyDWIMFix agg2 [] (Var exchange) [sExt64 i]-                          sIf-                            (bNot $ tvExp flag2 .==. statusA)-                            ( do-                                flag1 <-- tvExp flag2-                                used1 <-- tvExp used2-                                forM_ (zip3 agg1s tys agg2s) $ \(agg1, ty, agg2) ->-                                  agg1 <~~ toExp' ty (Var agg2)-                            )-                            ( do-                                used1 <-- tvExp used1 + tvExp used2-                                compileStms mempty (bodyStms $ lambdaBody scanOp'') $-                                  forM_ (zip3 agg1s tys $ map resSubExp $ bodyResult $ lambdaBody scanOp'') $-                                    \(agg1, ty, res) -> agg1 <~~ toExp' ty res-                            )-                        flag <-- tvExp flag1-                        used <-- tvExp used1-                        forM_ (zip3 aggrs tys agg1s) $ \(aggr, ty, agg1) ->-                          tvVar aggr <~~ toExp' ty (Var agg1)-                    )-                    -- else-                    ( forM_ (zip aggrs exchanges) $ \(aggr, exchange) ->-                        copyDWIMFix (tvVar aggr) [] (Var exchange) [sExt64 warpSize - 1]-                    )-                  -- end sIf-                  sIf-                    (tvExp flag .==. statusP)-                    (readOffset <-- loopStop)-                    (readOffset <-- tvExp readOffset - zExt32 (tvExp used))-                  copyDWIMFix sharedReadOffset [0] (tvSize readOffset) []-                  scanOp''' <- renameLambda scanOp'-                  let (xs, ys) = splitAt (length tys) $ map paramName $ lambdaParams scanOp'''+                -- update prefix if flag different than STATUS_X:+                sWhen (tvExp flag .>. (statusX :: Imp.TExp Int8)) $ do+                  lam <- renameLambda scanOp'+                  let (xs, ys) = splitAt (length tys) $ map paramName $ lambdaParams lam                   forM_ (zip xs aggrs) $ \(x, aggr) -> dPrimV_ x (tvExp aggr)                   forM_ (zip ys prefixes) $ \(y, prefix) -> dPrimV_ y (tvExp prefix)-                  compileStms mempty (bodyStms $ lambdaBody scanOp''') $-                    forM_ (zip3 prefixes tys $ map resSubExp $ bodyResult $ lambdaBody scanOp''') $+                  compileStms mempty (bodyStms $ lambdaBody lam) $+                    forM_ (zip3 prefixes tys $ map resSubExp $ bodyResult $ lambdaBody lam) $                       \(prefix, ty, res) -> prefix <-- TPrimExp (toExp' ty res)-                -- end sWhen                 sOp localFence-                copyDWIMFix (tvVar readOffset) [] (Var sharedReadOffset) [0]           )         -- end sWhile         -- end sIf@@ -532,33 +602,28 @@               \(dest, res) ->                 copyDWIMFix dest [i] res [] -    sComment "Transpose scan output" $ do-      forM_ (zip transposedArrays privateArrays) $ \(trans, priv) -> do-        sOp localBarrier+    sComment "Transpose scan output and Write it to global memory in coalesced fashion" $ do+      forM_ (zip3 transposedArrays privateArrays $ map patElemName all_pes) $ \(locmem, priv, dest) -> do+        --sOp localBarrier         sFor "i" m $ \i -> do           sharedIdx <-             dPrimV "sharedIdx" $               sExt64 (kernelLocalThreadId constants * m) + i-          copyDWIMFix trans [tvExp sharedIdx] (Var priv) [i]+          copyDWIMFix locmem [tvExp sharedIdx] (Var priv) [i]         sOp localBarrier         sFor "i" m $ \i -> do-          sharedIdx <--            dPrimV "sharedIdx" $-              kernelLocalThreadId constants-                + sExt32 (kernelGroupSize constants * i)-          copyDWIMFix priv [i] (Var trans) [sExt64 $ tvExp sharedIdx]-      sOp localBarrier--    sComment "Write block scan results to global memory" $-      sFor "i" m $ \i -> do-        flat_idx <--          dPrimVE "flat_idx" $-            tvExp blockOff + kernelGroupSize constants * i-              + sExt64 (kernelLocalThreadId constants)-        dIndexSpace (zip gtids dims') flat_idx-        sWhen (flat_idx .<. n) $ do-          forM_ (zip (map patElemName all_pes) privateArrays) $ \(dest, src) ->-            copyDWIMFix dest (map Imp.vi64 gtids) (Var src) [i]+          flat_idx <-+            dPrimVE "flat_idx" $+              tvExp blockOff + kernelGroupSize constants * i+                + sExt64 (kernelLocalThreadId constants)+          dIndexSpace (zip gtids dims') flat_idx+          sWhen (flat_idx .<. n) $ do+            copyDWIMFix+              dest+              (map Imp.vi64 gtids)+              (Var locmem)+              [sExt64 $ flat_idx - tvExp blockOff]+        sOp localBarrier      sComment "If this is the last block, reset the dynamicId" $       sWhen (tvExp dynamicId .==. unCount num_groups - 1) $

@@ -121,18 +121,19 @@  readCarries ::   Imp.TExp Int64 ->+  Imp.TExp Int64 ->   [Imp.TExp Int64] ->   [Imp.TExp Int64] ->   [PatElem GPUMem] ->   SegBinOp GPUMem ->   InKernelGen ()-readCarries chunk_offset dims' vec_is pes scan+readCarries chunk_id chunk_offset dims' vec_is pes scan   | shapeRank (segBinOpShape scan) > 0 = do     ltid <- kernelLocalThreadId . kernelConstants <$> askEnv     -- We may have to reload the carries from the output of the     -- previous chunk.     sIf-      (chunk_offset .>. 0 .&&. ltid .==. 0)+      (chunk_id .>. 0 .&&. ltid .==. 0)       ( do           let is = unflattenIndex dims' $ chunk_offset - 1           forM_ (zip (xParams scan) pes) $ \(p, pe) ->@@ -218,6 +219,9 @@       sComment "threads in bounds read input" $         sWhen in_bounds when_in_bounds +      unless (all (null . segBinOpShape) scans) $+        sOp $ Imp.Barrier Imp.FenceGlobal+       forM_ (zip3 per_scan_pes scans all_local_arrs) $         \(pes, scan@(SegBinOp _ scan_op nes vec_shape), local_arrs) ->           sComment "do one intra-group scan operation" $ do@@ -233,7 +237,7 @@                   in_bounds                   ( do                       readToScanValues (map Imp.vi64 gtids ++ vec_is) pes scan-                      readCarries (tvExp chunk_offset) dims' vec_is pes scan+                      readCarries j (tvExp chunk_offset) dims' vec_is pes scan                   )                   ( forM_ (zip (yParams scan) (segBinOpNeutral scan)) $ \(p, ne) ->                       copyDWIMFix (paramName p) [] ne []

@@ -34,7 +34,7 @@     (map (setParamSpace space) inputs)     (setCodeSpace space body)     (map (setExtValueSpace space) results)-    (map (setExtValueSpace space) args)+    (map (fmap $ setExtValueSpace space) args)  setParamSpace :: Space -> Param -> Param setParamSpace space (MemParam name DefaultSpace) =

@@ -47,10 +47,11 @@     ExternalError s -> do       T.hPutStrLn stderr $ prettyText s       T.hPutStrLn stderr ""-      T.hPutStrLn stderr "If you find this error message confusing, uninformative, or wrong, please open an issue at\nhttps://github.com/diku-dk/futhark/issues."+      T.hPutStrLn stderr "If you find this error message confusing, uninformative, or wrong, please open an issue:"+      T.hPutStrLn stderr "  https://github.com/diku-dk/futhark/issues"     InternalError s info CompilerBug -> do-      T.hPutStrLn stderr "Internal compiler error."-      T.hPutStrLn stderr "Please report this at https://github.com/diku-dk/futhark/issues."+      T.hPutStrLn stderr "Internal compiler error.  Please report this:"+      T.hPutStrLn stderr "  https://github.com/diku-dk/futhark/issues"       report s info     InternalError s info CompilerLimitation -> do       T.hPutStrLn stderr "Known compiler limitation encountered.  Sorry."

@@ -50,7 +50,9 @@ pName :: Parser Name pName =   lexeme . fmap nameFromString $-    (:) <$> satisfy isAlpha <*> many (satisfy constituent)+    (:) <$> satisfy leading <*> many (satisfy constituent)+  where+    leading c = isAlpha c || c == '_'  pVName :: Parser VName pVName = lexeme $ do@@ -528,10 +530,11 @@ pEntry =   parens $     (,,) <$> (nameFromString <$> pStringLiteral)-      <* pComma <*> pEntryPointTypes+      <* pComma <*> pEntryPointInputs       <* pComma <*> pEntryPointTypes   where     pEntryPointTypes = braces (pEntryPointType `sepBy` pComma)+    pEntryPointInputs = braces (pEntryPointInput `sepBy` pComma)     pEntryPointType = do       u <- pUniqueness       choice@@ -539,6 +542,8 @@           "unsigned" $> TypeUnsigned u,           "opaque" *> parens (TypeOpaque u <$> pStringLiteral <* pComma <*> pInt)         ]+    pEntryPointInput =+      EntryParam <$> pName <* pColon <*> pEntryPointType  pFunDef :: PR rep -> Parser (FunDef rep) pFunDef pr = do

@@ -313,6 +313,9 @@   ppr (TypeUnsigned u) = ppr u <> "unsigned"   ppr (TypeOpaque u desc n) = ppr u <> "opaque" <> apply [ppr (show desc), ppr n] +instance Pretty EntryParam where+  ppr (EntryParam name t) = ppr name <> colon <+> ppr t+ instance PrettyRep rep => Pretty (FunDef rep) where   ppr (FunDef entry attrs name rettype fparams body) =     annot (attrAnnots attrs) $

@@ -155,6 +155,7 @@     LParam,     FunDef (..),     EntryPoint,+    EntryParam (..),     EntryPointType (..),     Prog (..), @@ -546,11 +547,6 @@  deriving instance RepTypes rep => Ord (FunDef rep) --- | Information about the parameters and return value of an entry--- point.  The first element is for parameters, the second for return--- value.-type EntryPoint = (Name, [EntryPointType], [EntryPointType])- -- | Every entry point argument and return value has an annotation -- indicating how it maps to the original source program type. data EntryPointType@@ -564,6 +560,17 @@   | -- | Maps directly.     TypeDirect Uniqueness   deriving (Eq, Show, Ord)++-- | An entry point parameter, comprising its name and original type.+data EntryParam = EntryParam+  { entryParamName :: Name,+    entryParamType :: EntryPointType+  }+  deriving (Eq, Show, Ord)++-- | Information about the inputs and outputs (return value) of an entry+-- point.+type EntryPoint = (Name, [EntryParam], [EntryPointType])  -- | An entire Futhark program. data Prog rep = Prog

@@ -99,11 +99,11 @@     zeroExts ts = generaliseExtTypes ts ts  generateEntryPoint :: E.EntryPoint -> E.ValBind -> InternaliseM ()-generateEntryPoint (E.EntryPoint e_paramts e_rettype) vb = localConstsScope $ do+generateEntryPoint (E.EntryPoint e_params e_rettype) vb = localConstsScope $ do   let (E.ValBind _ ofname _ (Info (rettype, _)) tparams params _ _ attrs loc) = vb   bindingFParams tparams params $ \shapeparams params' -> do     entry_rettype <- internaliseEntryReturnType rettype-    let entry' = entryPoint (baseName ofname) (zip e_paramts params') (e_rettype, entry_rettype)+    let entry' = entryPoint (baseName ofname) (zip e_params params') (e_rettype, entry_rettype)         args = map (I.Var . I.paramName) $ concat params'      (entry_body, ctx_ts) <- buildBody $ do@@ -131,40 +131,45 @@  entryPoint ::   Name ->-  [(E.EntryType, [I.FParam])] ->+  [(E.EntryParam, [I.FParam])] ->   ( E.EntryType,     [[I.TypeBase ExtShape Uniqueness]]   ) ->   I.EntryPoint entryPoint name params (eret, crets) =   ( name,-    concatMap (entryPointType . preParam) params,+    map onParam params,     case ( isTupleRecord $ entryType eret,            entryAscribed eret          ) of       (Just ts, Just (E.TETuple e_ts _)) ->-        concatMap entryPointType $-          zip (zipWith E.EntryType ts (map Just e_ts)) crets+        zipWith+          entryPointType+          (zipWith E.EntryType ts (map Just e_ts))+          crets       (Just ts, Nothing) ->-        concatMap entryPointType $-          zip (map (`E.EntryType` Nothing) ts) crets+        zipWith+          entryPointType+          (map (`E.EntryType` Nothing) ts)+          crets       _ ->-        entryPointType (eret, concat crets)+        [entryPointType eret $ concat crets]   )   where-    preParam (e_t, ps) = (e_t, staticShapes $ map I.paramDeclType ps)+    onParam (E.EntryParam e_p e_t, ps) =+      I.EntryParam e_p $ entryPointType e_t $ staticShapes $ map I.paramDeclType ps -    entryPointType (t, ts)+    entryPointType t ts       | E.Scalar (E.Prim E.Unsigned {}) <- E.entryType t =-        [I.TypeUnsigned u]+        I.TypeUnsigned u       | E.Array _ _ (E.Prim E.Unsigned {}) _ <- E.entryType t =-        [I.TypeUnsigned u]+        I.TypeUnsigned u       | E.Scalar E.Prim {} <- E.entryType t =-        [I.TypeDirect u]+        I.TypeDirect u       | E.Array _ _ E.Prim {} _ <- E.entryType t =-        [I.TypeDirect u]+        I.TypeDirect u       | otherwise =-        [I.TypeOpaque u desc $ length ts]+        I.TypeOpaque u desc $ length ts       where         u = foldl max Nonunique $ map I.uniqueness ts         desc = maybe (prettyOneLine t') typeExpOpaqueName $ E.entryAscribed t@@ -177,6 +182,7 @@             ++ "_"             ++ show (1 + d)             ++ "d"+    typeExpOpaqueName (TEUnique te _) = prettyOneLine te     typeExpOpaqueName te = prettyOneLine te      withoutDims (TEArray te _ _) =@@ -727,8 +733,8 @@     _ ->       pure e'   where-    traceRes tag' e' =-      mapM (letSubExp desc . BasicOp . Opaque (OpaqueTrace tag')) e'+    traceRes tag' =+      mapM (letSubExp desc . BasicOp . Opaque (OpaqueTrace tag'))     attr' = internaliseAttr attr     f env       | attr' == "unsafe",@@ -1767,12 +1773,12 @@       internaliseOperation desc e $ \v -> do         r <- I.arrayRank <$> lookupType v         let zero = intConst Int64 0-            offsets = offset' : replicate (r -1) zero+            offsets = offset' : replicate (r - 1) zero         return $ I.Rotate offsets v     handleRest [e] "transpose" = Just $ \desc ->       internaliseOperation desc e $ \v -> do         r <- I.arrayRank <$> lookupType v-        return $ I.Rearrange ([1, 0] ++ [2 .. r -1]) v+        return $ I.Rearrange ([1, 0] ++ [2 .. r - 1]) v     handleRest [TupLit [x, y] _] "zip" = Just $ \desc ->       mapM (letSubExp "zip_copy" . BasicOp . Copy)         =<< ( (++)

@@ -0,0 +1,248 @@+{-# LANGUAGE ConstraintKinds #-}+{-# LANGUAGE FlexibleContexts #-}+{-# LANGUAGE NamedFieldPuns #-}+{-# LANGUAGE TypeFamilies #-}++-- | This module implements an optimization that tries to statically reuse+-- kernel-level allocations. The goal is to lower the static memory usage, which+-- might allow more programs to run using intra-group parallelism.+module Futhark.Optimise.MemoryBlockMerging (optimise) where++import Control.Exception+import Control.Monad.Reader+import Control.Monad.State.Strict+import Data.Function ((&))+import Data.Map (Map, (!))+import qualified Data.Map as M+import Data.Set (Set)+import qualified Data.Set as S+import qualified Futhark.Analysis.Interference as Interference+import qualified Futhark.Analysis.LastUse as LastUse+import Futhark.Builder.Class+import Futhark.Construct+import Futhark.IR.GPUMem+import qualified Futhark.Optimise.MemoryBlockMerging.GreedyColoring as GreedyColoring+import Futhark.Pass (Pass (..), PassM)+import qualified Futhark.Pass as Pass+import Futhark.Util (invertMap)++-- | A mapping from allocation names to their size and space.+type Allocs = Map VName (SubExp, Space)++getAllocsStm :: Stm GPUMem -> Allocs+getAllocsStm (Let (Pat [PatElem name _]) _ (Op (Alloc se sp))) =+  M.singleton name (se, sp)+getAllocsStm (Let _ _ (Op (Alloc _ _))) = error "impossible"+getAllocsStm (Let _ _ (If _ then_body else_body _)) =+  foldMap getAllocsStm (bodyStms then_body)+    <> foldMap getAllocsStm (bodyStms else_body)+getAllocsStm (Let _ _ (DoLoop _ _ body)) =+  foldMap getAllocsStm (bodyStms body)+getAllocsStm _ = mempty++getAllocsSegOp :: SegOp lvl GPUMem -> Allocs+getAllocsSegOp (SegMap _ _ _ body) =+  foldMap getAllocsStm (kernelBodyStms body)+getAllocsSegOp (SegRed _ _ _ _ body) =+  foldMap getAllocsStm (kernelBodyStms body)+getAllocsSegOp (SegScan _ _ _ _ body) =+  foldMap getAllocsStm (kernelBodyStms body)+getAllocsSegOp (SegHist _ _ _ _ body) =+  foldMap getAllocsStm (kernelBodyStms body)++setAllocsStm :: Map VName SubExp -> Stm GPUMem -> Stm GPUMem+setAllocsStm m stm@(Let (Pat [PatElem name _]) _ (Op (Alloc _ _)))+  | Just s <- M.lookup name m =+    stm {stmExp = BasicOp $ SubExp s}+setAllocsStm _ stm@(Let _ _ (Op (Alloc _ _))) = stm+setAllocsStm m stm@(Let _ _ (Op (Inner (SegOp segop)))) =+  stm {stmExp = Op $ Inner $ SegOp $ setAllocsSegOp m segop}+setAllocsStm m stm@(Let _ _ (If cse then_body else_body dec)) =+  stm+    { stmExp =+        If+          cse+          (then_body {bodyStms = setAllocsStm m <$> bodyStms then_body})+          (else_body {bodyStms = setAllocsStm m <$> bodyStms else_body})+          dec+    }+setAllocsStm m stm@(Let _ _ (DoLoop merge form body)) =+  stm+    { stmExp =+        DoLoop merge form (body {bodyStms = setAllocsStm m <$> bodyStms body})+    }+setAllocsStm _ stm = stm++setAllocsSegOp ::+  Map VName SubExp ->+  SegOp lvl GPUMem ->+  SegOp lvl GPUMem+setAllocsSegOp m (SegMap lvl sp tps body) =+  SegMap lvl sp tps $+    body {kernelBodyStms = setAllocsStm m <$> kernelBodyStms body}+setAllocsSegOp m (SegRed lvl sp segbinops tps body) =+  SegRed lvl sp segbinops tps $+    body {kernelBodyStms = setAllocsStm m <$> kernelBodyStms body}+setAllocsSegOp m (SegScan lvl sp segbinops tps body) =+  SegScan lvl sp segbinops tps $+    body {kernelBodyStms = setAllocsStm m <$> kernelBodyStms body}+setAllocsSegOp m (SegHist lvl sp segbinops tps body) =+  SegHist lvl sp segbinops tps $+    body {kernelBodyStms = setAllocsStm m <$> kernelBodyStms body}++maxSubExp :: MonadBuilder m => Set SubExp -> m SubExp+maxSubExp = helper . S.toList+  where+    helper (s1 : s2 : sexps) = do+      z <- letSubExp "maxSubHelper" $ BasicOp $ BinOp (UMax Int64) s1 s2+      helper (z : sexps)+    helper [s] =+      return s+    helper [] = error "impossible"++definedInExp :: Exp GPUMem -> Set VName+definedInExp (Op (Inner (SegOp segop))) =+  definedInSegOp segop+definedInExp (If _ then_body else_body _) =+  foldMap definedInStm (bodyStms then_body)+    <> foldMap definedInStm (bodyStms else_body)+definedInExp (DoLoop _ _ body) =+  foldMap definedInStm $ bodyStms body+definedInExp _ = mempty++definedInStm :: Stm GPUMem -> Set VName+definedInStm Let {stmPat = Pat merge, stmExp} =+  let definedInside = merge & fmap patElemName & S.fromList+   in definedInExp stmExp <> definedInside++definedInSegOp :: SegOp lvl GPUMem -> Set VName+definedInSegOp (SegMap _ _ _ body) =+  foldMap definedInStm $ kernelBodyStms body+definedInSegOp (SegRed _ _ _ _ body) =+  foldMap definedInStm $ kernelBodyStms body+definedInSegOp (SegScan _ _ _ _ body) =+  foldMap definedInStm $ kernelBodyStms body+definedInSegOp (SegHist _ _ _ _ body) =+  foldMap definedInStm $ kernelBodyStms body++isKernelInvariant :: SegOp lvl GPUMem -> (SubExp, space) -> Bool+isKernelInvariant segop (Var vname, _) =+  not $ vname `S.member` definedInSegOp segop+isKernelInvariant _ _ = True++onKernelBodyStms ::+  MonadBuilder m =>+  SegOp lvl GPUMem ->+  (Stms GPUMem -> m (Stms GPUMem)) ->+  m (SegOp lvl GPUMem)+onKernelBodyStms (SegMap lvl space ts body) f = do+  stms <- f $ kernelBodyStms body+  return $ SegMap lvl space ts $ body {kernelBodyStms = stms}+onKernelBodyStms (SegRed lvl space binops ts body) f = do+  stms <- f $ kernelBodyStms body+  return $ SegRed lvl space binops ts $ body {kernelBodyStms = stms}+onKernelBodyStms (SegScan lvl space binops ts body) f = do+  stms <- f $ kernelBodyStms body+  return $ SegScan lvl space binops ts $ body {kernelBodyStms = stms}+onKernelBodyStms (SegHist lvl space binops ts body) f = do+  stms <- f $ kernelBodyStms body+  return $ SegHist lvl space binops ts $ body {kernelBodyStms = stms}++-- | This is the actual optimiser. Given an interference graph and a @SegOp@,+-- replace allocations and references to memory blocks inside with a (hopefully)+-- reduced number of allocations.+optimiseKernel ::+  (MonadBuilder m, Rep m ~ GPUMem) =>+  Interference.Graph VName ->+  SegOp lvl GPUMem ->+  m (SegOp lvl GPUMem)+optimiseKernel graph segop0 = do+  segop <- onKernelBodyStms segop0 $ onKernels $ optimiseKernel graph+  let allocs = M.filter (isKernelInvariant segop) $ getAllocsSegOp segop+      (colorspaces, coloring) =+        GreedyColoring.colorGraph+          (fmap snd allocs)+          graph+  (maxes, maxstms) <-+    invertMap coloring+      & M.elems+      & mapM (maxSubExp . S.map (fst . (allocs !)))+      & collectStms+  (colors, stms) <-+    assert (length maxes == M.size colorspaces) maxes+      & zip [0 ..]+      & mapM (\(i, x) -> letSubExp "color" $ Op $ Alloc x $ colorspaces ! i)+      & collectStms+  let segop' = setAllocsSegOp (fmap (colors !!) coloring) segop+  return $ case segop' of+    SegMap lvl sp tps body ->+      SegMap lvl sp tps $+        body {kernelBodyStms = maxstms <> stms <> kernelBodyStms body}+    SegRed lvl sp binops tps body ->+      SegRed lvl sp binops tps $+        body {kernelBodyStms = maxstms <> stms <> kernelBodyStms body}+    SegScan lvl sp binops tps body ->+      SegScan lvl sp binops tps $+        body {kernelBodyStms = maxstms <> stms <> kernelBodyStms body}+    SegHist lvl sp binops tps body ->+      SegHist lvl sp binops tps $+        body {kernelBodyStms = maxstms <> stms <> kernelBodyStms body}++-- | Helper function that modifies kernels found inside some statements.+onKernels ::+  LocalScope GPUMem m =>+  (SegOp SegLevel GPUMem -> m (SegOp SegLevel GPUMem)) ->+  Stms GPUMem ->+  m (Stms GPUMem)+onKernels f =+  mapM helper+  where+    helper stm@Let {stmExp = Op (Inner (SegOp segop))} =+      inScopeOf stm $ do+        exp' <- f segop+        return $ stm {stmExp = Op $ Inner $ SegOp exp'}+    helper stm@Let {stmExp = If c then_body else_body dec} =+      inScopeOf stm $ do+        then_body_stms <- f `onKernels` bodyStms then_body+        else_body_stms <- f `onKernels` bodyStms else_body+        return $+          stm+            { stmExp =+                If+                  c+                  (then_body {bodyStms = then_body_stms})+                  (else_body {bodyStms = else_body_stms})+                  dec+            }+    helper stm@Let {stmExp = DoLoop merge form body} =+      inScopeOf stm $ do+        stms <- f `onKernels` bodyStms body+        return $ stm {stmExp = DoLoop merge form (body {bodyStms = stms})}+    helper stm =+      inScopeOf stm $ return stm++-- | Perform the reuse-allocations optimization.+optimise :: Pass GPUMem GPUMem+optimise =+  Pass "reuse allocations" "reuse allocations" $ \prog ->+    let (lumap, _) = LastUse.analyseProg prog+        graph =+          foldMap+            ( \f ->+                runReader+                  ( Interference.analyseGPU lumap $+                      bodyStms $ funDefBody f+                  )+                  $ scopeOf f+            )+            $ progFuns prog+     in Pass.intraproceduralTransformation (onStms graph) prog+  where+    onStms ::+      Interference.Graph VName ->+      Scope GPUMem ->+      Stms GPUMem ->+      PassM (Stms GPUMem)+    onStms graph scope stms = do+      let m = localScope scope $ optimiseKernel graph `onKernels` stms+      fmap fst $ modifyNameSource $ runState (runBuilderT m mempty)

@@ -0,0 +1,59 @@+-- | Provides a greedy graph-coloring algorithm.+module Futhark.Optimise.MemoryBlockMerging.GreedyColoring (colorGraph, Coloring) where++import Data.Function ((&))+import qualified Data.Map as M+import Data.Maybe (fromMaybe)+import qualified Data.Set as S+import qualified Futhark.Analysis.Interference as Interference++-- | A map of values to their color, identified by an integer.+type Coloring a = M.Map a Int++-- | A map of values to the set "neighbors" in the graph+type Neighbors a = M.Map a (S.Set a)++-- | Computes the neighbor map of a graph.+neighbors :: Ord a => Interference.Graph a -> Neighbors a+neighbors =+  S.foldr+    ( \(x, y) acc ->+        acc+          & M.insertWith S.union x (S.singleton y)+          & M.insertWith S.union y (S.singleton x)+    )+    M.empty++firstAvailable :: Eq space => M.Map Int space -> S.Set Int -> Int -> space -> (M.Map Int space, Int)+firstAvailable spaces xs i sp =+  case (i `S.member` xs, spaces M.!? i) of+    (False, Just sp') | sp' == sp -> (spaces, i)+    (False, Nothing) -> (M.insert i sp spaces, i)+    _ -> firstAvailable spaces xs (i + 1) sp++colorNode ::+  (Ord a, Eq space) =>+  Neighbors a ->+  (a, space) ->+  (M.Map Int space, Coloring a) ->+  (M.Map Int space, Coloring a)+colorNode nbs (x, sp) (spaces, coloring) =+  let nb_colors =+        foldMap (maybe S.empty S.singleton . (coloring M.!?)) $+          fromMaybe mempty (nbs M.!? x)+      (spaces', color) = firstAvailable spaces nb_colors 0 sp+   in (spaces', M.insert x color coloring)++-- | Graph coloring that takes into account the @space@ of values. Two values+-- can only share the same color if they live in the same space. The result is+-- map from each color to a space and a map from each value in the input graph+-- to it's new color.+colorGraph ::+  (Ord a, Ord space) =>+  M.Map a space ->+  Interference.Graph a ->+  (M.Map Int space, Coloring a)+colorGraph spaces graph =+  let nodes = S.fromList $ M.toList spaces+      nbs = neighbors graph+   in S.foldr (colorNode nbs) mempty nodes

@@ -1,248 +0,0 @@-{-# LANGUAGE ConstraintKinds #-}-{-# LANGUAGE FlexibleContexts #-}-{-# LANGUAGE NamedFieldPuns #-}-{-# LANGUAGE TypeFamilies #-}---- | This module implements an optimization that tries to statically reuse--- kernel-level allocations. The goal is to lower the static memory usage, which--- might allow more programs to run using intra-group parallelism.-module Futhark.Optimise.ReuseAllocations (optimise) where--import Control.Exception-import Control.Monad.Reader-import Control.Monad.State.Strict-import Data.Function ((&))-import Data.Map (Map, (!))-import qualified Data.Map as M-import Data.Set (Set)-import qualified Data.Set as S-import qualified Futhark.Analysis.Interference as Interference-import qualified Futhark.Analysis.LastUse as LastUse-import Futhark.Builder.Class-import Futhark.Construct-import Futhark.IR.GPUMem-import qualified Futhark.Optimise.ReuseAllocations.GreedyColoring as GreedyColoring-import Futhark.Pass (Pass (..), PassM)-import qualified Futhark.Pass as Pass-import Futhark.Util (invertMap)---- | A mapping from allocation names to their size and space.-type Allocs = Map VName (SubExp, Space)--getAllocsStm :: Stm GPUMem -> Allocs-getAllocsStm (Let (Pat [PatElem name _]) _ (Op (Alloc se sp))) =-  M.singleton name (se, sp)-getAllocsStm (Let _ _ (Op (Alloc _ _))) = error "impossible"-getAllocsStm (Let _ _ (If _ then_body else_body _)) =-  foldMap getAllocsStm (bodyStms then_body)-    <> foldMap getAllocsStm (bodyStms else_body)-getAllocsStm (Let _ _ (DoLoop _ _ body)) =-  foldMap getAllocsStm (bodyStms body)-getAllocsStm _ = mempty--getAllocsSegOp :: SegOp lvl GPUMem -> Allocs-getAllocsSegOp (SegMap _ _ _ body) =-  foldMap getAllocsStm (kernelBodyStms body)-getAllocsSegOp (SegRed _ _ _ _ body) =-  foldMap getAllocsStm (kernelBodyStms body)-getAllocsSegOp (SegScan _ _ _ _ body) =-  foldMap getAllocsStm (kernelBodyStms body)-getAllocsSegOp (SegHist _ _ _ _ body) =-  foldMap getAllocsStm (kernelBodyStms body)--setAllocsStm :: Map VName SubExp -> Stm GPUMem -> Stm GPUMem-setAllocsStm m stm@(Let (Pat [PatElem name _]) _ (Op (Alloc _ _)))-  | Just s <- M.lookup name m =-    stm {stmExp = BasicOp $ SubExp s}-setAllocsStm _ stm@(Let _ _ (Op (Alloc _ _))) = stm-setAllocsStm m stm@(Let _ _ (Op (Inner (SegOp segop)))) =-  stm {stmExp = Op $ Inner $ SegOp $ setAllocsSegOp m segop}-setAllocsStm m stm@(Let _ _ (If cse then_body else_body dec)) =-  stm-    { stmExp =-        If-          cse-          (then_body {bodyStms = setAllocsStm m <$> bodyStms then_body})-          (else_body {bodyStms = setAllocsStm m <$> bodyStms else_body})-          dec-    }-setAllocsStm m stm@(Let _ _ (DoLoop merge form body)) =-  stm-    { stmExp =-        DoLoop merge form (body {bodyStms = setAllocsStm m <$> bodyStms body})-    }-setAllocsStm _ stm = stm--setAllocsSegOp ::-  Map VName SubExp ->-  SegOp lvl GPUMem ->-  SegOp lvl GPUMem-setAllocsSegOp m (SegMap lvl sp tps body) =-  SegMap lvl sp tps $-    body {kernelBodyStms = setAllocsStm m <$> kernelBodyStms body}-setAllocsSegOp m (SegRed lvl sp segbinops tps body) =-  SegRed lvl sp segbinops tps $-    body {kernelBodyStms = setAllocsStm m <$> kernelBodyStms body}-setAllocsSegOp m (SegScan lvl sp segbinops tps body) =-  SegScan lvl sp segbinops tps $-    body {kernelBodyStms = setAllocsStm m <$> kernelBodyStms body}-setAllocsSegOp m (SegHist lvl sp segbinops tps body) =-  SegHist lvl sp segbinops tps $-    body {kernelBodyStms = setAllocsStm m <$> kernelBodyStms body}--maxSubExp :: MonadBuilder m => Set SubExp -> m SubExp-maxSubExp = helper . S.toList-  where-    helper (s1 : s2 : sexps) = do-      z <- letSubExp "maxSubHelper" $ BasicOp $ BinOp (UMax Int64) s1 s2-      helper (z : sexps)-    helper [s] =-      return s-    helper [] = error "impossible"--definedInExp :: Exp GPUMem -> Set VName-definedInExp (Op (Inner (SegOp segop))) =-  definedInSegOp segop-definedInExp (If _ then_body else_body _) =-  foldMap definedInStm (bodyStms then_body)-    <> foldMap definedInStm (bodyStms else_body)-definedInExp (DoLoop _ _ body) =-  foldMap definedInStm $ bodyStms body-definedInExp _ = mempty--definedInStm :: Stm GPUMem -> Set VName-definedInStm Let {stmPat = Pat merge, stmExp} =-  let definedInside = merge & fmap patElemName & S.fromList-   in definedInExp stmExp <> definedInside--definedInSegOp :: SegOp lvl GPUMem -> Set VName-definedInSegOp (SegMap _ _ _ body) =-  foldMap definedInStm $ kernelBodyStms body-definedInSegOp (SegRed _ _ _ _ body) =-  foldMap definedInStm $ kernelBodyStms body-definedInSegOp (SegScan _ _ _ _ body) =-  foldMap definedInStm $ kernelBodyStms body-definedInSegOp (SegHist _ _ _ _ body) =-  foldMap definedInStm $ kernelBodyStms body--isKernelInvariant :: SegOp lvl GPUMem -> (SubExp, space) -> Bool-isKernelInvariant segop (Var vname, _) =-  not $ vname `S.member` definedInSegOp segop-isKernelInvariant _ _ = True--onKernelBodyStms ::-  MonadBuilder m =>-  SegOp lvl GPUMem ->-  (Stms GPUMem -> m (Stms GPUMem)) ->-  m (SegOp lvl GPUMem)-onKernelBodyStms (SegMap lvl space ts body) f = do-  stms <- f $ kernelBodyStms body-  return $ SegMap lvl space ts $ body {kernelBodyStms = stms}-onKernelBodyStms (SegRed lvl space binops ts body) f = do-  stms <- f $ kernelBodyStms body-  return $ SegRed lvl space binops ts $ body {kernelBodyStms = stms}-onKernelBodyStms (SegScan lvl space binops ts body) f = do-  stms <- f $ kernelBodyStms body-  return $ SegScan lvl space binops ts $ body {kernelBodyStms = stms}-onKernelBodyStms (SegHist lvl space binops ts body) f = do-  stms <- f $ kernelBodyStms body-  return $ SegHist lvl space binops ts $ body {kernelBodyStms = stms}---- | This is the actual optimiser. Given an interference graph and a @SegOp@,--- replace allocations and references to memory blocks inside with a (hopefully)--- reduced number of allocations.-optimiseKernel ::-  (MonadBuilder m, Rep m ~ GPUMem) =>-  Interference.Graph VName ->-  SegOp lvl GPUMem ->-  m (SegOp lvl GPUMem)-optimiseKernel graph segop0 = do-  segop <- onKernelBodyStms segop0 $ onKernels $ optimiseKernel graph-  let allocs = M.filter (isKernelInvariant segop) $ getAllocsSegOp segop-      (colorspaces, coloring) =-        GreedyColoring.colorGraph-          (fmap snd allocs)-          graph-  (maxes, maxstms) <--    invertMap coloring-      & M.elems-      & mapM (maxSubExp . S.map (fst . (allocs !)))-      & collectStms-  (colors, stms) <--    assert (length maxes == M.size colorspaces) maxes-      & zip [0 ..]-      & mapM (\(i, x) -> letSubExp "color" $ Op $ Alloc x $ colorspaces ! i)-      & collectStms-  let segop' = setAllocsSegOp (fmap (colors !!) coloring) segop-  return $ case segop' of-    SegMap lvl sp tps body ->-      SegMap lvl sp tps $-        body {kernelBodyStms = maxstms <> stms <> kernelBodyStms body}-    SegRed lvl sp binops tps body ->-      SegRed lvl sp binops tps $-        body {kernelBodyStms = maxstms <> stms <> kernelBodyStms body}-    SegScan lvl sp binops tps body ->-      SegScan lvl sp binops tps $-        body {kernelBodyStms = maxstms <> stms <> kernelBodyStms body}-    SegHist lvl sp binops tps body ->-      SegHist lvl sp binops tps $-        body {kernelBodyStms = maxstms <> stms <> kernelBodyStms body}---- | Helper function that modifies kernels found inside some statements.-onKernels ::-  LocalScope GPUMem m =>-  (SegOp SegLevel GPUMem -> m (SegOp SegLevel GPUMem)) ->-  Stms GPUMem ->-  m (Stms GPUMem)-onKernels f =-  mapM helper-  where-    helper stm@Let {stmExp = Op (Inner (SegOp segop))} =-      inScopeOf stm $ do-        exp' <- f segop-        return $ stm {stmExp = Op $ Inner $ SegOp exp'}-    helper stm@Let {stmExp = If c then_body else_body dec} =-      inScopeOf stm $ do-        then_body_stms <- f `onKernels` bodyStms then_body-        else_body_stms <- f `onKernels` bodyStms else_body-        return $-          stm-            { stmExp =-                If-                  c-                  (then_body {bodyStms = then_body_stms})-                  (else_body {bodyStms = else_body_stms})-                  dec-            }-    helper stm@Let {stmExp = DoLoop merge form body} =-      inScopeOf stm $ do-        stms <- f `onKernels` bodyStms body-        return $ stm {stmExp = DoLoop merge form (body {bodyStms = stms})}-    helper stm =-      inScopeOf stm $ return stm---- | Perform the reuse-allocations optimization.-optimise :: Pass GPUMem GPUMem-optimise =-  Pass "reuse allocations" "reuse allocations" $ \prog ->-    let (lumap, _) = LastUse.analyseProg prog-        graph =-          foldMap-            ( \f ->-                runReader-                  ( Interference.analyseGPU lumap $-                      bodyStms $ funDefBody f-                  )-                  $ scopeOf f-            )-            $ progFuns prog-     in Pass.intraproceduralTransformation (onStms graph) prog-  where-    onStms ::-      Interference.Graph VName ->-      Scope GPUMem ->-      Stms GPUMem ->-      PassM (Stms GPUMem)-    onStms graph scope stms = do-      let m = localScope scope $ optimiseKernel graph `onKernels` stms-      fmap fst $ modifyNameSource $ runState (runBuilderT m mempty)

@@ -1,59 +0,0 @@--- | Provides a greedy graph-coloring algorithm.-module Futhark.Optimise.ReuseAllocations.GreedyColoring (colorGraph, Coloring) where--import Data.Function ((&))-import qualified Data.Map as M-import Data.Maybe (fromMaybe)-import qualified Data.Set as S-import qualified Futhark.Analysis.Interference as Interference---- | A map of values to their color, identified by an integer.-type Coloring a = M.Map a Int---- | A map of values to the set "neighbors" in the graph-type Neighbors a = M.Map a (S.Set a)---- | Computes the neighbor map of a graph.-neighbors :: Ord a => Interference.Graph a -> Neighbors a-neighbors =-  S.foldr-    ( \(x, y) acc ->-        acc-          & M.insertWith S.union x (S.singleton y)-          & M.insertWith S.union y (S.singleton x)-    )-    M.empty--firstAvailable :: Eq space => M.Map Int space -> S.Set Int -> Int -> space -> (M.Map Int space, Int)-firstAvailable spaces xs i sp =-  case (i `S.member` xs, spaces M.!? i) of-    (False, Just sp') | sp' == sp -> (spaces, i)-    (False, Nothing) -> (M.insert i sp spaces, i)-    _ -> firstAvailable spaces xs (i + 1) sp--colorNode ::-  (Ord a, Eq space) =>-  Neighbors a ->-  (a, space) ->-  (M.Map Int space, Coloring a) ->-  (M.Map Int space, Coloring a)-colorNode nbs (x, sp) (spaces, coloring) =-  let nb_colors =-        foldMap (maybe S.empty S.singleton . (coloring M.!?)) $-          fromMaybe mempty (nbs M.!? x)-      (spaces', color) = firstAvailable spaces nb_colors 0 sp-   in (spaces', M.insert x color coloring)---- | Graph coloring that takes into account the @space@ of values. Two values--- can only share the same color if they live in the same space. The result is--- map from each color to a space and a map from each value in the input graph--- to it's new color.-colorGraph ::-  (Ord a, Ord space) =>-  M.Map a space ->-  Interference.Graph a ->-  (M.Map Int space, Coloring a)-colorGraph spaces graph =-  let nodes = S.fromList $ M.toList spaces-      nbs = neighbors graph-   in S.foldr (colorNode nbs) mempty nodes

@@ -25,7 +25,7 @@ import Futhark.Optimise.Fusion import Futhark.Optimise.InPlaceLowering import Futhark.Optimise.InliningDeadFun-import qualified Futhark.Optimise.ReuseAllocations as ReuseAllocations+import qualified Futhark.Optimise.MemoryBlockMerging as MemoryBlockMerging import Futhark.Optimise.Sink import Futhark.Optimise.TileLoops import Futhark.Optimise.Unstream@@ -104,7 +104,7 @@         simplifyGPUMem,         doubleBufferGPU,         simplifyGPUMem,-        ReuseAllocations.optimise,+        MemoryBlockMerging.optimise,         simplifyGPUMem,         expandAllocations,         simplifyGPUMem

@@ -78,6 +78,7 @@     ValBindBase (..),     EntryPoint (..),     EntryType (..),+    EntryParam (..),     Liftedness (..),     TypeBindBase (..),     TypeParamBase (..),@@ -971,13 +972,20 @@   }   deriving (Show) +-- | A parameter of an entry point.+data EntryParam = EntryParam+  { entryParamName :: Name,+    entryParamType :: EntryType+  }+  deriving (Show)+ -- | Information about the external interface exposed by an entry -- point.  The important thing is that that we remember the original -- source-language types, without desugaring them at all.  The -- annoying thing is that we do not require type annotations on entry -- points, so the types can be either ascribed or inferred. data EntryPoint = EntryPoint-  { entryParams :: [EntryType],+  { entryParams :: [EntryParam],     entryReturn :: EntryType   }   deriving (Show)

@@ -591,22 +591,30 @@ entryPoint params orig_ret_te orig_ret =   EntryPoint (map patternEntry params ++ more_params) rettype'   where-    (more_params, rettype') =-      onRetType orig_ret_te orig_ret+    (more_params, rettype') = onRetType orig_ret_te orig_ret      patternEntry (PatParens p _) =       patternEntry p-    patternEntry (PatAscription _ tdecl _) =-      EntryType (unInfo (expandedType tdecl)) (Just (declaredType tdecl))+    patternEntry (PatAscription p tdecl _) =+      EntryParam (patternName p) $+        EntryType (unInfo (expandedType tdecl)) (Just (declaredType tdecl))     patternEntry p =-      EntryType (patternStructType p) Nothing+      EntryParam (patternName p) $+        EntryType (patternStructType p) Nothing -    onRetType (Just (TEArrow _ t1_te t2_te _)) (Scalar (Arrow _ _ t1 t2)) =+    patternName (Id x _ _) = baseName x+    patternName (PatParens p _) = patternName p+    patternName _ = "_"++    pname (Named v) = baseName v+    pname Unnamed = "_"++    onRetType (Just (TEArrow p t1_te t2_te _)) (Scalar (Arrow _ _ t1 t2)) =       let (xs, y) = onRetType (Just t2_te) t2-       in (EntryType t1 (Just t1_te) : xs, y)-    onRetType _ (Scalar (Arrow _ _ t1 t2)) =+       in (EntryParam (maybe "_" baseName p) (EntryType t1 (Just t1_te)) : xs, y)+    onRetType _ (Scalar (Arrow _ p t1 t2)) =       let (xs, y) = onRetType Nothing t2-       in (EntryType t1 Nothing : xs, y)+       in (EntryParam (pname p) (EntryType t1 Nothing) : xs, y)     onRetType te t =       ([], EntryType t te) @@ -698,6 +706,7 @@ niceTypeExp (TEVar (QualName [] _) _) = True niceTypeExp (TEApply te TypeArgExpDim {} _) = niceTypeExp te niceTypeExp (TEArray te _ _) = niceTypeExp te+niceTypeExp (TEUnique te _) = niceTypeExp te niceTypeExp _ = False  checkOneDec :: DecBase NoInfo Name -> TypeM (TySet, Env, DecBase Info VName)

@@ -32,6 +32,7 @@ import qualified Data.Map.Strict as M import Data.Maybe import qualified Data.Set as S+import qualified Data.Version as Version import Futhark.IR.Primitive (intByteSize) import Futhark.Util (nubOrd) import Futhark.Util.Pretty hiding (bool, group, space)@@ -44,6 +45,7 @@ import Language.Futhark.TypeChecker.Types hiding (checkTypeDecl) import qualified Language.Futhark.TypeChecker.Types as Types import Language.Futhark.TypeChecker.Unify hiding (Usage)+import qualified Paths_futhark import Prelude hiding (mod)  --- Uniqueness@@ -105,7 +107,7 @@ combineOccurences name (Observed rloc) (Consumed wloc) =   useAfterConsume name rloc wloc combineOccurences name (Consumed loc1) (Consumed loc2) =-  consumeAfterConsume name (max loc1 loc2) (min loc1 loc2)+  useAfterConsume name (max loc1 loc2) (min loc1 loc2)  checkOccurences :: Occurences -> TermTypeM () checkOccurences = void . M.traverseWithKey comb . usageMap@@ -692,20 +694,30 @@  --- Errors +errorIndexUrl :: Doc+errorIndexUrl = version_url <> "error-index.html"+  where+    version = Paths_futhark.version+    base_url = "https://futhark.readthedocs.io/en/"+    version_url+      | last (Version.versionBranch version) == 0 = base_url <> "latest/"+      | otherwise = base_url <> "v" <> text (Version.showVersion version) <> "/"++withIndexLink :: Doc -> Doc -> Doc+withIndexLink href msg =+  stack+    [ msg,+      "\nFor more information, see:",+      indent 2 (ppr errorIndexUrl <> "#" <> href)+    ]+ useAfterConsume :: VName -> SrcLoc -> SrcLoc -> TermTypeM a useAfterConsume name rloc wloc = do   name' <- describeVar rloc name-  typeError rloc mempty $+  typeError rloc mempty . withIndexLink "use-after-consume" $     "Using" <+> name' <> ", but this was consumed at"       <+> text (locStrRel rloc wloc) <> ".  (Possibly through aliasing.)" -consumeAfterConsume :: VName -> SrcLoc -> SrcLoc -> TermTypeM a-consumeAfterConsume name loc1 loc2 = do-  name' <- describeVar loc1 name-  typeError loc2 mempty $-    "Consuming" <+> name' <> ", but this was previously consumed at"-      <+> text (locStrRel loc2 loc1) <> "."- badLetWithValue :: (Pretty arr, Pretty src) => arr -> src -> SrcLoc -> TermTypeM a badLetWithValue arre vale loc =   typeError loc mempty $@@ -717,17 +729,17 @@  returnAliased :: Name -> Name -> SrcLoc -> TermTypeM () returnAliased fname name loc =-  typeError loc mempty $-    "Unique return value of" <+> pquote (pprName fname)+  typeError loc mempty . withIndexLink "return-aliased" $+    "Unique-typed return value of" <+> pquote (pprName fname)       <+> "is aliased to"-      <+> pquote (pprName name) <> ", which is not consumed."+      <+> pquote (pprName name) <> ", which is not consumable."  uniqueReturnAliased :: Name -> SrcLoc -> TermTypeM a uniqueReturnAliased fname loc =-  typeError loc mempty $-    "A unique tuple element of return value of"+  typeError loc mempty . withIndexLink "unique-return-aliased" $+    "A unique-typed component of the return value of"       <+> pquote (pprName fname)-      <+> "is aliased to some other tuple component."+      <+> "is aliased to some other component."  unexpectedType :: MonadTypeChecker m => SrcLoc -> StructType -> [StructType] -> m a unexpectedType loc _ [] =@@ -740,6 +752,16 @@       <+> commasep (map ppr ts) <> ", but is"       <+> ppr t <> "." +notConsumable :: MonadTypeChecker m => SrcLoc -> Doc -> m b+notConsumable loc v =+  typeError loc mempty . withIndexLink "not-consumable" $+    "Would consume" <+> v <> ", which is not consumable."++unusedSize :: (MonadTypeChecker m) => SizeBinder VName -> m a+unusedSize p =+  typeError p mempty . withIndexLink "unused-size" $+    "Size" <+> ppr p <+> "unused in pattern."+ --- Basic checking  -- | Determine if the two types of identical, ignoring uniqueness.@@ -1117,8 +1139,7 @@     let used_sizes = typeDimNames $ patternStructType p'     case filter ((`S.notMember` used_sizes) . sizeName) sizes of       [] -> m p'-      size : _ ->-        typeError size mempty $ "Size" <+> ppr size <+> "unused in pattern."+      size : _ -> unusedSize size  patternDims :: Pat -> [Ident] patternDims (PatParens p _) = patternDims p@@ -1362,11 +1383,10 @@       maybe_sloc <- gets $ M.lookup f       case maybe_sloc of         Just sloc ->-          lift $-            typeError rloc mempty $-              "Field" <+> pquote (ppr f)-                <+> "previously defined at"-                <+> text (locStrRel rloc sloc) <> "."+          lift . typeError rloc mempty $+            "Field" <+> pquote (ppr f)+              <+> "previously defined at"+              <+> text (locStrRel rloc sloc) <> "."         Nothing -> return () checkExp (ArrayLit all_es _ loc) =   -- Construct the result type and unify all elements with it.  We@@ -1620,21 +1640,7 @@      (elemt, _) <- sliceShape (Just (loc, Nonrigid)) slice' =<< normTypeFully t -    unless (unique src_t) $-      typeError loc mempty $-        "Source" <+> pquote (pprName (identName src))-          <+> "has type"-          <+> ppr src_t <> ", which is not unique."-    vtable <- asks $ scopeVtable . termScope-    forM_ (aliases src_t) $ \v ->-      case aliasVar v `M.lookup` vtable of-        Just (BoundV Local _ v_t)-          | not $ unique v_t ->-            typeError loc mempty $-              "Source" <+> pquote (pprName (identName src))-                <+> "aliases"-                <+> pquote (pprName (aliasVar v)) <> ", which is not consumable."-        _ -> return ()+    unless (unique src_t) $ notConsumable loc $ pquote $ pprName $ identName src      sequentially (unifies "type of target array" (toStruct elemt) =<< checkExp ve) $ \ve' _ -> do       ve_t <- expTypeFully ve'@@ -1655,12 +1661,9 @@   sequentially (checkExp ve >>= unifies "type of target array" elemt) $ \ve' _ ->     sequentially (checkExp src >>= unifies "type of target array" t) $ \src' _ -> do       src_t <- expTypeFully src'-      unless (unique src_t) $-        typeError loc mempty $-          "Source" <+> pquote (ppr src)-            <+> "has type"-            <+> ppr src_t <> ", which is not unique." +      unless (unique src_t) $ notConsumable loc $ pquote $ ppr src+       let src_als = aliases src_t       ve_t <- expTypeFully ve'       unless (S.null $ src_als `S.intersection` aliases ve_t) $ badLetWithValue src ve loc@@ -2601,7 +2604,7 @@      causality what loc d dloc t =       Left $-        TypeError loc mempty $+        TypeError loc mempty . withIndexLink "causality-check" $           "Causality check: size" <+/> pquote (pprName d)             <+/> "needed for type of"             <+> what <> colon@@ -3113,10 +3116,8 @@             rettype'' : map patternStructType params     case filter ((`S.notMember` used_sizes) . typeParamName) $       filter isSizeParam tparams' of-      [] -> return ()-      tp : _ ->-        typeError defloc mempty $-          "Size parameter" <+> pquote (ppr tp) <+> "unused."+      [] -> pure ()+      tp : _ -> unusedSize $ SizeBinder (typeParamName tp) (srclocOf tp)      -- We keep those type variables that were not closed over by     -- let-generalisation.@@ -3192,15 +3193,13 @@         Just (BoundV Local _ t)           | arrayRank t > 0 -> unique t           | Scalar TypeVar {} <- t -> unique t+          | Scalar Arrow {} <- t -> False           | otherwise -> True         Just (BoundV Global _ _) -> False         _ -> True   -- The sort ensures that AliasBound vars are shown before AliasFree.   case map aliasVar $ sort $ filter (not . consumable . aliasVar) $ S.toList als of-    v : _ -> do-      v' <- describeVar loc v-      typeError loc mempty $-        "Would consume" <+> v' <> ", which is not allowed."+    v : _ -> notConsumable loc =<< describeVar loc v     [] -> return ()  -- | Proclaim that we have written to the given variable.@@ -3214,7 +3213,8 @@ -- computation. consuming :: Ident -> TermTypeM a -> TermTypeM a consuming (Ident name (Info t) loc) m = do-  consume loc $ AliasBound name `S.insert` aliases t+  t' <- normTypeFully t+  consume loc $ AliasBound name `S.insert` aliases t'   localScope consume' m   where     consume' scope =

@@ -0,0 +1,70 @@+module Futhark.Optimise.MemoryBlockMerging.GreedyColoringTests+  ( tests,+  )+where++import Control.Arrow ((***))+import Data.Function ((&))+import qualified Data.Map as M+import qualified Data.Set as S+import qualified Futhark.Optimise.MemoryBlockMerging.GreedyColoring as GreedyColoring+import Test.Tasty+import Test.Tasty.HUnit++tests :: TestTree+tests =+  testGroup+    "GreedyColoringTests"+    [psumTest, allIntersect, emptyGraph, noIntersections, differentSpaces]++psumTest :: TestTree+psumTest =+  testCase "psumTest" $+    assertEqual+      "Color simple 1-2-3 using two colors"+      ([(0, "local"), (1, "local")], [(1 :: Int, 0), (2, 1), (3, 0)])+      $ (M.toList *** M.toList) $+        GreedyColoring.colorGraph+          (M.fromList [(1, "local"), (2, "local"), (3, "local")])+          $ S.fromList [(1, 2), (2, 3)]++allIntersect :: TestTree+allIntersect =+  testCase "allIntersect" $+    assertEqual+      "Color a graph where all values intersect"+      ([(0, "local"), (1, "local"), (2, "local")], [(1 :: Int, 2), (2, 1), (3, 0)])+      $ (M.toList *** M.toList) $+        GreedyColoring.colorGraph+          (M.fromList [(1, "local"), (2, "local"), (3, "local")])+          $ S.fromList [(1, 2), (2, 3), (1, 3)]++emptyGraph :: TestTree+emptyGraph =+  testCase "emptyGraph" $+    assertEqual+      "Color an empty graph"+      ([] :: [(Int, Char)], [] :: [(Int, Int)])+      $ (M.toList *** M.toList) $ GreedyColoring.colorGraph M.empty $ S.fromList []++noIntersections :: TestTree+noIntersections =+  GreedyColoring.colorGraph+    (M.fromList [(1, "local"), (2, "local"), (3, "local")])+    (S.fromList [])+    & M.toList *** M.toList+    & assertEqual+      "Color nodes with no intersections"+      ([(0, "local")], [(1, 0), (2, 0), (3, 0)] :: [(Int, Int)])+    & testCase "noIntersections"++differentSpaces :: TestTree+differentSpaces =+  GreedyColoring.colorGraph+    (M.fromList [(1, "a"), (2, "b"), (3, "c")])+    (S.fromList [])+    & M.toList *** M.toList+    & assertEqual+      "Color nodes with no intersections but in different spaces"+      ([(0, "c"), (1, "b"), (2, "a")], [(1, 2), (2, 1), (3, 0)] :: [(Int, Int)])+    & testCase "differentSpaces"

@@ -1,70 +0,0 @@-module Futhark.Optimise.ReuseAllocations.GreedyColoringTests-  ( tests,-  )-where--import Control.Arrow ((***))-import Data.Function ((&))-import qualified Data.Map as M-import qualified Data.Set as S-import qualified Futhark.Optimise.ReuseAllocations.GreedyColoring as GreedyColoring-import Test.Tasty-import Test.Tasty.HUnit--tests :: TestTree-tests =-  testGroup-    "GreedyColoringTests"-    [psumTest, allIntersect, emptyGraph, noIntersections, differentSpaces]--psumTest :: TestTree-psumTest =-  testCase "psumTest" $-    assertEqual-      "Color simple 1-2-3 using two colors"-      ([(0, "local"), (1, "local")], [(1 :: Int, 0), (2, 1), (3, 0)])-      $ (M.toList *** M.toList) $-        GreedyColoring.colorGraph-          (M.fromList [(1, "local"), (2, "local"), (3, "local")])-          $ S.fromList [(1, 2), (2, 3)]--allIntersect :: TestTree-allIntersect =-  testCase "allIntersect" $-    assertEqual-      "Color a graph where all values intersect"-      ([(0, "local"), (1, "local"), (2, "local")], [(1 :: Int, 2), (2, 1), (3, 0)])-      $ (M.toList *** M.toList) $-        GreedyColoring.colorGraph-          (M.fromList [(1, "local"), (2, "local"), (3, "local")])-          $ S.fromList [(1, 2), (2, 3), (1, 3)]--emptyGraph :: TestTree-emptyGraph =-  testCase "emptyGraph" $-    assertEqual-      "Color an empty graph"-      ([] :: [(Int, Char)], [] :: [(Int, Int)])-      $ (M.toList *** M.toList) $ GreedyColoring.colorGraph M.empty $ S.fromList []--noIntersections :: TestTree-noIntersections =-  GreedyColoring.colorGraph-    (M.fromList [(1, "local"), (2, "local"), (3, "local")])-    (S.fromList [])-    & M.toList *** M.toList-    & assertEqual-      "Color nodes with no intersections"-      ([(0, "local")], [(1, 0), (2, 0), (3, 0)] :: [(Int, Int)])-    & testCase "noIntersections"--differentSpaces :: TestTree-differentSpaces =-  GreedyColoring.colorGraph-    (M.fromList [(1, "a"), (2, "b"), (3, "c")])-    (S.fromList [])-    & M.toList *** M.toList-    & assertEqual-      "Color nodes with no intersections but in different spaces"-      ([(0, "c"), (1, "b"), (2, "a")], [(1, 2), (2, 1), (3, 0)] :: [(Int, Int)])-    & testCase "differentSpaces"

@@ -5,7 +5,7 @@ import qualified Futhark.IR.PrimitiveTests import qualified Futhark.IR.PropTests import qualified Futhark.IR.Syntax.CoreTests-import qualified Futhark.Optimise.ReuseAllocations.GreedyColoringTests+import qualified Futhark.Optimise.MemoryBlockMerging.GreedyColoringTests import qualified Futhark.Pkg.SolveTests import qualified Language.Futhark.SyntaxTests import Test.Tasty@@ -21,7 +21,7 @@       Futhark.Pkg.SolveTests.tests,       Futhark.IR.Mem.IxFunTests.tests,       Futhark.IR.PrimitiveTests.tests,-      Futhark.Optimise.ReuseAllocations.GreedyColoringTests.tests+      Futhark.Optimise.MemoryBlockMerging.GreedyColoringTests.tests     ]  main :: IO ()