futhark 0.9.1 → 0.10.1

raw patch · 115 files changed

+3602/−8917 lines, 115 filesdep +utf8-stringdep −HUnitdep −bifunctorsdep −extradep ~basedep ~containersdep ~megaparsecPVP ok

version bump matches the API change (PVP)

Dependencies added: utf8-string

Dependencies removed: HUnit, bifunctors, extra, raw-strings-qq, semigroups, th-lift-instances

Dependency ranges changed: base, containers, megaparsec, mtl, parser-combinators, text

API changes (from Hackage documentation)

Files

• futhark.cabal +26/−838
• futlib/prelude.fut +7/−0
• futlib/soacs.fut +4/−4
• rts/c/cuda.h +35/−11
• rts/c/opencl.h +38/−28
• rts/c/tuning.h +40/−0
• rts/csharp/opencl.cs +38/−12
• rts/python/opencl.py +21/−1
• rts/python/tuning.py +11/−0
• src/Futhark/CLI/Bench.hs +15/−3
• src/Futhark/CLI/Datacmp.hs +2/−4
• src/Futhark/CLI/Dev.hs +22/−10
• src/Futhark/CLI/Misc.hs +13/−0
• src/Futhark/CLI/Pkg.hs +4/−1
• src/Futhark/CLI/REPL.hs +1/−1
• src/Futhark/CLI/Test.hs +14/−7
• src/Futhark/CodeGen/Backends/CCUDA.hs +34/−11
• src/Futhark/CodeGen/Backends/CCUDA/Boilerplate.hs +16/−1
• src/Futhark/CodeGen/Backends/COpenCL.hs +40/−16
• src/Futhark/CodeGen/Backends/COpenCL/Boilerplate.hs +39/−2
• src/Futhark/CodeGen/Backends/CSOpenCL.hs +9/−3
• src/Futhark/CodeGen/Backends/CSOpenCL/Boilerplate.hs +1/−0
• src/Futhark/CodeGen/Backends/GenericC.hs +37/−27
• src/Futhark/CodeGen/Backends/GenericC/Options.hs +23/−8
• src/Futhark/CodeGen/Backends/GenericCSharp.hs +12/−5
• src/Futhark/CodeGen/Backends/GenericCSharp/AST.hs +3/−2
• src/Futhark/CodeGen/Backends/GenericPython.hs +21/−8
• src/Futhark/CodeGen/Backends/GenericPython/Definitions.hs +4/−0
• src/Futhark/CodeGen/Backends/GenericPython/Options.hs +6/−4
• src/Futhark/CodeGen/Backends/PyOpenCL.hs +62/−13
• src/Futhark/CodeGen/Backends/PyOpenCL/Boilerplate.hs +2/−0
• src/Futhark/CodeGen/Backends/SequentialPython.hs +1/−1
• src/Futhark/CodeGen/ImpCode.hs +16/−2
• src/Futhark/CodeGen/ImpCode/Kernels.hs +3/−3
• src/Futhark/CodeGen/ImpGen.hs +42/−24
• src/Futhark/CodeGen/ImpGen/Kernels.hs +9/−88
• src/Futhark/CodeGen/ImpGen/Kernels/Base.hs +188/−52
• src/Futhark/CodeGen/ImpGen/Kernels/SegGenRed.hs +275/−0
• src/Futhark/CodeGen/ImpGen/Kernels/SegRed.hs +70/−105
• src/Futhark/CodeGen/ImpGen/Kernels/ToOpenCL.hs +45/−35
• src/Futhark/CodeGen/OpenCL/Kernels.hs +3/−1
• src/Futhark/Construct.hs +17/−11
• src/Futhark/Doc/Generator.hs +3/−4
• src/Futhark/Internalise.hs +22/−91
• src/Futhark/Internalise/AccurateSizes.hs +19/−6
• src/Futhark/Internalise/Defunctionalise.hs +29/−65
• src/Futhark/Internalise/Monomorphise.hs +106/−63
• src/Futhark/Optimise/CSE.hs +32/−16
• src/Futhark/Optimise/InliningDeadFun.hs +4/−3
• src/Futhark/Optimise/MemoryBlockMerging.hs +0/−28
• src/Futhark/Optimise/MemoryBlockMerging/ActualVariables.hs +0/−358
• src/Futhark/Optimise/MemoryBlockMerging/AllExpVars.hs +0/−96
• src/Futhark/Optimise/MemoryBlockMerging/AuxiliaryInfo.hs +0/−63
• src/Futhark/Optimise/MemoryBlockMerging/Coalescing.hs +0/−31
• src/Futhark/Optimise/MemoryBlockMerging/Coalescing/AllocationMovingUp.hs +0/−94
• src/Futhark/Optimise/MemoryBlockMerging/Coalescing/Core.hs +0/−624
• src/Futhark/Optimise/MemoryBlockMerging/Coalescing/Exps.hs +0/−70
• src/Futhark/Optimise/MemoryBlockMerging/Coalescing/SafetyCondition2.hs +0/−110
• src/Futhark/Optimise/MemoryBlockMerging/Coalescing/SafetyCondition3.hs +0/−136
• src/Futhark/Optimise/MemoryBlockMerging/Coalescing/SafetyCondition5.hs +0/−120
• src/Futhark/Optimise/MemoryBlockMerging/CrudeMovingUp.hs +0/−263
• src/Futhark/Optimise/MemoryBlockMerging/Existentials.hs +0/−80
• src/Futhark/Optimise/MemoryBlockMerging/Liveness/FirstUse.hs +0/−198
• src/Futhark/Optimise/MemoryBlockMerging/Liveness/Interference.hs +0/−520
• src/Futhark/Optimise/MemoryBlockMerging/Liveness/LastUse.hs +0/−281
• src/Futhark/Optimise/MemoryBlockMerging/MemoryAliases.hs +0/−162
• src/Futhark/Optimise/MemoryBlockMerging/MemoryUpdater.hs +0/−418
• src/Futhark/Optimise/MemoryBlockMerging/Miscellaneous.hs +0/−263
• src/Futhark/Optimise/MemoryBlockMerging/PrimExps.hs +0/−105
• src/Futhark/Optimise/MemoryBlockMerging/Reuse.hs +0/−30
• src/Futhark/Optimise/MemoryBlockMerging/Reuse/AllocationSizeMovingUp.hs +0/−32
• src/Futhark/Optimise/MemoryBlockMerging/Reuse/AllocationSizeUses.hs +0/−127
• src/Futhark/Optimise/MemoryBlockMerging/Reuse/AllocationSizes.hs +0/−132
• src/Futhark/Optimise/MemoryBlockMerging/Reuse/Core.hs +0/−747
• src/Futhark/Optimise/MemoryBlockMerging/Types.hs +0/−89
• src/Futhark/Optimise/MemoryBlockMerging/VariableAliases.hs +0/−82
• src/Futhark/Optimise/MemoryBlockMerging/VariableMemory.hs +0/−99
• src/Futhark/Pass/ExpandAllocations.hs +182/−81
• src/Futhark/Pass/ExplicitAllocations.hs +7/−0
• src/Futhark/Pass/ExtractKernels.hs +27/−107
• src/Futhark/Pass/ExtractKernels/BlockedKernel.hs +24/−142
• src/Futhark/Pass/ExtractKernels/Segmented.hs +1/−1
• src/Futhark/Pass/ExtractKernels/Split.hs +0/−41
• src/Futhark/Passes.hs +2/−50
• src/Futhark/Representation/AST/Attributes/Types.hs +11/−10
• src/Futhark/Representation/Aliases.hs +0/−1
• src/Futhark/Representation/ExplicitMemory.hs +5/−3
• src/Futhark/Representation/ExplicitMemory/IndexFunction.hs +789/−447
• src/Futhark/Representation/ExplicitMemory/Lmad.hs +0/−761
• src/Futhark/Representation/Kernels/Kernel.hs +95/−0
• src/Futhark/Representation/Kernels/Simplify.hs +61/−11
• src/Futhark/Test.hs +25/−6
• src/Futhark/Transform/FirstOrderTransform.hs +26/−13
• src/Futhark/Version.hs +7/−2
• src/Language/Futhark/Attributes.hs +0/−3
• src/Language/Futhark/Interpreter.hs +120/−107
• src/Language/Futhark/Parser/Parser.y +11/−12
• src/Language/Futhark/Pretty.hs +1/−4
• src/Language/Futhark/Syntax.hs +0/−10
• src/Language/Futhark/Traversals.hs +0/−5
• src/Language/Futhark/TypeChecker.hs +11/−5
• src/Language/Futhark/TypeChecker/Monad.hs +16/−4
• src/Language/Futhark/TypeChecker/Terms.hs +58/−71
• src/Language/Futhark/TypeChecker/Types.hs +9/−8
• src/Language/Futhark/TypeChecker/Unify.hs +2/−4
• src/futhark.hs +1/−0
• src/futharki.hs +0/−19
• src/wrapper.hs +0/−29
• unittests/Futhark/Analysis/ScalExpTests.hs +2/−6
• unittests/Futhark/Optimise/AlgSimplifyTests.hs +5/−5
• unittests/Futhark/Representation/AST/AttributesTests.hs +0/−1
• unittests/Futhark/Representation/ExplicitMemory/IndexFunction/Alg.hs +178/−0
• unittests/Futhark/Representation/ExplicitMemory/IndexFunctionTests.hs +371/−0
• unittests/Futhark/Representation/ExplicitMemory/IndexFunctionWrapper.hs +65/−0
• unittests/futhark_tests.hs +6/−0

@@ -2,12 +2,19 @@ -- -- see: https://github.com/sol/hpack ----- hash: d926568e0952a1cecfa1efa7cf687ddc7cd27d1c642a929d85ca5b0e8465b408+-- hash: ca599dc882c25837255ce48c8e6ab8c5a9f1ff05234e05f6d73705f5cf9fe1c6  name:           futhark-version:        0.9.1+version:        0.10.1 synopsis:       An optimising compiler for a functional, array-oriented language.-description:    See the website at https://futhark-lang.org+description:    Futhark is a small programming language designed to be compiled to+                efficient parallel code. It is a statically typed, data-parallel,+                and purely functional array language in the ML family, and comes+                with a heavily optimising ahead-of-time compiler that presently+                generates GPU code via CUDA and OpenCL, although the language itself+                is hardware-agnostic.+                .+                For more information, see the website at https://futhark-lang.org category:       Language homepage:       https://futhark-lang.org bug-reports:    https://github.com/diku-dk/futhark/issues@@ -29,6 +36,7 @@     rts/c/opencl.h     rts/c/panic.h     rts/c/timing.h+    rts/c/tuning.h     rts/c/values.h     rts/csharp/exceptions.cs     rts/csharp/functions.cs@@ -43,6 +51,7 @@     rts/python/opencl.py     rts/python/panic.py     rts/python/scalar.py+    rts/python/tuning.py     rts/python/values.py  source-repository head@@ -58,7 +67,6 @@     , ansi-terminal >=0.6.3.1     , array >=0.4     , base >=4 && <5-    , bifunctors >=5.4.2     , binary >=0.8.3     , blaze-html >=0.9.0.1     , bytestring >=0.10.8@@ -67,7 +75,6 @@     , directory >=1.3.0.0     , directory-tree >=0.12.1     , dlist >=0.6.0.1-    , extra >=1.5.3     , file-embed >=0.0.9     , filepath >=1.4.1.1     , free >=4.12.4@@ -87,15 +94,14 @@     , process >=1.4.3.0     , process-extras >=0.7.2     , random-    , raw-strings-qq >=1.1     , regex-tdfa >=1.2     , srcloc >=0.4     , template-haskell >=2.11.1     , temporary     , text >=1.2.2.2-    , th-lift-instances >=0.1.11     , time >=1.6.0.1     , transformers >=0.3+    , utf8-string >=1     , vector >=0.12     , vector-binary-instances >=0.2.2.0     , versions >=3.3.1@@ -104,9 +110,6 @@   build-tools:       alex     , happy-  if !impl(ghc >= 8.0)-    build-depends:-        semigroups ==0.18.*   exposed-modules:       Futhark.Actions       Futhark.Analysis.AlgSimplify@@ -174,6 +177,7 @@       Futhark.CodeGen.ImpGen.CUDA       Futhark.CodeGen.ImpGen.Kernels       Futhark.CodeGen.ImpGen.Kernels.Base+      Futhark.CodeGen.ImpGen.Kernels.SegGenRed       Futhark.CodeGen.ImpGen.Kernels.SegRed       Futhark.CodeGen.ImpGen.Kernels.ToOpenCL       Futhark.CodeGen.ImpGen.Kernels.Transpose@@ -209,34 +213,6 @@       Futhark.Optimise.InPlaceLowering       Futhark.Optimise.InPlaceLowering.LowerIntoStm       Futhark.Optimise.InPlaceLowering.SubstituteIndices-      Futhark.Optimise.MemoryBlockMerging-      Futhark.Optimise.MemoryBlockMerging.ActualVariables-      Futhark.Optimise.MemoryBlockMerging.AllExpVars-      Futhark.Optimise.MemoryBlockMerging.AuxiliaryInfo-      Futhark.Optimise.MemoryBlockMerging.Coalescing-      Futhark.Optimise.MemoryBlockMerging.Coalescing.AllocationMovingUp-      Futhark.Optimise.MemoryBlockMerging.Coalescing.Core-      Futhark.Optimise.MemoryBlockMerging.Coalescing.Exps-      Futhark.Optimise.MemoryBlockMerging.Coalescing.SafetyCondition2-      Futhark.Optimise.MemoryBlockMerging.Coalescing.SafetyCondition3-      Futhark.Optimise.MemoryBlockMerging.Coalescing.SafetyCondition5-      Futhark.Optimise.MemoryBlockMerging.CrudeMovingUp-      Futhark.Optimise.MemoryBlockMerging.Existentials-      Futhark.Optimise.MemoryBlockMerging.Liveness.FirstUse-      Futhark.Optimise.MemoryBlockMerging.Liveness.Interference-      Futhark.Optimise.MemoryBlockMerging.Liveness.LastUse-      Futhark.Optimise.MemoryBlockMerging.MemoryAliases-      Futhark.Optimise.MemoryBlockMerging.MemoryUpdater-      Futhark.Optimise.MemoryBlockMerging.Miscellaneous-      Futhark.Optimise.MemoryBlockMerging.PrimExps-      Futhark.Optimise.MemoryBlockMerging.Reuse-      Futhark.Optimise.MemoryBlockMerging.Reuse.AllocationSizeMovingUp-      Futhark.Optimise.MemoryBlockMerging.Reuse.AllocationSizes-      Futhark.Optimise.MemoryBlockMerging.Reuse.AllocationSizeUses-      Futhark.Optimise.MemoryBlockMerging.Reuse.Core-      Futhark.Optimise.MemoryBlockMerging.Types-      Futhark.Optimise.MemoryBlockMerging.VariableAliases-      Futhark.Optimise.MemoryBlockMerging.VariableMemory       Futhark.Optimise.Simplify       Futhark.Optimise.Simplify.ClosedForm       Futhark.Optimise.Simplify.Engine@@ -257,7 +233,6 @@       Futhark.Pass.ExtractKernels.ISRWIM       Futhark.Pass.ExtractKernels.Kernelise       Futhark.Pass.ExtractKernels.Segmented-      Futhark.Pass.ExtractKernels.Split       Futhark.Pass.FirstOrderTransform       Futhark.Pass.KernelBabysitting       Futhark.Pass.ResolveAssertions@@ -288,7 +263,6 @@       Futhark.Representation.AST.Traversals       Futhark.Representation.ExplicitMemory       Futhark.Representation.ExplicitMemory.IndexFunction-      Futhark.Representation.ExplicitMemory.Lmad       Futhark.Representation.ExplicitMemory.Simplify       Futhark.Representation.Kernels       Futhark.Representation.Kernels.Kernel@@ -341,757 +315,13 @@   main-is: src/futhark.hs   ghc-options: -Wall -Wcompat -Wredundant-constraints -Wincomplete-record-updates -Wmissing-export-lists -threaded -rtsopts "-with-rtsopts=-N -qg"   build-depends:-      aeson-    , ansi-terminal >=0.6.3.1-    , array >=0.4-    , base >=4 && <5-    , bifunctors >=5.4.2-    , binary >=0.8.3-    , blaze-html >=0.9.0.1-    , bytestring >=0.10.8-    , containers >=0.5-    , data-binary-ieee754 >=0.1-    , directory >=1.3.0.0-    , directory-tree >=0.12.1-    , dlist >=0.6.0.1-    , extra >=1.5.3-    , file-embed >=0.0.9-    , filepath >=1.4.1.1-    , free >=4.12.4+      base     , futhark-    , gitrev >=1.2.0-    , haskeline-    , http-client >=0.5.7.0-    , http-client-tls >=0.3.5.1-    , http-conduit >=2.2.4-    , language-c-quote >=0.12-    , mainland-pretty >=0.6.1-    , markdown >=0.1.16-    , megaparsec >=7.0.1-    , mtl >=2.2.1-    , neat-interpolation >=0.3-    , parallel >=3.2.1.0-    , parser-combinators >=1.0.0-    , process >=1.4.3.0-    , process-extras >=0.7.2-    , random-    , raw-strings-qq >=1.1-    , regex-tdfa >=1.2-    , srcloc >=0.4-    , template-haskell >=2.11.1-    , temporary-    , text >=1.2.2.2-    , th-lift-instances >=0.1.11-    , time >=1.6.0.1-    , transformers >=0.3-    , vector >=0.12-    , vector-binary-instances >=0.2.2.0-    , versions >=3.3.1-    , zip-archive >=0.3.1.1-    , zlib >=0.6.1.2-  if !impl(ghc >= 8.0)-    build-depends:-        semigroups ==0.18.*+    , text   other-modules:       Paths_futhark   default-language: Haskell2010 -executable futhark-bench-  main-is: src/wrapper.hs-  other-modules:-      Paths_futhark-  ghc-options: -Wall -Wcompat -Wredundant-constraints -Wincomplete-record-updates -Wmissing-export-lists-  build-depends:-      aeson-    , ansi-terminal >=0.6.3.1-    , array >=0.4-    , base >=4 && <5-    , bifunctors >=5.4.2-    , binary >=0.8.3-    , blaze-html >=0.9.0.1-    , bytestring >=0.10.8-    , containers >=0.5-    , data-binary-ieee754 >=0.1-    , directory >=1.3.0.0-    , directory-tree >=0.12.1-    , dlist >=0.6.0.1-    , extra >=1.5.3-    , file-embed >=0.0.9-    , filepath >=1.4.1.1-    , free >=4.12.4-    , gitrev >=1.2.0-    , haskeline-    , http-client >=0.5.7.0-    , http-client-tls >=0.3.5.1-    , http-conduit >=2.2.4-    , language-c-quote >=0.12-    , mainland-pretty >=0.6.1-    , markdown >=0.1.16-    , megaparsec >=7.0.1-    , mtl >=2.2.1-    , neat-interpolation >=0.3-    , parallel >=3.2.1.0-    , parser-combinators >=1.0.0-    , process >=1.4.3.0-    , process-extras >=0.7.2-    , random-    , raw-strings-qq >=1.1-    , regex-tdfa >=1.2-    , srcloc >=0.4-    , template-haskell >=2.11.1-    , temporary-    , text >=1.2.2.2-    , th-lift-instances >=0.1.11-    , time >=1.6.0.1-    , transformers >=0.3-    , vector >=0.12-    , vector-binary-instances >=0.2.2.0-    , versions >=3.3.1-    , zip-archive >=0.3.1.1-    , zlib >=0.6.1.2-  if !impl(ghc >= 8.0)-    build-depends:-        semigroups ==0.18.*-  default-language: Haskell2010--executable futhark-c-  main-is: src/wrapper.hs-  other-modules:-      Paths_futhark-  ghc-options: -Wall -Wcompat -Wredundant-constraints -Wincomplete-record-updates -Wmissing-export-lists-  build-depends:-      aeson-    , ansi-terminal >=0.6.3.1-    , array >=0.4-    , base >=4 && <5-    , bifunctors >=5.4.2-    , binary >=0.8.3-    , blaze-html >=0.9.0.1-    , bytestring >=0.10.8-    , containers >=0.5-    , data-binary-ieee754 >=0.1-    , directory >=1.3.0.0-    , directory-tree >=0.12.1-    , dlist >=0.6.0.1-    , extra >=1.5.3-    , file-embed >=0.0.9-    , filepath >=1.4.1.1-    , free >=4.12.4-    , gitrev >=1.2.0-    , haskeline-    , http-client >=0.5.7.0-    , http-client-tls >=0.3.5.1-    , http-conduit >=2.2.4-    , language-c-quote >=0.12-    , mainland-pretty >=0.6.1-    , markdown >=0.1.16-    , megaparsec >=7.0.1-    , mtl >=2.2.1-    , neat-interpolation >=0.3-    , parallel >=3.2.1.0-    , parser-combinators >=1.0.0-    , process >=1.4.3.0-    , process-extras >=0.7.2-    , random-    , raw-strings-qq >=1.1-    , regex-tdfa >=1.2-    , srcloc >=0.4-    , template-haskell >=2.11.1-    , temporary-    , text >=1.2.2.2-    , th-lift-instances >=0.1.11-    , time >=1.6.0.1-    , transformers >=0.3-    , vector >=0.12-    , vector-binary-instances >=0.2.2.0-    , versions >=3.3.1-    , zip-archive >=0.3.1.1-    , zlib >=0.6.1.2-  if !impl(ghc >= 8.0)-    build-depends:-        semigroups ==0.18.*-  default-language: Haskell2010--executable futhark-cs-  main-is: src/wrapper.hs-  other-modules:-      Paths_futhark-  ghc-options: -Wall -Wcompat -Wredundant-constraints -Wincomplete-record-updates -Wmissing-export-lists-  build-depends:-      aeson-    , ansi-terminal >=0.6.3.1-    , array >=0.4-    , base >=4 && <5-    , bifunctors >=5.4.2-    , binary >=0.8.3-    , blaze-html >=0.9.0.1-    , bytestring >=0.10.8-    , containers >=0.5-    , data-binary-ieee754 >=0.1-    , directory >=1.3.0.0-    , directory-tree >=0.12.1-    , dlist >=0.6.0.1-    , extra >=1.5.3-    , file-embed >=0.0.9-    , filepath >=1.4.1.1-    , free >=4.12.4-    , gitrev >=1.2.0-    , haskeline-    , http-client >=0.5.7.0-    , http-client-tls >=0.3.5.1-    , http-conduit >=2.2.4-    , language-c-quote >=0.12-    , mainland-pretty >=0.6.1-    , markdown >=0.1.16-    , megaparsec >=7.0.1-    , mtl >=2.2.1-    , neat-interpolation >=0.3-    , parallel >=3.2.1.0-    , parser-combinators >=1.0.0-    , process >=1.4.3.0-    , process-extras >=0.7.2-    , random-    , raw-strings-qq >=1.1-    , regex-tdfa >=1.2-    , srcloc >=0.4-    , template-haskell >=2.11.1-    , temporary-    , text >=1.2.2.2-    , th-lift-instances >=0.1.11-    , time >=1.6.0.1-    , transformers >=0.3-    , vector >=0.12-    , vector-binary-instances >=0.2.2.0-    , versions >=3.3.1-    , zip-archive >=0.3.1.1-    , zlib >=0.6.1.2-  if !impl(ghc >= 8.0)-    build-depends:-        semigroups ==0.18.*-  default-language: Haskell2010--executable futhark-csopencl-  main-is: src/wrapper.hs-  other-modules:-      Paths_futhark-  ghc-options: -Wall -Wcompat -Wredundant-constraints -Wincomplete-record-updates -Wmissing-export-lists-  build-depends:-      aeson-    , ansi-terminal >=0.6.3.1-    , array >=0.4-    , base >=4 && <5-    , bifunctors >=5.4.2-    , binary >=0.8.3-    , blaze-html >=0.9.0.1-    , bytestring >=0.10.8-    , containers >=0.5-    , data-binary-ieee754 >=0.1-    , directory >=1.3.0.0-    , directory-tree >=0.12.1-    , dlist >=0.6.0.1-    , extra >=1.5.3-    , file-embed >=0.0.9-    , filepath >=1.4.1.1-    , free >=4.12.4-    , gitrev >=1.2.0-    , haskeline-    , http-client >=0.5.7.0-    , http-client-tls >=0.3.5.1-    , http-conduit >=2.2.4-    , language-c-quote >=0.12-    , mainland-pretty >=0.6.1-    , markdown >=0.1.16-    , megaparsec >=7.0.1-    , mtl >=2.2.1-    , neat-interpolation >=0.3-    , parallel >=3.2.1.0-    , parser-combinators >=1.0.0-    , process >=1.4.3.0-    , process-extras >=0.7.2-    , random-    , raw-strings-qq >=1.1-    , regex-tdfa >=1.2-    , srcloc >=0.4-    , template-haskell >=2.11.1-    , temporary-    , text >=1.2.2.2-    , th-lift-instances >=0.1.11-    , time >=1.6.0.1-    , transformers >=0.3-    , vector >=0.12-    , vector-binary-instances >=0.2.2.0-    , versions >=3.3.1-    , zip-archive >=0.3.1.1-    , zlib >=0.6.1.2-  if !impl(ghc >= 8.0)-    build-depends:-        semigroups ==0.18.*-  default-language: Haskell2010--executable futhark-dataset-  main-is: src/wrapper.hs-  other-modules:-      Paths_futhark-  ghc-options: -Wall -Wcompat -Wredundant-constraints -Wincomplete-record-updates -Wmissing-export-lists-  build-depends:-      aeson-    , ansi-terminal >=0.6.3.1-    , array >=0.4-    , base >=4 && <5-    , bifunctors >=5.4.2-    , binary >=0.8.3-    , blaze-html >=0.9.0.1-    , bytestring >=0.10.8-    , containers >=0.5-    , data-binary-ieee754 >=0.1-    , directory >=1.3.0.0-    , directory-tree >=0.12.1-    , dlist >=0.6.0.1-    , extra >=1.5.3-    , file-embed >=0.0.9-    , filepath >=1.4.1.1-    , free >=4.12.4-    , gitrev >=1.2.0-    , haskeline-    , http-client >=0.5.7.0-    , http-client-tls >=0.3.5.1-    , http-conduit >=2.2.4-    , language-c-quote >=0.12-    , mainland-pretty >=0.6.1-    , markdown >=0.1.16-    , megaparsec >=7.0.1-    , mtl >=2.2.1-    , neat-interpolation >=0.3-    , parallel >=3.2.1.0-    , parser-combinators >=1.0.0-    , process >=1.4.3.0-    , process-extras >=0.7.2-    , random-    , raw-strings-qq >=1.1-    , regex-tdfa >=1.2-    , srcloc >=0.4-    , template-haskell >=2.11.1-    , temporary-    , text >=1.2.2.2-    , th-lift-instances >=0.1.11-    , time >=1.6.0.1-    , transformers >=0.3-    , vector >=0.12-    , vector-binary-instances >=0.2.2.0-    , versions >=3.3.1-    , zip-archive >=0.3.1.1-    , zlib >=0.6.1.2-  if !impl(ghc >= 8.0)-    build-depends:-        semigroups ==0.18.*-  default-language: Haskell2010--executable futhark-doc-  main-is: src/wrapper.hs-  other-modules:-      Paths_futhark-  ghc-options: -Wall -Wcompat -Wredundant-constraints -Wincomplete-record-updates -Wmissing-export-lists-  build-depends:-      aeson-    , ansi-terminal >=0.6.3.1-    , array >=0.4-    , base >=4 && <5-    , bifunctors >=5.4.2-    , binary >=0.8.3-    , blaze-html >=0.9.0.1-    , bytestring >=0.10.8-    , containers >=0.5-    , data-binary-ieee754 >=0.1-    , directory >=1.3.0.0-    , directory-tree >=0.12.1-    , dlist >=0.6.0.1-    , extra >=1.5.3-    , file-embed >=0.0.9-    , filepath >=1.4.1.1-    , free >=4.12.4-    , gitrev >=1.2.0-    , haskeline-    , http-client >=0.5.7.0-    , http-client-tls >=0.3.5.1-    , http-conduit >=2.2.4-    , language-c-quote >=0.12-    , mainland-pretty >=0.6.1-    , markdown >=0.1.16-    , megaparsec >=7.0.1-    , mtl >=2.2.1-    , neat-interpolation >=0.3-    , parallel >=3.2.1.0-    , parser-combinators >=1.0.0-    , process >=1.4.3.0-    , process-extras >=0.7.2-    , random-    , raw-strings-qq >=1.1-    , regex-tdfa >=1.2-    , srcloc >=0.4-    , template-haskell >=2.11.1-    , temporary-    , text >=1.2.2.2-    , th-lift-instances >=0.1.11-    , time >=1.6.0.1-    , transformers >=0.3-    , vector >=0.12-    , vector-binary-instances >=0.2.2.0-    , versions >=3.3.1-    , zip-archive >=0.3.1.1-    , zlib >=0.6.1.2-  if !impl(ghc >= 8.0)-    build-depends:-        semigroups ==0.18.*-  default-language: Haskell2010--executable futhark-opencl-  main-is: src/wrapper.hs-  other-modules:-      Paths_futhark-  ghc-options: -Wall -Wcompat -Wredundant-constraints -Wincomplete-record-updates -Wmissing-export-lists-  build-depends:-      aeson-    , ansi-terminal >=0.6.3.1-    , array >=0.4-    , base >=4 && <5-    , bifunctors >=5.4.2-    , binary >=0.8.3-    , blaze-html >=0.9.0.1-    , bytestring >=0.10.8-    , containers >=0.5-    , data-binary-ieee754 >=0.1-    , directory >=1.3.0.0-    , directory-tree >=0.12.1-    , dlist >=0.6.0.1-    , extra >=1.5.3-    , file-embed >=0.0.9-    , filepath >=1.4.1.1-    , free >=4.12.4-    , gitrev >=1.2.0-    , haskeline-    , http-client >=0.5.7.0-    , http-client-tls >=0.3.5.1-    , http-conduit >=2.2.4-    , language-c-quote >=0.12-    , mainland-pretty >=0.6.1-    , markdown >=0.1.16-    , megaparsec >=7.0.1-    , mtl >=2.2.1-    , neat-interpolation >=0.3-    , parallel >=3.2.1.0-    , parser-combinators >=1.0.0-    , process >=1.4.3.0-    , process-extras >=0.7.2-    , random-    , raw-strings-qq >=1.1-    , regex-tdfa >=1.2-    , srcloc >=0.4-    , template-haskell >=2.11.1-    , temporary-    , text >=1.2.2.2-    , th-lift-instances >=0.1.11-    , time >=1.6.0.1-    , transformers >=0.3-    , vector >=0.12-    , vector-binary-instances >=0.2.2.0-    , versions >=3.3.1-    , zip-archive >=0.3.1.1-    , zlib >=0.6.1.2-  if !impl(ghc >= 8.0)-    build-depends:-        semigroups ==0.18.*-  default-language: Haskell2010--executable futhark-pkg-  main-is: src/wrapper.hs-  other-modules:-      Paths_futhark-  ghc-options: -Wall -Wcompat -Wredundant-constraints -Wincomplete-record-updates -Wmissing-export-lists-  build-depends:-      aeson-    , ansi-terminal >=0.6.3.1-    , array >=0.4-    , base >=4 && <5-    , bifunctors >=5.4.2-    , binary >=0.8.3-    , blaze-html >=0.9.0.1-    , bytestring >=0.10.8-    , containers >=0.5-    , data-binary-ieee754 >=0.1-    , directory >=1.3.0.0-    , directory-tree >=0.12.1-    , dlist >=0.6.0.1-    , extra >=1.5.3-    , file-embed >=0.0.9-    , filepath >=1.4.1.1-    , free >=4.12.4-    , gitrev >=1.2.0-    , haskeline-    , http-client >=0.5.7.0-    , http-client-tls >=0.3.5.1-    , http-conduit >=2.2.4-    , language-c-quote >=0.12-    , mainland-pretty >=0.6.1-    , markdown >=0.1.16-    , megaparsec >=7.0.1-    , mtl >=2.2.1-    , neat-interpolation >=0.3-    , parallel >=3.2.1.0-    , parser-combinators >=1.0.0-    , process >=1.4.3.0-    , process-extras >=0.7.2-    , random-    , raw-strings-qq >=1.1-    , regex-tdfa >=1.2-    , srcloc >=0.4-    , template-haskell >=2.11.1-    , temporary-    , text >=1.2.2.2-    , th-lift-instances >=0.1.11-    , time >=1.6.0.1-    , transformers >=0.3-    , vector >=0.12-    , vector-binary-instances >=0.2.2.0-    , versions >=3.3.1-    , zip-archive >=0.3.1.1-    , zlib >=0.6.1.2-  if !impl(ghc >= 8.0)-    build-depends:-        semigroups ==0.18.*-  default-language: Haskell2010--executable futhark-py-  main-is: src/wrapper.hs-  other-modules:-      Paths_futhark-  ghc-options: -Wall -Wcompat -Wredundant-constraints -Wincomplete-record-updates -Wmissing-export-lists-  build-depends:-      aeson-    , ansi-terminal >=0.6.3.1-    , array >=0.4-    , base >=4 && <5-    , bifunctors >=5.4.2-    , binary >=0.8.3-    , blaze-html >=0.9.0.1-    , bytestring >=0.10.8-    , containers >=0.5-    , data-binary-ieee754 >=0.1-    , directory >=1.3.0.0-    , directory-tree >=0.12.1-    , dlist >=0.6.0.1-    , extra >=1.5.3-    , file-embed >=0.0.9-    , filepath >=1.4.1.1-    , free >=4.12.4-    , gitrev >=1.2.0-    , haskeline-    , http-client >=0.5.7.0-    , http-client-tls >=0.3.5.1-    , http-conduit >=2.2.4-    , language-c-quote >=0.12-    , mainland-pretty >=0.6.1-    , markdown >=0.1.16-    , megaparsec >=7.0.1-    , mtl >=2.2.1-    , neat-interpolation >=0.3-    , parallel >=3.2.1.0-    , parser-combinators >=1.0.0-    , process >=1.4.3.0-    , process-extras >=0.7.2-    , random-    , raw-strings-qq >=1.1-    , regex-tdfa >=1.2-    , srcloc >=0.4-    , template-haskell >=2.11.1-    , temporary-    , text >=1.2.2.2-    , th-lift-instances >=0.1.11-    , time >=1.6.0.1-    , transformers >=0.3-    , vector >=0.12-    , vector-binary-instances >=0.2.2.0-    , versions >=3.3.1-    , zip-archive >=0.3.1.1-    , zlib >=0.6.1.2-  if !impl(ghc >= 8.0)-    build-depends:-        semigroups ==0.18.*-  default-language: Haskell2010--executable futhark-pyopencl-  main-is: src/wrapper.hs-  other-modules:-      Paths_futhark-  ghc-options: -Wall -Wcompat -Wredundant-constraints -Wincomplete-record-updates -Wmissing-export-lists-  build-depends:-      aeson-    , ansi-terminal >=0.6.3.1-    , array >=0.4-    , base >=4 && <5-    , bifunctors >=5.4.2-    , binary >=0.8.3-    , blaze-html >=0.9.0.1-    , bytestring >=0.10.8-    , containers >=0.5-    , data-binary-ieee754 >=0.1-    , directory >=1.3.0.0-    , directory-tree >=0.12.1-    , dlist >=0.6.0.1-    , extra >=1.5.3-    , file-embed >=0.0.9-    , filepath >=1.4.1.1-    , free >=4.12.4-    , gitrev >=1.2.0-    , haskeline-    , http-client >=0.5.7.0-    , http-client-tls >=0.3.5.1-    , http-conduit >=2.2.4-    , language-c-quote >=0.12-    , mainland-pretty >=0.6.1-    , markdown >=0.1.16-    , megaparsec >=7.0.1-    , mtl >=2.2.1-    , neat-interpolation >=0.3-    , parallel >=3.2.1.0-    , parser-combinators >=1.0.0-    , process >=1.4.3.0-    , process-extras >=0.7.2-    , random-    , raw-strings-qq >=1.1-    , regex-tdfa >=1.2-    , srcloc >=0.4-    , template-haskell >=2.11.1-    , temporary-    , text >=1.2.2.2-    , th-lift-instances >=0.1.11-    , time >=1.6.0.1-    , transformers >=0.3-    , vector >=0.12-    , vector-binary-instances >=0.2.2.0-    , versions >=3.3.1-    , zip-archive >=0.3.1.1-    , zlib >=0.6.1.2-  if !impl(ghc >= 8.0)-    build-depends:-        semigroups ==0.18.*-  default-language: Haskell2010--executable futhark-test-  main-is: src/wrapper.hs-  other-modules:-      Paths_futhark-  ghc-options: -Wall -Wcompat -Wredundant-constraints -Wincomplete-record-updates -Wmissing-export-lists-  build-depends:-      aeson-    , ansi-terminal >=0.6.3.1-    , array >=0.4-    , base >=4 && <5-    , bifunctors >=5.4.2-    , binary >=0.8.3-    , blaze-html >=0.9.0.1-    , bytestring >=0.10.8-    , containers >=0.5-    , data-binary-ieee754 >=0.1-    , directory >=1.3.0.0-    , directory-tree >=0.12.1-    , dlist >=0.6.0.1-    , extra >=1.5.3-    , file-embed >=0.0.9-    , filepath >=1.4.1.1-    , free >=4.12.4-    , gitrev >=1.2.0-    , haskeline-    , http-client >=0.5.7.0-    , http-client-tls >=0.3.5.1-    , http-conduit >=2.2.4-    , language-c-quote >=0.12-    , mainland-pretty >=0.6.1-    , markdown >=0.1.16-    , megaparsec >=7.0.1-    , mtl >=2.2.1-    , neat-interpolation >=0.3-    , parallel >=3.2.1.0-    , parser-combinators >=1.0.0-    , process >=1.4.3.0-    , process-extras >=0.7.2-    , random-    , raw-strings-qq >=1.1-    , regex-tdfa >=1.2-    , srcloc >=0.4-    , template-haskell >=2.11.1-    , temporary-    , text >=1.2.2.2-    , th-lift-instances >=0.1.11-    , time >=1.6.0.1-    , transformers >=0.3-    , vector >=0.12-    , vector-binary-instances >=0.2.2.0-    , versions >=3.3.1-    , zip-archive >=0.3.1.1-    , zlib >=0.6.1.2-  if !impl(ghc >= 8.0)-    build-depends:-        semigroups ==0.18.*-  default-language: Haskell2010--executable futharki-  main-is: src/futharki.hs-  other-modules:-      Paths_futhark-  ghc-options: -Wall -Wcompat -Wredundant-constraints -Wincomplete-record-updates -Wmissing-export-lists-  build-depends:-      aeson-    , ansi-terminal >=0.6.3.1-    , array >=0.4-    , base >=4 && <5-    , bifunctors >=5.4.2-    , binary >=0.8.3-    , blaze-html >=0.9.0.1-    , bytestring >=0.10.8-    , containers >=0.5-    , data-binary-ieee754 >=0.1-    , directory >=1.3.0.0-    , directory-tree >=0.12.1-    , dlist >=0.6.0.1-    , extra >=1.5.3-    , file-embed >=0.0.9-    , filepath >=1.4.1.1-    , free >=4.12.4-    , gitrev >=1.2.0-    , haskeline-    , http-client >=0.5.7.0-    , http-client-tls >=0.3.5.1-    , http-conduit >=2.2.4-    , language-c-quote >=0.12-    , mainland-pretty >=0.6.1-    , markdown >=0.1.16-    , megaparsec >=7.0.1-    , mtl >=2.2.1-    , neat-interpolation >=0.3-    , parallel >=3.2.1.0-    , parser-combinators >=1.0.0-    , process >=1.4.3.0-    , process-extras >=0.7.2-    , random-    , raw-strings-qq >=1.1-    , regex-tdfa >=1.2-    , srcloc >=0.4-    , template-haskell >=2.11.1-    , temporary-    , text >=1.2.2.2-    , th-lift-instances >=0.1.11-    , time >=1.6.0.1-    , transformers >=0.3-    , vector >=0.12-    , vector-binary-instances >=0.2.2.0-    , versions >=3.3.1-    , zip-archive >=0.3.1.1-    , zlib >=0.6.1.2-  if !impl(ghc >= 8.0)-    build-depends:-        semigroups ==0.18.*-  default-language: Haskell2010- test-suite unit   type: exitcode-stdio-1.0   main-is: futhark_tests.hs@@ -1099,62 +329,17 @@       unittests   ghc-options: -Wall -Wcompat -Wredundant-constraints -Wincomplete-record-updates -Wmissing-export-lists   build-depends:-      HUnit-    , QuickCheck >=2.8-    , aeson-    , ansi-terminal >=0.6.3.1-    , array >=0.4-    , base >=4 && <5-    , bifunctors >=5.4.2-    , binary >=0.8.3-    , blaze-html >=0.9.0.1-    , bytestring >=0.10.8-    , containers >=0.5-    , data-binary-ieee754 >=0.1-    , directory >=1.3.0.0-    , directory-tree >=0.12.1-    , dlist >=0.6.0.1-    , extra >=1.5.3-    , file-embed >=0.0.9-    , filepath >=1.4.1.1-    , free >=4.12.4+      QuickCheck >=2.8+    , base+    , containers     , futhark-    , gitrev >=1.2.0-    , haskeline-    , http-client >=0.5.7.0-    , http-client-tls >=0.3.5.1-    , http-conduit >=2.2.4-    , language-c-quote >=0.12-    , mainland-pretty >=0.6.1-    , markdown >=0.1.16-    , megaparsec >=7.0.1-    , mtl >=2.2.1-    , neat-interpolation >=0.3-    , parallel >=3.2.1.0-    , parser-combinators >=1.0.0-    , process >=1.4.3.0-    , process-extras >=0.7.2-    , random-    , raw-strings-qq >=1.1-    , regex-tdfa >=1.2-    , srcloc >=0.4+    , megaparsec+    , mtl+    , parser-combinators     , tasty     , tasty-hunit     , tasty-quickcheck-    , template-haskell >=2.11.1-    , temporary-    , text >=1.2.2.2-    , th-lift-instances >=0.1.11-    , time >=1.6.0.1-    , transformers >=0.3-    , vector >=0.12-    , vector-binary-instances >=0.2.2.0-    , versions >=3.3.1-    , zip-archive >=0.3.1.1-    , zlib >=0.6.1.2-  if !impl(ghc >= 8.0)-    build-depends:-        semigroups ==0.18.*+    , text   other-modules:       Futhark.Analysis.ScalExpTests       Futhark.Optimise.AlgSimplifyTests@@ -1164,6 +349,9 @@       Futhark.Representation.AST.AttributesTests       Futhark.Representation.AST.Syntax.CoreTests       Futhark.Representation.AST.SyntaxTests+      Futhark.Representation.ExplicitMemory.IndexFunction.Alg+      Futhark.Representation.ExplicitMemory.IndexFunctionTests+      Futhark.Representation.ExplicitMemory.IndexFunctionWrapper       Futhark.Representation.PrimitiveTests       Language.Futhark.CoreTests       Language.Futhark.SyntaxTests

@@ -16,6 +16,13 @@ -- | Create integer from double-precision float. let t64 (x: f64): i32 = i32.f64 x +-- | Semantically just identity, but serves as an optimisation+-- inhibitor.  The compiler will treat this function as a black box.+-- You can use this to work around optimisation deficiencies (or+-- bugs), although it should hopefully rarely be necessary.+let opaque 't (x: t): t =+  intrinsics.opaque x+ -- | Semantically just identity, but in `futharki` the argument value -- will be printed. let trace 't (x: t): t =

@@ -87,10 +87,10 @@ let reduce 'a (op: a -> a -> a) (ne: a) (as: []a): a =   intrinsics.reduce (op, ne, as) --- | As `reduce`, but the operator must also be commutative.  This--- is potentially faster than `reduce`.  For simple built-in--- operators, like addition, the compiler already knows that the--- operator is associative.+-- | As `reduce`, but the operator must also be commutative.  This is+-- potentially faster than `reduce`.  For simple built-in operators,+-- like addition, the compiler already knows that the operator is+-- commutative, so plain `reduce`@term will work just as well. -- -- **Work:** *O(n)* --

@@ -242,15 +242,31 @@   return x[chosen].arch_str; } -static char *cuda_nvrtc_build(struct cuda_context *ctx, const char *src)+static char *cuda_nvrtc_build(struct cuda_context *ctx, const char *src,+                              const char *extra_opts[]) {   nvrtcProgram prog;   NVRTC_SUCCEED(nvrtcCreateProgram(&prog, src, "futhark-cuda", 0, NULL, NULL));+  int arch_set = 0, num_extra_opts; -  size_t n_opts, i = 0, i_dyn, n_opts_alloc = 20 + ctx->cfg.num_sizes;+  // nvrtc cannot handle multiple -arch options.  Hence, if one of the+  // extra_opts is -arch, we have to be careful not to do our usual+  // automatic generation.+  for (num_extra_opts = 0; extra_opts[num_extra_opts] != NULL; num_extra_opts++) {+    if (strstr(extra_opts[num_extra_opts], "-arch")+        == extra_opts[num_extra_opts] ||+        strstr(extra_opts[num_extra_opts], "--gpu-architecture")+        == extra_opts[num_extra_opts]) {+      arch_set = 1;+    }+  }++  size_t n_opts, i = 0, i_dyn, n_opts_alloc = 20 + num_extra_opts + ctx->cfg.num_sizes;   const char **opts = malloc(n_opts_alloc * sizeof(const char *));-  opts[i++] = "-arch";-  opts[i++] = cuda_nvrtc_get_arch(ctx->dev);+  if (!arch_set) {+    opts[i++] = "-arch";+    opts[i++] = cuda_nvrtc_get_arch(ctx->dev);+  }   opts[i++] = "-default-device";   if (ctx->cfg.debugging) {     opts[i++] = "-G";@@ -265,6 +281,13 @@   }   opts[i++] = msgprintf("-DLOCKSTEP_WIDTH=%zu", ctx->lockstep_width);   opts[i++] = msgprintf("-DMAX_THREADS_PER_BLOCK=%zu", ctx->max_block_size);++  // It is crucial that the extra_opts are last, so that the free()+  // logic below does not cause problems.+  for (int j = 0; extra_opts[j] != NULL; j++) {+    opts[i++] = extra_opts[j];+  }+   n_opts = i;    if (ctx->cfg.debugging) {@@ -288,7 +311,7 @@     NVRTC_SUCCEED(res);   } -  for (i = i_dyn; i < n_opts; i++) { free((char *)opts[i]); }+  for (i = i_dyn; i < n_opts-num_extra_opts; i++) { free((char *)opts[i]); }   free(opts);    char *ptx;@@ -356,7 +379,7 @@     if (*size_value == 0) {       *size_value = default_value;     } else if (max_value > 0 && *size_value > max_value) {-      fprintf(stderr, "Note: Device limits %zu to %zu (down from %zu)\n",+      fprintf(stderr, "Note: Device limits %s to %zu (down from %zu)\n",               size_name, max_value, *size_value);       *size_value = max_value;     }@@ -394,16 +417,17 @@ }  static void cuda_module_setup(struct cuda_context *ctx,-    const char *src_fragments[])+                              const char *src_fragments[],+                              const char *extra_opts[]) {   char *ptx = NULL, *src = NULL;    if (ctx->cfg.load_ptx_from == NULL && ctx->cfg.load_program_from == NULL) {     src = concat_fragments(src_fragments);-    ptx = cuda_nvrtc_build(ctx, src);+    ptx = cuda_nvrtc_build(ctx, src, extra_opts);   } else if (ctx->cfg.load_ptx_from == NULL) {     load_string_from_file(ctx->cfg.load_program_from, &src, NULL);-    ptx = cuda_nvrtc_build(ctx, src);+    ptx = cuda_nvrtc_build(ctx, src, extra_opts);   } else {     if (ctx->cfg.load_program_from != NULL) {       fprintf(stderr,@@ -432,7 +456,7 @@   } } -void cuda_setup(struct cuda_context *ctx, const char *src_fragments[])+void cuda_setup(struct cuda_context *ctx, const char *src_fragments[], const char *extra_opts[]) {   CUDA_SUCCEED(cuInit(0)); @@ -450,7 +474,7 @@   ctx->lockstep_width = device_query(ctx->dev, WARP_SIZE);    cuda_size_setup(ctx);-  cuda_module_setup(ctx, src_fragments);+  cuda_module_setup(ctx, src_fragments, extra_opts); }  CUresult cuda_free_all(struct cuda_context *ctx);

@@ -70,8 +70,7 @@   cfg->default_group_size = 0;   cfg->default_num_groups = 0;   cfg->default_tile_size = 0;--  cfg->default_threshold = 32*1024;+  cfg->default_threshold = 0;    cfg->default_group_size_changed = 0;   cfg->default_tile_size_changed = 0;@@ -397,31 +396,29 @@ }  static cl_build_status build_opencl_program(cl_program program, cl_device_id device, const char* options) {-  cl_int ret_val = clBuildProgram(program, 1, &device, options, NULL, NULL);+  cl_int clBuildProgram_error = clBuildProgram(program, 1, &device, options, NULL, NULL);    // Avoid termination due to CL_BUILD_PROGRAM_FAILURE-  if (ret_val != CL_SUCCESS && ret_val != CL_BUILD_PROGRAM_FAILURE) {-    assert(ret_val == 0);+  if (clBuildProgram_error != CL_SUCCESS &&+      clBuildProgram_error != CL_BUILD_PROGRAM_FAILURE) {+    OPENCL_SUCCEED_FATAL(clBuildProgram_error);   }    cl_build_status build_status;-  ret_val = clGetProgramBuildInfo(program,-                                  device,-                                  CL_PROGRAM_BUILD_STATUS,-                                  sizeof(cl_build_status),-                                  &build_status,-                                  NULL);-  assert(ret_val == 0);+  OPENCL_SUCCEED_FATAL(clGetProgramBuildInfo(program,+                                             device,+                                             CL_PROGRAM_BUILD_STATUS,+                                             sizeof(cl_build_status),+                                             &build_status,+                                             NULL));    if (build_status != CL_SUCCESS) {     char *build_log;     size_t ret_val_size;-    ret_val = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);-    assert(ret_val == 0);+    OPENCL_SUCCEED_FATAL(clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size));      build_log = malloc(ret_val_size+1);-    clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);-    assert(ret_val == 0);+    OPENCL_SUCCEED_FATAL(clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL));      // The spec technically does not say whether the build log is zero-terminated, so let's be careful.     build_log[ret_val_size] = '\0';@@ -445,7 +442,8 @@ static cl_program setup_opencl_with_command_queue(struct opencl_context *ctx,                                                   cl_command_queue queue,                                                   const char *srcs[],-                                                  int required_types) {+                                                  int required_types,+                                                  const char *extra_build_opts[]) {   int error;    ctx->queue = queue;@@ -478,7 +476,7 @@     OPENCL_SUCCEED_FATAL(clGetDeviceInfo(device_option.device, CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE,                                    sizeof(cl_uint), &supported, NULL));     if (!supported) {-      panic(1, "Program uses double-precision floats, but this is not supported on the chosen device: %s",+      panic(1, "Program uses double-precision floats, but this is not supported on the chosen device: %s\n",             device_option.device_name);     }   }@@ -578,7 +576,7 @@   }    cl_program prog;-  error = 0;+  error = CL_SUCCESS;   const char* src_ptr[] = {fut_opencl_src};    if (ctx->cfg.dump_program_to != NULL) {@@ -590,12 +588,18 @@    if (ctx->cfg.load_binary_from == NULL) {     prog = clCreateProgramWithSource(ctx->ctx, 1, src_ptr, &src_size, &error);-    assert(error == 0);+    OPENCL_SUCCEED_FATAL(error);      int compile_opts_size = 1024;+     for (int i = 0; i < ctx->cfg.num_sizes; i++) {       compile_opts_size += strlen(ctx->cfg.size_names[i]) + 20;     }++    for (int i = 0; extra_build_opts[i] != NULL; i++) {+      compile_opts_size += strlen(extra_build_opts[i] + 1);+    }+     char *compile_opts = malloc(compile_opts_size);      int w = snprintf(compile_opts, compile_opts_size,@@ -609,6 +613,11 @@                     (int)ctx->cfg.size_values[i]);     } +    for (int i = 0; extra_build_opts[i] != NULL; i++) {+      w += snprintf(compile_opts+w, compile_opts_size-w,+                    "%s ", extra_build_opts[i]);+    }+     OPENCL_SUCCEED_FATAL(build_opencl_program(prog, device_option.device, compile_opts));      free(compile_opts);@@ -650,7 +659,8 @@  static cl_program setup_opencl(struct opencl_context *ctx,                                const char *srcs[],-                               int required_types) {+                               int required_types,+                               const char *extra_build_opts[]) {    ctx->lockstep_width = 0; // Real value set later. @@ -669,15 +679,15 @@     0   }; -  cl_int error;--  ctx->ctx = clCreateContext(properties, 1, &device_option.device, NULL, NULL, &error);-  assert(error == 0);+  cl_int clCreateContext_error;+  ctx->ctx = clCreateContext(properties, 1, &device_option.device, NULL, NULL, &clCreateContext_error);+  OPENCL_SUCCEED_FATAL(clCreateContext_error); -  cl_command_queue queue = clCreateCommandQueue(ctx->ctx, device_option.device, 0, &error);-  assert(error == 0);+  cl_int clCreateCommandQueue_error;+  cl_command_queue queue = clCreateCommandQueue(ctx->ctx, device_option.device, 0, &clCreateCommandQueue_error);+  OPENCL_SUCCEED_FATAL(clCreateCommandQueue_error); -  return setup_opencl_with_command_queue(ctx, queue, srcs, required_types);+  return setup_opencl_with_command_queue(ctx, queue, srcs, required_types, extra_build_opts); }  // Allocate memory from driver. The problem is that OpenCL may perform

@@ -0,0 +1,40 @@+// Start of tuning.h++static char* load_tuning_file(const char *fname,+                              void *cfg,+                              int (*set_size)(void*, const char*, size_t)) {+  const int max_line_len = 1024;+  char* line = malloc(max_line_len);++  FILE *f = fopen(fname, "r");++  if (f == NULL) {+    snprintf(line, max_line_len, "Cannot open file: %s", strerror(errno));+    return line;+  }++  int lineno = 0;+  while (fgets(line, max_line_len, f) != NULL) {+    lineno++;+    char *eql = strstr(line, "=");+    if (eql) {+      *eql = 0;+      int value = atoi(eql+1);+      if (set_size(cfg, line, value) != 0) {+        strncpy(eql+1, line, max_line_len-strlen(line)-1);+        snprintf(line, max_line_len, "Unknown name '%s' on line %d.", eql+1, lineno);+        return line;+      }+    } else {+      snprintf(line, max_line_len, "Invalid line %d (must be of form 'name=int').",+               lineno);+      return line;+    }+  }++  free(line);++  return NULL;+}++// End of tuning.h

@@ -374,10 +374,13 @@     cfg.DumpProgramTo = null;     cfg.LoadProgramFrom = null; -    cfg.DefaultGroupSize = 256;-    cfg.DefaultNumGroups = 128;-    cfg.DefaultTileSize = 32;-    cfg.DefaultThreshold = 32*1024;+    // The following are dummy sizes that mean the concrete defaults+    // will be set during initialisation via hardware-inspection-based+    // heuristics.+    cfg.DefaultGroupSize = 0;+    cfg.DefaultNumGroups = 0;+    cfg.DefaultTileSize = 0;+    cfg.DefaultThreshold = 0;      cfg.NumSizes = num_sizes;     cfg.SizeNames = size_names;@@ -774,6 +777,9 @@      int MaxTileSize = (int) Math.Sqrt(MaxGroupSize); +    // Make sure this function is defined.+    PostOpenCLSetup(ref ctx, ref device_option);+     if (MaxGroupSize < ctx.OpenCL.Cfg.DefaultGroupSize) {         Console.Error.WriteLine("Note: Device limits default group size to {0} (down from {1}).\n",                                 MaxGroupSize, ctx.OpenCL.Cfg.DefaultGroupSize);@@ -798,16 +804,16 @@         string size_name = ctx.OpenCL.Cfg.SizeNames[i];         int max_value, default_value;         max_value = default_value = 0;-        if (size_class == "group_size") {+        if (size_class.StartsWith("group_size")) {             max_value = MaxGroupSize;             default_value = ctx.OpenCL.Cfg.DefaultGroupSize;-        } else if (size_class == "num_groups") {+        } else if (size_class.StartsWith("num_groups")) {             max_value = MaxGroupSize; // Futhark assumes this constraint.             default_value = ctx.OpenCL.Cfg.DefaultNumGroups;-        } else if (size_class == "tile_size"){+        } else if (size_class.StartsWith("tile_size")){             max_value = (int) Math.Sqrt(MaxGroupSize);             default_value = ctx.OpenCL.Cfg.DefaultTileSize;-        } else if (size_class == "threshold") {+        } else if (size_class.StartsWith("threshold")) {             max_value = 0; // No limit.             default_value = ctx.OpenCL.Cfg.DefaultThreshold;         } else {@@ -835,9 +841,6 @@     ctx.OpenCL.Queue = CL10.CreateCommandQueue(ctx.OpenCL.Context, device, 0, out error);     Debug.Assert(error == 0); -    // Make sure this function is defined.-    PostOpenCLSetup(ref ctx, ref device_option);-     if (ctx.Debugging) {         Console.Error.WriteLine("Lockstep width: {0}\n", (int)ctx.OpenCL.LockstepWidth);         Console.Error.WriteLine("Default group size: {0}\n", (int)ctx.OpenCL.Cfg.DefaultGroupSize);@@ -916,8 +919,31 @@      var name = name_and_value[0];     var value = Convert.ToInt32(name_and_value[1]);-    if (!FutharkContextConfigSetSize(ref config, name, value))+    if (name == "default_num_groups") {+        config.OpenCL.DefaultNumGroups = value;+    }+    else if (name == "default_group_size") {+        config.OpenCL.DefaultGroupSize = value;+    }+    else if (name == "default_tile_size") {+        config.OpenCL.DefaultTileSize = value;+    }+    else if (name == "default_threshold") {+        config.OpenCL.DefaultThreshold = value;+    }+    else if (!FutharkContextConfigSetSize(ref config, name, value))     {         panic(1, "Unknown size: {0}", name);     }+}++private void FutharkConfigLoadTuning(ref FutharkContextConfig config, string fname)+{+    StreamReader file = new StreamReader(fname);+    String line;+    while((line = file.ReadLine()) != null)+    {+        FutharkConfigSetSize(ref config, line);+    }+    file.Close(); }

@@ -55,6 +55,10 @@                 device_matches += 1     raise Exception('No OpenCL platform and device matching constraints found.') +def size_assignment(s):+    name, value = s.split('=')+    return (name, int(value))+ def check_types(self, required_types):     if 'f64' in required_types:         if self.device.get_info(cl.device_info.PREFERRED_VECTOR_WIDTH_DOUBLE) == 0:@@ -107,6 +111,22 @@     self.max_num_groups = 0     self.free_list = {} +    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']+        del sizes['default_num_groups']++    if 'default_tile_size' in sizes:+        default_tile_size = sizes['default_tile_size']+        del sizes['default_tile_size']++    if 'default_threshold' in sizes:+        default_threshold = sizes['default_threshold']+        del sizes['default_threshold']+     default_group_size_set = default_group_size != None     default_tile_size_set = default_tile_size != None     default_sizes = apply_size_heuristics(self, size_heuristics,@@ -137,7 +157,7 @@         if k in all_sizes:             all_sizes[k]['value'] = v         else:-            raise Exception('Unknown size: {}'.format(k))+            raise Exception('Unknown size: {}\nKnown sizes: {}'.format(k, ' '.join(all_sizes.keys())))      self.sizes = {}     for (k,v) in all_sizes.items():

@@ -0,0 +1,11 @@+### start of tuning.py+###+### Reading the .tuning file.++def read_tuning_file(kvs, f):+    for line in f.read().splitlines():+        size, value = line.split('=')+        kvs[size] = int(value)+    return kvs++### end of tuning.py

@@ -45,11 +45,12 @@                    , optExcludeCase :: [String]                    , optIgnoreFiles :: [Regex]                    , optEntryPoint :: Maybe String+                   , optTuning :: Maybe String                    }  initialBenchOptions :: BenchOptions initialBenchOptions = BenchOptions "c" "futhark" "" 10 [] Nothing (-1) False-                      ["nobench", "disable"] [] Nothing+                      ["nobench", "disable"] [] Nothing (Just "tuning")  -- | The name we use for compiled programs. binaryName :: FilePath -> FilePath@@ -158,9 +159,13 @@ runBenchmark :: BenchOptions -> (FilePath, [InputOutputs]) -> IO [BenchResult] runBenchmark opts (program, cases) = mapM forInputOutputs $ filter relevant cases   where forInputOutputs (InputOutputs entry_name runs) = do-          putStr $ "Results for " ++ program' ++ ":\n"+          (tuning_opts, tuning_desc) <- determineTuning (optTuning opts) program++          putStr $ "Results for " ++ program' ++ tuning_desc ++ ":\n"+          let opts' = opts { optExtraOptions =+                               optExtraOptions opts ++ tuning_opts }           BenchResult program' . catMaybes <$>-            mapM (runBenchmarkCase opts program entry_name pad_to) runs+            mapM (runBenchmarkCase opts' program entry_name pad_to) runs           where program' = if entry_name == "main"                            then program                            else program ++ ":" ++ T.unpack entry_name@@ -368,6 +373,13 @@                 config { optEntryPoint = Just s })       "NAME")     "Only run this entry point."+  , Option [] ["tuning"]+    (ReqArg (\s -> Right $ \config -> config { optTuning = Just s })+    "EXTENSION")+    "Look for tuning files with this extension (defaults: .tuning)."+  , Option [] ["no-tuning"]+    (NoArg $ Right $ \config -> config { optTuning = Nothing })+    "Do not load tuning files."   ]   where max_timeout :: Int         max_timeout = maxBound `div` 1000000

@@ -3,7 +3,6 @@  import qualified Data.ByteString.Lazy.Char8 as BS import System.Exit-import System.IO  import Futhark.Test.Values import Futhark.Util.Options@@ -21,9 +20,8 @@             (Just vs_a, Just vs_b) ->               case compareValues vs_a vs_b of                 [] -> return ()-                es -> do-                  mapM_ (hPrint stderr) es-                  exitWith $ ExitFailure 2+                es -> do mapM_ print es+                         exitWith $ ExitFailure 2          f _ _ =           Nothing

@@ -66,6 +66,9 @@                      -- ^ Nothing is distinct from a empty pipeline -                      -- it means we don't even run the internaliser.                      , futharkAction :: UntypedAction+                     , futharkPrintAST :: Bool+                     -- ^ If true, prints programs as raw ASTs instead+                     -- of their prettyprinted form.                      }  @@ -121,7 +124,8 @@   representation PolyAction{} = "<any>"  newConfig :: Config-newConfig = Config newFutharkConfig (Pipeline []) $ PolyAction printAction printAction printAction+newConfig = Config newFutharkConfig (Pipeline [])+            (PolyAction printAction printAction printAction) False  changeFutharkConfig :: (FutharkConfig -> FutharkConfig)                     -> Config -> Config@@ -251,7 +255,7 @@   , Option "t" ["type-check"]     (NoArg $ Right $ \opts ->         opts { futharkPipeline = TypeCheck })-    "Type-check the program and print errors on standard error."+    "Print on standard output the type-checked program."    , Option [] ["pretty-print"]     (NoArg $ Right $ \opts ->@@ -284,6 +288,9 @@   , Option [] ["defunctionalise"]     (NoArg $ Right $ \opts -> opts { futharkPipeline = Defunctionalise })     "Defunctionalise the program."+  , Option [] ["ast"]+    (NoArg $ Right $ \opts -> opts { futharkPrintAST = True })+    "Output ASTs instead of prettyprinted programs."   , typedPassOption soacsProg Kernels firstOrderTransform "f"   , soacsPassOption fuseSOACs "o"   , soacsPassOption inlineAndRemoveDeadFunctions []@@ -336,27 +343,32 @@           Nothing         m file config =           case futharkPipeline config of-            TypeCheck -> do-              -- No pipeline; just read the program and type check-              (warnings, _, _) <- readProgram file-              liftIO $ hPutStr stderr $ show warnings             PrettyPrint -> liftIO $ do               maybe_prog <- parseFuthark file <$> T.readFile file               case maybe_prog of                 Left err  -> fail $ show err-                Right prog -> putStrLn $ pretty prog+                Right prog | futharkPrintAST config -> print prog+                           | otherwise -> putStrLn $ pretty prog+            TypeCheck -> do+              (_, imports, _) <- readProgram file+              liftIO $ forM_ (map snd imports) $ \fm ->+                putStrLn $ if futharkPrintAST config+                           then show $ fileProg fm+                           else pretty $ fileProg fm             Defunctorise -> do               (_, imports, src) <- readProgram file-              liftIO $ mapM_ (putStrLn . pretty) $+              liftIO $ mapM_ (putStrLn . if futharkPrintAST config then show else pretty) $                 evalState (Defunctorise.transformProg imports) src             Monomorphise -> do               (_, imports, src) <- readProgram file-              liftIO $ mapM_ (putStrLn . pretty) $ flip evalState src $+              liftIO $ mapM_ (putStrLn . if futharkPrintAST config then show else pretty) $+                flip evalState src $                 Defunctorise.transformProg imports                 >>= Monomorphise.transformProg             Defunctionalise -> do               (_, imports, src) <- readProgram file-              liftIO $ mapM_ (putStrLn . pretty) $ flip evalState src $+              liftIO $ mapM_ (putStrLn . if futharkPrintAST config then show else pretty) $+                flip evalState src $                 Defunctorise.transformProg imports                 >>= Monomorphise.transformProg                 >>= Defunctionalise.transformProg

@@ -2,9 +2,11 @@ -- Various small subcommands that are too simple to deserve their own file. module Futhark.CLI.Misc   ( mainCheck+  , mainImports   ) where +import Data.List (isPrefixOf) import Control.Monad.State import System.IO import System.Exit@@ -29,3 +31,14 @@     _ -> Nothing   where check file = do (warnings, _, _) <- readProgram file                         liftIO $ hPutStr stderr $ show warnings++mainImports :: String -> [String] -> IO ()+mainImports = mainWithOptions () [] "program" $ \args () ->+  case args of+    [file] -> Just $ runFutharkM' $ findImports file+    _ -> Nothing+  where findImports file = do+          (_, prog_imports, _) <- readProgram file+          liftIO $ putStr $ unlines $ map (++ ".fut")+            $ filter (\f -> not ("futlib/" `isPrefixOf` f))+            $ map fst prog_imports

@@ -283,7 +283,7 @@         Nothing ->           liftIO $ T.putStrLn $ "Added new required package " <> p <> " " <> prettySemVer v <> "."       putPkgManifest m'-      liftIO $ T.putStrLn "Remember to run 'futhark-pkg sync'."+      liftIO $ T.putStrLn "Remember to run 'futhark pkg sync'."  doRemove :: String -> [String] -> IO () doRemove = cmdMain "PKGPATH" $ \args cfg ->@@ -326,6 +326,9 @@       m <- getPkgManifest       rs <- traverse (mapM (traverse upgrade)) $ manifestRequire m       putPkgManifest m { manifestRequire = rs }+      if rs == manifestRequire m+        then liftIO $ T.putStrLn "Nothing to upgrade."+        else liftIO $ T.putStrLn "Remember to run 'futhark pkg sync'."     _ -> Nothing   where upgrade req = do           v <- lookupNewestRev $ requiredPkg req

@@ -91,7 +91,7 @@  confirmQuit :: Haskeline.InputT IO Bool confirmQuit = do-  c <- Haskeline.getInputChar "Quit futharki? (y/n) "+  c <- Haskeline.getInputChar "Quit REPL? (y/n) "   case c of     Nothing -> return True -- EOF     Just 'y' -> return True

@@ -169,9 +169,13 @@         context ("Compiling with --backend=" <> T.pack backend) $ do           compileTestProgram extra_options futhark backend program warnings           mapM_ (testMetrics progs program) structures-          unless (mode == Compile) $+          unless (mode == Compile) $ do+            (tuning_opts, _) <-+              liftIO $ determineTuning (configTuning progs) program+            let progs' = progs { configExtraOptions =+                                 tuning_opts ++ configExtraOptions progs }             context "Running compiled program" $-            accErrors_ $ map (runCompiledEntry program progs) ios+              accErrors_ $ map (runCompiledEntry program progs') ios       unless (mode == Compile || mode == Compiled) $         context "Interpreting" $           accErrors_ $ map (runInterpretedEntry futhark program) ios@@ -380,8 +384,10 @@   _ <- forkIO $ mapM_ (putMVar testmvar) included   isTTY <- (&& not (configLineOutput config)) <$> hIsTerminalDevice stdout -  let report = if isTTY then reportTable else reportText-      clear  = if isTTY then clearFromCursorToScreenEnd else putStr "\n"+  let report | isTTY = reportTable+             | otherwise = reportText+      clear | isTTY = clearFromCursorToScreenEnd+            |otherwise = putStr "\n"        numTestCases tc =         case testAction $ testCaseTest tc of@@ -442,9 +448,8 @@   -- Removes "Now testing" output.   when isTTY $ cursorUpLine 1 >> clearLine -  let excluded_str = if null excluded-                     then ""-                     else " (" ++ show (length excluded) ++ " program(s) excluded).\n"+  let excluded_str | null excluded = ""+                   | otherwise = " (" ++ show (length excluded) ++ " program(s) excluded).\n"   putStr excluded_str   exitWith $ case testStatusFail ts of 0 -> ExitSuccess                                        _ -> ExitFailure 1@@ -473,6 +478,7 @@                              , configRunner = ""                              , configExtraOptions = []                              , configExtraCompilerOptions = []+                             , configTuning = Just "tuning"                              }                            , configLineOutput = False                            }@@ -482,6 +488,7 @@                   , configFuthark :: FilePath                   , configRunner :: FilePath                   , configExtraCompilerOptions :: [String]+                  , configTuning :: Maybe String                   , configExtraOptions :: [String]                   -- ^ Extra options passed to the programs being run.                   }

@@ -51,24 +51,31 @@ cliOptions :: [Option] cliOptions = [ Option { optionLongName = "dump-cuda"                       , optionShortName = Nothing-                      , optionArgument = RequiredArgument-                      , optionAction = [C.cstm|futhark_context_config_dump_program_to(cfg, optarg);|]+                      , optionArgument = RequiredArgument "FILE"+                      , optionAction = [C.cstm|{futhark_context_config_dump_program_to(cfg, optarg);+                                                entry_point = NULL;}|]                       }              , Option { optionLongName = "load-cuda"                       , optionShortName = Nothing-                      , optionArgument = RequiredArgument+                      , optionArgument = RequiredArgument "FILE"                       , optionAction = [C.cstm|futhark_context_config_load_program_from(cfg, optarg);|]                       }              , Option { optionLongName = "dump-ptx"                       , optionShortName = Nothing-                      , optionArgument = RequiredArgument-                      , optionAction = [C.cstm|futhark_context_config_dump_ptx_to(cfg, optarg);|]+                      , optionArgument = RequiredArgument "FILE"+                      , optionAction = [C.cstm|{futhark_context_config_dump_ptx_to(cfg, optarg);+                                                entry_point = NULL;}|]                       }              , Option { optionLongName = "load-ptx"                       , optionShortName = Nothing-                      , optionArgument = RequiredArgument+                      , optionArgument = RequiredArgument "FILE"                       , optionAction = [C.cstm|futhark_context_config_load_ptx_from(cfg, optarg);|]                       }+             , Option { optionLongName = "nvrtc-option"+                      , optionShortName = Nothing+                      , optionArgument = RequiredArgument "OPT"+                      , optionAction = [C.cstm|futhark_context_config_add_nvrtc_option(cfg, optarg);|]+                      }              , Option { optionLongName = "print-sizes"                       , optionShortName = Nothing                       , optionArgument = NoArgument@@ -81,6 +88,17 @@                           exit(0);                         }|]                       }+             , Option { optionLongName = "tuning"+                 , optionShortName = Nothing+                 , optionArgument = RequiredArgument "FILE"+                 , optionAction = [C.cstm|{+                     char *fname = optarg;+                     char *ret = load_tuning_file(optarg, cfg, (int(*)(void*, const char*, size_t))+                                                               futhark_context_config_set_size);+                     if (ret != NULL) {+                       panic(1, "When loading tuning from '%s': %s\n", optarg, ret);+                     }}|]+                 }              ]  writeCUDAScalar :: GC.WriteScalar OpenCL ()@@ -138,12 +156,17 @@                            ++ "' from '" ++ show srcSpace ++ "'."  staticCUDAArray :: GC.StaticArray OpenCL ()-staticCUDAArray name "device" t vals = do+staticCUDAArray name "device" t vs = do   let ct = GC.primTypeToCType t-      vals' = [[C.cinit|$exp:v|] | v <- map GC.compilePrimValue vals]-      num_elems = length vals   name_realtype <- newVName $ baseString name ++ "_realtype"-  GC.libDecl [C.cedecl|static $ty:ct $id:name_realtype[$int:num_elems] = {$inits:vals'};|]+  num_elems <- case vs of+    ArrayValues vs' -> do+      let vs'' = [[C.cinit|$exp:v|] | v <- map GC.compilePrimValue vs']+      GC.libDecl [C.cedecl|static $ty:ct $id:name_realtype[$int:(length vs'')] = {$inits:vs''};|]+      return $ length vs''+    ArrayZeros n -> do+      GC.libDecl [C.cedecl|static $ty:ct $id:name_realtype[$int:n];|]+      return n   -- Fake a memory block.   GC.contextField (pretty name) [C.cty|struct memblock_device|] Nothing   -- During startup, copy the data to where we need it.@@ -274,4 +297,4 @@     printSizes =       intercalate [[C.cstm|fprintf(stderr, ", ");|]] . map printSize     printSize e =-      [[C.cstm|fprintf(stderr, "%zu", $exp:e);|]]+      [[C.cstm|fprintf(stderr, "%d", $exp:e);|]]

@@ -77,6 +77,8 @@     ([C.cedecl|struct $id:s;|],      [C.cedecl|struct $id:s { struct cuda_config cu_cfg;                               size_t sizes[$int:num_sizes];+                              int num_nvrtc_opts;+                              const char **nvrtc_opts;                             };|])    let size_value_inits = map (\i -> [C.cstm|cfg->sizes[$int:i] = 0;|])@@ -89,6 +91,9 @@                            return NULL;                          } +                         cfg->num_nvrtc_opts = 0;+                         cfg->nvrtc_opts = malloc(sizeof(const char*));+                         cfg->nvrtc_opts[0] = NULL;                          $stms:size_value_inits                          cuda_config_init(&cfg->cu_cfg, $int:num_sizes,                                           size_names, size_vars,@@ -99,9 +104,19 @@   GC.publicDef_ "context_config_free" GC.InitDecl $ \s ->     ([C.cedecl|void $id:s(struct $id:cfg* cfg);|],      [C.cedecl|void $id:s(struct $id:cfg* cfg) {+                         free(cfg->nvrtc_opts);                          free(cfg);                        }|]) +  GC.publicDef_ "context_config_add_nvrtc_option" GC.InitDecl $ \s ->+    ([C.cedecl|void $id:s(struct $id:cfg* cfg, const char *opt);|],+     [C.cedecl|void $id:s(struct $id:cfg* cfg, const char *opt) {+                         cfg->nvrtc_opts[cfg->num_nvrtc_opts] = opt;+                         cfg->num_nvrtc_opts++;+                         cfg->nvrtc_opts = realloc(cfg->nvrtc_opts, (cfg->num_nvrtc_opts+1) * sizeof(const char*));+                         cfg->nvrtc_opts[cfg->num_nvrtc_opts] = NULL;+                       }|])+   GC.publicDef_ "context_config_set_debugging" GC.InitDecl $ \s ->     ([C.cedecl|void $id:s(struct $id:cfg* cfg, int flag);|],      [C.cedecl|void $id:s(struct $id:cfg* cfg, int flag) {@@ -222,7 +237,7 @@                           create_lock(&ctx->lock);                           $stms:init_fields -                          cuda_setup(&ctx->cuda, cuda_program);+                          cuda_setup(&ctx->cuda, cuda_program, cfg->nvrtc_opts);                           $stms:(map (loadKernelByName) kernel_names)                            $stms:final_inits

@@ -45,6 +45,7 @@                      }         include_opencl_h = unlines ["#define CL_USE_DEPRECATED_OPENCL_1_2_APIS",                                     "#ifdef __APPLE__",+                                    "#define CL_SILENCE_DEPRECATION",                                     "#include <OpenCL/cl.h>",                                     "#else",                                     "#include <CL/cl.h>",@@ -53,54 +54,61 @@ cliOptions :: [Option] cliOptions = [ Option { optionLongName = "platform"                       , optionShortName = Just 'p'-                      , optionArgument = RequiredArgument+                      , optionArgument = RequiredArgument "NAME"                       , optionAction = [C.cstm|futhark_context_config_set_platform(cfg, optarg);|]                       }              , Option { optionLongName = "device"                       , optionShortName = Just 'd'-                      , optionArgument = RequiredArgument+                      , optionArgument = RequiredArgument "NAME"                       , optionAction = [C.cstm|futhark_context_config_set_device(cfg, optarg);|]                       }              , Option { optionLongName = "default-group-size"                       , optionShortName = Nothing-                      , optionArgument = RequiredArgument+                      , optionArgument = RequiredArgument "INT"                       , optionAction = [C.cstm|futhark_context_config_set_default_group_size(cfg, atoi(optarg));|]                       }              , Option { optionLongName = "default-num-groups"                       , optionShortName = Nothing-                      , optionArgument = RequiredArgument+                      , optionArgument = RequiredArgument "INT"                       , optionAction = [C.cstm|futhark_context_config_set_default_num_groups(cfg, atoi(optarg));|]                       }              , Option { optionLongName = "default-tile-size"                       , optionShortName = Nothing-                      , optionArgument = RequiredArgument+                      , optionArgument = RequiredArgument "INT"                       , optionAction = [C.cstm|futhark_context_config_set_default_tile_size(cfg, atoi(optarg));|]                       }              , Option { optionLongName = "default-threshold"                       , optionShortName = Nothing-                      , optionArgument = RequiredArgument+                      , optionArgument = RequiredArgument "INT"                       , optionAction = [C.cstm|futhark_context_config_set_default_threshold(cfg, atoi(optarg));|]                       }              , Option { optionLongName = "dump-opencl"                       , optionShortName = Nothing-                      , optionArgument = RequiredArgument-                      , optionAction = [C.cstm|futhark_context_config_dump_program_to(cfg, optarg);|]+                      , optionArgument = RequiredArgument "FILE"+                      , optionAction = [C.cstm|{futhark_context_config_dump_program_to(cfg, optarg);+                                                entry_point = NULL;}|]                       }              , Option { optionLongName = "load-opencl"                       , optionShortName = Nothing-                      , optionArgument = RequiredArgument+                      , optionArgument = RequiredArgument "FILE"                       , optionAction = [C.cstm|futhark_context_config_load_program_from(cfg, optarg);|]                       }              , Option { optionLongName = "dump-opencl-binary"                       , optionShortName = Nothing-                      , optionArgument = RequiredArgument-                      , optionAction = [C.cstm|futhark_context_config_dump_binary_to(cfg, optarg);|]+                      , optionArgument = RequiredArgument "FILE"+                      , optionAction = [C.cstm|{futhark_context_config_dump_binary_to(cfg, optarg);+                                                entry_point = NULL;}|]                       }              , Option { optionLongName = "load-opencl-binary"                       , optionShortName = Nothing-                      , optionArgument = RequiredArgument+                      , optionArgument = RequiredArgument "FILE"                       , optionAction = [C.cstm|futhark_context_config_load_binary_from(cfg, optarg);|]                       }+             , Option { optionLongName = "build-option"+                      , optionShortName = Nothing+                      , optionArgument = RequiredArgument "OPT"+                      , optionAction = [C.cstm|futhark_context_config_add_build_option(cfg, optarg);|]+                      }              , Option { optionLongName = "print-sizes"                       , optionShortName = Nothing                       , optionArgument = NoArgument@@ -115,7 +123,7 @@                       }              , Option { optionLongName = "size"                       , optionShortName = Nothing-                      , optionArgument = RequiredArgument+                      , optionArgument = RequiredArgument "NAME=INT"                       , optionAction = [C.cstm|{                           char *name = optarg;                           char *equals = strstr(optarg, "=");@@ -130,6 +138,17 @@                             panic(1, "Invalid argument for size option: %s\n", optarg);                           }}|]                       }+             , Option { optionLongName = "tuning"+                      , optionShortName = Nothing+                      , optionArgument = RequiredArgument "FILE"+                      , optionAction = [C.cstm|{+                          char *fname = optarg;+                          char *ret = load_tuning_file(optarg, cfg, (int(*)(void*, const char*, size_t))+                                                                    futhark_context_config_set_size);+                          if (ret != NULL) {+                            panic(1, "When loading tuning from '%s': %s\n", optarg, ret);+                          }}|]+                      }              ]  -- We detect the special case of writing a constant and turn it into a@@ -240,10 +259,15 @@ staticOpenCLArray :: GC.StaticArray OpenCL () staticOpenCLArray name "device" t vs = do   let ct = GC.primTypeToCType t-      vs' = [[C.cinit|$exp:v|] | v <- map GC.compilePrimValue vs]-      num_elems = length vs   name_realtype <- newVName $ baseString name ++ "_realtype"-  GC.libDecl [C.cedecl|static $ty:ct $id:name_realtype[$int:num_elems] = {$inits:vs'};|]+  num_elems <- case vs of+    ArrayValues vs' -> do+      let vs'' = [[C.cinit|$exp:v|] | v <- map GC.compilePrimValue vs']+      GC.libDecl [C.cedecl|static $ty:ct $id:name_realtype[$int:(length vs'')] = {$inits:vs''};|]+      return $ length vs''+    ArrayZeros n -> do+      GC.libDecl [C.cedecl|static $ty:ct $id:name_realtype[$int:n];|]+      return n   -- Fake a memory block.   GC.contextField (pretty name) [C.cty|struct memblock_device|] Nothing   -- During startup, copy the data to where we need it.

@@ -61,6 +61,8 @@     ([C.cedecl|struct $id:s;|],      [C.cedecl|struct $id:s { struct opencl_config opencl;                               size_t sizes[$int:num_sizes];+                              int num_build_opts;+                              const char **build_opts;                             };|])    let size_value_inits = map (\i -> [C.cstm|cfg->sizes[$int:i] = 0;|]) [0..M.size sizes-1]@@ -72,6 +74,9 @@                            return NULL;                          } +                         cfg->num_build_opts = 0;+                         cfg->build_opts = malloc(sizeof(const char*));+                         cfg->build_opts[0] = NULL;                          $stms:size_value_inits                          opencl_config_init(&cfg->opencl, $int:num_sizes,                                             size_names, size_vars,@@ -82,9 +87,19 @@   GC.publicDef_ "context_config_free" GC.InitDecl $ \s ->     ([C.cedecl|void $id:s(struct $id:cfg* cfg);|],      [C.cedecl|void $id:s(struct $id:cfg* cfg) {+                         free(cfg->build_opts);                          free(cfg);                        }|]) +  GC.publicDef_ "context_config_add_build_option" GC.InitDecl $ \s ->+    ([C.cedecl|void $id:s(struct $id:cfg* cfg, const char *opt);|],+     [C.cedecl|void $id:s(struct $id:cfg* cfg, const char *opt) {+                         cfg->build_opts[cfg->num_build_opts] = opt;+                         cfg->num_build_opts++;+                         cfg->build_opts = realloc(cfg->build_opts, (cfg->num_build_opts+1) * sizeof(const char*));+                         cfg->build_opts[cfg->num_build_opts] = NULL;+                       }|])+   GC.publicDef_ "context_config_set_debugging" GC.InitDecl $ \s ->     ([C.cedecl|void $id:s(struct $id:cfg* cfg, int flag);|],      [C.cedecl|void $id:s(struct $id:cfg* cfg, int flag) {@@ -170,6 +185,27 @@                              return 0;                            }                          }++                         if (strcmp(size_name, "default_group_size") == 0) {+                           cfg->opencl.default_group_size = size_value;+                           return 0;+                         }++                         if (strcmp(size_name, "default_num_groups") == 0) {+                           cfg->opencl.default_num_groups = size_value;+                           return 0;+                         }++                         if (strcmp(size_name, "default_threshold") == 0) {+                           cfg->opencl.default_threshold = size_value;+                           return 0;+                         }++                         if (strcmp(size_name, "default_tile_size") == 0) {+                           cfg->opencl.default_tile_size = size_value;+                           return 0;+                         }+                          return 1;                        }|]) @@ -231,7 +267,7 @@                           $stms:set_required_types                            init_context_early(cfg, ctx);-                          typename cl_program prog = setup_opencl(&ctx->opencl, opencl_program, required_types);+                          typename cl_program prog = setup_opencl(&ctx->opencl, opencl_program, required_types, cfg->build_opts);                           init_context_late(cfg, ctx, prog);                           return ctx;                        }|])@@ -248,7 +284,7 @@                           $stms:set_required_types                            init_context_early(cfg, ctx);-                          typename cl_program prog = setup_opencl_with_command_queue(&ctx->opencl, queue, opencl_program, required_types);+                          typename cl_program prog = setup_opencl_with_command_queue(&ctx->opencl, queue, opencl_program, required_types, cfg->build_opts);                           init_context_late(cfg, ctx, prog);                           return ctx;                        }|])@@ -403,6 +439,7 @@                    NumGroups -> [C.cexp|ctx->cfg.default_num_groups|]                    GroupSize -> [C.cexp|ctx->cfg.default_group_size|]                    TileSize -> [C.cexp|ctx->cfg.default_tile_size|]+                   Threshold -> [C.cexp|ctx->cfg.default_threshold|]          get_size = case what of                      HeuristicConst x ->

@@ -337,11 +337,17 @@   -- Create host-side C# array with intended values.   tmp_arr <- newVName' "tmpArr"   let t' = CS.compilePrimTypeToAST t-  CS.staticMemDecl $ AssignTyped (Composite $ ArrayT t') (Var tmp_arr) (Just $ CreateArray t' $ map CS.compilePrimValue vs)+  CS.staticMemDecl $ AssignTyped (Composite $ ArrayT t') (Var tmp_arr) $ Just $+    case vs of Imp.ArrayValues vs' ->+                 CreateArray (CS.compilePrimTypeToAST t) $ Right $ map CS.compilePrimValue vs'+               Imp.ArrayZeros n ->+                 CreateArray (CS.compilePrimTypeToAST t) $ Left n    -- Create memory block on the device.   ptr <- newVName' "ptr"-  let size = Integer $ genericLength vs * Imp.primByteSize t+  let num_elems = case vs of Imp.ArrayValues vs' -> length vs'+                             Imp.ArrayZeros n -> n+      size = Integer $ toInteger num_elems * Imp.primByteSize t    CS.staticMemAlloc $ Reassign (Var name') (CS.simpleCall "EmptyMemblock" [Var "Ctx.EMPTY_MEM_HANDLE"])   errcode <- CS.compileName <$> newVName "errCode"@@ -375,7 +381,7 @@   let createTuple = "createTuple_"++ pretty bt'    return $ CS.simpleCall createTuple [ memblockFromMem mem, Var "Ctx.OpenCL.Queue", nbytes-                                     , CreateArray (Primitive $ CSInt Int64T) dims']+                                     , CreateArray (Primitive $ CSInt Int64T) $ Right dims']   where dims' = map CS.compileDim dims  packArrayOutput _ sid _ _ _ =

@@ -271,6 +271,7 @@                    NumGroups ->     Var "Ctx.OpenCL.Cfg.DefaultNumGroups"                    GroupSize ->     Var "Ctx.OpenCL.Cfg.DefaultGroupSize"                    TileSize ->      Var "Ctx.OpenCL.Cfg.DefaultTileSize"+                   Threshold ->     Var "Ctx.OpenCL.Cfg.DefaultThreshold"          get_size = case what of                      HeuristicConst x ->

@@ -161,7 +161,7 @@ type Deallocate op s = C.Exp -> C.Exp -> SpaceId -> CompilerM op s ()  -- | Create a static array of values - initialised at load time.-type StaticArray op s = VName -> SpaceId -> PrimType -> [PrimValue] -> CompilerM op s ()+type StaticArray op s = VName -> SpaceId -> PrimType -> ArrayContents -> CompilerM op s ()  -- | Copy from one memory block to another. type Copy op s = C.Exp -> C.Exp -> Space ->@@ -1221,12 +1221,12 @@ benchmarkOptions =    [ Option { optionLongName = "write-runtime-to"             , optionShortName = Just 't'-            , optionArgument = RequiredArgument+            , optionArgument = RequiredArgument "FILE"             , optionAction = set_runtime_file             }    , Option { optionLongName = "runs"             , optionShortName = Just 'r'-            , optionArgument = RequiredArgument+            , optionArgument = RequiredArgument "INT"             , optionAction = set_num_runs             }    , Option { optionLongName = "debugging"@@ -1241,8 +1241,8 @@             }    , Option { optionLongName = "entry-point"             , optionShortName = Just 'e'-            , optionArgument = RequiredArgument-            , optionAction = [C.cstm|entry_point = optarg;|]+            , optionArgument = RequiredArgument "NAME"+            , optionAction = [C.cstm|if (entry_point != NULL) entry_point = optarg;|]             }    , Option { optionLongName = "binary-output"             , optionShortName = Just 'b'@@ -1356,8 +1356,11 @@ static typename FILE *runtime_file; static int perform_warmup = 0; static int num_runs = 1;+// If the entry point is NULL, the program will terminate after doing initialisation and such. static const char *entry_point = "main"; +$esc:tuning_h+ $func:option_parser  $edecls:cli_entry_point_decls@@ -1390,31 +1393,33 @@   struct futhark_context *ctx = futhark_context_new(cfg);   assert (ctx != NULL); -  int num_entry_points = sizeof(entry_points) / sizeof(entry_points[0]);-  entry_point_fun *entry_point_fun = NULL;-  for (int i = 0; i < num_entry_points; i++) {-    if (strcmp(entry_points[i].name, entry_point) == 0) {-      entry_point_fun = entry_points[i].fun;-      break;+  if (entry_point != NULL) {+    int num_entry_points = sizeof(entry_points) / sizeof(entry_points[0]);+    entry_point_fun *entry_point_fun = NULL;+    for (int i = 0; i < num_entry_points; i++) {+      if (strcmp(entry_points[i].name, entry_point) == 0) {+        entry_point_fun = entry_points[i].fun;+        break;+      }     }-  } -  if (entry_point_fun == NULL) {-    fprintf(stderr, "No entry point '%s'.  Select another with --entry-point.  Options are:\n",-                    entry_point);-    for (int i = 0; i < num_entry_points; i++) {-      fprintf(stderr, "%s\n", entry_points[i].name);+    if (entry_point_fun == NULL) {+      fprintf(stderr, "No entry point '%s'.  Select another with --entry-point.  Options are:\n",+                      entry_point);+      for (int i = 0; i < num_entry_points; i++) {+        fprintf(stderr, "%s\n", entry_points[i].name);+      }+      return 1;     }-    return 1;-  } -  entry_point_fun(ctx);+    entry_point_fun(ctx); -  if (runtime_file != NULL) {-    fclose(runtime_file);-  }+    if (runtime_file != NULL) {+      fclose(runtime_file);+    } -  futhark_debugging_report(ctx);+    futhark_debugging_report(ctx);+  }    futhark_context_free(ctx);   futhark_context_config_free(cfg);@@ -1489,6 +1494,7 @@         values_h = $(embedStringFile "rts/c/values.h")         timing_h = $(embedStringFile "rts/c/timing.h")         lock_h = $(embedStringFile "rts/c/lock.h")+        tuning_h = $(embedStringFile "rts/c/tuning.h")  compileFun :: (Name, Function op) -> CompilerM op s (C.Definition, C.Func) compileFun (fname, Function _ outputs inputs body _ _) = do@@ -1805,10 +1811,14 @@   decl [C.cdecl|$ty:ct $id:name;|]  compileCode (DeclareArray name DefaultSpace t vs) = do-  let ct = primTypeToCType t-      vs' = [[C.cinit|$exp:(compilePrimValue v)|] | v <- vs]   name_realtype <- newVName $ baseString name ++ "_realtype"-  libDecl [C.cedecl|static $ty:ct $id:name_realtype[$int:(length vs)] = {$inits:vs'};|]+  let ct = primTypeToCType t+  case vs of+    ArrayValues vs' -> do+      let vs'' = [[C.cinit|$exp:(compilePrimValue v)|] | v <- vs']+      libDecl [C.cedecl|static $ty:ct $id:name_realtype[$int:(length vs')] = {$inits:vs''};|]+    ArrayZeros n ->+      libDecl [C.cedecl|static $ty:ct $id:name_realtype[$int:n];|]   -- Fake a memory block.   contextField (pretty name)     [C.cty|struct memblock|] $

@@ -28,7 +28,8 @@  -- | Whether an option accepts an argument. data OptionArgument = NoArgument-                    | RequiredArgument+                    | RequiredArgument String+                    -- ^ The 'String' becomes part of the help text.                     | OptionalArgument  -- | Generate an option parser as a function of the given name, that@@ -52,7 +53,8 @@            panic(-1, "Missing argument for option %s\n", argv[optind-1]);          }          if ($id:chosen_option == '?') {-           panic(-1, "Unknown option %s\n", argv[optind-1]);+           fprintf(stderr, "Usage: %s: %s\n", fut_progname, $string:option_descriptions);+           panic(1, "Unknown option: %s\n", argv[optind-1]);          }        }        return optind;@@ -62,15 +64,28 @@         option_string = ':' : optionString options         option_applications = optionApplications chosen_option options         option_fields = optionFields options+        option_descriptions = unwords $ map describeOption options +describeOption :: Option -> String+describeOption opt =+  concat [ "["+         , maybe "" (\c -> "-" ++ [c] ++ "/") $ optionShortName opt+         , "--" ++ optionLongName opt+         , case optionArgument opt of+             NoArgument -> ""+             RequiredArgument what -> " " ++ what+             OptionalArgument -> " [ARG]"+         , "]"+         ]+ optionFields :: [Option] -> [C.Initializer] optionFields = zipWith field [(1::Int)..]   where field i option =           [C.cinit| { $string:(optionLongName option), $id:arg, NULL, $int:i } |]           where arg = case optionArgument option of-                        NoArgument       -> "no_argument"-                        RequiredArgument -> "required_argument"-                        OptionalArgument -> "optional_argument"+                        NoArgument         -> "no_argument"+                        RequiredArgument _ -> "required_argument"+                        OptionalArgument   -> "optional_argument"  optionApplications :: String -> [Option] -> [C.Stm] optionApplications chosen_option = zipWith check [(1::Int)..]@@ -86,6 +101,6 @@           short <- optionShortName option           return $ short :             case optionArgument option of-              NoArgument       -> ""-              RequiredArgument -> ":"-              OptionalArgument -> "::"+              NoArgument         -> ""+              RequiredArgument _ -> ":"+              OptionalArgument   -> "::"

@@ -114,7 +114,7 @@                  CompilerM op s ()  -- | Create a static array of values - initialised at load time.-type StaticArray op s = VName -> Imp.SpaceId -> PrimType -> [PrimValue] -> CompilerM op s ()+type StaticArray op s = VName -> Imp.SpaceId -> PrimType -> Imp.ArrayContents -> CompilerM op s ()  -- | Construct the C# array being returned from an entry point. type EntryOutput op s = VName -> Imp.SpaceId ->@@ -381,6 +381,11 @@          , optionAction =            [ Reassign (Var "EntryPoint") $ Var "optarg" ]          }+  , Option { optionLongName = "tuning"+           , optionShortName = Nothing+           , optionArgument = RequiredArgument+           , optionAction = [Escape "FutharkConfigLoadTuning(ref Cfg, optarg);"]+           }   ]  -- | The class generated by the code generator must have a@@ -644,7 +649,7 @@ entryPointOutput (Imp.TransparentValue (Imp.ArrayValue mem _ Imp.DefaultSpace bt ept dims)) = do   let src = Var $ compileName mem   let createTuple = "createTuple_" ++ compilePrimTypeExt bt ept-  return $ simpleCall createTuple [src, CreateArray (Primitive $ CSInt Int64T) $ map compileDim dims]+  return $ simpleCall createTuple [src, CreateArray (Primitive $ CSInt Int64T) $ Right $ map compileDim dims]  entryPointOutput (Imp.TransparentValue (Imp.ArrayValue mem _ (Imp.Space sid) bt ept dims)) = do   unRefMem mem (Imp.Space sid)@@ -804,7 +809,7 @@ printStm (Imp.ArrayValue mem memsize space bt ept (outer:shape)) ind e = do   ptr <- newVName "shapePtr"   first <- newVName "printFirst"-  let size = callMethod (CreateArray (Primitive $ CSInt Int32T) $ map compileDim $ outer:shape)+  let size = callMethod (CreateArray (Primitive $ CSInt Int32T) $ Right $ map compileDim $ outer:shape)                  "Aggregate" [ Integer 1                              , Lambda (Tuple [Var "acc", Var "val"])                                       [Exp $ BinOp "*" (Var "acc") (Var "val")]@@ -1257,13 +1262,15 @@   stms [Assign (Var $ "init_"++name') $         simpleCall "unwrapArray"          [-           CreateArray (compilePrimTypeToAST t) (map compilePrimValue vs)+           case vs of Imp.ArrayValues vs' ->+                        CreateArray (compilePrimTypeToAST t) $ Right $ map compilePrimValue vs'+                      Imp.ArrayZeros n ->+                        CreateArray (compilePrimTypeToAST t) $ Left n          , simpleCall "sizeof" [Var $ compilePrimType t]          ]        , Assign (Var name') $ Var ("init_"++name')        ]   where name' = compileName name-  compileCode (Imp.DeclareArray name (Space space) t vs) =   join $ asks envStaticArray <*>

@@ -143,7 +143,7 @@            | Call CSExp [CSArg]            | CallMethod CSExp CSExp [CSArg]            | CreateObject CSExp [CSArg]-           | CreateArray CSType [CSExp]+           | CreateArray CSType (Either Int [CSExp])            | CreateSystemTuple [CSExp]            | AllocArray CSType CSExp            | Cast CSType CSExp@@ -182,7 +182,8 @@   ppr (Call fun args) = ppr fun <> parens(commasep $ map ppr args)   ppr (CallMethod obj method args) = ppr obj <> dot <> ppr method <> parens(commasep $ map ppr args)   ppr (CreateObject className args) = text "new" <+> ppr className <> parens(commasep $ map ppr args)-  ppr (CreateArray t vs) = text "new" <+> ppr t <> text "[]" <+> braces(commasep $ map ppr vs)+  ppr (CreateArray t (Left n)) = text "new" <+> ppr t <> brackets (ppr n)+  ppr (CreateArray t (Right vs)) = text "new" <+> ppr t <> text "[]" <+> braces(commasep $ map ppr vs)   ppr (CreateSystemTuple exps) = text "Tuple.Create" <> parens(commasep $ map ppr exps)   ppr (Tuple exps) = parens(commasep $ map ppr exps)   ppr (Array exps) = braces(commasep $ map ppr exps) -- uhoh is this right?

@@ -91,7 +91,7 @@                  CompilerM op s ()  -- | Create a static array of values - initialised at load time.-type StaticArray op s = VName -> Imp.SpaceId -> PrimType -> [PrimValue] -> CompilerM op s ()+type StaticArray op s = VName -> Imp.SpaceId -> PrimType -> Imp.ArrayContents -> CompilerM op s ()  -- | Construct the Python array being returned from an entry point. type EntryOutput op s = VName -> Imp.SpaceId ->@@ -249,7 +249,7 @@ standardOptions = [   Option { optionLongName = "write-runtime-to"          , optionShortName = Just 't'-         , optionArgument = RequiredArgument+         , optionArgument = RequiredArgument "str"          , optionAction =            [              If (Var "runtime_file")@@ -260,7 +260,7 @@          },   Option { optionLongName = "runs"          , optionShortName = Just 'r'-         , optionArgument = RequiredArgument+         , optionArgument = RequiredArgument "str"          , optionAction =            [ Assign (Var "num_runs") $ Var "optarg"            , Assign (Var "do_warmup_run") $ Bool True@@ -268,7 +268,7 @@          },   Option { optionLongName = "entry-point"          , optionShortName = Just 'e'-         , optionArgument = RequiredArgument+         , optionArgument = RequiredArgument "str"          , optionAction =            [ Assign (Var "entry_point") $ Var "optarg" ]          },@@ -277,6 +277,11 @@          , optionShortName = Just 'b'          , optionArgument = NoArgument          , optionAction = [Pass]+         },+  Option { optionLongName = "tuning"+         , optionShortName = Nothing+         , optionArgument = RequiredArgument "open"+         , optionAction = [Exp $ simpleCall "read_tuning_file" [Var "sizes", Var "optarg"]]          }   ] @@ -312,7 +317,9 @@                             Just _  -> ""   return $ maybe_shebang ++     pretty (PyProg $ imports ++-            [Import "argparse" Nothing] +++            [ Import "argparse" Nothing+            , Assign (Var "sizes") $ Dict []+            ] ++             defines ++             [Escape pyUtility] ++             prog')@@ -879,9 +886,15 @@   -- to prevent it from going "out-of-scope" before calling   -- unwrapArray (which internally uses the .ctype method); see   -- https://docs.scipy.org/doc/numpy/reference/generated/numpy.ndarray.ctypes.html-  atInit $ Assign (Field (Var "self") arr_name) $ Call (Var "np.array")-    [Arg $ List $ map compilePrimValue vs,-     ArgKeyword "dtype" $ Var $ compilePrimToNp t]+  atInit $ Assign (Field (Var "self") arr_name) $ case vs of+    Imp.ArrayValues vs' ->+      Call (Var "np.array")+      [Arg $ List $ map compilePrimValue vs',+       ArgKeyword "dtype" $ Var $ compilePrimToNp t]+    Imp.ArrayZeros n ->+      Call (Var "np.zeros")+      [Arg $ Integer $ fromIntegral n,+       ArgKeyword "dtype" $ Var $ compilePrimToNp t]   atInit $     Assign (Field (Var "self") name') $     simpleCall "unwrapArray" [Field (Var "self") arr_name]

@@ -4,6 +4,7 @@   , pyUtility   , pyValues   , pyPanic+  , pyTuning   ) where  import Data.FileEmbed@@ -19,3 +20,6 @@  pyPanic :: String pyPanic = $(embedStringFile "rts/python/panic.py")++pyTuning :: String+pyTuning = $(embedStringFile "rts/python/tuning.py")

@@ -24,7 +24,7 @@  -- | Whether an option accepts an argument. data OptionArgument = NoArgument-                    | RequiredArgument+                    | RequiredArgument String                     | OptionalArgument  -- | Generate option parsing code that accepts the given command line options.  Will read from @sys.argv@.@@ -43,13 +43,15 @@   map executeOption options   where parseOption option =           Exp $ Call (Var "parser.add_argument") $-          map (Arg . String) name_args ++ argument_args+          map (Arg . String) name_args +++          argument_args           where name_args = maybe id ((:) . ('-':) . (:[])) (optionShortName option)                             ["--" ++ optionLongName option]                 argument_args = case optionArgument option of-                  RequiredArgument ->+                  RequiredArgument t ->                     [ArgKeyword "action" (String "append"),-                     ArgKeyword "default" $ List []]+                     ArgKeyword "default" $ List [],+                     ArgKeyword "type" $ Var t]                    NoArgument ->                     [ArgKeyword "action" (String "append_const"),

@@ -34,10 +34,15 @@             [Assign (Var "synchronous") $ Bool False,              Assign (Var "preferred_platform") None,              Assign (Var "preferred_device") None,+             Assign (Var "default_threshold") None,+             Assign (Var "default_group_size") None,+             Assign (Var "default_num_groups") None,+             Assign (Var "default_tile_size") None,              Assign (Var "fut_opencl_src") $ RawStringLiteral $ opencl_prelude ++ opencl_code,              Escape pyValues,              Escape pyFunctions,-             Escape pyPanic]+             Escape pyPanic,+             Escape pyTuning]       let imports = [Import "sys" Nothing,                      Import "numpy" $ Just "np",                      Import "ctypes" $ Just "ct",@@ -50,23 +55,59 @@                                        , "interactive=False"                                        , "platform_pref=preferred_platform"                                        , "device_pref=preferred_device"-                                       , "default_group_size=None"-                                       , "default_num_groups=None"-                                       , "default_tile_size=None"-                                       , "sizes={}"]+                                       , "default_group_size=default_group_size"+                                       , "default_num_groups=default_num_groups"+                                       , "default_tile_size=default_tile_size"+                                       , "default_threshold=default_threshold"+                                       , "sizes=sizes"]                         [Escape $ openClInit types assign sizes]           options = [ Option { optionLongName = "platform"                              , optionShortName = Just 'p'-                             , optionArgument = RequiredArgument+                             , optionArgument = RequiredArgument "str"                              , optionAction =                                [ Assign (Var "preferred_platform") $ Var "optarg" ]                              }                     , Option { optionLongName = "device"                              , optionShortName = Just 'd'-                             , optionArgument = RequiredArgument+                             , optionArgument = RequiredArgument "str"                              , optionAction =                                [ Assign (Var "preferred_device") $ Var "optarg" ]-                             }]+                             }+                    , Option { optionLongName = "default-threshold"+                             , optionShortName = Nothing+                             , optionArgument = RequiredArgument "int"+                             , optionAction =+                               [ Assign (Var "default_threshold") $ Var "optarg" ]+                             }+                    , Option { optionLongName = "default-group-size"+                             , optionShortName = Nothing+                             , optionArgument = RequiredArgument "int"+                             , optionAction =+                               [ Assign (Var "default_group_size") $ Var "optarg" ]+                             }+                    , Option { optionLongName = "default-num-groups"+                             , optionShortName = Nothing+                             , optionArgument = RequiredArgument "int"+                             , optionAction =+                               [ Assign (Var "default_num_groups") $ Var "optarg" ]+                             }+                    , Option { optionLongName = "default-tile-size"+                             , optionShortName = Nothing+                             , optionArgument = RequiredArgument "int"+                             , optionAction =+                               [ Assign (Var "default_tile_size") $ Var "optarg" ]+                             }+                    , Option { optionLongName = "size"+                             , optionShortName = Nothing+                             , optionArgument = RequiredArgument "size_assignment"+                             , optionAction =+                                 [Assign (Index (Var "sizes")+                                          (IdxExp (Index (Var "optarg")+                                                   (IdxExp (Integer 0)))))+                                   (Index (Var "optarg") (IdxExp (Integer 1)))+                                 ]+                             }+                    ]        Right <$> Py.compileProg module_name constructor imports defines operations ()         [Exp $ Py.simpleCall "self.queue.finish" []] options prog'@@ -216,14 +257,22 @@ staticOpenCLArray name "device" t vs = do   mapM_ Py.atInit <=< Py.collect $ do     -- Create host-side Numpy array with intended values.-    Py.stm $ Assign (Var name') $-      Call (Var "np.array")-      [Arg $ List $ map Py.compilePrimValue vs,-       ArgKeyword "dtype" $ Var $ Py.compilePrimToNp t]+    Py.stm $ Assign (Var name') $ case vs of+      Imp.ArrayValues vs' ->+        Call (Var "np.array")+        [Arg $ List $ map Py.compilePrimValue vs',+         ArgKeyword "dtype" $ Var $ Py.compilePrimToNp t]+      Imp.ArrayZeros n ->+        Call (Var "np.zeros")+        [Arg $ Integer $ fromIntegral n,+         ArgKeyword "dtype" $ Var $ Py.compilePrimToNp t] +    let num_elems = case vs of Imp.ArrayValues vs' -> length vs'+                               Imp.ArrayZeros n -> n+     -- Create memory block on the device.     static_mem <- newVName "static_mem"-    let size = Integer $ fromIntegral (length vs) * Imp.primByteSize t+    let size = Integer $ toInteger num_elems * Imp.primByteSize t     allocateOpenCLBuffer static_mem size "device"      -- Copy Numpy array to the device memory block.

@@ -35,6 +35,7 @@                                    default_group_size=default_group_size,                                    default_num_groups=default_num_groups,                                    default_tile_size=default_tile_size,+                                   default_threshold=default_threshold,                                    size_heuristics=size_heuristics,                                    required_types=$types',                                    user_sizes=sizes,@@ -68,6 +69,7 @@                                        NumGroups     -> String "num_groups"                                        GroupSize     -> String "group_size"                                        TileSize      -> String "tile_size"+                                       Threshold     -> String "threshold"                  what' = case what of                           HeuristicConst x -> Integer $ toInteger x

@@ -27,7 +27,7 @@                    Import "numpy" $ Just "np",                    Import "ctypes" $ Just "ct",                    Import "time" Nothing]-        defines = [Escape pyValues, Escape pyFunctions, Escape pyPanic]+        defines = [Escape pyValues, Escape pyFunctions, Escape pyPanic, Escape pyTuning]         operations :: GenericPython.Operations Imp.Sequential ()         operations = GenericPython.defaultOperations                      { GenericPython.opsCompiler = const $ return ()

@@ -34,6 +34,7 @@   , index   , ErrorMsg(..)   , ErrorMsgPart(..)+  , ArrayContents(..)      -- * Typed enumerations   , Count (..)@@ -137,13 +138,22 @@ -- | Type alias for namespace control. type Function = FunctionT +-- | The contents of a statically declared constant array.  Such+-- arrays are always unidimensional, and reshaped if necessary in the+-- code that uses them.+data ArrayContents = ArrayValues [PrimValue]+                     -- ^ Precisely these values.+                   | ArrayZeros Int+                     -- ^ This many zeroes.+                     deriving (Show)+ data Code a = Skip             | Code a :>>: Code a             | For VName IntType Exp (Code a)             | While Exp (Code a)             | DeclareMem VName Space             | DeclareScalar VName PrimType-            | DeclareArray VName Space PrimType [PrimValue]+            | DeclareArray VName Space PrimType ArrayContents               -- ^ Create an array containing the given values.  The               -- lifetime of the array will be the entire application.               -- This is mostly used for constant arrays, but also for@@ -295,6 +305,10 @@   ppr (ConstSize x) = ppr x   ppr (VarSize v)   = ppr v +instance Pretty ArrayContents where+  ppr (ArrayValues vs) = braces (commasep $ map ppr vs)+  ppr (ArrayZeros n) = braces (text "0") <+> text "*" <+> ppr n+ instance Pretty op => Pretty (Code op) where   ppr (Op op) = ppr op   ppr Skip   = text "skip"@@ -313,7 +327,7 @@     text "var" <+> ppr name <> text ":" <+> ppr t   ppr (DeclareArray name space t vs) =     text "array" <+> ppr name <> text "@" <> ppr space <+> text ":" <+> ppr t <+>-    equals <+> braces (commasep $ map ppr vs)+    equals <+> ppr vs   ppr (Allocate name e space) =     ppr name <+> text "<-" <+> text "malloc" <> parens (ppr e) <> ppr space   ppr (Free name space) =

@@ -103,11 +103,11 @@  instance Pretty KernelUse where   ppr (ScalarUse name t) =-    text "scalar_copy" <> parens (commasep [ppr name, ppr t])+    oneLine $ text "scalar_copy" <> parens (commasep [ppr name, ppr t])   ppr (MemoryUse name) =-    text "mem_copy" <> parens (commasep [ppr name])+    oneLine $ text "mem_copy" <> parens (commasep [ppr name])   ppr (ConstUse name e) =-    text "const" <> parens (commasep [ppr name, ppr e])+    oneLine $ text "const" <> parens (commasep [ppr name, ppr e])  instance Pretty HostOp where   ppr (GetSize dest key size_class) =

@@ -78,9 +78,9 @@   , sComment   , sIf, sWhen, sUnless   , sOp-  , sAlloc+  , sDeclareMem, sAlloc, sAlloc_   , sArray, sAllocArray, sStaticArray-  , sWrite+  , sWrite, sUpdate   , (<--)   )   where@@ -675,11 +675,8 @@ defCompileBasicOp _ Index{} =   return () -defCompileBasicOp (Pattern _ [pe]) (Update _ slice se) = do-  MemLocation mem shape ixfun <- entryArrayLocation <$> lookupArray (patElemName pe)-  let memdest = sliceArray (MemLocation mem shape ixfun) $-                map (fmap (compileSubExpOfType int32)) slice-  copyDWIMDest (ArrayDestination $ Just memdest) [] se []+defCompileBasicOp (Pattern _ [pe]) (Update _ slice se) =+  sUpdate (patElemName pe) (map (fmap (compileSubExpOfType int32)) slice) se  defCompileBasicOp (Pattern _ [pe]) (Replicate (Shape ds) se) = do   ds' <- mapM compileSubExp ds@@ -736,7 +733,7 @@       dest_space <- entryMemSpace <$> lookupMemory (memLocationName dest_mem)       let t = primValueType v       static_array <- newVName "static_array"-      emit $ Imp.DeclareArray static_array dest_space t vs+      emit $ Imp.DeclareArray static_array dest_space t $ Imp.ArrayValues vs       let static_src = MemLocation static_array [Imp.ConstSize $ fromIntegral $ length es] $                        IxFun.iota [fromIntegral $ length es]           num_bytes = Imp.ConstSize $ fromIntegral (length es) * primByteSize t@@ -1237,19 +1234,33 @@ sOp :: op -> ImpM lore op () sOp = emit . Imp.Op -sAlloc :: String -> Count Bytes -> Space -> ImpM lore op VName-sAlloc name size space = do+dSize :: Imp.Count u -> ImpM lore op Imp.Size+dSize size =+  case Imp.innerExp size of+    Imp.LeafExp (Imp.ScalarVar size') _ -> return $ Imp.VarSize size'+    Imp.ValueExp (IntValue (Int64Value v)) -> return $ Imp.ConstSize v+    _ -> do size_var <- dPrim "local_buf_size" int32+            size_var <-- Imp.innerExp size+            return $ Imp.VarSize size_var++sDeclareMem :: String -> Count Bytes -> Space -> ImpM lore op (VName, Imp.MemSize)+sDeclareMem name size space = do+   name' <- newVName name-  size' <- case Imp.innerExp size of-             Imp.LeafExp (Imp.ScalarVar size') _ -> return $ Imp.VarSize size'-             Imp.ValueExp (IntValue (Int64Value v)) -> return $ Imp.ConstSize v-             _ -> do size_var <- dPrim "local_buf_size" int32-                     size_var <-- Imp.innerExp size-                     return $ Imp.VarSize size_var+  size' <- dSize size   emit $ Imp.DeclareMem name' space-  fake <- asks $ elem space . envFakeMemory-  unless fake $ emit $ Imp.Allocate name' size space   addVar name' $ MemVar Nothing $ MemEntry size' space+  return (name', size')++sAlloc_ :: VName -> Imp.MemSize -> Space -> ImpM lore op ()+sAlloc_ name' size' space = do+  fake <- asks $ elem space . envFakeMemory+  unless fake $ emit $ Imp.Allocate name' (Imp.memSizeToExp size') space++sAlloc :: String -> Count Bytes -> Space -> ImpM lore op VName+sAlloc name size space = do+  (name', size') <- sDeclareMem name size space+  sAlloc_ name' size' space   return name'  sArray :: String -> PrimType -> ShapeBase SubExp -> MemBind -> ImpM lore op VName@@ -1268,21 +1279,28 @@     ArrayIn mem $ IxFun.iota $ map (primExpFromSubExp int32) $ shapeDims shape  -- | Uses linear/iota index function.-sStaticArray :: String -> Space -> PrimType -> [PrimValue] -> ImpM lore op VName+sStaticArray :: String -> Space -> PrimType -> Imp.ArrayContents -> ImpM lore op VName sStaticArray name space pt vs = do-  let shape = Shape [constant $ length vs]-      size = Imp.ConstSize $ fromIntegral (length vs) * primByteSize pt+  let num_elems = case vs of Imp.ArrayValues vs' -> genericLength vs'+                             Imp.ArrayZeros n -> fromIntegral n+      shape = Shape [constant num_elems]+      mem_size = Imp.ConstSize $ num_elems * primByteSize pt   mem <- newVName $ name ++ "_mem"   emit $ Imp.DeclareArray mem space pt vs-  addVar mem $ MemVar Nothing $ MemEntry size space-  sArray name pt shape $-    ArrayIn mem $ IxFun.iota $ map (primExpFromSubExp int32) $ shapeDims shape+  addVar mem $ MemVar Nothing $ MemEntry mem_size space+  sArray name pt shape $ ArrayIn mem $ IxFun.iota [fromIntegral num_elems]  sWrite :: VName -> [Imp.Exp] -> PrimExp Imp.ExpLeaf -> ImpM lore op () sWrite arr is v = do   (mem, space, offset) <- fullyIndexArray arr is   vol <- asks envVolatility   emit $ Imp.Write mem offset (primExpType v) space vol v++sUpdate :: VName -> Slice Imp.Exp -> SubExp -> ImpM lore op ()+sUpdate arr slice v = do+  MemLocation mem shape ixfun <- entryArrayLocation <$> lookupArray arr+  let memdest = sliceArray (MemLocation mem shape ixfun) slice+  copyDWIMDest (ArrayDestination $ Just memdest) [] v []  -- | ASsignment. (<--) :: VName -> Imp.Exp -> ImpM lore op ()

@@ -23,9 +23,9 @@ import qualified Futhark.CodeGen.ImpGen as ImpGen import Futhark.CodeGen.ImpGen.Kernels.Base import Futhark.CodeGen.ImpGen.Kernels.SegRed+import Futhark.CodeGen.ImpGen.Kernels.SegGenRed import Futhark.CodeGen.ImpGen (sFor, sWhen,-                               sOp,-                               dPrim, dPrim_, dPrimV)+                               sOp) import Futhark.CodeGen.ImpGen.Kernels.Transpose import qualified Futhark.Representation.ExplicitMemory.IndexFunction as IxFun import Futhark.CodeGen.SetDefaultSpace@@ -108,14 +108,12 @@                                show (baseTag $ kernelGlobalThreadIdVar constants)             } --- First handle the simple case of a non-segmented reduction.  Our--- strategy is the conventional approach of generating two kernels:--- one where each group is given a chunk of the total input and--- produces a partial result per group, and then a final kernel that--- combines the per-group partial results. kernelCompiler pat (SegRed space comm red_op nes _ body) =   compileSegRed pat space comm red_op nes body +kernelCompiler pat (SegGenRed space ops _ body) =+  compileSegGenRed pat space ops body+ kernelCompiler pat e =   compilerBugS $ "ImpGen.kernelCompiler: Invalid pattern\n  " ++   pretty pat ++ "\nfor expression\n  " ++ pretty e@@ -127,33 +125,11 @@   n' <- ImpGen.compileSubExp n   x' <- ImpGen.compileSubExp x   s' <- ImpGen.compileSubExp s-  destloc <- ImpGen.entryArrayLocation <$> ImpGen.lookupArray (patElemName pe)-  (constants, set_constants) <- simpleKernelConstants n' "iota" -  sKernel constants "iota" $ do-    set_constants-    let gtid = kernelGlobalThreadId constants-    sWhen (kernelThreadActive constants) $ do-      (destmem, destspace, destidx) <--        ImpGen.fullyIndexArray' destloc [gtid] (IntType et)--      ImpGen.emit $-        Imp.Write destmem destidx (IntType et) destspace Imp.Nonvolatile $-        Imp.ConvOpExp (SExt Int32 et) gtid * s' + x'--expCompiler (Pattern _ [pe]) (BasicOp (Replicate (Shape ds) se)) = do-  t <- subExpType se--  dims <- mapM ImpGen.compileSubExp $ ds ++ arrayDims t-  (constants, set_constants) <--    simpleKernelConstants (product dims) "replicate"--  let is' = unflattenIndex dims $ kernelGlobalThreadId constants+  sIota (patElemName pe) n' x' s' et -  sKernel constants "replicate" $ do-    set_constants-    sWhen (kernelThreadActive constants) $-      ImpGen.copyDWIM (patElemName pe) is' se $ drop (length ds) is'+expCompiler (Pattern _ [pe]) (BasicOp (Replicate shape se)) =+  sReplicate (patElemName pe) shape se  -- Allocation in the "local" space is just a placeholder. expCompiler _ (Op (Alloc _ (Space "local"))) =@@ -195,29 +171,7 @@           srcmem (bytes srcoffset) srcspace $           (n * row_size) `Imp.withElemType` bt -  | otherwise = do--  -- Note that the shape of the destination and the source are-  -- necessarily the same.-  let shape = map Imp.sizeToExp srcshape-      shape_se = map (Imp.innerExp . ImpGen.dimSizeToExp) srcshape-      kernel_size = Imp.innerExp n * product (drop 1 shape)--  (constants, set_constants) <- simpleKernelConstants kernel_size "copy"--  sKernel constants "copy" $ do-    set_constants--    let gtid = kernelGlobalThreadId constants-        dest_is = unflattenIndex shape_se gtid-        src_is = dest_is--    (_, destspace, destidx) <- ImpGen.fullyIndexArray' destloc dest_is bt-    (_, srcspace, srcidx) <- ImpGen.fullyIndexArray' srcloc src_is bt--    sWhen (gtid .<. kernel_size) $ ImpGen.emit $-      Imp.Write destmem destidx bt destspace Imp.Nonvolatile $-      Imp.index srcmem srcidx bt srcspace Imp.Nonvolatile+  | otherwise = sCopy bt destloc srcloc n  mapTransposeForType :: PrimType -> ImpGen.ImpM ExplicitMemory Imp.HostOp Name mapTransposeForType bt = do@@ -372,39 +326,6 @@           let (mapped, notmapped) = splitAt r1 shape               (pretrans, posttrans) = f $ splitAt r2 notmapped           in (product mapped, product pretrans, product posttrans)--simpleKernelConstants :: Imp.Exp -> String-                      -> CallKernelGen (KernelConstants, ImpGen.ImpM InKernel Imp.KernelOp ())-simpleKernelConstants kernel_size desc = do-  thread_gtid <- newVName $ desc ++ "_gtid"-  thread_ltid <- newVName $ desc ++ "_ltid"-  group_id <- newVName $ desc ++ "_gid"-  (group_size, num_groups) <- computeMapKernelGroups kernel_size-  let set_constants = do-        dPrim_ thread_gtid int32-        dPrim_ thread_ltid int32-        dPrim_ group_id int32-        sOp (Imp.GetGlobalId thread_gtid 0)-        sOp (Imp.GetLocalId thread_ltid 0)-        sOp (Imp.GetGroupId group_id 0)--  return (KernelConstants-          (Imp.var thread_gtid int32) (Imp.var thread_ltid int32) (Imp.var group_id int32)-          thread_gtid thread_ltid group_id-          group_size num_groups (group_size*num_groups) 0-          [] (Imp.var thread_gtid int32 .<. kernel_size) mempty,--          set_constants)--computeMapKernelGroups :: Imp.Exp -> CallKernelGen (Imp.Exp, Imp.Exp)-computeMapKernelGroups kernel_size = do-  group_size <- dPrim "group_size" int32-  fname <- asks ImpGen.envFunction-  let group_size_var = Imp.var group_size int32-      group_size_key = keyWithEntryPoint fname $ nameFromString $ pretty group_size-  sOp $ Imp.GetSize group_size group_size_key Imp.SizeGroup-  num_groups <- dPrimV "num_groups" $ kernel_size `quotRoundingUp` Imp.ConvOpExp (SExt Int32 Int32) group_size_var-  return (Imp.var group_size int32, Imp.var num_groups int32)  compileKernelBody :: Pattern InKernel                   -> KernelConstants

@@ -19,6 +19,14 @@   , groupScan   , isActive   , sKernel+  , sReplicate+  , sIota+  , sCopy++  , atomicUpdate+  , atomicUpdateLocking+  , Locking(..)+  , AtomicUpdate   )   where @@ -30,7 +38,7 @@ import qualified Data.Set as S import Data.List -import Prelude hiding (quot)+import Prelude hiding (quot, rem)  import Futhark.Error import Futhark.MonadFreshNames@@ -45,7 +53,7 @@                                dPrim, dPrim_, dPrimV) import Futhark.Tools (partitionChunkedKernelLambdaParameters) import Futhark.Util.IntegralExp (quotRoundingUp, quot, rem, IntegralExp)-import Futhark.Util (splitAt3, maybeNth)+import Futhark.Util (splitAt3, maybeNth, takeLast)  type CallKernelGen = ImpGen.ImpM ExplicitMemory Imp.HostOp type InKernelGen = ImpGen.ImpM InKernel Imp.KernelOp@@ -271,12 +279,17 @@   -- Check if bucket is in-bounds   bucket' <- mapM ImpGen.compileSubExp bucket   w' <- mapM ImpGen.compileSubExp w-  sWhen (indexInBounds bucket' w') $-    atomicUpdate arrs bucket op values locking+  num_locks <- ImpGen.compileSubExpOfType int32 . arraySize 0 <$> lookupType locks+  let locking = Locking locks 0 1 0 $ (`rem` num_locks) . sum+      values_params = takeLast (length values) $ lambdaParams op++  sWhen (indexInBounds bucket' w') $ do+    forM_ (zip values_params values) $ \(p, v) ->+      ImpGen.copyDWIM (paramName p) [] v []+    atomicUpdate arrs bucket' op locking   where indexInBounds inds bounds =           foldl1 (.&&.) $ zipWith checkBound inds bounds           where checkBound ind bound = 0 .<=. ind .&&. ind .<. bound-        locking = Locking locks 0 1 0  compileKernelExp _ dest e =   compilerBugS $ unlines ["Invalid target", "  " ++ show dest,@@ -290,35 +303,57 @@   ImpGen.emit =<< ImpGen.subImpM_ (inKernelOperations constants') m  -- | Locking strategy used for an atomic update.-data Locking = Locking { lockingArray :: VName -- ^ Array containing the lock.-                       , lockingIsUnlocked :: Imp.Exp -- ^ Value for us to consider the lock free.-                       , lockingToLock :: Imp.Exp -- ^ What to write when we lock it.-                       , lockingToUnlock :: Imp.Exp -- ^ What to write when we unlock it.-                       }+data Locking =+  Locking { lockingArray :: VName+            -- ^ Array containing the lock.+          , lockingIsUnlocked :: Imp.Exp+            -- ^ Value for us to consider the lock free.+          , lockingToLock :: Imp.Exp+            -- ^ What to write when we lock it.+          , lockingToUnlock :: Imp.Exp+            -- ^ What to write when we unlock it.+          , lockingMapping :: [Imp.Exp] -> Imp.Exp+            -- ^ A transformation from the logical lock index to the+            -- physical position in the array.  This can also be used+            -- to make the lock array smaller.+          } +-- | A function for generating code for an atomic update.  Assumes+-- that the bucket is in-bounds.+type AtomicUpdate lore =+  [VName] -> [Imp.Exp] -> ImpGen.ImpM lore Imp.KernelOp ()+ atomicUpdate :: ExplicitMemorish lore =>-                [VName] -> [SubExp] -> Lambda lore -> [SubExp] -> Locking+                [VName] -> [Imp.Exp] -> Lambda lore -> Locking              -> ImpGen.ImpM lore Imp.KernelOp ()+atomicUpdate arrs bucket lam locking =+  case atomicUpdateLocking lam of+    Left f -> f arrs bucket+    Right f -> f locking arrs bucket -atomicUpdate arrs bucket lam values _+-- | 'atomicUpdate', but where it is explicitly visible whether a+-- locking strategy is necessary.+atomicUpdateLocking :: ExplicitMemorish lore =>+                       Lambda lore+                    -> Either (AtomicUpdate lore) (Locking -> AtomicUpdate lore)++atomicUpdateLocking lam   | Just ops_and_ts <- splitOp lam,-    all ((==32) . primBitSize . snd) ops_and_ts =+    all (\(_, t, _) -> primBitSize t == 32) ops_and_ts = Left $ \arrs bucket ->   -- If the operator is a vectorised binary operator on 32-bit values,   -- we can use a particularly efficient implementation. If the   -- operator has an atomic implementation we use that, otherwise it   -- is still a binary operator which can be implemented by atomic   -- compare-and-swap if 32 bits.-  forM_ (zip3 arrs ops_and_ts values) $ \(a, (op, t), val) -> do+  forM_ (zip arrs ops_and_ts) $ \(a, (op, t, val)) -> do    -- Common variables.   old <- dPrim "old" t-  bucket' <- mapM ImpGen.compileSubExp bucket -  (arr', _a_space, bucket_offset) <- ImpGen.fullyIndexArray a bucket'+  (arr', _a_space, bucket_offset) <- ImpGen.fullyIndexArray a bucket -  val' <- ImpGen.compileSubExp val   case opHasAtomicSupport old arr' bucket_offset op of-    Just f -> sOp $ f val'+    Just f -> sOp $ f val      Nothing -> do       -- Code generation target:@@ -331,7 +366,7 @@       -- } while(assumed != old);       assumed <- dPrim "assumed" t       run_loop <- dPrimV "run_loop" 1-      ImpGen.copyDWIM old [] (Var a) bucket'+      ImpGen.copyDWIM old [] (Var a) bucket        -- Critical section       x <- dPrim "x" t@@ -344,7 +379,7 @@                       _                 -> (id, id)       sWhile (Imp.var run_loop int32) $ do         assumed <-- Imp.var old t-        x <-- val'+        x <-- val         y <-- Imp.var assumed t         x <-- Imp.BinOpExp op (Imp.var x t) (Imp.var y t)         old_bits <- dPrim "old_bits" int32@@ -359,69 +394,69 @@           let atomic f = Imp.Atomic . f old arr' bucket'           atomic <$> Imp.atomicBinOp bop -atomicUpdate arrs bucket op values locking = do+atomicUpdateLocking op = Right $ \locking arrs bucket -> do   old <- dPrim "old" int32   continue <- dPrimV "continue" true -  -- Check if bucket is in-bounds-  bucket' <- mapM ImpGen.compileSubExp bucket-   -- Correctly index into locks.   (locks', _locks_space, locks_offset) <--    ImpGen.fullyIndexArray (lockingArray locking) bucket'--  -- Preparing parameters-  let (acc_params, arr_params) =-        splitAt (length values) $ lambdaParams op+    ImpGen.fullyIndexArray (lockingArray locking) [lockingMapping locking bucket]    -- Critical section   let try_acquire_lock =         sOp $ Imp.Atomic $         Imp.AtomicCmpXchg old locks' locks_offset (lockingIsUnlocked locking) (lockingToLock locking)       lock_acquired = Imp.var old int32 .==. lockingIsUnlocked locking-      release_lock = ImpGen.everythingVolatile $-                     ImpGen.sWrite (lockingArray locking) bucket' $ lockingToUnlock locking+      -- Even the releasing is done with an atomic rather than a+      -- simple write, for memory coherency reasons.+      release_lock =+        sOp $ Imp.Atomic $+        Imp.AtomicCmpXchg old locks' locks_offset (lockingToLock locking) (lockingToUnlock locking)       break_loop = continue <-- false -  -- We copy the current value and the new value to the parameters.-  -- It is important that the right-hand-side is bound first for the-  -- (rare) case when we are dealing with arrays.-  let bind_acc_params =+  -- Preparing parameters. It is assumed that the caller has already+  -- filled the arr_params. We copy the current value to the+  -- accumulator parameters.+  --+  -- Note the use of 'everythingVolatile' when reading and writing the+  -- buckets.  This was necessary to ensure correct execution on a+  -- newer NVIDIA GPU (RTX 2080).  The 'volatile' modifiers likely+  -- make the writes pass through the (SM-local) L1 cache, which is+  -- necessary here, because we are really doing device-wide+  -- synchronisation without atomics (naughty!).+  let (acc_params, _arr_params) = splitAt (length arrs) $ lambdaParams op+      bind_acc_params =+        ImpGen.everythingVolatile $         ImpGen.sComment "bind lhs" $         forM_ (zip acc_params arrs) $ \(acc_p, arr) ->-        ImpGen.copyDWIM (paramName acc_p) [] (Var arr) bucket'--  let bind_arr_params =-        ImpGen.sComment "bind rhs" $-        forM_ (zip arr_params values) $ \(arr_p, val) ->-        ImpGen.copyDWIM (paramName arr_p) [] val []+        ImpGen.copyDWIM (paramName acc_p) [] (Var arr) bucket    let op_body = ImpGen.sComment "execute operation" $                 ImpGen.compileBody' acc_params $ lambdaBody op -      do_gen_reduce = ImpGen.sComment "update global result" $-                      zipWithM_ (writeArray bucket') arrs $ map (Var . paramName) acc_params+      do_gen_reduce =+        ImpGen.everythingVolatile $+        ImpGen.sComment "update global result" $+        zipWithM_ (writeArray bucket) arrs $ map (Var . paramName) acc_params    -- While-loop: Try to insert your value   sWhile (Imp.var continue Bool) $ do     try_acquire_lock     sWhen lock_acquired $ do-      ImpGen.dLParams $ lambdaParams op-      bind_arr_params+      ImpGen.dLParams acc_params       bind_acc_params       op_body       do_gen_reduce+      sOp Imp.MemFence       release_lock       break_loop     sOp Imp.MemFence-  where writeArray bucket' arr val =-          ImpGen.copyDWIM arr bucket' val []+  where writeArray bucket arr val = ImpGen.copyDWIM arr bucket val []  -- | Horizontally fission a lambda that models a binary operator.-splitOp :: Attributes lore => Lambda lore -> Maybe [(BinOp, PrimType)]+splitOp :: Attributes lore => Lambda lore -> Maybe [(BinOp, PrimType, Imp.Exp)] splitOp lam = mapM splitStm $ bodyResult $ lambdaBody lam   where n = length $ lambdaReturnType lam-        splitStm :: SubExp -> Maybe (BinOp, PrimType)         splitStm (Var res) = do           Let (Pattern [] [pe]) _ (BasicOp (BinOp op (Var x) (Var y))) <-             find (([res]==) . patternNames . stmPattern) $@@ -432,7 +467,7 @@           guard $ paramName xp == x           guard $ paramName yp == y           Prim t <- Just $ patElemType pe-          return (op, t)+          return (op, t, Imp.var (paramName yp) t)         splitStm _ = Nothing  computeKernelUses :: FreeIn a =>@@ -494,10 +529,10 @@         hasExp (ImpGen.ScalarVar e _) = e         hasExp (ImpGen.MemVar e _) = e --- | Only some constant expressions quality as *static* expressions,+-- | Only some constant expressions qualify as *static* expressions, -- which we can use for static memory allocation.  This is a bit of a -- hack, as it is primarly motivated by what you can put as the size--- when daring an array in C.+-- when declaring an array in C. isStaticExp :: Imp.KernelConstExp -> Bool isStaticExp LeafExp{} = True isStaticExp ValueExp{} = True@@ -853,6 +888,8 @@   sComment "restore correct values for first block" $     sWhen is_first_block write_final_result +  sOp Imp.LocalBarrier+ inBlockScan :: Maybe (Imp.Exp -> Imp.Exp -> Imp.Exp)             -> Imp.Exp             -> Imp.Exp@@ -944,6 +981,39 @@         collapse l@((g,_):_) =           (g, map snd l) +computeMapKernelGroups :: Imp.Exp -> CallKernelGen (Imp.Exp, Imp.Exp)+computeMapKernelGroups kernel_size = do+  group_size <- dPrim "group_size" int32+  fname <- asks ImpGen.envFunction+  let group_size_var = Imp.var group_size int32+      group_size_key = keyWithEntryPoint fname $ nameFromString $ pretty group_size+  sOp $ Imp.GetSize group_size group_size_key Imp.SizeGroup+  num_groups <- dPrimV "num_groups" $ kernel_size `quotRoundingUp` Imp.ConvOpExp (SExt Int32 Int32) group_size_var+  return (Imp.var group_size int32, Imp.var num_groups int32)++simpleKernelConstants :: Imp.Exp -> String+                      -> CallKernelGen (KernelConstants, ImpGen.ImpM InKernel Imp.KernelOp ())+simpleKernelConstants kernel_size desc = do+  thread_gtid <- newVName $ desc ++ "_gtid"+  thread_ltid <- newVName $ desc ++ "_ltid"+  group_id <- newVName $ desc ++ "_gid"+  (group_size, num_groups) <- computeMapKernelGroups kernel_size+  let set_constants = do+        dPrim_ thread_gtid int32+        dPrim_ thread_ltid int32+        dPrim_ group_id int32+        sOp (Imp.GetGlobalId thread_gtid 0)+        sOp (Imp.GetLocalId thread_ltid 0)+        sOp (Imp.GetGroupId group_id 0)++  return (KernelConstants+          (Imp.var thread_gtid int32) (Imp.var thread_ltid int32) (Imp.var group_id int32)+          thread_gtid thread_ltid group_id+          group_size num_groups (group_size*num_groups) 0+          [] (Imp.var thread_gtid int32 .<. kernel_size) mempty,++          set_constants)+ sKernel :: KernelConstants -> String -> ImpGen.ImpM InKernel Imp.KernelOp a -> CallKernelGen () sKernel constants name m = do   body <- makeAllMemoryGlobal $@@ -958,3 +1028,69 @@     , Imp.kernelName =         nameFromString $ name ++ "_" ++ show (baseTag $ kernelGlobalThreadIdVar constants)     }++-- | Perform a Replicate with a kernel.+sReplicate :: VName -> Shape -> SubExp+           -> CallKernelGen ()+sReplicate arr (Shape ds) se = do+  t <- subExpType se++  dims <- mapM ImpGen.compileSubExp $ ds ++ arrayDims t+  (constants, set_constants) <-+    simpleKernelConstants (product dims) "replicate"++  let is' = unflattenIndex dims $ kernelGlobalThreadId constants++  sKernel constants "replicate" $ do+    set_constants+    sWhen (kernelThreadActive constants) $+      ImpGen.copyDWIM arr is' se $ drop (length ds) is'++-- | Perform an Iota with a kernel.+sIota :: VName -> Imp.Exp -> Imp.Exp -> Imp.Exp -> IntType+      -> CallKernelGen ()+sIota arr n x s et = do+  destloc <- ImpGen.entryArrayLocation <$> ImpGen.lookupArray arr+  (constants, set_constants) <- simpleKernelConstants n "iota"++  sKernel constants "iota" $ do+    set_constants+    let gtid = kernelGlobalThreadId constants+    sWhen (kernelThreadActive constants) $ do+      (destmem, destspace, destidx) <-+        ImpGen.fullyIndexArray' destloc [gtid] (IntType et)++      ImpGen.emit $+        Imp.Write destmem destidx (IntType et) destspace Imp.Nonvolatile $+        Imp.ConvOpExp (SExt Int32 et) gtid * s + x++sCopy :: PrimType+      -> ImpGen.MemLocation+      -> ImpGen.MemLocation+      -> Imp.Count Imp.Elements+      -> CallKernelGen ()+sCopy bt+  destloc@(ImpGen.MemLocation destmem _ _)+  srcloc@(ImpGen.MemLocation srcmem srcshape _)+  n = do+  -- Note that the shape of the destination and the source are+  -- necessarily the same.+  let shape = map Imp.sizeToExp srcshape+      shape_se = map (Imp.innerExp . ImpGen.dimSizeToExp) srcshape+      kernel_size = Imp.innerExp n * product (drop 1 shape)++  (constants, set_constants) <- simpleKernelConstants kernel_size "copy"++  sKernel constants "copy" $ do+    set_constants++    let gtid = kernelGlobalThreadId constants+        dest_is = unflattenIndex shape_se gtid+        src_is = dest_is++    (_, destspace, destidx) <- ImpGen.fullyIndexArray' destloc dest_is bt+    (_, srcspace, srcidx) <- ImpGen.fullyIndexArray' srcloc src_is bt++    sWhen (gtid .<. kernel_size) $ ImpGen.emit $+      Imp.Write destmem destidx bt destspace Imp.Nonvolatile $+      Imp.index srcmem srcidx bt srcspace Imp.Nonvolatile

@@ -0,0 +1,275 @@+{-# LANGUAGE TypeFamilies #-}+{-# LANGUAGE FlexibleContexts #-}+-- | Our compilation strategy for 'SegGenRed' is based around avoiding+-- bin conflicts.  We do this by splitting the input into chunks, and+-- for each chunk computing a single subhistogram.  Then we combine+-- the subhistograms using an ordinary segmented reduction ('SegRed').+--+-- There are some branches around to efficiently handle the case where+-- we use only a single subhistogram (because it's large), so that we+-- respect the asymptotics, and do not copy the destination array.+module Futhark.CodeGen.ImpGen.Kernels.SegGenRed+  ( compileSegGenRed )+  where++import Control.Monad.Except+import Data.Maybe+import Data.List++import Prelude hiding (quot, rem)++import Futhark.MonadFreshNames+import Futhark.Representation.ExplicitMemory+import qualified Futhark.Representation.ExplicitMemory.IndexFunction as IxFun+import Futhark.Pass.ExplicitAllocations()+import qualified Futhark.CodeGen.ImpCode.Kernels as Imp+import qualified Futhark.CodeGen.ImpGen as ImpGen+import Futhark.CodeGen.ImpGen ((<--),+                               sFor, sComment, sIf, sWhen, sArray,+                               dPrim_, dPrimV)+import Futhark.CodeGen.ImpGen.Kernels.SegRed (compileSegRed')+import Futhark.CodeGen.ImpGen.Kernels.Base+import Futhark.Util.IntegralExp (quotRoundingUp, quot, rem)+import Futhark.Util (chunks, mapAccumLM, splitFromEnd, takeLast)+import Futhark.Construct (fullSliceNum)++prepareAtomicUpdate :: Maybe Locking -> [VName] -> Lambda InKernel+                    -> CallKernelGen (Maybe Locking,+                                      [Imp.Exp] -> ImpGen.ImpM InKernel Imp.KernelOp ())+prepareAtomicUpdate l dests lam =+  -- We need a separate lock array if the opterators are not all of a+  -- particularly simple form that permits pure atomic operations.+  case (l, atomicUpdateLocking lam) of+    (_, Left f) -> return (l, f dests)+    (Just l', Right f) -> return (l, f l' dests)+    (Nothing, Right f) -> do+      -- The number of locks used here is too low, but since we are+      -- currently forced to inline a huge list, I'm keeping it down+      -- for now.  Some quick experiments suggested that it has little+      -- impact anyway (maybe the locking case is just too slow).+      --+      -- A fun solution would also be to use a simple hashing+      -- algorithm to ensure good distribution of locks.+      let num_locks = 10000+      locks <-+        ImpGen.sStaticArray "genred_locks" (Space "device") int32 $+        Imp.ArrayZeros num_locks+      let l' = Locking locks 0 1 0 ((`rem` fromIntegral num_locks) . sum)+      return (Just l', f l' dests)++prepareIntermediateArrays :: [SubExp] -> Imp.Exp -> [GenReduceOp InKernel]+                          -> CallKernelGen+                             [(VName,+                               [VName],+                               [Imp.Exp] -> ImpGen.ImpM InKernel Imp.KernelOp ())]+prepareIntermediateArrays segment_dims num_threads = fmap snd . mapAccumLM onOp Nothing+  where+    onOp l op = do+      -- Determining the degree of cooperation (heuristic):+      -- coop_lvl   := size of histogram (Cooperation level)+      -- num_histos := (threads / coop_lvl) (Number of histograms)+      -- threads    := min(physical_threads, segment_size)+      --+      -- Careful to avoid division by zero when genReduceWidth==0.+      num_histos <- dPrimV "num_histos" $ num_threads `quotRoundingUp`+                    BinOpExp (SMax Int32) 1 (ImpGen.compileSubExpOfType int32 (genReduceWidth op))++      ImpGen.emit $ Imp.DebugPrint "num_histograms" int32 $ Imp.var num_histos int32++      -- Initialise sub-histograms.+      --+      -- If num_histos is 1, then we just reuse the original+      -- destination.  The idea is to avoid a copy if we are writing a+      -- small number of values into a very large prior histogram.++      dests <- forM (zip (genReduceDest op) (genReduceNeutral op)) $ \(dest, ne) -> do+        dest_t <- lookupType dest+        dest_mem <- ImpGen.entryArrayLocation <$> ImpGen.lookupArray dest+        let num_elems = foldl' (*) (Imp.var num_histos int32) $+                        map (ImpGen.compileSubExpOfType int32) $+                        arrayDims dest_t+        let size = Imp.elements num_elems `Imp.withElemType` int32++        (sub_mem, size') <-+          ImpGen.sDeclareMem "subhistogram_mem" size $ Space "device"++        let num_segments = length segment_dims+            sub_shape = Shape (segment_dims++[Var num_histos]) <>+                        stripDims num_segments (arrayShape dest_t)+            sub_membind = ArrayIn sub_mem $ IxFun.iota $+                          map (primExpFromSubExp int32) $ shapeDims sub_shape+        subhisto <- sArray "genred_dest" (elemType dest_t) sub_shape sub_membind++        let unitHistoCase =+              ImpGen.emit $+              Imp.SetMem sub_mem (ImpGen.memLocationName dest_mem) $+              Space "device"++            multiHistoCase = do+              ImpGen.sAlloc_ sub_mem size' $ Space "device"+              sReplicate subhisto (Shape $ segment_dims ++ [Var num_histos, genReduceWidth op]) ne+              subhisto_t <- lookupType subhisto+              let slice = fullSliceNum (map (ImpGen.compileSubExpOfType int32) $ arrayDims subhisto_t) $+                          map (unitSlice 0 . ImpGen.compileSubExpOfType int32) segment_dims +++                          [DimFix 0]+              ImpGen.sUpdate subhisto slice $ Var dest++        sIf (Imp.var num_histos int32 .==. 1) unitHistoCase multiHistoCase++        return subhisto++      (l', do_op) <- prepareAtomicUpdate l dests $ genReduceOp op++      return (l', (num_histos, dests, do_op))++genRedKernel :: [PatElem ExplicitMemory]+             -> KernelSpace+             -> [GenReduceOp InKernel]+             -> Body InKernel+             -> CallKernelGen [(VName, [VName])]+genRedKernel map_pes space ops body = do+  (base_constants, init_constants) <- kernelInitialisationSetSpace space $ return ()+  let constants = base_constants { kernelThreadActive = true }+      (space_is, space_sizes) = unzip $ spaceDimensions space+      i32_to_i64 = ConvOpExp (SExt Int32 Int64)+      space_sizes_64 = map (i32_to_i64 . ImpGen.compileSubExpOfType int32) space_sizes+      total_w_64 = product space_sizes_64++  histograms <- prepareIntermediateArrays (init space_sizes) (kernelNumThreads constants) ops++  elems_per_thread_64 <- dPrimV "elems_per_thread_64" $+                         total_w_64 `quotRoundingUp`+                         ConvOpExp (SExt Int32 Int64) (kernelNumThreads constants)++  sKernel constants "seggenred" $ allThreads constants $ do+    init_constants++    i <- newVName "i"++    -- Compute subhistogram index for each thread, per histogram.+    subhisto_inds <- forM histograms $ \(num_histograms, _, _) ->+      dPrimV "subhisto_ind" $+      kernelGlobalThreadId constants `quot`+      (kernelNumThreads constants `quotRoundingUp` Imp.var num_histograms int32)++    sFor i Int64 (Imp.var elems_per_thread_64 int64) $ do+      -- Compute the offset into the input and output.  To this a+      -- thread can add its local ID to figure out which element it is+      -- responsible for.  The calculation is done with 64-bit+      -- integers to avoid overflow, but the final segment indexes are+      -- 32 bit.+      offset <- dPrimV "offset" $+                (i32_to_i64 (kernelGroupId constants) *+                 (Imp.var elems_per_thread_64 int64 *+                  i32_to_i64 (kernelGroupSize constants)))+                + (Imp.var i int64 * i32_to_i64 (kernelGroupSize constants))++      j <- dPrimV "j" $ Imp.var offset int64 + i32_to_i64 (kernelLocalThreadId constants)++      -- Construct segment indices.+      let setIndex v e = do dPrim_ v int32+                            v <-- e+      zipWithM_ setIndex space_is $+        map (ConvOpExp (SExt Int64 Int32)) . unflattenIndex space_sizes_64 $ Imp.var j int64++      -- We execute the bucket function once and update each histogram serially.+      -- We apply the bucket function if j=offset+ltid is less than+      -- num_elements.  This also involves writing to the mapout+      -- arrays.+      let input_in_bounds = Imp.var j int32 .<. total_w_64++      sWhen input_in_bounds $ ImpGen.compileStms mempty (stmsToList $ bodyStms body) $ do+        let (red_res, map_res) = splitFromEnd (length map_pes) $ bodyResult body++        sComment "save map-out results" $+          forM_ (zip map_pes map_res) $ \(pe, se) ->+          ImpGen.copyDWIM (patElemName pe)+          (map ((`Imp.var` int32) . fst) $ kernelDimensions constants) se []++        let (buckets, vs) = splitAt (length ops) red_res+            perOp = chunks $ map (length . genReduceDest) ops++        sComment "perform atomic updates" $+          forM_ (zip5 ops histograms buckets (perOp vs) subhisto_inds) $+          \(GenReduceOp dest_w _ _ shape lam,+            (_, _, do_op), bucket, vs', subhisto_ind) -> do++            let bucket' = ImpGen.compileSubExpOfType int32 bucket+                dest_w' = ImpGen.compileSubExpOfType int32 dest_w+                bucket_in_bounds = 0 .<=. bucket' .&&. bucket' .<. dest_w'+                bucket_is = map (`Imp.var` int32) (init space_is) +++                            [Imp.var subhisto_ind int32, bucket']+                vs_params = takeLast (length vs') $ lambdaParams lam++            sWhen bucket_in_bounds $ do+              ImpGen.dLParams vs_params+              vectorLoops [] (shapeDims shape) $ \is -> do+                forM_ (zip vs_params vs') $ \(p, v) ->+                  ImpGen.copyDWIM (paramName p) [] v is+                do_op (bucket_is ++ is)++  let histogramInfo (num_histos, dests, _) = (num_histos, dests)+  return $ map histogramInfo histograms++vectorLoops :: [Imp.Exp] -> [SubExp]+            -> ([Imp.Exp] -> ImpGen.ImpM lore op ())+            -> ImpGen.ImpM lore op ()+vectorLoops is [] f = f $ reverse is+vectorLoops is (d:ds) f = do+  i <- newVName "i"+  d' <- ImpGen.compileSubExp d+  ImpGen.sFor i Int32 d' $ vectorLoops (Imp.var i int32:is) ds f++compileSegGenRed :: Pattern ExplicitMemory+                 -> KernelSpace+                 -> [GenReduceOp InKernel]+                 -> Body InKernel+                 -> CallKernelGen ()+compileSegGenRed (Pattern _ pes) genred_space ops body = do+  let num_red_res = length ops + sum (map (length . genReduceNeutral) ops)+      (all_red_pes, map_pes) = splitAt num_red_res pes++  infos <- genRedKernel map_pes genred_space ops body+  let pes_per_op = chunks (map (length . genReduceDest) ops) all_red_pes++  forM_ (zip3 infos pes_per_op ops) $ \((num_histos, subhistos), red_pes, op) -> do+    let unitHistoCase =+          -- This is OK because the memory blocks are at least as+          -- large as the ones we are supposed to use for the result.+          forM_ (zip red_pes subhistos) $ \(pe, subhisto) -> do+            pe_mem <- ImpGen.memLocationName . ImpGen.entryArrayLocation <$>+                      ImpGen.lookupArray (patElemName pe)+            subhisto_mem <- ImpGen.memLocationName . ImpGen.entryArrayLocation <$>+                            ImpGen.lookupArray subhisto+            ImpGen.emit $ Imp.SetMem pe_mem subhisto_mem $ Space "device"++    sIf (Imp.var num_histos int32 .==. 1) unitHistoCase $ do+      -- For the segmented reduction, we keep the segment dimensions+      -- unchanged.  To this, we add two dimensions: one over the number+      -- of buckets, and one over the number of subhistograms.  This+      -- inner dimension is the one that is collapsed in the reduction.+      let segment_dims = init $ spaceDimensions genred_space+          num_buckets = genReduceWidth op++      bucket_id <- newVName "bucket_id"+      subhistogram_id <- newVName "subhistogram_id"+      vector_ids <- mapM (const $ newVName "vector_id") $+                    shapeDims $ genReduceShape op+      gtid <- newVName $ baseString $ spaceGlobalId genred_space+      let lam = genReduceOp op+          segred_space =+            genred_space+            { spaceStructure =+                FlatThreadSpace $+                segment_dims +++                [(bucket_id, num_buckets)] +++                zip vector_ids (shapeDims $ genReduceShape op) +++                [(subhistogram_id, Var num_histos)]+            , spaceGlobalId = gtid+            }++      compileSegRed' (Pattern [] red_pes) segred_space+        Commutative lam (genReduceNeutral op) $ \red_dests _ ->+        forM_ (zip red_dests subhistos) $ \((d, is), subhisto) ->+          ImpGen.copyDWIM d is (Var subhisto) $ map (`Imp.var` int32) $+          map fst segment_dims ++ [subhistogram_id, bucket_id] ++ vector_ids

@@ -43,12 +43,12 @@ --   have any element to read, which becomes highly inefficient. module Futhark.CodeGen.ImpGen.Kernels.SegRed   ( compileSegRed+  , compileSegRed'   )   where  import Control.Monad.Except import Data.Maybe-import qualified Data.Set as S import Data.List  import Prelude hiding (quot, rem)@@ -83,14 +83,29 @@   i <- newVName "i"   sFor i Int32 iterations $ m $ group_id + Imp.var i int32 * kernelNumGroups constants --- Compile 'SegRed' instance to host-level code with calls to various--- kernels.+-- | Compile 'SegRed' instance to host-level code with calls to+-- various kernels. compileSegRed :: Pattern ExplicitMemory               -> KernelSpace               -> Commutativity -> Lambda InKernel -> [SubExp]               -> Body InKernel               -> CallKernelGen ()-compileSegRed pat space comm red_op nes body+compileSegRed pat space comm red_op nes body =+  compileSegRed' pat space comm red_op nes $ \red_dests map_dests ->+  ImpGen.compileStms mempty (stmsToList $ bodyStms body) $ do+  let (red_res, map_res) = splitAt (length nes) $ bodyResult body+  sComment "save results to be reduced" $+    forM_ (zip red_dests red_res) $ \((d,is), se) -> ImpGen.copyDWIM d is se []+  sComment "save map-out results" $+    forM_ (zip map_dests map_res) $ \((d,is), se) -> ImpGen.copyDWIM d is se []++-- | Like 'compileSegRed', but where the body is a monadic action.+compileSegRed' :: Pattern ExplicitMemory+               -> KernelSpace+               -> Commutativity -> Lambda InKernel -> [SubExp]+               -> ([(VName, [Imp.Exp])] -> [(VName, [Imp.Exp])] -> InKernelGen ())+               -> CallKernelGen ()+compileSegRed' pat space comm red_op nes body   | [(_, Constant (IntValue (Int32Value 1))), _] <- spaceDimensions space =       nonsegmentedReduction pat space comm red_op nes body   | otherwise = do@@ -98,52 +113,14 @@         ImpGen.compileSubExp $ last $ map snd $ spaceDimensions space       group_size <- ImpGen.compileSubExp $ spaceGroupSize space       let use_small_segments = segment_size * 2 .<. group_size-      sIf (segment_size .==. 1)-        (unitSegmentsReduction pat space nes body) $-        sIf use_small_segments+      sIf use_small_segments         (smallSegmentsReduction pat space red_op nes body)         (largeSegmentsReduction pat space comm red_op nes body) --- Handle degenerate case where segments are of size 1, meaning--- that it is really just a 'map' in disguise.-unitSegmentsReduction :: Pattern ExplicitMemory-                      -> KernelSpace-                      -> [SubExp]-                      -> Body InKernel-                      -> CallKernelGen ()-unitSegmentsReduction (Pattern _ segred_pes) space nes body = do-  (constants, init_constants) <- kernelInitialisationSetSpace space $ return ()--  let (gtids, dims) = unzip $ spaceDimensions space-      (redout_pes, mapout_pes) = splitAt (length nes) segred_pes--  dims' <- mapM ImpGen.compileSubExp dims--  let num_segments = product $ init dims'-      required_groups = num_segments `quotRoundingUp` kernelGroupSize constants--  ImpGen.emit $ Imp.DebugPrint "num_segments" int32 num_segments-  ImpGen.emit $ Imp.DebugPrint "required_groups" int32 required_groups--  sKernel constants "segred_mapseg" $ do-    init_constants-    virtualiseGroups constants required_groups $ \group_id -> do-      setSpaceIndices (group_id * kernelGroupSize constants + kernelLocalThreadId constants) space-      ImpGen.compileStms mempty (stmsToList $ bodyStms body) $-        sWhen (kernelThreadActive constants) $ do-        let (redout_ses, mapout_ses) = splitAt (length nes) $ bodyResult body-        forM_ (zip redout_pes redout_ses) $ \(pe, se) ->-          ImpGen.copyDWIM (patElemName pe)-          (map (`Imp.var` int32) (init gtids)) se []--        forM_ (zip mapout_pes mapout_ses) $ \(pe, se) ->-          ImpGen.copyDWIM (patElemName pe)-          (map (`Imp.var` int32) gtids) se []- nonsegmentedReduction :: Pattern ExplicitMemory                       -> KernelSpace                       -> Commutativity -> Lambda InKernel -> [SubExp]-                      -> Body InKernel+                      -> ([(VName, [Imp.Exp])] -> [(VName, [Imp.Exp])] -> InKernelGen ())                       -> CallKernelGen () nonsegmentedReduction segred_pat space comm red_op nes body = do   (base_constants, init_constants) <- kernelInitialisationSetSpace space $ return ()@@ -166,7 +143,7 @@    counter <-     ImpGen.sStaticArray "counter" (Space "device") int32 $-    replicate 1 $ IntValue $ Int32Value 0+    Imp.ArrayValues $ replicate 1 $ IntValue $ Int32Value 0    group_res_arrs <- forM (lambdaReturnType red_op) $ \t -> do     let pt = elemType t@@ -198,14 +175,10 @@       (kernelNumGroups constants) group_result_params red_acc_params red_op_renamed nes       1 counter sync_arr group_res_arrs red_arrs -hasMemoryAccesses :: Body InKernel -> ImpGen.ImpM InKernel Imp.KernelOp Bool-hasMemoryAccesses body = or <$> mapM isArray (S.toList $ freeInBody body)-  where isArray = fmap (not . primType) . lookupType- smallSegmentsReduction :: Pattern ExplicitMemory                        -> KernelSpace                        -> Lambda InKernel -> [SubExp]-                       -> Body InKernel+                       -> ([(VName, [Imp.Exp])] -> [(VName, [Imp.Exp])] -> InKernelGen ())                        -> CallKernelGen () smallSegmentsReduction (Pattern _ segred_pes) space red_op nes body = do   (base_constants, init_constants) <- kernelInitialisationSetSpace space $ return ()@@ -215,8 +188,12 @@   dims' <- mapM ImpGen.compileSubExp dims    let segment_size = last dims'+  -- Careful to avoid division by zero now.+  segment_size_nonzero_v <- dPrimV "segment_size_nonzero" $+                            BinOpExp (SMax Int32) 1 segment_size+  let segment_size_nonzero = Imp.var segment_size_nonzero_v int32       num_segments = product $ init dims'-      segments_per_group = kernelGroupSize constants `quot` segment_size+      segments_per_group = kernelGroupSize constants `quot` segment_size_nonzero       required_groups = num_segments `quotRoundingUp` segments_per_group    let red_op_params = lambdaParams red_op@@ -248,44 +225,37 @@       -- the segment ID, and are computed from the group id.  The inner       -- is computed from the local thread id, and may be out-of-bounds.       let ltid = kernelLocalThreadId constants-          segment_index = (ltid `quot` segment_size) + (group_id' * segments_per_group)+          segment_index = (ltid `quot` segment_size_nonzero) + (group_id' * segments_per_group)           index_within_segment = ltid `rem` segment_size        zipWithM_ (<--) (init gtids) $ unflattenIndex (init dims') segment_index       last gtids <-- index_within_segment        let toLocalMemory ses =-            forM_ (zip red_arrs ses) $ \(arr, se) -> do-            se_t <- subExpType se-            when (primType se_t) $-              ImpGen.copyDWIM arr [ltid] se []+            forM_ (zip red_arrs ses) $ \(arr, se) ->+            ImpGen.copyDWIM arr [ltid] se []            in_bounds =-            ImpGen.compileStms mempty (stmsToList $ bodyStms body) $ do-              let (red_res, map_res) = splitAt (length nes) $ bodyResult body--              sComment "save results to be reduced" $-                toLocalMemory red_res--              sComment "save map-out results" $-                forM_ (zip (drop (length nes) segred_pes) map_res) $ \(pe, se) ->-                ImpGen.copyDWIM (patElemName pe) (map (`Imp.var` int32) gtids) se []+            body (zip red_arrs $ repeat [ltid])+                 (zip (map patElemName $ drop (length nes) segred_pes) $+                  repeat $ map (`Imp.var` int32) gtids)        sComment "apply map function if in bounds" $-        sIf (isActive (init $ zip gtids dims) .&&.+        sIf (segment_size .>. 0 .&&.+             isActive (init $ zip gtids dims) .&&.              ltid .<. segment_size * segments_per_group) in_bounds (toLocalMemory nes)        sOp Imp.LocalBarrier        index_i <- newVName "index_i"       index_j <- newVName "index_j"-      let crossesSegment from to =-            (to-from) .>. (to `rem` segment_size)+      let crossesSegment from to = (to-from) .>. (to `rem` segment_size)           red_op' = red_op { lambdaParams = Param index_i (MemPrim int32) :                                             Param index_j (MemPrim int32) :                                             lambdaParams red_op } -      sComment "perform segmented scan to imitate reduction" $+      sWhen (segment_size .>. 0) $+        sComment "perform segmented scan to imitate reduction" $         groupScan constants (Just crossesSegment) (segment_size*segments_per_group) red_op' red_arrs        sOp Imp.LocalBarrier@@ -298,12 +268,16 @@         let flat_segment_index = group_id' * segments_per_group + ltid             gtids' = unflattenIndex (init dims') flat_segment_index         ImpGen.copyDWIM (patElemName pe) gtids'-                        (Var arr) [(ltid+1) * segment_size - 1]+                        (Var arr) [(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.+      sOp Imp.LocalBarrier+ largeSegmentsReduction :: Pattern ExplicitMemory                        -> KernelSpace                        -> Commutativity -> Lambda InKernel -> [SubExp]-                       -> Body InKernel+                       -> ([(VName, [Imp.Exp])] -> [(VName, [Imp.Exp])] -> InKernelGen ())                        -> CallKernelGen () largeSegmentsReduction segred_pat space comm red_op nes body = do   (base_constants, init_constants) <- kernelInitialisationSetSpace space $ return ()@@ -368,7 +342,7 @@   let num_counters = 1024   counter <-     ImpGen.sStaticArray "counter" (Space "device") int32 $-    replicate num_counters $ IntValue $ Int32Value 0+    Imp.ArrayZeros num_counters    sync_arr <- ImpGen.sAllocArray "sync_arr" Bool (Shape [intConst Int32 1]) $ Space "local" @@ -408,11 +382,12 @@      sIf (groups_per_segment .==. 1) one_group_per_segment multiple_groups_per_segment +-- Careful to avoid division by zero here. groupsPerSegmentAndElementsPerThread :: Imp.Exp -> Imp.Exp -> Imp.Exp -> Imp.Exp                                      -> (Imp.Exp, Imp.Count Imp.Elements) groupsPerSegmentAndElementsPerThread segment_size num_segments num_groups_hint group_size =   let groups_per_segment =-        num_groups_hint `quotRoundingUp` num_segments+        num_groups_hint `quotRoundingUp` BinOpExp (SMax Int32) 1 num_segments       elements_per_thread =         segment_size `quotRoundingUp` (group_size * groups_per_segment)   in (groups_per_segment, Imp.elements elements_per_thread)@@ -427,7 +402,7 @@                   -> LambdaT InKernel                   -> [SubExp]                   -> [VName]-                  -> Body InKernel+                  -> ([(VName, [Imp.Exp])] -> [(VName, [Imp.Exp])] -> InKernelGen ())                   -> InKernelGen ([LParam InKernel], Lambda InKernel) reductionStageOne constants (Pattern _ segred_pes) num_elements global_tid elems_per_thread threads_per_segment comm red_op nes red_arrs body = do @@ -436,13 +411,11 @@       (gtids, _dims) = unzip $ kernelDimensions constants       gtid = last gtids       local_tid = kernelLocalThreadId constants-      index_in_segment = global_tid `quot` kernelGroupSize constants    -- Figure out how many elements this thread should process.   chunk_size <- dPrim "chunk_size" int32   let ordering = case comm of Commutative -> SplitStrided $ Var threads_per_segment                               Noncommutative -> SplitContiguous-  accesses_memory <- hasMemoryAccesses body   computeThreadChunkSize ordering global_tid elems_per_thread num_elements chunk_size    ImpGen.dScope Nothing $ scopeOfLParams $ lambdaParams red_op@@ -462,6 +435,8 @@          groupReduce constants (kernelGroupSize constants) red_op_renamed red_arrs +        sOp Imp.LocalBarrier+   i <- newVName "i"   -- If this is a non-commutative reduction, each thread must run the   -- loop the same number of iterations, because we will be performing@@ -478,31 +453,24 @@         Commutative ->           global_tid +           Imp.var threads_per_segment int32 * Imp.var i int32-        Noncommutative | accesses_memory ->-          local_tid +-          (index_in_segment * Imp.innerExp elems_per_thread + Imp.var i int32) *-          kernelGroupSize constants         Noncommutative ->-          Imp.var i int32 +-          global_tid * Imp.innerExp elems_per_thread--    check_bounds $ sComment "apply map function" $-      ImpGen.compileStms mempty (stmsToList $ bodyStms body) $ do-        let (red_res, map_res) = splitAt (length nes) $ bodyResult body+          let index_in_segment = global_tid `quot` kernelGroupSize constants+          in local_tid ++             (index_in_segment * Imp.innerExp elems_per_thread + Imp.var i int32) *+             kernelGroupSize constants -        sComment "save results to be reduced" $-          forM_ (zip red_next_params red_res) $ \(p, se) ->-          ImpGen.copyDWIM (paramName p) [] se []+    let red_dests = zip (map paramName red_next_params) $ repeat []+        map_dests = zip (map patElemName $ drop (length nes) segred_pes) $+                    repeat $ map (`Imp.var` int32) gtids -        sComment "save map-out results" $-          forM_ (zip (drop (length nes) segred_pes) map_res) $ \(pe, se) ->-          ImpGen.copyDWIM (patElemName pe) (map (`Imp.var` int32) gtids) se []+    check_bounds $ sComment "apply map function" $ do+      body red_dests map_dests -        sComment "apply reduction operator" $-          ImpGen.compileBody' red_acc_params $ lambdaBody red_op+      sComment "apply reduction operator" $+        ImpGen.compileBody' red_acc_params $ lambdaBody red_op      case comm of-      Noncommutative | accesses_memory -> do+      Noncommutative -> do         doTheReduction         sComment "first thread takes carry-out; others neutral element" $ do           let carry_out =@@ -512,17 +480,12 @@                 forM_ (zip red_acc_params nes) $ \(p, ne) ->                 ImpGen.copyDWIM (paramName p) [] ne []           sIf (local_tid .==. 0) carry_out reset_to_neutral-      _ ->-        return ()--  group_result_params <- case comm of-    Noncommutative | accesses_memory ->-      return red_acc_params--    _ -> do-      doTheReduction+      _ -> return () -      return $ lambdaParams red_op_renamed+  group_result_params <-+    case comm of Noncommutative -> return red_acc_params+                 _              -> do doTheReduction+                                      return $ lambdaParams red_op_renamed    return (group_result_params, red_op_renamed) @@ -590,6 +553,8 @@           load_group_result load_neutral_element         when (primType $ paramType p) $           ImpGen.copyDWIM arr [local_tid] (Var $ paramName p) []++    sOp Imp.LocalBarrier      sComment "reduce the per-group results" $ do       groupReduce constants group_size red_op_renamed red_arrs

@@ -29,7 +29,7 @@ import qualified Futhark.CodeGen.ImpCode.OpenCL as ImpOpenCL import Futhark.MonadFreshNames import Futhark.Util (zEncodeString)-import Futhark.Util.Pretty (pretty, prettyOneLine)+import Futhark.Util.Pretty (pretty)  kernelsToCUDA, kernelsToOpenCL :: ImpKernels.Program                                -> Either InternalError ImpOpenCL.Program@@ -65,17 +65,15 @@  type UsedFunctions = [(String,C.Func)] -- The ordering is important! -data OpenClRequirements =-  OpenClRequirements { kernelUsedTypes :: S.Set PrimType-                     , _kernelConstants :: [(VName, KernelConstExp)]-                     }+newtype OpenClRequirements =+  OpenClRequirements { kernelUsedTypes :: S.Set PrimType }  instance Semigroup OpenClRequirements where-  OpenClRequirements ts1 consts1 <> OpenClRequirements ts2 consts2 =-    OpenClRequirements (ts1 <> ts2) (consts1 <> consts2)+  OpenClRequirements ts1 <> OpenClRequirements ts2 =+    OpenClRequirements (ts1 <> ts2)  instance Monoid OpenClRequirements where-  mempty = OpenClRequirements mempty mempty+  mempty = OpenClRequirements mempty  data ToOpenCL = ToOpenCL { clExtraFuns :: M.Map Name ImpOpenCL.Function                          , clKernels :: M.Map KernelName C.Func@@ -140,16 +138,17 @@        params = perm_params ++ catMaybes local_memory_params ++ use_params +      const_defs = mapMaybe constDef $ kernelUses kernel+   tell mempty { clExtraFuns = mempty               , clKernels = M.singleton name                             [C.cfun|__kernel void $id:name ($params:params) {+                                $items:const_defs                                 $items:block_dim_init                                 $items:local_memory_init                                 $items:kernel_body                                 }|]-              , clRequirements = OpenClRequirements-                                 (typesInKernel kernel)-                                 (mapMaybe useAsConst $ kernelUses kernel)+              , clRequirements = OpenClRequirements (typesInKernel kernel)               }    return $ LaunchKernel name (kernelArgs kernel) num_groups group_size@@ -186,9 +185,11 @@ useAsParam ConstUse{} =   Nothing -useAsConst :: KernelUse -> Maybe (VName, KernelConstExp)-useAsConst (ConstUse v e) = Just (v,e)-useAsConst _ = Nothing+constDef :: KernelUse -> Maybe C.BlockItem+constDef (ConstUse v e) = Just [C.citem|const $ty:t $id:v = $exp:e';|]+  where t = GenericC.primTypeToCType $ primExpType e+        e' = compilePrimExp e+constDef _ = Nothing  openClCode :: [C.Func] -> String openClCode kernels =@@ -198,9 +199,13 @@            kernel_func <- kernels ]  genOpenClPrelude :: OpenClRequirements -> [C.Definition]-genOpenClPrelude (OpenClRequirements ts consts) =+genOpenClPrelude (OpenClRequirements ts) =   -- Clang-based OpenCL implementations need this for 'static' to work.-  [C.cedecl|$esc:("#pragma OPENCL EXTENSION cl_clang_storage_class_specifiers : enable")|] :+  [ [C.cedecl|$esc:("#ifdef cl_clang_storage_class_specifiers")|]+  , [C.cedecl|$esc:("#pragma OPENCL EXTENSION cl_clang_storage_class_specifiers : enable")|]+  , [C.cedecl|$esc:("#endif")|]+  , [C.cedecl|$esc:("#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable")|]]+  ++   [[C.cedecl|$esc:("#pragma OPENCL EXTENSION cl_khr_fp64 : enable")|] | uses_float64] ++   [C.cunit| /* Some OpenCL programs dislike empty progams, or programs with no kernels.@@ -222,10 +227,24 @@ typedef ulong uint64_t;  $esc:("#define ALIGNED_LOCAL_MEMORY(m,size) __local unsigned char m[size] __attribute__ ((align))")++// NVIDIAs OpenCL does not create device-wide memory fences (see #734), so we+// use inline assembly if we detect we are on an NVIDIA GPU.+$esc:("#ifdef cl_nv_pragma_unroll")+static inline void mem_fence_global() {+  asm("membar.gl;");+}+$esc:("#else")+static inline void mem_fence_global() {+  mem_fence(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);+}+$esc:("#endif")+static inline void mem_fence_local() {+  mem_fence(CLK_LOCAL_MEM_FENCE);+} |] ++   cIntOps ++ cFloat32Ops ++ cFloat32Funs ++-  (if uses_float64 then cFloat64Ops ++ cFloat64Funs ++ cFloatConvOps else []) ++-  [ [C.cedecl|$esc:def|] | def <- map constToDefine consts ]+  (if uses_float64 then cFloat64Ops ++ cFloat64Funs ++ cFloatConvOps else [])   where uses_float64 = FloatType Float64 `S.member` ts  @@ -254,11 +273,10 @@                 }|] | t <- types]  genCUDAPrelude :: OpenClRequirements -> [C.Definition]-genCUDAPrelude (OpenClRequirements _ consts) =-  cudafy ++ cudaAtomicOps ++ defs ++ ops+genCUDAPrelude (OpenClRequirements _) =+  cudafy ++ cudaAtomicOps ++ ops   where ops = cIntOps ++ cFloat32Ops ++ cFloat32Funs ++ cFloat64Ops                 ++ cFloat64Funs ++ cFloatConvOps-        defs = [ [C.cedecl|$esc:def|] | def <- map constToDefine consts ]         cudafy = [CUDAC.cunit| typedef char int8_t; typedef short int16_t;@@ -349,25 +367,17 @@ {   __syncthreads(); }-static inline void mem_fence(int x)-{-  if (x == CLK_LOCAL_MEM_FENCE) {-    __threadfence_block();-  } else {-    __threadfence();-  }+static inline void mem_fence_local() {+  __threadfence_block(); }+static inline void mem_fence_global() {+  __threadfence();+} $esc:("#define NAN (0.0/0.0)") $esc:("#define INFINITY (1.0/0.0)") extern volatile __shared__ char shared_mem[]; |] -constToDefine :: (VName, KernelConstExp) -> String-constToDefine (name, e) =-  let e' = compilePrimExp e-  in unwords ["#define", zEncodeString (pretty name), "("++prettyOneLine e'++")"]-- compilePrimExp :: PrimExp KernelConst -> C.Exp compilePrimExp e = runIdentity $ GenericC.compilePrimExp compileKernelConst e   where compileKernelConst (SizeConst key) =@@ -413,7 +423,7 @@         kernelOps GlobalBarrier =           GenericC.stm [C.cstm|barrier(CLK_GLOBAL_MEM_FENCE);|]         kernelOps MemFence =-          GenericC.stm [C.cstm|mem_fence(CLK_GLOBAL_MEM_FENCE);|]+          GenericC.stm [C.cstm|mem_fence_global();|]         kernelOps (Atomic aop) = atomicOps aop          atomicOps (AtomicAdd old arr ind val) = do

@@ -28,7 +28,7 @@                     | HeuristicDeviceInfo String  -- | A size that can be assigned a default.-data WhichSize = LockstepWidth | NumGroups | GroupSize | TileSize+data WhichSize = LockstepWidth | NumGroups | GroupSize | TileSize | Threshold  -- | A heuristic for setting the default value for something. data SizeHeuristic =@@ -47,9 +47,11 @@   , SizeHeuristic "" DeviceGPU NumGroups $ HeuristicConst 256   , SizeHeuristic "" DeviceGPU GroupSize $ HeuristicConst 256   , SizeHeuristic "" DeviceGPU TileSize $ HeuristicConst 32+  , SizeHeuristic "" DeviceGPU Threshold $ HeuristicConst $ 32*1024    , SizeHeuristic "" DeviceCPU LockstepWidth $ HeuristicConst 1   , SizeHeuristic "" DeviceCPU NumGroups $ HeuristicDeviceInfo "MAX_COMPUTE_UNITS"   , SizeHeuristic "" DeviceCPU GroupSize $ HeuristicConst 32   , SizeHeuristic "" DeviceCPU TileSize $ HeuristicConst 4+  , SizeHeuristic "" DeviceCPU Threshold $ HeuristicDeviceInfo "MAX_COMPUTE_UNITS"   ]

@@ -30,6 +30,7 @@   , eSliceArray   , eSplitArray +  , eOutOfBounds   , eWriteArray    , asIntZ, asIntS@@ -285,16 +286,13 @@           offset' <- letSubExp "offset" $ BasicOp $ BinOp (Add Int32) offset size           return (offset', offset) --- | Write to an index of the array, if within bounds.  Otherwise,--- nothing.  Produces the updated array.-eWriteArray :: (MonadBinder m, BranchType (Lore m) ~ ExtType) =>-               VName -> [m (Exp (Lore m))] -> m (Exp (Lore m))-            -> m (Exp (Lore m))-eWriteArray arr is v = do+-- | Are these indexes out-of-bounds for the array?+eOutOfBounds :: MonadBinder m =>+                VName -> [m (Exp (Lore m))] -> m (Exp (Lore m))+eOutOfBounds arr is = do   arr_t <- lookupType arr   let ws = arrayDims arr_t   is' <- mapM (letSubExp "write_i") =<< sequence is-  v' <- letSubExp "write_v" =<< v   let checkDim w i = do         less_than_zero <- letSubExp "less_than_zero" $           BasicOp $ CmpOp (CmpSlt Int32) i (constant (0::Int32))@@ -302,11 +300,19 @@           BasicOp $ CmpOp (CmpSle Int32) w i         letSubExp "outside_bounds_dim" $           BasicOp $ BinOp LogOr less_than_zero greater_than_size+  foldBinOp LogOr (constant False) =<< zipWithM checkDim ws is' -  outside_bounds <--    letSubExp "outside_bounds" =<<-    foldBinOp LogOr (constant False) =<<-    zipWithM checkDim ws is'+-- | Write to an index of the array, if within bounds.  Otherwise,+-- nothing.  Produces the updated array.+eWriteArray :: (MonadBinder m, BranchType (Lore m) ~ ExtType) =>+               VName -> [m (Exp (Lore m))] -> m (Exp (Lore m))+            -> m (Exp (Lore m))+eWriteArray arr is v = do+  arr_t <- lookupType arr+  is' <- mapM (letSubExp "write_i") =<< sequence is+  v' <- letSubExp "write_v" =<< v++  outside_bounds <- letSubExp "outside_bounds" =<< eOutOfBounds arr is    outside_bounds_branch <- insertStmsM $ resultBodyM [Var arr] 

@@ -715,10 +715,9 @@ describeDec _ ImportDec{} = Nothing  valBindWhat :: ValBind -> IndexWhat-valBindWhat vb =-  if null (valBindParams vb) && orderZero (unInfo (valBindRetType vb))-  then IndexValue-  else IndexFunction+valBindWhat vb | null (valBindParams vb),+                 orderZero (unInfo (valBindRetType vb)) = IndexValue+               | otherwise                              = IndexFunction  describeSpecs :: [Spec] -> DocM Html describeSpecs specs =

@@ -459,11 +459,11 @@      mergeinit_ts' <- mapM subExpType mergeinit' -    let ctxinit = argShapes-                  (map I.paramName shapepat)-                  (map I.paramType mergepat')-                  mergeinit_ts'-        ctxmerge = zip shapepat ctxinit+    ctxinit <- argShapes+               (map I.paramName shapepat)+               (map I.paramType mergepat')+               mergeinit_ts'+    let ctxmerge = zip shapepat ctxinit         valmerge = zip mergepat' mergeinit'         merge = ctxmerge ++ valmerge         dropCond = case form of E.While{} -> drop 1@@ -487,10 +487,10 @@       inScopeOf form' $ internaliseBodyStms loopbody $ \ses -> do       sets <- mapM subExpType ses       let mergepat' = concat nested_mergepat-          shapeargs = argShapes-                      (map I.paramName shapepat)-                      (map I.paramType mergepat')-                      sets+      shapeargs <- argShapes+                   (map I.paramName shapepat)+                   (map I.paramType mergepat')+                   sets       return (resultBody $ shapeargs ++ ses,               (form',                shapepat,@@ -536,10 +536,10 @@         -- next iteration?).  This is safe, as the type rules for         -- the external language guarantees that 'cond' does not         -- consume anything.-        let shapeinit = argShapes-                        (map I.paramName shapepat)-                        (map I.paramType mergepat')-                        mergeinit_ts+        shapeinit <- argShapes+                     (map I.paramName shapepat)+                     (map I.paramType mergepat')+                     mergeinit_ts          (loop_initial_cond, init_loop_cond_bnds) <- collectStms $ do           forM_ (zip shapepat shapeinit) $ \(p, se) ->@@ -555,10 +555,10 @@         internaliseBodyStms loopbody $ \ses -> do           sets <- mapM subExpType ses           loop_while <- newParam "loop_while" $ I.Prim I.Bool-          let shapeargs = argShapes-                          (map I.paramName shapepat)-                          (map I.paramType mergepat')-                          sets+          shapeargs <- argShapes+                       (map I.paramName shapepat)+                       (map I.paramType mergepat')+                       sets            -- Careful not to clobber anything.           loop_end_cond_body <- renameBody <=< insertStmsM $ do@@ -618,9 +618,6 @@           return $ bef ++ src'' ++ aft         replace _ _ ve' _ = return ve' -internaliseExp desc (E.Unzip e _ _) =-  internaliseExp desc e- internaliseExp desc (E.Unsafe e _) =   local (\env -> env { envDoBoundsChecks = False }) $   internaliseExp desc e@@ -636,33 +633,6 @@           letBindNames_ [v'] $ I.BasicOp $ I.SubExp v           return $ I.Var v' -internaliseExp _ (E.Zip _ e es _ loc) = do-  e' <- internaliseExpToVars "zip_arg" $ TupLit (e:es) loc-  case e' of-    e_key:es_unchecked -> do-      -- We will reshape all of es_unchecked' to have the same outer-      -- size as ts.  We will not change any of the inner dimensions.-      -- This will cause a runtime error if the outer sizes do not match,-      -- thus preserving the semantics of zip().-      w <- arraySize 0 <$> lookupType e_key-      let reshapeToOuter e_unchecked' = do-            unchecked_t <- lookupType e_unchecked'-            case I.arrayDims unchecked_t of-              outer:inner | w /= outer -> do-                cmp <- letSubExp "zip_cmp" $ I.BasicOp $-                       I.CmpOp (I.CmpEq I.int32) w outer-                c   <- assertingOne $-                       letExp "zip_assert" $ I.BasicOp $-                       I.Assert cmp "arrays differ in length" (loc, mempty)-                certifying c $ letExp (postfix e_unchecked' "_zip_res") $-                  shapeCoerce (w:inner) e_unchecked'-              _ -> return e_unchecked'-      es' <- mapM reshapeToOuter es_unchecked-      return $ map I.Var $ e_key : es'-    [] -> return []--  where postfix i s = baseString i ++ s- internaliseExp desc (E.Map lam arr _ _) = do   arr' <- internaliseExpToVars "map_arr" arr   lam' <- internaliseMapLambda internaliseLambda lam $ map I.Var arr'@@ -1420,48 +1390,9 @@       ressize <- foldM sumdims outer_size =<<                  mapM (fmap (arraysSize 0) . mapM lookupType) [ys] -      let conc xarr yarr = do-            -- All dimensions except the outermost must match.  An-            -- empty array matches anything.-            xt <- lookupType xarr-            yt <- lookupType yarr-            let matches n m =-                  letSubExp "match" $-                  I.BasicOp $ I.CmpOp (I.CmpEq I.int32) n m--                emptyRow arr_t =-                  letSubExp "empty_row" =<<-                  foldBinOp I.LogOr (constant False) =<<-                  mapM (matches (intConst Int32 0)) (arrayDims $ rowType arr_t)--            all_match <- letSubExp "all_match" =<<-                         foldBinOp I.LogAnd (constant True) =<<-                         zipWithM matches-                         (arrayDims (rowType xt)) (arrayDims (rowType yt))-            xarr_empty <- emptyRow xt-            yarr_empty <- emptyRow yt-            either_empty <- letSubExp "either_empty" $-                            I.BasicOp $ I.BinOp I.LogOr xarr_empty yarr_empty-            matchcs <- assertingOne $ letExp "concat_ok" =<<-                       eAssert (pure $ I.BasicOp $ I.BinOp I.LogOr either_empty all_match)-                       "row sizes do not match when concatenating" loc--            let updims (j, xd, yd)-                  | j == 0    =-                      return (xd, yd)-                  | otherwise = do-                      d <- letSubExp "dim" $ I.BasicOp $ I.BinOp (SMax Int32) xd yd-                      return (d, d)--            (xdims, ydims) <- unzip <$>-              mapM updims (zip3 [(0::Int)..] (I.arrayDims xt) (I.arrayDims yt))--            xarr' <- certifying matchcs $ letExp "concat_x_reshaped" $-                     shapeCoerce xdims xarr-            yarr' <- certifying matchcs $ letExp "concat_y_reshaped" $-                     shapeCoerce ydims yarr-            return $ I.BasicOp $ I.Concat 0 xarr' [yarr'] ressize-      letSubExps desc =<< zipWithM conc xs ys+      let conc xarr yarr =+            I.BasicOp $ I.Concat 0 xarr [yarr] ressize+      letSubExps desc $ zipWith conc xs ys      handle [TupLit [offset, e] _] "rotate" = Just $ \desc -> do       offset' <- internaliseExp1 "rotation_offset" offset@@ -1635,8 +1566,8 @@   (constargs, const_ds, _) <- unzip3 <$> constFunctionArgs loc constparams   argts <- mapM subExpType args   closure_ts <- mapM lookupType closure-  let shapeargs = argShapes shapes value_paramts argts-      diets = const_ds ++ replicate (length closure + length shapeargs) I.Observe +++  shapeargs <- argShapes shapes value_paramts argts+  let diets = const_ds ++ replicate (length closure + length shapeargs) I.Observe ++               map I.diet value_paramts       constOrShape = const $ I.Prim int32       paramts = map constOrShape constargs ++ closure_ts ++

@@ -15,6 +15,7 @@ import Control.Monad import Data.Loc import qualified Data.Map.Strict as M+import qualified Data.Set as S  import Futhark.Construct import Futhark.Representation.AST@@ -26,7 +27,7 @@   runBodyBinder $ do     ses <- bodyBind body     sets <- mapM subExpType ses-    return $ resultBody $ argShapes shapenames ts sets+    resultBody <$> argShapes shapenames ts sets  annotateArrayShape :: ArrayShape shape =>                       TypeBase shape u -> [Int] -> TypeBase Shape u@@ -34,13 +35,25 @@   t `setArrayShape` Shape (take (arrayRank t) $                            map (intConst Int32 . toInteger) $ newshape ++ repeat 0) -argShapes :: [VName] -> [TypeBase Shape u0] -> [TypeBase Shape u1] -> [SubExp]+-- Some trickery is needed here to predict sensible values for+-- dimensions that are used exclusively as the inner dimension of an+-- array.  The issue is that the dimension may be inside an empty+-- array.  In this case, the dimension inside the empty array should+-- not count, as it will be zero.  The solution we use is to take the+-- maximum of such sizes; this will effectively disregard the zeroes.+argShapes :: MonadBinder m =>+             [VName] -> [TypeBase Shape u0] -> [TypeBase Shape u1] -> m [SubExp] argShapes shapes valts valargts =-  map addShape shapes+  mapM addShape shapes   where mapping = shapeMapping valts valargts-        addShape name-          | Just se <- M.lookup name mapping = se-          | otherwise                        = intConst Int32 0+        outer_dims = map (arraySize 0) valts+        addShape name =+          case M.lookup name mapping of+            Just s | x:xs <- S.toList s ->+                       if Var name `elem` outer_dims+                       then return x+                       else letSubExp "size" =<< foldBinOp (SMax Int32) x xs+            _ -> return $ intConst Int32 0  ensureResultShape :: MonadBinder m =>                      (m Certificates -> m Certificates)

@@ -363,15 +363,6 @@   arr' <- defuncExp' arr   return (Stream form' lam' arr' loc, Dynamic $ typeOf e) -defuncExp e@(Zip i e1 es t loc) = do-  e1' <- defuncExp' e1-  es' <- mapM defuncExp' es-  return (Zip i e1' es' t loc, Dynamic $ typeOf e)--defuncExp e@(Unzip e0 tps loc) = do-  e0' <- defuncExp' e0-  return (Unzip e0' tps loc, Dynamic $ typeOf e)- defuncExp (Unsafe e1 loc) = do   (e1', sv) <- defuncExp e1   return (Unsafe e1' loc, sv)@@ -497,67 +488,42 @@       let closure_pat = buildEnvPattern closure_env           pat' = updatePattern pat sv2 -      -- Inline certain trivial lifted functions immediately.  This is-      -- purely an optimisation to avoid having the rest of the-      -- compiler spend a lot of time processing them (they will end-      -- up being inlined later anyway).  We also try to simplify away-      -- some let-bindings, to make the generated code look slightly-      -- more comprehensible.-      ---      -- If you are debugging the defunctionaliser, you may want to-      -- turn this off to see the original structure of the generated-      -- code instead.-      let letPat (RecordPattern [] _) _ pbody = pbody-          letPat (PatternParens pp _) pe pbody = letPat pp pe pbody-          letPat (Id v1 _ _) pe (RecordLit [RecordFieldExplicit f (Var v2 _ _) floc] rloc)-            | v1 == qualLeaf v2 =-                RecordLit [RecordFieldExplicit f pe floc] rloc-          letPat (RecordPattern [(pf,p)] _) (RecordLit [RecordFieldExplicit f pe _] _) pbody-            | pf == f =-                letPat p pe pbody-          letPat pp pe pbody = LetPat [] pp pe pbody noLoc--          inline RecordLit{} = True-          inline TupLit{} = True-          inline (Apply x y _ _ _) = inline x && inline y-          inline (BinOp _ _ (x, _) (y, _) _ _) = inline x && inline y-          inline Var{} = True-          inline Literal{} = True-          inline (LetPat _ _ x y _) = inline x && inline y-          inline Negate{} = True-          inline _ = False-      if inline e0' && null dims-        then return (letPat closure_pat e1' $ letPat pat' e2' e0', sv)-        else do-          -- Lift lambda to top-level function definition.-          let params = [closure_pat, pat']-              rettype = buildRetType closure_env params e0_t $ typeOf e0'+      -- Lift lambda to top-level function definition.  We put in+      -- a lot of effort to try to infer the uniqueness attributes+      -- of the lifted function, but this is ultimately all a sham+      -- and a hack.  There is some piece we're missing.+      let params = [closure_pat, pat']+          params_for_rettype = params ++ svParams sv1 ++ svParams sv2+          svParams (LambdaSV _ sv_pat _ _ _) = [sv_pat]+          svParams _                         = []+          rettype = buildRetType closure_env params_for_rettype e0_t $ typeOf e0' -              -- Embed some information about the original function-              -- into the name of the lifted function, to make the-              -- result slightly more human-readable.-              liftedName i (Var f _ _) =-                "lifted_" ++ show i ++ "_" ++ baseString (qualLeaf f)-              liftedName i (Apply f _ _ _ _) =-                liftedName (i+1) f-              liftedName _ _ = "lifted"-          fname <- newNameFromString $ liftedName (0::Int) e1-          liftValDec fname rettype dims params e0'+          -- Embed some information about the original function+          -- into the name of the lifted function, to make the+          -- result slightly more human-readable.+          liftedName i (Var f _ _) =+            "lifted_" ++ show i ++ "_" ++ baseString (qualLeaf f)+          liftedName i (Apply f _ _ _ _) =+            liftedName (i+1) f+          liftedName _ _ = "lifted"+      fname <- newNameFromString $ liftedName (0::Int) e1+      liftValDec fname rettype dims params e0' -          let t1 = vacuousShapeAnnotations . toStruct $ typeOf e1'-              t2 = vacuousShapeAnnotations . toStruct $ typeOf e2'-              fname' = qualName fname-          return (Parens (Apply (Apply (Var fname' (Info (Arrow mempty Nothing (fromStruct t1) $-                                                          Arrow mempty Nothing (fromStruct t2) rettype)) loc)-                                 e1' (Info Observe) (Info $ Arrow mempty Nothing (fromStruct t2) rettype) loc)-                          e2' d (Info rettype) loc) noLoc, sv)+      let t1 = vacuousShapeAnnotations . toStruct $ typeOf e1'+          t2 = vacuousShapeAnnotations . toStruct $ typeOf e2'+          fname' = qualName fname+      return (Parens (Apply (Apply (Var fname' (Info (Arrow mempty Nothing (fromStruct t1) $+                                                      Arrow mempty Nothing (fromStruct t2) rettype)) loc)+                             e1' (Info Observe) (Info $ Arrow mempty Nothing (fromStruct t2) rettype) loc)+                      e2' d (Info rettype) loc) noLoc, sv)      -- If e1 is a dynamic function, we just leave the application in place,     -- but we update the types since it may be partially applied or return     -- a higher-order term.     DynamicFun _ sv ->       let (argtypes', rettype) = dynamicFunType sv argtypes-          apply_e = Apply e1' e2' d (Info $ foldFunType argtypes' rettype) loc+          apply_e = Apply e1' e2' d (Info $ foldFunType argtypes' rettype+                                    `setAliases` aliases ret) loc       in return (apply_e, sv)      -- Propagate the 'IntrinsicsSV' until we reach the outermost application,@@ -681,7 +647,7 @@         comb (Record fs_annot) (Record fs_got) =           Record $ M.intersectionWith comb fs_annot fs_got         comb Arrow{} t = vacuousShapeAnnotations $ descend t-        comb got _ = fromStruct got+        comb got et = descend $ fromStruct got `setUniqueness` uniqueness et `setAliases` aliases et          descend t@Array{}           | any (problematic . aliasVar) (aliases t) = t `setUniqueness` Nonunique@@ -865,8 +831,6 @@     where freeInForm (RedLike _ _ e) = freeVars e           freeInForm _ = mempty -  Zip _ e es _ _      -> freeVars e <> foldMap freeVars es-  Unzip e _ _         -> freeVars e   Unsafe e _          -> freeVars e   Assert e1 e2 _ _    -> freeVars e1 <> freeVars e2   VConstr0{}          -> mempty

@@ -1,3 +1,4 @@+{-# LANGUAGE FlexibleContexts #-} -- | This monomorphization module converts a well-typed, polymorphic, -- module-free Futhark program into an equivalent monomorphic program. --@@ -29,15 +30,19 @@  import           Control.Monad.RWS import           Control.Monad.State+import           Control.Monad.Writer+import           Data.Bitraversable+import           Data.Bifunctor import           Data.Loc import qualified Data.Map.Strict as M+import qualified Data.Set as S import qualified Data.Sequence as Seq import           Data.Foldable  import           Futhark.MonadFreshNames import           Language.Futhark import           Language.Futhark.Traversals-import           Language.Futhark.TypeChecker.Monad (TypeBinding(..))+import           Language.Futhark.Semantic (TypeBinding(..)) import           Language.Futhark.TypeChecker.Types  -- | The monomorphization monad reads 'PolyBinding's and writes 'ValBinding's.@@ -135,11 +140,24 @@         (Nothing, Nothing) -> return fname         -- A polymorphic function.         (Nothing, Just funbind) -> do-          (fname', funbind') <- monomorphizeBinding funbind t+          (fname', funbind') <- monomorphizeBinding False funbind t           tell $ Seq.singleton (fname, funbind')           addLifted fname t fname'           return fname' +-- | This carries out record replacements in the alias information of a type.+transformType :: TypeBase dim Aliasing -> MonoM (TypeBase dim Aliasing)+transformType t = do+  rrs <- asks envRecordReplacements+  let replace (AliasBound v) | Just d <- M.lookup v rrs =+                                 S.fromList $ map (AliasBound . fst) $ M.elems d+      replace x = S.singleton x+  -- As an attempt at an optimisation, only transform the aliases if+  -- they refer to a variable we have record-replaced.+  return $ if any ((`M.member` rrs) . aliasVar) $ aliases t+           then bimap id (mconcat . map replace . S.toList) t+           else t+ -- | Monomorphization of expressions. transformExp :: Exp -> MonoM Exp transformExp e@Literal{} = return e@@ -159,9 +177,10 @@   RecordLit <$> mapM transformField fs <*> pure loc   where transformField (RecordFieldExplicit name e loc') =           RecordFieldExplicit name <$> transformExp e <*> pure loc'-        transformField (RecordFieldImplicit v t _) =+        transformField (RecordFieldImplicit v t _) = do+          t' <- traverse transformType t           transformField $ RecordFieldExplicit (baseName v)-          (Var (qualName v) (vacuousShapeAnnotations <$> t) loc) loc+            (Var (qualName v) (vacuousShapeAnnotations <$> t') loc) loc  transformExp (ArrayLit es tp loc) =   ArrayLit <$> mapM transformExp es <*> pure tp <*> pure loc@@ -176,13 +195,15 @@   maybe_fs <- lookupRecordReplacement fname   case maybe_fs of     Just fs -> do-      let toField (f, (f_v, f_t)) =-            let f_v' = Var (qualName f_v) (Info $ vacuousShapeAnnotations f_t) loc-            in RecordFieldExplicit f f_v' loc-      return $ RecordLit (map toField $ M.toList fs) loc+      let toField (f, (f_v, f_t)) = do+            f_t' <- transformType f_t+            let f_v' = Var (qualName f_v) (Info $ vacuousShapeAnnotations f_t') loc+            return $ RecordFieldExplicit f f_v' loc+      RecordLit <$> mapM toField (M.toList fs) <*> pure loc     Nothing -> do       fname' <- transformFName fname (toStructural t)-      return $ Var (QualName qs fname') (Info t) loc+      t' <- transformType t+      return $ Var (QualName qs fname') (Info t') loc  transformExp (Ascript e tp loc) =   Ascript <$> transformExp e <*> pure tp <*> pure loc@@ -213,7 +234,8 @@   e1' <- transformExp e1   e2' <- transformExp e2   e3' <- transformExp e3-  return $ If e1' e2' e3' tp loc+  tp' <- traverse transformType tp+  return $ If e1' e2' e3' tp' loc  transformExp (Apply e1 e2 d tp loc) =   -- We handle on an ad-hoc basis certain polymorphic higher-order@@ -254,7 +276,8 @@     _ -> do       e1' <- transformExp e1       e2' <- transformExp e2-      return $ Apply e1' e2' d tp loc+      tp' <- traverse transformType tp+      return $ Apply e1' e2' d tp' loc   where intrinsic s (QualName _ v) =           baseTag v <= maxIntrinsicTag && baseName v == nameFromString s @@ -364,14 +387,6 @@     <*> transformExp e5 -- input image     <*> pure loc -transformExp (Zip i e1 es t loc) = do-  e1' <- transformExp e1-  es' <- mapM transformExp es-  return $ Zip i e1' es' t loc--transformExp (Unzip e0 tps loc) =-  Unzip <$> transformExp e0 <*> pure tps <*> pure loc- transformExp (Unsafe e1 loc) =   Unsafe <$> transformExp e1 <*> pure loc @@ -380,7 +395,7 @@  transformExp e@VConstr0{} = return e transformExp (Match e cs t loc) =-  Match <$> transformExp e <*> mapM transformCase cs <*> pure t <*> pure loc+  Match <$> transformExp e <*> mapM transformCase cs <*> traverse transformType t <*> pure loc  transformCase :: Case -> MonoM Case transformCase (CasePat p e loc) = do@@ -447,7 +462,7 @@ expandRecordPattern (Id v (Info (Record fs)) loc) = do   let fs' = M.toList fs   (fs_ks, fs_ts) <- fmap unzip $ forM fs' $ \(f, ft) ->-    (,) <$> newVName (nameToString f) <*> pure ft+    (,) <$> newVName (nameToString f) <*> transformType ft   return (RecordPattern (zip (map fst fs')                              (zipWith3 Id fs_ks (map Info fs_ts) $ repeat loc))                         loc,@@ -463,7 +478,9 @@ expandRecordPattern (PatternParens pat loc) = do   (pat', rr) <- expandRecordPattern pat   return (PatternParens pat' loc, rr)-expandRecordPattern (Wildcard t loc) = return (Wildcard t loc, mempty)+expandRecordPattern (Wildcard t loc) = do+  t' <- traverse transformType t+  return (Wildcard t' loc, mempty) expandRecordPattern (PatternAscription pat td loc) = do   (pat', rr) <- expandRecordPattern pat   return (PatternAscription pat' td loc, rr)@@ -472,24 +489,27 @@ -- | Monomorphize a polymorphic function at the types given in the instance -- list. Monomorphizes the body of the function as well. Returns the fresh name -- of the generated monomorphic function and its 'ValBind' representation.-monomorphizeBinding :: PolyBinding -> TypeBase () () -> MonoM (VName, ValBind)-monomorphizeBinding (PolyBinding rr (name, tparams, params, retdecl, rettype, body, loc)) t =+monomorphizeBinding :: Bool -> PolyBinding -> TypeBase () () -> MonoM (VName, ValBind)+monomorphizeBinding entry (PolyBinding rr (name, tparams, params, retdecl, rettype, body, loc)) t =   replaceRecordReplacements rr $ do   t' <- removeTypeVariablesInType t   let bind_t = foldFunType (map (toStructural . patternType) params) $                toStructural rettype-      substs = M.map Subst $ typeSubsts bind_t t'-      rettype' = applySubst (`M.lookup` substs) rettype-      params' = map (substPattern $ applySubst (`M.lookup` substs)) params+  (substs, t_shape_params) <- typeSubstsM loc bind_t t'+  let substs' = M.map Subst substs+      rettype' = substTypesAny (`M.lookup` substs') rettype+      substPatternType =+        substTypesAny (fmap (fmap fromStruct) . (`M.lookup` substs'))+      params' = map (substPattern entry substPatternType) params    (params'', rrs) <- unzip <$> mapM expandRecordPattern params'    mapM_ noticeDims $ rettype : map patternStructType params'' -  body' <- updateExpTypes (`M.lookup` substs) body+  body' <- updateExpTypes (`M.lookup` M.map (fmap toStructural) substs') body   body'' <- withRecordReplacements (mconcat rrs) $ transformExp body'   name' <- if null tparams then return name else newName name-  return (name', toValBinding name' params'' rettype' body'')+  return (name', toValBinding t_shape_params name' params'' rettype' body'')    where shape_params = filter (not . isTypeParam) tparams @@ -503,45 +523,69 @@                                   , mapOnPatternType = pure . applySubst substs                                   } -        toValBinding name' params'' rettype' body'' =+        toValBinding t_shape_params name' params'' rettype' body'' =           ValBind { valBindEntryPoint = False                   , valBindName       = name'                   , valBindRetDecl    = retdecl                   , valBindRetType    = Info rettype'-                  , valBindTypeParams = shape_params+                  , valBindTypeParams = shape_params ++ t_shape_params                   , valBindParams     = params''                   , valBindBody       = body''                   , valBindDoc        = Nothing                   , valBindLocation   = loc                   } -typeSubsts :: TypeBase () () -> TypeBase () ()-           -> M.Map VName (TypeBase () ())-typeSubsts (Record fields1) (Record fields2) =-  mconcat $ zipWith typeSubsts-  (map snd $ sortFields fields1) (map snd $ sortFields fields2)-typeSubsts (TypeVar _ _ v _) t =-  M.singleton (typeLeaf v) t-typeSubsts Prim{} Prim{} = mempty-typeSubsts (Arrow _ _ t1a t1b) (Arrow _ _ t2a t2b) =-  typeSubsts t1a t2a <> typeSubsts t1b t2b-typeSubsts t1@Array{} t2@Array{}-  | Just t1' <- peelArray (arrayRank t1) t1,-    Just t2' <- peelArray (arrayRank t1) t2 =-      typeSubsts t1' t2'-typeSubsts Enum{} Enum{} = mempty-typeSubsts t1 t2 = error $ unlines ["typeSubsts: mismatched types:", pretty t1, pretty t2]+-- Careful not to introduce size parameters for non-positive positions+-- (i.e. function parameters).+typeSubstsM :: MonadFreshNames m =>+               SrcLoc -> TypeBase () () -> TypeBase () ()+            -> m (M.Map VName StructType, [TypeParam])+typeSubstsM loc orig_t1 orig_t2 =+  let (t1_pts, t1_rt) = unfoldFunType orig_t1+      (t2_pts, t2_rt) = unfoldFunType orig_t2+      m = do zipWithM_ (sub True) t1_pts t2_pts+             sub False t1_rt t2_rt+  in runWriterT $ execStateT m mempty +  where sub pos (TypeVar _ _ v _) t = addSubst pos v t+        sub pos (Record fields1) (Record fields2) =+          zipWithM_ (sub pos)+          (map snd $ sortFields fields1) (map snd $ sortFields fields2)+        sub _ Prim{} Prim{} = return ()+        sub _ Enum{} Enum{} = return ()+        sub pos t1@Array{} t2@Array{}+          | Just t1' <- peelArray (arrayRank t1) t1,+            Just t2' <- peelArray (arrayRank t1) t2 =+              sub pos t1' t2'+        sub pos (Arrow _ _ t1a t1b) (Arrow _ _ t2a t2b) = do+          sub False t1a t2a+          sub pos t1b t2b++        sub _ t1 t2 = error $ unlines ["typeSubstsM: mismatched types:", pretty t1, pretty t2]++        addSubst pos (TypeName _ v) t = do+          exists <- gets $ M.member v+          unless exists $ do+            t' <- if pos+                  then bitraverse onDim pure t+                  else pure $ vacuousShapeAnnotations t+            modify $ M.insert v t'++        onDim () = do d <- lift $ lift $ newVName "d"+                      tell [TypeParamDim d loc]+                      return $ NamedDim $ qualName d+ -- | Perform a given substitution on the types in a pattern.-substPattern :: (PatternType -> PatternType) -> Pattern -> Pattern-substPattern f pat = case pat of-  TuplePattern pats loc       -> TuplePattern (map (substPattern f) pats) loc+substPattern :: Bool -> (PatternType -> PatternType) -> Pattern -> Pattern+substPattern entry f pat = case pat of+  TuplePattern pats loc       -> TuplePattern (map (substPattern entry f) pats) loc   RecordPattern fs loc        -> RecordPattern (map substField fs) loc-    where substField (n, p) = (n, substPattern f p)-  PatternParens p loc         -> PatternParens (substPattern f p) loc+    where substField (n, p) = (n, substPattern entry f p)+  PatternParens p loc         -> PatternParens (substPattern entry f p) loc   Id vn (Info tp) loc         -> Id vn (Info $ f tp) loc   Wildcard (Info tp) loc      -> Wildcard (Info $ f tp) loc-  PatternAscription p td loc  -> PatternAscription (substPattern f p) td loc+  PatternAscription p td loc | entry     -> PatternAscription (substPattern False f p) td loc+                             | otherwise -> substPattern False f p   PatternLit e (Info tp) loc  -> PatternLit e (Info $ f tp) loc  toPolyBinding :: ValBind -> PolyBinding@@ -549,27 +593,26 @@   PolyBinding mempty (name, tparams, params, retdecl, rettype, body, loc)  -- | Remove all type variables and type abbreviations from a value binding.-removeTypeVariables :: ValBind -> MonoM ValBind-removeTypeVariables valbind@(ValBind _ _ _ (Info rettype) _ pats body _ _) = do+removeTypeVariables :: Bool -> ValBind -> MonoM ValBind+removeTypeVariables entry valbind@(ValBind _ _ _ (Info rettype) _ pats body _ _) = do   subs <- asks $ M.map TypeSub . envTypeBindings-  let substPatternType = fromStruct . substituteTypes subs . toStruct-      mapper = ASTMapper {+  let mapper = ASTMapper {           mapOnExp         = astMap mapper         , mapOnName        = pure         , mapOnQualName    = pure         , mapOnType        = pure . removeShapeAnnotations .                              substituteTypes subs . vacuousShapeAnnotations-        , mapOnCompType    = pure . fromStruct . removeShapeAnnotations .+        , mapOnCompType    = pure . removeShapeAnnotations .                              substituteTypes subs .-                             vacuousShapeAnnotations . toStruct+                             vacuousShapeAnnotations         , mapOnStructType  = pure . substituteTypes subs-        , mapOnPatternType = pure . substPatternType+        , mapOnPatternType = pure . substituteTypes subs         }    body' <- astMap mapper body    return valbind { valBindRetType = Info $ substituteTypes subs rettype-                 , valBindParams  = map (substPattern substPatternType) pats+                 , valBindParams  = map (substPattern entry $ substituteTypes subs) pats                  , valBindBody    = body'                  } @@ -580,12 +623,12 @@  transformValBind :: ValBind -> MonoM Env transformValBind valbind = do-  valbind' <- toPolyBinding <$> removeTypeVariables valbind+  valbind' <- toPolyBinding <$> removeTypeVariables (valBindEntryPoint valbind) valbind   when (valBindEntryPoint valbind) $ do     t <- removeTypeVariablesInType $ removeShapeAnnotations $ foldFunType          (map patternStructType (valBindParams valbind)) $          unInfo $ valBindRetType valbind-    (name, valbind'') <- monomorphizeBinding valbind' t+    (name, valbind'') <- monomorphizeBinding True valbind' t     tell $ Seq.singleton (name, valbind'' { valBindEntryPoint = True})     addLifted (valBindName valbind) t name   return mempty { envPolyBindings = M.singleton (valBindName valbind) valbind' }

@@ -47,8 +47,9 @@ import qualified Futhark.Representation.ExplicitMemory as ExplicitMemory import Futhark.Transform.Substitute import Futhark.Pass+import Futhark.Util (takeLast) --- | Perform CSE on every functioon in a program.+-- | Perform CSE on every function in a program. performCSE :: (Attributes lore, CanBeAliased (Op lore),                CSEInOp (OpWithAliases (Op lore))) =>               Bool -> Pass lore lore@@ -61,22 +62,27 @@                Bool -> FunDef lore -> FunDef lore cseInFunDef cse_arrays fundec =   fundec { funDefBody =-              runReader (cseInBody $ funDefBody fundec) $ newCSEState cse_arrays+              runReader (cseInBody ds $ funDefBody fundec) $ newCSEState cse_arrays          }+  where ds = map (diet . declExtTypeOf) $ funDefRetType fundec  type CSEM lore = Reader (CSEState lore)  cseInBody :: (Attributes lore, Aliased lore, CSEInOp (Op lore)) =>-             Body lore -> CSEM lore (Body lore)-cseInBody (Body bodyattr bnds res) =-  cseInStms (consumedInStms bnds res) (stmsToList bnds) $ do+             [Diet] -> Body lore -> CSEM lore (Body lore)+cseInBody ds (Body bodyattr bnds res) =+  cseInStms (res_cons <> consumedInStms bnds res) (stmsToList bnds) $ do     CSEState (_, nsubsts) _ <- ask     return $ Body bodyattr mempty $ substituteNames nsubsts res+  where res_cons = mconcat $ zipWith consumeResult ds $+                   takeLast (length ds) res+        consumeResult Consume se = freeIn se+        consumeResult _ _ = mempty  cseInLambda :: (Attributes lore, Aliased lore, CSEInOp (Op lore)) =>                Lambda lore -> CSEM lore (Lambda lore) cseInLambda lam = do-  body' <- cseInBody $ lambdaBody lam+  body' <- cseInBody (map (const Observe) $ lambdaReturnType lam) $ lambdaBody lam   return lam { lambdaBody = body' }  cseInStms :: (Attributes lore, Aliased lore, CSEInOp (Op lore)) =>@@ -90,12 +96,17 @@     bnd'' <- mapM nestedCSE bnd'     return $ Body bodyattr (stmsFromList bnd''<>bnds') es   where nestedCSE bnd' = do-          e <- mapExpM cse $ stmExp bnd'+          let ds = map patElemDiet $ patternValueElements $ stmPattern bnd'+          e <- mapExpM (cse ds) $ stmExp bnd'           return bnd' { stmExp = e }-        cse = identityMapper { mapOnBody = const cseInBody-                             , mapOnOp = cseInOp-                             } +        cse ds = identityMapper { mapOnBody = const $ cseInBody ds+                                , mapOnOp = cseInOp+                                }++        patElemDiet pe | patElemName pe `S.member` consumed = Consume+                       | otherwise                          = Observe+ cseInStm :: Attributes lore =>             Names -> Stm lore          -> ([Stm lore] -> CSEM lore a)@@ -169,29 +180,34 @@ instance (Attributes lore, Aliased lore, CSEInOp (Op lore)) => CSEInOp (Kernel.Kernel lore) where   cseInOp = subCSE .             Kernel.mapKernelM-            (Kernel.KernelMapper return cseInLambda cseInBody+            (Kernel.KernelMapper return cseInLambda+             (\b -> cseInBody (map (const Observe) $ bodyResult b) b)              return return cseInKernelBody)  cseInKernelBody :: (Attributes lore, Aliased lore, CSEInOp (Op lore)) =>                    Kernel.KernelBody lore -> CSEM lore (Kernel.KernelBody lore) cseInKernelBody (Kernel.KernelBody bodyattr bnds res) = do-  Body _ bnds' _ <- cseInBody $ Body bodyattr bnds []+  Body _ bnds' _ <- cseInBody (map (const Observe) res) $ Body bodyattr bnds []   return $ Kernel.KernelBody bodyattr bnds' res  instance (Attributes lore, Aliased lore, CSEInOp (Op lore)) => CSEInOp (KernelExp.KernelExp lore) where   cseInOp (KernelExp.Combine cspace ts active body) =-    subCSE $ KernelExp.Combine cspace ts active <$> cseInBody body+    subCSE $ KernelExp.Combine cspace ts active <$>+    cseInBody (map (const Observe) ts) body   cseInOp (KernelExp.GroupReduce w lam input) =-    subCSE $ KernelExp.GroupReduce w <$> cseInLambda lam <*> pure input+    subCSE $ KernelExp.GroupReduce w <$>+    cseInLambda lam <*> pure input   cseInOp (KernelExp.GroupStream w max_chunk lam nes arrs) =-    subCSE $ KernelExp.GroupStream w max_chunk <$> cseInGroupStreamLambda lam <*> pure nes <*> pure arrs+    subCSE $ KernelExp.GroupStream w max_chunk <$>+    cseInGroupStreamLambda lam <*> pure nes <*> pure arrs   cseInOp op = return op  cseInGroupStreamLambda :: (Attributes lore, Aliased lore, CSEInOp (Op lore)) =>                           KernelExp.GroupStreamLambda lore                        -> CSEM lore (KernelExp.GroupStreamLambda lore) cseInGroupStreamLambda lam = do-  body' <- cseInBody $ KernelExp.groupStreamLambdaBody lam+  body' <- cseInBody (map (const Observe) $ KernelExp.groupStreamAccParams lam) $+           KernelExp.groupStreamLambdaBody lam   return lam { KernelExp.groupStreamLambdaBody = body' }  

@@ -63,13 +63,14 @@ inlineInBody inlcallees (Body attr stms res) = Body attr stms' res   where stms' = stmsFromList (concatMap inline $ stmsToList stms) -        inline (Let pat _ (Apply fname args _ (safety,loc,locs)))+        inline (Let pat aux (Apply fname args _ (safety,loc,locs)))           | fun:_ <- filter ((== fname) . funDefName) inlcallees =               let param_stms = zipWith reshapeIfNecessary (map paramIdent $ funDefParams fun) (map fst args)                   body_stms = stmsToList $ addLocations safety                               (filter notNoLoc (loc:locs)) $ bodyStms $ funDefBody fun-                  res_stms = zipWith reshapeIfNecessary (patternIdents pat)-                             (bodyResult $ funDefBody fun)+                  res_stms = map (certify $ stmAuxCerts aux) $+                             zipWith reshapeIfNecessary (patternIdents pat) $+                             bodyResult $ funDefBody fun               in param_stms ++ body_stms ++ res_stms         inline stm = [inlineInStm inlcallees stm] 

@@ -1,28 +0,0 @@--- | Merge memory blocks.-module Futhark.Optimise.MemoryBlockMerging-  ( memoryBlockMergingCoalescing-  , memoryBlockMergingReuse-  ) where--import Futhark.Pass-import Futhark.Representation.ExplicitMemory (ExplicitMemory)--import Futhark.Optimise.MemoryBlockMerging.Coalescing (coalesceInProg)-import Futhark.Optimise.MemoryBlockMerging.Reuse (reuseInProg)----- | Apply the coalescing part of the memory block merging optimisation.-memoryBlockMergingCoalescing :: Pass ExplicitMemory ExplicitMemory-memoryBlockMergingCoalescing =-  Pass-  "Memory block merging (coalescing)"-  "Coalesce the memory blocks of arrays"-  coalesceInProg---- | Apply the reuse part of the memory block merging optimisation.-memoryBlockMergingReuse :: Pass ExplicitMemory ExplicitMemory-memoryBlockMergingReuse =-  Pass-  "Memory block merging (reuse)"-  "Reuse the memory blocks of arrays"-  reuseInProg

@@ -1,358 +0,0 @@-{-# LANGUAGE GeneralizedNewtypeDeriving #-}-{-# LANGUAGE TypeFamilies #-}-{-# LANGUAGE FlexibleContexts #-}-{-# LANGUAGE ConstraintKinds #-}--- | Find the actual variables that need updating when a variable attribute--- needs updating.  This is different than variable aliasing: Variable aliasing--- is a theoretical concept, while this module has the practical purpose of--- finding any extra variables that also need a change when a variable has a--- change of memory block.------ If and DoLoop statements have special requirements, as do some aliasing--- expressions.  We don't want to (just) use the obvious statement variable;--- sometimes updating the memory block of one variable actually means updating--- the memory block of other variables as well.--module Futhark.Optimise.MemoryBlockMerging.ActualVariables-  ( findActualVariables-  ) where--import qualified Data.Set as S-import qualified Data.Map.Strict as M-import qualified Data.List as L-import Data.Maybe (fromMaybe, mapMaybe, catMaybes)-import Control.Monad-import Control.Monad.RWS--import Futhark.Representation.AST-import Futhark.Representation.ExplicitMemory (-  ExplicitMemorish, ExplicitMemory, InKernel)-import qualified Futhark.Representation.ExplicitMemory as ExpMem-import Futhark.Representation.Kernels.Kernel--import Futhark.Optimise.MemoryBlockMerging.Miscellaneous-import Futhark.Optimise.MemoryBlockMerging.Types-import Futhark.Optimise.MemoryBlockMerging.AllExpVars---data Context = Context-  { ctxVarToMem :: VarMemMappings MemorySrc-  , ctxFirstUses :: FirstUses-  }-  deriving (Show)--newtype FindM lore a = FindM { unFindM :: RWS Context () ActualVariables a }-  deriving (Monad, Functor, Applicative,-            MonadReader Context,-            MonadState ActualVariables)--type LoreConstraints lore = (ExplicitMemorish lore,-                             FullWalk lore,-                             LookInKernelExp lore)--coerce :: FindM flore a -> FindM tlore a-coerce = FindM . unFindM--recordActuals :: VName -> Names -> FindM lore ()-recordActuals stmt_var more_actuals = do-  -- If S.empty has already been recorded, keep it at that.  This is because the-  -- ActualVariables system is currently also used for disabling memory block-  -- optimisations -- if a variables resolves to the empty set, don't touch it.-  -- This keeps some edge cases simple.  FIXME at some point.-  current_actuals <- M.lookup stmt_var <$> get-  case S.null <$> current_actuals of-    Just True -> return ()-    _ -> modify (insertOrUpdateMany stmt_var more_actuals)---- Find all the actual variables in a function definition.-findActualVariables :: VarMemMappings MemorySrc -> FirstUses ->-                       FunDef ExplicitMemory -> ActualVariables-findActualVariables var_mem_mappings first_uses fundef =-  let context = Context var_mem_mappings first_uses-      m = unFindM $ lookInBody $ funDefBody fundef-      actual_variables = fst $ execRWS m context M.empty-  in actual_variables--lookInFParam :: FParam lore -> FindM lore ()-lookInFParam (Param v _) =-  recordActuals v $ S.singleton v--lookInLParam :: LParam lore -> FindM lore ()-lookInLParam (Param v _) =-  recordActuals v $ S.singleton v--lookInLambda :: LoreConstraints lore => Lambda lore -> FindM lore ()-lookInLambda (Lambda params body _) = do-  forM_ params lookInLParam-  lookInBody body--lookInBody :: LoreConstraints lore =>-              Body lore -> FindM lore ()-lookInBody (Body _ bnds _res) =-  mapM_ lookInStm bnds--lookInKernelBody :: LoreConstraints lore =>-                    KernelBody lore -> FindM lore ()-lookInKernelBody (KernelBody _ bnds _res) =-  mapM_ lookInStm bnds--lookInStm :: LoreConstraints lore =>-             Stm lore -> FindM lore ()-lookInStm stm@(Let (Pattern patctxelems patvalelems) _ e) = do-  case (patvalelems, e) of-    ([PatElem var _], BasicOp (Update orig _ _)) -> do-      let actuals = S.fromList [var, orig]-      -- When coalescing an in-place update statement, also look at the original-      -- array.-      recordActuals var actuals-      -- When reusing a previous memory block, make sure to also update related-      -- in-place updates.-      recordActuals orig actuals-    _ -> return ()--  -- Ignore the existential memory blocks.-  let bodyResult' = drop (length patctxelems) . bodyResult--  -- Special handling of loops, ifs, etc.-  case e of-    DoLoop _mergectxparams mergevalparams loopform body -> do-      let body_vars0 = mapMaybe (subExpVar . snd) mergevalparams-          body_vars1 = map (paramName . fst) mergevalparams-          body_vars2 = S.toList $ findAllExpVars e-          body_vars = body_vars0 ++ body_vars1 ++ body_vars2-      forM_ patvalelems $ \(PatElem var membound) -> do-        case membound of-          ExpMem.MemArray _ _ _ (ExpMem.ArrayIn mem _) -> do-            -- If mem is existential, we need to find the return memory that it-            -- refers to.  We cannot just look at its memory aliases, since it-            -- likely aliases both the initial memory and the final memory.--            let zipped = zip patctxelems (bodyResult body)-                mem_search = case L.find ((== mem) . patElemName . fst) zipped of-                  Just (_, Var res_mem) -> res_mem-                  _ -> mem-            -- Find the ones using the same memory as the result of the loop-            -- expression.-            body_vars' <- filterM (lookupGivesMem mem_search) body_vars-            -- Not only the result variable needs to change its memory block in-            -- case of a future memory merging with it; also the variables-            -- extracted above.-            let actuals = var : body_vars'-            forM_ actuals $ \a -> recordActuals a (S.fromList actuals)-            -- Some of these can be changed later on to have an actual variable-            -- set of S.empty, e.g. if one of the variables using the memory is-            -- a rearrange operation.  This is fine, and will occur in the walk-            -- later on.--            -- If you extend this loop handling, make sure not to target existential-            -- memory blocks.  We want those to stay.-          _ -> return ()--        -- It seems wrong to change the memory of merge variables, so we disable-        -- it.  If we were to accept it, we would need to record what other-        -- variables to change as well.  Seems hard.-        recordActuals var S.empty--      case loopform of-        ForLoop _ _ _ loop_vars ->-          -- Link 'array' to 'lvar' in 'for lvar in array' loop expressions.-          forM_ loop_vars $ \(Param lvar _, array) ->-            aliasOpHandleVar array lvar-        WhileLoop _ -> return ()--    If _se body_then body_else _types ->-      -- We don't want to coalesce the existiential memory block of the if.-      -- However, if a branch result has a memory block that is firstly used-      -- inside the branch, it is okay to coalesce that in a future statement.-      forM_ (zip3 patvalelems (bodyResult' body_then) (bodyResult' body_else))-        $ \(PatElem var membound, res_then, res_else) -> do-        let body_vars = S.toList $ findAllExpVars e-        case membound of-          ExpMem.MemArray _ _ _ (ExpMem.ArrayIn mem _) ->-            if mem `L.elem` map patElemName patctxelems-              then-              -- If the memory block is existential, we say that the If result-              -- refers to all results in the If.-              recordActuals var-              $ S.fromList (var : catMaybes [subExpVar res_then, subExpVar res_else])--              else do-              -- If the memory block is not existential, we need to find all the-              -- variables in any sub-bodies using the same memory block (like-              -- with loops).-              body_vars' <- filterM (lookupGivesMem mem) body_vars--              first_uses <- asks ctxFirstUses-              case filter ((mem `S.member`) . (`lookupEmptyable` first_uses)) body_vars' of-                [] ->-                  -- Not just the result variable needs to change its memory-                  -- block in case of a future memory block merging with it;-                  -- also the variables extracted above.-                  recordActuals var $ S.fromList (var : body_vars')-                _ ->-                  -- If we come across a non-existential If which can be said to-                  -- create a new array *and* which has one or more bodies which-                  -- can also be said to create a new array *in the same memory*-                  -- (i.e. has first memory uses), then we disable it.  This is-                  -- not at all an impossible case to handle, but such an If is-                  -- weird, since it would make more sense if it had existential-                  -- memory, so maybe something needs to be done somewhere else-                  -- in the compiler?  If this is naively enabled, we can get an-                  -- error because the sub-body results are first uses while the-                  -- main result is not.  This can be "fixed" by stating that-                  -- the If as a whole is also a first use of the memory, but-                  -- this seems too conservative.  FIXME.-                  forM_ (var : body_vars') $ \v -> recordActuals v S.empty--          _ -> return ()--    BasicOp (Index orig _) -> do-      let ielem = head patvalelems -- Should be okay.-          var = patElemName ielem-      case patElemAttr ielem of-        ExpMem.MemArray{} ->-          -- Disable merging for index expressions that return arrays.  Maybe-          -- too restrictive.  Make sure the source also updates the memory of-          -- the index when updated.  The array might be an aliasing operation,-          -- in which case we try to find the original array.-          aliasOpHandleVar orig var-        _ -> return ()--    -- Support reusing the memory of reshape operations by recording the origin-    -- array that is being reshaped.  Only partial support for reshape-    -- operations: If the shape is more than one-dimensional, mark the statement-    -- as disabled for memory merging operations.-    BasicOp (Reshape shapechange_var orig) ->-      forM_ (map patElemName patvalelems) $ \var -> do-        orig' <- aliasOpRoot' orig-        mem_orig <- M.lookup orig' <$> asks ctxVarToMem-        case (shapechange_var, mem_orig) of-          ([_], Just (MemorySrc _ _ (Shape [_]))) ->-            recordActuals var $ S.fromList [var, orig]-            -- Works, but only in limited cases where the reshape is not even-            -- that useful to begin with; mostly cases where a reshape was-            -- inserted by the compiler in an assert-like manner.-          _ ->-            recordActuals var S.empty-            -- FIXME: The problem with these more complex cases with more than-            -- one dimension is that a slice is relative to the shape of the-            -- reshaped array, and not the original array.  Disabled for now.-        recordActuals orig' $ S.fromList [orig', var]--    -- For the other aliasing operations, disable their use for now.  If the-    -- source has a change of memory block, make sure to change this as well.-    BasicOp (Rearrange _ orig) ->-      aliasOpHandle orig patvalelems--    BasicOp (Rotate _ orig) ->-      aliasOpHandle orig patvalelems--    BasicOp (Opaque (Var orig)) ->-      aliasOpHandle orig patvalelems--    _ -> forM_ patvalelems $ \(PatElem var membound) -> do-      let body_vars = S.toList $ findAllExpVars e-      case membound of-        ExpMem.MemArray _ _ _ (ExpMem.ArrayIn mem _) -> do-          body_vars' <- filterM (lookupGivesMem mem) body_vars-          recordActuals var $ S.fromList (var : body_vars')-        _ -> return ()--  -- If we are inside a kernel, check for actual variables in the KernelExp of-  -- the statement.-  lookInKernelExp stm--  -- Recurse over any sub-bodies.-  fullWalkExpM walker walker_kernel e-  where walker = identityWalker-          { walkOnBody = lookInBody-          , walkOnFParam = lookInFParam-          , walkOnLParam = lookInLParam-          }-        walker_kernel = identityKernelWalker-          { walkOnKernelBody = coerce . lookInBody-          , walkOnKernelKernelBody = coerce . lookInKernelBody-          , walkOnKernelLambda = coerce . lookInLambda-          , walkOnKernelLParam = lookInLParam-          }---- If we have a rotate or similar, we want to find the original array and--- associate *that* with this aliasing array, so that changes to the original--- array will affect this one as well.-aliasOpHandle :: VName -> [PatElem lore] -> FindM lore ()-aliasOpHandle orig patvalelems =-  forM_ (map patElemName patvalelems) $ aliasOpHandleVar orig--aliasOpHandleVar :: VName -> VName -> FindM lore ()-aliasOpHandleVar orig var = do-  recordActuals var S.empty-  orig' <- aliasOpRoot' orig-  recordActuals orig' $ S.fromList [orig', var]--aliasOpRoot :: VName -> FindM lore (Maybe VName)-aliasOpRoot orig = do-  current_actuals <- get-  return $ case S.null <$> M.lookup orig current_actuals of-    -- If the original array is itself an aliasing operation, find the *actual*-    -- original array.  There can be more than one reference.  We just pick the-    -- first one -- any one should do, since there is a transitive closure-    -- calculation later on.-    Just True -> case M.keys (M.filter (orig `S.member`) current_actuals) of-      orig' : _ -> Just orig'-      _ -> Nothing-    -- Else, just return orig.-    _ -> Just orig--aliasOpRoot' :: VName -> FindM lore VName-aliasOpRoot' orig =-  fromJust ("at some point there will have been a proper statement: "-            ++ pretty orig) <$> aliasOpRoot orig---- Is the memory block of 'v' the same as 'mem'?-lookupGivesMem :: MName -> VName -> FindM lore Bool-lookupGivesMem mem v = do-  m <- M.lookup v <$> asks ctxVarToMem-  return (Just mem == (memSrcName <$> m))--class LookInKernelExp lore where-  -- Find actual vars in 'KernelExp's.-  lookInKernelExp :: Stm lore -> FindM lore ()--instance LookInKernelExp ExplicitMemory where-  lookInKernelExp (Let (Pattern _ patvalelems) _ e) = case e of-    Op (ExpMem.Inner (Kernel _ _ _ (KernelBody _ _ ress))) ->-      zipWithM_ (\(PatElem var _) res -> case res of-                    WriteReturn _ arr _ ->-                      recordActuals arr $ S.singleton var-                    _ -> return ()-                ) patvalelems ress-    _ -> return ()--instance LookInKernelExp InKernel where-  lookInKernelExp (Let _ _ e) = case e of-    Op (ExpMem.Inner ke) -> case ke of-      ExpMem.GroupReduce _ _ input -> do-        let arrs = map snd input-        extendActualVarsInKernel e arrs-      ExpMem.GroupScan _ _ input -> do-        let arrs = map snd input-        extendActualVarsInKernel e arrs-      ExpMem.GroupStream _ _ _ _ arrs ->-        extendActualVarsInKernel e arrs-      _ -> return ()-    _ -> return ()---- Record actual variables for input arrays to 'KernelExp's.-extendActualVarsInKernel :: Exp InKernel -> [VName] -> FindM InKernel ()-extendActualVarsInKernel e arrs = forM_ arrs $ \var -> do-  -- The array might be an aliasing operation, in which case we try to find the-  -- original array.-  var' <- fromMaybe var <$> aliasOpRoot var-  varmem <- M.lookup var <$> asks ctxVarToMem-  case varmem of-    Just mem -> do-      let body_vars = findAllExpVars e-      body_vars' <- filterSetM (lookupGivesMem $ memSrcName mem) body_vars-      let actuals = S.insert var' body_vars'-      recordActuals var' actuals-    Nothing -> return ()

@@ -1,96 +0,0 @@-{-# LANGUAGE GeneralizedNewtypeDeriving #-}-{-# LANGUAGE TypeFamilies #-}-{-# LANGUAGE FlexibleContexts #-}-{-# LANGUAGE ConstraintKinds #-}-{-# LANGUAGE LambdaCase #-}--- | Find all variables in a statement.-module Futhark.Optimise.MemoryBlockMerging.AllExpVars-  ( findAllExpVars-  ) where--import qualified Data.Set as S-import Control.Monad-import Control.Monad.Writer--import Futhark.Representation.AST-import Futhark.Representation.ExplicitMemory (ExplicitMemorish)-import Futhark.Representation.Kernels.Kernel--import Futhark.Optimise.MemoryBlockMerging.Miscellaneous---newtype FindM lore a = FindM { unFindM :: Writer Names a }-  deriving (Monad, Functor, Applicative,-            MonadWriter Names)--type LoreConstraints lore = (ExplicitMemorish lore,-                             FullWalk lore)--coerce :: FindM flore a -> FindM tlore a-coerce = FindM . unFindM---- Find all the variables (both free and bound) that occur in a statement and--- any nested bodies.  We use this to record which extra variables need to have--- their memory blocks updated when some variable needs updating.  The result--- might be an empty set, but in the case of If, DoLoop, and kernels, the result--- might be nonempty.  We cannot just find all variables in the program and look--- through them every time we need to, since a memory block can (at least in--- theory) be present in two different places (which also means by two different--- variable sets) in a program, so we should limit ourselves to looking in the--- statement declaring a new current use of the memory.-findAllExpVars :: LoreConstraints lore =>-                  Exp lore -> Names-findAllExpVars e =-  let m = unFindM $ lookInExp e-  in execWriter m--lookInExp :: LoreConstraints lore =>-             Exp lore -> FindM lore ()-lookInExp = fullWalkExpM walker walker_kernel-  where walker = identityWalker-          { walkOnBody = lookInBody-          , walkOnFParam = lookInFParam-          , walkOnLParam = lookInLParam-          }-        walker_kernel = identityKernelWalker-          { walkOnKernelBody = coerce . lookInBody-          , walkOnKernelKernelBody = coerce . lookInKernelBody-          , walkOnKernelLambda = coerce . lookInLambda-          , walkOnKernelLParam = lookInLParam-          }--lookInFParam :: FParam lore -> FindM lore ()-lookInFParam (Param x _) =-  tell $ S.singleton x--lookInLParam :: LParam lore -> FindM lore ()-lookInLParam (Param x _) =-  tell $ S.singleton x--lookInBody :: LoreConstraints lore =>-              Body lore -> FindM lore ()-lookInBody (Body _ bnds _res) =-  mapM_ lookInStm bnds--lookInKernelBody :: LoreConstraints lore =>-                    KernelBody lore -> FindM lore ()-lookInKernelBody (KernelBody _ bnds res) = do-  mapM_ lookInStm bnds-  forM_ res $ \case-    ThreadsReturn{} -> return ()-    WriteReturn _ arr _ -> tell $ S.singleton arr-    ConcatReturns{} -> return ()-    KernelInPlaceReturn v -> tell $ S.singleton v--lookInStm :: LoreConstraints lore =>-             Stm lore -> FindM lore ()-lookInStm (Let (Pattern _ patvalelems) _ e) = do-  forM_ patvalelems $ \(PatElem x _) ->-    tell $ S.singleton x-  lookInExp e--lookInLambda :: LoreConstraints lore =>-                Lambda lore -> FindM lore ()-lookInLambda (Lambda params body _) = do-  forM_ params lookInLParam-  lookInBody body

@@ -1,63 +0,0 @@-{-# LANGUAGE TypeFamilies #-}-{-# LANGUAGE FlexibleContexts #-}--- | Helper information for the main optimisation passes.-module Futhark.Optimise.MemoryBlockMerging.AuxiliaryInfo-  ( AuxiliaryInfo(..), getAuxiliaryInfo)-where--import Futhark.Representation.AST-import Futhark.Representation.ExplicitMemory (ExplicitMemory)--import Futhark.Optimise.MemoryBlockMerging.Types--import Futhark.Optimise.MemoryBlockMerging.VariableMemory-import Futhark.Optimise.MemoryBlockMerging.MemoryAliases-import Futhark.Optimise.MemoryBlockMerging.VariableAliases-import Futhark.Optimise.MemoryBlockMerging.Liveness.FirstUse-import Futhark.Optimise.MemoryBlockMerging.Liveness.LastUse-import Futhark.Optimise.MemoryBlockMerging.Liveness.Interference-import Futhark.Optimise.MemoryBlockMerging.ActualVariables-import Futhark.Optimise.MemoryBlockMerging.Existentials---- Information needed by multiple transformations.-data AuxiliaryInfo = AuxiliaryInfo-  { auxName :: Name -- For debugging.-  , auxVarMemMappings :: VarMemMappings MemorySrc-  , auxMemAliases :: MemAliases-  , auxVarAliases :: VarAliases-  , auxFirstUses :: FirstUses-  , auxLastUses :: LastUses-  , auxInterferences :: Interferences-  , auxPotentialKernelDataRaceInterferences-    :: PotentialKernelDataRaceInterferences-  , auxActualVariables :: ActualVariables-  , auxExistentials :: Names-  }-  deriving (Show)--getAuxiliaryInfo :: FunDef ExplicitMemory -> AuxiliaryInfo-getAuxiliaryInfo fundef =-  let name = funDefName fundef-      var_to_mem = findVarMemMappings fundef-      mem_aliases = findMemAliases fundef var_to_mem-      var_aliases = findVarAliases fundef-      first_uses = findFirstUses var_to_mem mem_aliases fundef-      last_uses = findLastUses var_to_mem mem_aliases first_uses existentials-                  fundef-      (interferences, potential_kernel_interferences) =-        findInterferences var_to_mem mem_aliases first_uses last_uses-        existentials fundef-      actual_variables = findActualVariables var_to_mem first_uses fundef-      existentials = findExistentials fundef-  in AuxiliaryInfo-     { auxName = name-     , auxVarMemMappings = var_to_mem-     , auxMemAliases = mem_aliases-     , auxVarAliases = var_aliases-     , auxFirstUses = first_uses-     , auxLastUses = last_uses-     , auxInterferences = interferences-     , auxPotentialKernelDataRaceInterferences = potential_kernel_interferences-     , auxActualVariables = actual_variables-     , auxExistentials = existentials-     }

@@ -1,31 +0,0 @@--- | Coalesce the memory blocks of arrays.------ Enable by setting the environment variable MEMORY_BLOCK_MERGING_COALESCING=1.-module Futhark.Optimise.MemoryBlockMerging.Coalescing-  ( coalesceInProg-  ) where--import Futhark.Pass--import Futhark.MonadFreshNames-import Futhark.Representation.AST-import Futhark.Representation.ExplicitMemory (ExplicitMemory)--import Futhark.Optimise.MemoryBlockMerging.AuxiliaryInfo-import Futhark.Optimise.MemoryBlockMerging.Coalescing.AllocationMovingUp-import Futhark.Optimise.MemoryBlockMerging.Coalescing.Core---coalesceInProg :: Prog ExplicitMemory -> PassM (Prog ExplicitMemory)-coalesceInProg = intraproceduralTransformation coalesceInFunDef--coalesceInFunDef :: MonadFreshNames m-                 => FunDef ExplicitMemory-                 -> m (FunDef ExplicitMemory)-coalesceInFunDef fundef0 = do-  let fundef1 = moveUpAllocsFunDef fundef0-      aux1 = getAuxiliaryInfo fundef1-  coreCoalesceFunDef fundef1-    (auxVarMemMappings aux1) (auxMemAliases aux1)-    (auxVarAliases aux1) (auxFirstUses aux1) (auxLastUses aux1)-    (auxActualVariables aux1) (auxExistentials aux1)

@@ -1,94 +0,0 @@--- | Move allocation statements upwards in the bodies of a program to enable--- more memory block coalescings.------ This should be run *before* the coalescing pass, as it enables more--- optimisations.-module Futhark.Optimise.MemoryBlockMerging.Coalescing.AllocationMovingUp-  ( moveUpAllocsFunDef-  ) where--import qualified Data.Set as S-import Data.Maybe (mapMaybe)--import Futhark.Representation.AST-import Futhark.Representation.ExplicitMemory (ExplicitMemory)-import qualified Futhark.Representation.ExplicitMemory as ExpMem--import Futhark.Optimise.MemoryBlockMerging.CrudeMovingUp---findAllocHoistees :: Body ExplicitMemory -> Maybe [FParam ExplicitMemory]-                  -> [VName]-findAllocHoistees body params =-  let all_found = mapMaybe findThemStm stms-                  ++ maybe [] (mapMaybe findThemFParam) params-      extras = concatMap snd all_found-      allocs = map fst all_found-      -- We must hoist the alloc expressions in the end.  If we hoist an alloc-      -- before we hoist one of its array creations (in case of in-place-      -- updates), that array creation might in turn hoist something depending-      -- on another memory block mem_y further up than the allocation of memory-      -- block mem_x.  This will become a problem if mem_y can get coalesced-      -- into mem_x.-      ---      -- Maybe there is a nicer way to guarantee that this does not happen, but-      -- this seems to work for now.-      ---      -- We reverse the non-alloc dependencies to ensure (sloppily) that they do-      -- not change positions internally compared to the original program: For-      -- example, if a statement x is located before a statement y, and both x-      -- and y need to be hoisted, then we need to hoist x in the end, so that-      -- it can be hoisted further than y, which might have been hoisted to-      -- before x.  A better solution is welcome!-      in reverse extras ++ reverse allocs--  where stms :: [Stm ExplicitMemory]-        stms = stmsToList $ bodyStms body--        findThemStm :: Stm ExplicitMemory -> Maybe (VName, [VName])-        findThemStm (Let (Pattern _ [PatElem xmem _]) _ (Op ExpMem.Alloc{})) =-          usedByCopyOrConcat xmem-        findThemStm _ = Nothing--        -- A function paramater can be a unique memory block.  While we cannot-        -- hoist that, we may have to hoist an index in an in-place update that-        -- uses the memory.-        findThemFParam :: FParam ExplicitMemory -> Maybe (VName, [VName])-        findThemFParam (Param xmem ExpMem.MemMem{}) = usedByCopyOrConcat xmem-        findThemFParam _ = Nothing--        -- Is the allocated memory used by either Copy or Concat in the function-        -- body?  Those are the only kinds of memory we care about, since those-        -- are the cases handled by coalescing.  Also find the names used by-        -- in-place updates, since those also need to be hoisted (as an example-        -- of this, consider the 'copy/pos1.fut' test where the replicate-        -- expression needs to be hoisted as well as its memory allocation).-        usedByCopyOrConcat :: VName -> Maybe (VName, [VName])-        usedByCopyOrConcat xmem_alloc =-          let vs = mapMaybe checkStm stms-              vs' = if null vs then Nothing else Just (xmem_alloc, concat vs)-          in vs'--          where checkStm :: Stm ExplicitMemory -> Maybe [VName]-                checkStm (Let-                          (Pattern _-                           [PatElem _ (ExpMem.MemArray _ _ _ (ExpMem.ArrayIn xmem_pat _))])-                           _-                           (BasicOp bop))-                  | xmem_pat == xmem_alloc =-                    case bop of-                      Update v slice _ ->-                        -- The source array must also be hoisted so that it-                        -- is initialized before it is used by the-                        -- coalesced party.  Any index variables are also-                        -- hoisted.-                        Just $ v : S.toList (freeIn slice)-                      Copy{} -> Just []-                      Concat{} -> Just []-                      _ -> Nothing-                checkStm _ = Nothing--moveUpAllocsFunDef :: FunDef ExplicitMemory-                  -> FunDef ExplicitMemory-moveUpAllocsFunDef fundef =-  moveUpInFunDef fundef findAllocHoistees

@@ -1,624 +0,0 @@-{-# LANGUAGE GeneralizedNewtypeDeriving #-}-{-# LANGUAGE TypeFamilies #-}-{-# LANGUAGE FlexibleContexts #-}-{-# LANGUAGE ConstraintKinds #-}-module Futhark.Optimise.MemoryBlockMerging.Coalescing.Core-  ( coreCoalesceFunDef-  ) where--import qualified Data.Set as S-import qualified Data.List as L-import qualified Data.Map.Strict as M-import Data.Maybe (maybe, fromMaybe, mapMaybe, isJust)-import Control.Monad-import Control.Monad.RWS--import Futhark.MonadFreshNames-import Futhark.Representation.AST-import Futhark.Representation.ExplicitMemory (-  ExplicitMemory, ExplicitMemorish)-import qualified Futhark.Representation.ExplicitMemory as ExpMem-import Futhark.Representation.Kernels.Kernel-import qualified Futhark.Representation.ExplicitMemory.IndexFunction as IxFun-import Futhark.Tools--import Futhark.Optimise.MemoryBlockMerging.Miscellaneous-import Futhark.Optimise.MemoryBlockMerging.Types-import Futhark.Optimise.MemoryBlockMerging.MemoryUpdater--import Futhark.Optimise.MemoryBlockMerging.PrimExps (findPrimExpsFunDef)-import Futhark.Optimise.MemoryBlockMerging.Coalescing.Exps-import Futhark.Optimise.MemoryBlockMerging.Coalescing.SafetyCondition2-import Futhark.Optimise.MemoryBlockMerging.Coalescing.SafetyCondition3-import Futhark.Optimise.MemoryBlockMerging.Coalescing.SafetyCondition5-import Futhark.Optimise.MemoryBlockMerging.Reuse.AllocationSizes----- Some of these attributes could be split into separate Coalescing helper--- modules if it becomes confusing.  Their computations are fairly independent.-data Current = Current-  { -- Coalescings state.  Also save offsets and slices in the case that an-    -- optimistic coalescing later becomes part of a chain of coalescings, where-    -- it is offset yet again, and where it should maintain its old relative-    -- offset.  FIXME: This works, but is inefficient in the long run, as we-    -- need to update it whenever we come across a coalescing that also affects-    -- previous coalescings.  The directions of the coalescings is inherently-    -- bottom-up, but our algorithm is top-down.  It should be possible to-    -- rewrite it.-    curCoalescedIntos :: CoalescedIntos-  , curMemsCoalesced :: MemsCoalesced-  }-  deriving (Show)--type CoalescedIntos = M.Map VName (S.Set (VName, PrimExp VName,-                                          [Slice (PrimExp VName)]))-type MemsCoalesced = M.Map VName MemoryLoc--emptyCurrent :: Current-emptyCurrent = Current-  { curCoalescedIntos = M.empty-  , curMemsCoalesced = M.empty-  }--data Context = Context-  { ctxFunDef :: FunDef ExplicitMemory-    -- ^ Keep the entire function definition around for lookup purposes.-  , ctxVarToMem :: VarMemMappings MemorySrc-    -- ^ From the module VariableMemory.-  , ctxMemAliases :: MemAliases-    -- ^ From the module MemoryAliases.-  , ctxVarAliases :: VarAliases-    -- ^ From the module VariableAliases.-  , ctxFirstUses :: FirstUses-    -- ^ From the module FirstUses.-  , ctxLastUses :: LastUses-    -- ^ From the module LastUses.-  , ctxActualVars :: M.Map VName Names-    -- ^ From the module ActualVariables.-  , ctxExistentials :: Names-    -- ^ From the module Existentials.-  , ctxVarPrimExps :: M.Map VName (PrimExp VName)-    -- ^ From the module PrimExps.-  , ctxVarExps :: M.Map VName Exp'-    -- ^ Statement-name-to-expression mappins for the entire function.-  , ctxAllocatedBlocksBeforeCreation :: M.Map VName MNames-    -- ^ Safety condition 2.-  , ctxVarsInUseBeforeMem :: M.Map MName Names-    -- ^ Safety condition 5.-  , ctxCurSnapshot :: Current-    -- ^ Keep a snapshot (used in 'tryCoalesce' for Concat).-  }-  deriving (Show)--newtype FindM lore a = FindM { unFindM :: RWS Context () Current a }-  deriving (Monad, Functor, Applicative,-            MonadReader Context,-            MonadState Current)--type LoreConstraints lore = (ExplicitMemorish lore,-                             FullWalk lore)--coerce :: FindM flore a -> FindM tlore a-coerce = FindM . unFindM--modifyCurCoalescedIntos :: (CoalescedIntos -> CoalescedIntos) -> FindM lore ()-modifyCurCoalescedIntos f =-  modify $ \c -> c { curCoalescedIntos = f $ curCoalescedIntos c }--modifyCurMemsCoalesced :: (MemsCoalesced -> MemsCoalesced) -> FindM lore ()-modifyCurMemsCoalesced f =-  modify $ \c -> c { curMemsCoalesced = f $ curMemsCoalesced c }--ifExp :: MonadReader Context m =>-         VName -> m (Maybe Exp')-ifExp var = do-  var_exp <- M.lookup var <$> asks ctxVarExps-  return $ case var_exp of-    Just e@(Exp _ _ If{}) -> Just e-    _ -> Nothing--isIfExp :: MonadReader Context m =>-           VName -> m Bool-isIfExp var = isJust <$> ifExp var--isLoopExp :: MonadReader Context m =>-             VName -> m Bool-isLoopExp var = do-  var_exp <- M.lookup var <$> asks ctxVarExps-  return $ case var_exp of-    Just (Exp _ _ DoLoop{}) -> True-    _ -> False--isReshapeExp :: MonadReader Context m =>-                VName -> m Bool-isReshapeExp var = do-  var_exp <- M.lookup var <$> asks ctxVarExps-  return $ case var_exp of-    Just (Exp _ _ (BasicOp Reshape{})) -> True-    _ -> False---- Lookup the memory block statically associated with a variable.-lookupVarMem :: MonadReader Context m =>-                VName -> m MemorySrc-lookupVarMem var =-  -- This should always be called from a place where it is certain that 'var'-  -- refers to a statement with an array expression.-  fromJust ("lookup memory block from " ++ pretty var) . M.lookup var-  <$> asks ctxVarToMem--lookupActualVars :: MonadReader Context m =>-                    VName -> m Names-lookupActualVars var = do-  actual_vars <- asks ctxActualVars-  -- Do this recursively.-  let actual_vars' = expandWithAliases actual_vars actual_vars-  return $ fromMaybe (S.singleton var) $ M.lookup var actual_vars'---- Lookup the memory block currenty associated with a variable.  In most cases--- (maybe all) this could probably replace 'lookupVarMem', though it would not--- always be necessary.-lookupCurrentVarMem :: VName -> FindM lore (Maybe VName)-lookupCurrentVarMem var = do-        -- Current result...-        mem_cur <- M.lookup var . curMemsCoalesced <$> asks ctxCurSnapshot-        -- ... or original result.-        ---        -- This is why we save the variables after creation, not the memory-        -- blocks: Variables stay the same, but memory blocks may change, which-        -- is relevant in the case of a chain of coalescings.-        mem_orig <- M.lookup var <$> asks ctxVarToMem-        return $ case (mem_cur, mem_orig) of-          (Just m, _) -> Just (memLocName m) -- priority choice-          (_, Just m) -> Just (memSrcName m)-          _ -> Nothing--withMemAliases :: MonadReader Context m =>-                  VName -> m Names-withMemAliases mem =-  -- The only memory blocks with memory aliases are the existiential ones, so-  -- using a static ctxMemAliases should be okay, as they will not change during-  -- the transformation in this module.-  S.union (S.singleton mem) . lookupEmptyable mem-  <$> asks ctxMemAliases--data Bindage = BindInPlace VName (Slice SubExp)-             | BindVar--recordOptimisticCoalescing :: VName -> PrimExp VName-                           -> [Slice (PrimExp VName)]-                           -> VName -> MemoryLoc -> Bindage -> FindM lore ()-recordOptimisticCoalescing src offset ixfun_slices dst dst_memloc bindage = do-  modifyCurCoalescedIntos $ insertOrUpdate dst (src, offset, ixfun_slices)--  -- If this is an in-place operation, we future-proof future coalescings by-  -- recording that they also need to take a look at the original array, not-  -- just the result of an in-place update into it.-  case bindage of-    BindVar -> return ()-    BindInPlace orig _ ->-      modifyCurCoalescedIntos $ insertOrUpdate dst (orig, zeroOffset, [])--  modifyCurMemsCoalesced $ M.insert src dst_memloc--coreCoalesceFunDef :: MonadFreshNames m =>-                      FunDef ExplicitMemory -> VarMemMappings MemorySrc-                   -> MemAliases -> VarAliases -> FirstUses -> LastUses-                   -> ActualVariables -> Names -> m (FunDef ExplicitMemory)-coreCoalesceFunDef fundef var_to_mem mem_aliases var_aliases first_uses-  last_uses actual_vars existentials = do-  let primexps = findPrimExpsFunDef fundef-      exps = findExpsFunDef fundef-      cond2 = findSafetyCondition2FunDef fundef-      cond5 = findSafetyCondition5FunDef fundef first_uses-      context = Context { ctxFunDef = fundef-                        , ctxVarToMem = var_to_mem-                        , ctxMemAliases = mem_aliases-                        , ctxVarAliases = var_aliases-                        , ctxFirstUses = first_uses-                        , ctxLastUses = last_uses-                        , ctxActualVars = actual_vars-                        , ctxExistentials = existentials-                        , ctxVarPrimExps = primexps-                        , ctxVarExps = exps-                        , ctxAllocatedBlocksBeforeCreation = cond2-                        , ctxVarsInUseBeforeMem = cond5-                        , ctxCurSnapshot = emptyCurrent-                        }-      m = unFindM $ lookInBody $ funDefBody fundef-      var_to_mem_res = curMemsCoalesced $ fst $ execRWS m context emptyCurrent-      sizes = memBlockSizesFunDef fundef-  transformFromVarMemMappings var_to_mem_res (M.map memSrcName var_to_mem) (M.map fst sizes) (M.map fst sizes) False fundef--lookInBody :: LoreConstraints lore =>-              Body lore -> FindM lore ()-lookInBody (Body _ bnds _res) =-  mapM_ lookInStm bnds--lookInKernelBody :: LoreConstraints lore =>-                    KernelBody lore -> FindM lore ()-lookInKernelBody (KernelBody _ bnds _res) =-  mapM_ lookInStm bnds--zeroOffset :: PrimExp VName-zeroOffset = primExpFromSubExp (IntType Int32) (constant (0 :: Int32))--lookInStm :: LoreConstraints lore =>-             Stm lore -> FindM lore ()-lookInStm (Let (Pattern _patctxelems patvalelems) _ e) = do-  -- COALESCING-SPECIFIC HANDLING for Copy and Concat.-  case patvalelems of-    [PatElem dst ExpMem.MemArray{}] -> do-      -- We create a function and pass it around instead of just applying it to-      -- the memory of the MemBound.  We do this, since any source variables-      -- might have more actual variables with different index functions that-      -- also need to be fixed -- e.g. in the case of reshape, where both the-      -- reshaped array and the original array need to get their index functions-      -- updated.-      ---      -- We take a snapshot of the current state of the curCoalescedIntos state-      -- field.  We need this feature to avoid having fewer coalescings just-      -- because of the placement of the sources.  For example, for-      ---      --     let b = ...-      --     let a = ...-      --     let c = concat a b-      ---      -- the coalescing pass will first coalesce m_a into m_c, which will-      -- succeed.  Then it will to coalesce m_b into m_c, which will (naively)-      -- fail because of safety condition 3 arguing that m_c is now in use after-      -- the creation of 'b' and before its use, since 'a' now uses m_c.-      ---      -- (Alternatively, we could do some more general index function analysis-      -- to check for things that will never overlap in merged memory, but this-      -- seems easier.)-      cur_snapshot <- get-      var_to_mem <- asks ctxVarToMem-      local (\ctx -> ctx { ctxCurSnapshot = cur_snapshot })-        $ case e of-            -- In-place update.-            BasicOp (Update orig slice (Var src)) ->-              case M.lookup src var_to_mem of-                Just _ ->-                  let ixfun_slices =-                        let slice' = map (primExpFromSubExp (IntType Int32) <$>) slice-                        in [slice']-                      bindage = BindInPlace orig slice-                  in tryCoalesce dst ixfun_slices bindage src zeroOffset-                Nothing ->-                  return ()--            -- Copy.-            BasicOp (Copy src) ->-              tryCoalesce dst [] BindVar src zeroOffset--            -- Concat.-            BasicOp (Concat 0 src0 src0s _) -> do-              let srcs = src0 : src0s-              shapes <- mapM ((memSrcShape <$>) . lookupVarMem) srcs-              let getOffsets offset_prev shape =-                    let se = head (shapeDims shape) -- Should work.-                        len = primExpFromSubExp (IntType Int32) se-                        offset_new = offset_prev + len-                    in offset_new-                  offsets = init (scanl getOffsets zeroOffset shapes)-              zipWithM_ (tryCoalesce dst [] BindVar) srcs offsets-            _ -> return ()-    _ -> return ()---  -- RECURSIVE BODY WALK.-  fullWalkExpM walker walker_kernel e-  where walker = identityWalker-          { walkOnBody = lookInBody }-        walker_kernel = identityKernelWalker-          { walkOnKernelBody = coerce . lookInBody-          , walkOnKernelKernelBody = coerce . lookInKernelBody-          , walkOnKernelLambda = coerce . lookInBody . lambdaBody-          }--tryCoalesce :: VName -> [Slice (PrimExp VName)] -> Bindage ->-               VName -> PrimExp VName -> FindM lore ()-tryCoalesce dst ixfun_slices bindage src offset = do-  mem_dst <- lookupVarMem dst--  -- For ifs and loops and some aliasing expressions (e.g. reshape), this tells-  -- us what non-existential source variables actually need to have assigned the-  -- new memory block.-  src's <- S.toList <$> lookupActualVars src--  -- From earlier optimistic coalescings.  Remember to also get the coalescings-  -- from the actual variables in e.g. loops.-  coalesced_intos <- curCoalescedIntos <$> asks ctxCurSnapshot-  let (src0s, offset0s, ixfun_slice0ss) =-        unzip3 $ S.toList $ S.unions-        $ map (`lookupEmptyable` coalesced_intos) (src : src's)--  var_to_pe <- asks ctxVarPrimExps--  let srcs = src's ++ src0s-                -- The same number of base offsets as in src's.-      offsets = replicate (length src's) offset-                -- The offsets of any previously optimistically coalesced src0s must be-                -- re-offset relative to the offset of the newest coalescing.-                ++ map (\o0 -> if o0 == zeroOffset && offset == zeroOffset-                                    -- This should not be necessary, and maybe it-                                    -- is not (but there were some problems).-                               then zeroOffset-                               else offset + o0) offset0s-      ixfun_slicess = replicate (length src's) ixfun_slices-                -- Same as above, kind of.-                ++ map (\slices0 -> ixfun_slices ++ slices0) ixfun_slice0ss--  let ixfuns' = zipWith (\offset_local islices ->-                           let ixfun0 = memSrcIxFun mem_dst-                               ixfun1 = foldl IxFun.slice ixfun0 islices--                               -- 'ixfun_slices' contain the slices that are the-                               -- result of a new coalescing, contrary to the-                               -- slices in 'ixfun_slice0ss' which contain-                               -- previously registered slices.-                               -- 'offsetIndexDWIM' handles the case that we-                               -- want to offset a DimFix if it is the result of-                               -- a previous coalescing, and not the current-                               -- one.  We do that by counting the number of-                               -- 'DimFix'es that originate in the new-                               -- coalescing, and then ignore those for our-                               -- heuristic.  This is a hack.-                               initial_dimfixes = L.takeWhile (isJust . dimFix) (concat ixfun_slices)-                               ixfun2 = if offset_local == zeroOffset-                                        then ixfun1 -- Should not be necessary,-                                                    -- but it makes the type-                                                    -- checker happy for now.-                                        else IxFun.offsetIndexDWIM (length initial_dimfixes) ixfun1 offset_local-                               ixfun3 = expandIxFun var_to_pe ixfun2-                           in ixfun3-                        ) offsets ixfun_slicess--  -- Not everything supported yet.  This dials back the optimisation on areas-  -- where it fails.-  existentials <- asks ctxExistentials-  let currentlyDisabled src_local = do-        -- This case covers the problem described in several programs in-        -- tests/coalescing/wip/loop/ (for programs where it is overly-        -- conservative) and tests/coalescing/loop/replicate-in-loop.fut (where-        -- it is absolutely needed to keep the program correct).  It is a-        -- conservative requirement and could likely be loosened up.--        src_local_is_loop <- isLoopExp src_local--        -- if the source contains the result a loop expression, and that result-        -- is an array with existential memory, don't coalesce.  Since memory-        -- can be allocated inside loops, coalescing with no further rules might-        -- end up having the same arrays use memory allocated outside the loop,-        -- which is not always okay.-        let res = src_local_is_loop-                  && src_local `L.elem` existentials-        return res--  safe0 <- not . or <$> mapM currentlyDisabled srcs--  -- Safety condition 1 is the same for all eventual previous arrays from srcs-  -- that also need to be coalesced into dst, so we check it here instead of-  -- checking it independently for every sub src.  This also ensures that we-  -- check that the destination memory is lastly used in *just* this statement,-  -- not also in any previous statement that uses the same memory block, which-  -- could very well fail.-  mem_src_base <- lookupVarMem src-  safe1 <- safetyCond1 dst mem_src_base--  when (safe0 && safe1) $ do-    safes <- zipWithM (canBeCoalesced dst) srcs ixfuns'-    when (and safes) $ do-      -- Any previous src0s coalescings must be deleted.-      modifyCurCoalescedIntos $ M.delete src-      -- The rest will be overwritten below.--      -- We then need to record that, from what we currently know, src and any-      -- nested src0s can all use the memory of dst with the new index functions.-      forM_ (L.zip4 srcs offsets ixfun_slicess ixfuns')-        $ \(src_local, offset_local, ixfun_slices_local, ixfun_local) -> do-        denotes_existential <- S.member src_local <$> asks ctxExistentials-        is_if <- isIfExp src_local-        dst_memloc <--          if denotes_existential && not is_if-          then do-            -- Only use the new index function.  Keep the existential memory-            -- block.  This means we have to make fewer changes to the program.-            ---            -- FIXME: However, if we are at an If expression with an existential-            -- memory block, we ignore it.  This is due to some special handling-            -- of If in MemoryUpdater, which is again due to branches having-            -- explicit return types.  This might not be correct.-            mem_src <- lookupVarMem src_local-            return $ MemoryLoc (memSrcName mem_src) ixfun_local-          else-            -- Use both the new memory block and the new index function.-            return $ MemoryLoc (memSrcName mem_dst) ixfun_local-        recordOptimisticCoalescing-          src_local offset_local ixfun_slices_local-          dst dst_memloc bindage--canBeCoalesced :: VName -> VName -> ExpMem.IxFun -> FindM lore Bool-canBeCoalesced dst src ixfun = do-  mem_dst <- lookupVarMem dst-  mem_src <- lookupVarMem src--  safe2 <- safetyCond2 src mem_dst-  safe3 <- safetyCond3 src dst mem_dst-  safe4 <- safetyCond4 src-  safe5 <- safetyCond5 mem_src ixfun--  safe_if <- safetyIf src dst--  let safe_all = safe2 && safe3 && safe4 && safe5 && safe_if-  return safe_all---- Safety conditions for each statement with a Copy or Concat:------ 1. mem_src is not used beyond the statement.  Handle by checking LastUses for---    the statement.------ 2. The allocation of mem_dst occurs before the creation of src, i.e. the---    first use of mem_src.  Handle by checking---    ctxAllocatedBlocksBeforeCreation.------ 3. There is no use or creation of mem_dst after the creation of src and---    before the current statement.  Handle by calling getVarUsesBetween and---    looking at both the original var-mem mappings *and* the new, temporary---    ones.------ 4. src (the variable, not the memory) does not alias anything.  Handle by---    checking VarAliases.------ 5. The new index function of src only uses variables declared prior to the---    first use of mem_src.  Handle by first using curVarPrimExps and---    ExpMem.substituteInIxFun to create a (possibly larger) index function that---    uses earlier variables.  Then use ctxVarsInUseBeforeMem to check that all---    the variables in the new index function are available before the creation---    of mem_src.------ If an array src0 has been coalesced into mem_src, handle that by *also*--- checking src0 and mem_src0 where src and mem_src are checked.  We choose to--- coalesce in a top-down fashion, even though that might exclude some potential--- coalescings -- however, doing it differently might exclude some other--- potentials, so we just make a choice.------ We only coalesce src into dst if all eventual src0 can also be coalesced into--- dst.  It does not make sense to coalesce only part of them, since in that--- case both memory blocks and related allocations will still be around.--safetyCond1 :: MonadReader Context m =>-               VName -> MemorySrc -> m Bool-safetyCond1 dst mem_src = do-  last_uses <- lookupEmptyable (FromStm dst) <$> asks ctxLastUses-  let res = S.member (memSrcName mem_src) last_uses-  return res--safetyCond2 :: MonadReader Context m =>-               VName -> MemorySrc -> m Bool-safetyCond2 src mem_dst = do-  allocs_before_src <- lookupEmptyable src-                       <$> asks ctxAllocatedBlocksBeforeCreation-  let res = S.member (memSrcName mem_dst) allocs_before_src-  return res--safetyCond3 :: VName -> VName -> MemorySrc -> FindM lore Bool-safetyCond3 src dst mem_dst = do-  fundef <- asks ctxFunDef-  let uses_after_src_vars = S.toList $ getVarUsesBetween fundef src dst-  uses_after_src <- mapM (maybe (return S.empty) withMemAliases-                          <=< lookupCurrentVarMem) uses_after_src_vars-  return $ not $ S.member (memSrcName mem_dst) (S.unions uses_after_src)--safetyCond4 :: MonadReader Context m =>-               VName -> m Bool-safetyCond4 src = do-  -- Special If handling: An If can have aliases, but that can be okay and is-  -- checked in safe If: It is okay for it to have one alias (one of the-  -- branches), while two aliases are wrong.-  if_handling <- isIfExp src--  -- Special Reshape handling: If a reshape has variables associated with it, it-  -- is okay to use it.-  src_actuals <- lookupEmptyable src <$> asks ctxActualVars-  reshape_handling <- isReshapeExp src <&&> pure (not (S.null src_actuals))--  -- This needs to be extended if support for e.g. reshape coalescing is wanted:-  -- Some operations can be aliasing, but still be okay to coalesce if you also-  -- coalesce their aliased sources.-  src_aliases <- lookupEmptyable src <$> asks ctxVarAliases-  let res = if_handling || reshape_handling || S.null src_aliases-  return res--safetyCond5 :: MonadReader Context m =>-               MemorySrc -> ExpMem.IxFun -> m Bool-safetyCond5 mem_src ixfun = do-  in_use_before_mem_src <- lookupEmptyable (memSrcName mem_src)-                           <$> asks ctxVarsInUseBeforeMem-  let used_vars = freeIn ixfun-      res = all (`S.member` in_use_before_mem_src) $ S.toList used_vars-  return res--safetyIf :: VName -> VName -> FindM lore Bool-safetyIf src dst = do-  -- Special handling: If src refers to an If expression, we need to check that-  -- not just is mem_dst not used after src and before dst, but neither is any-  -- other memory that will be merged after the coalescing.  Normally this is-  -- not an issue, since a coalescing means changing just one memory block ---  -- but in the case of an If expression, each branch can have its own memory-  -- block, and both of them will try to be coalesced.  This extra test only-  -- applies to the actual memory blocks in the branches, not any existential-  -- memory block in the If, which in any case will be "used" in both branches.-  ---  -- See tests/coalescing/if/if-neg-3.fut for an example of where this should-  -- fail.-  mem_src <- lookupVarMem src-  actual_srcs <- S.toList <$> lookupActualVars src-  existentials <- asks ctxExistentials-  var_to_mem <- asks ctxVarToMem-  first_uses_all <- asks ctxFirstUses--  -- Find all variables that have 'src' as an actual var, and then check if one-  -- of those is an If expression.-  reverse_actual_srcs <--    S.toList . S.unions . M.elems . M.filter (src `S.member`)-    <$> asks ctxActualVars-  outer <- mapMaybeM ifExp reverse_actual_srcs-  let (is_in_if,-       if_branch_results_from_outer,-       at_least_one_creation_inside) = case outer of-        -- This is the if expression of which we are currently looking at one of-        -- its branch results.-        [Exp nctx nthpat (If _ body0 body1 _)] ->-          let results_from_outer = S.fromList $ mapMaybe subExpVar-                                   $ concatMap (drop nctx . bodyResult)-                                   $ filter (null . bodyStms) [body0, body1]--              resultCreatedInside body se = fromMaybe False $ do-                res <- subExpVar se-                res_mem <- memSrcName <$> M.lookup res var_to_mem-                let body_vars = concatMap (map patElemName . patternValueElements-                                           . stmPattern) $ bodyStms body-                    body_first_uses = S.unions $ map (`lookupEmptyable` first_uses_all)-                                      body_vars-                return $ S.member res_mem body_first_uses--              at_least = resultCreatedInside body0 (bodyResult body0 !! (nctx + nthpat))-                         || resultCreatedInside body1 (bodyResult body1 !! (nctx + nthpat))-          in (True, results_from_outer, at_least)-        _ -> (False, S.empty, False)--  -- This success requirement is independent of whichever branch we are in right-  -- now.  We say that the results of an if-expression can be coalesced if the-  -- branch-specific requirements hold *and* this general rule holds: Either the-  -- If has no existentials (e.g. if it does in-place updates), or it has-  -- existentials and at least one of the branches returns an array that was-  -- created inside the branch.-  let res_general = not is_in_if || (not (any (`S.member` existentials) actual_srcs)-                                     || at_least_one_creation_inside)--  -- Check if the branch described by 'src' needs special handling.-  let if_handling =-        -- We are sure this is an if.  This might not actually be necessary.-        is_in_if-        -- This does not refer to the result of a branch where the array is-        -- created outside the if.  It is a requirement that there is at most-        -- one such branch.  The extra safety here only relates to branches-        -- whose result arrays are created inside.-        && not (any (`S.member` if_branch_results_from_outer) actual_srcs)-        -- Ignore existentials as well.-        && not (src `S.member` existentials)--  -- This success requirement is part is specific to this branch.-  res_current <--    if if_handling-    then do-      -- Get the memory used in the other branch.  Use a reverse lookup.-      mem_actual_srcs <- L.nub <$> mapM lookupVarMem reverse_actual_srcs-      let mem_actual_srcs_cur = L.delete mem_src mem_actual_srcs-      and <$> mapM (safetyCond3 src dst) mem_actual_srcs_cur-    else return True--  -- The full result.-  let res = res_general && res_current-  return res

@@ -1,70 +0,0 @@-{-# LANGUAGE GeneralizedNewtypeDeriving #-}-{-# LANGUAGE FlexibleContexts #-}-{-# LANGUAGE TypeFamilies #-}-{-# LANGUAGE ConstraintKinds #-}-{-# LANGUAGE ExistentialQuantification #-}--- | Get a mapping from statement patterns to statement expression for all--- statements.-module Futhark.Optimise.MemoryBlockMerging.Coalescing.Exps-  ( Exp'(..)-  , findExpsFunDef-  ) where--import qualified Data.Map.Strict as M-import Control.Monad-import Control.Monad.Writer--import Futhark.Representation.AST-import Futhark.Representation.ExplicitMemory (ExplicitMemorish)-import Futhark.Representation.Kernels.Kernel--import Futhark.Optimise.MemoryBlockMerging.Miscellaneous---- | Describes the nth pattern and the statement expression.-data Exp' = forall lore. Annotations lore => Exp Int Int (Exp lore)-instance Show Exp' where-  show (Exp _nctxpatters _nthvalpattern e) = show e--type Exps = M.Map VName Exp'--newtype FindM lore a = FindM { unFindM :: Writer Exps a }-  deriving (Monad, Functor, Applicative,-            MonadWriter Exps)--type LoreConstraints lore = (ExplicitMemorish lore,-                             FullWalk lore)--coerce :: FindM flore a -> FindM tlore a-coerce = FindM . unFindM--findExpsFunDef :: LoreConstraints lore =>-                  FunDef lore -> Exps-findExpsFunDef fundef =-  let m = unFindM $ lookInBody $ funDefBody fundef-      res = execWriter m-  in res--lookInBody :: LoreConstraints lore =>-              Body lore -> FindM lore ()-lookInBody (Body _ bnds _res) =-  mapM_ lookInStm bnds--lookInKernelBody :: LoreConstraints lore =>-                    KernelBody lore -> FindM lore ()-lookInKernelBody (KernelBody _ bnds _res) =-  mapM_ lookInStm bnds--lookInStm :: LoreConstraints lore =>-             Stm lore -> FindM lore ()-lookInStm (Let (Pattern patctxelems patvalelems) _ e) = do-  forM_ (zip patvalelems [0..]) $ \(PatElem var _, i) ->-    tell $ M.singleton var $ Exp (length patctxelems) i e--  -- Recursive body walk.-  fullWalkExpM walker walker_kernel e-  where walker = identityWalker-          { walkOnBody = lookInBody }-        walker_kernel = identityKernelWalker-          { walkOnKernelBody = coerce . lookInBody-          , walkOnKernelKernelBody = coerce . lookInKernelBody-          }

@@ -1,110 +0,0 @@-{-# LANGUAGE GeneralizedNewtypeDeriving #-}-{-# LANGUAGE FlexibleContexts #-}-{-# LANGUAGE TypeFamilies #-}-{-# LANGUAGE ConstraintKinds #-}--- | Find safety condition 2 for all statements.-module Futhark.Optimise.MemoryBlockMerging.Coalescing.SafetyCondition2-  ( findSafetyCondition2FunDef-  ) where--import qualified Data.Map.Strict as M-import qualified Data.Set as S-import Control.Monad-import Control.Monad.RWS--import Futhark.Representation.AST-import Futhark.Representation.ExplicitMemory (-  ExplicitMemory, InKernel, ExplicitMemorish)-import qualified Futhark.Representation.ExplicitMemory as ExpMem-import Futhark.Representation.Kernels.Kernel--import Futhark.Optimise.MemoryBlockMerging.Types-import Futhark.Optimise.MemoryBlockMerging.Miscellaneous---type CurrentAllocatedBlocks = MNames-type AllocatedBlocksBeforeCreation = M.Map VName MNames--newtype FindM lore a = FindM { unFindM :: RWS ()-                               AllocatedBlocksBeforeCreation CurrentAllocatedBlocks a }-  deriving (Monad, Functor, Applicative,-            MonadWriter AllocatedBlocksBeforeCreation,-            MonadState CurrentAllocatedBlocks)--type LoreConstraints lore = (ExplicitMemorish lore,-                             IsAlloc lore,-                             FullWalk lore)--coerce :: FindM flore a -> FindM tlore a-coerce = FindM . unFindM--findSafetyCondition2FunDef :: FunDef ExplicitMemory-                           -> AllocatedBlocksBeforeCreation-findSafetyCondition2FunDef fundef =-  let m = unFindM $ do-        forM_ (funDefParams fundef) lookInFParam-        lookInBody $ funDefBody fundef-      res = snd $ evalRWS m () S.empty-  in res--lookInFParam :: FParam ExplicitMemory -> FindM lore ()-lookInFParam (Param _ membound) =-  -- Unique array function parameters also count as "allocations" in which-  -- memory can be coalesced.-  case membound of-    ExpMem.MemArray _ _ Unique (ExpMem.ArrayIn mem _) ->-      modify $ S.insert mem-    _ -> return ()--lookInBody :: LoreConstraints lore =>-              Body lore -> FindM lore ()-lookInBody (Body _ bnds _res) =-  mapM_ lookInStm bnds--lookInKernelBody :: LoreConstraints lore =>-                    KernelBody lore -> FindM lore ()-lookInKernelBody (KernelBody _ bnds _res) =-  mapM_ lookInStm bnds--lookInStm :: LoreConstraints lore =>-             Stm lore -> FindM lore ()-lookInStm (Let (Pattern patctxelems patvalelems) _ e) = do-  let new_decls0 = map patElemName (patctxelems ++ patvalelems)-      new_decls1 = case e of-        DoLoop _mergectxparams mergevalparams _loopform _body ->-          -- Technically not a declaration for the current expression, but very-          -- close, and hopefully okay to consider it as one.-          map (paramName . fst) mergevalparams-        _ -> []-      new_decls = new_decls0 ++ new_decls1--  cur_allocated_blocks <- get-  forM_ new_decls $ \x ->-    tell $ M.singleton x cur_allocated_blocks--  case patvalelems of-    [PatElem mem _] ->-      when (isAlloc e) $ modify $ S.insert mem-    _ -> return ()--  -- RECURSIVE BODY WALK.-  fullWalkExpM walker walker_kernel e-  where walker = identityWalker-          { walkOnBody = lookInBody-          , walkOnFParam = lookInFParam-          }-        walker_kernel = identityKernelWalker-          { walkOnKernelBody = coerce . lookInBody-          , walkOnKernelKernelBody = coerce . lookInKernelBody-          , walkOnKernelLambda = coerce . lookInBody . lambdaBody-          }--class IsAlloc lore where-  isAlloc :: Exp lore -> Bool--instance IsAlloc ExplicitMemory where-  isAlloc (Op ExpMem.Alloc{}) = True-  isAlloc _ = False--instance IsAlloc InKernel where-  isAlloc _ = False

@@ -1,136 +0,0 @@-{-# LANGUAGE GeneralizedNewtypeDeriving #-}-{-# LANGUAGE TypeFamilies #-}-{-# LANGUAGE FlexibleContexts #-}-{-# LANGUAGE ConstraintKinds #-}--- | Safety condition 3 verification.-module Futhark.Optimise.MemoryBlockMerging.Coalescing.SafetyCondition3-  ( getVarUsesBetween-  ) where--import qualified Data.Set as S-import qualified Data.List as L-import Control.Monad-import Control.Monad.RWS--import Futhark.Representation.AST-import Futhark.Representation.ExplicitMemory (-  ExplicitMemory, ExplicitMemorish)-import Futhark.Representation.Kernels.Kernel--import Futhark.Optimise.MemoryBlockMerging.Miscellaneous---data Context = Context-  { ctxSource :: VName-  , ctxDestination :: VName-  }-  deriving (Show)--data Current = Current-  { curHasReachedSource :: Bool-  , curHasReachedDestination :: Bool-  , curVars :: Names-  }-  deriving (Show)--newtype FindM lore a = FindM { unFindM :: RWS Context () Current a }-  deriving (Monad, Functor, Applicative,-            MonadReader Context,-            MonadState Current)--type LoreConstraints lore = (ExplicitMemorish lore,-                             FullWalk lore)--coerce :: FindM flore a -> FindM tlore a-coerce = FindM . unFindM--modifyCurVars :: (Names -> Names) -> FindM lore ()-modifyCurVars f = modify $ \c -> c { curVars = f $ curVars c }---- Find all the variables present between the creations of two variables (not--- inclusive).-getVarUsesBetween :: FunDef ExplicitMemory-                  -> VName -> VName-                  -> Names-getVarUsesBetween fundef src dst =-  let context = Context src dst-      m = unFindM $ lookInBody $ funDefBody fundef-      res = curVars $ fst $ execRWS m context (Current False False S.empty)-  in res--lookInBody :: LoreConstraints lore =>-              Body lore -> FindM lore ()-lookInBody (Body _ bnds _res) =-  mapM_ lookInStm bnds--lookInKernelBody :: LoreConstraints lore =>-                    KernelBody lore -> FindM lore ()-lookInKernelBody (KernelBody _ bnds _res) =-  mapM_ lookInStm bnds--lookInStm :: LoreConstraints lore =>-             Stm lore -> FindM lore ()-lookInStm stm@(Let _ _ e) = do-  let new_decls = newDeclarationsStm stm--  dst <- asks ctxDestination-  when (dst `L.elem` new_decls)-    $ modify $ \c -> c { curHasReachedDestination = True }--  is_after_source <- gets curHasReachedSource-  is_before_destination <- gets curHasReachedDestination--  unless is_before_destination $ do-    let e_free_vars = freeInExp e-        e_used_vars = S.union e_free_vars (S.fromList new_decls)--    -- If the source has been created, add the newly used variables.-    ---    -- Note that "used after creation" refers both to used in subsequent-    -- statements AND any statements in any sub-bodies (if and loop).-    when is_after_source-      $ modifyCurVars $ S.union e_used_vars--    -- If the source is present in the declarations, state that it has been-    -- created.-    src <- asks ctxSource-    when (src `L.elem` new_decls)-      $ modify $ \c -> c { curHasReachedSource = True  }--    -- RECURSIVE BODY WALK.-    case e of-      If _ body0 body1 _ -> do-        -- This is not very If-specific, but rather specific to expressions with-        -- multiple, independent bodies, where If is just the only such-        -- expression.-        ---        -- We do not want the state (for safety condition 3) after traversing-        -- the first branch to be present when traversing the second branch,-        -- since they really will never both be run, so we compute them-        -- independently and then merge them at the end.-        before <- get-        lookInBody body0-        after0 <- get-        put Current { curHasReachedSource = curHasReachedSource before-                    , curHasReachedDestination = curHasReachedDestination after0-                    , curVars = curVars before-                    }-        lookInBody body1-        after1 <- get-        put Current { curHasReachedSource =-                      curHasReachedSource after0 || curHasReachedSource after1-                    , curHasReachedDestination =-                      curHasReachedDestination after0 || curHasReachedDestination after1-                    , curVars =-                      S.union (curVars after0) (curVars after1)-                    }-      _ -> do-        -- In the general case, just look through any 'Body' you can find.  (This-        -- is the case for loops.)-        let walker = identityWalker { walkOnBody = lookInBody }-            walker_kernel = identityKernelWalker-              { walkOnKernelBody = coerce . lookInBody-              , walkOnKernelKernelBody = coerce . lookInKernelBody-              , walkOnKernelLambda = coerce . lookInBody . lambdaBody-              }-        fullWalkExpM walker walker_kernel e

@@ -1,120 +0,0 @@-{-# LANGUAGE GeneralizedNewtypeDeriving #-}-{-# LANGUAGE FlexibleContexts #-}-{-# LANGUAGE TypeFamilies #-}-{-# LANGUAGE ConstraintKinds #-}--- | Find safety condition 5 for all statements.-module Futhark.Optimise.MemoryBlockMerging.Coalescing.SafetyCondition5-  ( findSafetyCondition5FunDef-  ) where--import qualified Data.Map.Strict as M-import qualified Data.Set as S-import Control.Monad-import Control.Monad.RWS--import Futhark.Representation.AST-import Futhark.Representation.ExplicitMemory (-  InKernel, ExplicitMemory, ExplicitMemorish)-import qualified Futhark.Representation.ExplicitMemory as ExpMem-import Futhark.Representation.Kernels.Kernel--import Futhark.Optimise.MemoryBlockMerging.Types-import Futhark.Optimise.MemoryBlockMerging.Miscellaneous---type DeclarationsSoFar = Names-type VarsInUseBeforeMem = M.Map MName Names--newtype FindM lore a = FindM { unFindM :: RWS FirstUses-                               VarsInUseBeforeMem DeclarationsSoFar a }-  deriving (Monad, Functor, Applicative,-            MonadReader FirstUses,-            MonadWriter VarsInUseBeforeMem,-            MonadState DeclarationsSoFar)--type LoreConstraints lore = (ExplicitMemorish lore,-                             ExtractKernelDefVars lore,-                             FullWalk lore)--coerce :: FindM flore a -> FindM tlore a-coerce = FindM . unFindM--findSafetyCondition5FunDef :: FunDef ExplicitMemory -> FirstUses-                           -> VarsInUseBeforeMem-findSafetyCondition5FunDef fundef first_uses =-  let m = unFindM $ do-        forM_ (funDefParams fundef) lookInFParam-        lookInBody $ funDefBody fundef-      res = snd $ evalRWS m first_uses S.empty-  in res--lookInFParam :: FParam lore -> FindM lore ()-lookInFParam (Param x _) =-  modify $ S.insert x--lookInLParam :: LParam lore -> FindM lore ()-lookInLParam (Param x _) =-  modify $ S.insert x--lookInBody :: LoreConstraints lore => Body lore -> FindM lore ()-lookInBody (Body _ bnds _res) =-  mapM_ lookInStm bnds--lookInKernelBody :: LoreConstraints lore => KernelBody lore -> FindM lore ()-lookInKernelBody (KernelBody _ bnds _res) =-  mapM_ lookInStm bnds--lookInStm :: LoreConstraints lore => Stm lore -> FindM lore ()-lookInStm stm@(Let _ _ e) = do-  let new_decls = newDeclarationsStm stm--  first_uses <- ask-  declarations_so_far <- get-  forM_ (S.toList $ S.unions $ map (`lookupEmptyable` first_uses) new_decls) $ \mem ->-    tell $ M.singleton mem declarations_so_far--  forM_ new_decls $ \x ->-    modify $ S.insert x--  -- Special loop handling: Extract useful variables that are in use.-  case e of-    DoLoop _ _ loopform _ ->-      case loopform of-        ForLoop i _ _ _ -> modify $ S.insert i-        WhileLoop c -> modify $ S.insert c-    _ -> return ()--  modify $ S.union (extractKernelDefVars e)--  -- Recursive body walk.-  fullWalkExpM walker walker_kernel e-  where walker = identityWalker-          { walkOnBody = lookInBody-          , walkOnFParam = lookInFParam-          , walkOnLParam = lookInLParam-          }-        walker_kernel = identityKernelWalker-          { walkOnKernelBody = coerce . lookInBody-          , walkOnKernelKernelBody = coerce . lookInKernelBody-          , walkOnKernelLambda = coerce . lookInLambda-          , walkOnKernelLParam = lookInLParam-          }--lookInLambda :: LoreConstraints lore =>-                Lambda lore -> FindM lore ()-lookInLambda (Lambda params body _) = do-  forM_ params lookInLParam-  lookInBody body--class ExtractKernelDefVars lore where-  -- Extract variables from a kernel definition.-  extractKernelDefVars :: Exp lore -> Names--instance ExtractKernelDefVars ExplicitMemory where-  extractKernelDefVars (Op (ExpMem.Inner (Kernel _ kernelspace _ _))) =-    S.fromList $ map ($ kernelspace)-    [spaceGlobalId, spaceLocalId, spaceGroupId]-  extractKernelDefVars _ = S.empty--instance ExtractKernelDefVars InKernel where-  extractKernelDefVars _ = S.empty

@@ -1,263 +0,0 @@--- | Move variables as much as possible upwards in a program.-module Futhark.Optimise.MemoryBlockMerging.CrudeMovingUp-  ( moveUpInFunDef-  ) where--import qualified Data.Set as S-import qualified Data.List as L-import qualified Data.Map.Strict as M-import Data.Maybe (mapMaybe)-import Control.Monad-import Control.Monad.RWS-import Control.Monad.Writer--import Futhark.Representation.AST-import Futhark.Representation.ExplicitMemory (ExplicitMemory)-import qualified Futhark.Representation.ExplicitMemory as ExpMem--import Futhark.Optimise.MemoryBlockMerging.Miscellaneous--import Control.Monad.State-import Control.Monad.Identity---type Line = Int-data Origin = FromFParam-            | FromLine Line (Exp ExplicitMemory)-  deriving (Eq, Ord, Show)---- The dependencies and the location.-data PrimBinding = PrimBinding { pbFrees :: Names-                               , _pbConsumed :: Names-                               , pbOrigin :: Origin-                               }-  deriving (Show)---- A mapping from names to PrimBinding.  The key is a collection of names, since--- a statement can have multiple patterns.-type BindingMap = [(Names, PrimBinding)]---- | Call 'findHoistees' for every body, and then hoist every one of the found--- hoistees (variables).-moveUpInFunDef :: FunDef ExplicitMemory-               -> (Body ExplicitMemory -> Maybe [FParam ExplicitMemory] -> [VName])-               -> FunDef ExplicitMemory-moveUpInFunDef fundef findHoistees =-  let scope_new = scopeOf fundef-      bindingmap_cur = []-      body' = hoistInBody scope_new bindingmap_cur-              (Just (funDefParams fundef)) findHoistees (funDefBody fundef)-      fundef' = fundef { funDefBody = body' }-  in fundef'--lookupPrimBinding :: VName -> State BindingMap PrimBinding-lookupPrimBinding vname =-  gets $ snd . fromJust (pretty vname ++ " was not found in BindingMap."-                         ++ "  This should not happen!")-  . L.find ((vname `S.member`) . fst)--namesDependingOn :: VName -> State BindingMap Names-namesDependingOn v =-  gets $ S.unions . map fst . filter (\(_, pb) -> v `S.member` pbFrees pb)--scopeBindingMap :: (VName, NameInfo ExplicitMemory)-                -> BindingMap-scopeBindingMap (x, _) = [(S.singleton x, PrimBinding S.empty S.empty FromFParam)]---- Find all variables bound in a KernelSpace.-boundInKernelSpace :: ExpMem.KernelSpace -> Names-boundInKernelSpace space =-  -- This might do too much.-  S.fromList ([ ExpMem.spaceGlobalId space-              , ExpMem.spaceLocalId space-              , ExpMem.spaceGroupId space]-              ++ (case ExpMem.spaceStructure space of-                    ExpMem.FlatThreadSpace ts ->-                      map fst ts ++ mapMaybe (subExpVar . snd) ts-                    ExpMem.NestedThreadSpace ts ->-                      map (\(x, _, _, _) -> x) ts-                      ++ mapMaybe (subExpVar . (\(_, x, _, _) -> x)) ts-                      ++ map (\(_, _, x, _) -> x) ts-                      ++ mapMaybe (subExpVar . (\(_, _, _, x) -> x)) ts-                 ))---- FIXME: The results of this should maybe go in the core 'freeIn' function, or--- perhaps the ExplicitMemory module, instead of this arbitrary module.-boundInExpExtra :: Exp ExplicitMemory -> Names-boundInExpExtra = execWriter . inExp-  where inExp :: Exp ExplicitMemory -> Writer Names ()-        inExp e = case e of-          Op (ExpMem.Inner (ExpMem.Kernel _ space _ _)) ->-            tell $ boundInKernelSpace space-          _ -> walkExpM walker e--        walker = identityWalker {-          walkOnBody = mapM_ (inExp . stmExp) . bodyStms-          }--bodyBindingMap :: [Stm ExplicitMemory] -> BindingMap-bodyBindingMap stms =-  concatMap createBindingStmt $ zip [0..] stms-  -- We do not need to run this recursively on any sub-bodies, since this will-  -- be run for every call to hoistInBody, which *does* run recursively on-  -- sub-bodies.--  where createBindingStmt :: (Line, Stm ExplicitMemory)-                          -> BindingMap-        createBindingStmt (line, stmt@(Let (Pattern patctxelems patvalelems) _ e)) =-          let stmt_vars = S.fromList (map patElemName (patctxelems ++ patvalelems))-              frees = freeInStm stmt-              consumed = case e of BasicOp (Update src _ _) -> S.singleton src-                                   _ -> mempty-              bound_extra = boundInExpExtra e-              frees' = frees `S.difference` bound_extra-              vars_binding = (stmt_vars, PrimBinding frees' consumed (FromLine line e))--              -- Some variables exist only in a shape declaration.-              shape_sizes = S.fromList $ concatMap shapeSizes (patctxelems ++ patvalelems)-              sizes_binding = (shape_sizes, PrimBinding frees' consumed (FromLine line e))--              -- Some expressions contain special identifiers that are used in a-              -- body.  This should go somewhere else than here.-              param_vars = case e of-                Op (ExpMem.Inner (ExpMem.Kernel _ space _ _)) ->-                  boundInKernelSpace space-                _ -> S.empty-              params_binding = (param_vars, PrimBinding S.empty S.empty FromFParam)--              bmap = [vars_binding, sizes_binding, params_binding]-          in bmap--        shapeSizes (PatElem _ (ExpMem.MemArray _ shape _ _)) =-          mapMaybe subExpVar $ shapeDims shape-        shapeSizes _ = []--hoistInBody :: Scope ExplicitMemory-            -> BindingMap-            -> Maybe [FParam ExplicitMemory]-            -> (Body ExplicitMemory -> Maybe [FParam ExplicitMemory] -> [VName])-            -> Body ExplicitMemory-            -> Body ExplicitMemory-hoistInBody scope_new bindingmap_old params findHoistees body =-  let hoistees = findHoistees body params--      -- We use the possibly non-empty scope to extend our BindingMap.-      bindingmap_fromscope = concatMap scopeBindingMap $ M.toList scope_new-      bindingmap_body = bodyBindingMap $ stmsToList $ bodyStms body-      bindingmap = bindingmap_old ++ bindingmap_fromscope ++ bindingmap_body--      -- Create a new body where all hoistees have been moved as much upwards in-      -- the statement list as possible.-      (Body () bnds res, bindingmap') =-        foldl (\(body0, lbindingmap) -> hoist lbindingmap body0)-        (body, bindingmap) hoistees--      -- Touch upon any subbodies.-      bnds' = fmap (hoistRecursivelyStm bindingmap' findHoistees) bnds-      body' = Body () bnds' res--  in body'--hoistRecursivelyStm :: BindingMap-                    -> (Body ExplicitMemory -> Maybe [FParam ExplicitMemory] -> [VName])-                    -> Stm ExplicitMemory-                    -> Stm ExplicitMemory-hoistRecursivelyStm bindingmap findHoistees (Let pat aux e) =-  runIdentity (Let pat aux <$> mapExpM transform e)--  where transform = identityMapper { mapOnBody = mapper }-        mapper scope_new = return . hoistInBody scope_new bindingmap' Nothing findHoistees-        -- The nested body cannot move to any of its locations of its parent's-        -- body, so we say that all its parent's bindings are parameters.-        bindingmap' = map (\(ns, PrimBinding frees consumed _) ->-                             (ns, PrimBinding frees consumed FromFParam))-                      bindingmap---- Hoist the statement denoted by 'hoistee' as much upwards as possible in--- 'body', and return the new body.-hoist :: BindingMap-      -> Body ExplicitMemory-      -> VName-      -> (Body ExplicitMemory, BindingMap)-hoist bindingmap_cur body hoistee =-  let bindingmap = bindingmap_cur <> bodyBindingMap (stmsToList $ bodyStms body)--      body' = runState (moveLetUpwards hoistee body) bindingmap--  in body'---- Move a statement as much up as possible.-moveLetUpwards :: VName -> Body ExplicitMemory-               -> State BindingMap (Body ExplicitMemory)-moveLetUpwards letname body = do-  PrimBinding deps consumed letorig <- lookupPrimBinding letname--  -- Extend the dependencies with all those statements that use the consumed-  -- variables of this statement, except the current statement.-  deps' <- S.delete letname-           <$> (S.union deps-                <$> (S.unions <$> mapM namesDependingOn (S.toList consumed)))--  case letorig of-    FromFParam -> return body-    FromLine line_cur exp_cur ->-      case exp_cur of-        -- We do not want to change the structure of the program too much, so we-        -- restrict the aggressive hoister to *stop* and not hoist loops and-        -- kernels, as hoisting these expressions might actually make a-        -- hoisting-dependent optimisation *poorer* because of some assumptions-        -- about the structure.  FIXME: Do this nicer in a way where it is easy-        -- to argue for it.-        DoLoop{} -> return body-        Op ExpMem.Inner{} -> return body-        _ -> do-          -- Sort by how close they are to the beginning of the body.  The closest-          -- one should be the first one to hoist, so that the other ones can maybe-          -- exploit it.-          deps'' <- sortByKeyM (fmap pbOrigin . lookupPrimBinding)-                    $ S.toList deps'-          body' <- foldM (flip moveLetUpwards) body deps''-          origins <- mapM (fmap pbOrigin . lookupPrimBinding) deps''-          let line_dest = case foldl max FromFParam origins of-                FromFParam -> 0-                FromLine n _e -> n + 1--          PrimBinding _ _ letorig' <- lookupPrimBinding letname-          when (letorig' /= letorig) $ error "Assertion: This should not happen."--          stms' <- moveLetToLine letname line_cur line_dest $ stmsToList $ bodyStms body'--          return body' { bodyStms = stmsFromList stms' }---- Both move the statement to the new line and update the BindingMap.-moveLetToLine :: VName -> Line -> Line -> [Stm ExplicitMemory]-              -> State BindingMap [Stm ExplicitMemory]-moveLetToLine stm_cur_name line_cur line_dest stms-  | line_cur == line_dest = return stms-  | otherwise = do--  let stm_cur = stms !! line_cur-      stms1 = take line_cur stms ++ drop (line_cur + 1) stms-      stms2 = take line_dest stms1 ++ [stm_cur] ++ drop line_dest stms1--  modify $ map (\t@(ns, PrimBinding frees consumed orig) ->-                   case orig of-                     FromFParam -> t-                     FromLine l e -> if l >= line_dest && l < line_cur-                                     then (ns, PrimBinding frees consumed-                                               (FromLine (l + 1) e))-                                     else t)--  r <- lookupPrimBinding stm_cur_name-  case r of-    PrimBinding frees consumed (FromLine _ exp_cur) ->-      modify $ replaceWhere stm_cur_name (PrimBinding frees consumed-                                          (FromLine line_dest exp_cur))-    _ -> error "moveLetToLine: unhandled case" -- fixme-  return stms2--replaceWhere :: VName -> PrimBinding -> BindingMap -> BindingMap-replaceWhere n pb1 =-  map (\(ns, pb) -> (ns, if n `S.member` ns-                         then pb1-                         else pb))

@@ -1,80 +0,0 @@-{-# LANGUAGE GeneralizedNewtypeDeriving #-}-{-# LANGUAGE TypeFamilies #-}-{-# LANGUAGE FlexibleContexts #-}-{-# LANGUAGE ConstraintKinds #-}--- | Find all existential variables.-module Futhark.Optimise.MemoryBlockMerging.Existentials-  ( findExistentials-  ) where--import qualified Data.Set as S-import qualified Data.List as L-import Control.Monad-import Control.Monad.Writer--import Futhark.Representation.AST-import Futhark.Representation.ExplicitMemory (ExplicitMemorish)-import qualified Futhark.Representation.ExplicitMemory as ExpMem-import Futhark.Representation.Kernels.Kernel--import Futhark.Optimise.MemoryBlockMerging.Miscellaneous---newtype FindM lore a = FindM { unFindM :: Writer Names a }-  deriving (Monad, Functor, Applicative,-            MonadWriter Names)--type LoreConstraints lore = (ExplicitMemorish lore,-                             FullWalk lore)--record :: VName -> FindM lore ()-record = tell . S.singleton--coerce :: FindM flore a -> FindM tlore a-coerce = FindM . unFindM--findExistentials :: LoreConstraints lore =>-                    FunDef lore -> Names-findExistentials fundef =-  let m = unFindM $ lookInBody $ funDefBody fundef-      existentials = execWriter m-  in existentials--lookInBody :: LoreConstraints lore =>-              Body lore -> FindM lore ()-lookInBody (Body _ bnds _res) =-  mapM_ lookInStm bnds--lookInKernelBody :: LoreConstraints lore =>-                    KernelBody lore -> FindM lore ()-lookInKernelBody (KernelBody _ bnds _res) =-  mapM_ lookInStm bnds--lookInStm :: LoreConstraints lore =>-             Stm lore -> FindM lore ()-lookInStm (Let (Pattern patctxelems patvalelems) _ e) = do-  forM_ patvalelems $ \(PatElem var membound) ->-    case membound of-      ExpMem.MemArray _ _ _ (ExpMem.ArrayIn mem _) ->-        when (mem `L.elem` map patElemName patctxelems)-        $ record var-      _ -> return ()--  case e of-    DoLoop mergectxparams mergevalparams _loopform _body ->-      forM_ mergevalparams $ \(Param var membound, _) ->-        case membound of-          ExpMem.MemArray _ _ _ (ExpMem.ArrayIn mem _) ->-            when (mem `L.elem` map (paramName . fst) mergectxparams)-            $ record var-          _ -> return ()-    _ -> return ()--  fullWalkExpM walker walker_kernel e-  where walker = identityWalker-          { walkOnBody = lookInBody }-        walker_kernel = identityKernelWalker-          { walkOnKernelBody = coerce . lookInBody-          , walkOnKernelKernelBody = coerce . lookInKernelBody-          , walkOnKernelLambda = coerce . lookInBody . lambdaBody-          }