diff --git a/docs/man/futhark.rst b/docs/man/futhark.rst
--- a/docs/man/futhark.rst
+++ b/docs/man/futhark.rst
@@ -58,6 +58,12 @@
 
 Print all non-builtin imported Futhark files to stdout, one per line.
 
+futhark query PROGRAM LINE COL
+------------------------------
+
+Print information about the variable at the given position in the
+program.
+
 SEE ALSO
 ========
 
diff --git a/futhark.cabal b/futhark.cabal
--- a/futhark.cabal
+++ b/futhark.cabal
@@ -1,13 +1,13 @@
 cabal-version: 1.12
 
--- This file has been generated from package.yaml by hpack version 0.33.0.
+-- This file has been generated from package.yaml by hpack version 0.31.2.
 --
 -- see: https://github.com/sol/hpack
 --
--- hash: 9e36798f8f4bb7d1cf26a2305298c0f022e8d80c05d3f4ab961a06976a711aeb
+-- hash: b3f4de758564b29b93d0b51ac0628666279014b6a327dfe8d83b743ad959b249
 
 name:           futhark
-version:        0.13.1
+version:        0.13.2
 synopsis:       An optimising compiler for a functional, array-oriented language.
 description:    Futhark is a small programming language designed to be compiled to
                 efficient parallel code. It is a statically typed, data-parallel,
@@ -25,7 +25,6 @@
 license-file:   LICENSE
 build-type:     Simple
 extra-source-files:
-    rts/python/__init__.py
     rts/python/memory.py
     rts/python/opencl.py
     rts/python/panic.py
@@ -117,6 +116,7 @@
       Futhark.CLI.Pkg
       Futhark.CLI.PyOpenCL
       Futhark.CLI.Python
+      Futhark.CLI.Query
       Futhark.CLI.REPL
       Futhark.CLI.Run
       Futhark.CLI.Test
@@ -269,6 +269,7 @@
       Language.Futhark.Interpreter
       Language.Futhark.Parser
       Language.Futhark.Pretty
+      Language.Futhark.Query
       Language.Futhark.Semantic
       Language.Futhark.Syntax
       Language.Futhark.Traversals
@@ -286,9 +287,6 @@
   hs-source-dirs:
       src
   ghc-options: -Wall -Wcompat -Wredundant-constraints -Wincomplete-record-updates -Wmissing-export-lists
-  build-tools:
-      alex
-    , happy
   build-depends:
       aeson >=1.0.0.0
     , ansi-terminal >=0.6.3.1
@@ -336,6 +334,9 @@
     , versions >=3.3.1
     , zip-archive >=0.3.1.1
     , zlib >=0.6.1.2
+  build-tools:
+      alex
+    , happy
   default-language: Haskell2010
 
 executable futhark
diff --git a/package.yaml b/package.yaml
--- a/package.yaml
+++ b/package.yaml
@@ -1,5 +1,5 @@
 name: futhark
-version: "0.13.1"
+version: "0.13.2"
 synopsis: An optimising compiler for a functional, array-oriented language.
 description: |
   Futhark is a small programming language designed to be compiled to
diff --git a/rts/c/cuda.h b/rts/c/cuda.h
--- a/rts/c/cuda.h
+++ b/rts/c/cuda.h
@@ -96,6 +96,7 @@
   size_t max_tile_size;
   size_t max_threshold;
   size_t max_shared_memory;
+  size_t max_bespoke;
 
   size_t lockstep_width;
 };
@@ -213,7 +214,8 @@
     { 6, 1, "compute_61" },
     { 6, 2, "compute_62" },
     { 7, 0, "compute_70" },
-    { 7, 2, "compute_72" }
+    { 7, 2, "compute_72" },
+    { 7, 5, "compute_75" }
   };
 
   int major = device_query(dev, COMPUTE_CAPABILITY_MAJOR);
@@ -231,6 +233,13 @@
   if (chosen == -1) {
     panic(-1, "Unsupported compute capability %d.%d\n", major, minor);
   }
+
+  if (x[chosen].major != major || x[chosen].minor != minor) {
+    fprintf(stderr,
+            "Warning: device compute capability is %d.%d, but newest supported by Futhark is %d.%d.\n",
+            major, minor, x[chosen].major, x[chosen].minor);
+  }
+
   return x[chosen].arch_str;
 }
 
@@ -364,7 +373,8 @@
       max_value = ctx->max_threshold;
       default_value = ctx->cfg.default_threshold;
     } else {
-      panic(1, "Unknown size class for size '%s': %s\n", size_name, size_class);
+      // Bespoke sizes have no limit or default.
+      max_value = 0;
     }
 
     if (*size_value == 0) {
@@ -459,6 +469,7 @@
   ctx->max_grid_size = device_query(ctx->dev, MAX_GRID_DIM_X);
   ctx->max_tile_size = sqrt(ctx->max_block_size);
   ctx->max_threshold = 0;
+  ctx->max_bespoke = 0;
   ctx->lockstep_width = device_query(ctx->dev, WARP_SIZE);
 
   cuda_size_setup(ctx);
diff --git a/rts/c/opencl.h b/rts/c/opencl.h
--- a/rts/c/opencl.h
+++ b/rts/c/opencl.h
@@ -593,7 +593,8 @@
       max_value = 0; // No limit.
       default_value = ctx->cfg.default_threshold;
     } else {
-      panic(1, "Unknown size class for size '%s': %s\n", size_name, size_class);
+      // Bespoke sizes have no limit or default.
+      max_value = 0;
     }
     if (*size_value == 0) {
       *size_value = default_value;
diff --git a/rts/csharp/opencl.cs b/rts/csharp/opencl.cs
--- a/rts/csharp/opencl.cs
+++ b/rts/csharp/opencl.cs
@@ -402,6 +402,7 @@
    public int MaxTileSize;
    public int MaxThreshold;
    public int MaxLocalMemory;
+   public int MaxBespoke;
 
    public int LockstepWidth;
 }
@@ -807,6 +808,7 @@
     ctx.OpenCL.MaxTileSize = MaxTileSize; // No limit.
     ctx.OpenCL.MaxThreshold = ctx.OpenCL.MaxNumGroups; // No limit.
     ctx.OpenCL.MaxLocalMemory = MaxLocalMemory;
+    ctx.OpenCL.MaxBespoke = 0; // No limit.
 
     // Now we go through all the sizes, clamp them to the valid range,
     // or set them to the default.
@@ -829,7 +831,8 @@
             max_value = 0; // No limit.
             default_value = ctx.OpenCL.Cfg.DefaultThreshold;
         } else {
-            panic(1, "Unknown size class for size '{0}': {1}\n", size_name, size_class);
+            // Bespoke sizes have no limit or default.
+            max_value = 0;
         }
         if (size_value == 0) {
             ctx.OpenCL.Cfg.SizeValues[i] = default_value;
diff --git a/rts/python/__init__.py b/rts/python/__init__.py
deleted file mode 100644
--- a/rts/python/__init__.py
+++ /dev/null
diff --git a/rts/python/opencl.py b/rts/python/opencl.py
--- a/rts/python/opencl.py
+++ b/rts/python/opencl.py
@@ -175,7 +175,8 @@
             max_value = None
             default_value = default_threshold
         else:
-            raise Exception('Unknown size class for size \'{}\': {}'.format(k, v['class']))
+            # Bespoke sizes have no limit or default.
+            max_value = None
         if v['value'] == None:
             self.sizes[k] = default_value
         elif max_value != None and v['value'] > max_value:
diff --git a/src/Futhark/CLI/Dataset.hs b/src/Futhark/CLI/Dataset.hs
--- a/src/Futhark/CLI/Dataset.hs
+++ b/src/Futhark/CLI/Dataset.hs
@@ -147,7 +147,7 @@
   d' <- constantDim d
   ValueType ds t' <- toValueType t
   return $ ValueType (d':ds) t'
-  where constantDim (ConstDim k) = Right k
+  where constantDim (DimExpConst k _) = Right k
         constantDim _ = Left "Array has non-constant dimension declaration."
 toValueType (TEVar (QualName [] v) _)
   | Just t <- M.lookup v namesToPrimTypes = Right $ ValueType [] t
diff --git a/src/Futhark/CLI/Misc.hs b/src/Futhark/CLI/Misc.hs
--- a/src/Futhark/CLI/Misc.hs
+++ b/src/Futhark/CLI/Misc.hs
@@ -17,36 +17,25 @@
 
 import Futhark.Compiler
 import Futhark.Util.Options
-import Futhark.Pipeline
 import Futhark.Test
 
-runFutharkM' :: FutharkM () -> IO ()
-runFutharkM' m = do
-  res <- runFutharkM m NotVerbose
-  case res of
-    Left err -> do
-      dumpError newFutharkConfig err
-      exitWith $ ExitFailure 2
-    Right () -> return ()
-
 mainCheck :: String -> [String] -> IO ()
 mainCheck = mainWithOptions () [] "program" $ \args () ->
   case args of
-    [file] -> Just $ runFutharkM' $ check file
+    [file] -> Just $ do
+      (warnings, _, _) <- readProgramOrDie file
+      liftIO $ hPutStr stderr $ show warnings
     _ -> Nothing
-  where check file = do (warnings, _, _) <- readProgram file
-                        liftIO $ hPutStr stderr $ show warnings
 
 mainImports :: String -> [String] -> IO ()
 mainImports = mainWithOptions () [] "program" $ \args () ->
   case args of
-    [file] -> Just $ runFutharkM' $ findImports file
+    [file] -> Just $ do
+      (_, prog_imports, _) <- readProgramOrDie file
+      liftIO $ putStr $ unlines $ map (++ ".fut")
+        $ filter (\f -> not ("futlib/" `isPrefixOf` f))
+        $ map fst prog_imports
     _ -> Nothing
-  where findImports file = do
-          (_, prog_imports, _) <- readProgram file
-          liftIO $ putStr $ unlines $ map (++ ".fut")
-            $ filter (\f -> not ("futlib/" `isPrefixOf` f))
-            $ map fst prog_imports
 
 mainDataget :: String -> [String] -> IO ()
 mainDataget = mainWithOptions () [] "program dataset" $ \args () ->
diff --git a/src/Futhark/CLI/Query.hs b/src/Futhark/CLI/Query.hs
new file mode 100644
--- /dev/null
+++ b/src/Futhark/CLI/Query.hs
@@ -0,0 +1,36 @@
+module Futhark.CLI.Query (main) where
+
+import Text.Read (readMaybe)
+
+import Data.Loc
+import Futhark.Compiler
+import Futhark.Util.Options
+import Language.Futhark.Query
+import Language.Futhark.Syntax
+
+main :: String -> [String] -> IO ()
+main = mainWithOptions () [] "program line:col" $ \args () ->
+  case args of
+    [file, line, col] -> do
+      line' <- readMaybe line
+      col' <- readMaybe col
+      Just $ do
+        (_, imports, _) <- readProgramOrDie file
+        -- The 'offset' part of the Pos is not used and can be arbitrary.
+        case atPos imports $ Pos file line' col' 0 of
+          Nothing -> putStrLn "No information available."
+          Just (AtName qn def loc) -> do
+            putStrLn $ "Name: " ++ pretty qn
+            putStrLn $ "Position: " ++ locStr (srclocOf loc)
+            case def of
+              Nothing -> return ()
+              Just (BoundTerm t defloc) -> do
+                putStrLn $ "Type: " ++ pretty t
+                putStrLn $ "Definition: " ++ locStr (srclocOf defloc)
+              Just (BoundType defloc) ->
+                putStrLn $ "Definition: " ++ locStr (srclocOf defloc)
+              Just (BoundModule defloc) ->
+                putStrLn $ "Definition: " ++ locStr (srclocOf defloc)
+              Just (BoundModuleType defloc) ->
+                putStrLn $ "Definition: " ++ locStr (srclocOf defloc)
+    _ -> Nothing
diff --git a/src/Futhark/CodeGen/Backends/CCUDA.hs b/src/Futhark/CodeGen/Backends/CCUDA.hs
--- a/src/Futhark/CodeGen/Backends/CCUDA.hs
+++ b/src/Futhark/CodeGen/Backends/CCUDA.hs
@@ -223,6 +223,7 @@
     cudaSizeClass SizeNumGroups = "grid_size"
     cudaSizeClass SizeTile = "tile_size"
     cudaSizeClass SizeLocalMemory = "shared_memory"
+    cudaSizeClass (SizeBespoke x _) = pretty x
 callKernel (LaunchKernel name args num_blocks block_size) = do
   args_arr <- newVName "kernel_args"
   time_start <- newVName "time_start"
diff --git a/src/Futhark/CodeGen/Backends/CCUDA/Boilerplate.hs b/src/Futhark/CodeGen/Backends/CCUDA/Boilerplate.hs
--- a/src/Futhark/CodeGen/Backends/CCUDA/Boilerplate.hs
+++ b/src/Futhark/CodeGen/Backends/CCUDA/Boilerplate.hs
@@ -81,8 +81,10 @@
                               const char **nvrtc_opts;
                             };|])
 
-  let size_value_inits = map (\i -> [C.cstm|cfg->sizes[$int:i] = 0;|])
-                           [0..M.size sizes-1]
+  let size_value_inits = zipWith sizeInit [0..M.size sizes-1] (M.elems sizes)
+      sizeInit i size = [C.cstm|cfg->sizes[$int:i] = $int:val;|]
+         where val = case size of SizeBespoke _ x -> x
+                                  _               -> 0
   GC.publicDef_ "context_config_new" GC.InitDecl $ \s ->
     ([C.cedecl|struct $id:cfg* $id:s(void);|],
      [C.cedecl|struct $id:cfg* $id:s(void) {
diff --git a/src/Futhark/CodeGen/Backends/COpenCL/Boilerplate.hs b/src/Futhark/CodeGen/Backends/COpenCL/Boilerplate.hs
--- a/src/Futhark/CodeGen/Backends/COpenCL/Boilerplate.hs
+++ b/src/Futhark/CodeGen/Backends/COpenCL/Boilerplate.hs
@@ -65,7 +65,10 @@
                               const char **build_opts;
                             };|])
 
-  let size_value_inits = map (\i -> [C.cstm|cfg->sizes[$int:i] = 0;|]) [0..M.size sizes-1]
+  let size_value_inits = zipWith sizeInit [0..M.size sizes-1] (M.elems sizes)
+      sizeInit i size = [C.cstm|cfg->sizes[$int:i] = $int:val;|]
+         where val = case size of SizeBespoke _ x -> x
+                                  _               -> 0
   GC.publicDef_ "context_config_new" GC.InitDecl $ \s ->
     ([C.cedecl|struct $id:cfg* $id:s(void);|],
      [C.cedecl|struct $id:cfg* $id:s(void) {
diff --git a/src/Futhark/CodeGen/Backends/CSOpenCL.hs b/src/Futhark/CodeGen/Backends/CSOpenCL.hs
--- a/src/Futhark/CodeGen/Backends/CSOpenCL.hs
+++ b/src/Futhark/CodeGen/Backends/CSOpenCL.hs
@@ -133,6 +133,7 @@
                      Imp.SizeTile -> "MaxTileSize"
                      Imp.SizeThreshold{} -> "MaxThreshold"
                      Imp.SizeLocalMemory -> "MaxLocalMemory"
+                     Imp.SizeBespoke{} -> "MaxBespoke"
 
 callKernel (Imp.LaunchKernel name args num_workgroups workgroup_size) = do
   num_workgroups' <- mapM CS.compileExp num_workgroups
diff --git a/src/Futhark/CodeGen/Backends/CSOpenCL/Boilerplate.hs b/src/Futhark/CodeGen/Backends/CSOpenCL/Boilerplate.hs
--- a/src/Futhark/CodeGen/Backends/CSOpenCL/Boilerplate.hs
+++ b/src/Futhark/CodeGen/Backends/CSOpenCL/Boilerplate.hs
@@ -71,9 +71,12 @@
                          , (intArrayT, "Sizes")]
 
   let tmp_cfg = Var "tmp_cfg"
+      sizeInit (SizeBespoke _ x) = Integer $ toInteger x
+      sizeInit _ = Integer 0
+
   CS.stm $ CS.privateFunDef new_cfg (CustomT cfg) []
     [ Assign tmp_cfg $ CS.simpleInitClass cfg []
-    , Reassign (Field tmp_cfg "Sizes") (Collection "int[]" (replicate (M.size sizes) (Integer 0)))
+    , Reassign (Field tmp_cfg "Sizes") (Collection "int[]" (map sizeInit (M.elems sizes)))
     , Exp $ CS.simpleCall "OpenCLConfigInit" [ Out $ Field tmp_cfg "OpenCL", (Integer . toInteger) $ M.size sizes
                                              , Var "SizeNames", Var "SizeVars", Field tmp_cfg "Sizes", Var "SizeClasses" ]
     , Return tmp_cfg
diff --git a/src/Futhark/CodeGen/Backends/PyOpenCL/Boilerplate.hs b/src/Futhark/CodeGen/Backends/PyOpenCL/Boilerplate.hs
--- a/src/Futhark/CodeGen/Backends/PyOpenCL/Boilerplate.hs
+++ b/src/Futhark/CodeGen/Backends/PyOpenCL/Boilerplate.hs
@@ -52,7 +52,9 @@
   where f (size_name, size_class) =
           (String $ pretty size_name,
            Dict [(String "class", String $ pretty size_class),
-                 (String "value", None)])
+                 (String "value", defValue size_class)])
+        defValue (SizeBespoke _ x) = Integer $ toInteger x
+        defValue _ = None
 
 sizeHeuristicsToPython :: [SizeHeuristic] -> PyExp
 sizeHeuristicsToPython = List . map f
diff --git a/src/Futhark/CodeGen/ImpGen/Kernels/Base.hs b/src/Futhark/CodeGen/ImpGen/Kernels/Base.hs
--- a/src/Futhark/CodeGen/ImpGen/Kernels/Base.hs
+++ b/src/Futhark/CodeGen/ImpGen/Kernels/Base.hs
@@ -134,13 +134,16 @@
 -- multidimensional loops, use 'groupCoverSpace'.
 kernelLoop :: Imp.Exp -> Imp.Exp -> Imp.Exp
            -> (Imp.Exp -> InKernelGen ()) -> InKernelGen ()
-kernelLoop tid num_threads n f = do
-  -- Compute how many elements this thread is responsible for.
-  -- Formula: (n - tid) / num_threads (rounded up).
-  let elems_for_this = (n - tid) `quotRoundingUp` num_threads
+kernelLoop tid num_threads n f =
+  if n == num_threads then
+    f tid
+  else do
+    -- Compute how many elements this thread is responsible for.
+    -- Formula: (n - tid) / num_threads (rounded up).
+    let elems_for_this = (n - tid) `quotRoundingUp` num_threads
 
-  sFor "i" elems_for_this $ \i -> f $
-    i * num_threads + tid
+    sFor "i" elems_for_this $ \i -> f $
+      i * num_threads + tid
 
 -- | Assign iterations of a for-loop to threads in the workgroup.  The
 -- passed-in function is invoked with the (symbolic) iteration.  For
@@ -1037,8 +1040,12 @@
   let iterations = (required_groups - Imp.vi32 phys_group_id) `quotRoundingUp`
                    kernelNumGroups constants
 
-  sFor "i" iterations $ \i ->
+  sFor "i" iterations $ \i -> do
     m =<< dPrimV "virt_group_id" (Imp.vi32 phys_group_id + i * kernelNumGroups constants)
+    -- Make sure the virtual group is actually done before we let
+    -- another virtual group have its way with it.
+    sOp Imp.GlobalBarrier
+    sOp Imp.LocalBarrier
 
 sKernelThread, sKernelGroup :: String
                             -> Count NumGroups Imp.Exp -> Count GroupSize Imp.Exp
diff --git a/src/Futhark/CodeGen/ImpGen/Kernels/SegHist.hs b/src/Futhark/CodeGen/ImpGen/Kernels/SegHist.hs
--- a/src/Futhark/CodeGen/ImpGen/Kernels/SegHist.hs
+++ b/src/Futhark/CodeGen/ImpGen/Kernels/SegHist.hs
@@ -161,7 +161,7 @@
       --
       -- A fun solution would also be to use a simple hashing
       -- algorithm to ensure good distribution of locks.
-      let num_locks = 10000
+      let num_locks = 100151
           dims = map (toExp' int32) $
                  shapeDims (histShape (slugOp slug)) ++
                  [ Var (slugNumSubhistos slug)
@@ -172,90 +172,143 @@
       let l' = Locking locks 0 1 0 (pure . (`rem` fromIntegral num_locks) . flattenIndex dims)
       return (Just l', f l' (Space "global") dests)
 
-infoPrints :: Imp.Exp -> VName -> Imp.Exp -> ImpM lore op ()
-infoPrints hist_H hist_M hist_C = do
-  emit $ Imp.DebugPrint "Histogram size (H)" $ Just hist_H
-  emit $ Imp.DebugPrint "Multiplication degree (M)" $ Just $ Imp.vi32 hist_M
-  emit $ Imp.DebugPrint "Cooperation level (C)" $ Just hist_C
+-- | Some kernel bodies are not safe (or efficient) to execute
+-- multiple times.
+data Passage = MustBeSinglePass | MayBeMultiPass deriving (Eq, Ord)
 
-prepareIntermediateArraysGlobal :: Imp.Exp -> Imp.Exp -> [SegHistSlug]
+bodyPassage :: KernelBody ExplicitMemory -> Passage
+bodyPassage kbody
+  | mempty == consumedInKernelBody (aliasAnalyseKernelBody kbody) =
+      MayBeMultiPass
+  | otherwise =
+      MustBeSinglePass
+
+prepareIntermediateArraysGlobal :: Passage -> Imp.Exp -> Imp.Exp -> [SegHistSlug]
                                 -> CallKernelGen
-                                   [[Imp.Exp] -> InKernelGen ()]
-prepareIntermediateArraysGlobal hist_T hist_N =
-  fmap snd . mapAccumLM onOp Nothing
-  where
-    onOp l slug@(SegHistSlug op num_subhistos subhisto_info do_op) = do
-      hist_H <- toExp $ histWidth op
+                                   (Imp.Exp,
+                                    [[Imp.Exp] -> InKernelGen ()])
+prepareIntermediateArraysGlobal passage hist_T hist_N slugs = do
+  -- The paper formulae assume there is only one histogram, but in our
+  -- implementation there can be multiple that have been horisontally
+  -- fused.  We do a bit of trickery with summings and averages to
+  -- pretend there is really only one.  For the case of a single
+  -- histogram, the actual calculations should be the same as in the
+  -- paper.
 
-      hist_u <- dPrimVE "hist_u" $
-                case do_op of
-                  AtomicPrim{} -> 1
-                  _            -> 2
+  -- The sum of all Hs.
+  hist_H <- dPrimVE "hist_H" . sum =<< mapM (toExp . histWidth . slugOp) slugs
 
-      hist_RF <- toExp $ histRaceFactor op
+  hist_RF <- dPrimVE "hist_RF" $
+    sum (map (r64 . toExp' int32 . histRaceFactor . slugOp) slugs)
+    / r64 (genericLength slugs)
 
-      let hist_k_RF = 0.75
-          hist_L2 = 16384
-          hist_L2_ln_sz = 64
-          hist_F_L2 = 0.4
+  hist_el_size <- dPrimVE "hist_el_size" $ sum $ map slugElAvgSize slugs
 
-      let r64 = ConvOpExp (SIToFP Int32 Float64)
-          t64 = ConvOpExp (FPToSI Float64 Int32)
+  hist_C_max <- dPrimVE "hist_C_max" $
+    Imp.BinOpExp (SMin Int32) hist_T $ hist_H `quot` hist_k_ct_min
 
-      let hist_el_size =
-            case do_op of
-              AtomicLocking{} ->
-                unCount
-                (sum $ map (typeSize . (`arrayOfShape` histShape op)) $
-                 Prim int32 : lambdaReturnType (histOp op))
-                `quot` genericLength (lambdaReturnType (histOp op))
-              _ ->
-                unCount $ sum $
-                map (typeSize . (`arrayOfShape` histShape op)) $
-                lambdaReturnType (histOp op)
+  hist_M_min <- dPrimVE "hist_M_min" $
+    Imp.BinOpExp (SMax Int32) 1 $ hist_T `quot` hist_C_max
 
-      hist_RACE_exp <- dPrimVE "hist_RACE_exp" $
-        Imp.BinOpExp (FMax Float64) 1 $
-        (hist_k_RF * r64 hist_RF) /
-        Imp.BinOpExp (FMin Float64) 1 (r64 hist_L2_ln_sz  / r64 hist_el_size)
+  -- Querying L2 cache size is not reliable.  Instead we provide a
+  -- tunable knob with a hopefully sane default.
+  let hist_L2_def = 5632 * 1024
+  hist_L2 <- dPrim "L2_size" int32
+  -- Equivalent to F_L2*L2 in paper.
+  sOp $ Imp.GetSize hist_L2 (nameFromString $ pretty hist_L2) $
+    Imp.SizeBespoke (nameFromString "L2_for_histogram") hist_L2_def
 
-      -- Hardcode a single pass for now.
-      let hist_S = 1
+  let hist_L2_ln_sz = 16*4 -- L2 cache line size approximation
 
+  hist_RACE_exp <- dPrimVE "hist_RACE_exp" $
+    Imp.BinOpExp (FMax Float64) 1 $
+    (hist_k_RF * hist_RF) /
+    (hist_L2_ln_sz / r64 hist_el_size)
+
+  hist_S <- dPrim "hist_S" int32
+
+  -- For sparse histograms (H exceeds N) we only want a single chunk.
+  sIf (hist_N .<. hist_H)
+    (hist_S <-- 1) $
+    hist_S <--
+    case passage of
+      MayBeMultiPass ->
+        (hist_M_min * hist_H * hist_el_size) `quotRoundingUp`
+        t64 (hist_F_L2 * r64 (Imp.vi32 hist_L2) * hist_RACE_exp)
+      MustBeSinglePass ->
+        1
+
+  emit $ Imp.DebugPrint "Race expansion factor (RACE^exp)" $ Just hist_RACE_exp
+  emit $ Imp.DebugPrint "Number of chunks (S)" $ Just $ Imp.vi32 hist_S
+
+  histograms <- snd <$> mapAccumLM (onOp (Imp.vi32 hist_L2) hist_M_min (Imp.vi32 hist_S) hist_RACE_exp) Nothing slugs
+
+  return (Imp.vi32 hist_S, histograms)
+  where
+    hist_k_ct_min = 2 -- Chosen experimentally
+    hist_k_RF = 0.75 -- Chosen experimentally
+    hist_F_L2 = 0.4 -- Chosen experimentally
+
+    r64 = ConvOpExp (SIToFP Int32 Float64)
+    t64 = ConvOpExp (FPToSI Float64 Int32)
+
+    -- "Average element size" as computed by a formula that also takes
+    -- locking into account.
+    slugElAvgSize slug@(SegHistSlug op _ _ do_op) =
+      case do_op of
+        AtomicLocking{} ->
+          slugElSize slug `quot` (1+genericLength (lambdaReturnType (histOp op)))
+        _ ->
+          slugElSize slug `quot` 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{} ->
+          unCount
+          (sum $ map (typeSize . (`arrayOfShape` histShape op)) $
+           Prim int32 : lambdaReturnType (histOp op))
+        _ ->
+          unCount $ sum $
+          map (typeSize . (`arrayOfShape` histShape 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
+      hist_H <- toExp $ histWidth op
+
       hist_H_chk <- dPrimVE "hist_H_chk" $
                     hist_H `quotRoundingUp` hist_S
 
+      emit $ Imp.DebugPrint "Chunk size (H_chk)" $ Just hist_H_chk
 
       hist_k_max <- dPrimVE "hist_k_max" $
-        Imp.BinOpExp (SMin Int32)
-        (t64 (hist_F_L2 * hist_L2 * hist_RACE_exp)) hist_N
-        `quotRoundingUp` hist_T
-
-      hist_C <- dPrimVE "hist_C" $
-                Imp.BinOpExp (SMin Int32) hist_T $
-                (hist_u * hist_H_chk) `quotRoundingUp` hist_k_max
-
-      -- Minimal sequential chunking factor.
-      let q_small = 2
-      work_asymp_M_max <- dPrimVE "work_asymp_M_max" $
-                          hist_N `quot` (q_small * hist_H)
+        Imp.BinOpExp (FMin Float64)
+        (hist_F_L2 * (r64 hist_L2 / r64 (slugElSize slug)) * hist_RACE_exp)
+        (r64 hist_N)
+        / r64 hist_T
 
-      let glb_k_min = 2
-      coop_min <- dPrimVE "coop_min" $
-                  Imp.BinOpExp (SMin Int32) hist_T $ hist_H `quot` glb_k_min
+      hist_u <- dPrimVE "hist_u" $
+                case do_op of
+                  AtomicPrim{} -> 2
+                  _            -> 1
 
-      hist_M_min <- dPrimVE "hist_M_min" $
-                    Imp.BinOpExp (SMax Int32) 1 $
-                    Imp.BinOpExp (SMin Int32) work_asymp_M_max $
-                    hist_T `quot` coop_min
+      hist_C <- dPrimVE "hist_C" $
+                Imp.BinOpExp (FMin Float64) (r64 hist_T) $
+                r64 (hist_u * hist_H_chk) / hist_k_max
 
       -- Number of subhistograms per result histogram.
-      hist_M <- dPrimV "hist_M" $
-                Imp.BinOpExp (SMin Int32) hist_M_min $
-                Imp.BinOpExp (SMax Int32) 1 $ hist_T `quot` hist_C
+      hist_M <- dPrimVE "hist_M" $
+                Imp.BinOpExp (SMax Int32) hist_M_min $
+                t64 $ r64 hist_T / hist_C
 
+      emit $ Imp.DebugPrint "Elements/thread in L2 cache (k_max)" $ Just hist_k_max
+      emit $ Imp.DebugPrint "Multiplication degree (M)" $ Just hist_M
+      emit $ Imp.DebugPrint "Cooperation level (C)" $ Just hist_C
+
       -- num_subhistos is the variable we use to communicate back.
-      num_subhistos <-- Imp.vi32 hist_M
+      num_subhistos <-- hist_M
 
       -- Initialise sub-histograms.
       --
@@ -276,7 +329,7 @@
 
             multiHistoCase = subhistosAlloc info
 
-        sIf (Imp.var hist_M int32 .==. 1) unitHistoCase multiHistoCase
+        sIf (hist_M .==. 1) unitHistoCase multiHistoCase
 
         return $ subhistosArray info
 
@@ -284,27 +337,27 @@
 
       return (l', do_op')
 
-histKernelGlobal :: [PatElem ExplicitMemory]
-                 -> Count NumGroups SubExp -> Count GroupSize SubExp
-                 -> SegSpace
-                 -> [SegHistSlug]
-                 -> KernelBody ExplicitMemory
-                 -> CallKernelGen ()
-histKernelGlobal map_pes num_groups group_size space slugs kbody = do
-  num_groups' <- traverse toExp num_groups
-  group_size' <- traverse toExp group_size
+histKernelGlobalPass :: [PatElem ExplicitMemory]
+                     -> Count NumGroups Imp.Exp
+                     -> Count GroupSize Imp.Exp
+                     -> SegSpace
+                     -> [SegHistSlug]
+                     -> KernelBody ExplicitMemory
+                     -> [[Imp.Exp] -> ImpM ExplicitMemory Imp.KernelOp ()]
+                     -> Imp.Exp
+                     -> Imp.Exp
+                     -> ImpM ExplicitMemory Imp.HostOp ()
+histKernelGlobalPass map_pes num_groups group_size space slugs kbody histograms hist_S chk_i = do
+
   let (space_is, space_sizes) = unzip $ unSegSpace space
       space_sizes_64 = map (i32Toi64 . toExp' int32) space_sizes
       total_w_64 = product space_sizes_64
-      num_threads = unCount num_groups' * unCount group_size'
 
-  emit $ Imp.DebugPrint "## Using global memory" Nothing
-
-  histograms <-
-    prepareIntermediateArraysGlobal num_threads
-    (toExp' int32 $ last space_sizes) slugs
+  hist_H_chks <- forM (map (histWidth . slugOp) slugs) $ \w -> do
+    w' <- toExp w
+    dPrimVE "hist_H_chk" $ w' `quotRoundingUp` hist_S
 
-  sKernelThread "seghist_global" num_groups' group_size' (segFlat space) $ \constants -> do
+  sKernelThread "seghist_global" num_groups group_size (segFlat space) $ \constants -> do
     -- Compute subhistogram index for each thread, per histogram.
     subhisto_inds <- forM slugs $ \slug ->
       dPrimVE "subhisto_ind" $
@@ -315,7 +368,8 @@
     -- calculation is done with 64-bit integers to avoid overflow,
     -- but the final unflattened segment indexes are 32 bit.
     let gtid = i32Toi64 $ kernelGlobalThreadId constants
-    kernelLoop gtid (i32Toi64 num_threads) total_w_64 $ \offset -> do
+        num_threads = i32Toi64 $ kernelNumThreads constants
+    kernelLoop gtid num_threads total_w_64 $ \offset -> do
 
       -- Construct segment indices.
       let setIndex v e = do dPrim_ v int32
@@ -342,24 +396,50 @@
             perOp = chunks $ map (length . histDest . slugOp) slugs
 
         sComment "perform atomic updates" $
-          forM_ (zip5 (map slugOp slugs) histograms buckets (perOp vs) subhisto_inds) $
+          forM_ (zip6 (map slugOp slugs) histograms buckets (perOp vs) subhisto_inds hist_H_chks) $
           \(HistOp dest_w _ _ _ shape lam,
-            do_op, bucket, vs', subhisto_ind) -> do
+            do_op, bucket, vs', subhisto_ind, hist_H_chk) -> do
 
-            let bucket' = toExp' int32 $ kernelResultSubExp bucket
+            let chk_beg = chk_i * hist_H_chk
+                bucket' = toExp' int32 $ kernelResultSubExp bucket
                 dest_w' = toExp' int32 dest_w
-                bucket_in_bounds = 0 .<=. bucket' .&&. bucket' .<. dest_w'
-                bucket_is = map Imp.vi32 (init space_is) ++
-                            [subhisto_ind, bucket']
+                bucket_in_bounds = chk_beg .<=. bucket' .&&.
+                                   bucket' .<. (chk_beg + hist_H_chk) .&&.
+                                   bucket' .<. dest_w'
                 vs_params = takeLast (length vs') $ lambdaParams lam
 
             sWhen bucket_in_bounds $ do
+              let bucket_is = map Imp.vi32 (init space_is) ++
+                              [subhisto_ind, bucket']
               dLParams $ lambdaParams lam
               sLoopNest shape $ \is -> do
                 forM_ (zip vs_params vs') $ \(p, res) ->
                   copyDWIM (paramName p) [] (kernelResultSubExp res) is
                 do_op (bucket_is ++ is)
 
+
+histKernelGlobal :: [PatElem ExplicitMemory]
+                 -> Count NumGroups SubExp -> Count GroupSize SubExp
+                 -> SegSpace
+                 -> [SegHistSlug]
+                 -> KernelBody ExplicitMemory
+                 -> CallKernelGen ()
+histKernelGlobal map_pes num_groups group_size space slugs kbody = do
+  num_groups' <- traverse toExp num_groups
+  group_size' <- traverse toExp group_size
+  let (_space_is, space_sizes) = unzip $ unSegSpace space
+      num_threads = unCount num_groups' * unCount group_size'
+
+  emit $ Imp.DebugPrint "## Using global memory" Nothing
+
+  (hist_S, histograms) <-
+    prepareIntermediateArraysGlobal (bodyPassage kbody)
+    num_threads (toExp' int32 $ last space_sizes) slugs
+
+  sFor "chk_i" hist_S $ \chk_i ->
+    histKernelGlobalPass map_pes num_groups' group_size' space slugs kbody
+    histograms hist_S chk_i
+
 prepareIntermediateArraysLocal :: VName
                                -> Count NumGroups SubExp
                                -> SegSpace -> [SegHistSlug]
@@ -601,25 +681,28 @@
               copyDWIM global_dest global_is (Var $ paramName xp) []
 
 localMemoryCase :: [PatElem ExplicitMemory]
-                -> Count NumGroups SubExp -> Count GroupSize SubExp
+                -> Imp.Exp
                 -> SegSpace
                 -> Imp.Exp -> Imp.Exp -> Imp.Exp -> Imp.Exp
                 -> [SegHistSlug]
                 -> KernelBody ExplicitMemory
                 -> CallKernelGen (Imp.Exp, ImpM ExplicitMemory Imp.HostOp ())
-localMemoryCase map_pes num_groups group_size space hist_H hist_el_size hist_N hist_RF slugs kbody = do
-  num_groups' <- traverse toExp num_groups
-  group_size' <- traverse toExp group_size
-
-  -- Maximum group size (or actual, in this case).
-  let hist_B = unCount group_size'
-      space_sizes = segSpaceDims space
+localMemoryCase map_pes hist_T space hist_H hist_el_size hist_N hist_RF slugs kbody = do
+  let space_sizes = segSpaceDims space
       segment_dims = init space_sizes
       segmented = not $ null segment_dims
 
   hist_L <- dPrim "hist_L" int32
   sOp $ Imp.GetSizeMax hist_L Imp.SizeLocalMemory
 
+  max_group_size <- dPrim "max_group_size" int32
+  sOp $ Imp.GetSizeMax max_group_size Imp.SizeGroup
+  let group_size = Imp.Count $ Var max_group_size
+  num_groups <- fmap (Imp.Count . Var) $ dPrimV "num_groups" $
+                hist_T `quotRoundingUp` toExp' int32 (unCount group_size)
+  let num_groups' = toExp' int32 <$> num_groups
+      group_size' = toExp' int32 <$> group_size
+
   let r64 = ConvOpExp (SIToFP Int32 Float64)
       t64 = ConvOpExp (FPToSI Float64 Int32)
       f64ceil x = t64 $ FunExp "round64" [x] $ FloatType Float64
@@ -641,6 +724,7 @@
               Imp.BinOpExp (FMin Float64) (r64 hist_RF) $
               r64 hist_W * Imp.BinOpExp (FPow Float64) (r64 hist_RF / r64 hist_W) (1.0 / 3.0)
 
+  let hist_B = unCount group_size'
   hist_f' <- dPrimVE "hist_f_prime" $
              (r64 hist_B * hist_RFC) /
              (hist_m * hist_m * r64 hist_H)
@@ -687,8 +771,6 @@
   work_asymp_M_max <-
     if segmented then do
 
-      hist_T <- dPrimVE "hist_T" $ unCount num_groups' * unCount group_size'
-
       hist_T_hist_min <- dPrimVE "hist_T_hist_min" $
                          Imp.BinOpExp (SMin Int32) (hist_Nin * hist_Nout) hist_T
                          `quotRoundingUp`
@@ -742,8 +824,11 @@
 
       run = do
         emit $ Imp.DebugPrint "## Using local memory" Nothing
-        infoPrints hist_H hist_M hist_C
-        histKernelLocal hist_M groups_per_segment map_pes num_groups group_size space slugs kbody
+        emit $ Imp.DebugPrint "Histogram size (H)" $ Just hist_H
+        emit $ Imp.DebugPrint "Multiplication degree (M)" $ Just $ Imp.vi32 hist_M
+        emit $ Imp.DebugPrint "Cooperation level (C)" $ Just hist_C
+        histKernelLocal hist_M groups_per_segment map_pes
+          num_groups group_size space slugs kbody
 
   return (pick_local, run)
 
@@ -758,6 +843,7 @@
                -> KernelBody ExplicitMemory
                -> CallKernelGen ()
 compileSegHist (Pattern _ pes) num_groups group_size space ops kbody = do
+  num_groups' <- traverse toExp num_groups
   group_size' <- traverse toExp group_size
 
   dims <- mapM toExp $ segSpaceDims space
@@ -797,7 +883,9 @@
   -- Check for emptyness to avoid division-by-zero.
   sUnless (seg_h .==. 0) $ do
 
+    let hist_T = unCount num_groups' * unCount group_size'
     emit $ Imp.DebugPrint "\n# SegHist" Nothing
+    emit $ Imp.DebugPrint "Number of threads (T)" $ Just hist_T
     emit $ Imp.DebugPrint "Memory per set of subhistograms per segment" $ Just h
     emit $ Imp.DebugPrint "Memory per set of subhistograms times segments" $ Just seg_h
     emit $ Imp.DebugPrint "Input elements per histogram (N)" $ Just hist_N
@@ -807,7 +895,7 @@
     emit $ Imp.DebugPrint "Race factor (RF)" $ Just hist_RF
 
     (use_local_memory, run_in_local_memory) <-
-      localMemoryCase map_pes num_groups group_size space hist_H hist_el_size hist_N hist_RF slugs kbody
+      localMemoryCase map_pes hist_T space hist_H hist_el_size hist_N hist_RF slugs kbody
 
     sIf use_local_memory run_in_local_memory $
       histKernelGlobal map_pes num_groups group_size space slugs kbody
diff --git a/src/Futhark/Compiler.hs b/src/Futhark/Compiler.hs
--- a/src/Futhark/Compiler.hs
+++ b/src/Futhark/Compiler.hs
@@ -12,6 +12,7 @@
        , module Futhark.Compiler.Program
        , readProgram
        , readLibrary
+       , readProgramOrDie
        )
 where
 
@@ -137,3 +138,13 @@
 readLibrary :: (MonadError CompilerError m, MonadIO m) =>
                [FilePath] -> m (Warnings, Imports, VNameSource)
 readLibrary = readLibraryWithBasis emptyBasis
+
+-- | Not verbose, and terminates process on error.
+readProgramOrDie :: MonadIO m => FilePath -> m (Warnings, Imports, VNameSource)
+readProgramOrDie file = liftIO $ do
+  res <- runFutharkM (readProgram file) NotVerbose
+  case res of
+    Left err -> do
+      dumpError newFutharkConfig err
+      exitWith $ ExitFailure 2
+    Right res' -> return res'
diff --git a/src/Futhark/Doc/Generator.hs b/src/Futhark/Doc/Generator.hs
--- a/src/Futhark/Doc/Generator.hs
+++ b/src/Futhark/Doc/Generator.hs
@@ -428,7 +428,7 @@
 
 synopsisSigExp :: SigExpBase Info VName -> DocM Html
 synopsisSigExp e = case e of
-  SigVar v _ -> qualNameHtml v
+  SigVar v _ _ -> qualNameHtml v
   SigParens e' _ -> parens <$> synopsisSigExp e'
   SigSpecs ss _ -> braces . (H.table ! A.class_ "specs") . mconcat <$> mapM synopsisSpec ss
   SigWith s (TypeRef v ps t _) _ -> do
@@ -495,7 +495,7 @@
   TEUnique t _  -> ("*"<>) <$> typeExpHtml t
   TEArray at d _ -> do
     at' <- typeExpHtml at
-    d' <- dimDeclHtml d
+    d' <- dimExpHtml d
     return $ brackets d' <> at'
   TETuple ts _ -> parens . commas <$> mapM typeExpHtml ts
   TERecord fs _ -> braces . commas <$> mapM ppField fs
@@ -567,8 +567,13 @@
 dimDeclHtml (NamedDim v) = qualNameHtml v
 dimDeclHtml (ConstDim n) = return $ toHtml (show n)
 
+dimExpHtml :: DimExp VName -> DocM Html
+dimExpHtml DimExpAny = return mempty
+dimExpHtml (DimExpNamed v _) = qualNameHtml v
+dimExpHtml (DimExpConst n _) = return $ toHtml (show n)
+
 typeArgExpHtml :: TypeArgExp VName -> DocM Html
-typeArgExpHtml (TypeArgExpDim d _) = dimDeclHtml d
+typeArgExpHtml (TypeArgExpDim d _) = dimExpHtml d
 typeArgExpHtml (TypeArgExpType d) = typeExpHtml d
 
 typeParamHtml :: TypeParam -> Html
diff --git a/src/Futhark/Internalise.hs b/src/Futhark/Internalise.hs
--- a/src/Futhark/Internalise.hs
+++ b/src/Futhark/Internalise.hs
@@ -508,6 +508,7 @@
         merge_ts = existentialiseExtTypes merge_names $
                    staticShapes $ map (I.paramType . fst) merge
     loopbody'' <- localScope (scopeOfFParams $ map fst merge) $
+                  inScopeOf form' $
                   ensureResultExtShapeNoCtx asserting
                   "shape of loop result does not match shapes in loop parameters"
                   loc merge_ts loopbody'
@@ -1758,7 +1759,7 @@
   return [ErrorString $ pretty qn]
 typeExpForError cm (E.TEUnique te _) = ("*":) <$> typeExpForError cm te
 typeExpForError cm (E.TEArray te d _) = do
-  d' <- dimDeclForError cm d
+  d' <- dimExpForError cm d
   te' <- typeExpForError cm te
   return $ ["[", d', "]"] ++ te'
 typeExpForError cm (E.TETuple tes _) = do
@@ -1775,7 +1776,7 @@
 typeExpForError cm (E.TEApply t arg _) = do
   t' <- typeExpForError cm t
   arg' <- case arg of TypeArgExpType argt -> typeExpForError cm argt
-                      TypeArgExpDim d _   -> pure <$> dimDeclForError cm d
+                      TypeArgExpDim d _   -> pure <$> dimExpForError cm d
   return $ t' ++ [" "] ++ arg'
 typeExpForError cm (E.TESum cs _) = do
   cs' <- mapM (onClause . snd) cs
@@ -1784,8 +1785,8 @@
           c' <- mapM (typeExpForError cm) c
           return $ intercalate [" "] c'
 
-dimDeclForError :: ConstParams -> E.DimDecl VName -> InternaliseM (ErrorMsgPart SubExp)
-dimDeclForError cm (NamedDim d) = do
+dimExpForError :: ConstParams -> E.DimExp VName -> InternaliseM (ErrorMsgPart SubExp)
+dimExpForError cm (DimExpNamed d _) = do
   substs <- asks $ M.lookup (E.qualLeaf d) . envSubsts
   let fname = nameFromString $ pretty (E.qualLeaf d) ++ "f"
   d' <- case (substs, lookup fname cm) of
@@ -1793,9 +1794,9 @@
           (_, Just v)   -> return $ I.Var v
           _             -> return $ I.Var $ E.qualLeaf d
   return $ ErrorInt32 d'
-dimDeclForError _ (ConstDim d) =
+dimExpForError _ (DimExpConst d _) =
   return $ ErrorString $ pretty d
-dimDeclForError _ AnyDim = return ""
+dimExpForError _ DimExpAny = return ""
 
 -- A smart constructor that compacts neighbouring literals for easier
 -- reading in the IR.
diff --git a/src/Futhark/Internalise/Defunctionalise.hs b/src/Futhark/Internalise/Defunctionalise.hs
--- a/src/Futhark/Internalise/Defunctionalise.hs
+++ b/src/Futhark/Internalise/Defunctionalise.hs
@@ -531,8 +531,17 @@
     -- Propagate the 'IntrinsicsSV' until we reach the outermost application,
     -- where we construct a dynamic static value with the appropriate type.
     IntrinsicSV
-      | depth == 0 -> return (e', Dynamic $ typeOf e)
-      | otherwise  -> return (e', IntrinsicSV)
+      | depth == 0 ->
+          -- If the intrinsic is fully applied, then we are done.
+          -- Otherwise we need to eta-expand it and recursively
+          -- defunctionalise. XXX: might it be better to simply
+          -- eta-expand immediately any time we encounter a
+          -- non-fully-applied intrinsic?
+          if null argtypes
+            then return (e', Dynamic $ typeOf e)
+            else do (pats, body, tp) <- etaExpand e'
+                    defuncExp $ Lambda pats body Nothing (Info (mempty, tp)) noLoc
+      | otherwise -> return (e', IntrinsicSV)
 
     _ -> error $ "Application of an expression that is neither a static lambda "
               ++ "nor a dynamic function, but has static value: " ++ show sv1
diff --git a/src/Futhark/Internalise/Defunctorise.hs b/src/Futhark/Internalise/Defunctorise.hs
--- a/src/Futhark/Internalise/Defunctorise.hs
+++ b/src/Futhark/Internalise/Defunctorise.hs
@@ -213,7 +213,7 @@
         onExp scope e =
           -- One expression is tricky, because it interacts with scoping rules.
           case e of
-            QualParens mn e' _ ->
+            QualParens (mn, _) e' _ ->
               case lookupMod' mn scope of
                 Left err -> error err
                 Right mod ->
diff --git a/src/Futhark/Optimise/DoubleBuffer.hs b/src/Futhark/Optimise/DoubleBuffer.hs
--- a/src/Futhark/Optimise/DoubleBuffer.hs
+++ b/src/Futhark/Optimise/DoubleBuffer.hs
@@ -221,7 +221,14 @@
               bound = MemArray bt shape NoUniqueness $ ArrayIn mem v_ixfun
           tell [Let (Pattern [] [PatElem v_copy bound]) (defAux ()) $
                 BasicOp $ Copy v]
-          return (f, Var v_copy)
+          -- It is important that we treat this as a consumption, to
+          -- avoid the Copy from being hoisted out of any enclosing
+          -- loops.  Since we re-use (=overwrite) memory in the loop,
+          -- the copy is critical for initialisation.  See issue #816.
+          let uniqueMemInfo (MemArray pt pshape _ ret) =
+                MemArray pt pshape Unique ret
+              uniqueMemInfo info = info
+          return (uniqueMemInfo <$> f, Var v_copy)
         allocation (f, se) _ =
           return (f, se)
 
diff --git a/src/Futhark/Optimise/Simplify/Rules.hs b/src/Futhark/Optimise/Simplify/Rules.hs
--- a/src/Futhark/Optimise/Simplify/Rules.hs
+++ b/src/Futhark/Optimise/Simplify/Rules.hs
@@ -625,8 +625,10 @@
           Just $ do
             i_offset' <- asIntS to_it i_offset
             i_stride' <- asIntS to_it i_stride
-            i_offset'' <- letSubExp "iota_offset" $
-                          BasicOp $ BinOp (Add Int32) x i_offset'
+            i_offset'' <- letSubExp "iota_offset" <=< toExp $
+                          primExpFromSubExp (IntType to_it) x +
+                          primExpFromSubExp (IntType to_it) s *
+                          primExpFromSubExp (IntType to_it) i_offset'
             i_stride'' <- letSubExp "iota_offset" $
                           BasicOp $ BinOp (Mul Int32) s i_stride'
             fmap (SubExpResult cs) $ letSubExp "slice_iota" $
diff --git a/src/Futhark/Representation/Kernels/Kernel.hs b/src/Futhark/Representation/Kernels/Kernel.hs
--- a/src/Futhark/Representation/Kernels/Kernel.hs
+++ b/src/Futhark/Representation/Kernels/Kernel.hs
@@ -13,6 +13,8 @@
        , SegRedOp(..)
        , segRedResults
        , KernelBody(..)
+       , aliasAnalyseKernelBody
+       , consumedInKernelBody
        , KernelResult(..)
        , kernelResultSubExp
        , SplitOrdering(..)
diff --git a/src/Futhark/Representation/Kernels/Sizes.hs b/src/Futhark/Representation/Kernels/Sizes.hs
--- a/src/Futhark/Representation/Kernels/Sizes.hs
+++ b/src/Futhark/Representation/Kernels/Sizes.hs
@@ -7,6 +7,7 @@
   )
   where
 
+import Data.Int (Int32)
 import Data.Traversable
 
 import Futhark.Util.Pretty
@@ -27,6 +28,8 @@
                | SizeLocalMemory
                -- ^ Likely not useful on its own, but querying the
                -- maximum can be handy.
+               | SizeBespoke Name Int32
+               -- ^ A bespoke size with a default.
                deriving (Eq, Ord, Show)
 
 instance Pretty SizeClass where
@@ -37,6 +40,7 @@
   ppr SizeNumGroups = text "num_groups"
   ppr SizeTile = text "tile_size"
   ppr SizeLocalMemory = text "local_memory"
+  ppr (SizeBespoke k _) = ppr k
 
 -- | A wrapper supporting a phantom type for indicating what we are counting.
 newtype Count u e = Count { unCount :: e }
diff --git a/src/Language/Futhark/Attributes.hs b/src/Language/Futhark/Attributes.hs
--- a/src/Language/Futhark/Attributes.hs
+++ b/src/Language/Futhark/Attributes.hs
@@ -833,7 +833,7 @@
 
         onModParam = onSigExp . modParamType
 
-        onSigExp (SigVar v _) = S.singleton $ qualLeaf v
+        onSigExp (SigVar v _ _) = S.singleton $ qualLeaf v
         onSigExp (SigParens e _) = onSigExp e
         onSigExp SigSpecs{} = mempty
         onSigExp (SigWith e _ _) = onSigExp e
diff --git a/src/Language/Futhark/Parser/Parser.y b/src/Language/Futhark/Parser/Parser.y
--- a/src/Language/Futhark/Parser/Parser.y
+++ b/src/Language/Futhark/Parser/Parser.y
@@ -222,7 +222,7 @@
 ;
 
 SigExp :: { UncheckedSigExp }
-        : QualName            { let (v, loc) = $1 in SigVar v loc }
+        : QualName            { let (v, loc) = $1 in SigVar v NoInfo loc }
         | '{' Specs '}'  { SigSpecs $2 (srcspan $1 $>) }
         | SigExp with TypeRef { SigWith $1 $3 (srcspan $1 $>) }
         | '(' SigExp ')'      { SigParens $2 (srcspan $1 $>) }
@@ -267,7 +267,7 @@
             | '{' Decs '}' { ModDecs $2 (srcspan $1 $>) }
 
 SimpleSigExp :: { UncheckedSigExp }
-             : QualName            { let (v, loc) = $1 in SigVar v loc }
+             : QualName            { let (v, loc) = $1 in SigVar v NoInfo loc }
              | '(' SigExp ')'      { $2 }
 
 ModBind :: { ModBindBase NoInfo Name }
@@ -423,17 +423,17 @@
 TypeExpTerm :: { UncheckedTypeExp }
          : '*' TypeExpTerm
            { TEUnique $2 (srcspan $1 $>) }
-         | '[' DimDecl ']' TypeExpTerm %prec indexprec
-           { TEArray $4 (fst $2) (srcspan $1 $>) }
+         | '[' DimExp ']' TypeExpTerm %prec indexprec
+           { TEArray $4 $2 (srcspan $1 $>) }
          | '['  ']' TypeExpTerm %prec indexprec
-           { TEArray $3 AnyDim (srcspan $1 $>) }
+           { TEArray $3 DimExpAny (srcspan $1 $>) }
          | TypeExpApply %prec sumprec { $1 }
 
          -- Errors
-         | '[' DimDecl ']' %prec bottom
+         | '[' DimExp ']' %prec bottom
            {% parseErrorAt (srcspan $1 $>) $ Just $
                 unlines ["missing array row type.",
-                         "Did you mean []"  ++ pretty (fst $2) ++ "?"]
+                         "Did you mean []"  ++ pretty $2 ++ "?"]
            }
 
 SumType :: { UncheckedTypeExp }
@@ -455,12 +455,12 @@
 TypeExpApply :: { UncheckedTypeExp }
               : TypeExpApply TypeArg
                 { TEApply $1 $2 (srcspan $1 $>) }
-              | 'id[' DimDecl ']'
+              | 'id[' DimExp ']'
                 { let L loc (INDEXING v) = $1
-                  in TEApply (TEVar (qualName v) loc) (TypeArgExpDim (fst $2) loc) (srcspan $1 $>) }
-              | 'qid[' DimDecl ']'
+                  in TEApply (TEVar (qualName v) loc) (TypeArgExpDim $2 loc) (srcspan $1 $>) }
+              | 'qid[' DimExp ']'
                 { let L loc (QUALINDEXING qs v) = $1
-                  in TEApply (TEVar (QualName qs v) loc) (TypeArgExpDim (fst $2) loc) (srcspan $1 $>) }
+                  in TEApply (TEVar (QualName qs v) loc) (TypeArgExpDim $2 loc) (srcspan $1 $>) }
               | TypeExpAtom
                 { $1 }
 
@@ -477,8 +477,8 @@
         : constructor { let L _ (CONSTRUCTOR c) = $1 in (c, srclocOf $1) }
 
 TypeArg :: { TypeArgExp Name }
-         : '[' DimDecl ']' { TypeArgExpDim (fst $2) (srcspan $1 $>) }
-         | '[' ']'         { TypeArgExpDim AnyDim (srcspan $1 $>) }
+         : '[' DimExp ']' { TypeArgExpDim $2 (srcspan $1 $>) }
+         | '[' ']'         { TypeArgExpDim DimExpAny (srcspan $1 $>) }
          | TypeExpAtom     { TypeArgExpType $1 }
 
 FieldType :: { (Name, UncheckedTypeExp) }
@@ -492,12 +492,12 @@
             : TypeExp                { [$1] }
             | TypeExp ',' TupleTypes { $1 : $3 }
 
-DimDecl :: { (DimDecl Name, SrcLoc) }
+DimExp :: { DimExp Name }
         : QualName
-          { (NamedDim (fst $1), snd $1) }
+          { DimExpNamed (fst $1) (snd $1) }
         | intlit
           { let L loc (INTLIT n) = $1
-            in (ConstDim (fromIntegral n), loc) }
+            in DimExpConst (fromIntegral n) loc }
 
 FunParam :: { PatternBase NoInfo Name }
 FunParam : InnerPattern { $1 }
@@ -628,15 +628,16 @@
      | '['       ']'                  { ArrayLit [] NoInfo (srcspan $1 $>) }
 
      | QualVarSlice FieldAccesses
-       { let (v,slice,loc) = $1
-         in foldl (\x (y, _) -> Project y x NoInfo (srclocOf x))
-                  (Index (Var v NoInfo loc) slice NoInfo loc)
+       { let ((v, vloc),slice,loc) = $1
+         in foldl (\x (y, _) -> Project y x NoInfo (srcspan x (srclocOf x)))
+                  (Index (Var v NoInfo vloc) slice NoInfo (srcspan vloc loc))
                   $2 }
      | QualName
        { Var (fst $1) NoInfo (snd $1) }
      | '{' Fields '}' { RecordLit $2 (srcspan $1 $>) }
      | 'qid.(' Exp ')'
-       { let L loc (QUALPAREN qs name) = $1 in QualParens (QualName qs name) $2 loc }
+       { let L loc (QUALPAREN qs name) = $1 in
+         QualParens (QualName qs name, loc) $2 (srcspan $1 $>) }
 
      -- Operator sections.
      | '(' UnOp ')'
@@ -717,7 +718,7 @@
          in LetFun name ($3, fst $4 : snd $4, (fmap declaredType $5), NoInfo, $7) $8 (srcspan $1 $>) }
 
      | let VarSlice '=' Exp LetBody
-                      { let (v,slice,loc) = $2; ident = Ident v NoInfo loc
+                      { let ((v,_),slice,loc) = $2; ident = Ident v NoInfo loc
                         in LetWith ident ident slice $4 $5 NoInfo (srcspan $1 $>) }
 
 LetBody :: { UncheckedExp }
@@ -797,16 +798,17 @@
          | while Exp
            { While $2 }
 
-VarSlice :: { (Name, [UncheckedDimIndex], SrcLoc) }
+VarSlice :: { ((Name, SrcLoc), [UncheckedDimIndex], SrcLoc) }
           : 'id[' DimIndices ']'
-              { let L _ (INDEXING v) = $1
-                in (v, $2, srcspan $1 $>) }
+              { let L vloc (INDEXING v) = $1
+                in ((v, vloc), $2, srcspan $1 $>) }
 
-QualVarSlice :: { (QualName Name, [UncheckedDimIndex], SrcLoc) }
+QualVarSlice :: { ((QualName Name, SrcLoc), [UncheckedDimIndex], SrcLoc) }
               : VarSlice
-                { let (x, y, z) = $1 in (qualName x, y, z) }
+                { let ((v, vloc), y, loc) = $1 in ((qualName v, vloc), y, loc) }
               | 'qid[' DimIndices ']'
-                { let L _ (QUALINDEXING qs v) = $1 in (QualName qs v, $2, srcspan $1 $>) }
+                { let L vloc (QUALINDEXING qs v) = $1
+                  in ((QualName qs v, vloc), $2, srcspan $1 $>) }
 
 DimIndex :: { UncheckedDimIndex }
          : Exp2                   { DimFix $1 }
diff --git a/src/Language/Futhark/Pretty.hs b/src/Language/Futhark/Pretty.hs
--- a/src/Language/Futhark/Pretty.hs
+++ b/src/Language/Futhark/Pretty.hs
@@ -103,6 +103,10 @@
   ppr (NamedDim v) = ppr v
   ppr (ConstDim n) = ppr n
 
+instance IsName vn => Pretty (DimExp vn) where
+  ppr DimExpAny         = mempty
+  ppr (DimExpNamed v _) = ppr v
+  ppr (DimExpConst n _) = ppr n
 
 instance IsName vn => Pretty (ShapeDecl (DimDecl vn)) where
   ppr (ShapeDecl ds) = mconcat (map (brackets . ppr) ds)
@@ -117,8 +121,8 @@
   ppr = pprPrec 0
   pprPrec _ (Prim et) = ppr et
   pprPrec p (TypeVar _ u et targs) =
-    parensIf (not (null targs) && p > 0) $
-    ppr u <> ppr (qualNameFromTypeName et) <+> spread (map (pprPrec 1) targs)
+    parensIf (not (null targs) && p > 3) $
+    ppr u <> ppr (qualNameFromTypeName et) <+> spread (map (pprPrec 3) targs)
   pprPrec _ (Record fs)
     | Just ts <- areTupleFields fs =
         parens $ commasep $ map ppr ts
@@ -151,7 +155,7 @@
 
 instance (Eq vn, IsName vn) => Pretty (TypeExp vn) where
   ppr (TEUnique t _) = text "*" <> ppr t
-  ppr (TEArray at d _) = ppr (ShapeDecl [d]) <> ppr at
+  ppr (TEArray at d _) = brackets (ppr d) <> ppr at
   ppr (TETuple ts _) = parens $ commasep $ map ppr ts
   ppr (TERecord fs _) = braces $ commasep $ map ppField fs
     where ppField (name, t) = text (nameToString name) <> colon <+> ppr t
@@ -165,8 +169,8 @@
     where ppConstr (name, fs) = text "#" <> ppr name <+> sep (map ppr fs)
 
 instance (Eq vn, IsName vn) => Pretty (TypeArgExp vn) where
-  ppr (TypeArgExpDim d _) = ppr $ ShapeDecl [d]
-  ppr (TypeArgExpType d) = ppr d
+  ppr (TypeArgExpDim d _) = brackets $ ppr d
+  ppr (TypeArgExpType t) = ppr t
 
 instance (Eq vn, IsName vn, Annot f) => Pretty (TypeDeclBase f vn) where
   ppr x = pprAnnot (declaredType x) (expandedType x)
@@ -204,7 +208,7 @@
   ppr = pprPrec (-1)
   pprPrec _ (Var name _ _) = ppr name
   pprPrec _ (Parens e _) = align $ parens $ ppr e
-  pprPrec _ (QualParens v e _) = ppr v <> text "." <> align (parens $ ppr e)
+  pprPrec _ (QualParens (v, _) e _) = ppr v <> text "." <> align (parens $ ppr e)
   pprPrec p (Ascript e t _ _) =
     parensIf (p /= -1) $ pprPrec 0 e <+> colon <+> pprPrec 0 t
   pprPrec _ (Literal v _) = ppr v
@@ -392,7 +396,7 @@
     text "include" <+> ppr e
 
 instance (Eq vn, IsName vn, Annot f) => Pretty (SigExpBase f vn) where
-  ppr (SigVar v _) = ppr v
+  ppr (SigVar v _ _) = ppr v
   ppr (SigParens e _) = parens $ ppr e
   ppr (SigSpecs ss _) = nestedBlock "{" "}" (stack $ punctuate line $ map ppr ss)
   ppr (SigWith s (TypeRef v ps td _) _) =
diff --git a/src/Language/Futhark/Query.hs b/src/Language/Futhark/Query.hs
new file mode 100644
--- /dev/null
+++ b/src/Language/Futhark/Query.hs
@@ -0,0 +1,383 @@
+{-# LANGUAGE FlexibleContexts #-}
+-- | Facilities for answering queries about a program, such as "what
+-- appears at this source location", or "where is this name bound".
+-- The intent is that this is used as a building block for IDE-like
+-- functionality.
+module Language.Futhark.Query
+  ( BoundTo(..)
+  , boundLoc
+  , allBindings
+  , AtPos(..)
+  , atPos
+  , Pos(..)
+  )
+  where
+
+import Control.Monad
+import Control.Monad.State
+import Data.List (find)
+import Data.Loc (Pos(..), Located(..), Loc(..))
+import qualified Data.Map as M
+import qualified System.FilePath.Posix as Posix
+
+import Language.Futhark
+import Language.Futhark.Semantic
+import Language.Futhark.Traversals
+
+-- | What a name is bound to.
+data BoundTo = BoundTerm StructType Loc
+             | BoundModule Loc
+             | BoundModuleType Loc
+             | BoundType Loc
+  deriving (Eq, Show)
+
+data Def = DefBound BoundTo | DefIndirect VName
+  deriving (Eq, Show)
+
+type Defs = M.Map VName Def
+
+-- | Where was a bound variable actually bound?  That is, what is the
+-- location of its definition?
+boundLoc :: BoundTo -> Loc
+boundLoc (BoundTerm _ loc) = loc
+boundLoc (BoundModule loc) = loc
+boundLoc (BoundModuleType loc) = loc
+boundLoc (BoundType loc) = loc
+
+patternDefs :: Pattern -> Defs
+patternDefs (Id vn (Info t) loc) =
+  M.singleton vn $ DefBound $ BoundTerm (toStruct t) (locOf loc)
+patternDefs (TuplePattern pats _) =
+  mconcat $ map patternDefs pats
+patternDefs (RecordPattern fields _) =
+  mconcat $ map (patternDefs . snd) fields
+patternDefs (PatternParens pat _) =
+  patternDefs pat
+patternDefs Wildcard{} = mempty
+patternDefs PatternLit{} = mempty
+patternDefs (PatternAscription pat _ _) =
+  patternDefs pat
+patternDefs (PatternConstr _ _ pats _) =
+  mconcat $ map patternDefs pats
+
+typeParamDefs :: TypeParamBase VName -> Defs
+typeParamDefs (TypeParamDim vn loc) =
+  M.singleton vn $ DefBound $ BoundTerm (Scalar $ Prim $ Signed Int32) (locOf loc)
+typeParamDefs (TypeParamType _ vn loc) =
+  M.singleton vn $ DefBound $ BoundType $ locOf loc
+
+expDefs :: Exp -> Defs
+expDefs e =
+  execState (astMap mapper e) extra
+  where mapper = ASTMapper { mapOnExp = onExp
+                           , mapOnName = pure
+                           , mapOnQualName = pure
+                           , mapOnStructType = pure
+                           , mapOnPatternType = pure
+                           }
+        onExp e' = do
+          modify (<>expDefs e')
+          return e'
+
+        identDefs (Ident v (Info vt) vloc) =
+          M.singleton v $ DefBound $ BoundTerm (toStruct vt) $ locOf vloc
+
+        extra =
+          case e of
+            LetPat pat _ _ _ _ ->
+              patternDefs pat
+            Lambda params _ _ _ _ ->
+              mconcat (map patternDefs params)
+            LetFun name (tparams, params, _, Info ret, _) _ loc ->
+              let name_t = foldFunType (map patternStructType params) ret
+              in M.singleton name (DefBound $ BoundTerm name_t (locOf loc)) <>
+                 mconcat (map typeParamDefs tparams) <>
+                 mconcat (map patternDefs params)
+            LetWith v _ _ _ _ _ _ ->
+              identDefs v
+            DoLoop merge _ form _ _ ->
+              patternDefs merge <>
+              case form of
+                For i _ -> identDefs i
+                ForIn pat _ -> patternDefs pat
+                While{} -> mempty
+            _ ->
+              mempty
+
+valBindDefs :: ValBind -> Defs
+valBindDefs vbind =
+  M.insert (valBindName vbind) (DefBound $ BoundTerm vbind_t (locOf vbind)) $
+  mconcat (map typeParamDefs (valBindTypeParams vbind)) <>
+  mconcat (map patternDefs (valBindParams vbind)) <>
+  expDefs (valBindBody vbind)
+  where vbind_t =
+          foldFunType (map patternStructType (valBindParams vbind)) $
+          unInfo $ valBindRetType vbind
+
+typeBindDefs :: TypeBind -> Defs
+typeBindDefs tbind =
+  M.singleton (typeAlias tbind) $ DefBound $ BoundType $ locOf tbind
+
+modParamDefs :: ModParam -> Defs
+modParamDefs (ModParam p se _ loc) =
+  M.singleton p (DefBound $ BoundModule $ locOf loc) <>
+  sigExpDefs se
+
+modExpDefs :: ModExp -> Defs
+modExpDefs ModVar{} =
+  mempty
+modExpDefs (ModParens me _) =
+  modExpDefs me
+modExpDefs ModImport{} =
+  mempty
+modExpDefs (ModDecs decs _) =
+  mconcat $ map decDefs decs
+modExpDefs (ModApply e1 e2 _ (Info substs) _) =
+  modExpDefs e1 <> modExpDefs e2 <> M.map DefIndirect substs
+modExpDefs (ModAscript e _ (Info substs) _) =
+  modExpDefs e <> M.map DefIndirect substs
+modExpDefs (ModLambda p _ e _) =
+  modParamDefs p <> modExpDefs e
+
+modBindDefs :: ModBind -> Defs
+modBindDefs mbind =
+  M.singleton (modName mbind) (DefBound $ BoundModule $ locOf mbind) <>
+  mconcat (map modParamDefs (modParams mbind)) <>
+  modExpDefs (modExp mbind) <>
+  case modSignature mbind of
+    Nothing -> mempty
+    Just (_, Info substs) ->
+      M.map DefIndirect substs
+
+specDefs :: Spec -> Defs
+specDefs spec =
+  case spec of
+    ValSpec v tparams tdecl _ loc ->
+      let vdef = DefBound $ BoundTerm (unInfo $ expandedType tdecl) (locOf loc)
+      in M.insert v vdef $ mconcat (map typeParamDefs tparams)
+    TypeAbbrSpec tbind -> typeBindDefs tbind
+    TypeSpec _ v _ _ loc ->
+      M.singleton v $ DefBound $ BoundType $ locOf loc
+    ModSpec v se _ loc ->
+      M.singleton v (DefBound $ BoundModuleType $ locOf loc) <>
+      sigExpDefs se
+    IncludeSpec se _ -> sigExpDefs se
+
+sigExpDefs :: SigExp -> Defs
+sigExpDefs se =
+  case se of
+    SigVar _ (Info substs) _ -> M.map DefIndirect substs
+    SigParens e _ -> sigExpDefs e
+    SigSpecs specs _ -> mconcat $ map specDefs specs
+    SigWith e _ _ -> sigExpDefs e
+    SigArrow _ e1 e2 _ -> sigExpDefs e1 <> sigExpDefs e2
+
+sigBindDefs :: SigBind -> Defs
+sigBindDefs sbind =
+  M.singleton (sigName sbind) (DefBound $ BoundModuleType $ locOf sbind) <>
+  sigExpDefs (sigExp sbind)
+
+decDefs :: Dec -> Defs
+decDefs (ValDec vbind) = valBindDefs vbind
+decDefs (TypeDec vbind) = typeBindDefs vbind
+decDefs (ModDec mbind) = modBindDefs mbind
+decDefs (SigDec mbind) = sigBindDefs mbind
+decDefs (OpenDec me _) = modExpDefs me
+decDefs (LocalDec dec _) = decDefs dec
+decDefs ImportDec{} = mempty
+
+-- | All bindings of everything in the program.
+progDefs :: Prog -> Defs
+progDefs = mconcat . map decDefs . progDecs
+
+allBindings :: Imports -> M.Map VName BoundTo
+allBindings imports = M.mapMaybe forward defs
+  where defs = mconcat $ map (progDefs . fileProg . snd) imports
+        forward (DefBound x) = Just x
+        forward (DefIndirect v) = forward =<< M.lookup v defs
+
+data RawAtPos = RawAtName (QualName VName) Loc
+
+contains :: Located a => a -> Pos -> Bool
+contains a pos =
+  case locOf a of
+    Loc start end -> pos >= start && pos <= end
+    NoLoc -> False
+
+atPosInTypeExp :: TypeExp VName -> Pos -> Maybe RawAtPos
+atPosInTypeExp te pos =
+  case te of
+    TEVar qn loc -> do
+      guard $ loc `contains` pos
+      Just $ RawAtName qn $ locOf loc
+    TETuple es _ ->
+      msum $ map (`atPosInTypeExp` pos) es
+    TERecord fields _ ->
+      msum $ map ((`atPosInTypeExp` pos) . snd) fields
+    TEArray te' dim _ ->
+      atPosInTypeExp te' pos `mplus` inDim dim
+    TEUnique te' _ ->
+      atPosInTypeExp te' pos
+    TEApply e1 arg _ ->
+      atPosInTypeExp e1 pos `mplus` inArg arg
+    TEArrow _ e1 e2 _ ->
+      atPosInTypeExp e1 pos `mplus` atPosInTypeExp e2 pos
+    TESum cs _ ->
+      msum $ map (`atPosInTypeExp` pos) $ concatMap snd cs
+  where inArg (TypeArgExpDim dim _) = inDim dim
+        inArg (TypeArgExpType e2) = atPosInTypeExp e2 pos
+        inDim (DimExpNamed qn loc) = do
+          guard $ loc `contains` pos
+          Just $ RawAtName qn $ locOf loc
+        inDim _ = Nothing
+
+atPosInPattern :: Pattern -> Pos -> Maybe RawAtPos
+atPosInPattern (Id vn _ loc) pos = do
+  guard $ loc `contains` pos
+  Just $ RawAtName (qualName vn) $ locOf loc
+atPosInPattern (TuplePattern pats _) pos =
+  msum $ map (`atPosInPattern` pos) pats
+atPosInPattern (RecordPattern fields _) pos =
+  msum $ map ((`atPosInPattern` pos) . snd) fields
+atPosInPattern (PatternParens pat _) pos =
+  atPosInPattern pat pos
+atPosInPattern (PatternAscription pat tdecl _) pos =
+  atPosInPattern pat pos `mplus` atPosInTypeExp (declaredType tdecl) pos
+atPosInPattern (PatternConstr _ _ pats _) pos =
+  msum $ map (`atPosInPattern` pos) pats
+atPosInPattern PatternLit{} _ = Nothing
+atPosInPattern Wildcard{} _ = Nothing
+
+atPosInExp :: Exp -> Pos -> Maybe RawAtPos
+atPosInExp (Var qn _ loc) pos = do
+  guard $ loc `contains` pos
+  Just $ RawAtName qn $ locOf loc
+atPosInExp (QualParens (qn, loc) _ _) pos
+  | loc `contains` pos = Just $ RawAtName qn $ locOf loc
+
+-- All the value cases are TODO - we need another RawAtPos constructor.
+atPosInExp Literal{} _ = Nothing
+atPosInExp IntLit{} _ = Nothing
+atPosInExp FloatLit{} _ = Nothing
+
+atPosInExp (LetPat pat _ _ _ _) pos
+  | pat `contains` pos = atPosInPattern pat pos
+
+atPosInExp (LetWith a b _ _ _ _ _) pos
+  | a `contains` pos = Just $ RawAtName (qualName $ identName a) (locOf a)
+  | b `contains` pos = Just $ RawAtName (qualName $ identName b) (locOf b)
+
+atPosInExp (DoLoop merge _ _ _ _) pos
+  | merge `contains` pos = atPosInPattern merge pos
+
+atPosInExp (Ascript _ tdecl _ _) pos
+  | tdecl `contains` pos = atPosInTypeExp (declaredType tdecl) pos
+
+atPosInExp e pos = do
+  guard $ e `contains` pos
+  -- Use the Either monad for short-circuiting for efficiency reasons.
+  -- The first hit is going to be the only one.
+  case astMap mapper e of
+    Left atpos -> Just atpos
+    Right _ -> Nothing
+  where mapper = ASTMapper { mapOnExp = onExp
+                           , mapOnName = pure
+                           , mapOnQualName = pure
+                           , mapOnStructType = pure
+                           , mapOnPatternType = pure
+                           }
+        onExp e' =
+          case atPosInExp e' pos of
+            Just atpos -> Left atpos
+            Nothing -> Right e'
+
+atPosInModExp :: ModExp -> Pos -> Maybe RawAtPos
+atPosInModExp (ModVar qn loc) pos = do
+  guard $ loc `contains` pos
+  Just $ RawAtName qn $ locOf loc
+atPosInModExp (ModParens me _) pos =
+  atPosInModExp me pos
+atPosInModExp ModImport{} _ =
+  Nothing
+atPosInModExp (ModDecs decs _) pos =
+  msum $ map (`atPosInDec` pos) decs
+atPosInModExp (ModApply e1 e2 _ _ _) pos =
+  atPosInModExp e1 pos `mplus` atPosInModExp e2 pos
+atPosInModExp (ModAscript e _ _ _) pos =
+  atPosInModExp e pos
+atPosInModExp (ModLambda _ _ e _) pos =
+  atPosInModExp e pos
+
+atPosInSpec :: Spec -> Pos -> Maybe RawAtPos
+atPosInSpec spec pos =
+  case spec of
+    ValSpec _ _ tdecl _ _ -> atPosInTypeExp (declaredType tdecl) pos
+    TypeAbbrSpec tbind -> atPosInTypeBind tbind pos
+    TypeSpec{} -> Nothing
+    ModSpec _ se _ _ -> atPosInSigExp se pos
+    IncludeSpec se _ -> atPosInSigExp se pos
+
+atPosInSigExp :: SigExp -> Pos -> Maybe RawAtPos
+atPosInSigExp se pos =
+  case se of
+    SigVar qn _ loc -> do guard $ loc `contains` pos
+                          Just $ RawAtName qn $ locOf loc
+    SigParens e _ -> atPosInSigExp e pos
+    SigSpecs specs _ -> msum $ map (`atPosInSpec` pos) specs
+    SigWith e _ _ -> atPosInSigExp e pos
+    SigArrow _ e1 e2 _ -> atPosInSigExp e1 pos `mplus` atPosInSigExp e2 pos
+
+atPosInValBind :: ValBind -> Pos -> Maybe RawAtPos
+atPosInValBind vbind pos =
+  msum (map (`atPosInPattern` pos) (valBindParams vbind)) `mplus`
+  atPosInExp (valBindBody vbind) pos `mplus`
+  join (atPosInTypeExp <$> valBindRetDecl vbind <*> pure pos)
+
+atPosInTypeBind :: TypeBind -> Pos -> Maybe RawAtPos
+atPosInTypeBind = atPosInTypeExp . declaredType . typeExp
+
+atPosInModBind :: ModBind -> Pos -> Maybe RawAtPos
+atPosInModBind (ModBind _ params sig e _ _) pos =
+  msum (map inParam params) `mplus`
+  atPosInModExp e pos `mplus`
+  case sig of Nothing -> Nothing
+              Just (se, _) -> atPosInSigExp se pos
+  where inParam (ModParam _ se _ _) = atPosInSigExp se pos
+
+atPosInSigBind :: SigBind -> Pos -> Maybe RawAtPos
+atPosInSigBind = atPosInSigExp . sigExp
+
+atPosInDec :: Dec -> Pos -> Maybe RawAtPos
+atPosInDec dec pos = do
+  guard $ dec `contains` pos
+  case dec of
+    ValDec vbind -> atPosInValBind vbind pos
+    TypeDec tbind -> atPosInTypeBind tbind pos
+    ModDec mbind -> atPosInModBind mbind pos
+    SigDec sbind -> atPosInSigBind sbind pos
+    OpenDec e _ -> atPosInModExp e pos
+    LocalDec dec' _ -> atPosInDec dec' pos
+    ImportDec{} -> Nothing
+
+atPosInProg :: Prog -> Pos -> Maybe RawAtPos
+atPosInProg prog pos =
+  msum $ map (`atPosInDec` pos) (progDecs prog)
+
+containingModule :: Imports -> Pos -> Maybe FileModule
+containingModule imports (Pos file _ _ _) =
+  snd <$> find ((==file') . fst) imports
+  where file' = includeToString $ mkInitialImport $
+                fst $ Posix.splitExtension file
+
+-- | Information about what is at the given source location.
+data AtPos = AtName (QualName VName) (Maybe BoundTo) Loc
+  deriving (Eq, Show)
+
+-- | Information about what's at the given source position.  Returns
+-- 'Nothing' if there is nothing there, including if the source
+-- position is invalid.
+atPos :: Imports -> Pos -> Maybe AtPos
+atPos imports pos = do
+  prog <- fileProg <$> containingModule imports pos
+  RawAtName qn loc <- atPosInProg prog pos
+  Just $ AtName qn (qualLeaf qn `M.lookup` allBindings imports) loc
diff --git a/src/Language/Futhark/Syntax.hs b/src/Language/Futhark/Syntax.hs
--- a/src/Language/Futhark/Syntax.hs
+++ b/src/Language/Futhark/Syntax.hs
@@ -27,6 +27,7 @@
   , qualNameFromTypeName
   , TypeBase(..)
   , TypeArg(..)
+  , DimExp(..)
   , TypeExp(..)
   , TypeArgExp(..)
   , PName(..)
@@ -387,12 +388,26 @@
 -- | A value type contains full, manifest size information.
 type ValueType = TypeBase Int32 ()
 
+-- | A dimension declaration expression for use in a 'TypeExp'.
+data DimExp vn = DimExpNamed (QualName vn) SrcLoc
+                 -- ^ The size of the dimension is this name, which
+                 -- must be in scope.
+               | DimExpConst Int SrcLoc
+                  -- ^ The size is a constant.
+               | DimExpAny
+                  -- ^ No dimension declaration.
+                deriving Show
+deriving instance Eq (DimExp Name)
+deriving instance Eq (DimExp VName)
+deriving instance Ord (DimExp Name)
+deriving instance Ord (DimExp VName)
+
 -- | An unstructured type with type variables and possibly shape
 -- declarations - this is what the user types in the source program.
 data TypeExp vn = TEVar (QualName vn) SrcLoc
                 | TETuple [TypeExp vn] SrcLoc
                 | TERecord [(Name, TypeExp vn)] SrcLoc
-                | TEArray (TypeExp vn) (DimDecl vn) SrcLoc
+                | TEArray (TypeExp vn) (DimExp vn) SrcLoc
                 | TEUnique (TypeExp vn) SrcLoc
                 | TEApply (TypeExp vn) (TypeArgExp vn) SrcLoc
                 | TEArrow (Maybe vn) (TypeExp vn) (TypeExp vn) SrcLoc
@@ -413,7 +428,7 @@
   locOf (TEArrow _ _ _ loc) = locOf loc
   locOf (TESum _ loc)      = locOf loc
 
-data TypeArgExp vn = TypeArgExpDim (DimDecl vn) SrcLoc
+data TypeArgExp vn = TypeArgExpDim (DimExp vn) SrcLoc
                    | TypeArgExpType (TypeExp vn)
                 deriving (Show)
 deriving instance Eq (TypeArgExp Name)
@@ -590,7 +605,7 @@
             | Parens (ExpBase f vn) SrcLoc
             -- ^ A parenthesized expression.
 
-            | QualParens (QualName vn) (ExpBase f vn) SrcLoc
+            | QualParens (QualName vn, SrcLoc) (ExpBase f vn) SrcLoc
 
             | TupLit    [ExpBase f vn] SrcLoc
             -- ^ Tuple literals, e.g., @{1+3, {x, y+z}}@.
@@ -859,7 +874,7 @@
   locOf (ModSpec _ _ _ loc)    = locOf loc
   locOf (IncludeSpec _ loc)    = locOf loc
 
-data SigExpBase f vn = SigVar (QualName vn) SrcLoc
+data SigExpBase f vn = SigVar (QualName vn) (f (M.Map VName VName)) SrcLoc
                      | SigParens (SigExpBase f vn) SrcLoc
                      | SigSpecs [SpecBase f vn] SrcLoc
                      | SigWith (SigExpBase f vn) (TypeRefBase f vn) SrcLoc
@@ -874,7 +889,7 @@
   locOf (TypeRef _ _ _ loc) = locOf loc
 
 instance Located (SigExpBase f vn) where
-  locOf (SigVar _ loc)       = locOf loc
+  locOf (SigVar _ _ loc)     = locOf loc
   locOf (SigParens _ loc)    = locOf loc
   locOf (SigSpecs _ loc)     = locOf loc
   locOf (SigWith _ _ loc)    = locOf loc
diff --git a/src/Language/Futhark/Traversals.hs b/src/Language/Futhark/Traversals.hs
--- a/src/Language/Futhark/Traversals.hs
+++ b/src/Language/Futhark/Traversals.hs
@@ -59,8 +59,9 @@
     FloatLit val <$> traverse (mapOnPatternType tv) t <*> pure loc
   astMap tv (Parens e loc) =
     Parens <$> mapOnExp tv e <*> pure loc
-  astMap tv (QualParens name e loc) =
-    QualParens <$> mapOnQualName tv name <*> mapOnExp tv e <*> pure loc
+  astMap tv (QualParens (name, nameloc) e loc) =
+    QualParens <$> ((,) <$> mapOnQualName tv name <*> pure nameloc) <*>
+    mapOnExp tv e <*> pure loc
   astMap tv (TupLit els loc) =
     TupLit <$> mapM (mapOnExp tv) els <*> pure loc
   astMap tv (RecordLit fields loc) =
@@ -112,7 +113,7 @@
     Project field <$> mapOnExp tv e <*> traverse (mapOnPatternType tv) t <*> pure loc
   astMap tv (Index arr idxexps t loc) =
     Index <$>
-    astMap tv arr <*>
+    mapOnExp tv arr <*>
     mapM (astMap tv) idxexps <*>
     traverse (mapOnPatternType tv) t <*>
     pure loc
@@ -156,8 +157,8 @@
 
 instance ASTMappable (LoopFormBase Info VName) where
   astMap tv (For i bound) = For <$> astMap tv i <*> astMap tv bound
-  astMap tv (ForIn pat e) = ForIn <$> astMap tv pat <*> astMap tv e
-  astMap tv (While e)     = While <$> astMap tv e
+  astMap tv (ForIn pat e) = ForIn <$> astMap tv pat <*> mapOnExp tv e
+  astMap tv (While e)     = While <$> mapOnExp tv e
 
 instance ASTMappable (TypeExp VName) where
   astMap tv (TEVar qn loc) = TEVar <$> mapOnQualName tv qn <*> pure loc
@@ -179,6 +180,12 @@
     TypeArgExpDim <$> astMap tv dim <*> pure loc
   astMap tv (TypeArgExpType te) =
     TypeArgExpType <$> astMap tv te
+
+instance ASTMappable (DimExp VName) where
+  astMap tv (DimExpNamed vn loc) =
+    DimExpNamed <$> mapOnQualName tv vn <*> pure loc
+  astMap _ (DimExpConst k loc) = pure $ DimExpConst k loc
+  astMap _ DimExpAny = pure DimExpAny
 
 instance ASTMappable (DimDecl VName) where
   astMap tv (NamedDim vn) = NamedDim <$> mapOnQualName tv vn
diff --git a/src/Language/Futhark/TypeChecker.hs b/src/Language/Futhark/TypeChecker.hs
--- a/src/Language/Futhark/TypeChecker.hs
+++ b/src/Language/Futhark/TypeChecker.hs
@@ -252,10 +252,10 @@
 checkSigExp (SigParens e loc) = do
   (mty, e') <- checkSigExp e
   return (mty, SigParens e' loc)
-checkSigExp (SigVar name loc) = do
+checkSigExp (SigVar name NoInfo loc) = do
   (name', mty) <- lookupMTy loc name
-  (mty', _) <- newNamesForMTy mty
-  return (mty', SigVar name' loc)
+  (mty', substs) <- newNamesForMTy mty
+  return (mty', SigVar name' (Info substs) loc)
 checkSigExp (SigSpecs specs loc) = do
   checkForDuplicateSpecs specs
   (abstypes, env, specs') <- checkSpecs specs
diff --git a/src/Language/Futhark/TypeChecker/Terms.hs b/src/Language/Futhark/TypeChecker/Terms.hs
--- a/src/Language/Futhark/TypeChecker/Terms.hs
+++ b/src/Language/Futhark/TypeChecker/Terms.hs
@@ -930,12 +930,12 @@
 checkExp (Parens e loc) =
   Parens <$> checkExp e <*> pure loc
 
-checkExp (QualParens modname e loc) = do
+checkExp (QualParens (modname, modnameloc) e loc) = do
   (modname',mod) <- lookupMod loc modname
   case mod of
     ModEnv env -> local (`withEnv` qualifyEnv modname' env) $ do
       e' <- checkExp e
-      return $ QualParens modname' e' loc
+      return $ QualParens (modname', modnameloc) e' loc
     ModFun{} ->
       typeError loc $ "Module " ++ pretty modname ++ " is a parametric module."
   where qualifyEnv modname' env =
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
@@ -139,17 +139,18 @@
           foldl' max Unlifted ls)
 checkTypeExp (TEArray t d loc) = do
   (t', st, l) <- checkTypeExp t
-  d' <- checkDimDecl d
-  case (l, arrayOf st (ShapeDecl [d']) Nonunique) of
+  (d', d'') <- checkDimExp d
+  case (l, arrayOf st (ShapeDecl [d'']) Nonunique) of
     (Unlifted, st') -> return (TEArray t' d' loc, st', Unlifted)
     _ -> throwError $ TypeError loc $
          "Cannot create array with elements of type " ++ quote (pretty st) ++ " (might be functional)."
-  where checkDimDecl AnyDim =
-          return AnyDim
-        checkDimDecl (ConstDim k) =
-          return $ ConstDim k
-        checkDimDecl (NamedDim v) =
-          NamedDim <$> checkNamedDim loc v
+  where checkDimExp DimExpAny =
+          return (DimExpAny, AnyDim)
+        checkDimExp (DimExpConst k dloc) =
+          return (DimExpConst k dloc, ConstDim k)
+        checkDimExp (DimExpNamed v dloc) = do
+          v' <-  checkNamedDim loc v
+          return (DimExpNamed v' dloc, NamedDim v')
 checkTypeExp (TEUnique t loc) = do
   (t', st, l) <- checkTypeExp t
   unless (mayContainArray st) $
@@ -197,15 +198,15 @@
         rootAndArgs te' = throwError $ TypeError (srclocOf te') $
                           "Type '" ++ pretty te' ++ "' is not a type constructor."
 
-        checkArgApply (TypeParamDim pv _) (TypeArgExpDim (NamedDim v) loc) = do
+        checkArgApply (TypeParamDim pv _) (TypeArgExpDim (DimExpNamed v dloc) loc) = do
           v' <- checkNamedDim loc v
-          return (TypeArgExpDim (NamedDim v') loc,
+          return (TypeArgExpDim (DimExpNamed v' dloc) loc,
                   M.singleton pv $ DimSub $ NamedDim v')
-        checkArgApply (TypeParamDim pv _) (TypeArgExpDim (ConstDim x) loc) =
-          return (TypeArgExpDim (ConstDim x) loc,
+        checkArgApply (TypeParamDim pv _) (TypeArgExpDim (DimExpConst x dloc) loc) =
+          return (TypeArgExpDim (DimExpConst x dloc) loc,
                   M.singleton pv $ DimSub $ ConstDim x)
-        checkArgApply (TypeParamDim pv _) (TypeArgExpDim AnyDim loc) =
-          return (TypeArgExpDim AnyDim loc,
+        checkArgApply (TypeParamDim pv _) (TypeArgExpDim DimExpAny loc) =
+          return (TypeArgExpDim DimExpAny loc,
                   M.singleton pv $ DimSub AnyDim)
 
         checkArgApply (TypeParamType l pv _) (TypeArgExpType te) = do
@@ -340,19 +341,19 @@
 typeExpUses (TEVar _ _) = mempty
 typeExpUses (TETuple tes _) = foldMap typeExpUses tes
 typeExpUses (TERecord fs _) = foldMap (typeExpUses . snd) fs
-typeExpUses (TEArray te d _) = typeExpUses te <> dimDeclUses d
+typeExpUses (TEArray te d _) = typeExpUses te <> dimExpUses d
 typeExpUses (TEUnique te _) = typeExpUses te
 typeExpUses (TEApply te targ _) = typeExpUses te <> typeArgUses targ
-  where typeArgUses (TypeArgExpDim d _) = dimDeclUses d
+  where typeArgUses (TypeArgExpDim d _) = dimExpUses d
         typeArgUses (TypeArgExpType tae) = typeExpUses tae
 typeExpUses (TEArrow _ t1 t2 _) =
   let (pos, neg) = typeExpUses t1 <> typeExpUses t2
   in (mempty, pos <> neg)
 typeExpUses (TESum cs _) = foldMap (mconcat . fmap typeExpUses . snd) cs
 
-dimDeclUses :: DimDecl VName -> ([VName], [VName])
-dimDeclUses (NamedDim v) = ([qualLeaf v], [])
-dimDeclUses _ = mempty
+dimExpUses :: DimExp VName -> ([VName], [VName])
+dimExpUses (DimExpNamed v _) = ([qualLeaf v], [])
+dimExpUses _ = mempty
 
 data TypeSub = TypeSub TypeBinding
              | DimSub (DimDecl VName)
diff --git a/src/futhark.hs b/src/futhark.hs
--- a/src/futhark.hs
+++ b/src/futhark.hs
@@ -35,6 +35,7 @@
 import qualified Futhark.CLI.Run as Run
 import qualified Futhark.CLI.Misc as Misc
 import qualified Futhark.CLI.Autotune as Autotune
+import qualified Futhark.CLI.Query as Query
 
 type Command = String -> [String] -> IO ()
 
@@ -68,6 +69,7 @@
            , ("check", (Misc.mainCheck, "Type check a program."))
            , ("imports", (Misc.mainImports, "Print all non-builtin imported Futhark files."))
            , ("autotune", (Autotune.main, "Autotune threshold parameters."))
+           , ("query", (Query.main, "Query semantic information about program."))
            ]
 
 msg :: String