futhark 0.25.12 → 0.25.13

raw patch · 48 files changed

+1030/−705 lines, 48 filesdep ~futhark-manifestPVP: major bump suggested

API removals or changes: PVP suggests a major version bump

Dependency ranges changed: futhark-manifest

API changes (from Hackage documentation)

Files

• CHANGELOG.md +31/−1
• docs/conf.py +7/−0
• docs/index.rst +1/−0
• docs/man/futhark-literate.rst +6/−1
• docs/man/futhark-script.rst +79/−0
• futhark.cabal +15/−12
• rts/c/backends/cuda.h +1/−1
• rts/c/backends/hip.h +1/−1
• rts/c/backends/opencl.h +2/−2
• rts/python/opencl.py +34/−25
• src/Futhark/CLI/Literate.hs +81/−55
• src/Futhark/CLI/Main.hs +2/−0
• src/Futhark/CLI/REPL.hs +6/−1
• src/Futhark/CLI/Script.hs +62/−0
• src/Futhark/CodeGen/Backends/GenericC/Code.hs +3/−3
• src/Futhark/CodeGen/Backends/PyOpenCL.hs +11/−11
• src/Futhark/CodeGen/Backends/PyOpenCL/Boilerplate.hs +4/−4
• src/Futhark/CodeGen/ImpCode.hs +1/−1
• src/Futhark/CodeGen/ImpCode/GPU.hs +4/−4
• src/Futhark/CodeGen/ImpGen/GPU.hs +2/−2
• src/Futhark/CodeGen/ImpGen/GPU/Base.hs +9/−12
• src/Futhark/CodeGen/ImpGen/GPU/Block.hs +18/−2
• src/Futhark/CodeGen/ImpGen/GPU/SegHist.hs +10/−10
• src/Futhark/CodeGen/ImpGen/GPU/SegRed.hs +9/−9
• src/Futhark/CodeGen/ImpGen/GPU/SegScan/TwoPass.hs +1/−1
• src/Futhark/CodeGen/ImpGen/GPU/ToOpenCL.hs +2/−2
• src/Futhark/IR/GPUMem.hs +6/−6
• src/Futhark/IR/Mem/Simplify.hs +2/−78
• src/Futhark/IR/Parse.hs +41/−1
• src/Futhark/IR/SegOp.hs +1/−1
• src/Futhark/Internalise/Exps.hs +3/−0
• src/Futhark/Optimise/ArrayShortCircuiting/ArrayCoalescing.hs +1/−1
• src/Futhark/Optimise/BlkRegTiling.hs +3/−3
• src/Futhark/Optimise/TileLoops.hs +1/−1
• src/Futhark/Pass/ExpandAllocations.hs +47/−23
• src/Futhark/Pass/ExplicitAllocations.hs +1/−1
• src/Futhark/Pass/ExtractKernels.hs +58/−28
• src/Futhark/Pass/ExtractKernels/DistributeNests.hs +6/−4
• src/Futhark/Pass/ExtractKernels/Intrablock.hs +345/−0
• src/Futhark/Pass/ExtractKernels/Intragroup.hs +0/−345
• src/Futhark/Script.hs +26/−15
• src/Language/Futhark/TypeChecker/Consumption.hs +2/−0
• src/Language/Futhark/TypeChecker/Terms.hs +4/−5
• src/Language/Futhark/TypeChecker/Terms/Monad.hs +2/−2
• src/Language/Futhark/TypeChecker/Terms/Pat.hs +20/−22
• unittests/Futhark/IR/GPUTests.hs +20/−0
• unittests/Futhark/IR/MCTests.hs +20/−0
• unittests/Futhark/IR/SyntaxTests.hs +19/−9

@@ -5,6 +5,36 @@ 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.13]++### Added++* Incremental flattening of `map`-`scan` compositions with nested+  parallelism (similar to the logic for `map`-`reduce` compositions+  that we have had for years).++* `futhark script`, for running FutharkScript expressions from the+  command line.++* `futhark repl` now prints out a message when it ignores a breakpoint+  during initialisation. (#2098)++### Fixed++* Flattening of `scatter` with multi-dimensional elements (#2089).++* Some instances of not-actually-irregular allocations were mistakenly+  interpreted as irregular. Fixing this was a dividend of the memory+  representation simplifications of 0.25.12.++* Obscure issue related to expansion of shared memory allocations (#2092).++* A crash in alias checking under some rare circumstances (#2096).++* Mishandling of existential sizes for top level constants. (#2099)++* Compiler crash when generating code for copying nothing at all. (#2100)+ ## [0.25.12]  ### Added@@ -17,7 +47,7 @@ * The C API now allows destruction and construction of sum types (with   some caveats). (#2074) -* An overall reduction in memory copies, though simplifying the+* An overall reduction in memory copies, through simplifying the   internal representation.  ### Fixed

@@ -519,6 +519,13 @@         1,     ),     (+        "man/futhark-script",+        "futhark-script",+        "execute FutharkScript expression",+        [],+        1,+    ),+    (         "man/futhark-profile",         "futhark-profile",         "profile Futhark programs",

@@ -54,6 +54,7 @@    man/futhark-hip.rst    man/futhark-ispc.rst    man/futhark-literate.rst+   man/futhark-script.rst    man/futhark-multicore.rst    man/futhark-opencl.rst    man/futhark-pkg.rst

@@ -278,6 +278,11 @@   soundfile. Most common audio-formats are supported, including mp3, ogg, wav,   flac and opus. +FutharkScript supports a form of automatic uncurrying. If a function+taking *n* parameters is applied to a single argument that is an+*n*-element tuple, the function is applied to the elements of the+tuple as individual arguments.+ SAFETY ====== @@ -297,4 +302,4 @@ SEE ALSO ======== -:ref:`futhark-test(1)`, :ref:`futhark-bench(1)`+:ref:`futhark-script(1)`, :ref:`futhark-test(1)`, :ref:`futhark-bench(1)`

@@ -0,0 +1,79 @@+.. role:: ref(emphasis)++.. _futhark-script(1):++================+futhark-script+================++SYNOPSIS+========++futhark script [options...] program expression++DESCRIPTION+===========++The command ``futhark script foo.fut expr`` will compile ``foo.fut``,+run the provided FutharkScript expression ``expr``, and finally print+the result to stdout. It is essentially a simpler way to access the+evaluation facilities of :ref:`futhark-literate(1)`, and provides the+same FutharkScript facilities.++If the provided program does not have a ``.fut`` extension, it is+assumed to be a previously compiled server-mode program, and simply+run directly.++OPTIONS+=======++--backend=name++  The backend used when compiling Futhark programs (without leading+  ``futhark``, e.g. just ``opencl``).  Defaults to ``c``.++-D, --debug++  Pass ``-D`` to the executable and show debug prints.++--futhark=program++  The program used to perform operations (eg. compilation). Defaults+  to the binary running ``futhark script`` itself.++-L, --log++  Pass ``-L`` to the executable and show debug prints.++--pass-option=opt++  Pass an option to benchmark programs that are being run.++--pass-compiler-option=opt++  Pass an extra option to the compiler when compiling the programs.++--skip-compilation++  Do not run the compiler, and instead assume that the program has+  already been compiled.  Use with caution.++-v, --verbose++  Print verbose information on stderr about directives as they are+  executing.  This is also needed to see ``#[trace]`` output.++BUGS+====++FutharkScript expressions can only refer to names defined in the file+passed to ``futhark script``, not any names in imported files.++If the result of the expression does not have an external+representation (e.g. is an array of tuples), the value that is printed+is misleading and somewhat nonsensical.++SEE ALSO+========++:ref:`futhark-test(1)`, :ref:`futhark-bench(1)`, :ref:`futhark-literate(1)`

@@ -1,6 +1,6 @@ cabal-version: 2.4 name:           futhark-version:        0.25.12+version:        0.25.13 synopsis:       An optimising compiler for a functional, array-oriented language.  description:    Futhark is a small programming language designed to be compiled to@@ -170,6 +170,7 @@       Futhark.CLI.Query       Futhark.CLI.REPL       Futhark.CLI.Run+      Futhark.CLI.Script       Futhark.CLI.Test       Futhark.CLI.WASM       Futhark.CodeGen.Backends.CCUDA@@ -353,7 +354,7 @@       Futhark.Pass.ExtractKernels.Distribution       Futhark.Pass.ExtractKernels.ISRWIM       Futhark.Pass.ExtractKernels.Interchange-      Futhark.Pass.ExtractKernels.Intragroup+      Futhark.Pass.ExtractKernels.Intrablock       Futhark.Pass.ExtractKernels.StreamKernel       Futhark.Pass.ExtractKernels.ToGPU       Futhark.Pass.ExtractMulticore@@ -497,26 +498,28 @@   hs-source-dirs: unittests   other-modules:       Futhark.AD.DerivativesTests-      Futhark.BenchTests-      Futhark.ProfileTests-      Futhark.Pkg.SolveTests       Futhark.Analysis.AlgSimplifyTests-      Futhark.Internalise.TypesValuesTests+      Futhark.BenchTests+      Futhark.IR.GPUTests+      Futhark.IR.MCTests+      Futhark.IR.Mem.IntervalTests+      Futhark.IR.Mem.IxFun.Alg+      Futhark.IR.Mem.IxFunTests+      Futhark.IR.Mem.IxFunWrapper       Futhark.IR.Prop.RearrangeTests       Futhark.IR.Prop.ReshapeTests       Futhark.IR.PropTests       Futhark.IR.Syntax.CoreTests       Futhark.IR.SyntaxTests-      Futhark.IR.Mem.IntervalTests-      Futhark.IR.Mem.IxFun.Alg-      Futhark.IR.Mem.IxFunTests-      Futhark.IR.Mem.IxFunWrapper+      Futhark.Internalise.TypesValuesTests+      Futhark.Optimise.MemoryBlockMerging.GreedyColoringTests+      Futhark.Pkg.SolveTests+      Futhark.ProfileTests       Language.Futhark.CoreTests       Language.Futhark.PrimitiveTests       Language.Futhark.SyntaxTests-      Language.Futhark.TypeCheckerTests       Language.Futhark.TypeChecker.TypesTests-      Futhark.Optimise.MemoryBlockMerging.GreedyColoringTests+      Language.Futhark.TypeCheckerTests       Paths_futhark   build-depends:       QuickCheck >=2.8

@@ -1051,7 +1051,7 @@               msgprintf("Kernel %s with\n"                         "  grid=(%d,%d,%d)\n"                         "  block=(%d,%d,%d)\n"-                        "  local memory=%d",+                        "  shared memory=%d",                         name,                         grid[0], grid[1], grid[2],                         block[0], block[1], block[2],

@@ -907,7 +907,7 @@               msgprintf("Kernel %s with\n"                         "  grid=(%d,%d,%d)\n"                         "  block=(%d,%d,%d)\n"-                        "  local memory=%d",+                        "  shared memory=%d",                         name,                         grid[0], grid[1], grid[2],                         block[0], block[1], block[2],

@@ -1315,7 +1315,7 @@               msgprintf("Kernel %s with\n"                         "  grid=(%d,%d,%d)\n"                         "  block=(%d,%d,%d)\n"-                        "  local memory=%d",+                        "  shared memory=%d",                         name,                         grid[0], grid[1], grid[2],                         block[0], block[1], block[2],@@ -1328,7 +1328,7 @@     time_start = get_wall_time();   } -  // Some implementations do not work with 0-byte local memory.+  // Some implementations do not work with 0-byte shared memory.   if (shared_mem_bytes == 0) {     shared_mem_bytes = 4;   }

@@ -120,7 +120,7 @@     interactive=False,     platform_pref=None,     device_pref=None,-    default_tblock_size=None,+    default_group_size=None,     default_num_groups=None,     default_tile_size=None,     default_reg_tile_size=None,@@ -146,25 +146,32 @@      check_types(self, required_types) -    max_tblock_size = int(self.device.max_work_tblock_size)-    max_tile_size = int(np.sqrt(self.device.max_work_tblock_size))+    max_group_size = int(self.device.max_work_group_size)+    max_tile_size = int(np.sqrt(self.device.max_work_group_size)) -    self.max_tblock_size = max_tblock_size+    self.max_thread_block_size = max_group_size     self.max_tile_size = max_tile_size     self.max_threshold = 0-    self.max_num_groups = 0+    self.max_grid_size = 0 -    self.max_local_memory = int(self.device.local_mem_size)+    self.max_shared_memory = int(self.device.local_mem_size)      # Futhark reserves 4 bytes of local memory for its own purposes.-    self.max_local_memory -= 4+    self.max_shared_memory -= 4      # See comment in rts/c/opencl.h.     if self.platform.name.find("NVIDIA CUDA") >= 0:-        self.max_local_memory -= 12+        self.max_shared_memory -= 12     elif self.platform.name.find("AMD") >= 0:-        self.max_local_memory -= 16+        self.max_shared_memory -= 16 +    self.max_registers = int(2**16)  # Not sure how to query for this.++    self.max_cache = self.device.get_info(cl.device_info.GLOBAL_MEM_CACHE_SIZE)++    if self.max_cache == 0:+        self.max_cache = 1024 * 1024+     self.free_list = {}      self.global_failure = self.pool.allocate(np.int32().itemsize)@@ -176,9 +183,9 @@     )     self.failure_is_an_option = np.int32(0) -    if "default_tblock_size" in sizes:-        default_tblock_size = sizes["default_tblock_size"]-        del sizes["default_tblock_size"]+    if "default_group_size" in sizes:+        default_group_size = sizes["default_group_size"]+        del sizes["default_group_size"]      if "default_num_groups" in sizes:         default_num_groups = sizes["default_num_groups"]@@ -196,13 +203,13 @@         default_threshold = sizes["default_threshold"]         del sizes["default_threshold"] -    default_tblock_size_set = default_tblock_size != None+    default_group_size_set = default_group_size != None     default_tile_size_set = default_tile_size != None     default_sizes = apply_size_heuristics(         self,         size_heuristics,         {-            "tblock_size": default_tblock_size,+            "group_size": default_group_size,             "tile_size": default_tile_size,             "reg_tile_size": default_reg_tile_size,             "num_groups": default_num_groups,@@ -210,21 +217,21 @@             "threshold": default_threshold,         },     )-    default_tblock_size = default_sizes["tblock_size"]+    default_group_size = default_sizes["group_size"]     default_num_groups = default_sizes["num_groups"]     default_threshold = default_sizes["threshold"]     default_tile_size = default_sizes["tile_size"]     default_reg_tile_size = default_sizes["reg_tile_size"]     lockstep_width = default_sizes["lockstep_width"] -    if default_tblock_size > max_tblock_size:-        if default_tblock_size_set:+    if default_group_size > max_group_size:+        if default_group_size_set:             sys.stderr.write(                 "Note: Device limits group size to {} (down from {})\n".format(-                    max_tile_size, default_tblock_size+                    max_tile_size, default_group_size                 )             )-        default_tblock_size = max_tblock_size+        default_group_size = max_group_size      if default_tile_size > max_tile_size:         if default_tile_size_set:@@ -247,11 +254,11 @@      self.sizes = {}     for k, v in all_sizes.items():-        if v["class"] == "tblock_size":-            max_value = max_tblock_size-            default_value = default_tblock_size-        elif v["class"] == "num_groups":-            max_value = max_tblock_size  # Intentional!+        if v["class"] == "thread_block_size":+            max_value = max_group_size+            default_value = default_group_size+        elif v["class"] == "grid_size":+            max_value = max_group_size  # Intentional!             default_value = default_num_groups         elif v["class"] == "tile_size":             max_value = max_tile_size@@ -283,7 +290,9 @@     if len(program_src) >= 0:         build_options += ["-DLOCKSTEP_WIDTH={}".format(lockstep_width)] -        build_options += ["-D{}={}".format("max_tblock_size", max_tblock_size)]+        build_options += [+            "-D{}={}".format("max_thread_block_size", max_group_size)+        ]          build_options += [             "-D{}={}".format(

@@ -1,5 +1,14 @@ -- | @futhark literate@-module Futhark.CLI.Literate (main) where+--+-- Also contains various utility definitions used by "Futhark.CLI.Script".+module Futhark.CLI.Literate+  ( main,+    Options (..),+    initialOptions,+    scriptCommandLineOptions,+    prepareServer,+  )+where  import Codec.BMP qualified as BMP import Control.Monad@@ -651,6 +660,8 @@ literateBuiltin f vs =   scriptBuiltin "." f vs +-- | Some of these only make sense for @futhark literate@, but enough+-- are also sensible for @futhark script@ that we can share them. data Options = Options   { scriptBackend :: String,     scriptFuthark :: Maybe FilePath,@@ -1061,8 +1072,8 @@   cleanupImgDir env $ mconcat files   pure (foldl' min Success failures, T.intercalate "\n" outputs) -commandLineOptions :: [FunOptDescr Options]-commandLineOptions =+scriptCommandLineOptions :: [FunOptDescr Options]+scriptCommandLineOptions =   [ Option       []       ["backend"]@@ -1110,74 +1121,89 @@       "v"       ["verbose"]       (NoArg $ Right $ \config -> config {scriptVerbose = scriptVerbose config + 1})-      "Enable logging. Pass multiple times for more.",-    Option-      "o"-      ["output"]-      (ReqArg (\opt -> Right $ \config -> config {scriptOutput = Just opt}) "FILE")-      "Override output file. Image directory is set to basename appended with -img/.",-    Option-      []-      ["stop-on-error"]-      (NoArg $ Right $ \config -> config {scriptStopOnError = True})-      "Stop and do not produce output file if any directive fails."+      "Enable logging. Pass multiple times for more."   ] +commandLineOptions :: [FunOptDescr Options]+commandLineOptions =+  scriptCommandLineOptions+    <> [ Option+           "o"+           ["output"]+           (ReqArg (\opt -> Right $ \config -> config {scriptOutput = Just opt}) "FILE")+           "Override output file. Image directory is set to basename appended with -img/.",+         Option+           []+           ["stop-on-error"]+           (NoArg $ Right $ \config -> config {scriptStopOnError = True})+           "Stop and do not produce output file if any directive fails."+       ]++prepareServer :: FilePath -> Options -> (ScriptServer -> IO a) -> IO a+prepareServer prog opts f = do+  futhark <- maybe getExecutablePath pure $ scriptFuthark opts++  let is_fut = takeExtension prog == ".fut"++  unless (scriptSkipCompilation opts || not is_fut) $ do+    let compile_options = "--server" : scriptCompilerOptions opts+    when (scriptVerbose opts > 0) $+      T.hPutStrLn stderr $+        "Compiling " <> T.pack prog <> "..."+    when (scriptVerbose opts > 1) $+      T.hPutStrLn stderr $+        T.pack $+          unwords compile_options++    let onError err = do+          mapM_ (T.hPutStrLn stderr) err+          exitFailure++    void $+      either onError pure <=< runExceptT $+        compileProgram compile_options (FutharkExe futhark) (scriptBackend opts) prog++  let run_options = scriptExtraOptions opts+      onLine "call" l = T.putStrLn l+      onLine _ _ = pure ()+      prog' = if is_fut then dropExtension prog else prog+      cfg =+        (futharkServerCfg ("." </> prog') run_options)+          { cfgOnLine =+              if scriptVerbose opts > 0+                then onLine+                else const . const $ pure ()+          }++  withScriptServer cfg f+ -- | Run @futhark literate@. main :: String -> [String] -> IO () main = mainWithOptions initialOptions commandLineOptions "program" $ \args opts ->   case args of     [prog] -> Just $ do       futhark <- maybe getExecutablePath pure $ scriptFuthark opts--      script <- parseProgFile prog--      unless (scriptSkipCompilation opts) $ do-        let entryOpt v = "--entry-point=" ++ T.unpack v-            compile_options =-              "--server"-                : map entryOpt (S.toList (varsInScripts script))-                ++ scriptCompilerOptions opts-        when (scriptVerbose opts > 0) $-          T.hPutStrLn stderr $-            "Compiling " <> T.pack prog <> "..."-        when (scriptVerbose opts > 1) $-          T.hPutStrLn stderr $-            T.pack $-              unwords compile_options--        let onError err = do-              mapM_ (T.hPutStrLn stderr) err-              exitFailure-        void $-          either onError pure <=< runExceptT $-            compileProgram compile_options (FutharkExe futhark) (scriptBackend opts) prog-       let onError err = do             T.hPutStrLn stderr err             exitFailure       proghash <-         either onError pure <=< runExceptT $           system futhark ["hash", prog] mempty--      let mdfile = fromMaybe (prog `replaceExtension` "md") $ scriptOutput opts-          prog_dir = takeDirectory prog-          imgdir = dropExtension (takeFileName mdfile) <> "-img"-          run_options = scriptExtraOptions opts-          onLine "call" l = T.putStrLn l-          onLine _ _ = pure ()-          cfg =-            (futharkServerCfg ("." </> dropExtension prog) run_options)-              { cfgOnLine =-                  if scriptVerbose opts > 0-                    then onLine-                    else const . const $ pure ()-              }+      script <- parseProgFile prog        orig_dir <- getCurrentDirectory--      withScriptServer cfg $ \server -> do-        let env =+      let entryOpt v = "--entry-point=" ++ T.unpack v+          opts' =+            opts+              { scriptCompilerOptions =+                  map entryOpt (S.toList (varsInScripts script))+                    <> scriptCompilerOptions opts+              }+      prepareServer prog opts' $ \server -> do+        let mdfile = fromMaybe (prog `replaceExtension` "md") $ scriptOutput opts+            prog_dir = takeDirectory prog+            imgdir = dropExtension (takeFileName mdfile) <> "-img"+            env =               Env                 { envServer = server,                   envOpts = opts,

@@ -32,6 +32,7 @@ import Futhark.CLI.Query qualified as Query import Futhark.CLI.REPL qualified as REPL import Futhark.CLI.Run qualified as Run+import Futhark.CLI.Script qualified as Script import Futhark.CLI.Test qualified as Test import Futhark.CLI.WASM qualified as WASM import Futhark.Error@@ -79,6 +80,7 @@       ("defs", (Defs.main, "Show location and name of all definitions.")),       ("query", (Query.main, "Query semantic information about program.")),       ("literate", (Literate.main, "Process a literate Futhark program.")),+      ("script", (Script.main, "Run FutharkScript expressions.")),       ("lsp", (LSP.main, "Run LSP server.")),       ("thanks", (Misc.mainThanks, "Express gratitude.")),       ("tokens", (Misc.mainTokens, "Print tokens from Futhark file.")),

@@ -377,7 +377,12 @@     intOp (I.ExtOpTrace w v c) = do       liftIO $ putDocLn $ pretty w <> ":" <+> align (unAnnotate v)       c-    intOp (I.ExtOpBreak _ _ _ c) = c+    intOp (I.ExtOpBreak _ I.BreakNaN _ c) = c+    intOp (I.ExtOpBreak w _ _ c) = do+      liftIO $+        T.putStrLn $+          locText w <> ": " <> "ignoring breakpoint when computating constant."+      c  type Command = T.Text -> FutharkiM () 

@@ -0,0 +1,62 @@+module Futhark.CLI.Script (main) where++import Control.Monad.Except+import Data.Text qualified as T+import Data.Text.IO qualified as T+import Futhark.CLI.Literate+  ( Options (..),+    initialOptions,+    prepareServer,+    scriptCommandLineOptions,+  )+import Futhark.Script+import Futhark.Util.Options+import Futhark.Util.Pretty (prettyText)+import System.Exit+import System.IO++commandLineOptions :: [FunOptDescr Options]+commandLineOptions =+  scriptCommandLineOptions+    ++ [ Option+           "D"+           ["debug"]+           ( NoArg $ Right $ \config ->+               config+                 { scriptExtraOptions = "-D" : scriptExtraOptions config,+                   scriptVerbose = scriptVerbose config + 1+                 }+           )+           "Enable debugging.",+         Option+           "L"+           ["log"]+           ( NoArg $ Right $ \config ->+               config+                 { scriptExtraOptions = "-L" : scriptExtraOptions config,+                   scriptVerbose = scriptVerbose config + 1+                 }+           )+           "Enable logging."+       ]++-- | Run @futhark script@.+main :: String -> [String] -> IO ()+main = mainWithOptions initialOptions commandLineOptions "program script" $ \args opts ->+  case args of+    [prog, script] -> Just $ do+      script' <- case parseExpFromText "command line argument" $ T.pack script of+        Left e -> do+          T.hPutStrLn stderr e+          exitFailure+        Right x -> pure x+      prepareServer prog opts $ \s -> do+        r <-+          runExceptT $ getExpValue s =<< evalExp (scriptBuiltin ".") s script'+        case r of+          Left e -> do+            T.hPutStrLn stderr e+            exitFailure+          Right v ->+            T.putStrLn $ prettyText v+    _ -> Nothing

@@ -329,9 +329,7 @@ compileCode (Copy t shape (dst, dstspace) (dstoffset, dststrides) (src, srcspace) (srcoffset, srcstrides)) = do   cp <- asks $ M.lookup (dstspace, srcspace) . opsCopies . envOperations   case cp of-    Nothing ->-      compileCopy t shape (dst, dstspace) (dstoffset, dststrides) (src, srcspace) (srcoffset, srcstrides)-    Just cp' -> do+    Just cp' | t /= Unit -> do       shape' <- traverse (traverse (compileExp . untyped)) shape       dst' <- rawMem dst       src' <- rawMem src@@ -340,6 +338,8 @@       srcoffset' <- traverse (compileExp . untyped) srcoffset       srcstrides' <- traverse (traverse (compileExp . untyped)) srcstrides       cp' CopyBarrier t shape' dst' (dstoffset', dststrides') src' (srcoffset', srcstrides')+    _ ->+      compileCopy t shape (dst, dstspace) (dstoffset, dststrides) (src, srcspace) (srcoffset, srcstrides) compileCode (Write _ _ Unit _ _ _) = pure () compileCode (Write dst (Count idx) elemtype space vol elemexp) = do   dst' <- rawMem dst

@@ -59,8 +59,8 @@           Assign (Var "build_options") $ List [],           Assign (Var "preferred_device") None,           Assign (Var "default_threshold") None,-          Assign (Var "default_tblock_size") None,-          Assign (Var "default_num_tblocks") None,+          Assign (Var "default_group_size") None,+          Assign (Var "default_num_groups") None,           Assign (Var "default_tile_size") None,           Assign (Var "default_reg_tile_size") None,           Assign (Var "fut_opencl_src") $ RawStringLiteral $ opencl_prelude <> opencl_code@@ -83,8 +83,8 @@             "interactive=False",             "platform_pref=preferred_platform",             "device_pref=preferred_device",-            "default_tblock_size=default_tblock_size",-            "default_num_tblocks=default_num_tblocks",+            "default_group_size=default_group_size",+            "default_num_groups=default_num_groups",             "default_tile_size=default_tile_size",             "default_reg_tile_size=default_reg_tile_size",             "default_threshold=default_threshold",@@ -128,14 +128,14 @@               optionShortName = Nothing,               optionArgument = RequiredArgument "int",               optionAction =-                [Assign (Var "default_tblock_size") $ Var "optarg"]+                [Assign (Var "default_group_size") $ Var "optarg"]             },           Option             { optionLongName = "default-num-groups",               optionShortName = Nothing,               optionArgument = RequiredArgument "int",               optionAction =-                [Assign (Var "default_num_tblocks") $ Var "optarg"]+                [Assign (Var "default_num_groups") $ Var "optarg"]             },           Option             { optionLongName = "default-tile-size",@@ -227,14 +227,14 @@ callKernel (Imp.GetSizeMax v size_class) = do   v' <- compileVar v   stm $ Assign v' $ kernelConstToExp $ Imp.SizeMaxConst size_class-callKernel (Imp.LaunchKernel safety name shared_memory args num_threadblocks worktblock_size) = do+callKernel (Imp.LaunchKernel safety name shared_memory args num_threadblocks workgroup_size) = do   num_threadblocks' <- mapM (fmap asLong . compileExp) num_threadblocks-  worktblock_size' <- mapM compileBlockDim worktblock_size-  let kernel_size = zipWith mult_exp num_threadblocks' worktblock_size'+  workgroup_size' <- mapM compileBlockDim workgroup_size+  let kernel_size = zipWith mult_exp num_threadblocks' workgroup_size'       total_elements = foldl mult_exp (Integer 1) kernel_size       cond = BinOp "!=" total_elements (Integer 0)   shared_memory' <- compileExp $ Imp.untyped $ Imp.unCount shared_memory-  body <- collect $ launchKernel name safety kernel_size worktblock_size' shared_memory' args+  body <- collect $ launchKernel name safety kernel_size workgroup_size' shared_memory' args   stm $ If cond body []    when (safety >= Imp.SafetyFull) $@@ -266,7 +266,7 @@           ]   stm . Exp $     simpleCall (T.unpack $ kernel_name' <> ".set_args") $-      [simpleCall "cl.SharedMemory" [simpleCall "max" [shared_memory, Integer 1]]]+      [simpleCall "cl.LocalMemory" [simpleCall "max" [shared_memory, Integer 1]]]         ++ failure_args         ++ args'   stm . Exp $

@@ -59,8 +59,8 @@                                    interactive=interactive,                                    platform_pref=platform_pref,                                    device_pref=device_pref,-                                   default_tblock_size=default_tblock_size,-                                   default_num_tblocks=default_num_tblocks,+                                   default_group_size=default_group_size,+                                   default_num_groups=default_num_groups,                                    default_tile_size=default_tile_size,                                    default_reg_tile_size=default_reg_tile_size,                                    default_threshold=default_threshold,@@ -128,8 +128,8 @@          which' = case which of           LockstepWidth -> String "lockstep_width"-          NumBlocks -> String "num_tblocks"-          BlockSize -> String "tblock_size"+          NumBlocks -> String "num_groups"+          BlockSize -> String "group_size"           TileSize -> String "tile_size"           RegTileSize -> String "reg_tile_size"           Threshold -> String "threshold"

@@ -326,7 +326,7 @@     -- statement.     DebugPrint String (Maybe Exp)   | -- | Log the given message, *without* a trailing linebreak (unless-    -- part of the mssage).+    -- part of the message).     TracePrint (ErrorMsg Exp)   | -- | Perform an extensible operation.     Op a

@@ -69,7 +69,7 @@     -- kernels are examples of this.     kernelFailureTolerant :: Bool,     -- | If true, multi-versioning branches will consider this kernel-    -- when considering the local memory requirements. Set this to+    -- when considering the shared memory requirements. Set this to     -- false for kernels that do their own checking.     kernelCheckSharedMemory :: Bool   }@@ -170,7 +170,7 @@   | Atomic Space AtomicOp   | Barrier Fence   | MemFence Fence-  | LocalAlloc VName (Count Bytes (TExp Int64))+  | SharedAlloc VName (Count Bytes (TExp Int64))   | -- | Perform a barrier and also check whether any     -- threads have failed an assertion.  Make sure all     -- threads would reach all 'ErrorSync's if any of them@@ -238,8 +238,8 @@     "mem_fence_local()"   pretty (MemFence FenceGlobal) =     "mem_fence_global()"-  pretty (LocalAlloc name size) =-    pretty name <+> equals <+> "local_alloc" <> parens (pretty size)+  pretty (SharedAlloc name size) =+    pretty name <+> equals <+> "shared_alloc" <> parens (pretty size)   pretty (ErrorSync FenceLocal) =     "error_sync_local()"   pretty (ErrorSync FenceGlobal) =

@@ -165,7 +165,7 @@   compilerBugS $ "segOpCompiler: unexpected " ++ prettyString (segLevel segop) ++ " for rhs of pattern " ++ prettyString pat  -- Create boolean expression that checks whether all kernels in the--- enclosed code do not use more local memory than we have available.+-- enclosed code do not use more shared memory than we have available. -- We look at *all* the kernels here, even those that might be -- otherwise protected by their own multi-versioning branches deeper -- down.  Currently the compiler will not generate multi-versioning@@ -193,7 +193,7 @@     getKernel _ = []      localAllocSizes = foldMap localAllocSize-    localAllocSize (Imp.LocalAlloc _ size) = [size]+    localAllocSize (Imp.SharedAlloc _ size) = [size]     localAllocSize _ = []      -- These allocations will actually be padded to an 8-byte aligned

@@ -21,7 +21,6 @@     defKernelAttrs,     lvlKernelAttrs,     allocLocal,-    kernelAlloc,     compileThreadResult,     virtualiseBlocks,     kernelLoop,@@ -115,22 +114,20 @@  allocLocal :: AllocCompiler GPUMem r Imp.KernelOp allocLocal mem size =-  sOp $ Imp.LocalAlloc mem size+  sOp $ Imp.SharedAlloc mem size -kernelAlloc ::+threadAlloc ::   Pat LetDecMem ->   SubExp ->   Space ->   InKernelGen ()-kernelAlloc (Pat [_]) _ ScalarSpace {} =+threadAlloc (Pat [_]) _ ScalarSpace {} =   -- Handled by the declaration of the memory block, which is then   -- translated to an actual scalar variable during C code generation.   pure ()-kernelAlloc (Pat [mem]) size (Space "shared") =-  allocLocal (patElemName mem) $ Imp.bytes $ pe64 size-kernelAlloc (Pat [mem]) _ _ =-  compilerLimitationS $ "Cannot allocate memory block " ++ prettyString mem ++ " in kernel."-kernelAlloc dest _ _ =+threadAlloc (Pat [mem]) _ _ =+  compilerLimitationS $ "Cannot allocate memory block " ++ prettyString mem ++ " in kernel thread."+threadAlloc dest _ _ =   error $ "Invalid target for in-kernel allocation: " ++ show dest  updateAcc :: VName -> [SubExp] -> [SubExp] -> InKernelGen ()@@ -663,7 +660,7 @@  compileThreadOp :: OpCompiler GPUMem KernelEnv Imp.KernelOp compileThreadOp pat (Alloc size space) =-  kernelAlloc pat size space+  threadAlloc pat size space compileThreadOp pat _ =   compilerBugS $ "compileThreadOp: cannot compile rhs of binding " ++ prettyString pat @@ -1000,7 +997,7 @@     active i = (Imp.le64 i .<.)  -- | Change every memory block to be in the global address space,--- except those who are in the local memory space.  This only affects+-- except those who are in the shared memory space.  This only affects -- generated code - we still need to make sure that the memory is -- actually present on the device (and declared as variables in the -- kernel).@@ -1121,7 +1118,7 @@ data KernelAttrs = KernelAttrs   { -- | Can this kernel execute correctly even if previous kernels failed?     kAttrFailureTolerant :: Bool,-    -- | Does whatever launch this kernel check for local memory capacity itself?+    -- | Does whatever launch this kernel check for shared memory capacity itself?     kAttrCheckSharedMemory :: Bool,     -- | Number of blocks.     kAttrNumBlocks :: Count NumBlocks SubExp,

@@ -336,9 +336,25 @@     doSync Match {} = True     doSync _ = False +blockAlloc ::+  Pat LetDecMem ->+  SubExp ->+  Space ->+  InKernelGen ()+blockAlloc (Pat [_]) _ ScalarSpace {} =+  -- Handled by the declaration of the memory block, which is then+  -- translated to an actual scalar variable during C code generation.+  pure ()+blockAlloc (Pat [mem]) size (Space "shared") =+  allocLocal (patElemName mem) $ Imp.bytes $ pe64 size+blockAlloc (Pat [mem]) _ _ =+  compilerLimitationS $ "Cannot allocate memory block " ++ prettyString mem ++ " in kernel block."+blockAlloc dest _ _ =+  error $ "Invalid target for in-kernel allocation: " ++ show dest+ compileBlockOp :: OpCompiler GPUMem KernelEnv Imp.KernelOp compileBlockOp pat (Alloc size space) =-  kernelAlloc pat size space+  blockAlloc pat size space compileBlockOp pat (Inner (SegOp (SegMap lvl space _ body))) = do   compileFlatId space @@ -671,7 +687,7 @@       localOps threadOperations $         sWhen (kernelLocalThreadId constants .==. 0) $           copyDWIMFix (patElemName pe) gids what []-    else -- If the result of the block is an array in local memory, we+    else -- If the result of the block is an array in shared memory, we     -- store it by collective copying among all the threads of the     -- block.  TODO: also do this if the array is in global memory     -- (but this is a bit more tricky, synchronisation-wise).

@@ -10,29 +10,29 @@ -- respect the asymptotics, and do not copy the destination array. -- -- We also use a heuristic strategy for computing subhistograms in--- local memory when possible.  Given:+-- shared memory when possible.  Given: -- -- H: total size of histograms in bytes, including any lock arrays. -- -- G: block size ----- T: number of bytes of local memory each thread can be given without+-- T: number of bytes of shared memory each thread can be given without -- impacting occupancy (determined experimentally, e.g. 32). ----- LMAX: maximum amount of local memory per threadblock (hard limit).+-- LMAX: maximum amount of shared memory per threadblock (hard limit). -- -- We wish to compute: -- -- COOP: cooperation level (number of threads per subhistogram) ----- LH: number of local memory subhistograms+-- LH: number of shared memory subhistograms -- -- We do this as: -- -- COOP = ceil(H / T) -- LH = ceil((G*T)/H) -- if COOP <= G && H <= LMAX then---   use local memory+--   use shared memory -- else --   use global memory module Futhark.CodeGen.ImpGen.GPU.SegHist (compileSegHist) where@@ -706,7 +706,7 @@                           local_bucket_is                           global_bucket_is -        sComment "initialize histograms in local memory" $+        sComment "initialize histograms in shared memory" $           onAllHistograms $ \dest_local dest_global op ne local_subhisto_i global_subhisto_i local_bucket_is global_bucket_is ->             sComment "First subhistogram is initialised from global memory; others with neutral element." $ do               let global_is = map Imp.le64 segment_is ++ [0] ++ global_bucket_is@@ -772,7 +772,7 @@          sOp $ Imp.ErrorSync Imp.FenceGlobal -        sComment "Compact the multiple local memory subhistograms to result in global memory" $+        sComment "Compact the multiple shared memory subhistograms to result in global memory" $           onSlugs $ \slug dests hist_H_chk histo_dims _histo_size bins_per_thread -> do             trunc_H <-               dPrimV "trunc_H" . sMin64 hist_H_chk $@@ -988,7 +988,7 @@   emit $ Imp.DebugPrint "local B" $ Just $ untyped hist_B   emit $ Imp.DebugPrint "local M" $ Just $ untyped $ tvExp hist_M   emit $-    Imp.DebugPrint "local memory needed" $+    Imp.DebugPrint "shared memory needed" $       Just $         untyped $           hist_H * hist_el_size * sExt64 (tvExp hist_M)@@ -1016,7 +1016,7 @@             unCount num_tblocks' `divUp` hist_Nout       else pure num_tblocks' -  -- We only use local memory if the number of updates per histogram+  -- We only use shared memory if the number of updates per histogram   -- at least matches the histogram size, as otherwise it is not   -- asymptotically efficient.  This mostly matters for the segmented   -- case.@@ -1031,7 +1031,7 @@           .>. 0        run = do-        emit $ Imp.DebugPrint "## Using local memory" Nothing+        emit $ Imp.DebugPrint "## Using shared memory" Nothing         emit $ Imp.DebugPrint "Histogram size (H)" $ Just $ untyped hist_H         emit $ Imp.DebugPrint "Multiplication degree (M)" $ Just $ untyped $ tvExp hist_M         emit $ Imp.DebugPrint "Cooperation level (C)" $ Just $ untyped hist_C

@@ -45,7 +45,7 @@ -- An optimization specfically targeted at non-segmented and large-segments -- segmented reductions with non-commutative is made: The stage one main loop is -- essentially stripmined by a factor *chunk*, inserting collective copies via--- local memory of each reduction parameter going into the intra-block (partial)+-- shared memory of each reduction parameter going into the intra-block (partial) -- reductions. This saves a factor *chunk* number of intra-block reductions at -- the cost of some overhead in collective copies. module Futhark.CodeGen.ImpGen.GPU.SegRed@@ -173,7 +173,7 @@         else Imp.ValueExp $ IntValue $ intValue Int64 (1 :: Int64)  -- | Prepare intermediate arrays for the reduction.  Prim-typed--- arguments go in local memory (so we need to do the allocation of+-- arguments go in shared memory (so we need to do the allocation of -- those arrays inside the kernel), while array-typed arguments go in -- global memory.  Allocations for the latter have already been -- performed.  This policy is baked into how the allocations are done@@ -478,7 +478,7 @@             [(ltid + 1) * segment_size_nonzero - 1]        -- Finally another barrier, because we will be writing to the-      -- local memory array first thing in the next iteration.+      -- shared memory array first thing in the next iteration.       sOp $ Imp.Barrier Imp.FenceLocal  largeSegmentsReduction :: DoCompileSegRed@@ -858,7 +858,7 @@      sOp $ Imp.ErrorSync Imp.FenceLocal -    sComment "effectualize collective copies in local memory" $ do+    sComment "effectualize collective copies in shared memory" $ do       forM_ slugs $ \slug -> do         let coll_copy_arrs = slugCollCopyArrs slug         let priv_chunks = slugPrivChunks slug@@ -1031,7 +1031,7 @@ -- Intermediate memory for the nonsegmented and large segments non-commutative -- reductions with all primitive parameters: -----   These kernels need local memory for 1) the initial collective copy, 2) the+--   These kernels need shared memory for 1) the initial collective copy, 2) the --   (virtualized) block reductions, and (TODO: this one not implemented yet!) --   3) the final single-block collective copy. There are no dependencies --   between these three stages, so we can reuse the same pool of local mem for@@ -1046,7 +1046,7 @@ --   requires `tblock_size * CHUNK * max elem_sizes`, since the collective copies --   are performed in sequence (ie. inputs to different reduction operators need --   not be held in local mem simultaneously).---   2) The intra-block reductions of local memory held per-thread results+--   2) The intra-block reductions of shared memory held per-thread results --   require `tblock_size * sum elem_sizes` bytes, since per-thread results for --   all fused reductions are block-reduced simultaneously. --   3) If tblock_size < num_tblocks, then after the final single-block collective@@ -1065,12 +1065,12 @@ --   `num_tblocks > tblock_size * CHUNK`, but this is unlikely, in which case 2) --   and 3), respectively, will dominate. -----   Aside from local memory, these kernels also require a CHUNK-sized array of+--   Aside from shared memory, these kernels also require a CHUNK-sized array of --   thread-private register memory per reduction operator. -- -- For all other reductions, ie. commutative reductions, reductions with at -- least one non-primitive operator, and small segments reductions: -----   These kernels use local memory only for the intra-block reductions, and+--   These kernels use shared memory only for the intra-block reductions, and --   since they do not use chunking or CHUNK, they all require onlly `tblock_size---   * max elem_sizes` bytes of local memory and no thread-private register mem.+--   * max elem_sizes` bytes of shared memory and no thread-private register mem.

@@ -242,7 +242,7 @@                       copyDWIMFix (paramName p) [] ne []                   ) -              sComment "combine with carry and write to local memory" $+              sComment "combine with carry and write to shared memory" $                 compileStms mempty (bodyStms $ lambdaBody scan_op) $                   forM_ (zip3 rets local_arrs $ map resSubExp $ bodyResult $ lambdaBody scan_op) $                     \(t, arr, se) ->

@@ -370,7 +370,7 @@    -- The local_failure variable is an int despite only really storing   -- a single bit of information, as some OpenCL implementations-  -- (e.g. AMD) does not like byte-sized local memory (and the others+  -- (e.g. AMD) does not like byte-sized shared memory (and the others   -- likely pad to a whole word anyway).   let (safety, error_init)         -- We conservatively assume that any called function can fail.@@ -629,7 +629,7 @@       GC.stm [C.cstm|mem_fence_local();|]     kernelOps (MemFence FenceGlobal) =       GC.stm [C.cstm|mem_fence_global();|]-    kernelOps (LocalAlloc name size) = do+    kernelOps (SharedAlloc name size) = do       name' <- newVName $ prettyString name ++ "_backing"       GC.modifyUserState $ \s ->         s {kernelSharedMemory = (name', size) : kernelSharedMemory s}

@@ -101,17 +101,17 @@   simpleGeneric usage $ simplifyKernelOp $ const $ pure (NoOp, mempty)   where     -- Slightly hackily and very inefficiently, we look at the inside-    -- of SegOps to figure out the sizes of local memory allocations,+    -- of SegOps to figure out the sizes of shared memory allocations,     -- and add usages for those sizes.  This is necessary so the     -- simplifier will hoist those sizes out as far as possible (most     -- importantly, past the versioning If, but see also #1569).     usage (SegOp (SegMap _ _ _ kbody)) = localAllocs kbody     usage _ = mempty-    localAllocs = foldMap stmLocalAlloc . kernelBodyStms-    stmLocalAlloc = expLocalAlloc . stmExp-    expLocalAlloc (Op (Alloc (Var v) _)) =+    localAllocs = foldMap stmSharedAlloc . kernelBodyStms+    stmSharedAlloc = expSharedAlloc . stmExp+    expSharedAlloc (Op (Alloc (Var v) _)) =       UT.sizeUsage v-    expLocalAlloc (Op (Inner (SegOp (SegMap _ _ _ kbody)))) =+    expSharedAlloc (Op (Inner (SegOp (SegMap _ _ _ kbody)))) =       localAllocs kbody-    expLocalAlloc _ =+    expSharedAlloc _ =       mempty

@@ -9,13 +9,10 @@   ) where -import Control.Monad-import Data.List (find) import Futhark.Analysis.SymbolTable qualified as ST import Futhark.Analysis.UsageTable qualified as UT import Futhark.Construct import Futhark.IR.Mem-import Futhark.IR.Mem.LMAD qualified as LMAD import Futhark.IR.Prop.Aliases (AliasedOp) import Futhark.Optimise.Simplify qualified as Simplify import Futhark.Optimise.Simplify.Engine qualified as Engine@@ -24,7 +21,6 @@ import Futhark.Optimise.Simplify.Rules import Futhark.Pass import Futhark.Pass.ExplicitAllocations (simplifiable)-import Futhark.Util  -- | Some constraints that must hold for the simplification rules to work. type SimplifyMemory rep inner =@@ -83,12 +79,7 @@   m (Stms rep) simplifyStmsGeneric rules ops stms = do   scope <- askScope-  Simplify.simplifyStms-    ops-    rules-    blockers-    scope-    stms+  Simplify.simplifyStms ops rules blockers scope stms  isResultAlloc :: (OpC rep ~ MemOp op) => Engine.BlockPred rep isResultAlloc _ usage (Let (Pat [pe]) _ (Op Alloc {})) =@@ -115,76 +106,9 @@ memRuleBook =   standardRules     <> ruleBook-      [ RuleMatch unExistentialiseMemory,-        RuleOp decertifySafeAlloc+      [ RuleOp decertifySafeAlloc       ]       []---- | If a branch is returning some existential memory, but the size of--- the array is not existential, and the index function of the array--- does not refer to any names in the pattern, then we can create a--- block of the proper size and always return there.-unExistentialiseMemory :: (SimplifyMemory rep inner) => TopDownRuleMatch (Wise rep)-unExistentialiseMemory vtable pat _ (cond, cases, defbody, ifdec)-  | ST.simplifyMemory vtable,-    fixable <- foldl hasConcretisableMemory mempty $ patElems pat,-    not $ null fixable = Simplify $ do-      -- Create non-existential memory blocks big enough to hold the-      -- arrays.-      (arr_to_mem, oldmem_to_mem) <--        fmap unzip $-          forM fixable $ \(arr_pe, mem_size, oldmem, space) -> do-            size <- toSubExp "unext_mem_size" mem_size-            mem <- letExp "unext_mem" $ Op $ Alloc size space-            pure ((patElemName arr_pe, mem), (oldmem, mem))--      -- Update the branches to contain Copy expressions putting the-      -- arrays where they are expected.-      let updateBody body = buildBody_ $ do-            res <- bodyBind body-            zipWithM updateResult (patElems pat) res-          updateResult pat_elem (SubExpRes cs (Var v))-            | Just mem <- lookup (patElemName pat_elem) arr_to_mem,-              (_, MemArray pt shape u (ArrayIn _ lmad)) <- patElemDec pat_elem = do-                v_copy <- newVName $ baseString v <> "_nonext_copy"-                let v_pat =-                      Pat [PatElem v_copy $ MemArray pt shape u $ ArrayIn mem lmad]-                addStm $ mkWiseStm v_pat (defAux ()) $ BasicOp $ Replicate mempty $ Var v-                pure $ SubExpRes cs $ Var v_copy-            | Just mem <- lookup (patElemName pat_elem) oldmem_to_mem =-                pure $ SubExpRes cs $ Var mem-          updateResult _ se =-            pure se-      cases' <- mapM (traverse updateBody) cases-      defbody' <- updateBody defbody-      letBind pat $ Match cond cases' defbody' ifdec-  where-    onlyUsedIn name here =-      not . any ((name `nameIn`) . freeIn) . filter ((/= here) . patElemName) $-        patElems pat-    knownSize Constant {} = True-    knownSize (Var v) = not $ inContext v-    inContext = (`elem` patNames pat)--    hasConcretisableMemory fixable pat_elem-      | (_, MemArray pt shape _ (ArrayIn mem lmad)) <- patElemDec pat_elem,-        Just (j, Mem space) <--          fmap patElemType-            <$> find-              ((mem ==) . patElemName . snd)-              (zip [(0 :: Int) ..] $ patElems pat),-        Just cases_ses <- mapM (maybeNth j . bodyResult . caseBody) cases,-        Just defbody_se <- maybeNth j $ bodyResult defbody,-        mem `onlyUsedIn` patElemName pat_elem,-        all knownSize (shapeDims shape),-        not $ freeIn lmad `namesIntersect` namesFromList (patNames pat),-        any (defbody_se /=) cases_ses,-        LMAD.offset lmad == 0 =-          let mem_size = untyped $ primByteSize pt * (1 + LMAD.range lmad)-           in (pat_elem, mem_size, mem, space) : fixable-      | otherwise =-          fixable-unExistentialiseMemory _ _ _ _ = Skip  -- If an allocation is statically known to be safe, then we can remove -- the certificates on it.  This can help hoist things that would

@@ -10,8 +10,16 @@     parseSeqMem,      -- * Fragments+    parseType,     parseDeclExtType,     parseDeclType,+    parseVName,+    parseSubExp,+    parseSubExpRes,+    parseBodyGPU,+    parseBodyMC,+    parseStmGPU,+    parseStmMC,   ) where @@ -668,7 +676,13 @@ pOpaqueTypes = keyword "types" $> OpaqueTypes <*> braces (many pOpaqueType)  pProg :: PR rep -> Parser (Prog rep)-pProg pr = Prog <$> pOpaqueTypes <*> pStms pr <*> many (pFunDef pr)+pProg pr =+  Prog+    <$> (fromMaybe noTypes <$> optional pOpaqueTypes)+    <*> pStms pr+    <*> many (pFunDef pr)+  where+    noTypes = OpaqueTypes mempty  pSOAC :: PR rep -> Parser (SOAC.SOAC rep) pSOAC pr =@@ -1137,8 +1151,34 @@ parseMCMem :: FilePath -> T.Text -> Either T.Text (Prog MCMem) parseMCMem = parseRep prMCMem +--- Fragment parsers++parseType :: FilePath -> T.Text -> Either T.Text Type+parseType = parseFull pType+ parseDeclExtType :: FilePath -> T.Text -> Either T.Text DeclExtType parseDeclExtType = parseFull pDeclExtType  parseDeclType :: FilePath -> T.Text -> Either T.Text DeclType parseDeclType = parseFull pDeclType++parseVName :: FilePath -> T.Text -> Either T.Text VName+parseVName = parseFull pVName++parseSubExp :: FilePath -> T.Text -> Either T.Text SubExp+parseSubExp = parseFull pSubExp++parseSubExpRes :: FilePath -> T.Text -> Either T.Text SubExpRes+parseSubExpRes = parseFull pSubExpRes++parseBodyGPU :: FilePath -> T.Text -> Either T.Text (Body GPU)+parseBodyGPU = parseFull $ pBody prGPU++parseStmGPU :: FilePath -> T.Text -> Either T.Text (Stm GPU)+parseStmGPU = parseFull $ pStm prGPU++parseBodyMC :: FilePath -> T.Text -> Either T.Text (Body MC)+parseBodyMC = parseFull $ pBody prMC++parseStmMC :: FilePath -> T.Text -> Either T.Text (Stm MC)+parseStmMC = parseFull $ pStm prMC

@@ -1255,7 +1255,7 @@  -- If a SegRed contains two reduction operations that have the same -- vector shape, merge them together.  This saves on communication--- overhead, but can in principle lead to more local memory usage.+-- overhead, but can in principle lead to more shared memory usage. topDownSegOp _ (Pat pes) _ (SegRed lvl space ops ts kbody)   | length ops > 1,     op_groupings <-

@@ -58,6 +58,9 @@         first zeroExts . unzip . internaliseReturnType (map (fmap paramDeclType) all_params) rettype           <$> mapM subExpType body_res +      when (null params') $+        bindExtSizes (E.AppRes (E.toStruct $ E.retType rettype) (E.retDims rettype)) body_res+       body_res' <-         ensureResultExtShape msg loc (map I.fromDecl rettype') $ subExpsRes body_res       let num_ctx = length (shapeContext rettype')

@@ -420,7 +420,7 @@                 (aggSummaryMapPartial (scalarTable td_env) $ unSegSpace space0)                 (S.toList s)             Undeterminable -> pure Undeterminable-        -- Do not allow short-circuiting from a segop-local memory+        -- Do not allow short-circuiting from a segop-shared memory         -- block (not in the topdown scope) to an outer memory block.         if dstmem entry `M.member` scope td_env           && noMemOverlap td_env destination_uses source_writes

@@ -79,7 +79,7 @@   epilogue = do     let (map_t1, map_t2) = (pt_A, pt_B)     kk <- letExp "kk" =<< toExp (le64 kk0 * pe64 tk)-    -- copy A to local memory+    -- copy A to shared memory     (a_loc, aCopyLoc2Reg) <-       copyGlb2ShMem kk (gtid_y, iii, map_t1, height_A, inp_A, load_A, a_loc_init') @@ -100,9 +100,9 @@           ( do               reg_mem <- segMap2D "reg_mem" segthd_lvl ResultPrivate (ty, tx) $                 \(ltid_y, ltid_x) -> do-                  -- copy A from local memory to registers+                  -- copy A from shared memory to registers                   asss <- aCopyLoc2Reg k ltid_y-                  -- copy B from local memory to registers+                  -- copy B from shared memory to registers                   bsss <- bCopyLoc2Reg k ltid_x                   pure $ varsRes [asss, bsss]               let [asss, bsss] = reg_mem

@@ -2,7 +2,7 @@ {-# LANGUAGE TypeFamilies #-}  -- | Perform a restricted form of loop tiling within SegMaps.  We only--- tile primitive types, to avoid excessive local memory use.+-- tile primitive types, to avoid excessive shared memory use. module Futhark.Optimise.TileLoops (tileLoops) where  import Control.Monad

@@ -13,6 +13,7 @@ import Data.List (find, foldl') import Data.Map.Strict qualified as M import Data.Maybe+import Data.Sequence qualified as Seq import Futhark.Analysis.Alias as Alias import Futhark.Analysis.SymbolTable qualified as ST import Futhark.Error@@ -176,7 +177,7 @@       let num_is = shapeRank shape           is = map paramName $ take num_is $ lambdaParams op_lam       (alloc_stms, alloc_offsets) <--        genericExpandedInvariantAllocations (const (shape, map le64 is)) invariant_allocs+        genericExpandedInvariantAllocations (const $ const (shape, map le64 is)) invariant_allocs        scope <- askScope       let scope' = scopeOf op_lam <> scope <> scopeOf alloc_stms@@ -260,6 +261,10 @@ boundInKernelBody :: KernelBody GPUMem -> Names boundInKernelBody = namesFromList . M.keys . scopeOf . kernelBodyStms +addStmsToKernelBody :: Stms GPUMem -> KernelBody GPUMem -> KernelBody GPUMem+addStmsToKernelBody stms kbody =+  kbody {kernelBodyStms = stms <> kernelBodyStms kbody}+ allocsForBody ::   Extraction ->   Extraction ->@@ -278,11 +283,25 @@       variant_allocs       invariant_allocs +  -- We assume that any shared memory allocations can be inserted back+  -- into kbody'. This would not work if we had SegRed/SegScan+  -- operations that performed shared memory allocations. We don't+  -- currently, and if we would in the future, we would need to be+  -- more careful about summarising the allocations in+  -- transformScanRed.+  let (alloc_stms_dev, alloc_stms_shared) =+        Seq.partition (not . isSharedAlloc) alloc_stms+   scope <- askScope   let scope' = scopeOfSegSpace space <> scope <> scopeOf alloc_stms   either throwError pure <=< runOffsetM scope' $ do-    kbody'' <- offsetMemoryInKernelBody alloc_offsets kbody'-    m alloc_offsets alloc_stms kbody''+    kbody'' <-+      addStmsToKernelBody alloc_stms_shared+        <$> offsetMemoryInKernelBody alloc_offsets kbody'+    m alloc_offsets alloc_stms_dev kbody''+  where+    isSharedAlloc (Let _ _ (Op (Alloc _ (Space "shared")))) = True+    isSharedAlloc _ = False  memoryRequirements ::   KernelGrid ->@@ -383,10 +402,12 @@             stmsToList (get_stms body)    in (set_stms (stmsFromList stms) body, allocs) -expandable, notScalar :: Space -> Bool-expandable (Space "shared") = False-expandable ScalarSpace {} = False-expandable _ = True+expandable :: User -> Space -> Bool+expandable (SegBlock {}, _) (Space "shared") = False+expandable _ ScalarSpace {} = False+expandable _ _ = True++notScalar :: Space -> Bool notScalar ScalarSpace {} = False notScalar _ = True @@ -397,7 +418,7 @@   Stm GPUMem ->   Writer Extraction (Maybe (Stm GPUMem)) extractStmAllocations user bound_outside bound_kernel (Let (Pat [patElem]) _ (Op (Alloc size space)))-  | expandable space && expandableSize size+  | expandable user space && expandableSize size       -- FIXME: the '&& notScalar space' part is a hack because we       -- don't otherwise hoist the sizes out far enough, and we       -- promise to be super-duper-careful about not having variant@@ -449,7 +470,7 @@       pure lam {lambdaBody = body}  genericExpandedInvariantAllocations ::-  (User -> (Shape, [Exp64])) -> Extraction -> ExpandM (Stms GPUMem, RebaseMap)+  (User -> Space -> (Shape, [Exp64])) -> Extraction -> ExpandM (Stms GPUMem, RebaseMap) genericExpandedInvariantAllocations getNumUsers invariant_allocs = do   -- We expand the invariant allocations by adding an inner dimension   -- equal to the number of kernel threads.@@ -458,25 +479,25 @@   pure (alloc_stms, mconcat rebases)   where     expand (mem, (user, per_thread_size, space)) = do-      let num_users = fst $ getNumUsers user+      let num_users = fst $ getNumUsers user space           allocpat = Pat [PatElem mem $ MemMem space]       total_size <-         letExp "total_size" <=< toExp . product $           pe64 per_thread_size : map pe64 (shapeDims num_users)       letBind allocpat $ Op $ Alloc (Var total_size) space-      pure $ M.singleton mem $ newBase user+      pure $ M.singleton mem $ newBase user space -    newBaseThread user _old_shape =-      let (users_shape, user_ids) = getNumUsers user+    newBaseThread user space _old_shape =+      let (users_shape, user_ids) = getNumUsers user space           dims = map pe64 (shapeDims users_shape)        in ( flattenIndex dims user_ids,             product dims           ) -    newBase user@(SegThreadInBlock {}, _) = newBaseThread user-    newBase user@(SegThread {}, _) = newBaseThread user-    newBase user@(SegBlock {}, _) = \_old_shape ->-      let (users_shape, user_ids) = getNumUsers user+    newBase user@(SegThreadInBlock {}, _) space = newBaseThread user space+    newBase user@(SegThread {}, _) space = newBaseThread user space+    newBase user@(SegBlock {}, _) space = \_old_shape ->+      let (users_shape, user_ids) = getNumUsers user space           dims = map pe64 (shapeDims users_shape)        in ( flattenIndex dims user_ids,             product dims@@ -491,12 +512,15 @@ expandedInvariantAllocations num_threads (Count num_tblocks) (Count tblock_size) =   genericExpandedInvariantAllocations getNumUsers   where-    getNumUsers (SegThread {}, [gtid]) = (Shape [num_threads], [gtid])-    getNumUsers (SegThread {}, [gid, ltid]) = (Shape [num_tblocks, tblock_size], [gid, ltid])-    getNumUsers (SegThreadInBlock {}, [gtid]) = (Shape [num_threads], [gtid])-    getNumUsers (SegThreadInBlock {}, [gid, ltid]) = (Shape [num_tblocks, tblock_size], [gid, ltid])-    getNumUsers (SegBlock {}, [gid]) = (Shape [num_tblocks], [gid])-    getNumUsers user = error $ "getNumUsers: unhandled " ++ show user+    getNumUsers (SegThread {}, [gtid]) _ = (Shape [num_threads], [gtid])+    getNumUsers (SegThread {}, [gid, ltid]) _ = (Shape [num_tblocks, tblock_size], [gid, ltid])+    getNumUsers (SegThreadInBlock {}, [gtid]) _ = (Shape [num_threads], [gtid])+    getNumUsers (SegThreadInBlock {}, [_gid, ltid]) (Space "shared") =+      (Shape [tblock_size], [ltid])+    getNumUsers (SegThreadInBlock {}, [gid, ltid]) (Space "device") =+      (Shape [num_tblocks, tblock_size], [gid, ltid])+    getNumUsers (SegBlock {}, [gid]) _ = (Shape [num_tblocks], [gid])+    getNumUsers user space = error $ "getNumUsers: unhandled " ++ show (user, space)  expandedVariantAllocations ::   SubExp ->

@@ -76,7 +76,7 @@ data AllocEnv fromrep torep = AllocEnv   { -- | When allocating memory, put it in this memory space.     -- This is primarily used to ensure that group-wide-    -- statements store their results in local memory.+    -- statements store their results in shared memory.     allocSpace :: Space,     -- | The set of names that are known to be constants at     -- kernel compile time.

@@ -167,7 +167,7 @@ import Futhark.Pass.ExtractKernels.DistributeNests import Futhark.Pass.ExtractKernels.Distribution import Futhark.Pass.ExtractKernels.ISRWIM-import Futhark.Pass.ExtractKernels.Intragroup+import Futhark.Pass.ExtractKernels.Intrablock import Futhark.Pass.ExtractKernels.StreamKernel import Futhark.Pass.ExtractKernels.ToGPU import Futhark.Tools@@ -344,6 +344,29 @@   Lambda params ret     <$> localScope (scopeOfLParams params) (transformBody path body) +versionScanRed ::+  KernelPath ->+  Pat Type ->+  StmAux () ->+  SubExp ->+  Lambda SOACS ->+  DistribM (Stms GPU) ->+  DistribM (Body GPU) ->+  ([(Name, Bool)] -> DistribM (Body GPU)) ->+  DistribM (Stms GPU)+versionScanRed path pat aux w map_lam paralleliseOuter outerParallelBody innerParallelBody =+  if not (lambdaContainsParallelism map_lam)+    || ("sequential_inner" `inAttrs` stmAuxAttrs aux)+    then paralleliseOuter+    else do+      ((outer_suff, outer_suff_key), suff_stms) <-+        sufficientParallelism "suff_outer_screma" [w] path Nothing++      outer_stms <- outerParallelBody+      inner_stms <- innerParallelBody ((outer_suff_key, False) : path)++      (suff_stms <>) <$> kernelAlternatives pat inner_stms [(outer_suff, outer_stms)]+ transformStm :: KernelPath -> Stm SOACS -> DistribM GPUStms transformStm _ stm   | "sequential" `inAttrs` stmAuxAttrs (stmAux stm) =@@ -370,21 +393,39 @@ transformStm path (Let pat aux (Op (Screma w arrs form)))   | Just lam <- isMapSOAC form =       onMap path $ MapLoop pat aux w lam arrs-transformStm path (Let res_pat (StmAux cs _ _) (Op (Screma w arrs form)))+transformStm path (Let pat aux@(StmAux cs _ _) (Op (Screma w arrs form)))   | Just scans <- isScanSOAC form,     Scan scan_lam nes <- singleScan scans,-    Just do_iswim <- iswim res_pat w scan_lam $ zip nes arrs = do+    Just do_iswim <- iswim pat w scan_lam $ zip nes arrs = do       types <- asksScope scopeForSOACs       transformStms path . stmsToList . snd =<< runBuilderT (certifying cs do_iswim) types-  | Just (scans, map_lam) <- isScanomapSOAC form = runBuilder_ $ do-      scan_ops <- forM scans $ \(Scan scan_lam nes) -> do-        (scan_lam', nes', shape) <- determineReduceOp scan_lam nes-        let scan_lam'' = soacsLambdaToGPU scan_lam'-        pure $ SegBinOp Noncommutative scan_lam'' nes' shape-      let map_lam_sequential = soacsLambdaToGPU map_lam-      lvl <- segThreadCapped [w] "segscan" $ NoRecommendation SegNoVirt-      addStms . fmap (certify cs)-        =<< segScan lvl res_pat mempty w scan_ops map_lam_sequential arrs [] []+  | Just (scans, map_lam) <- isScanomapSOAC form = do+      let paralleliseOuter = runBuilder_ $ do+            scan_ops <- forM scans $ \(Scan scan_lam nes) -> do+              (scan_lam', nes', shape) <- determineReduceOp scan_lam nes+              let scan_lam'' = soacsLambdaToGPU scan_lam'+              pure $ SegBinOp Noncommutative scan_lam'' nes' shape+            let map_lam_sequential = soacsLambdaToGPU map_lam+            lvl <- segThreadCapped [w] "segscan" $ NoRecommendation SegNoVirt+            addStms . fmap (certify cs)+              =<< segScan lvl pat mempty w scan_ops map_lam_sequential arrs [] []++          outerParallelBody =+            renameBody+              =<< (mkBody <$> paralleliseOuter <*> pure (varsRes (patNames pat)))++          paralleliseInner path' = do+            (mapstm, scanstm) <-+              scanomapToMapAndScan pat (w, scans, map_lam, arrs)+            types <- asksScope scopeForSOACs+            transformStms path' . stmsToList <=< (`runBuilderT_` types) $+              addStms =<< simplifyStms (stmsFromList [certify cs mapstm, certify cs scanstm])++          innerParallelBody path' =+            renameBody+              =<< (mkBody <$> paralleliseInner path' <*> pure (varsRes (patNames pat)))++      versionScanRed path pat aux w map_lam paralleliseOuter outerParallelBody innerParallelBody transformStm path (Let res_pat aux (Op (Screma w arrs form)))   | Just [Reduce comm red_fun nes] <- isReduceSOAC form,     let comm'@@ -424,18 +465,7 @@             renameBody               =<< (mkBody <$> paralleliseInner path' <*> pure (varsRes (patNames pat))) -      if not (lambdaContainsParallelism map_lam)-        || "sequential_inner"-          `inAttrs` stmAuxAttrs aux-        then paralleliseOuter-        else do-          ((outer_suff, outer_suff_key), suff_stms) <--            sufficientParallelism "suff_outer_redomap" [w] path Nothing--          outer_stms <- outerParallelBody-          inner_stms <- innerParallelBody ((outer_suff_key, False) : path)--          (suff_stms <>) <$> kernelAlternatives pat inner_stms [(outer_suff, outer_stms)]+      versionScanRed path pat aux w map_lam paralleliseOuter outerParallelBody innerParallelBody transformStm path (Let pat (StmAux cs _ _) (Op (Screma w arrs form))) = do   -- This screma is too complicated for us to immediately do   -- anything, so split it up and try again.@@ -547,8 +577,8 @@ -- | Intra-group parallelism is worthwhile if the lambda contains more -- than one instance of non-map nested parallelism, or any nested -- parallelism inside a loop.-worthIntraGroup :: Lambda SOACS -> Bool-worthIntraGroup lam = bodyInterest (lambdaBody lam) > 1+worthIntrablock :: Lambda SOACS -> Bool+worthIntrablock lam = bodyInterest (lambdaBody lam) > 1   where     bodyInterest body =       sum $ interest <$> bodyStms body@@ -707,11 +737,11 @@   types <- askScope    let only_intra = onlyExploitIntra (stmAuxAttrs aux)-      may_intra = worthIntraGroup lam && mayExploitIntra attrs+      may_intra = worthIntrablock lam && mayExploitIntra attrs    intra <-     if only_intra || may_intra-      then flip runReaderT types $ intraGroupParallelise loopnest lam+      then flip runReaderT types $ intrablockParallelise loopnest lam       else pure Nothing    case intra of

@@ -793,13 +793,15 @@    let grouped = groupScatterResults (zip3 as_ws as_ns as_inps) (is' ++ vs)       (_, dest_arrs, _) = unzip3 grouped-      k_body = KernelBody () k_body_stms (map (inPlaceReturn ispace) grouped)++  dest_arrs_ts <- mapM (lookupType . kernelInputArray) dest_arrs++  let k_body = KernelBody () k_body_stms (zipWith (inPlaceReturn ispace) dest_arrs_ts grouped)       -- Remove unused kernel inputs, since some of these might       -- reference the array we are scattering into.       kernel_inps' =         filter ((`nameIn` freeIn k_body) . kernelInputName) kernel_inps -  dest_arrs_ts <- mapM (lookupType . kernelInputArray) dest_arrs   (k, k_stms) <- mapKernel mk_lvl ispace kernel_inps' dest_arrs_ts k_body    traverse renameStm <=< runBuilder_ $ do@@ -817,13 +819,13 @@       maybe bad pure $ find ((== a) . kernelInputName) kernel_inps     bad = error "Ill-typed nested scatter encountered." -    inPlaceReturn ispace (_aw, inp, is_vs) =+    inPlaceReturn ispace arr_t (_, inp, is_vs) =       WriteReturns         ( foldMap (foldMap resCerts . fst) is_vs             <> foldMap (resCerts . snd) is_vs         )         (kernelInputArray inp)-        [ (Slice $ map DimFix $ map Var (init gtids) ++ map resSubExp is, resSubExp v)+        [ (fullSlice arr_t $ map DimFix $ map Var (init gtids) ++ map resSubExp is, resSubExp v)           | (is, v) <- is_vs         ]       where

@@ -0,0 +1,345 @@+{-# LANGUAGE TypeFamilies #-}++-- | Extract limited nested parallelism for execution inside+-- individual kernel threadblocks.+module Futhark.Pass.ExtractKernels.Intrablock (intrablockParallelise) where++import Control.Monad+import Control.Monad.RWS+import Control.Monad.Trans.Maybe+import Data.Map.Strict qualified as M+import Data.Set qualified as S+import Futhark.Analysis.PrimExp.Convert+import Futhark.IR.GPU hiding (HistOp)+import Futhark.IR.GPU.Op qualified as GPU+import Futhark.IR.SOACS+import Futhark.MonadFreshNames+import Futhark.Pass.ExtractKernels.BlockedKernel+import Futhark.Pass.ExtractKernels.DistributeNests+import Futhark.Pass.ExtractKernels.Distribution+import Futhark.Pass.ExtractKernels.ToGPU+import Futhark.Tools+import Futhark.Transform.FirstOrderTransform qualified as FOT+import Futhark.Util.Log+import Prelude hiding (log)++-- | Convert the statements inside a map nest to kernel statements,+-- attempting to parallelise any remaining (top-level) parallel+-- statements.  Anything that is not a map, scan or reduction will+-- simply be sequentialised.  This includes sequential loops that+-- contain maps, scans or reduction.  In the future, we could probably+-- do something more clever.  Make sure that the amount of parallelism+-- to be exploited does not exceed the group size.  Further, as a hack+-- we also consider the size of all intermediate arrays as+-- "parallelism to be exploited" to avoid exploding shared memory.+--+-- We distinguish between "minimum group size" and "maximum+-- exploitable parallelism".+intrablockParallelise ::+  (MonadFreshNames m, LocalScope GPU m) =>+  KernelNest ->+  Lambda SOACS ->+  m+    ( Maybe+        ( (SubExp, SubExp),+          SubExp,+          Log,+          Stms GPU,+          Stms GPU+        )+    )+intrablockParallelise knest lam = runMaybeT $ do+  (ispace, inps) <- lift $ flatKernel knest++  (num_tblocks, w_stms) <-+    lift $+      runBuilder $+        letSubExp "intra_num_tblocks"+          =<< foldBinOp (Mul Int64 OverflowUndef) (intConst Int64 1) (map snd ispace)++  let body = lambdaBody lam++  tblock_size <- newVName "computed_tblock_size"+  (wss_min, wss_avail, log, kbody) <-+    lift . localScope (scopeOfLParams $ lambdaParams lam) $+      intrablockParalleliseBody body++  outside_scope <- lift askScope+  -- outside_scope may also contain the inputs, even though those are+  -- not actually available outside the kernel.+  let available v =+        v `M.member` outside_scope+          && v `notElem` map kernelInputName inps+  unless (all available $ namesToList $ freeIn (wss_min ++ wss_avail)) $+    fail "Irregular parallelism"++  ((intra_avail_par, kspace, read_input_stms), prelude_stms) <- lift $+    runBuilder $ do+      let foldBinOp' _ [] = eSubExp $ intConst Int64 1+          foldBinOp' bop (x : xs) = foldBinOp bop x xs+      ws_min <-+        mapM (letSubExp "one_intra_par_min" <=< foldBinOp' (Mul Int64 OverflowUndef)) $+          filter (not . null) wss_min+      ws_avail <-+        mapM (letSubExp "one_intra_par_avail" <=< foldBinOp' (Mul Int64 OverflowUndef)) $+          filter (not . null) wss_avail++      -- The amount of parallelism available *in the worst case* is+      -- equal to the smallest parallel loop, or *at least* 1.+      intra_avail_par <-+        letSubExp "intra_avail_par" =<< foldBinOp' (SMin Int64) ws_avail++      -- The group size is either the maximum of the minimum parallelism+      -- exploited, or the desired parallelism (bounded by the max group+      -- size) in case there is no minimum.+      letBindNames [tblock_size]+        =<< if null ws_min+          then+            eBinOp+              (SMin Int64)+              (eSubExp =<< letSubExp "max_tblock_size" (Op $ SizeOp $ GetSizeMax SizeThreadBlock))+              (eSubExp intra_avail_par)+          else foldBinOp' (SMax Int64) ws_min++      let inputIsUsed input = kernelInputName input `nameIn` freeIn body+          used_inps = filter inputIsUsed inps++      addStms w_stms+      read_input_stms <- runBuilder_ $ mapM readGroupKernelInput used_inps+      space <- SegSpace <$> newVName "phys_tblock_id" <*> pure ispace+      pure (intra_avail_par, space, read_input_stms)++  let kbody' = kbody {kernelBodyStms = read_input_stms <> kernelBodyStms kbody}++  let nested_pat = loopNestingPat first_nest+      rts = map (length ispace `stripArray`) $ patTypes nested_pat+      grid = KernelGrid (Count num_tblocks) (Count $ Var tblock_size)+      lvl = SegBlock SegNoVirt (Just grid)+      kstm =+        Let nested_pat aux $ Op $ SegOp $ SegMap lvl kspace rts kbody'++  let intra_min_par = intra_avail_par+  pure+    ( (intra_min_par, intra_avail_par),+      Var tblock_size,+      log,+      prelude_stms,+      oneStm kstm+    )+  where+    first_nest = fst knest+    aux = loopNestingAux first_nest++readGroupKernelInput ::+  (DistRep (Rep m), MonadBuilder m) =>+  KernelInput ->+  m ()+readGroupKernelInput inp+  | Array {} <- kernelInputType inp = do+      v <- newVName $ baseString $ kernelInputName inp+      readKernelInput inp {kernelInputName = v}+      letBindNames [kernelInputName inp] $ BasicOp $ Replicate mempty $ Var v+  | otherwise =+      readKernelInput inp++data IntraAcc = IntraAcc+  { accMinPar :: S.Set [SubExp],+    accAvailPar :: S.Set [SubExp],+    accLog :: Log+  }++instance Semigroup IntraAcc where+  IntraAcc min_x avail_x log_x <> IntraAcc min_y avail_y log_y =+    IntraAcc (min_x <> min_y) (avail_x <> avail_y) (log_x <> log_y)++instance Monoid IntraAcc where+  mempty = IntraAcc mempty mempty mempty++type IntrablockM =+  BuilderT GPU (RWS () IntraAcc VNameSource)++instance MonadLogger IntrablockM where+  addLog log = tell mempty {accLog = log}++runIntrablockM ::+  (MonadFreshNames m, HasScope GPU m) =>+  IntrablockM () ->+  m (IntraAcc, Stms GPU)+runIntrablockM m = do+  scope <- castScope <$> askScope+  modifyNameSource $ \src ->+    let (((), kstms), src', acc) = runRWS (runBuilderT m scope) () src+     in ((acc, kstms), src')++parallelMin :: [SubExp] -> IntrablockM ()+parallelMin ws =+  tell+    mempty+      { accMinPar = S.singleton ws,+        accAvailPar = S.singleton ws+      }++intrablockBody :: Body SOACS -> IntrablockM (Body GPU)+intrablockBody body = do+  stms <- collectStms_ $ intrablockStms $ bodyStms body+  pure $ mkBody stms $ bodyResult body++intrablockLambda :: Lambda SOACS -> IntrablockM (Lambda GPU)+intrablockLambda lam =+  mkLambda (lambdaParams lam) $+    bodyBind =<< intrablockBody (lambdaBody lam)++intrablockWithAccInput :: WithAccInput SOACS -> IntrablockM (WithAccInput GPU)+intrablockWithAccInput (shape, arrs, Nothing) =+  pure (shape, arrs, Nothing)+intrablockWithAccInput (shape, arrs, Just (lam, nes)) = do+  lam' <- intrablockLambda lam+  pure (shape, arrs, Just (lam', nes))++intrablockStm :: Stm SOACS -> IntrablockM ()+intrablockStm stm@(Let pat aux e) = do+  scope <- askScope+  let lvl = SegThreadInBlock SegNoVirt++  case e of+    Loop merge form loopbody ->+      localScope (scopeOfLoopForm form <> scopeOfFParams (map fst merge)) $ do+        loopbody' <- intrablockBody loopbody+        certifying (stmAuxCerts aux) . letBind pat $+          Loop merge form loopbody'+    Match cond cases defbody ifdec -> do+      cases' <- mapM (traverse intrablockBody) cases+      defbody' <- intrablockBody defbody+      certifying (stmAuxCerts aux) . letBind pat $+        Match cond cases' defbody' ifdec+    WithAcc inputs lam -> do+      inputs' <- mapM intrablockWithAccInput inputs+      lam' <- intrablockLambda lam+      certifying (stmAuxCerts aux) . letBind pat $ WithAcc inputs' lam'+    Op soac+      | "sequential_outer" `inAttrs` stmAuxAttrs aux ->+          intrablockStms . fmap (certify (stmAuxCerts aux))+            =<< runBuilder_ (FOT.transformSOAC pat soac)+    Op (Screma w arrs form)+      | Just lam <- isMapSOAC form -> do+          let loopnest = MapNesting pat aux w $ zip (lambdaParams lam) arrs+              env =+                DistEnv+                  { distNest =+                      singleNesting $ Nesting mempty loopnest,+                    distScope =+                      scopeOfPat pat+                        <> scopeForGPU (scopeOf lam)+                        <> scope,+                    distOnInnerMap =+                      distributeMap,+                    distOnTopLevelStms =+                      liftInner . collectStms_ . intrablockStms,+                    distSegLevel = \minw _ _ -> do+                      lift $ parallelMin minw+                      pure lvl,+                    distOnSOACSStms =+                      pure . oneStm . soacsStmToGPU,+                    distOnSOACSLambda =+                      pure . soacsLambdaToGPU+                  }+              acc =+                DistAcc+                  { distTargets = singleTarget (pat, bodyResult $ lambdaBody lam),+                    distStms = mempty+                  }++          addStms+            =<< runDistNestT env (distributeMapBodyStms acc (bodyStms $ lambdaBody lam))+    Op (Screma w arrs form)+      | Just (scans, mapfun) <- isScanomapSOAC form,+        -- FIXME: Futhark.CodeGen.ImpGen.GPU.Block.compileGroupOp+        -- cannot handle multiple scan operators yet.+        Scan scanfun nes <- singleScan scans -> do+          let scanfun' = soacsLambdaToGPU scanfun+              mapfun' = soacsLambdaToGPU mapfun+          certifying (stmAuxCerts aux) $+            addStms =<< segScan lvl pat mempty w [SegBinOp Noncommutative scanfun' nes mempty] mapfun' arrs [] []+          parallelMin [w]+    Op (Screma w arrs form)+      | Just (reds, map_lam) <- isRedomapSOAC form -> do+          let onReduce (Reduce comm red_lam nes) =+                SegBinOp comm (soacsLambdaToGPU red_lam) nes mempty+              reds' = map onReduce reds+              map_lam' = soacsLambdaToGPU map_lam+          certifying (stmAuxCerts aux) $+            addStms =<< segRed lvl pat mempty w reds' map_lam' arrs [] []+          parallelMin [w]+    Op (Screma w arrs form) ->+      -- This screma is too complicated for us to immediately do+      -- anything, so split it up and try again.+      mapM_ intrablockStm . fmap (certify (stmAuxCerts aux)) . snd+        =<< runBuilderT (dissectScrema pat w form arrs) (scopeForSOACs scope)+    Op (Hist w arrs ops bucket_fun) -> do+      ops' <- forM ops $ \(HistOp num_bins rf dests nes op) -> do+        (op', nes', shape) <- determineReduceOp op nes+        let op'' = soacsLambdaToGPU op'+        pure $ GPU.HistOp num_bins rf dests nes' shape op''++      let bucket_fun' = soacsLambdaToGPU bucket_fun+      certifying (stmAuxCerts aux) $+        addStms =<< segHist lvl pat w [] [] ops' bucket_fun' arrs+      parallelMin [w]+    Op (Stream w arrs accs lam)+      | chunk_size_param : _ <- lambdaParams lam -> do+          types <- asksScope castScope+          ((), stream_stms) <-+            runBuilderT (sequentialStreamWholeArray pat w accs lam arrs) types+          let replace (Var v) | v == paramName chunk_size_param = w+              replace se = se+              replaceSets (IntraAcc x y log) =+                IntraAcc (S.map (map replace) x) (S.map (map replace) y) log+          censor replaceSets $ intrablockStms stream_stms+    Op (Scatter w ivs lam dests) -> do+      write_i <- newVName "write_i"+      space <- mkSegSpace [(write_i, w)]++      let lam' = soacsLambdaToGPU lam+          krets = do+            (_a_w, a, is_vs) <-+              groupScatterResults dests $ bodyResult $ lambdaBody lam'+            let cs =+                  foldMap (foldMap resCerts . fst) is_vs+                    <> foldMap (resCerts . snd) is_vs+                is_vs' = [(Slice $ map (DimFix . resSubExp) is, resSubExp v) | (is, v) <- is_vs]+            pure $ WriteReturns cs a is_vs'+          inputs = do+            (p, p_a) <- zip (lambdaParams lam') ivs+            pure $ KernelInput (paramName p) (paramType p) p_a [Var write_i]++      kstms <- runBuilder_ $+        localScope (scopeOfSegSpace space) $ do+          mapM_ readKernelInput inputs+          addStms $ bodyStms $ lambdaBody lam'++      certifying (stmAuxCerts aux) $ do+        let body = KernelBody () kstms krets+        letBind pat $ Op $ SegOp $ SegMap lvl space (patTypes pat) body++      parallelMin [w]+    _ ->+      addStm $ soacsStmToGPU stm++intrablockStms :: Stms SOACS -> IntrablockM ()+intrablockStms = mapM_ intrablockStm++intrablockParalleliseBody ::+  (MonadFreshNames m, HasScope GPU m) =>+  Body SOACS ->+  m ([[SubExp]], [[SubExp]], Log, KernelBody GPU)+intrablockParalleliseBody body = do+  (IntraAcc min_ws avail_ws log, kstms) <-+    runIntrablockM $ intrablockStms $ bodyStms body+  pure+    ( S.toList min_ws,+      S.toList avail_ws,+      log,+      KernelBody () kstms $ map ret $ bodyResult body+    )+  where+    ret (SubExpRes cs se) = Returns ResultMaySimplify cs se

@@ -1,345 +0,0 @@-{-# LANGUAGE TypeFamilies #-}---- | Extract limited nested parallelism for execution inside--- individual kernel threadblocks.-module Futhark.Pass.ExtractKernels.Intragroup (intraGroupParallelise) where--import Control.Monad-import Control.Monad.RWS-import Control.Monad.Trans.Maybe-import Data.Map.Strict qualified as M-import Data.Set qualified as S-import Futhark.Analysis.PrimExp.Convert-import Futhark.IR.GPU hiding (HistOp)-import Futhark.IR.GPU.Op qualified as GPU-import Futhark.IR.SOACS-import Futhark.MonadFreshNames-import Futhark.Pass.ExtractKernels.BlockedKernel-import Futhark.Pass.ExtractKernels.DistributeNests-import Futhark.Pass.ExtractKernels.Distribution-import Futhark.Pass.ExtractKernels.ToGPU-import Futhark.Tools-import Futhark.Transform.FirstOrderTransform qualified as FOT-import Futhark.Util.Log-import Prelude hiding (log)---- | Convert the statements inside a map nest to kernel statements,--- attempting to parallelise any remaining (top-level) parallel--- statements.  Anything that is not a map, scan or reduction will--- simply be sequentialised.  This includes sequential loops that--- contain maps, scans or reduction.  In the future, we could probably--- do something more clever.  Make sure that the amount of parallelism--- to be exploited does not exceed the group size.  Further, as a hack--- we also consider the size of all intermediate arrays as--- "parallelism to be exploited" to avoid exploding local memory.------ We distinguish between "minimum group size" and "maximum--- exploitable parallelism".-intraGroupParallelise ::-  (MonadFreshNames m, LocalScope GPU m) =>-  KernelNest ->-  Lambda SOACS ->-  m-    ( Maybe-        ( (SubExp, SubExp),-          SubExp,-          Log,-          Stms GPU,-          Stms GPU-        )-    )-intraGroupParallelise knest lam = runMaybeT $ do-  (ispace, inps) <- lift $ flatKernel knest--  (num_tblocks, w_stms) <--    lift $-      runBuilder $-        letSubExp "intra_num_tblocks"-          =<< foldBinOp (Mul Int64 OverflowUndef) (intConst Int64 1) (map snd ispace)--  let body = lambdaBody lam--  tblock_size <- newVName "computed_tblock_size"-  (wss_min, wss_avail, log, kbody) <--    lift . localScope (scopeOfLParams $ lambdaParams lam) $-      intraGroupParalleliseBody body--  outside_scope <- lift askScope-  -- outside_scope may also contain the inputs, even though those are-  -- not actually available outside the kernel.-  let available v =-        v `M.member` outside_scope-          && v `notElem` map kernelInputName inps-  unless (all available $ namesToList $ freeIn (wss_min ++ wss_avail)) $-    fail "Irregular parallelism"--  ((intra_avail_par, kspace, read_input_stms), prelude_stms) <- lift $-    runBuilder $ do-      let foldBinOp' _ [] = eSubExp $ intConst Int64 1-          foldBinOp' bop (x : xs) = foldBinOp bop x xs-      ws_min <--        mapM (letSubExp "one_intra_par_min" <=< foldBinOp' (Mul Int64 OverflowUndef)) $-          filter (not . null) wss_min-      ws_avail <--        mapM (letSubExp "one_intra_par_avail" <=< foldBinOp' (Mul Int64 OverflowUndef)) $-          filter (not . null) wss_avail--      -- The amount of parallelism available *in the worst case* is-      -- equal to the smallest parallel loop, or *at least* 1.-      intra_avail_par <--        letSubExp "intra_avail_par" =<< foldBinOp' (SMin Int64) ws_avail--      -- The group size is either the maximum of the minimum parallelism-      -- exploited, or the desired parallelism (bounded by the max group-      -- size) in case there is no minimum.-      letBindNames [tblock_size]-        =<< if null ws_min-          then-            eBinOp-              (SMin Int64)-              (eSubExp =<< letSubExp "max_tblock_size" (Op $ SizeOp $ GetSizeMax SizeThreadBlock))-              (eSubExp intra_avail_par)-          else foldBinOp' (SMax Int64) ws_min--      let inputIsUsed input = kernelInputName input `nameIn` freeIn body-          used_inps = filter inputIsUsed inps--      addStms w_stms-      read_input_stms <- runBuilder_ $ mapM readGroupKernelInput used_inps-      space <- SegSpace <$> newVName "phys_tblock_id" <*> pure ispace-      pure (intra_avail_par, space, read_input_stms)--  let kbody' = kbody {kernelBodyStms = read_input_stms <> kernelBodyStms kbody}--  let nested_pat = loopNestingPat first_nest-      rts = map (length ispace `stripArray`) $ patTypes nested_pat-      grid = KernelGrid (Count num_tblocks) (Count $ Var tblock_size)-      lvl = SegBlock SegNoVirt (Just grid)-      kstm =-        Let nested_pat aux $ Op $ SegOp $ SegMap lvl kspace rts kbody'--  let intra_min_par = intra_avail_par-  pure-    ( (intra_min_par, intra_avail_par),-      Var tblock_size,-      log,-      prelude_stms,-      oneStm kstm-    )-  where-    first_nest = fst knest-    aux = loopNestingAux first_nest--readGroupKernelInput ::-  (DistRep (Rep m), MonadBuilder m) =>-  KernelInput ->-  m ()-readGroupKernelInput inp-  | Array {} <- kernelInputType inp = do-      v <- newVName $ baseString $ kernelInputName inp-      readKernelInput inp {kernelInputName = v}-      letBindNames [kernelInputName inp] $ BasicOp $ Replicate mempty $ Var v-  | otherwise =-      readKernelInput inp--data IntraAcc = IntraAcc-  { accMinPar :: S.Set [SubExp],-    accAvailPar :: S.Set [SubExp],-    accLog :: Log-  }--instance Semigroup IntraAcc where-  IntraAcc min_x avail_x log_x <> IntraAcc min_y avail_y log_y =-    IntraAcc (min_x <> min_y) (avail_x <> avail_y) (log_x <> log_y)--instance Monoid IntraAcc where-  mempty = IntraAcc mempty mempty mempty--type IntraGroupM =-  BuilderT GPU (RWS () IntraAcc VNameSource)--instance MonadLogger IntraGroupM where-  addLog log = tell mempty {accLog = log}--runIntraGroupM ::-  (MonadFreshNames m, HasScope GPU m) =>-  IntraGroupM () ->-  m (IntraAcc, Stms GPU)-runIntraGroupM m = do-  scope <- castScope <$> askScope-  modifyNameSource $ \src ->-    let (((), kstms), src', acc) = runRWS (runBuilderT m scope) () src-     in ((acc, kstms), src')--parallelMin :: [SubExp] -> IntraGroupM ()-parallelMin ws =-  tell-    mempty-      { accMinPar = S.singleton ws,-        accAvailPar = S.singleton ws-      }--intraGroupBody :: Body SOACS -> IntraGroupM (Body GPU)-intraGroupBody body = do-  stms <- collectStms_ $ intraGroupStms $ bodyStms body-  pure $ mkBody stms $ bodyResult body--intraGroupLambda :: Lambda SOACS -> IntraGroupM (Lambda GPU)-intraGroupLambda lam =-  mkLambda (lambdaParams lam) $-    bodyBind =<< intraGroupBody (lambdaBody lam)--intraGroupWithAccInput :: WithAccInput SOACS -> IntraGroupM (WithAccInput GPU)-intraGroupWithAccInput (shape, arrs, Nothing) =-  pure (shape, arrs, Nothing)-intraGroupWithAccInput (shape, arrs, Just (lam, nes)) = do-  lam' <- intraGroupLambda lam-  pure (shape, arrs, Just (lam', nes))--intraGroupStm :: Stm SOACS -> IntraGroupM ()-intraGroupStm stm@(Let pat aux e) = do-  scope <- askScope-  let lvl = SegThread SegNoVirt Nothing--  case e of-    Loop merge form loopbody ->-      localScope (scopeOfLoopForm form <> scopeOfFParams (map fst merge)) $ do-        loopbody' <- intraGroupBody loopbody-        certifying (stmAuxCerts aux) . letBind pat $-          Loop merge form loopbody'-    Match cond cases defbody ifdec -> do-      cases' <- mapM (traverse intraGroupBody) cases-      defbody' <- intraGroupBody defbody-      certifying (stmAuxCerts aux) . letBind pat $-        Match cond cases' defbody' ifdec-    WithAcc inputs lam -> do-      inputs' <- mapM intraGroupWithAccInput inputs-      lam' <- intraGroupLambda lam-      certifying (stmAuxCerts aux) . letBind pat $ WithAcc inputs' lam'-    Op soac-      | "sequential_outer" `inAttrs` stmAuxAttrs aux ->-          intraGroupStms . fmap (certify (stmAuxCerts aux))-            =<< runBuilder_ (FOT.transformSOAC pat soac)-    Op (Screma w arrs form)-      | Just lam <- isMapSOAC form -> do-          let loopnest = MapNesting pat aux w $ zip (lambdaParams lam) arrs-              env =-                DistEnv-                  { distNest =-                      singleNesting $ Nesting mempty loopnest,-                    distScope =-                      scopeOfPat pat-                        <> scopeForGPU (scopeOf lam)-                        <> scope,-                    distOnInnerMap =-                      distributeMap,-                    distOnTopLevelStms =-                      liftInner . collectStms_ . intraGroupStms,-                    distSegLevel = \minw _ _ -> do-                      lift $ parallelMin minw-                      pure lvl,-                    distOnSOACSStms =-                      pure . oneStm . soacsStmToGPU,-                    distOnSOACSLambda =-                      pure . soacsLambdaToGPU-                  }-              acc =-                DistAcc-                  { distTargets = singleTarget (pat, bodyResult $ lambdaBody lam),-                    distStms = mempty-                  }--          addStms-            =<< runDistNestT env (distributeMapBodyStms acc (bodyStms $ lambdaBody lam))-    Op (Screma w arrs form)-      | Just (scans, mapfun) <- isScanomapSOAC form,-        -- FIXME: Futhark.CodeGen.ImpGen.GPU.Block.compileGroupOp-        -- cannot handle multiple scan operators yet.-        Scan scanfun nes <- singleScan scans -> do-          let scanfun' = soacsLambdaToGPU scanfun-              mapfun' = soacsLambdaToGPU mapfun-          certifying (stmAuxCerts aux) $-            addStms =<< segScan lvl pat mempty w [SegBinOp Noncommutative scanfun' nes mempty] mapfun' arrs [] []-          parallelMin [w]-    Op (Screma w arrs form)-      | Just (reds, map_lam) <- isRedomapSOAC form -> do-          let onReduce (Reduce comm red_lam nes) =-                SegBinOp comm (soacsLambdaToGPU red_lam) nes mempty-              reds' = map onReduce reds-              map_lam' = soacsLambdaToGPU map_lam-          certifying (stmAuxCerts aux) $-            addStms =<< segRed lvl pat mempty w reds' map_lam' arrs [] []-          parallelMin [w]-    Op (Screma w arrs form) ->-      -- This screma is too complicated for us to immediately do-      -- anything, so split it up and try again.-      mapM_ intraGroupStm . fmap (certify (stmAuxCerts aux)) . snd-        =<< runBuilderT (dissectScrema pat w form arrs) (scopeForSOACs scope)-    Op (Hist w arrs ops bucket_fun) -> do-      ops' <- forM ops $ \(HistOp num_bins rf dests nes op) -> do-        (op', nes', shape) <- determineReduceOp op nes-        let op'' = soacsLambdaToGPU op'-        pure $ GPU.HistOp num_bins rf dests nes' shape op''--      let bucket_fun' = soacsLambdaToGPU bucket_fun-      certifying (stmAuxCerts aux) $-        addStms =<< segHist lvl pat w [] [] ops' bucket_fun' arrs-      parallelMin [w]-    Op (Stream w arrs accs lam)-      | chunk_size_param : _ <- lambdaParams lam -> do-          types <- asksScope castScope-          ((), stream_stms) <--            runBuilderT (sequentialStreamWholeArray pat w accs lam arrs) types-          let replace (Var v) | v == paramName chunk_size_param = w-              replace se = se-              replaceSets (IntraAcc x y log) =-                IntraAcc (S.map (map replace) x) (S.map (map replace) y) log-          censor replaceSets $ intraGroupStms stream_stms-    Op (Scatter w ivs lam dests) -> do-      write_i <- newVName "write_i"-      space <- mkSegSpace [(write_i, w)]--      let lam' = soacsLambdaToGPU lam-          krets = do-            (_a_w, a, is_vs) <--              groupScatterResults dests $ bodyResult $ lambdaBody lam'-            let cs =-                  foldMap (foldMap resCerts . fst) is_vs-                    <> foldMap (resCerts . snd) is_vs-                is_vs' = [(Slice $ map (DimFix . resSubExp) is, resSubExp v) | (is, v) <- is_vs]-            pure $ WriteReturns cs a is_vs'-          inputs = do-            (p, p_a) <- zip (lambdaParams lam') ivs-            pure $ KernelInput (paramName p) (paramType p) p_a [Var write_i]--      kstms <- runBuilder_ $-        localScope (scopeOfSegSpace space) $ do-          mapM_ readKernelInput inputs-          addStms $ bodyStms $ lambdaBody lam'--      certifying (stmAuxCerts aux) $ do-        let body = KernelBody () kstms krets-        letBind pat $ Op $ SegOp $ SegMap lvl space (patTypes pat) body--      parallelMin [w]-    _ ->-      addStm $ soacsStmToGPU stm--intraGroupStms :: Stms SOACS -> IntraGroupM ()-intraGroupStms = mapM_ intraGroupStm--intraGroupParalleliseBody ::-  (MonadFreshNames m, HasScope GPU m) =>-  Body SOACS ->-  m ([[SubExp]], [[SubExp]], Log, KernelBody GPU)-intraGroupParalleliseBody body = do-  (IntraAcc min_ws avail_ws log, kstms) <--    runIntraGroupM $ intraGroupStms $ bodyStms body-  pure-    ( S.toList min_ws,-      S.toList avail_ws,-      log,-      KernelBody () kstms $ map ret $ bodyResult body-    )-  where-    ret (SubExpRes cs se) = Returns ResultMaySimplify cs se

@@ -214,7 +214,7 @@ -- | Parse a FutharkScript expression with normal whitespace handling. parseExpFromText :: FilePath -> T.Text -> Either T.Text Exp parseExpFromText f s =-  either (Left . T.pack . errorBundlePretty) Right $ parse (parseExp space) f s+  either (Left . T.pack . errorBundlePretty) Right $ parse (parseExp space <* eof) f s  readVar :: (MonadError T.Text m, MonadIO m) => Server -> VarName -> m V.Value readVar server v =@@ -514,25 +514,36 @@               throwError $                 "Function \""                   <> name-                  <> "\" expects arguments of types:\n"-                  <> prettyText (V.mkCompound $ map V.ValueAtom in_types)-                  <> "\nBut called with arguments of types:\n"-                  <> prettyText (V.mkCompound $ map V.ValueAtom es_types)+                  <> "\" expects "+                  <> prettyText (length in_types)+                  <> " argument(s) of types:\n"+                  <> T.intercalate "\n" (map prettyTextOneLine in_types)+                  <> "\nBut applied to "+                  <> prettyText (length es_types)+                  <> " argument(s) of types:\n"+                  <> T.intercalate "\n" (map prettyTextOneLine es_types) +            tryApply args = do+              arg_types <- zipWithM (interValToVar cannotApply) in_types args++              if length in_types == length arg_types+                then do+                  outs <- replicateM (length out_types) $ newVar "out"+                  void $ cmdEither $ cmdCall server name outs arg_types+                  pure $ V.mkCompound $ map V.ValueAtom $ zipWith SValue out_types $ map VVar outs+                else+                  pure . V.ValueAtom . SFun name in_types out_types $+                    zipWith SValue in_types $+                      map VVar arg_types+         -- Careful to not require saturated application, but do still         -- check for over-saturation.         when (length es_types > length in_types) cannotApply-        ins <- zipWithM (interValToVar cannotApply) in_types es' -        if length in_types == length ins-          then do-            outs <- replicateM (length out_types) $ newVar "out"-            void $ cmdEither $ cmdCall server name outs ins-            pure $ V.mkCompound $ map V.ValueAtom $ zipWith SValue out_types $ map VVar outs-          else-            pure . V.ValueAtom . SFun name in_types out_types $-              zipWith SValue in_types $-                map VVar ins+        -- Allow automatic uncurrying if applicable.+        case es' of+          [V.ValueTuple es''] | length es'' == length in_types -> tryApply es''+          _ -> tryApply es'       evalExp' _ (StringLit s) =         case V.putValue s of           Just s' ->

@@ -390,6 +390,8 @@   Array (als1 <> aliases t2) et1 shape1 combineAliases (Scalar (TypeVar als1 tv1 targs1)) t2 =   Scalar $ TypeVar (als1 <> aliases t2) tv1 targs1+combineAliases t1 (Scalar (TypeVar als2 tv2 targs2)) =+  Scalar $ TypeVar (als2 <> aliases t1) tv2 targs2 combineAliases (Scalar (Record ts1)) (Scalar (Record ts2))   | length ts1 == length ts2,     L.sort (M.keys ts1) == L.sort (M.keys ts2) =

@@ -406,12 +406,11 @@       e' <- checkExp e       et <- expType e'       es' <- mapM (unifies "type of first array element" et <=< checkExp) es-      et' <- normTypeFully et-      t <- arrayOfM loc et' (Shape [sizeFromInteger (genericLength all_es) mempty])+      t <- arrayOfM loc et (Shape [sizeFromInteger (genericLength all_es) mempty])       pure $ ArrayLit (e' : es') (Info t) loc checkExp (AppExp (Range start maybe_step end loc) _) = do   start' <- require "use in range expression" anySignedType =<< checkExp start-  start_t <- expTypeFully start'+  start_t <- expType start'   maybe_step' <- case maybe_step of     Nothing -> pure Nothing     Just step -> do@@ -850,12 +849,12 @@ checkExp (Constr name es NoInfo loc) = do   t <- newTypeVar loc "t"   es' <- mapM checkExp es-  ets <- mapM expTypeFully es'+  ets <- mapM expType es'   mustHaveConstr (mkUsage loc "use of constructor") name t ets   pure $ Constr name es' (Info t) loc checkExp (AppExp (Match e cs loc) _) = do   e' <- checkExp e-  mt <- expTypeFully e'+  mt <- expType e'   (cs', t, retext) <- checkCases mt cs   zeroOrderType     (mkUsage loc "being returned 'match'")

@@ -168,13 +168,13 @@     where       fs' = mconcat $ punctuate "." $ map pretty fs   pretty (CheckingRequired [expected] actual) =-    "Expression must must have type"+    "Expression must have type"       <+> pretty expected       <> "."         </> "Actual type:"         <+> align (pretty actual)   pretty (CheckingRequired expected actual) =-    "Type of expression must must be one of "+    "Type of expression must be one of "       <+> expected'       <> "."         </> "Actual type:"

@@ -194,17 +194,16 @@ checkPat' _ (Wildcard NoInfo loc) NoneInferred = do   t <- newTypeVar loc "t"   pure $ Wildcard (Info t) loc-checkPat' sizes (TuplePat ps loc) (Ascribed t)+checkPat' sizes p@(TuplePat ps loc) (Ascribed t)   | Just ts <- isTupleRecord t,     length ts == length ps =       TuplePat         <$> zipWithM (checkPat' sizes) ps (map Ascribed ts)         <*> pure loc-checkPat' sizes p@(TuplePat ps loc) (Ascribed t) = do-  ps_t <- replicateM (length ps) (newTypeVar loc "t")-  unify (mkUsage loc "matching a tuple pattern") (Scalar (tupleRecord ps_t)) (toStruct t)-  t' <- normTypeFully t-  checkPat' sizes p $ Ascribed t'+  | otherwise = do+      ps_t <- replicateM (length ps) (newTypeVar loc "t")+      unify (mkUsage loc "matching a tuple pattern") (Scalar (tupleRecord ps_t)) (toStruct t)+      checkPat' sizes p $ Ascribed $ toParam Observe $ Scalar $ tupleRecord ps_t checkPat' sizes (TuplePat ps loc) NoneInferred =   TuplePat <$> mapM (\p -> checkPat' sizes p NoneInferred) ps <*> pure loc checkPat' _ (RecordPat p_fs _) _@@ -214,23 +213,23 @@           </> "Did you mean"           <> dquotes (pretty (drop 1 (nameToString f)) <> "=_")           <> "?"-checkPat' sizes (RecordPat p_fs loc) (Ascribed (Scalar (Record t_fs)))-  | sort (map fst p_fs) == sort (M.keys t_fs) =-      RecordPat . M.toList <$> check <*> pure loc+checkPat' sizes p@(RecordPat p_fs loc) (Ascribed t)+  | Scalar (Record t_fs) <- t,+    sort (map fst p_fs) == sort (M.keys t_fs) =+      RecordPat . M.toList <$> check t_fs <*> pure loc+  | otherwise = do+      p_fs' <- traverse (const $ newTypeVar loc "t") $ M.fromList p_fs++      when (sort (M.keys p_fs') /= sort (map fst p_fs)) $+        typeError loc mempty $+          "Duplicate fields in record pattern" <+> pretty p <> "."++      unify (mkUsage loc "matching a record pattern") (Scalar (Record p_fs')) (toStruct t)+      checkPat' sizes p $ Ascribed $ toParam Observe $ Scalar (Record p_fs')   where-    check =+    check t_fs =       traverse (uncurry (checkPat' sizes)) $         M.intersectionWith (,) (M.fromList p_fs) (fmap Ascribed t_fs)-checkPat' sizes p@(RecordPat fields loc) (Ascribed t) = do-  fields' <- traverse (const $ newTypeVar loc "t") $ M.fromList fields--  when (sort (M.keys fields') /= sort (map fst fields)) $-    typeError loc mempty $-      "Duplicate fields in record pattern" <+> pretty p <> "."--  unify (mkUsage loc "matching a record pattern") (Scalar (Record fields')) (toStruct t)-  t' <- normTypeFully t-  checkPat' sizes p $ Ascribed t' checkPat' sizes (RecordPat fs loc) NoneInferred =   RecordPat . M.toList     <$> traverse (\p -> checkPat' sizes p NoneInferred) (M.fromList fs)@@ -279,8 +278,7 @@     checkPat' sizes p $ Ascribed p_t   mustHaveConstr usage n (toStruct t') (patternStructType <$> ps')   unify usage t' (toStruct t)-  t'' <- normTypeFully t-  pure $ PatConstr n (Info t'') ps' loc+  pure $ PatConstr n (Info t) ps' loc   where     usage = mkUsage loc "matching against constructor" checkPat' sizes (PatConstr n NoInfo ps loc) NoneInferred = do

@@ -0,0 +1,20 @@+{-# OPTIONS_GHC -fno-warn-orphans #-}++module Futhark.IR.GPUTests () where++import Data.String+import Futhark.IR.GPU+import Futhark.IR.Parse+import Futhark.IR.SyntaxTests (parseString)++-- There isn't anything to test in this module, but we define some+-- convenience instances.++instance IsString (Stm GPU) where+  fromString = parseString "Stm GPU" parseStmGPU++instance IsString (Body GPU) where+  fromString = parseString "Body GPU" parseBodyGPU++instance IsString (Prog GPU) where+  fromString = parseString "Prog GPU" parseGPU

@@ -0,0 +1,20 @@+{-# OPTIONS_GHC -fno-warn-orphans #-}++module Futhark.IR.MCTests () where++import Data.String+import Futhark.IR.MC+import Futhark.IR.Parse+import Futhark.IR.SyntaxTests (parseString)++-- There isn't anything to test in this module, but we define some+-- convenience instances.++instance IsString (Stm MC) where+  fromString = parseString "Stm MC" parseStmMC++instance IsString (Body MC) where+  fromString = parseString "Body MC" parseBodyMC++instance IsString (Prog MC) where+  fromString = parseString "Prog MC" parseMC

@@ -1,6 +1,6 @@ {-# OPTIONS_GHC -fno-warn-orphans #-} -module Futhark.IR.SyntaxTests () where+module Futhark.IR.SyntaxTests (parseString) where  import Data.String import Data.Text qualified as T@@ -10,14 +10,24 @@ -- There isn't anything to test in this module, but we define some -- convenience instances. +parseString :: String -> (FilePath -> T.Text -> Either T.Text a) -> String -> a+parseString desc p =+  either (error . T.unpack) id . p ("IsString " <> desc) . T.pack++instance IsString Type where+  fromString = parseString "Type" parseType+ instance IsString DeclExtType where-  fromString =-    either (error . T.unpack) id-      . parseDeclExtType "IsString DeclExtType"-      . T.pack+  fromString = parseString "DeclExtType" parseDeclExtType  instance IsString DeclType where-  fromString =-    either (error . T.unpack) id-      . parseDeclType "IsString DeclType"-      . T.pack+  fromString = parseString "DeclType" parseDeclType++instance IsString VName where+  fromString = parseString "VName" parseVName++instance IsString SubExp where+  fromString = parseString "SubExp" parseSubExp++instance IsString SubExpRes where+  fromString = parseString "SubExpRes" parseSubExpRes