diff --git a/CHANGELOG.md b/CHANGELOG.md
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -5,6 +5,23 @@
 The format is based on [Keep a Changelog](http://keepachangelog.com/en/1.0.0/)
 and this project adheres to [Semantic Versioning](http://semver.org/spec/v2.0.0.html).
 
+## [0.25.17]
+
+* Faster device-to-device copies on CUDA.
+
+* "More correctly" detect L2 cache size for OpenCL backend on AMD GPUs.
+
+### Fixed
+
+* Handling of `..` in `import` paths (again).
+
+* Detection of impossible loop parameter sizes (#2144).
+
+* Rare case where GPU histograms would use slightly too much shared
+  memory and fail at run-time.
+
+* Rare crash in layout optimisation.
+
 ## [0.25.16]
 
 ### Added
diff --git a/docs/man/futhark-bench.rst b/docs/man/futhark-bench.rst
--- a/docs/man/futhark-bench.rst
+++ b/docs/man/futhark-bench.rst
@@ -143,7 +143,8 @@
 --skip-compilation
 
   Do not run the compiler, and instead assume that each benchmark
-  program has already been compiled.  Use with caution.
+  program has already been compiled into a server-mode executable. Use
+  with caution.
 
 --spec-file=FILE
 
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.25.16
+version:        0.25.17
 synopsis:       An optimising compiler for a functional, array-oriented language.
 
 description:    Futhark is a small programming language designed to be compiled to
@@ -527,6 +527,7 @@
       Futhark.ProfileTests
       Language.Futhark.CoreTests
       Language.Futhark.PrimitiveTests
+      Language.Futhark.SemanticTests
       Language.Futhark.SyntaxTests
       Language.Futhark.TypeChecker.TypesTests
       Language.Futhark.TypeCheckerTests
diff --git a/rts/c/backends/cuda.h b/rts/c/backends/cuda.h
--- a/rts/c/backends/cuda.h
+++ b/rts/c/backends/cuda.h
@@ -979,7 +979,7 @@
               (event_report_fn)cuda_event_report);
     CUDA_SUCCEED_FATAL(cuEventRecord(event->start, ctx->stream));
   }
-  CUDA_SUCCEED_OR_RETURN(cuMemcpy(dst+dst_offset, src+src_offset, nbytes));
+  CUDA_SUCCEED_OR_RETURN(cuMemcpyAsync(dst+dst_offset, src+src_offset, nbytes, ctx->stream));
   if (event != NULL) {
     CUDA_SUCCEED_FATAL(cuEventRecord(event->end, ctx->stream));
   }
@@ -1101,7 +1101,10 @@
     CUDA_SUCCEED_FATAL(cuCtxSynchronize());
     time_end = get_wall_time();
     long int time_diff = time_end - time_start;
-    fprintf(ctx->log, "  runtime: %ldus\n\n", time_diff);
+    fprintf(ctx->log, "  runtime: %ldus\n", time_diff);
+  }
+  if (ctx->logging) {
+    fprintf(ctx->log, "\n");
   }
 
   return FUTHARK_SUCCESS;
diff --git a/rts/c/backends/hip.h b/rts/c/backends/hip.h
--- a/rts/c/backends/hip.h
+++ b/rts/c/backends/hip.h
@@ -958,7 +958,10 @@
     HIP_SUCCEED_FATAL(hipStreamSynchronize(ctx->stream));
     time_end = get_wall_time();
     long int time_diff = time_end - time_start;
-    fprintf(ctx->log, "  runtime: %ldus\n\n", time_diff);
+    fprintf(ctx->log, "  runtime: %ldus\n", time_diff);
+  }
+  if (ctx->logging) {
+    fprintf(ctx->log, "\n");
   }
 
   return FUTHARK_SUCCESS;
diff --git a/rts/c/backends/opencl.h b/rts/c/backends/opencl.h
--- a/rts/c/backends/opencl.h
+++ b/rts/c/backends/opencl.h
@@ -796,6 +796,9 @@
   // Futhark reserves 4 bytes for bookkeeping information.
   max_shared_memory -= 4;
 
+  bool is_amd = strstr(device_option.platform_name, "AMD") != NULL;
+  bool is_nvidia = strstr(device_option.platform_name, "NVIDIA CUDA") != NULL;
+
   // The OpenCL implementation may reserve some local memory bytes for
   // various purposes.  In principle, we should use
   // clGetKernelWorkGroupInfo() to figure out for each kernel how much
@@ -804,9 +807,9 @@
   // arbitrarily subtract some bytes, based on empirical measurements
   // (but which might be arbitrarily wrong).  Fortunately, we rarely
   // try to really push the local memory usage.
-  if (strstr(device_option.platform_name, "NVIDIA CUDA") != NULL) {
+  if (is_nvidia) {
     max_shared_memory -= 12;
-  } else if (strstr(device_option.platform_name, "AMD") != NULL) {
+  } else if (is_amd) {
     max_shared_memory -= 16;
   }
 
@@ -829,17 +832,36 @@
     ctx->cfg->default_tile_size = max_tile_size;
   }
 
+  // Some of the code generated by Futhark will use the L2 cache size
+  // to make very precise decisions about execution. OpenCL does not
+  // specify whether CL_DEVICE_GLOBAL_MEM_CACHE_SIZE is L1 or L2 cache
+  // (or maybe something else entirely). NVIDIA's implementation
+  // reports L2, but AMDs reports L1 (and provides no way to query for
+  // the L2 size). That means it is time to hack.
 
-  cl_ulong cache_size;
+  cl_ulong l2_cache_size;
+  cl_ulong opencl_cache_size;
   OPENCL_SUCCEED_FATAL(clGetDeviceInfo(device_option.device, CL_DEVICE_GLOBAL_MEM_CACHE_SIZE,
-                                       sizeof(cache_size), &cache_size, NULL));
+                                       sizeof(opencl_cache_size), &opencl_cache_size, NULL));
 
-  if (cache_size == 0) {
+  if (is_amd) {
+    // We multiply the L1 cache size with the number of compute units
+    // times 4 (number of SIMD units with GCN). Empirically this
+    // doesn't get us the right result, but it gets us fairly close.
+    cl_ulong compute_units;
+    OPENCL_SUCCEED_FATAL(clGetDeviceInfo(device_option.device, CL_DEVICE_MAX_COMPUTE_UNITS,
+                                         sizeof(compute_units), &compute_units, NULL));
+    l2_cache_size = opencl_cache_size * compute_units * 4;
+  } else {
+    l2_cache_size = opencl_cache_size;
+  }
+
+  if (l2_cache_size == 0) {
     // Some code assumes nonzero cache.
-    cache_size = 1024*1024;
+    l2_cache_size = 1024*1024;
   }
 
-  ctx->max_cache = cache_size;
+  ctx->max_cache = l2_cache_size;
 
   ctx->max_registers = 1<<16; // I cannot find a way to query for this.
 
diff --git a/src/Futhark/Analysis/AccessPattern.hs b/src/Futhark/Analysis/AccessPattern.hs
--- a/src/Futhark/Analysis/AccessPattern.hs
+++ b/src/Futhark/Analysis/AccessPattern.hs
@@ -463,20 +463,20 @@
 analyseBasicOp ctx expression pats =
   -- Construct a VariableInfo from the subexpressions
   let ctx_val = case expression of
-        SubExp se -> varInfoFromSubExpr se
-        Opaque _ se -> varInfoFromSubExpr se
+        SubExp se -> varInfoFromSubExp se
+        Opaque _ se -> varInfoFromSubExp se
         ArrayLit ses _t -> concatVariableInfos mempty ses
-        UnOp _ se -> varInfoFromSubExpr se
+        UnOp _ se -> varInfoFromSubExp se
         BinOp _ lsubexp rsubexp -> concatVariableInfos mempty [lsubexp, rsubexp]
         CmpOp _ lsubexp rsubexp -> concatVariableInfos mempty [lsubexp, rsubexp]
-        ConvOp _ se -> varInfoFromSubExpr se
-        Assert se _ _ -> varInfoFromSubExpr se
+        ConvOp _ se -> varInfoFromSubExp se
+        Assert se _ _ -> varInfoFromSubExp se
         Index name _ ->
           error $ "unhandled: Index (This should NEVER happen) into " ++ prettyString name
         Update _ name _slice _subexp ->
           error $ "unhandled: Update (This should NEVER happen) onto " ++ prettyString name
         -- Technically, do we need this case?
-        Concat _ _ length_subexp -> varInfoFromSubExpr length_subexp
+        Concat _ _ length_subexp -> varInfoFromSubExp length_subexp
         Manifest _dim name -> varInfoFromNames ctx $ oneName name
         Iota end start stride _ -> concatVariableInfos mempty [end, start, stride]
         Replicate (Shape shape) value' -> concatVariableInfos mempty (value' : shape)
@@ -491,10 +491,10 @@
    in (ctx', mempty)
   where
     concatVariableInfos ne nn =
-      varInfoFromNames ctx (ne <> mconcat (map (analyseSubExpr pats ctx) nn))
+      varInfoFromNames ctx (ne <> mconcat (map (analyseSubExp pats ctx) nn))
 
-    varInfoFromSubExpr (Constant _) = (varInfoFromNames ctx mempty) {variableType = ConstType}
-    varInfoFromSubExpr (Var v) =
+    varInfoFromSubExp (Constant _) = (varInfoFromNames ctx mempty) {variableType = ConstType}
+    varInfoFromSubExp (Var v) =
       case M.lookup v (assignments ctx) of
         Just _ -> (varInfoFromNames ctx $ oneName v) {variableType = Variable}
         Nothing ->
@@ -578,7 +578,7 @@
   where
     subexprsToContext =
       contextFromNames ctx (varInfoZeroDeps ctx)
-        . concatMap (namesToList . analyseSubExpr pats ctx)
+        . concatMap (namesToList . analyseSubExp pats ctx)
 
 -- | Analyse statements in a rep body.
 analyseGPUBody :: (Analyse rep) => Body rep -> Context rep -> (Context rep, IndexTable rep)
@@ -590,23 +590,9 @@
 
 -- | Returns an intmap of names, to be used as dependencies in construction of
 -- VariableInfos.
--- Throws an error if SubExp contains a name not in context. This behaviour
--- might be thrown out in the future, as it is mostly just a very verbose way to
--- ensure that we capture all necessary variables in the context at the moment
--- of development.
-analyseSubExpr :: [VName] -> Context rep -> SubExp -> Names
-analyseSubExpr _ _ (Constant _) = mempty
-analyseSubExpr pp ctx (Var v) =
-  case M.lookup v (assignments ctx) of
-    (Just _) -> oneName v
-    Nothing ->
-      error $
-        "Failed to lookup variable \""
-          ++ prettyString v
-          ++ "\npat: "
-          ++ prettyString pp
-          ++ "\n\nContext\n"
-          ++ show ctx
+analyseSubExp :: [VName] -> Context rep -> SubExp -> Names
+analyseSubExp _ _ (Constant _) = mempty
+analyseSubExp _ _ (Var v) = oneName v
 
 -- | Reduce a DimFix into its set of dependencies
 consolidate :: Context rep -> SubExp -> DimAccess rep
@@ -617,7 +603,7 @@
 reduceDependencies :: Context rep -> VName -> M.Map VName Dependency
 reduceDependencies ctx v =
   case M.lookup v (assignments ctx) of
-    Nothing -> error $ "Unable to find " ++ prettyString v
+    Nothing -> mempty -- Means a global.
     Just (VariableInfo deps lvl _parents t) ->
       -- We detect whether it is a threadID or loop counter by checking
       -- whether or not it has any dependencies
diff --git a/src/Futhark/Analysis/Alias.hs b/src/Futhark/Analysis/Alias.hs
--- a/src/Futhark/Analysis/Alias.hs
+++ b/src/Futhark/Analysis/Alias.hs
@@ -22,7 +22,7 @@
   )
 where
 
-import Data.List (foldl')
+import Data.List qualified as L
 import Data.Map qualified as M
 import Futhark.IR.Aliases
 
@@ -64,7 +64,7 @@
   Stms rep ->
   (Stms (Aliases rep), AliasesAndConsumed)
 analyseStms orig_aliases =
-  withoutBound . foldl' f (mempty, (orig_aliases, mempty)) . stmsToList
+  withoutBound . L.foldl' f (mempty, (orig_aliases, mempty)) . stmsToList
   where
     withoutBound (stms, (aliases, consumed)) =
       let bound = foldMap (namesFromList . patNames . stmPat) stms
diff --git a/src/Futhark/Analysis/UsageTable.hs b/src/Futhark/Analysis/UsageTable.hs
--- a/src/Futhark/Analysis/UsageTable.hs
+++ b/src/Futhark/Analysis/UsageTable.hs
@@ -30,7 +30,7 @@
 import Data.Bits
 import Data.Foldable qualified as Foldable
 import Data.IntMap.Strict qualified as IM
-import Data.List (foldl')
+import Data.List qualified as L
 import Futhark.IR
 import Futhark.IR.Prop.Aliases
 import Prelude hiding (lookup)
@@ -64,10 +64,10 @@
 
 -- | Expand the usage table based on aliasing information.
 expand :: (VName -> Names) -> UsageTable -> UsageTable
-expand look (UsageTable m) = UsageTable $ foldl' grow m $ IM.toList m
+expand look (UsageTable m) = UsageTable $ L.foldl' grow m $ IM.toList m
   where
     grow m' (k, v) =
-      foldl'
+      L.foldl'
         (grow'' $ v `withoutU` presentU)
         m'
         (namesIntMap $ look $ VName (nameFromString "") k)
diff --git a/src/Futhark/CLI/Bench.hs b/src/Futhark/CLI/Bench.hs
--- a/src/Futhark/CLI/Bench.hs
+++ b/src/Futhark/CLI/Bench.hs
@@ -527,7 +527,7 @@
       []
       ["skip-compilation"]
       (NoArg $ Right $ \config -> config {optSkipCompilation = True})
-      "Use already compiled program.",
+      "Use already compiled server-mode program.",
     Option
       []
       ["exclude-case"]
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
@@ -37,6 +37,7 @@
 import Futhark.Optimise.CSE
 import Futhark.Optimise.DoubleBuffer
 import Futhark.Optimise.Fusion
+import Futhark.Optimise.GenRedOpt
 import Futhark.Optimise.HistAccs
 import Futhark.Optimise.InliningDeadFun
 import Futhark.Optimise.MemoryBlockMerging qualified as MemoryBlockMerging
@@ -646,6 +647,7 @@
     soacsPassOption applyADInnermost [],
     kernelsPassOption optimiseArrayLayoutGPU [],
     mcPassOption optimiseArrayLayoutMC [],
+    kernelsPassOption optimiseGenRed [],
     kernelsPassOption tileLoops [],
     kernelsPassOption histAccsGPU [],
     unstreamOption [],
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
@@ -21,7 +21,7 @@
 import Data.Char
 import Data.Functor (($>))
 import Data.Int (Int64)
-import Data.List (foldl', transpose)
+import Data.List qualified as L
 import Data.Map qualified as M
 import Data.Maybe
 import Data.Set qualified as S
@@ -537,7 +537,7 @@
     prog' = "'" <> T.pack prog <> "'"
 
 formatDataForGnuplot :: [Value] -> T.Text
-formatDataForGnuplot = T.unlines . map line . transpose . map valueElems
+formatDataForGnuplot = T.unlines . map line . L.transpose . map valueElems
   where
     line = T.unwords . map prettyText
 
@@ -1071,7 +1071,7 @@
   (failures, outputs, files) <-
     unzip3 <$> mapM (processBlock env) script
   cleanupImgDir env $ mconcat files
-  pure (foldl' min Success failures, T.intercalate "\n" outputs)
+  pure (L.foldl' min Success failures, T.intercalate "\n" outputs)
 
 -- | Common command line options that transform 'Options'.
 scriptCommandLineOptions :: [FunOptDescr Options]
diff --git a/src/Futhark/CLI/REPL.hs b/src/Futhark/CLI/REPL.hs
--- a/src/Futhark/CLI/REPL.hs
+++ b/src/Futhark/CLI/REPL.hs
@@ -101,7 +101,7 @@
         Left err ->
           error $ "Failed to initialise interpreter state: " <> T.unpack (docText err)
         Right s -> do
-          liftIO $ putDoc prog_err
+          liftIO $ putDocLn prog_err
           pure s {futharkiLoaded = maybe_prog}
     Right s ->
       pure s
diff --git a/src/Futhark/CodeGen/Backends/HIP.hs b/src/Futhark/CodeGen/Backends/HIP.hs
--- a/src/Futhark/CodeGen/Backends/HIP.hs
+++ b/src/Futhark/CodeGen/Backends/HIP.hs
@@ -79,7 +79,7 @@
            { optionLongName = "build-option",
              optionShortName = Nothing,
              optionArgument = RequiredArgument "OPT",
-             optionDescription = "Add an additional build option to the string passed to NVRTC.",
+             optionDescription = "Add an additional build option to the string passed to HIPRTC.",
              optionAction = [C.cstm|futhark_context_config_add_build_option(cfg, optarg);|]
            }
        ]
diff --git a/src/Futhark/CodeGen/ImpCode/GPU.hs b/src/Futhark/CodeGen/ImpCode/GPU.hs
--- a/src/Futhark/CodeGen/ImpCode/GPU.hs
+++ b/src/Futhark/CodeGen/ImpCode/GPU.hs
@@ -183,9 +183,10 @@
     ErrorSync Fence
   deriving (Show)
 
--- | Atomic operations return the value stored before the update.
--- This old value is stored in the first 'VName'.  The second 'VName'
--- is the memory block to update.  The 'Exp' is the new value.
+-- | Atomic operations return the value stored before the update. This
+-- old value is stored in the first 'VName' (except for
+-- 'AtomicWrite'). The second 'VName' is the memory block to update.
+-- The 'Exp' is the new value.
 data AtomicOp
   = AtomicAdd IntType VName VName (Count Elements (TExp Int64)) Exp
   | AtomicFAdd FloatType VName VName (Count Elements (TExp Int64)) Exp
@@ -198,6 +199,9 @@
   | AtomicXor IntType VName VName (Count Elements (TExp Int64)) Exp
   | AtomicCmpXchg PrimType VName VName (Count Elements (TExp Int64)) Exp Exp
   | AtomicXchg PrimType VName VName (Count Elements (TExp Int64)) Exp
+  | -- | Corresponds to a write followed by a memory fence. The old
+    -- value is not read.
+    AtomicWrite PrimType VName (Count Elements (TExp Int64)) Exp
   deriving (Show)
 
 instance FreeIn AtomicOp where
@@ -212,6 +216,7 @@
   freeIn' (AtomicXor _ _ arr i x) = freeIn' arr <> freeIn' i <> freeIn' x
   freeIn' (AtomicCmpXchg _ _ arr i x y) = freeIn' arr <> freeIn' i <> freeIn' x <> freeIn' y
   freeIn' (AtomicXchg _ _ arr i x) = freeIn' arr <> freeIn' i <> freeIn' x
+  freeIn' (AtomicWrite _ arr i x) = freeIn' arr <> freeIn' i <> freeIn' x
 
 instance Pretty KernelOp where
   pretty (GetBlockId dest i) =
@@ -311,6 +316,10 @@
     pretty old
       <+> "<-"
       <+> "atomic_xchg"
+      <> pretty t
+      <> parens (commasep [pretty arr <> brackets (pretty ind), pretty x])
+  pretty (Atomic _ (AtomicWrite t arr ind x)) =
+    "atomic_write"
       <> pretty t
       <> parens (commasep [pretty arr <> brackets (pretty ind), pretty x])
 
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
@@ -44,13 +44,14 @@
     -- * Lookups
     lookupVar,
     lookupArray,
+    lookupArraySpace,
     lookupMemory,
     lookupAcc,
     askAttrs,
 
     -- * Building Blocks
     TV,
-    mkTV,
+    MkTV (..),
     tvSize,
     tvExp,
     tvVar,
@@ -82,6 +83,8 @@
     dScope,
     dArray,
     dPrim,
+    dPrimS,
+    dPrimSV,
     dPrimVol,
     dPrim_,
     dPrimV_,
@@ -851,9 +854,9 @@
 traceArray s t shape se = do
   emit . Imp.TracePrint $ ErrorMsg [ErrorString (s <> ": ")]
   sLoopNest shape $ \is -> do
-    arr_elem <- dPrim "arr_elem" t
-    copyDWIMFix (tvVar arr_elem) [] se is
-    emit . Imp.TracePrint $ ErrorMsg [ErrorVal t (untyped (tvExp arr_elem)), " "]
+    arr_elem <- dPrimS "arr_elem" t
+    copyDWIMFix arr_elem [] se is
+    emit . Imp.TracePrint $ ErrorMsg [ErrorVal t (toExp' t arr_elem), " "]
   emit . Imp.TracePrint $ ErrorMsg ["\n"]
 
 defCompileBasicOp ::
@@ -1072,15 +1075,28 @@
   emit $ Imp.DeclareScalar name Imp.Nonvolatile t
   addVar name $ ScalarVar Nothing $ ScalarEntry t
 
--- | The return type is polymorphic, so there is no guarantee it
--- actually matches the 'PrimType', but at least we have to use it
--- consistently.
-dPrim :: String -> PrimType -> ImpM rep r op (TV t)
-dPrim name t = do
+-- | Create variable of some provided dynamic type. You'll need this
+-- when you are compiling program code of Haskell-level unknown type.
+-- For other things, use other functions.
+dPrimS :: String -> PrimType -> ImpM rep r op VName
+dPrimS name t = do
   name' <- newVName name
   dPrim_ name' t
-  pure $ TV name' t
+  pure name'
 
+-- | Create 'TV' of some provided dynamic type. No guarantee that the
+-- dynamic type matches the inferred type.
+dPrimSV :: String -> PrimType -> ImpM rep r op (TV t)
+dPrimSV name t = TV <$> dPrimS name t <*> pure t
+
+-- | Create 'TV' of some fixed type.
+dPrim :: (MkTV t) => String -> ImpM rep r op (TV t)
+dPrim name = do
+  name' <- newVName name
+  let tv = mkTV name'
+  dPrim_ name' $ tvType tv
+  pure tv
+
 dPrimV_ :: VName -> Imp.TExp t -> ImpM rep r op ()
 dPrimV_ name e = do
   dPrim_ name t
@@ -1090,15 +1106,21 @@
 
 dPrimV :: String -> Imp.TExp t -> ImpM rep r op (TV t)
 dPrimV name e = do
-  name' <- dPrim name $ primExpType $ untyped e
-  name' <-- e
-  pure name'
+  name' <- dPrimS name pt
+  let tv = TV name' pt
+  tv <-- e
+  pure tv
+  where
+    pt = primExpType $ untyped e
 
 dPrimVE :: String -> Imp.TExp t -> ImpM rep r op (Imp.TExp t)
 dPrimVE name e = do
-  name' <- dPrim name $ primExpType $ untyped e
-  name' <-- e
-  pure $ tvExp name'
+  name' <- dPrimS name pt
+  let tv = TV name' pt
+  tv <-- e
+  pure $ tvExp tv
+  where
+    pt = primExpType $ untyped e
 
 memBoundToVarEntry ::
   Maybe (Exp rep) ->
@@ -1180,13 +1202,47 @@
 -- It is still easy to cheat when you need to.
 data TV t = TV VName PrimType
 
--- | Create a typed variable from a name and a dynamic type.  Note
--- that there is no guarantee that the dynamic type corresponds to the
--- inferred static type, but the latter will at least have to be used
--- consistently.
-mkTV :: VName -> PrimType -> TV t
-mkTV = TV
+-- | A type class that helps ensuring that the type annotation in a
+-- 'TV' is correct.
+class MkTV t where
+  -- | Create a typed variable from a name and a dynamic type.
+  mkTV :: VName -> TV t
 
+  -- | Extract type from a 'TV'.
+  tvType :: TV t -> PrimType
+
+instance MkTV Bool where
+  mkTV v = TV v Bool
+  tvType _ = Bool
+
+instance MkTV Int8 where
+  mkTV v = TV v (IntType Int8)
+  tvType _ = IntType Int8
+
+instance MkTV Int16 where
+  mkTV v = TV v (IntType Int16)
+  tvType _ = IntType Int16
+
+instance MkTV Int32 where
+  mkTV v = TV v (IntType Int32)
+  tvType _ = IntType Int32
+
+instance MkTV Int64 where
+  mkTV v = TV v (IntType Int64)
+  tvType _ = IntType Int64
+
+instance MkTV Half where
+  mkTV v = TV v (FloatType Float16)
+  tvType _ = FloatType Float16
+
+instance MkTV Float where
+  mkTV v = TV v (FloatType Float32)
+  tvType _ = FloatType Float32
+
+instance MkTV Double where
+  mkTV v = TV v (FloatType Float64)
+  tvType _ = FloatType Float64
+
 -- | Convert a typed variable to a size (a SubExp).
 tvSize :: TV t -> Imp.DimSize
 tvSize = Var . tvVar
@@ -1220,6 +1276,10 @@
   toExp' _ (Constant v) = Imp.ValueExp v
   toExp' t (Var v) = Imp.var v t
 
+instance ToExp VName where
+  toExp = toExp . Var
+  toExp' t = toExp' t . Var
+
 instance ToExp (PrimExp VName) where
   toExp = pure
   toExp' _ = id
@@ -1309,6 +1369,7 @@
     MemVar _ entry -> pure entry
     _ -> error $ "Unknown memory block: " ++ prettyString name
 
+-- | In which memory space is this array allocated?
 lookupArraySpace :: VName -> ImpM rep r op Space
 lookupArraySpace =
   fmap entryMemSpace . lookupMemory
@@ -1450,7 +1511,7 @@
           fullyIndexArray' srclocation srcis
         vol <- asks envVolatility
         collect $ do
-          tmp <- tvVar <$> dPrim "tmp" bt
+          tmp <- dPrimS "tmp" bt
           emit $ Imp.Read tmp srcmem srcoffset bt srcspace vol
           emit $ Imp.Write targetmem targetoffset bt destspace vol $ Imp.var tmp bt
     | otherwise = do
diff --git a/src/Futhark/CodeGen/ImpGen/GPU.hs b/src/Futhark/CodeGen/ImpGen/GPU.hs
--- a/src/Futhark/CodeGen/ImpGen/GPU.hs
+++ b/src/Futhark/CodeGen/ImpGen/GPU.hs
@@ -13,7 +13,7 @@
 where
 
 import Control.Monad
-import Data.List (foldl')
+import Data.List qualified as L
 import Data.Map qualified as M
 import Data.Maybe
 import Futhark.CodeGen.ImpCode.GPU qualified as Imp
@@ -110,7 +110,7 @@
   sOp $ Imp.GetSizeMax (patElemName pe) size_class
 opCompiler (Pat [pe]) (Inner (SizeOp (CalcNumBlocks w64 max_num_tblocks_key tblock_size))) = do
   fname <- askFunction
-  max_num_tblocks :: TV Int32 <- dPrim "max_num_tblocks" int32
+  max_num_tblocks :: TV Int64 <- dPrim "max_num_tblocks"
   sOp $
     Imp.GetSize (tvVar max_num_tblocks) (keyWithEntryPoint fname max_num_tblocks_key) $
       sizeClassWithEntryPoint fname SizeGrid
@@ -125,7 +125,7 @@
           sExt64 (tvExp max_num_tblocks)
   -- We also don't want zero blocks.
   let num_tblocks = sMax64 1 num_tblocks_maybe_zero
-  mkTV (patElemName pe) int32 <-- sExt32 num_tblocks
+  mkTV (patElemName pe) <-- sExt32 num_tblocks
 opCompiler dest (Inner (SegOp op)) =
   segOpCompiler dest op
 opCompiler (Pat pes) (Inner (GPUBody _ (Body _ stms res))) = do
@@ -179,14 +179,14 @@
   if not $ all in_scope $ namesToList $ freeIn alloc_sizes
     then pure Nothing
     else do
-      shared_memory_capacity :: TV Int32 <- dPrim "shared_memory_capacity" int32
+      shared_memory_capacity :: TV Int64 <- dPrim "shared_memory_capacity"
       sOp $ Imp.GetSizeMax (tvVar shared_memory_capacity) SizeSharedMemory
 
       let shared_memory_capacity_64 =
             sExt64 $ tvExp shared_memory_capacity
           fits size =
             unCount size .<=. shared_memory_capacity_64
-      pure $ Just $ foldl' (.&&.) true (map fits alloc_sizes)
+      pure $ Just $ L.foldl' (.&&.) true (map fits alloc_sizes)
   where
     getGPU = foldMap getKernel
     getKernel (Imp.CallKernel k) | Imp.kernelCheckSharedMemory k = [k]
diff --git a/src/Futhark/CodeGen/ImpGen/GPU/Base.hs b/src/Futhark/CodeGen/ImpGen/GPU/Base.hs
--- a/src/Futhark/CodeGen/ImpGen/GPU/Base.hs
+++ b/src/Futhark/CodeGen/ImpGen/GPU/Base.hs
@@ -3,6 +3,9 @@
 
 module Futhark.CodeGen.ImpGen.GPU.Base
   ( KernelConstants (..),
+    kernelGlobalThreadId,
+    kernelLocalThreadId,
+    kernelBlockId,
     threadOperations,
     keyWithEntryPoint,
     CallKernelGen,
@@ -42,11 +45,12 @@
     Locking (..),
     AtomicUpdate (..),
     DoAtomicUpdate,
+    writeAtomic,
   )
 where
 
 import Control.Monad
-import Data.List (foldl')
+import Data.List qualified as L
 import Data.Map.Strict qualified as M
 import Data.Maybe
 import Futhark.CodeGen.ImpCode.GPU qualified as Imp
@@ -88,12 +92,9 @@
 type InKernelGen = ImpM GPUMem KernelEnv Imp.KernelOp
 
 data KernelConstants = KernelConstants
-  { kernelGlobalThreadId :: Imp.TExp Int32,
-    kernelLocalThreadId :: Imp.TExp Int32,
-    kernelBlockId :: Imp.TExp Int32,
-    kernelGlobalThreadIdVar :: VName,
-    kernelLocalThreadIdVar :: VName,
-    kernelBlockIdVar :: VName,
+  { kernelGlobalThreadIdVar :: TV Int32,
+    kernelLocalThreadIdVar :: TV Int32,
+    kernelBlockIdVar :: TV Int32,
     kernelNumBlocksCount :: Count NumBlocks SubExp,
     kernelBlockSizeCount :: Count BlockSize SubExp,
     kernelNumBlocks :: Imp.TExp Int64,
@@ -108,6 +109,11 @@
     kernelChunkItersMap :: M.Map [SubExp] (Imp.TExp Int32)
   }
 
+kernelGlobalThreadId, kernelLocalThreadId, kernelBlockId :: KernelConstants -> Imp.TExp Int32
+kernelGlobalThreadId = tvExp . kernelGlobalThreadIdVar
+kernelLocalThreadId = tvExp . kernelLocalThreadIdVar
+kernelBlockId = tvExp . kernelBlockIdVar
+
 keyWithEntryPoint :: Maybe Name -> Name -> Name
 keyWithEntryPoint fname key =
   nameFromString $ maybe "" ((++ ".") . nameToString) fname ++ nameToString key
@@ -251,7 +257,7 @@
 
 -- | If we are touching these arrays, which kind of fence do we need?
 fenceForArrays :: [VName] -> InKernelGen Imp.Fence
-fenceForArrays = fmap (foldl' max Imp.FenceLocal) . mapM need
+fenceForArrays = fmap (L.foldl' max Imp.FenceLocal) . mapM need
   where
     need arr =
       fmap (fenceForSpace . entryMemSpace)
@@ -267,13 +273,13 @@
 kernelConstToExp = traverse f
   where
     f (Imp.SizeMaxConst c) = do
-      v <- dPrim (prettyString c) int64
-      sOp $ Imp.GetSizeMax (tvVar v) c
-      pure $ tvVar v
+      v <- dPrimS (prettyString c) int64
+      sOp $ Imp.GetSizeMax v c
+      pure v
     f (Imp.SizeConst k c) = do
-      v <- dPrim (nameToString k) int64
-      sOp $ Imp.GetSize (tvVar v) k c
-      pure $ tvVar v
+      v <- dPrimS (nameToString k) int64
+      sOp $ Imp.GetSize v k c
+      pure v
 
 -- | Given available register and a list of parameter types, compute
 -- the largest available chunk size given the parameters for which we
@@ -310,7 +316,7 @@
   Lambda GPUMem ->
   InKernelGen ()
 inChunkScan constants seg_flag arrs_full_size lockstep_width block_size active arrs barrier scan_lam = everythingVolatile $ do
-  skip_threads <- dPrim "skip_threads" int32
+  skip_threads <- dPrim "skip_threads"
   let actual_params = lambdaParams scan_lam
       (x_params, y_params) =
         splitAt (length actual_params `div` 2) actual_params
@@ -562,7 +568,7 @@
   [VName] ->
   InKernelGen ()
 blockReduce w lam arrs = do
-  offset <- dPrim "offset" int32
+  offset <- dPrim "offset"
   blockReduceWithOffset offset w lam arrs
 
 blockReduceWithOffset ::
@@ -668,6 +674,28 @@
 compileThreadOp pat _ =
   compilerBugS $ "compileThreadOp: cannot compile rhs of binding " ++ prettyString pat
 
+-- | Perform a scalar write followed by a fence.
+writeAtomic ::
+  VName ->
+  [Imp.TExp Int64] ->
+  SubExp ->
+  [Imp.TExp Int64] ->
+  InKernelGen ()
+writeAtomic dst dst_is src src_is = do
+  t <- stripArray (length dst_is) <$> lookupType dst
+  sLoopSpace (map pe64 (arrayDims t)) $ \is -> do
+    let pt = elemType t
+    (dst_mem, dst_space, dst_offset) <- fullyIndexArray dst (dst_is ++ is)
+    case src_is ++ is of
+      [] ->
+        sOp . Imp.Atomic dst_space $
+          Imp.AtomicWrite pt dst_mem dst_offset (toExp' pt src)
+      _ -> do
+        tmp <- dPrimSV "tmp" pt
+        copyDWIMFix (tvVar tmp) [] src (src_is ++ is)
+        sOp . Imp.Atomic dst_space $
+          Imp.AtomicWrite pt dst_mem dst_offset (untyped (tvExp tmp))
+
 -- | Locking strategy used for an atomic update.
 data Locking = Locking
   { -- | Array containing the lock.
@@ -721,14 +749,14 @@
         -- can be implemented by atomic compare-and-swap if 32/64 bits.
         forM_ (zip arrs ops_and_ts) $ \(a, (op, t, x, y)) -> do
           -- Common variables.
-          old <- dPrim "old" t
+          old <- dPrimS "old" t
 
           (arr', _a_space, bucket_offset) <- fullyIndexArray a bucket
 
-          case opHasAtomicSupport space (tvVar old) arr' bucket_offset op of
+          case opHasAtomicSupport space old arr' bucket_offset op of
             Just f -> sOp $ f $ Imp.var y t
             Nothing ->
-              atomicUpdateCAS space t a (tvVar old) bucket x $
+              atomicUpdateCAS space t a old bucket x $
                 x <~~ Imp.BinOpExp op (Imp.var x t) (Imp.var y t)
   where
     opHasAtomicSupport space old arr' bucket' bop = do
@@ -748,12 +776,11 @@
   | [Prim t] <- lambdaReturnType op,
     [xp, _] <- lambdaParams op,
     primBitSize t `elem` [32, 64] = AtomicCAS $ \space [arr] bucket -> do
-      old <- dPrim "old" t
-      atomicUpdateCAS space t arr (tvVar old) bucket (paramName xp) $
-        compileBody' [xp] $
-          lambdaBody op
+      old <- dPrimS "old" t
+      atomicUpdateCAS space t arr old bucket (paramName xp) $
+        compileBody' [xp] (lambdaBody op)
 atomicUpdateLocking _ op = AtomicLocking $ \locking space arrs bucket -> do
-  old <- dPrim "old" int32
+  old <- dPrim "old"
   continue <- dPrimVol "continue" Bool true
 
   -- Correctly index into locks.
@@ -814,8 +841,6 @@
             zipWithM_ (writeArray bucket) arrs $
               map (Var . paramName) acc_params
 
-      fence = sOp $ Imp.MemFence $ fenceForSpace space
-
   -- While-loop: Try to insert your value
   sWhile (tvExp continue) $ do
     try_acquire_lock
@@ -824,12 +849,10 @@
       bind_acc_params
       op_body
       do_hist
-      fence
       release_lock
       break_loop
-    fence
   where
-    writeArray bucket arr val = copyDWIMFix arr bucket val []
+    writeArray bucket arr val = writeAtomic arr bucket val []
 
 atomicUpdateCAS ::
   Space ->
@@ -849,7 +872,7 @@
   --   x = do_op(assumed, y);
   --   old = atomicCAS(&d_his[idx], assumed, tmp);
   -- } while(assumed != old);
-  assumed <- tvVar <$> dPrim "assumed" t
+  assumed <- dPrimS "assumed" t
   run_loop <- dPrimV "run_loop" true
 
   -- XXX: CUDA may generate really bad code if this is not a volatile
@@ -961,14 +984,11 @@
       tblock_size' = Imp.pe64 (unCount tblock_size)
       constants =
         KernelConstants
-          { kernelGlobalThreadId = Imp.le32 global_tid,
-            kernelLocalThreadId = Imp.le32 local_tid,
-            kernelBlockId = Imp.le32 tblock_id,
-            kernelGlobalThreadIdVar = global_tid,
-            kernelLocalThreadIdVar = local_tid,
+          { kernelGlobalThreadIdVar = mkTV global_tid,
+            kernelLocalThreadIdVar = mkTV local_tid,
+            kernelBlockIdVar = mkTV tblock_id,
             kernelNumBlocksCount = num_tblocks,
             kernelBlockSizeCount = tblock_size,
-            kernelBlockIdVar = tblock_id,
             kernelNumBlocks = num_tblocks',
             kernelBlockSize = tblock_size',
             kernelNumThreads = sExt32 (tblock_size' * num_tblocks'),
@@ -1020,7 +1040,7 @@
   Imp.TExp Int64 ->
   CallKernelGen (Imp.TExp Int32, Count NumBlocks SubExp, Count BlockSize SubExp)
 simpleKernelBlocks max_num_tblocks kernel_size = do
-  tblock_size <- dPrim "tblock_size" int64
+  tblock_size <- dPrim "tblock_size"
   fname <- askFunction
   let tblock_size_key = keyWithEntryPoint fname $ nameFromString $ prettyString $ tvVar tblock_size
   sOp $ Imp.GetSize (tvVar tblock_size) tblock_size_key Imp.SizeThreadBlock
@@ -1053,12 +1073,9 @@
 
       constants =
         KernelConstants
-          { kernelGlobalThreadId = Imp.le32 thread_gtid,
-            kernelLocalThreadId = Imp.le32 thread_ltid,
-            kernelBlockId = Imp.le32 tblock_id,
-            kernelGlobalThreadIdVar = thread_gtid,
-            kernelLocalThreadIdVar = thread_ltid,
-            kernelBlockIdVar = tblock_id,
+          { kernelGlobalThreadIdVar = mkTV thread_gtid,
+            kernelLocalThreadIdVar = mkTV thread_ltid,
+            kernelBlockIdVar = mkTV tblock_id,
             kernelNumBlocksCount = num_tblocks,
             kernelBlockSizeCount = tblock_size,
             kernelNumBlocks = num_tblocks',
@@ -1071,7 +1088,7 @@
 
       wrapKernel m = do
         dPrim_ thread_ltid int32
-        dPrim_ inner_tblock_size int64
+        dPrim_ inner_tblock_size int32
         dPrim_ tblock_id int32
         sOp (Imp.GetLocalId thread_ltid 0)
         sOp (Imp.GetLocalSize inner_tblock_size 0)
@@ -1100,7 +1117,7 @@
   InKernelGen ()
 virtualiseBlocks SegVirt required_blocks m = do
   constants <- kernelConstants <$> askEnv
-  phys_tblock_id <- dPrim "phys_tblock_id" int32
+  phys_tblock_id <- dPrim "phys_tblock_id"
   sOp $ Imp.GetBlockId (tvVar phys_tblock_id) 0
   iterations <-
     dPrimVE "iterations" $
@@ -1114,9 +1131,8 @@
     -- Make sure the virtual block is actually done before we let
     -- another virtual block have its way with it.
     sOp $ Imp.ErrorSync Imp.FenceGlobal
-virtualiseBlocks _ _ m = do
-  gid <- kernelBlockIdVar . kernelConstants <$> askEnv
-  m $ Imp.le32 gid
+virtualiseBlocks _ _ m =
+  m . tvExp . kernelBlockIdVar . kernelConstants =<< askEnv
 
 -- | Various extra configuration of the kernel being generated.
 data KernelAttrs = KernelAttrs
@@ -1151,7 +1167,7 @@
 
 getSize :: String -> SizeClass -> CallKernelGen (TV Int64)
 getSize desc size_class = do
-  v <- dPrim desc int64
+  v <- dPrim desc
   fname <- askFunction
   let v_key = keyWithEntryPoint fname $ nameFromString $ prettyString $ tvVar v
   sOp $ Imp.GetSize (tvVar v) v_key size_class
@@ -1178,7 +1194,7 @@
 
 sKernel ::
   Operations GPUMem KernelEnv Imp.KernelOp ->
-  (KernelConstants -> Imp.TExp Int32) ->
+  (KernelConstants -> Imp.TExp Int64) ->
   String ->
   VName ->
   KernelAttrs ->
@@ -1199,7 +1215,7 @@
   KernelAttrs ->
   InKernelGen () ->
   CallKernelGen ()
-sKernelThread = sKernel threadOperations kernelGlobalThreadId
+sKernelThread = sKernel threadOperations $ sExt64 . kernelGlobalThreadId
 
 sKernelOp ::
   KernelAttrs ->
@@ -1278,7 +1294,7 @@
   let name =
         keyWithEntryPoint fname $
           nameFromString $
-            "replicate_" ++ show (baseTag $ kernelGlobalThreadIdVar constants)
+            "replicate_" ++ show (baseTag $ tvVar $ kernelGlobalThreadIdVar constants)
 
   sKernelFailureTolerant True threadOperations constants name $
     virtualise $ \gtid -> do
@@ -1362,7 +1378,7 @@
             "iota_"
               ++ prettyString et
               ++ "_"
-              ++ show (baseTag $ kernelGlobalThreadIdVar constants)
+              ++ show (baseTag $ tvVar $ kernelGlobalThreadIdVar constants)
 
   sKernelFailureTolerant True threadOperations constants name $
     virtualise $ \gtid ->
diff --git a/src/Futhark/CodeGen/ImpGen/GPU/Block.hs b/src/Futhark/CodeGen/ImpGen/GPU/Block.hs
--- a/src/Futhark/CodeGen/ImpGen/GPU/Block.hs
+++ b/src/Futhark/CodeGen/ImpGen/GPU/Block.hs
@@ -196,7 +196,7 @@
 compileFlatId :: SegSpace -> InKernelGen ()
 compileFlatId space = do
   ltid <- kernelLocalThreadId . kernelConstants <$> askEnv
-  dPrimV_ (segFlat space) ltid
+  dPrimV_ (segFlat space) $ sExt64 ltid
 
 -- Construct the necessary lock arrays for an intra-block histogram.
 prepareIntraBlockSegHist ::
@@ -602,7 +602,7 @@
   KernelAttrs ->
   InKernelGen () ->
   CallKernelGen ()
-sKernelBlock = sKernel blockOperations kernelBlockId
+sKernelBlock = sKernel blockOperations $ sExt64 . kernelBlockId
 
 compileBlockResult ::
   SegSpace ->
diff --git a/src/Futhark/CodeGen/ImpGen/GPU/SegHist.hs b/src/Futhark/CodeGen/ImpGen/GPU/SegHist.hs
--- a/src/Futhark/CodeGen/ImpGen/GPU/SegHist.hs
+++ b/src/Futhark/CodeGen/ImpGen/GPU/SegHist.hs
@@ -38,7 +38,7 @@
 module Futhark.CodeGen.ImpGen.GPU.SegHist (compileSegHist) where
 
 import Control.Monad
-import Data.List (foldl', genericLength, zip5)
+import Data.List qualified as L
 import Data.Map qualified as M
 import Data.Maybe
 import Futhark.CodeGen.ImpCode.GPU qualified as Imp
@@ -97,7 +97,7 @@
 
   -- Create names for the intermediate array memory blocks,
   -- memory block sizes, arrays, and number of subhistograms.
-  num_subhistos <- dPrim "num_subhistos" int32
+  num_subhistos <- dPrim "num_subhistos"
   subhisto_infos <- forM (zip (histDest op) (histNeutral op)) $ \(dest, ne) -> do
     dest_t <- lookupType dest
     dest_mem <- entryArrayLoc <$> lookupArray dest
@@ -226,7 +226,7 @@
   hist_RF <-
     dPrimVE "hist_RF" $
       sum (map (r64 . pe64 . histRaceFactor . slugOp) slugs)
-        / genericLength slugs
+        / L.genericLength slugs
 
   hist_el_size <- dPrimVE "hist_el_size" $ sum $ map slugElAvgSize slugs
 
@@ -242,7 +242,7 @@
           t64 $
             r64 hist_T / hist_C_max
 
-  hist_L2 <- dPrim "L2_size" int32
+  hist_L2 :: TV Int32 <- dPrim "L2_size"
   -- Equivalent to F_L2*L2 in paper.
   sOp $ Imp.GetSizeMax (tvVar hist_L2) Imp.SizeCache
 
@@ -253,7 +253,7 @@
         (hist_k_RF * hist_RF)
           / (hist_L2_ln_sz / r64 hist_el_size)
 
-  hist_S <- dPrim "hist_S" int32
+  hist_S <- dPrim "hist_S"
 
   -- For sparse histograms (H exceeds N) we only want a single chunk.
   sIf
@@ -291,26 +291,21 @@
     slugElAvgSize slug@(SegHistSlug op _ _ do_op) =
       case do_op of
         AtomicLocking {} ->
-          slugElSize slug `quot` (1 + genericLength (lambdaReturnType (histOp op)))
+          slugElSize slug `quot` (1 + L.genericLength (lambdaReturnType (histOp op)))
         _ ->
-          slugElSize slug `quot` genericLength (lambdaReturnType (histOp op))
+          slugElSize slug `quot` L.genericLength (lambdaReturnType (histOp op))
 
     -- "Average element size" as computed by a formula that also takes
     -- locking into account.
     slugElSize (SegHistSlug op _ _ do_op) =
-      case do_op of
-        AtomicLocking {} ->
-          sExt32 $
-            unCount $
-              sum $
-                map (typeSize . (`arrayOfShape` histOpShape op)) $
-                  Prim int32 : lambdaReturnType (histOp op)
-        _ ->
-          sExt32 $
-            unCount $
-              sum $
-                map (typeSize . (`arrayOfShape` histOpShape op)) $
-                  lambdaReturnType (histOp op)
+      sExt32 . unCount . sum $
+        case do_op of
+          AtomicLocking {} ->
+            map (typeSize . (`arrayOfShape` histOpShape op)) $
+              Prim int32 : lambdaReturnType (histOp op)
+          _ ->
+            map (typeSize . (`arrayOfShape` histOpShape op)) $
+              lambdaReturnType (histOp op)
 
     onOp hist_L2 hist_M_min hist_S hist_RACE_exp l slug = do
       let SegHistSlug op num_subhistos subhisto_info do_op = slug
@@ -403,7 +398,7 @@
     -- Compute subhistogram index for each thread, per histogram.
     subhisto_inds <- forM slugs $ \slug ->
       dPrimVE "subhisto_ind" $
-        kernelGlobalThreadId constants
+        sExt32 (kernelGlobalThreadId constants)
           `quot` ( kernelNumThreads constants
                      `divUp` sExt32 (tvExp (slugNumSubhistos slug))
                  )
@@ -440,7 +435,7 @@
                   map kernelResultSubExp red_res
 
           sComment "perform atomic updates" $
-            forM_ (zip5 (map slugOp slugs) histograms red_res_split subhisto_inds hist_H_chks) $
+            forM_ (L.zip5 (map slugOp slugs) histograms red_res_split subhisto_inds hist_H_chks) $
               \( HistOp dest_shape _ _ _ shape lam,
                  do_op,
                  (bucket, vs'),
@@ -554,16 +549,17 @@
       -- Initialise local-memory sub-histograms.  These are
       -- represented as two-dimensional arrays.
       let init_local_subhistos hist_H_chk = do
-            local_subhistos <-
-              forM (histType op) $ \t -> do
-                let sub_local_shape =
-                      Shape [tvSize num_subhistos_per_block]
-                        <> setOuterDims (arrayShape t) (histRank op) (Shape [hist_H_chk])
-                sAllocArray
-                  "subhistogram_local"
-                  (elemType t)
-                  sub_local_shape
-                  (Space "shared")
+            local_subhistos <- forM (histType op) $ \t -> do
+              let subhisto_shape =
+                    setOuterDims
+                      (arrayShape t)
+                      (histRank op)
+                      (Shape [hist_H_chk])
+              sAllocArray
+                "subhistogram_local"
+                (elemType t)
+                (Shape [tvSize num_subhistos_per_block] <> subhisto_shape)
+                (Space "shared")
 
             do_op' <- mk_op hist_H_chk
 
@@ -608,22 +604,17 @@
         num_subhistos_per_block = tvExp num_subhistos_per_block_var
         segment_size' = pe64 segment_size
 
-    num_segments <-
-      dPrimVE "num_segments" $
-        product $
-          map pe64 segment_dims
+    num_segments <- dPrimVE "num_segments" $ product $ map pe64 segment_dims
 
     hist_H_chks <- forM (map slugOp slugs) $ \op ->
       dPrimV "hist_H_chk" $ histSize op `divUp` sExt64 hist_S
 
     histo_sizes <- forM (zip slugs hist_H_chks) $ \(slug, hist_H_chk) -> do
       let histo_dims =
-            tvExp hist_H_chk
-              : map pe64 (shapeDims (histOpShape (slugOp slug)))
+            tvExp hist_H_chk : map pe64 (shapeDims (histOpShape (slugOp slug)))
       histo_size <-
         dPrimVE "histo_size" $ product histo_dims
-      let block_hists_size =
-            sExt64 num_subhistos_per_block * histo_size
+      let block_hists_size = sExt64 num_subhistos_per_block * histo_size
       init_per_thread <-
         dPrimVE "init_per_thread" $ sExt32 $ block_hists_size `divUp` pe64 (unCount tblock_size)
       pure (histo_dims, histo_size, init_per_thread)
@@ -896,10 +887,10 @@
       segment_dims = init space_sizes
       segmented = not $ null segment_dims
 
-  hist_L <- dPrim "hist_L" int32
+  hist_L :: TV Int64 <- dPrim "hist_L"
   sOp $ Imp.GetSizeMax (tvVar hist_L) Imp.SizeSharedMemory
 
-  max_tblock_size <- dPrim "max_tblock_size" int32
+  max_tblock_size :: TV Int64 <- dPrim "max_tblock_size"
   sOp $ Imp.GetSizeMax (tvVar max_tblock_size) Imp.SizeThreadBlock
 
   -- XXX: we need to record for later use that max_tblock_size is the
@@ -969,7 +960,7 @@
         dPrimVE "work_asymp_M_max" $
           (hist_Nout * hist_N)
             `quot` ( (q_small * unCount num_tblocks' * hist_H)
-                       `quot` genericLength slugs
+                       `quot` L.genericLength slugs
                    )
 
   -- Number of subhistograms per result histogram.
@@ -1003,10 +994,13 @@
   local_mem_needed <-
     dPrimVE "local_mem_needed" $
       hist_el_size * sExt64 (tvExp hist_M)
+  -- We add one to the memory requirement because if the chunk
+  -- otherwise *exactly* fits, it might actually *not* fit in the case
+  -- of a multi-value operator, as we individually round up the sizes
+  -- of the component arrays. (Very rare edge case.)
   hist_S <-
-    dPrimVE "hist_S" $
-      sExt32 $
-        (hist_H * local_mem_needed) `divUp` tvExp hist_L
+    dPrimVE "hist_S" . sExt32 $
+      (hist_H * local_mem_needed + 1) `divUp` tvExp hist_L
   let max_S = case bodyPassage kbody of
         MustBeSinglePass -> 1
         MayBeMultiPass -> fromIntegral $ maxinum $ map slugMaxLocalMemPasses slugs
@@ -1102,7 +1096,7 @@
           _ -> Nothing
     hist_el_size <-
       dPrimVE "hist_el_size" $
-        foldl' (+) (h `divUp` hist_H) $
+        L.foldl' (+) (h `divUp` hist_H) $
           mapMaybe lockSize slugs
 
     -- Input elements contributing to each histogram.
@@ -1113,7 +1107,7 @@
       dPrimVE "hist_RF" $
         sExt32 $
           sum (map (pe64 . histRaceFactor . slugOp) slugs)
-            `quot` genericLength slugs
+            `quot` L.genericLength slugs
 
     let hist_T = sExt32 $ unCount num_tblocks' * unCount tblock_size'
     emit $ Imp.DebugPrint "\n# SegHist" Nothing
diff --git a/src/Futhark/CodeGen/ImpGen/GPU/SegRed.hs b/src/Futhark/CodeGen/ImpGen/GPU/SegRed.hs
--- a/src/Futhark/CodeGen/ImpGen/GPU/SegRed.hs
+++ b/src/Futhark/CodeGen/ImpGen/GPU/SegRed.hs
@@ -638,8 +638,8 @@
     mkAcc p block_res_arr
       | Prim t <- paramType p,
         shapeRank (segBinOpShape op) == 0 = do
-          block_res_acc <- dPrim (baseString (paramName p) <> "_block_res_acc") t
-          pure (tvVar block_res_acc, [])
+          block_res_acc <- dPrimS (baseString (paramName p) <> "_block_res_acc") t
+          pure (block_res_acc, [])
       -- if this is a non-primitive reduction, the global mem result array will
       -- double as accumulator.
       | otherwise =
@@ -687,7 +687,7 @@
   InKernelGen [Lambda GPUMem]
 reductionStageOne gtids n global_tid q chunk threads_per_segment slugs body_cont = do
   constants <- kernelConstants <$> askEnv
-  let glb_ind_var = mkTV (last gtids) int64
+  let glb_ind_var = mkTV (last gtids)
       ltid = sExt64 $ kernelLocalThreadId constants
 
   dScope Nothing $ scopeOfLParams $ concatMap slugParams slugs
@@ -920,7 +920,7 @@
       red_arrs = slugBlockRedArrs slug
       block_res_arrs = blockResArrs slug
 
-  old_counter <- dPrim "old_counter" int32
+  old_counter <- dPrim "old_counter"
   (counter_mem, _, counter_offset) <-
     fullyIndexArray
       counters
@@ -928,8 +928,7 @@
   sComment "first thread in block saves block result to global memory" $
     sWhen (ltid32 .==. 0) $ do
       forM_ (take (length nes) $ zip block_res_arrs (slugAccs slug)) $ \(v, (acc, acc_is)) ->
-        copyDWIMFix v [0, sExt64 tblock_id] (Var acc) acc_is
-      sOp $ Imp.MemFence Imp.FenceGlobal
+        writeAtomic v [0, sExt64 tblock_id] (Var acc) acc_is
       -- Increment the counter, thus stating that our result is
       -- available.
       sOp
@@ -942,11 +941,11 @@
         $ untyped (1 :: Imp.TExp Int32)
       -- Now check if we were the last block to write our result.  If
       -- so, it is our responsibility to produce the final result.
-      sWrite sync_arr [0] $ untyped $ tvExp old_counter .==. blocks_per_segment - 1
+      sWrite sync_arr [0] $ untyped $ tvExp old_counter .==. sExt32 (blocks_per_segment - 1)
 
   sOp $ Imp.Barrier Imp.FenceGlobal
 
-  is_last_block <- dPrim "is_last_block" Bool
+  is_last_block <- dPrim "is_last_block"
   copyDWIMFix (tvVar is_last_block) [] (Var sync_arr) [0]
   sWhen (tvExp is_last_block) $ do
     -- The final block has written its result (and it was
@@ -960,7 +959,7 @@
         Imp.Atomic DefaultSpace $
           Imp.AtomicAdd Int32 (tvVar old_counter) counter_mem counter_offset $
             untyped $
-              negate blocks_per_segment
+              sExt32 (negate blocks_per_segment)
 
     sLoopNest (slugShape slug) $ \vec_is -> do
       unless (null $ slugShape slug) $
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
@@ -116,8 +116,8 @@
   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
+  flg_x :: TV Int8 <- dPrim "flg_x"
+  flg_y :: TV Int8 <- dPrim "flg_y"
   let flg_param_x = Param mempty (tvVar flg_x) (MemPrim p_int8)
       flg_param_y = Param mempty (tvVar flg_y) (MemPrim p_int8)
       flg_y_exp = tvExp flg_y
@@ -126,7 +126,7 @@
 
   dLParams (lambdaParams scan_lam)
 
-  skip_threads <- dPrim "skip_threads" int32
+  skip_threads <- dPrim "skip_threads"
   let in_block_thread_active =
         tvExp skip_threads .<=. in_block_id
       actual_params = lambdaParams scan_lam
@@ -288,7 +288,7 @@
     -- We could use virtualiseBlocks, but this introduces a barrier which is
     -- redundant in this case, and also we don't need to base virtual block IDs
     -- on the loop variable, but rather on the dynamic IDs.
-    phys_block_id <- dPrim "phys_block_id" int32
+    phys_block_id <- dPrim "phys_block_id"
     sOp $ Imp.GetBlockId (tvVar phys_block_id) 0
     iters <-
       dPrimVE "virtloop_bound" $
@@ -296,7 +296,7 @@
           `divUp` num_phys_blocks_e
 
     sFor "virtloop_i" iters $ const $ do
-      dyn_id <- dPrim "dynamic_id" int32
+      dyn_id <- dPrim "dynamic_id"
       sComment "First thread in block fetches this block's dynamic_id" $ do
         sWhen (ltid32 .==. 0) $ do
           (globalIdMem, _, globalIdOff) <- fullyIndexArray global_id [0]
@@ -421,7 +421,7 @@
 
       scan_op1 <- renameLambda $ segBinOpLambda scan_op
 
-      accs <- mapM (dPrim "acc") tys
+      accs <- mapM (dPrimSV "acc") tys
       sComment "Scan results (with warp scan)" $ do
         blockScan
           crossesSegment
@@ -484,7 +484,7 @@
           -- sWhen
           sOp local_fence
 
-          status <- dPrim "status" int8 :: InKernelGen (TV Int8)
+          status :: TV Int8 <- dPrim "status"
           copyDWIMFix (tvVar status) [] (Var warpscan) [0]
 
           sIf
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
@@ -6,7 +6,7 @@
 
 import Control.Monad
 import Control.Monad.State
-import Data.List (delete, find, foldl', zip4)
+import Data.List qualified as L
 import Data.Maybe
 import Futhark.CodeGen.ImpCode.GPU qualified as Imp
 import Futhark.CodeGen.ImpGen
@@ -27,7 +27,7 @@
   InKernelGen [[VName]]
 makeLocalArrays (Count tblock_size) num_threads scans = do
   (arrs, mems_and_sizes) <- runStateT (mapM onScan scans) mempty
-  let maxSize sizes = Imp.bytes $ foldl' sMax64 1 $ map Imp.unCount sizes
+  let maxSize sizes = Imp.bytes $ L.foldl' sMax64 1 $ map Imp.unCount sizes
   forM_ mems_and_sizes $ \(sizes, mem) ->
     sAlloc_ mem (maxSize sizes) (Space "shared")
   pure arrs
@@ -56,9 +56,9 @@
     getMem pt shape = do
       let size = typeSize $ Array pt shape NoUniqueness
       mems <- get
-      case (find ((size `elem`) . fst) mems, mems) of
+      case (L.find ((size `elem`) . fst) mems, mems) of
         (Just mem, _) -> do
-          modify $ delete mem
+          modify $ L.delete mem
           pure mem
         (Nothing, (size', mem) : mems') -> do
           put mems'
@@ -353,7 +353,7 @@
     -- Construct segment indices.
     zipWithM_ dPrimV_ gtids $ unflattenIndex dims' $ tvExp flat_idx
 
-    forM_ (zip4 scans per_scan_local_arrs per_scan_rets per_scan_pes) $
+    forM_ (L.zip4 scans per_scan_local_arrs per_scan_rets per_scan_pes) $
       \(SegBinOp _ scan_op nes vec_shape, local_arrs, rets, pes) ->
         sLoopNest vec_shape $ \vec_is -> do
           let glob_is = map Imp.le64 gtids ++ vec_is
@@ -494,7 +494,7 @@
   -- Since stage 2 involves a group size equal to the number of groups
   -- used for stage 1, we have to cap this number to the maximum group
   -- size.
-  stage1_max_num_tblocks <- dPrim "stage1_max_num_tblocks" int64
+  stage1_max_num_tblocks <- dPrim "stage1_max_num_tblocks"
   sOp $ Imp.GetSizeMax (tvVar stage1_max_num_tblocks) SizeThreadBlock
 
   stage1_num_tblocks <-
diff --git a/src/Futhark/CodeGen/ImpGen/GPU/ToOpenCL.hs b/src/Futhark/CodeGen/ImpGen/GPU/ToOpenCL.hs
--- a/src/Futhark/CodeGen/ImpGen/GPU/ToOpenCL.hs
+++ b/src/Futhark/CodeGen/ImpGen/GPU/ToOpenCL.hs
@@ -734,6 +734,17 @@
       doAtomicCmpXchg s t old arr ind cmp val [C.cty|int|]
     atomicOps s (AtomicXchg t old arr ind val) =
       doAtomicXchg s t old arr ind val [C.cty|int|]
+    atomicOps s (AtomicWrite t arr ind val) = do
+      ind' <- GC.compileExp $ untyped $ unCount ind
+      val' <- toStorage t <$> GC.compileExp val
+      let quals = case s of
+            Space sid -> pointerQuals sid
+            _ -> pointerQuals "global"
+      GC.stm [C.cstm|(($tyquals:quals $ty:(primStorageType t)*)$id:arr)[$exp:ind'] = $exp:val';|]
+      GC.stm $
+        case s of
+          Space "shared" -> [C.cstm|mem_fence_local();|]
+          _ -> [C.cstm|mem_fence_global();|]
 
     cannotAllocate :: GC.Allocate KernelOp KernelState
     cannotAllocate _ =
diff --git a/src/Futhark/CodeGen/ImpGen/Multicore.hs b/src/Futhark/CodeGen/ImpGen/Multicore.hs
--- a/src/Futhark/CodeGen/ImpGen/Multicore.hs
+++ b/src/Futhark/CodeGen/ImpGen/Multicore.hs
@@ -44,7 +44,7 @@
       is <- dIndexSpace' "i" (map pe64 srcshape) i
       (_, destspace, destidx) <- fullyIndexArray' destloc is
       (_, srcspace, srcidx) <- fullyIndexArray' srcloc is
-      tmp <- tvVar <$> dPrim "tmp" pt
+      tmp <- dPrimS "tmp" pt
       emit $ Imp.Read tmp srcmem srcidx pt srcspace Imp.Nonvolatile
       emit $ Imp.Write destmem destidx pt destspace Imp.Nonvolatile $ Imp.var tmp pt
 
@@ -134,7 +134,7 @@
   dPrimV_ (segFlat space) (0 :: Imp.TExp Int64)
   iterations <- getIterationDomain op space
   seq_code <- collect $ localOps inThreadOps $ do
-    nsubtasks <- dPrim "nsubtasks" int32
+    nsubtasks <- dPrim "nsubtasks"
     sOp $ Imp.GetNumTasks $ tvVar nsubtasks
     emit =<< compileSegOp pat op nsubtasks
   retvals <- getReturnParams pat op
@@ -146,7 +146,7 @@
       let space' = getSpace nested_op
       dPrimV_ (segFlat space') (0 :: Imp.TExp Int64)
       par_code <- collect $ do
-        nsubtasks <- dPrim "nsubtasks" int32
+        nsubtasks <- dPrim "nsubtasks"
         sOp $ Imp.GetNumTasks $ tvVar nsubtasks
         emit =<< compileSegOp pat nested_op nsubtasks
       pure $ Just $ Imp.ParallelTask par_code
diff --git a/src/Futhark/CodeGen/ImpGen/Multicore/Base.hs b/src/Futhark/CodeGen/ImpGen/Multicore/Base.hs
--- a/src/Futhark/CodeGen/ImpGen/Multicore/Base.hs
+++ b/src/Futhark/CodeGen/ImpGen/Multicore/Base.hs
@@ -86,8 +86,8 @@
 
 getLoopBounds :: MulticoreGen (Imp.TExp Int64, Imp.TExp Int64)
 getLoopBounds = do
-  start <- dPrim "start" int64
-  end <- dPrim "end" int64
+  start <- dPrim "start"
+  end <- dPrim "end"
   emit $ Imp.Op $ Imp.GetLoopBounds (tvVar start) (tvVar end)
   pure (tvExp start, tvExp end)
 
@@ -278,7 +278,7 @@
         -- ISPC v1.17 does not support extract on f16 yet..
         -- Thus we do this stupid conversion to f32
         Prim (FloatType Float16) -> do
-          tv <- dPrim "hack_extract_f16" (FloatType Float32)
+          tv :: TV Float <- dPrim "hack_extract_f16"
           emit $ Imp.SetScalar (tvVar tv) e
           emit $ Imp.Op $ Imp.ExtractLane vname (untyped $ tvExp tv) ut_exp
         _ -> emit $ Imp.Op $ Imp.ExtractLane vname e ut_exp
@@ -383,14 +383,14 @@
         -- compare-and-swap if 32 bits.
         forM_ (zip arrs ops_and_ts) $ \(a, (op, t, x, y)) -> do
           -- Common variables.
-          old <- dPrim "old" t
+          old <- dPrimS "old" t
 
           (arr', _a_space, bucket_offset) <- fullyIndexArray a bucket
 
-          case opHasAtomicSupport (tvVar old) arr' (sExt32 <$> bucket_offset) op of
+          case opHasAtomicSupport old arr' (sExt32 <$> bucket_offset) op of
             Just f -> sOp $ f $ Imp.var y t
             Nothing ->
-              atomicUpdateCAS t a (tvVar old) bucket x $
+              atomicUpdateCAS t a old bucket x $
                 x <~~ Imp.BinOpExp op (Imp.var x t) (Imp.var y t)
   where
     opHasAtomicSupport old arr' bucket' bop = do
@@ -406,12 +406,12 @@
   | [Prim t] <- lambdaReturnType op,
     [xp, _] <- lambdaParams op,
     supportedPrims (primBitSize t) = AtomicCAS $ \[arr] bucket -> do
-      old <- dPrim "old" t
-      atomicUpdateCAS t arr (tvVar old) bucket (paramName xp) $
+      old <- dPrimS "old" t
+      atomicUpdateCAS t arr old bucket (paramName xp) $
         compileBody' [xp] $
           lambdaBody op
 atomicUpdateLocking _ op = AtomicLocking $ \locking arrs bucket -> do
-  old <- dPrim "old" int32
+  old <- dPrim "old"
   continue <- dPrimVol "continue" int32 (0 :: Imp.TExp Int32)
 
   -- Correctly index into locks.
@@ -512,7 +512,7 @@
 
   everythingVolatile $ copyDWIMFix old [] (Var arr) bucket
 
-  old_bits_v <- tvVar <$> dPrim "old_bits" int
+  old_bits_v <- dPrimS "old_bits" int
   old_bits_v <~~ toBits (Imp.var old t)
   let old_bits = Imp.var old_bits_v int
 
diff --git a/src/Futhark/CodeGen/ImpGen/Multicore/SegHist.hs b/src/Futhark/CodeGen/ImpGen/Multicore/SegHist.hs
--- a/src/Futhark/CodeGen/ImpGen/Multicore/SegHist.hs
+++ b/src/Futhark/CodeGen/ImpGen/Multicore/SegHist.hs
@@ -256,7 +256,7 @@
                   sComment "perform updates" $ do
                     -- Create new set of uniform buckets
                     -- That is extract each bucket from a SIMD vector lane
-                    extract_buckets <- mapM (dPrim "extract_bucket" . (primExpType . untyped)) bucket'
+                    extract_buckets <- mapM (dPrimSV "extract_bucket" . (primExpType . untyped)) bucket'
                     forM_ (zip extract_buckets bucket') $ \(x, y) ->
                       emit $ Imp.Op $ Imp.ExtractLane (tvVar x) (untyped y) (untyped j)
                     let bucket'' = map tvExp extract_buckets
@@ -269,11 +269,10 @@
                         forM_ (zip vs_params' vs') $ \(p, res) ->
                           ifPrimType (paramType p) $ \pt -> do
                             -- Hack to copy varying load into uniform result variable
-                            tmp <- dPrim "tmp" pt
-                            copyDWIMFix (tvVar tmp) [] res is'
-                            extractVectorLane j $
-                              pure $
-                                Imp.SetScalar (paramName p) (Imp.LeafExp (tvVar tmp) pt)
+                            tmp <- dPrimS "tmp" pt
+                            copyDWIMFix tmp [] res is'
+                            extractVectorLane j . pure $
+                              Imp.SetScalar (paramName p) (toExp' pt tmp)
                         updateHisto histop' histop_subhistograms (bucket'' ++ is') j acc_params'
 
     -- Copy the task-local subhistograms to the global subhistograms,
@@ -299,7 +298,7 @@
         segred_op = SegBinOp Noncommutative (histOp op) (histNeutral op) (histOpShape op)
 
     red_code <- collect $ do
-      nsubtasks <- dPrim "nsubtasks" int32
+      nsubtasks <- dPrim "nsubtasks"
       sOp $ Imp.GetNumTasks $ tvVar nsubtasks
       emit <=< compileSegRed' (Pat red_pes) segred_space [segred_op] nsubtasks $ \red_cont ->
         red_cont $
diff --git a/src/Futhark/CodeGen/ImpGen/Multicore/SegRed.hs b/src/Futhark/CodeGen/ImpGen/Multicore/SegRed.hs
--- a/src/Futhark/CodeGen/ImpGen/Multicore/SegRed.hs
+++ b/src/Futhark/CodeGen/ImpGen/Multicore/SegRed.hs
@@ -140,7 +140,7 @@
         case paramType p of
           Prim pt
             | shape == mempty ->
-                tvVar <$> dPrim "local_acc" pt
+                dPrimS "local_acc" pt
             | otherwise ->
                 sAllocArray "local_acc" pt shape DefaultSpace
           _ ->
@@ -360,7 +360,7 @@
   dScope Nothing $ scopeOfLParams $ concatMap slugParams slugs
 
   sFor "i" nsubtasks $ \i' -> do
-    mkTV (segFlat space) int64 <-- i'
+    mkTV (segFlat space) <-- i'
     sComment "Apply main thread reduction" $
       forM_ (zip slugs per_red_pes) $ \(slug, red_res) ->
         sLoopNest (slugShape slug) $ \vec_is -> do
@@ -422,7 +422,7 @@
         sFor "i" inner_bound $ \i -> do
           zipWithM_
             (<--)
-            (map (`mkTV` int64) $ init is)
+            (map mkTV $ init is)
             (unflattenIndex (init ns_64) (sExt64 n_segments))
           dPrimV_ (last is) i
           kbody $ \red_res' -> do
diff --git a/src/Futhark/CodeGen/OpenCL/Heuristics.hs b/src/Futhark/CodeGen/OpenCL/Heuristics.hs
--- a/src/Futhark/CodeGen/OpenCL/Heuristics.hs
+++ b/src/Futhark/CodeGen/OpenCL/Heuristics.hs
@@ -58,6 +58,7 @@
     SizeHeuristic "" DeviceGPU TileSize 16,
     SizeHeuristic "" DeviceGPU RegTileSize 4,
     SizeHeuristic "" DeviceGPU Threshold $ 32 * 1024,
+    --
     SizeHeuristic "" DeviceCPU LockstepWidth 1,
     SizeHeuristic "" DeviceCPU NumBlocks max_compute_units,
     SizeHeuristic "" DeviceCPU BlockSize 32,
diff --git a/src/Futhark/Construct.hs b/src/Futhark/Construct.hs
--- a/src/Futhark/Construct.hs
+++ b/src/Futhark/Construct.hs
@@ -123,7 +123,7 @@
 import Control.Monad
 import Control.Monad.Identity
 import Control.Monad.State
-import Data.List (foldl', sortOn, transpose)
+import Data.List qualified as L
 import Data.Map.Strict qualified as M
 import Futhark.Builder
 import Futhark.IR
@@ -213,8 +213,8 @@
 removeRedundantScrutinees :: [SubExp] -> [Case b] -> ([SubExp], [Case b])
 removeRedundantScrutinees ses cases =
   let (ses', vs) =
-        unzip $ filter interesting $ zip ses $ transpose (map casePat cases)
-   in (ses', zipWith Case (transpose vs) $ map caseBody cases)
+        unzip $ filter interesting $ zip ses $ L.transpose (map casePat cases)
+   in (ses', zipWith Case (L.transpose vs) $ map caseBody cases)
   where
     interesting = any (/= Nothing) . snd
 
@@ -230,7 +230,7 @@
   cases <- mapM (traverse insertStmsM) cases_m
   defbody <- insertStmsM defbody_m
   ts <-
-    foldl' generaliseExtTypes
+    L.foldl' generaliseExtTypes
       <$> bodyExtType defbody
       <*> mapM (bodyExtType . caseBody) cases
   cases' <- mapM (traverse $ addContextForBranch ts) cases
@@ -242,7 +242,7 @@
     addContextForBranch ts (Body _ stms val_res) = do
       body_ts <- extendedScope (traverse subExpResType val_res) stmsscope
       let ctx_res =
-            map snd $ sortOn fst $ M.toList $ shapeExtMapping ts body_ts
+            map snd $ L.sortOn fst $ M.toList $ shapeExtMapping ts body_ts
       mkBodyM stms $ subExpsRes ctx_res ++ val_res
       where
         stmsscope = scopeOf stms
diff --git a/src/Futhark/IR/Prop/Types.hs b/src/Futhark/IR/Prop/Types.hs
--- a/src/Futhark/IR/Prop/Types.hs
+++ b/src/Futhark/IR/Prop/Types.hs
@@ -67,7 +67,7 @@
 
 import Control.Monad
 import Control.Monad.State
-import Data.List (elemIndex, foldl')
+import Data.List qualified as L
 import Data.Map.Strict qualified as M
 import Data.Maybe
 import Data.Set qualified as S
@@ -502,7 +502,7 @@
     makeBoundShapesFree =
       modifyArrayShape $ fmap checkDim
     checkDim (Free (Var v))
-      | Just i <- v `elemIndex` inaccessible =
+      | Just i <- v `L.elemIndex` inaccessible =
           Ext i
     checkDim d = d
 
@@ -523,7 +523,7 @@
   [t2] ->
   res
 dimMapping getDims1 getDims2 f comb ts1 ts2 =
-  foldl' comb mempty $ concat $ zipWith (zipWith f) (map getDims1 ts1) (map getDims2 ts2)
+  L.foldl' comb mempty $ concat $ zipWith (zipWith f) (map getDims1 ts1) (map getDims2 ts2)
 
 -- | @IntType Int8@
 int8 :: PrimType
diff --git a/src/Futhark/Optimise/BlkRegTiling.hs b/src/Futhark/Optimise/BlkRegTiling.hs
--- a/src/Futhark/Optimise/BlkRegTiling.hs
+++ b/src/Futhark/Optimise/BlkRegTiling.hs
@@ -46,6 +46,23 @@
 se8 :: SubExp
 se8 = intConst Int64 8
 
+isInnerCoal :: Env -> VName -> Stm GPU -> Bool
+isInnerCoal (_, ixfn_env) slc_X (Let (Pat [pe]) _ (BasicOp (Index x _)))
+  | slc_X == patElemName pe,
+    Nothing <- M.lookup x ixfn_env =
+      True -- if not in the table, we assume not-transposed!
+isInnerCoal (_, ixfn_env) slc_X (Let (Pat [pe]) _ (BasicOp (Index x _)))
+  | slc_X == patElemName pe,
+    Just lmad <- M.lookup x ixfn_env =
+      innerHasStride1 lmad
+  where
+    innerHasStride1 lmad =
+      let lmad_dims = LMAD.dims lmad
+          stride = LMAD.ldStride $ last lmad_dims
+       in stride == pe64 (intConst Int64 1)
+isInnerCoal _ _ _ =
+  error "kkLoopBody.isInnerCoal: not an error, but I would like to know why!"
+
 scratch :: (MonadBuilder m) => String -> PrimType -> [SubExp] -> m VName
 scratch se_name t shape = letExp se_name $ BasicOp $ Scratch t shape
 
@@ -81,165 +98,139 @@
     kk <- letExp "kk" =<< toExp (le64 kk0 * pe64 tk)
     -- copy A to shared memory
     (a_loc, aCopyLoc2Reg) <-
-      copyGlb2ShMem kk (gtid_y, iii, map_t1, height_A, inp_A, load_A, a_loc_init')
+      copyGlb2ShMem False kk (gtid_y, iii, map_t1, height_A, inp_A, load_A, a_loc_init')
 
     -- copy B from global to shared memory
     (b_loc, bCopyLoc2Reg) <-
-      copyGlb2ShMem kk (gtid_x, jjj, map_t2, width_B, inp_B, load_B, b_loc_init')
+      copyGlb2ShMem True kk (gtid_x, jjj, map_t2, width_B, inp_B, load_B, b_loc_init')
 
     -- inner loop updating this thread's accumulator (loop k in mmm_kernels).
-    thd_acc <- forLoop tk [thd_res_merge] $ \k [acc_merge] ->
-      resultBodyM
-        =<< letTupExp' "foo"
-        =<< eIf
-          ( toExp $
-              if epilogue
-                then le64 kk + le64 k .<. pe64 common_dim
-                else true -- if in prologue, always compute redomap.
-          )
-          ( do
-              reg_mem <- segMap2D "reg_mem" segthd_lvl ResultPrivate (ty, tx) $
-                \(ltid_y, ltid_x) -> do
-                  -- copy A from shared memory to registers
-                  asss <- aCopyLoc2Reg k ltid_y
-                  -- copy B from shared memory to registers
-                  bsss <- bCopyLoc2Reg k ltid_x
-                  pure $ varsRes [asss, bsss]
-              let [asss, bsss] = reg_mem
-              mkRedomapOneTileBody acc_merge asss bsss True
-          )
-          (resultBodyM [Var acc_merge])
+    thd_acc <- mkRedomapOneTileBody kk thd_res_merge aCopyLoc2Reg bCopyLoc2Reg True
     pure [thd_acc, a_loc, b_loc]
     where
-      mk_ik is_coal (thd_y, thd_x) (i0, k0)
+      mk_ik is_B is_coal (thd_y, thd_x) (i0, k0)
         | is_coal = do
             -- not-transposed case (i.e., already coalesced)
             let (t_par, t_seq) = (tx, tk)
             k <- letExp "k" =<< toExp (le64 thd_x + le64 k0 * pe64 t_par)
             i <- letExp "i" =<< toExp (le64 thd_y + le64 i0 * pe64 t_par)
-            -- we have padded to minimize bank conflicts,
-            -- hence the length of inner dim is (t_seq + 1)
-            let e = le64 k + le64 i * (pe64 t_seq + pe64 se1)
+            -- to optimize bank conflicts, we use padding only for B
+            -- iff B has the last dimension permuted.
+            let pad_term = if is_B then pe64 se1 else pe64 se0
+            let e = le64 k + le64 i * (pe64 t_seq + pad_term)
             pure (i, k, e)
-      mk_ik _ (thd_y, thd_x) (i0, k0) = do
+      mk_ik _ _ (thd_y, thd_x) (i0, k0) = do
         -- matrix is transposed case (i.e., uncoalesced):
         let (t_par, tr_par) = (tx, tx_rx)
         k <- letExp "k" =<< toExp (le64 thd_y + le64 k0 * pe64 t_par)
         i <- letExp "i" =<< toExp (le64 thd_x + le64 i0 * pe64 t_par)
-        -- we have padded to minimize bank conflicts,
-        -- hence the length of inner dim is (tr_par + 1)
-        let e = le64 i + le64 k * (pe64 tr_par + pe64 se1)
+        -- no padding
+        let e = le64 i + le64 k * pe64 tr_par
         pure (i, k, e)
-      isInnerCoal :: Env -> VName -> Stm GPU -> Bool
-      isInnerCoal (_, ixfn_env) slc_X (Let pat _ (BasicOp (Index x _)))
-        | [slc_X'] <- patNames pat,
-          slc_X == slc_X',
-          Nothing <- M.lookup x ixfn_env =
-            True -- if not in the table, we assume not-transposed!
-      isInnerCoal (_, ixfn_env) slc_X (Let pat _ (BasicOp (Index x _)))
-        | [slc_X'] <- patNames pat,
-          slc_X == slc_X',
-          Just lmad <- M.lookup x ixfn_env =
-            innerHasStride1 lmad
-      isInnerCoal _ _ _ =
-        error "kkLoopBody.isInnerCoal: not an error, but I would like to know why!"
-      innerHasStride1 lmad =
-        let lmad_dims = LMAD.dims lmad
-            stride = LMAD.ldStride $ last lmad_dims
-         in stride == pe64 (intConst Int64 1)
       --
-      mkRedomapOneTileBody acc_merge asss bsss fits_ij = do
-        -- the actual redomap.
-        redomap_res <- segMap2D "redomap_res" segthd_lvl ResultPrivate (ty, tx) $
-          \(ltid_y, ltid_x) -> do
-            as <- index "as" asss [ltid_y, ltid_x]
-            bs <- index "bs" bsss [ltid_y, ltid_x]
-            css_init <- index "css_init" acc_merge [ltid_y, ltid_x]
+      --
+      mkCompLoopRxRy fits_ij css_init (a_idx_fn, b_idx_fn) (ltid_y, ltid_x) = do
+        css <- forLoop ry [css_init] $ \i [css_merge] -> do
+          css <- forLoop rx [css_merge] $ \j [css_merge'] ->
+            resultBodyM
+              =<< letTupExp' "foo"
+              =<< eIf
+                ( toExp $
+                    if fits_ij
+                      then true
+                      else -- this condition is never needed because
+                      -- if i and j are out of range than css[i,j]
+                      -- is garbage anyways and should not be written.
+                      -- so fits_ij should be always true!!!
 
-            css <- forLoop ry [css_init] $ \i [css_merge] -> do
-              css <- forLoop rx [css_merge] $ \j [css_merge'] ->
-                resultBodyM
-                  =<< letTupExp' "foo"
-                  =<< eIf
-                    ( toExp $
-                        if fits_ij
-                          then true
-                          else -- this condition is never needed because
-                          -- if i and j are out of range than css[i,j]
-                          -- is garbage anyways and should not be written.
-                          -- so fits_ij should be always true!!!
+                        le64 iii
+                          + le64 i
+                          + pe64 ry
+                            * le64 ltid_y
+                              .<. pe64 height_A
+                              .&&. le64 jjj
+                          + le64 j
+                          + pe64 rx
+                            * le64 ltid_x
+                              .<. pe64 width_B
+                )
+                ( do
+                    a <- a_idx_fn ltid_y i
+                    b <- b_idx_fn ltid_x j
+                    c <- index "c" css_merge' [i, j]
 
-                            le64 iii
-                              + le64 i
-                              + pe64 ry
-                                * le64 ltid_y
-                                  .<. pe64 height_A
-                                  .&&. le64 jjj
-                              + le64 j
-                              + pe64 rx
-                                * le64 ltid_x
-                                  .<. pe64 width_B
-                    )
-                    ( do
-                        a <- index "a" as [i]
-                        b <- index "b" bs [j]
-                        c <- index "c" css_merge' [i, j]
+                    map_lam' <- renameLambda map_lam
+                    red_lam' <- renameLambda red_lam
 
-                        map_lam' <- renameLambda map_lam
-                        red_lam' <- renameLambda red_lam
+                    -- the inputs to map are supposed to be permutted with the
+                    -- inverted permutation, so as to reach the original position;
+                    -- it just so happens that the inverse of [a,b] is [b,a]
+                    let map_inp_reg = if var_dims == [0, 1] then [a, b] else [b, a]
 
-                        -- the inputs to map are supposed to be permutted with the
-                        -- inverted permutation, so as to reach the original position;
-                        -- it just so happens that the inverse of [a,b] is [b,a]
-                        let map_inp_reg = if var_dims == [0, 1] then [a, b] else [b, a]
+                    map_res <- eLambda map_lam' (map (eSubExp . Var) map_inp_reg)
+                    ~[red_res] <- eLambda red_lam' (map eSubExp $ Var c : map resSubExp map_res)
+                    css <- update "css" css_merge' [i, j] (resSubExp red_res)
 
-                        map_res <- eLambda map_lam' (map (eSubExp . Var) map_inp_reg)
-                        ~[red_res] <- eLambda red_lam' (map eSubExp $ Var c : map resSubExp map_res)
-                        css <- update "css" css_merge' [i, j] (resSubExp red_res)
+                    resultBodyM [Var css]
+                )
+                (resultBodyM [Var css_merge'])
+          resultBodyM [Var css]
+        resultBodyM [Var css]
+      --
+      mkRedomapOneTileBody kk css_merge a_idx_fn b_idx_fn fits_ij = do
+        -- the actual redomap.
+        redomap_res <- segMap2D "redomap_res" segthd_lvl ResultPrivate (ty, tx) $
+          \(ltid_y, ltid_x) -> do
+            css_init <- index "css_init" css_merge [ltid_y, ltid_x]
 
-                        resultBodyM [Var css]
-                    )
-                    (resultBodyM [Var css_merge'])
-              resultBodyM [Var css]
+            css <- forLoop tk [css_init] $ \k [acc_merge] ->
+              resultBodyM
+                =<< letTupExp' "foo"
+                =<< eIf
+                  ( toExp $
+                      if epilogue
+                        then le64 kk + le64 k .<. pe64 common_dim
+                        else true -- if in prologue, always compute redomap.
+                  )
+                  (mkCompLoopRxRy fits_ij acc_merge (a_idx_fn k, b_idx_fn k) (ltid_y, ltid_x))
+                  (resultBodyM [Var acc_merge])
+
             pure [varRes css]
-        resultBodyM $ map Var redomap_res
+        pure $ head redomap_res
       --
       copyGlb2ShMem ::
+        Bool ->
         VName ->
         (VName, VName, PrimType, SubExp, VName, Stm GPU, VName) ->
-        Builder GPU (VName, VName -> VName -> Builder GPU VName)
-      copyGlb2ShMem kk (gtid, ii, ptp_X_el, parlen_X, inp_X, load_X, x_loc_init') = do
+        Builder GPU (VName, VName -> VName -> VName -> Builder GPU VName)
+      copyGlb2ShMem is_B kk (gtid, ii, ptp_X_el, parlen_X, inp_X, load_X, x_loc_init') = do
         let (t_par, r_par, tseq_div_tpar) = (tx, rx, tk_div_tx)
             is_inner_coal = isInnerCoal env inp_X load_X
             str_A = baseString inp_X
         x_loc <-
           segScatter2D (str_A ++ "_glb2loc") x_loc_init' [r_par, tseq_div_tpar] (t_par, t_par) $
             scatterFun is_inner_coal
-
-        pure (x_loc, copyLoc2Reg is_inner_coal str_A x_loc)
+        pure (x_loc, indexLocMem is_inner_coal str_A x_loc)
         where
-          copyLoc2Reg ::
+          indexLocMem ::
             Bool ->
             String ->
             VName ->
             VName ->
             VName ->
+            VName ->
             Builder GPU VName
-          copyLoc2Reg is_inner_coal str_A x_loc k ltid_yx = do
+          indexLocMem is_inner_coal str_A x_loc k ltid_yx ij = do
             let (r_par, t_seq, tr_par) = (rx, tk, tx_rx)
-            xsss_init <- scratch (str_A ++ "_init_regs") ptp_X_el [r_par]
-            forLoop r_par [xsss_init] $ \ij [xsss_merge] -> do
-              x_loc_ind <-
-                letExp (str_A ++ "_loc_ind")
-                  =<< toExp
-                    ( if is_inner_coal
-                        then le64 k + (le64 ltid_yx * pe64 r_par + le64 ij) * (pe64 t_seq + pe64 se1)
-                        else le64 ij + le64 ltid_yx * pe64 r_par + le64 k * (pe64 tr_par + pe64 se1)
-                    )
-              xsss <-
-                update (str_A ++ "_regs") xsss_merge [ij] . Var
-                  =<< index (str_A ++ "_loc_elem") x_loc [x_loc_ind]
-              resultBodyM [Var xsss]
+            let pad_term = if is_B then pe64 se1 else pe64 se0
+            x_loc_ind_32 <-
+              letExp (str_A ++ "_loc_ind_64")
+                =<< toExp
+                  ( if is_inner_coal -- ToDo: check this is correct + turn to i32
+                      then le64 k + (le64 ltid_yx * pe64 r_par + le64 ij) * (pe64 t_seq + pad_term)
+                      else le64 ij + le64 ltid_yx * pe64 r_par + le64 k * pe64 tr_par
+                  )
+            index (str_A ++ "_loc_elem") x_loc [x_loc_ind_32]
           --
           scatterFun ::
             Bool ->
@@ -249,7 +240,7 @@
           scatterFun is_inner_coal [i0, k0] (thd_y, thd_x) = do
             let str_A = baseString inp_X
                 t_seq = tk
-            (i, k, epx_loc_fi) <- mk_ik is_inner_coal (thd_y, thd_x) (i0, k0)
+            (i, k, epx_loc_fi) <- mk_ik is_B is_inner_coal (thd_y, thd_x) (i0, k0)
             letBindNames [gtid] =<< toExp (le64 ii + le64 i)
             a_seqdim_idx <- letExp (str_A ++ "_seqdim_idx") =<< toExp (le64 kk + le64 k)
 
@@ -312,11 +303,12 @@
       matchesBlkRegTile seg_space kstms,
     checkAccumulatesRedomapRes res_nm code2' redomap_orig_res = do
       -- Here we start the implementation --
+      let is_B_coal = isInnerCoal env inp_B load_B
       ---- in this binder: host code and outer seggroup (ie. the new kernel) ----
       (new_kernel, host_stms) <- runBuilder $ do
         -- host code
         (rx, ry, tx, ty, tk, tk_div_tx, tk_div_ty, tx_rx, ty_ry, a_loc_sz, b_loc_sz) <-
-          mkTileMemSizes height_A width_B common_dim
+          mkTileMemSizes height_A width_B common_dim is_B_coal
 
         rk <- letSubExp "rk" $ BasicOp $ SubExp $ intConst Int64 8 -- 16 and 8 seem good values
         tk_rk <- letSubExp "tk_rk" =<< toExp (pe64 tk * pe64 rk)
@@ -436,7 +428,7 @@
       foldl getAccumStm False $ reverse $ stmsToList acc_code
       where
         getAccumStm True _ = True
-        getAccumStm False (Let (Pat [pat_el]) _aux (BasicOp (UpdateAcc Safe _acc_nm _ind vals)))
+        getAccumStm False (Let (Pat [pat_el]) _aux (BasicOp (UpdateAcc _ _acc_nm _ind vals)))
           | [v] <- vals,
             patElemName pat_el == res_nm =
               v == Var redomap_orig_res
@@ -505,11 +497,12 @@
         ) <-
       matchesBlkRegTile seg_space kstms = do
       -- Here we start the implementation
+      let is_B_coal = isInnerCoal env inp_B load_B
       ---- in this binder: host code and outer seggroup (ie. the new kernel) ----
       (new_kernel, host_stms) <- runBuilder $ do
         -- host code
         (rx, ry, tx, ty, tk, tk_div_tx, tk_div_ty, tx_rx, ty_ry, a_loc_sz, b_loc_sz) <-
-          mkTileMemSizes height_A width_B common_dim
+          mkTileMemSizes height_A width_B common_dim is_B_coal
 
         gridDim_x <- letSubExp "gridDim_x" =<< ceilDiv width_B tx_rx
         gridDim_y <- letSubExp "gridDim_y" =<< ceilDiv height_A ty_ry
@@ -725,6 +718,7 @@
   SubExp ->
   SubExp ->
   SubExp ->
+  Bool ->
   Builder
     GPU
     ( SubExp,
@@ -739,15 +733,17 @@
       SubExp,
       SubExp
     )
-mkTileMemSizes height_A width_B common_dim = do
+mkTileMemSizes height_A _width_B common_dim is_B_not_transp = do
   tk_name <- nameFromString . prettyString <$> newVName "Tk"
-  tx_name <- nameFromString . prettyString <$> newVName "Tx"
   ty_name <- nameFromString . prettyString <$> newVName "Ty"
-  rx_name <- nameFromString . prettyString <$> newVName "Rx"
   ry_name <- nameFromString . prettyString <$> newVName "Ry"
 
+  -- until we change the copying to use lmads we need to
+  --   guarantee that Tx=Ty AND Rx = Ry AND Tx | Tk
+  -- for matrix multiplication it would be safe if they aren't
+  --   but not for any of the other three cases!
   (ty, ry) <- getParTiles ("Ty", "Ry") (ty_name, ry_name) height_A
-  (tx, rx) <- getParTiles ("Tx", "Rx") (tx_name, rx_name) width_B
+  let (tx, rx) = (ty, ry)
   tk <- getSeqTile "Tk" tk_name common_dim tx ty
 
   tk_div_tx <- letSubExp "tk_div_tx" =<< ceilDiv tk tx
@@ -756,19 +752,19 @@
   tx_rx <- letSubExp "TxRx" =<< toExp (pe64 tx * pe64 rx)
   ty_ry <- letSubExp "TyRy" =<< toExp (pe64 ty * pe64 ry)
 
-  let pad_term = sMax64 (pe64 tk) (pe64 ty * pe64 ry)
-  -- if A not transposed, its shmem should be [ty*ry][tk]
-  -- we pad to [ty*ry][tk+1] size to minimize bank conflicts
+  -- let pad_term = sMax64 (pe64 tk) (pe64 ty * pe64 ry)
+  let pad_term =
+        if is_B_not_transp
+          then pe64 ty * pe64 ry
+          else pe64 se0
   a_loc_sz <-
     letSubExp "a_loc_sz"
-      =<< toExp (pe64 ty * pe64 ry * pe64 tk + pad_term)
+      =<< toExp (pe64 ty * pe64 ry * pe64 tk)
   -- if B is transposed, its shmem should be [tk][tx*rx]
   -- we pad as above, by assuming tx*rx == ty*ry >= tk
-  -- ToDo: we can decrease the size by checking at this
-  --       point whether A and B are transposed (or not).
   b_loc_sz <-
     letSubExp "b_loc_sz"
-      =<< toExp (pe64 tx * pe64 rx * pe64 tk + pad_term) -- (pe64 tk * pe64 tx * pe64 rx)
+      =<< toExp (pe64 tx * pe64 rx * pe64 tk + pad_term)
   pure (rx, ry, tx, ty, tk, tk_div_tx, tk_div_ty, tx_rx, ty_ry, a_loc_sz, b_loc_sz)
 
 mkNewSegthdLvl ::
diff --git a/src/Futhark/Optimise/ReduceDeviceSyncs/MigrationTable/Graph.hs b/src/Futhark/Optimise/ReduceDeviceSyncs/MigrationTable/Graph.hs
--- a/src/Futhark/Optimise/ReduceDeviceSyncs/MigrationTable/Graph.hs
+++ b/src/Futhark/Optimise/ReduceDeviceSyncs/MigrationTable/Graph.hs
@@ -83,7 +83,7 @@
 
 import Data.IntMap.Strict qualified as IM
 import Data.IntSet qualified as IS
-import Data.List (foldl')
+import Data.List qualified as L
 import Data.Map.Strict qualified as M
 import Data.Maybe (fromJust)
 import Prelude hiding (lookup)
@@ -286,7 +286,7 @@
 -- in @srcs@. Returns the ids for the vertices connected to each found sink.
 routeMany :: [Id] -> Graph m -> ([Id], Graph m)
 routeMany srcs graph =
-  foldl' f ([], graph) srcs
+  L.foldl' f ([], graph) srcs
   where
     f (snks, g) src =
       case route src g of
@@ -489,7 +489,7 @@
       let (res, g0) = f g (IM.insert i d p)
        in case res of
             CycleDetected d' as _
-              | d == d' -> (DeadEnd, foldl' (\g1 a -> a g1) g0 as)
+              | d == d' -> (DeadEnd, L.foldl' (\g1 a -> a g1) g0 as)
             _ | otherwise -> (res, g0)
 
     routeAll rev v g0 p0 =
diff --git a/src/Futhark/Optimise/TileLoops.hs b/src/Futhark/Optimise/TileLoops.hs
--- a/src/Futhark/Optimise/TileLoops.hs
+++ b/src/Futhark/Optimise/TileLoops.hs
@@ -919,7 +919,7 @@
   let tblock_size = tile_size
 
   (grid, space) <- do
-    -- How many groups we need to exhaust the innermost dimension.
+    -- How many blocks we need to exhaust the innermost dimension.
     ldim <-
       letSubExp "ldim" . BasicOp $
         BinOp (SDivUp Int64 Unsafe) kdim tblock_size
diff --git a/src/Futhark/Pkg/Info.hs b/src/Futhark/Pkg/Info.hs
--- a/src/Futhark/Pkg/Info.hs
+++ b/src/Futhark/Pkg/Info.hs
@@ -22,7 +22,7 @@
 import Control.Monad.IO.Class
 import Data.ByteString qualified as BS
 import Data.IORef
-import Data.List (foldl', intersperse)
+import Data.List qualified as L
 import Data.Map qualified as M
 import Data.Maybe
 import Data.Text qualified as T
@@ -322,7 +322,7 @@
                 [] -> "Package " <> p <> " has no versions.  Invalid package path?"
                 ks ->
                   "Known versions: "
-                    <> T.concat (intersperse ", " $ map prettySemVer ks)
+                    <> T.concat (L.intersperse ", " $ map prettySemVer ks)
               major
                 | (_, vs) <- majorRevOfPkg p,
                   _svMajor v `notElem` vs =
@@ -356,4 +356,4 @@
     [] -> do
       logMsg $ "Package " <> p <> " has no released versions.  Using HEAD."
       fst <$> lookupPackageCommit cachedir p Nothing
-    v : vs -> pure $ foldl' max v vs
+    v : vs -> pure $ L.foldl' max v vs
diff --git a/src/Futhark/Test.hs b/src/Futhark/Test.hs
--- a/src/Futhark/Test.hs
+++ b/src/Futhark/Test.hs
@@ -417,24 +417,23 @@
   unless (null missing) $ do
     void $ compileProgram [] futhark compiler prog
 
-    res <- liftIO $
-      flip (pmapIO concurrency) missing $ \(entry, tr) -> do
-        (code, stdout, stderr) <- runProgram futhark "" ["-b"] prog entry $ runInput tr
-        case code of
-          ExitFailure e ->
-            pure $
-              Left
-                [ T.pack $
-                    "Reference dataset generation failed with exit code "
-                      ++ show e
-                      ++ " and stderr:\n"
-                      ++ map (chr . fromIntegral) (SBS.unpack stderr)
-                ]
-          ExitSuccess -> do
-            let f = file (entry, tr)
-            liftIO $ createDirectoryIfMissing True $ takeDirectory f
-            SBS.writeFile f stdout
-            pure $ Right ()
+    res <- liftIO . flip (pmapIO concurrency) missing $ \(entry, tr) -> do
+      (code, stdout, stderr) <- runProgram futhark "" ["-b"] prog entry $ runInput tr
+      case code of
+        ExitFailure e ->
+          pure $
+            Left
+              [ T.pack $
+                  "Reference dataset generation failed with exit code "
+                    ++ show e
+                    ++ " and stderr:\n"
+                    ++ map (chr . fromIntegral) (SBS.unpack stderr)
+              ]
+        ExitSuccess -> do
+          let f = file (entry, tr)
+          liftIO $ createDirectoryIfMissing True $ takeDirectory f
+          SBS.writeFile f stdout
+          pure $ Right ()
 
     case sequence_ res of
       Left err -> throwError err
diff --git a/src/Futhark/Test/Spec.hs b/src/Futhark/Test/Spec.hs
--- a/src/Futhark/Test/Spec.hs
+++ b/src/Futhark/Test/Spec.hs
@@ -27,7 +27,7 @@
 import Control.Monad
 import Data.Char
 import Data.Functor
-import Data.List (foldl', (\\))
+import Data.List qualified as L
 import Data.Map.Strict qualified as M
 import Data.Maybe
 import Data.Text qualified as T
@@ -180,7 +180,7 @@
 
 parseNatural :: Parser () -> Parser Int
 parseNatural sep =
-  lexeme sep $ foldl' addDigit 0 . map num <$> some digitChar
+  lexeme sep $ L.foldl' addDigit 0 . map num <$> some digitChar
   where
     addDigit acc x = acc * 10 + x
     num c = ord c - ord '0'
@@ -412,7 +412,7 @@
     noDups xs =
       let xs' = nubOrd xs
        in -- Works because \\ only removes first instance.
-          case xs \\ xs' of
+          case xs L.\\ xs' of
             [] -> Right ()
             x : _ -> Left $ path <> ": multiple datasets with name " <> show (T.unpack x)
 
diff --git a/src/Futhark/Util.hs b/src/Futhark/Util.hs
--- a/src/Futhark/Util.hs
+++ b/src/Futhark/Util.hs
@@ -52,6 +52,7 @@
     fixPoint,
     concatMapM,
     topologicalSort,
+    debugTraceM,
   )
 where
 
@@ -68,7 +69,7 @@
 import Data.Foldable (fold, toList)
 import Data.Function ((&))
 import Data.IntMap qualified as IM
-import Data.List (findIndex, foldl', genericDrop, genericSplitAt, sortBy)
+import Data.List qualified as L
 import Data.List.NonEmpty qualified as NE
 import Data.Map qualified as M
 import Data.Maybe
@@ -78,6 +79,7 @@
 import Data.Text.Encoding.Error qualified as T
 import Data.Time.Clock (UTCTime, getCurrentTime)
 import Data.Tuple (swap)
+import Debug.Trace
 import Numeric
 import System.Directory.Tree qualified as Dir
 import System.Environment
@@ -96,7 +98,7 @@
 
 -- | Like @nubBy@, but without the quadratic runtime.
 nubByOrd :: (a -> a -> Ordering) -> [a] -> [a]
-nubByOrd cmp = map NE.head . NE.groupBy eq . sortBy cmp
+nubByOrd cmp = map NE.head . NE.groupBy eq . L.sortBy cmp
   where
     eq x y = cmp x y == EQ
 
@@ -144,11 +146,11 @@
 
 -- | Like 'maximum', but returns zero for an empty list.
 maxinum :: (Num a, Ord a, Foldable f) => f a -> a
-maxinum = foldl' max 0
+maxinum = L.foldl' max 0
 
 -- | Like 'minimum', but returns zero for an empty list.
 mininum :: (Num a, Ord a, Foldable f) => f a -> a
-mininum xs = foldl' min (maxinum xs) xs
+mininum xs = L.foldl' min (maxinum xs) xs
 
 -- | @dropAt i n@ drops @n@ elements starting at element @i@.
 dropAt :: Int -> Int -> [a] -> [a]
@@ -179,7 +181,7 @@
 -- | Return the list element at the given index, if the index is valid.
 maybeNth :: (Integral int) => int -> [a] -> Maybe a
 maybeNth i l
-  | i >= 0, v : _ <- genericDrop i l = Just v
+  | i >= 0, v : _ <- L.genericDrop i l = Just v
   | otherwise = Nothing
 
 -- | Return the first element of the list, if it exists.
@@ -207,14 +209,14 @@
 -- valid, along with the elements before and after.
 focusNth :: (Integral int) => int -> [a] -> Maybe ([a], a, [a])
 focusNth i xs
-  | (bef, x : aft) <- genericSplitAt i xs = Just (bef, x, aft)
+  | (bef, x : aft) <- L.genericSplitAt i xs = Just (bef, x, aft)
   | otherwise = Nothing
 
 -- | Return the first list element that satisifes a predicate, along with the
 -- elements before and after.
 focusMaybe :: (a -> Maybe b) -> [a] -> Maybe ([a], b, [a])
 focusMaybe f xs = do
-  idx <- findIndex (isJust . f) xs
+  idx <- L.findIndex (isJust . f) xs
   (before, focus, after) <- focusNth idx xs
   res <- f focus
   pure (before, res, after)
@@ -512,3 +514,10 @@
         modify $ second $ IM.insert i True
         mapM_ sorting $ mapMaybe (depends_of node) nodes_idx
         modify $ bimap (node :) (IM.insert i False)
+
+-- | 'traceM', but only if @FUTHARK_COMPILER_DEBUGGING@ is set to to
+-- the appropriate level.
+debugTraceM :: (Monad m) => Int -> String -> m ()
+debugTraceM level
+  | isEnvVarAtLeast "FUTHARK_COMPILER_DEBUGGING" level = traceM
+  | otherwise = const $ pure ()
diff --git a/src/Language/Futhark/Semantic.hs b/src/Language/Futhark/Semantic.hs
--- a/src/Language/Futhark/Semantic.hs
+++ b/src/Language/Futhark/Semantic.hs
@@ -23,7 +23,7 @@
 
 import Data.Map.Strict qualified as M
 import Data.Text qualified as T
-import Futhark.Util (dropLast, fromPOSIX, toPOSIX)
+import Futhark.Util (fromPOSIX, toPOSIX)
 import Futhark.Util.Pretty
 import Language.Futhark
 import System.FilePath qualified as Native
@@ -41,15 +41,14 @@
 mkImportFrom :: ImportName -> String -> ImportName
 mkImportFrom (ImportName includer) includee
   | Posix.isAbsolute includee = ImportName includee
-  | otherwise = ImportName $ Posix.normalise $ Posix.joinPath $ includer' ++ includee'
+  | otherwise =
+      ImportName . Posix.normalise . Posix.joinPath . resolveDotDot [] $
+        init (Posix.splitPath includer) ++ Posix.splitPath includee
   where
-    (dotdots, includee') = span ("../" ==) $ Posix.splitPath includee
-    includer_parts = init $ Posix.splitPath includer
-    includer'
-      | length dotdots > length includer_parts =
-          replicate (length dotdots - length includer_parts) "../"
-      | otherwise =
-          dropLast (length dotdots) includer_parts
+    resolveDotDot parts [] = reverse parts
+    resolveDotDot parts@("../" : _) ("../" : todo) = resolveDotDot ("../" : parts) todo
+    resolveDotDot (_ : parts) ("../" : todo) = resolveDotDot parts todo
+    resolveDotDot parts (p : todo) = resolveDotDot (p : parts) todo
 
 -- | Create a @.fut@ file corresponding to an 'ImportName'.
 includeToFilePath :: ImportName -> Native.FilePath
diff --git a/src/Language/Futhark/TypeChecker/Terms/Loop.hs b/src/Language/Futhark/TypeChecker/Terms/Loop.hs
--- a/src/Language/Futhark/TypeChecker/Terms/Loop.hs
+++ b/src/Language/Futhark/TypeChecker/Terms/Loop.hs
@@ -108,6 +108,24 @@
 type CheckedLoop =
   ([VName], Pat ParamType, Exp, LoopFormBase Info VName, Exp)
 
+checkForImpossible :: Loc -> S.Set VName -> ParamType -> TermTypeM ()
+checkForImpossible loc known_before pat_t = do
+  cs <- getConstraints
+  let bad v = do
+        guard $ v `S.notMember` known_before
+        (_, UnknownSize v_loc _) <- M.lookup v cs
+        Just . typeError (srclocOf loc) mempty $
+          "Inferred type for loop parameter is"
+            </> indent 2 (pretty pat_t)
+            </> "but"
+            <+> dquotes (prettyName v)
+            <+> "is an existential size created inside the loop body at"
+            <+> pretty (locStrRel loc v_loc)
+            <> "."
+  case mapMaybe bad $ S.toList $ fvVars $ freeInType pat_t of
+    problem : _ -> problem
+    [] -> pure ()
+
 -- | Type-check a @loop@ expression, passing in a function for
 -- type-checking subexpressions.
 checkLoop ::
@@ -155,10 +173,13 @@
   -- dim handling (2)
   let checkLoopReturnSize mergepat' loopbody' = do
         loopbody_t <- expTypeFully loopbody'
-        pat_t <-
-          someDimsFreshInType loc "loop" new_dims
-            =<< normTypeFully (patternType mergepat')
+        mergepat_t <- normTypeFully (patternType mergepat')
 
+        let ok_names = known_before <> S.fromList new_dims
+        checkForImpossible (locOf mergepat) ok_names mergepat_t
+
+        pat_t <- someDimsFreshInType loc "loop" new_dims mergepat_t
+
         -- We are ignoring the dimensions here, because any mismatches
         -- should be turned into fresh size variables.
         onFailure (CheckingLoopBody (toStruct pat_t) (toStruct loopbody_t)) $
@@ -225,16 +246,15 @@
             =<< checkExp uboundexp
         bound_t <- expTypeFully uboundexp'
         bindingIdent i bound_t $ \i' ->
-          bindingPat [] mergepat merge_t $
-            \mergepat' -> incLevel $ do
-              loopbody' <- checkExp loopbody
-              (sparams, mergepat'') <- checkLoopReturnSize mergepat' loopbody'
-              pure
-                ( sparams,
-                  mergepat'',
-                  For i' uboundexp',
-                  loopbody'
-                )
+          bindingPat [] mergepat merge_t $ \mergepat' -> incLevel $ do
+            loopbody' <- checkExp loopbody
+            (sparams, mergepat'') <- checkLoopReturnSize mergepat' loopbody'
+            pure
+              ( sparams,
+                mergepat'',
+                For i' uboundexp',
+                loopbody'
+              )
       ForIn xpat e -> do
         (arr_t, _) <- newArrayType (mkUsage' (srclocOf e)) "e" 1
         e' <- unifies "being iterated in a 'for-in' loop" arr_t =<< checkExp e
@@ -243,16 +263,15 @@
           _
             | Just t' <- peelArray 1 t ->
                 bindingPat [] xpat t' $ \xpat' ->
-                  bindingPat [] mergepat merge_t $
-                    \mergepat' -> incLevel $ do
-                      loopbody' <- checkExp loopbody
-                      (sparams, mergepat'') <- checkLoopReturnSize mergepat' loopbody'
-                      pure
-                        ( sparams,
-                          mergepat'',
-                          ForIn (fmap toStruct xpat') e',
-                          loopbody'
-                        )
+                  bindingPat [] mergepat merge_t $ \mergepat' -> incLevel $ do
+                    loopbody' <- checkExp loopbody
+                    (sparams, mergepat'') <- checkLoopReturnSize mergepat' loopbody'
+                    pure
+                      ( sparams,
+                        mergepat'',
+                        ForIn (fmap toStruct xpat') e',
+                        loopbody'
+                      )
             | otherwise ->
                 typeError (srclocOf e) mempty $
                   "Iteratee of a for-in loop must be an array, but expression has type"
diff --git a/src/Language/Futhark/TypeChecker/Terms/Monad.hs b/src/Language/Futhark/TypeChecker/Terms/Monad.hs
--- a/src/Language/Futhark/TypeChecker/Terms/Monad.hs
+++ b/src/Language/Futhark/TypeChecker/Terms/Monad.hs
@@ -10,7 +10,7 @@
   ( TermTypeM,
     runTermTypeM,
     ValBinding (..),
-    SizeSource (SourceBound, SourceSlice),
+    SizeSource (SourceSlice),
     Inferred (..),
     Checking (..),
     withEnv,
@@ -241,7 +241,6 @@
 -- encountered in multiple locations.
 data SizeSource
   = SourceArg FName (ExpBase NoInfo VName)
-  | SourceBound (ExpBase NoInfo VName)
   | SourceSlice
       (Maybe Size)
       (Maybe (ExpBase NoInfo VName))
@@ -491,8 +490,6 @@
   let rsrc = case e of
         SourceArg (FName fname) e' ->
           RigidArg fname $ prettyTextOneLine e'
-        SourceBound e' ->
-          RigidBound $ prettyTextOneLine e'
         SourceSlice d i j s ->
           RigidSlice d $ prettyTextOneLine $ DimSlice i j s
   d <- newRigidDim loc rsrc "n"
diff --git a/src/Language/Futhark/TypeChecker/Types.hs b/src/Language/Futhark/TypeChecker/Types.hs
--- a/src/Language/Futhark/TypeChecker/Types.hs
+++ b/src/Language/Futhark/TypeChecker/Types.hs
@@ -20,7 +20,7 @@
 import Control.Monad.Identity
 import Control.Monad.State
 import Data.Bifunctor
-import Data.List (find, foldl', sort, unzip4, (\\))
+import Data.List qualified as L
 import Data.Map.Strict qualified as M
 import Data.Set qualified as S
 import Futhark.Util (nubOrd)
@@ -59,7 +59,7 @@
 mustBeExplicitInBinding bind_t =
   let (ts, ret) = unfoldFunType bind_t
       alsoRet = M.unionWith (&&) $ M.fromList $ map (,True) (S.toList (fvVars (freeInType ret)))
-   in S.fromList $ M.keys $ M.filter id $ alsoRet $ foldl' onType mempty $ map toStruct ts
+   in S.fromList $ M.keys $ M.filter id $ alsoRet $ L.foldl' onType mempty $ map toStruct ts
   where
     onType uses t = uses <> mustBeExplicitAux t -- Left-biased union.
 
@@ -103,18 +103,18 @@
   pure (TEParens te' loc, svars, ts, ls)
 --
 evalTypeExp df (TETuple ts loc) = do
-  (ts', svars, ts_s, ls) <- unzip4 <$> mapM (evalTypeExp df) ts
+  (ts', svars, ts_s, ls) <- L.unzip4 <$> mapM (evalTypeExp df) ts
   pure
     ( TETuple ts' loc,
       mconcat svars,
       RetType (foldMap retDims ts_s) $ Scalar $ tupleRecord $ map retType ts_s,
-      foldl' max Unlifted ls
+      L.foldl' max Unlifted ls
     )
 --
 evalTypeExp df t@(TERecord fs loc) = do
   -- Check for duplicate field names.
   let field_names = map fst fs
-  unless (sort field_names == sort (nubOrd field_names)) $
+  unless (L.sort field_names == L.sort (nubOrd field_names)) $
     typeError loc mempty $
       "Duplicate record fields in" <+> pretty t <> "."
 
@@ -127,7 +127,7 @@
     ( TERecord (M.toList fs') loc,
       fs_svars,
       RetType (foldMap retDims ts_s) $ Scalar $ Record $ M.map retType ts_s,
-      foldl' max Unlifted ls
+      L.foldl' max Unlifted ls
     )
 --
 evalTypeExp df (TEArray d t loc) = do
@@ -200,7 +200,7 @@
   bindDims dims $ do
     (t', svars, RetType t_dims st, l) <- evalTypeExp df t
     let (witnessed, _) = determineSizeWitnesses $ toStruct st
-    case find (`S.notMember` witnessed) dims of
+    case L.find (`S.notMember` witnessed) dims of
       Just d ->
         typeError loc mempty . withIndexLink "unused-existential" $
           "Existential size "
@@ -220,7 +220,7 @@
 --
 evalTypeExp df t@(TESum cs loc) = do
   let constructors = map fst cs
-  unless (sort constructors == sort (nubOrd constructors)) $
+  unless (L.sort constructors == L.sort (nubOrd constructors)) $
     typeError loc mempty $
       "Duplicate constructors in" <+> pretty t
 
@@ -239,7 +239,7 @@
         Scalar $
           Sum $
             M.map (map retType) ts_s,
-      foldl' max Unlifted ls
+      L.foldl' max Unlifted ls
     )
 evalTypeExp df ote@TEApply {} = do
   (tname, tname_loc, targs) <- rootAndArgs ote
@@ -495,7 +495,7 @@
     onRetType (RetType dims t) = do
       ext <- get
       let (t', ext') = runState (onType t) ext
-          new_ext = ext' \\ ext
+          new_ext = ext' L.\\ ext
       case t of
         Scalar Arrow {} -> do
           put ext'
diff --git a/src/Language/Futhark/TypeChecker/Unify.hs b/src/Language/Futhark/TypeChecker/Unify.hs
--- a/src/Language/Futhark/TypeChecker/Unify.hs
+++ b/src/Language/Futhark/TypeChecker/Unify.hs
@@ -31,7 +31,7 @@
 import Control.Monad
 import Control.Monad.Except
 import Control.Monad.State
-import Data.List (foldl', intersect)
+import Data.List qualified as L
 import Data.Map.Strict qualified as M
 import Data.Maybe
 import Data.Set qualified as S
@@ -165,20 +165,18 @@
     RigidArg (Maybe (QualName VName)) T.Text
   | -- | An existential return size.
     RigidRet (Maybe (QualName VName))
-  | RigidLoop
+  | -- | Similarly to 'RigidRet', but produce by a loop.
+    RigidLoop
   | -- | Produced by a complicated slice expression.
     RigidSlice (Maybe Size) T.Text
   | -- | Produced by a complicated range expression.
     RigidRange
-  | -- | Produced by a range expression with this bound.
-    RigidBound T.Text
   | -- | Mismatch in branches.
     RigidCond StructType StructType
   | -- | Invented during unification.
     RigidUnify
-  | RigidOutOfScope Loc VName
-  | -- | Blank dimension in coercion.
-    RigidCoerce
+  | -- | A name used in a size went out of scope.
+    RigidOutOfScope Loc VName
   deriving (Eq, Ord, Show)
 
 -- | The ridigity of a size variable.  All rigid sizes are tagged with
@@ -222,12 +220,6 @@
   "is unknown size of value returned at" <+> pretty (locStrRel ctx loc) <> "."
 prettySource ctx loc RigidRange =
   "is unknown length of range at" <+> pretty (locStrRel ctx loc) <> "."
-prettySource ctx loc (RigidBound bound) =
-  "generated from expression"
-    </> indent 2 (shorten (pretty bound))
-    </> "used in range at "
-    <> pretty (locStrRel ctx loc)
-    <> "."
 prettySource ctx loc (RigidOutOfScope boundloc v) =
   "is an unknown size arising from "
     <> dquotes (prettyName v)
@@ -237,10 +229,6 @@
       </> "Originally bound at "
     <> pretty (locStrRel ctx boundloc)
     <> "."
-prettySource ctx loc RigidCoerce =
-  "is an unknown size arising from empty dimension in coercion at"
-    <+> pretty (locStrRel ctx loc)
-    <> "."
 prettySource _ _ RigidUnify =
   "is an artificial size invented during unification of functions with anonymous sizes."
 prettySource ctx loc (RigidCond t1 t2) =
@@ -514,7 +502,7 @@
 
                 -- Delete the size variables we introduced to represent
                 -- the existential sizes.
-                modifyConstraints $ \m -> foldl' (flip M.delete) m (b1_dims <> b2_dims)
+                modifyConstraints $ \m -> L.foldl' (flip M.delete) m (b1_dims <> b2_dims)
             where
               (b1', b2') =
                 -- Replace one parameter name with the other in the
@@ -873,7 +861,7 @@
   vn_constraint <- M.lookup vn <$> getConstraints
   case vn_constraint of
     Just (lvl, Overloaded vn_ts vn_usage) ->
-      case ts `intersect` vn_ts of
+      case ts `L.intersect` vn_ts of
         [] ->
           unifyError usage mempty noBreadCrumbs $
             "Type constrained to one of"
diff --git a/unittests/Language/Futhark/SemanticTests.hs b/unittests/Language/Futhark/SemanticTests.hs
new file mode 100644
--- /dev/null
+++ b/unittests/Language/Futhark/SemanticTests.hs
@@ -0,0 +1,24 @@
+module Language.Futhark.SemanticTests (tests) where
+
+import Language.Futhark (ImportName (..))
+import Language.Futhark.Semantic
+import Test.Tasty
+import Test.Tasty.HUnit
+
+tests :: TestTree
+tests =
+  testGroup
+    "Semantic objects"
+    [ testCase "a" $
+        mkInitialImport "a" @?= ImportName "a",
+      testCase "./a" $
+        mkInitialImport "./a" @?= ImportName "a",
+      testCase "a/b -> ../c" $
+        mkImportFrom (mkInitialImport "a/b") "../c" @?= ImportName "c",
+      testCase "a/b -> ../../c" $
+        mkImportFrom (mkInitialImport "a/b") "../../c" @?= ImportName "../c",
+      testCase "../a -> b" $
+        mkImportFrom (mkInitialImport "../a") "b" @?= ImportName "../b",
+      testCase "../a -> ../b" $
+        mkImportFrom (mkInitialImport "../a") "../b" @?= ImportName "../../b"
+    ]
diff --git a/unittests/futhark_tests.hs b/unittests/futhark_tests.hs
--- a/unittests/futhark_tests.hs
+++ b/unittests/futhark_tests.hs
@@ -12,6 +12,7 @@
 import Futhark.Optimise.MemoryBlockMerging.GreedyColoringTests qualified
 import Futhark.Pkg.SolveTests qualified
 import Language.Futhark.PrimitiveTests qualified
+import Language.Futhark.SemanticTests qualified
 import Language.Futhark.SyntaxTests qualified
 import Language.Futhark.TypeCheckerTests qualified
 import Test.Tasty
@@ -33,6 +34,7 @@
       Futhark.Optimise.MemoryBlockMerging.GreedyColoringTests.tests,
       Futhark.Analysis.AlgSimplifyTests.tests,
       Language.Futhark.TypeCheckerTests.tests,
+      Language.Futhark.SemanticTests.tests,
       Futhark.Optimise.ArrayLayoutTests.tests
     ]