diff --git a/CHANGELOG.md b/CHANGELOG.md --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -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 diff --git a/docs/conf.py b/docs/conf.py --- a/docs/conf.py +++ b/docs/conf.py @@ -519,6 +519,13 @@ 1, ), ( + "man/futhark-script", + "futhark-script", + "execute FutharkScript expression", + [], + 1, + ), + ( "man/futhark-profile", "futhark-profile", "profile Futhark programs", diff --git a/docs/index.rst b/docs/index.rst --- a/docs/index.rst +++ b/docs/index.rst @@ -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 diff --git a/docs/man/futhark-literate.rst b/docs/man/futhark-literate.rst --- a/docs/man/futhark-literate.rst +++ b/docs/man/futhark-literate.rst @@ -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)` diff --git a/docs/man/futhark-script.rst b/docs/man/futhark-script.rst new file mode 100644 --- /dev/null +++ b/docs/man/futhark-script.rst @@ -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)` diff --git a/futhark.cabal b/futhark.cabal --- a/futhark.cabal +++ b/futhark.cabal @@ -1,6 +1,6 @@ cabal-version: 2.4 name: futhark -version: 0.25.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 diff --git a/rts/c/backends/cuda.h b/rts/c/backends/cuda.h --- a/rts/c/backends/cuda.h +++ b/rts/c/backends/cuda.h @@ -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], diff --git a/rts/c/backends/hip.h b/rts/c/backends/hip.h --- a/rts/c/backends/hip.h +++ b/rts/c/backends/hip.h @@ -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], diff --git a/rts/c/backends/opencl.h b/rts/c/backends/opencl.h --- a/rts/c/backends/opencl.h +++ b/rts/c/backends/opencl.h @@ -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; } diff --git a/rts/python/opencl.py b/rts/python/opencl.py --- a/rts/python/opencl.py +++ b/rts/python/opencl.py @@ -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( diff --git a/src/Futhark/CLI/Literate.hs b/src/Futhark/CLI/Literate.hs --- a/src/Futhark/CLI/Literate.hs +++ b/src/Futhark/CLI/Literate.hs @@ -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, diff --git a/src/Futhark/CLI/Main.hs b/src/Futhark/CLI/Main.hs --- a/src/Futhark/CLI/Main.hs +++ b/src/Futhark/CLI/Main.hs @@ -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.")), diff --git a/src/Futhark/CLI/REPL.hs b/src/Futhark/CLI/REPL.hs --- a/src/Futhark/CLI/REPL.hs +++ b/src/Futhark/CLI/REPL.hs @@ -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 () diff --git a/src/Futhark/CLI/Script.hs b/src/Futhark/CLI/Script.hs new file mode 100644 --- /dev/null +++ b/src/Futhark/CLI/Script.hs @@ -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 diff --git a/src/Futhark/CodeGen/Backends/GenericC/Code.hs b/src/Futhark/CodeGen/Backends/GenericC/Code.hs --- a/src/Futhark/CodeGen/Backends/GenericC/Code.hs +++ b/src/Futhark/CodeGen/Backends/GenericC/Code.hs @@ -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 diff --git a/src/Futhark/CodeGen/Backends/PyOpenCL.hs b/src/Futhark/CodeGen/Backends/PyOpenCL.hs --- a/src/Futhark/CodeGen/Backends/PyOpenCL.hs +++ b/src/Futhark/CodeGen/Backends/PyOpenCL.hs @@ -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 $ diff --git a/src/Futhark/CodeGen/Backends/PyOpenCL/Boilerplate.hs b/src/Futhark/CodeGen/Backends/PyOpenCL/Boilerplate.hs --- a/src/Futhark/CodeGen/Backends/PyOpenCL/Boilerplate.hs +++ b/src/Futhark/CodeGen/Backends/PyOpenCL/Boilerplate.hs @@ -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" diff --git a/src/Futhark/CodeGen/ImpCode.hs b/src/Futhark/CodeGen/ImpCode.hs --- a/src/Futhark/CodeGen/ImpCode.hs +++ b/src/Futhark/CodeGen/ImpCode.hs @@ -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 diff --git a/src/Futhark/CodeGen/ImpCode/GPU.hs b/src/Futhark/CodeGen/ImpCode/GPU.hs --- a/src/Futhark/CodeGen/ImpCode/GPU.hs +++ b/src/Futhark/CodeGen/ImpCode/GPU.hs @@ -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) = diff --git a/src/Futhark/CodeGen/ImpGen/GPU.hs b/src/Futhark/CodeGen/ImpGen/GPU.hs --- a/src/Futhark/CodeGen/ImpGen/GPU.hs +++ b/src/Futhark/CodeGen/ImpGen/GPU.hs @@ -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 diff --git a/src/Futhark/CodeGen/ImpGen/GPU/Base.hs b/src/Futhark/CodeGen/ImpGen/GPU/Base.hs --- a/src/Futhark/CodeGen/ImpGen/GPU/Base.hs +++ b/src/Futhark/CodeGen/ImpGen/GPU/Base.hs @@ -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, diff --git a/src/Futhark/CodeGen/ImpGen/GPU/Block.hs b/src/Futhark/CodeGen/ImpGen/GPU/Block.hs --- a/src/Futhark/CodeGen/ImpGen/GPU/Block.hs +++ b/src/Futhark/CodeGen/ImpGen/GPU/Block.hs @@ -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). diff --git a/src/Futhark/CodeGen/ImpGen/GPU/SegHist.hs b/src/Futhark/CodeGen/ImpGen/GPU/SegHist.hs --- a/src/Futhark/CodeGen/ImpGen/GPU/SegHist.hs +++ b/src/Futhark/CodeGen/ImpGen/GPU/SegHist.hs @@ -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 diff --git a/src/Futhark/CodeGen/ImpGen/GPU/SegRed.hs b/src/Futhark/CodeGen/ImpGen/GPU/SegRed.hs --- a/src/Futhark/CodeGen/ImpGen/GPU/SegRed.hs +++ b/src/Futhark/CodeGen/ImpGen/GPU/SegRed.hs @@ -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. diff --git a/src/Futhark/CodeGen/ImpGen/GPU/SegScan/TwoPass.hs b/src/Futhark/CodeGen/ImpGen/GPU/SegScan/TwoPass.hs --- a/src/Futhark/CodeGen/ImpGen/GPU/SegScan/TwoPass.hs +++ b/src/Futhark/CodeGen/ImpGen/GPU/SegScan/TwoPass.hs @@ -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) -> diff --git a/src/Futhark/CodeGen/ImpGen/GPU/ToOpenCL.hs b/src/Futhark/CodeGen/ImpGen/GPU/ToOpenCL.hs --- a/src/Futhark/CodeGen/ImpGen/GPU/ToOpenCL.hs +++ b/src/Futhark/CodeGen/ImpGen/GPU/ToOpenCL.hs @@ -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} diff --git a/src/Futhark/IR/GPUMem.hs b/src/Futhark/IR/GPUMem.hs --- a/src/Futhark/IR/GPUMem.hs +++ b/src/Futhark/IR/GPUMem.hs @@ -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 diff --git a/src/Futhark/IR/Mem/Simplify.hs b/src/Futhark/IR/Mem/Simplify.hs --- a/src/Futhark/IR/Mem/Simplify.hs +++ b/src/Futhark/IR/Mem/Simplify.hs @@ -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 diff --git a/src/Futhark/IR/Parse.hs b/src/Futhark/IR/Parse.hs --- a/src/Futhark/IR/Parse.hs +++ b/src/Futhark/IR/Parse.hs @@ -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 diff --git a/src/Futhark/IR/SegOp.hs b/src/Futhark/IR/SegOp.hs --- a/src/Futhark/IR/SegOp.hs +++ b/src/Futhark/IR/SegOp.hs @@ -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 <- diff --git a/src/Futhark/Internalise/Exps.hs b/src/Futhark/Internalise/Exps.hs --- a/src/Futhark/Internalise/Exps.hs +++ b/src/Futhark/Internalise/Exps.hs @@ -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') diff --git a/src/Futhark/Optimise/ArrayShortCircuiting/ArrayCoalescing.hs b/src/Futhark/Optimise/ArrayShortCircuiting/ArrayCoalescing.hs --- a/src/Futhark/Optimise/ArrayShortCircuiting/ArrayCoalescing.hs +++ b/src/Futhark/Optimise/ArrayShortCircuiting/ArrayCoalescing.hs @@ -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 diff --git a/src/Futhark/Optimise/BlkRegTiling.hs b/src/Futhark/Optimise/BlkRegTiling.hs --- a/src/Futhark/Optimise/BlkRegTiling.hs +++ b/src/Futhark/Optimise/BlkRegTiling.hs @@ -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 diff --git a/src/Futhark/Optimise/TileLoops.hs b/src/Futhark/Optimise/TileLoops.hs --- a/src/Futhark/Optimise/TileLoops.hs +++ b/src/Futhark/Optimise/TileLoops.hs @@ -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 diff --git a/src/Futhark/Pass/ExpandAllocations.hs b/src/Futhark/Pass/ExpandAllocations.hs --- a/src/Futhark/Pass/ExpandAllocations.hs +++ b/src/Futhark/Pass/ExpandAllocations.hs @@ -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 -> diff --git a/src/Futhark/Pass/ExplicitAllocations.hs b/src/Futhark/Pass/ExplicitAllocations.hs --- a/src/Futhark/Pass/ExplicitAllocations.hs +++ b/src/Futhark/Pass/ExplicitAllocations.hs @@ -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. diff --git a/src/Futhark/Pass/ExtractKernels.hs b/src/Futhark/Pass/ExtractKernels.hs --- a/src/Futhark/Pass/ExtractKernels.hs +++ b/src/Futhark/Pass/ExtractKernels.hs @@ -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 diff --git a/src/Futhark/Pass/ExtractKernels/DistributeNests.hs b/src/Futhark/Pass/ExtractKernels/DistributeNests.hs --- a/src/Futhark/Pass/ExtractKernels/DistributeNests.hs +++ b/src/Futhark/Pass/ExtractKernels/DistributeNests.hs @@ -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 diff --git a/src/Futhark/Pass/ExtractKernels/Intrablock.hs b/src/Futhark/Pass/ExtractKernels/Intrablock.hs new file mode 100644 --- /dev/null +++ b/src/Futhark/Pass/ExtractKernels/Intrablock.hs @@ -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 diff --git a/src/Futhark/Pass/ExtractKernels/Intragroup.hs b/src/Futhark/Pass/ExtractKernels/Intragroup.hs deleted file mode 100644 --- a/src/Futhark/Pass/ExtractKernels/Intragroup.hs +++ /dev/null @@ -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 diff --git a/src/Futhark/Script.hs b/src/Futhark/Script.hs --- a/src/Futhark/Script.hs +++ b/src/Futhark/Script.hs @@ -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' -> diff --git a/src/Language/Futhark/TypeChecker/Consumption.hs b/src/Language/Futhark/TypeChecker/Consumption.hs --- a/src/Language/Futhark/TypeChecker/Consumption.hs +++ b/src/Language/Futhark/TypeChecker/Consumption.hs @@ -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) = diff --git a/src/Language/Futhark/TypeChecker/Terms.hs b/src/Language/Futhark/TypeChecker/Terms.hs --- a/src/Language/Futhark/TypeChecker/Terms.hs +++ b/src/Language/Futhark/TypeChecker/Terms.hs @@ -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'") diff --git a/src/Language/Futhark/TypeChecker/Terms/Monad.hs b/src/Language/Futhark/TypeChecker/Terms/Monad.hs --- a/src/Language/Futhark/TypeChecker/Terms/Monad.hs +++ b/src/Language/Futhark/TypeChecker/Terms/Monad.hs @@ -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:" diff --git a/src/Language/Futhark/TypeChecker/Terms/Pat.hs b/src/Language/Futhark/TypeChecker/Terms/Pat.hs --- a/src/Language/Futhark/TypeChecker/Terms/Pat.hs +++ b/src/Language/Futhark/TypeChecker/Terms/Pat.hs @@ -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 diff --git a/unittests/Futhark/IR/GPUTests.hs b/unittests/Futhark/IR/GPUTests.hs new file mode 100644 --- /dev/null +++ b/unittests/Futhark/IR/GPUTests.hs @@ -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 diff --git a/unittests/Futhark/IR/MCTests.hs b/unittests/Futhark/IR/MCTests.hs new file mode 100644 --- /dev/null +++ b/unittests/Futhark/IR/MCTests.hs @@ -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 diff --git a/unittests/Futhark/IR/SyntaxTests.hs b/unittests/Futhark/IR/SyntaxTests.hs --- a/unittests/Futhark/IR/SyntaxTests.hs +++ b/unittests/Futhark/IR/SyntaxTests.hs @@ -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