diff --git a/docs/c-api.rst b/docs/c-api.rst
--- a/docs/c-api.rst
+++ b/docs/c-api.rst
@@ -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.
diff --git a/docs/error-index.rst b/docs/error-index.rst
new file mode 100644
--- /dev/null
+++ b/docs/error-index.rst
@@ -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.
diff --git a/docs/index.rst b/docs/index.rst
--- a/docs/index.rst
+++ b/docs/index.rst
@@ -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
diff --git a/docs/language-reference.rst b/docs/language-reference.rst
--- a/docs/language-reference.rst
+++ b/docs/language-reference.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
 ~~~~~~~~~~~~~~~~~~~~~
 
diff --git a/docs/man/futhark-c.rst b/docs/man/futhark-c.rst
--- a/docs/man/futhark-c.rst
+++ b/docs/man/futhark-c.rst
@@ -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
diff --git a/docs/man/futhark-cuda.rst b/docs/man/futhark-cuda.rst
--- a/docs/man/futhark-cuda.rst
+++ b/docs/man/futhark-cuda.rst
@@ -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
diff --git a/docs/man/futhark-literate.rst b/docs/man/futhark-literate.rst
--- a/docs/man/futhark-literate.rst
+++ b/docs/man/futhark-literate.rst
@@ -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``
 
diff --git a/docs/man/futhark-multicore.rst b/docs/man/futhark-multicore.rst
--- a/docs/man/futhark-multicore.rst
+++ b/docs/man/futhark-multicore.rst
@@ -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
diff --git a/docs/man/futhark-opencl.rst b/docs/man/futhark-opencl.rst
--- a/docs/man/futhark-opencl.rst
+++ b/docs/man/futhark-opencl.rst
@@ -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.
diff --git a/docs/usage.rst b/docs/usage.rst
--- a/docs/usage.rst
+++ b/docs/usage.rst
@@ -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
 
diff --git a/futhark.cabal b/futhark.cabal
--- a/futhark.cabal
+++ b/futhark.cabal
@@ -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
diff --git a/rts/c/cuda.h b/rts/c/cuda.h
--- a/rts/c/cuda.h
+++ b/rts/c/cuda.h
@@ -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.
diff --git a/src/Futhark/Actions.hs b/src/Futhark/Actions.hs
--- a/src/Futhark/Actions.hs
+++ b/src/Futhark/Actions.hs
@@ -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
diff --git a/src/Futhark/CLI/Dataset.hs b/src/Futhark/CLI/Dataset.hs
--- a/src/Futhark/CLI/Dataset.hs
+++ b/src/Futhark/CLI/Dataset.hs
@@ -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
diff --git a/src/Futhark/CLI/Dev.hs b/src/Futhark/CLI/Dev.hs
--- a/src/Futhark/CLI/Dev.hs
+++ b/src/Futhark/CLI/Dev.hs
@@ -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
diff --git a/src/Futhark/CLI/Literate.hs b/src/Futhark/CLI/Literate.hs
--- a/src/Futhark/CLI/Literate.hs
+++ b/src/Futhark/CLI/Literate.hs
@@ -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
diff --git a/src/Futhark/CodeGen/Backends/CCUDA/Boilerplate.hs b/src/Futhark/CodeGen/Backends/CCUDA/Boilerplate.hs
--- a/src/Futhark/CodeGen/Backends/CCUDA/Boilerplate.hs
+++ b/src/Futhark/CodeGen/Backends/CCUDA/Boilerplate.hs
@@ -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)));
diff --git a/src/Futhark/CodeGen/Backends/GenericC.hs b/src/Futhark/CodeGen/Backends/GenericC.hs
--- a/src/Futhark/CodeGen/Backends/GenericC.hs
+++ b/src/Futhark/CodeGen/Backends/GenericC.hs
@@ -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
diff --git a/src/Futhark/CodeGen/Backends/GenericC/CLI.hs b/src/Futhark/CodeGen/Backends/GenericC/CLI.hs
--- a/src/Futhark/CodeGen/Backends/GenericC/CLI.hs
+++ b/src/Futhark/CodeGen/Backends/GenericC/CLI.hs
@@ -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;
diff --git a/src/Futhark/CodeGen/Backends/GenericC/Manifest.hs b/src/Futhark/CodeGen/Backends/GenericC/Manifest.hs
new file mode 100644
--- /dev/null
+++ b/src/Futhark/CodeGen/Backends/GenericC/Manifest.hs
@@ -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
diff --git a/src/Futhark/CodeGen/Backends/GenericC/Server.hs b/src/Futhark/CodeGen/Backends/GenericC/Server.hs
--- a/src/Futhark/CodeGen/Backends/GenericC/Server.hs
+++ b/src/Futhark/CodeGen/Backends/GenericC/Server.hs
@@ -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>")
diff --git a/src/Futhark/CodeGen/Backends/GenericPython.hs b/src/Futhark/CodeGen/Backends/GenericPython.hs
--- a/src/Futhark/CodeGen/Backends/GenericPython.hs
+++ b/src/Futhark/CodeGen/Backends/GenericPython.hs
@@ -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)"]]
diff --git a/src/Futhark/CodeGen/Backends/MulticoreWASM.hs b/src/Futhark/CodeGen/Backends/MulticoreWASM.hs
--- a/src/Futhark/CodeGen/Backends/MulticoreWASM.hs
+++ b/src/Futhark/CodeGen/Backends/MulticoreWASM.hs
@@ -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
diff --git a/src/Futhark/CodeGen/Backends/SequentialWASM.hs b/src/Futhark/CodeGen/Backends/SequentialWASM.hs
--- a/src/Futhark/CodeGen/Backends/SequentialWASM.hs
+++ b/src/Futhark/CodeGen/Backends/SequentialWASM.hs
@@ -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
diff --git a/src/Futhark/CodeGen/Backends/SimpleRep.hs b/src/Futhark/CodeGen/Backends/SimpleRep.hs
--- a/src/Futhark/CodeGen/Backends/SimpleRep.hs
+++ b/src/Futhark/CodeGen/Backends/SimpleRep.hs
@@ -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
diff --git a/src/Futhark/CodeGen/ImpCode.hs b/src/Futhark/CodeGen/ImpCode.hs
--- a/src/Futhark/CodeGen/ImpCode.hs
+++ b/src/Futhark/CodeGen/ImpCode.hs
@@ -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
diff --git a/src/Futhark/CodeGen/ImpGen.hs b/src/Futhark/CodeGen/ImpGen.hs
--- a/src/Futhark/CodeGen/ImpGen.hs
+++ b/src/Futhark/CodeGen/ImpGen.hs
@@ -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
diff --git a/src/Futhark/CodeGen/ImpGen/GPU/SegScan/SinglePass.hs b/src/Futhark/CodeGen/ImpGen/GPU/SegScan/SinglePass.hs
--- a/src/Futhark/CodeGen/ImpGen/GPU/SegScan/SinglePass.hs
+++ b/src/Futhark/CodeGen/ImpGen/GPU/SegScan/SinglePass.hs
@@ -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) $
diff --git a/src/Futhark/CodeGen/ImpGen/GPU/SegScan/TwoPass.hs b/src/Futhark/CodeGen/ImpGen/GPU/SegScan/TwoPass.hs
--- a/src/Futhark/CodeGen/ImpGen/GPU/SegScan/TwoPass.hs
+++ b/src/Futhark/CodeGen/ImpGen/GPU/SegScan/TwoPass.hs
@@ -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 []
diff --git a/src/Futhark/CodeGen/SetDefaultSpace.hs b/src/Futhark/CodeGen/SetDefaultSpace.hs
--- a/src/Futhark/CodeGen/SetDefaultSpace.hs
+++ b/src/Futhark/CodeGen/SetDefaultSpace.hs
@@ -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) =
diff --git a/src/Futhark/Compiler.hs b/src/Futhark/Compiler.hs
--- a/src/Futhark/Compiler.hs
+++ b/src/Futhark/Compiler.hs
@@ -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."
diff --git a/src/Futhark/IR/Parse.hs b/src/Futhark/IR/Parse.hs
--- a/src/Futhark/IR/Parse.hs
+++ b/src/Futhark/IR/Parse.hs
@@ -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
diff --git a/src/Futhark/IR/Pretty.hs b/src/Futhark/IR/Pretty.hs
--- a/src/Futhark/IR/Pretty.hs
+++ b/src/Futhark/IR/Pretty.hs
@@ -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) $
diff --git a/src/Futhark/IR/Syntax.hs b/src/Futhark/IR/Syntax.hs
--- a/src/Futhark/IR/Syntax.hs
+++ b/src/Futhark/IR/Syntax.hs
@@ -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
diff --git a/src/Futhark/Internalise/Exps.hs b/src/Futhark/Internalise/Exps.hs
--- a/src/Futhark/Internalise/Exps.hs
+++ b/src/Futhark/Internalise/Exps.hs
@@ -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)
         =<< ( (++)
diff --git a/src/Futhark/Optimise/MemoryBlockMerging.hs b/src/Futhark/Optimise/MemoryBlockMerging.hs
new file mode 100644
--- /dev/null
+++ b/src/Futhark/Optimise/MemoryBlockMerging.hs
@@ -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)
diff --git a/src/Futhark/Optimise/MemoryBlockMerging/GreedyColoring.hs b/src/Futhark/Optimise/MemoryBlockMerging/GreedyColoring.hs
new file mode 100644
--- /dev/null
+++ b/src/Futhark/Optimise/MemoryBlockMerging/GreedyColoring.hs
@@ -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
diff --git a/src/Futhark/Optimise/ReuseAllocations.hs b/src/Futhark/Optimise/ReuseAllocations.hs
deleted file mode 100644
--- a/src/Futhark/Optimise/ReuseAllocations.hs
+++ /dev/null
@@ -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)
diff --git a/src/Futhark/Optimise/ReuseAllocations/GreedyColoring.hs b/src/Futhark/Optimise/ReuseAllocations/GreedyColoring.hs
deleted file mode 100644
--- a/src/Futhark/Optimise/ReuseAllocations/GreedyColoring.hs
+++ /dev/null
@@ -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
diff --git a/src/Futhark/Passes.hs b/src/Futhark/Passes.hs
--- a/src/Futhark/Passes.hs
+++ b/src/Futhark/Passes.hs
@@ -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
diff --git a/src/Language/Futhark/Syntax.hs b/src/Language/Futhark/Syntax.hs
--- a/src/Language/Futhark/Syntax.hs
+++ b/src/Language/Futhark/Syntax.hs
@@ -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)
diff --git a/src/Language/Futhark/TypeChecker.hs b/src/Language/Futhark/TypeChecker.hs
--- a/src/Language/Futhark/TypeChecker.hs
+++ b/src/Language/Futhark/TypeChecker.hs
@@ -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)
diff --git a/src/Language/Futhark/TypeChecker/Terms.hs b/src/Language/Futhark/TypeChecker/Terms.hs
--- a/src/Language/Futhark/TypeChecker/Terms.hs
+++ b/src/Language/Futhark/TypeChecker/Terms.hs
@@ -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 =
diff --git a/unittests/Futhark/Optimise/MemoryBlockMerging/GreedyColoringTests.hs b/unittests/Futhark/Optimise/MemoryBlockMerging/GreedyColoringTests.hs
new file mode 100644
--- /dev/null
+++ b/unittests/Futhark/Optimise/MemoryBlockMerging/GreedyColoringTests.hs
@@ -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"
diff --git a/unittests/Futhark/Optimise/ReuseAllocations/GreedyColoringTests.hs b/unittests/Futhark/Optimise/ReuseAllocations/GreedyColoringTests.hs
deleted file mode 100644
--- a/unittests/Futhark/Optimise/ReuseAllocations/GreedyColoringTests.hs
+++ /dev/null
@@ -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"
diff --git a/unittests/futhark_tests.hs b/unittests/futhark_tests.hs
--- a/unittests/futhark_tests.hs
+++ b/unittests/futhark_tests.hs
@@ -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 ()