futhark 0.25.16 → 0.25.17

raw patch · 45 files changed

+626/−468 lines, 45 filesPVP: major bump suggested

API removals or changes: PVP suggests a major version bump

API changes (from Hackage documentation)

Files

• CHANGELOG.md +17/−0
• docs/man/futhark-bench.rst +2/−1
• futhark.cabal +2/−1
• rts/c/backends/cuda.h +5/−2
• rts/c/backends/hip.h +4/−1
• rts/c/backends/opencl.h +29/−7
• src/Futhark/Analysis/AccessPattern.hs +14/−28
• src/Futhark/Analysis/Alias.hs +2/−2
• src/Futhark/Analysis/UsageTable.hs +3/−3
• src/Futhark/CLI/Bench.hs +1/−1
• src/Futhark/CLI/Dev.hs +2/−0
• src/Futhark/CLI/Literate.hs +3/−3
• src/Futhark/CLI/REPL.hs +1/−1
• src/Futhark/CodeGen/Backends/HIP.hs +1/−1
• src/Futhark/CodeGen/ImpCode/GPU.hs +12/−3
• src/Futhark/CodeGen/ImpGen.hs +84/−23
• src/Futhark/CodeGen/ImpGen/GPU.hs +5/−5
• src/Futhark/CodeGen/ImpGen/GPU/Base.hs +69/−53
• src/Futhark/CodeGen/ImpGen/GPU/Block.hs +2/−2
• src/Futhark/CodeGen/ImpGen/GPU/SegHist.hs +42/−48
• src/Futhark/CodeGen/ImpGen/GPU/SegRed.hs +8/−9
• src/Futhark/CodeGen/ImpGen/GPU/SegScan/SinglePass.hs +7/−7
• src/Futhark/CodeGen/ImpGen/GPU/SegScan/TwoPass.hs +6/−6
• src/Futhark/CodeGen/ImpGen/GPU/ToOpenCL.hs +11/−0
• src/Futhark/CodeGen/ImpGen/Multicore.hs +3/−3
• src/Futhark/CodeGen/ImpGen/Multicore/Base.hs +10/−10
• src/Futhark/CodeGen/ImpGen/Multicore/SegHist.hs +6/−7
• src/Futhark/CodeGen/ImpGen/Multicore/SegRed.hs +3/−3
• src/Futhark/CodeGen/OpenCL/Heuristics.hs +1/−0
• src/Futhark/Construct.hs +5/−5
• src/Futhark/IR/Prop/Types.hs +3/−3
• src/Futhark/Optimise/BlkRegTiling.hs +126/−130
• src/Futhark/Optimise/ReduceDeviceSyncs/MigrationTable/Graph.hs +3/−3
• src/Futhark/Optimise/TileLoops.hs +1/−1
• src/Futhark/Pkg/Info.hs +3/−3
• src/Futhark/Test.hs +17/−18
• src/Futhark/Test/Spec.hs +3/−3
• src/Futhark/Util.hs +16/−7
• src/Language/Futhark/Semantic.hs +8/−9
• src/Language/Futhark/TypeChecker/Terms/Loop.hs +42/−23
• src/Language/Futhark/TypeChecker/Terms/Monad.hs +1/−4
• src/Language/Futhark/TypeChecker/Types.hs +10/−10
• src/Language/Futhark/TypeChecker/Unify.hs +7/−19
• unittests/Language/Futhark/SemanticTests.hs +24/−0
• unittests/futhark_tests.hs +2/−0

@@ -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

@@ -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 

@@ -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

@@ -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;

@@ -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;

@@ -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. 

@@ -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

@@ -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

@@ -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)

@@ -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"]

@@ -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 [],

@@ -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]

@@ -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

@@ -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);|]            }        ]

@@ -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]) 

@@ -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

@@ -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]

@@ -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 ->

@@ -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 ->

@@ -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

@@ -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) $

@@ -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

@@ -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 <-

@@ -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 _ =

@@ -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

@@ -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 

@@ -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 $

@@ -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

@@ -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,

@@ -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

@@ -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

@@ -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 ::

@@ -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 =

@@ -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

@@ -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

@@ -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

@@ -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) 

@@ -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 ()

@@ -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

@@ -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"

@@ -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"

@@ -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'

@@ -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"

@@ -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"+    ]

@@ -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     ]