futhark 0.22.7 → 0.23.1

raw patch · 174 files changed

+5936/−6087 lines, 174 filesdep −zip-archivedep ~futhark-manifestPVP ok

version bump matches the API change (PVP)

Dependencies removed: zip-archive

Dependency ranges changed: futhark-manifest

API changes (from Hackage documentation)

Files

• docs/c-api.rst +6/−4
• docs/error-index.rst +26/−0
• docs/language-reference.rst +23/−9
• docs/man/futhark-dataset.rst +1/−1
• docs/man/futhark-literate.rst +5/−0
• docs/man/futhark-pkg.rst +4/−6
• docs/man/futhark-test.rst +7/−1
• docs/man/futhark.rst +13/−0
• docs/package-management.rst +25/−3
• docs/performance.rst +2/−1
• futhark.cabal +10/−8
• prelude/math.fut +5/−0
• prelude/soacs.fut +1/−1
• rts/c/backends/c.h +61/−0
• rts/c/backends/cuda.h +947/−0
• rts/c/backends/multicore.h +95/−0
• rts/c/backends/opencl.h +1232/−0
• rts/c/context.h +59/−2
• rts/c/context_prototypes.h +14/−5
• rts/c/cuda.h +0/−766
• rts/c/opencl.h +0/−1066
• rts/c/scalar.h +26/−0
• rts/c/scalar_f16.h +19/−0
• rts/c/server.h +33/−27
• rts/c/uniform.h +21/−0
• rts/c/util.h +13/−0
• rts/javascript/server.js +7/−7
• rts/python/scalar.py +9/−0
• rts/python/server.py +11/−6
• src/Futhark/AD/Derivatives.hs +6/−0
• src/Futhark/AD/Fwd.hs +47/−56
• src/Futhark/Actions.hs +2/−14
• src/Futhark/Analysis/Alias.hs +16/−10
• src/Futhark/Analysis/HORep/SOAC.hs +27/−18
• src/Futhark/Analysis/Interference.hs +17/−16
• src/Futhark/Analysis/LastUse.hs +407/−202
• src/Futhark/Analysis/MemAlias.hs +20/−6
• src/Futhark/Analysis/Metrics.hs +2/−2
• src/Futhark/Analysis/Rephrase.hs +0/−103
• src/Futhark/Analysis/SymbolTable.hs +10/−4
• src/Futhark/Analysis/UsageTable.hs +1/−1
• src/Futhark/Bench.hs +4/−4
• src/Futhark/CLI/Autotune.hs +45/−41
• src/Futhark/CLI/Bench.hs +10/−9
• src/Futhark/CLI/Benchcmp.hs +327/−0
• src/Futhark/CLI/Dataset.hs +2/−2
• src/Futhark/CLI/Dev.hs +3/−12
• src/Futhark/CLI/Eval.hs +130/−0
• src/Futhark/CLI/Literate.hs +32/−17
• src/Futhark/CLI/Main.hs +5/−1
• src/Futhark/CLI/Misc.hs +9/−8
• src/Futhark/CLI/Pkg.hs +72/−115
• src/Futhark/CLI/REPL.hs +20/−33
• src/Futhark/CLI/Run.hs +9/−21
• src/Futhark/CLI/Test.hs +5/−4
• src/Futhark/CodeGen/Backends/CCUDA.hs +6/−44
• src/Futhark/CodeGen/Backends/CCUDA/Boilerplate.hs +73/−403
• src/Futhark/CodeGen/Backends/COpenCL.hs +15/−57
• src/Futhark/CodeGen/Backends/COpenCL/Boilerplate.hs +136/−530
• src/Futhark/CodeGen/Backends/GenericC.hs +48/−31
• src/Futhark/CodeGen/Backends/GenericC/CLI.hs +3/−1
• src/Futhark/CodeGen/Backends/GenericC/Code.hs +7/−20
• src/Futhark/CodeGen/Backends/GenericC/EntryPoints.hs +1/−4
• src/Futhark/CodeGen/Backends/GenericC/Fun.hs +3/−2
• src/Futhark/CodeGen/Backends/GenericC/Monad.hs +36/−44
• src/Futhark/CodeGen/Backends/GenericPython.hs +47/−44
• src/Futhark/CodeGen/Backends/GenericWASM.hs +4/−3
• src/Futhark/CodeGen/Backends/MulticoreC.hs +30/−205
• src/Futhark/CodeGen/Backends/MulticoreC/Boilerplate.hs +23/−0
• src/Futhark/CodeGen/Backends/MulticoreISPC.hs +12/−13
• src/Futhark/CodeGen/Backends/MulticoreWASM.hs +2/−1
• src/Futhark/CodeGen/Backends/PyOpenCL.hs +3/−57
• src/Futhark/CodeGen/Backends/SequentialC/Boilerplate.hs +10/−125
• src/Futhark/CodeGen/Backends/SimpleRep.hs +18/−1
• src/Futhark/CodeGen/ImpCode.hs +19/−12
• src/Futhark/CodeGen/ImpGen.hs +53/−32
• src/Futhark/CodeGen/ImpGen/GPU.hs +3/−16
• src/Futhark/CodeGen/ImpGen/GPU/Base.hs +11/−0
• src/Futhark/CodeGen/ImpGen/GPU/SegHist.hs +3/−5
• src/Futhark/CodeGen/ImpGen/GPU/SegRed.hs +2/−6
• src/Futhark/CodeGen/ImpGen/GPU/SegScan/SinglePass.hs +1/−1
• src/Futhark/CodeGen/ImpGen/GPU/ToOpenCL.hs +27/−25
• src/Futhark/CodeGen/ImpGen/Multicore.hs +3/−4
• src/Futhark/CodeGen/ImpGen/Multicore/SegHist.hs +1/−1
• src/Futhark/CodeGen/ImpGen/Sequential.hs +1/−1
• src/Futhark/CodeGen/RTS/C.hs +24/−12
• src/Futhark/CodeGen/SetDefaultSpace.hs +0/−112
• src/Futhark/Doc/Generator.hs +2/−2
• src/Futhark/IR/Aliases.hs +48/−32
• src/Futhark/IR/GPU.hs +8/−4
• src/Futhark/IR/GPU/Op.hs +25/−33
• src/Futhark/IR/GPU/Simplify.hs +3/−3
• src/Futhark/IR/GPUMem.hs +13/−9
• src/Futhark/IR/MC.hs +2/−4
• src/Futhark/IR/MC/Op.hs +28/−43
• src/Futhark/IR/MCMem.hs +9/−11
• src/Futhark/IR/Mem.hs +26/−26
• src/Futhark/IR/Mem/Simplify.hs +12/−10
• src/Futhark/IR/Parse.hs +4/−4
• src/Futhark/IR/Pretty.hs +3/−1
• src/Futhark/IR/Prop.hs +8/−5
• src/Futhark/IR/Prop/Aliases.hs +53/−53
• src/Futhark/IR/Prop/Names.hs +3/−0
• src/Futhark/IR/Prop/TypeOf.hs +2/−24
• src/Futhark/IR/Rep.hs +17/−3
• src/Futhark/IR/Rephrase.hs +114/−0
• src/Futhark/IR/SOACS.hs +2/−4
• src/Futhark/IR/SOACS/SOAC.hs +29/−19
• src/Futhark/IR/SegOp.hs +44/−61
• src/Futhark/IR/Seq.hs +4/−7
• src/Futhark/IR/SeqMem.hs +6/−10
• src/Futhark/IR/Syntax.hs +2/−2
• src/Futhark/IR/TypeCheck.hs +7/−9
• src/Futhark/Internalise/Defunctionalise.hs +3/−2
• src/Futhark/Optimise/ArrayShortCircuiting.hs +21/−14
• src/Futhark/Optimise/ArrayShortCircuiting/ArrayCoalescing.hs +49/−43
• src/Futhark/Optimise/ArrayShortCircuiting/DataStructs.hs +1/−1
• src/Futhark/Optimise/ArrayShortCircuiting/LastUse.hs +0/−458
• src/Futhark/Optimise/ArrayShortCircuiting/MemRefAggreg.hs +3/−3
• src/Futhark/Optimise/ArrayShortCircuiting/TopdownAnalysis.hs +7/−7
• src/Futhark/Optimise/CSE.hs +19/−33
• src/Futhark/Optimise/Fusion.hs +40/−29
• src/Futhark/Optimise/Fusion/GraphRep.hs +34/−41
• src/Futhark/Optimise/Fusion/TryFusion.hs +8/−4
• src/Futhark/Optimise/InPlaceLowering.hs +2/−2
• src/Futhark/Optimise/InPlaceLowering/LowerIntoStm.hs +1/−1
• src/Futhark/Optimise/InliningDeadFun.hs +3/−4
• src/Futhark/Optimise/MergeGPUBodies.hs +1/−1
• src/Futhark/Optimise/ReduceDeviceSyncs.hs +1/−1
• src/Futhark/Optimise/Simplify/Engine.hs +6/−4
• src/Futhark/Optimise/Simplify/Rep.hs +28/−28
• src/Futhark/Optimise/Simplify/Rules/BasicOp.hs +4/−7
• src/Futhark/Optimise/Simplify/Rules/Index.hs +4/−0
• src/Futhark/Optimise/Sink.hs +2/−2
• src/Futhark/Pass/ExpandAllocations.hs +41/−31
• src/Futhark/Pass/ExplicitAllocations.hs +69/−44
• src/Futhark/Pass/ExplicitAllocations/GPU.hs +10/−11
• src/Futhark/Pass/ExplicitAllocations/MC.hs +1/−1
• src/Futhark/Pass/ExplicitAllocations/SegOp.hs +1/−1
• src/Futhark/Pass/ExplicitAllocations/Seq.hs +5/−1
• src/Futhark/Pass/ExtractKernels.hs +51/−43
• src/Futhark/Pass/ExtractKernels/BlockedKernel.hs +2/−2
• src/Futhark/Pass/ExtractKernels/StreamKernel.hs +1/−1
• src/Futhark/Pass/ExtractKernels/ToGPU.hs +2/−3
• src/Futhark/Pass/ExtractMulticore.hs +0/−1
• src/Futhark/Pass/LiftAllocations.hs +8/−5
• src/Futhark/Pass/LowerAllocations.hs +8/−5
• src/Futhark/Pipeline.hs +61/−19
• src/Futhark/Pkg/Info.hs +155/−277
• src/Futhark/Pkg/Solve.hs +3/−2
• src/Futhark/Script.hs +36/−0
• src/Futhark/Test.hs +6/−8
• src/Futhark/Test/Spec.hs +37/−16
• src/Futhark/Transform/FirstOrderTransform.hs +3/−3
• src/Futhark/Transform/Rename.hs +3/−0
• src/Futhark/Transform/Substitute.hs +3/−0
• src/Futhark/Util.hs +2/−2
• src/Futhark/Util/Log.hs +5/−0
• src/Language/Futhark/Interpreter.hs +7/−8
• src/Language/Futhark/Interpreter/Values.hs +6/−4
• src/Language/Futhark/Parser/Lexer.x +4/−4
• src/Language/Futhark/Parser/Parser.y +34/−22
• src/Language/Futhark/Pretty.hs +5/−3
• src/Language/Futhark/Primitive.hs +5/−0
• src/Language/Futhark/Prop.hs +4/−2
• src/Language/Futhark/Semantic.hs +5/−1
• src/Language/Futhark/Syntax.hs +8/−1
• src/Language/Futhark/Traversals.hs +2/−2
• src/Language/Futhark/TypeChecker.hs +58/−43
• src/Language/Futhark/TypeChecker/Terms.hs +2/−0
• src/Language/Futhark/TypeChecker/Terms/Monad.hs +3/−2
• src/Language/Futhark/TypeChecker/Terms/Pat.hs +4/−2
• src/Language/Futhark/TypeChecker/Unify.hs +36/−28
• unittests/Futhark/BenchTests.hs +1/−1

@@ -496,11 +496,13 @@    Immediately conduct an interactive dialogue on standard output to    select the platform and device from a list. -.. c:function:: struct futhark_context *futhark_context_new_with_command_queue(struct futhark_context_config *cfg, cl_command_queue queue)+.. c:function:: void futhark_context_config_set_command_queue(struct futhark_context_config *cfg, cl_command_queue queue) -   Construct a context that uses a pre-existing command queue.  This-   allows the caller to directly customise which device and platform-   is used.+   Use exactly this command queue for the context.  If this is set,+   all other device/platform configuration options are ignored.  Once+   the context is active, the command queue belongs to Futhark and+   should not be used by anything else.  This is useful for+   implementing custom device selection logic in application code.  .. c:function:: cl_command_queue futhark_context_get_command_queue(struct futhark_context *ctx) 

@@ -533,6 +533,32 @@ technique, but often necessary when writing advanced size-dependent code. +.. _unify-consuming-param:++"Parameter types *x* and *y* are incompatible regarding consuming their arguments+~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~++This error occurs when you provide a function that *does* consume its+argument in a context that expects a function that *does not* allow a+function that consumes its argument.++As a simple example, consider the following contrived function that+does consume its argument:++.. code-block:: futhark++   def f (xs: *[]f32) : f32 = 0f32++Now we define another function that is merely ``f``, but with a type+annotation that tries to hide the consumption:++.. code-block:: futhark++   def g : []f32 -> f32 = f++Allowing this would permit us to hide the fact that ``f`` consumes its+argument, which would not be sound, so the type checker complains.+ .. _ambiguous-size:  "Ambiguous size *x*"

@@ -542,12 +542,20 @@  * Function and type application, and prefix operators, bind more   tightly than any infix operator.  Note that the only prefix-  operators are ``!`` and ``-``, and more cannot be defined.+  operators are the builtin ``!`` and ``-``, and more cannot be+  defined.  In particular, a user-defined operator beginning with+  ``!`` binds as ``!=``, as on the table below, not as the prefix+  operator ``!``  * ``#foo #bar`` is interpreted as a constructor with a ``#bar``   payload, not as applying ``#foo`` to ``#bar`` (the latter would be   semantically invalid anyway). +* A type application ``pt [n]t`` is parsed as an application of the+  type constructor ``pt`` to the size argument ``[n]`` and the type+  ``t``.  To pass a single array-typed parameter, enclose it in+  parens.+ * The following table describes the precedence and associativity of   infix operators in both expressions and type expressions.  All   operators in the same row have the same precedence.  The rows are@@ -563,7 +571,7 @@   left               ```op```   left               ``||``   left               ``&&``-  left               ``<=`` ``>=`` ``>`` ``<`` ``==`` ``!=``+  left               ``<=`` ``>=`` ``>`` ``<`` ``==`` ``!=`` ``!`` ``=``   left               ``&`` ``^`` ``|``   left               ``<<`` ``>>``   left               ``+`` ``-``@@ -767,8 +775,8 @@  Apply an operator to ``x`` and ``y``.  Operators are functions like any other, and can be user-defined.  Futhark pre-defines certain-"magical" *overloaded* operators that work on many different types.-Overloaded functions cannot be defined by the user.  Both operands+"magical" *overloaded* operators that work on several types.+Overloaded operators cannot be defined by the user.  Both operands must have the same type.  The predefined operators and their semantics are: @@ -801,16 +809,16 @@    ``==``, ``!=`` -      Compare any two values of builtin or compound type for equality.+    Compare any two values of builtin or compound type for equality.    ``<``, ``<=``.  ``>``, ``>=`` -      Company any two values of numeric type for equality.+    Company any two values of numeric type for equality.    ```op``` -      Use ``op``, which may be any non-operator function name, as an-      infix operator.+    Use ``op``, which may be any non-operator function name, as an+    infix operator.  ``x && y`` ..........@@ -927,7 +935,9 @@  Evaluate ``e`` and bind the result to the irrefutable pattern ``pat`` (see :ref:`patterns`) while evaluating ``body``.  The ``in`` keyword-is optional if ``body`` is a ``let`` expression.+is optional if ``body`` is a ``let`` expression.  The binding is not+let-generalised, meaning it has a monomorphic type.  This can be+significant if ``e`` is of functional type.  ``let [n] pat = e in body`` ...........................@@ -1072,6 +1082,10 @@ definition site.  Specifically, if a top-level function uses overloaded arithmetic operators, the resolution of those overloads cannot be influenced by later uses of the function.++Local bindings made with ``let`` are not made polymorphic through+let-generalisation *unless* they are syntactically functions, meaning+they have at least one named parameter.  .. _size-types: 

@@ -42,7 +42,7 @@   generated.  -s int-  Set the seed used for the RNG.  Zero by default.+  Set the seed used for the RNG.  1 by default.  --T-bounds=<min:max>   Set inclusive lower and upper bounds on generated values of type

@@ -270,6 +270,11 @@   of values is returned, which should be destructured before use. For example:   ``let (a, b) = $loaddata "foo.in" in bar a b``. +* ``$loadaudio "file"`` reads audio from the given file and returns it as a+  ``[][]f64``, where each row corresponds to a channel of the original+  soundfile. Most common audio-formats are supported, including mp3, ogg, wav,+  flac and opus.+ SAFETY ====== 

@@ -53,8 +53,8 @@ Most commands take a ``-v``/``--verbose`` option that makes ``futhark pkg`` write running diagnostics to stderr. -Network requests (exclusively HTTP GETs) are done via ``curl``, which-must be available on the ``PATH``.+Packages must correspond to Git repositories, and all interactions are+done by invoking ``git``.  COMMANDS ========@@ -159,10 +159,8 @@ the dependencies of its dependencies, without the ``futhark.pkg`` file reflecting this. -There is no caching of zipballs and version lists between invocations,-so the network traffic can be rather heavy.--Only GitHub and GitLab are supported as code hosting sites.+There is no caching of package metadata between invocations, so the+network traffic can be rather heavy.  SEE ALSO ========

@@ -37,9 +37,14 @@    [tags { tags... }]   [entry: names...]-  [compiled|nobench|random|script] input ({ values... } | @ filename)+  ["name..."] [compiled|nobench|random|script] input ({ values... } | @ filename)   output { values... } | auto output | error: regex +If a test case begins with a quoted string, that string is reported as+the dataset name, including in the JSON file produced by+:ref:`futhark-bench(1)`.  If no name is provided, one is automatically+generated.  The name must be unique across all test cases.+ If ``compiled`` is present before the ``input`` keyword, this test case will never be passed to the interpreter.  This is useful for test cases that are annoyingly slow to interpret.  The ``nobench`` keyword@@ -60,6 +65,7 @@ It must use only functions explicitly declared as entry points.  If the expression produces an *n*-element tuple, it will be unpacked and its components passed as *n* distinct arguments to the test function.+The only builtin function supported is ``$loaddata``.  If ``input`` is followed by an ``@`` and a file name (which must not contain any whitespace) instead of curly braces, values or

@@ -24,6 +24,13 @@ COMMANDS ======== +futhark benchcmp FILE_A FILE_B+------------------------------++Compares two Futhark benchmarks and reports changes in performance.+The files must be formatted in the same manner as a JSON file returned+from :ref:`futhark-bench(1)`.+ futhark check [-w] PROGRAM -------------------------- @@ -65,6 +72,12 @@ A Futhark compiler development command, intentionally undocumented and intended for use in developing the Futhark compiler, not for programmers writing in Futhark.++futhark eval [-f FILE] [-w] <exprs...>+--------------------------------------++Evaluates expressions given as command-line arguments. Optionally +allows a file import using ``-f``.  futhark hash PROGRAM --------------------

@@ -14,9 +14,11 @@  A package is uniquely identified with a *package path*, which is similar to a URL, except without a protocol.  At the moment, package-paths are always links to Git repositories hosted on GitHub or GitLab.-In the future, this will become more flexible.  As an example, a-package path may be ``github.com/athas/fut-foo``.+paths must be something that can be passed to ``git clone``.  In+particular, this includes paths to repositories on major code hosting+sites such as GitLab and GitHub. In the future, this will become more+flexible.  As an example, a package path may be+``github.com/athas/fut-foo``.  Packages are versioned with `semantic version numbers <https://semver.org/>`_ of the form ``X.Y.Z``.  Whenever versions are@@ -359,3 +361,23 @@ (but the ``--safe`` compiler option may help).  However, this is not any worse than calling external code in a conventional impure language, which generally can perform any conceivable harmful action.++Private repositories+--------------------++The Futhark package manager is intentionally very simple - perhaps+even simplistic.  The key philosophy is that if you can ``git clone``+a repository from the command line, then ``futhark pkg`` can also+access it.  However, ``futhark pkg`` always uses the ``https://``+protocol when converting package paths to the URLs that are passed to+``git``, which is sometimes inconvenient for self-hosted or private+repositories.  As a workaround, you can modify your Git configuration+file to transparently replace ``https://`` with ``ssh://`` for certain+repositories.  For example, you can add the following entry+``$HOME/.gitconfig``::++  [url "ssh://git@github.com/sturluson"]+	insteadOf = https://github.com/sturluson++This will make all interactions with repositories owned by the+``sturluson`` user on GitHub use SSH instead of HTTPS.

@@ -242,12 +242,13 @@     (u8, [n]i32) +However the type  .. code-block:: futhark     #foo [n]i32 | #bar [n]f32 -becomes+is represented as  .. code-block:: futhark 

@@ -1,6 +1,6 @@ cabal-version: 2.4 name:           futhark-version:        0.22.7+version:        0.23.1 synopsis:       An optimising compiler for a functional, array-oriented language.  description:    Futhark is a small programming language designed to be compiled to@@ -39,6 +39,10 @@     rts/c/atomics.h     rts/c/context.h     rts/c/context_prototypes.h+    rts/c/backends/c.h+    rts/c/backends/cuda.h+    rts/c/backends/multicore.h+    rts/c/backends/opencl.h     rts/c/lock.h     rts/c/timing.h     rts/c/errors.h@@ -46,8 +50,6 @@     rts/c/tuning.h     rts/c/values.h     rts/c/half.h-    rts/c/opencl.h-    rts/c/cuda.h     rts/c/cache.h     rts/c/ispc_util.h     rts/c/scalar.h@@ -172,7 +174,6 @@       Futhark.Analysis.PrimExp.Convert       Futhark.Analysis.PrimExp.Parse       Futhark.Analysis.PrimExp.Simplify-      Futhark.Analysis.Rephrase       Futhark.Analysis.SymbolTable       Futhark.Analysis.UsageTable       Futhark.Bench@@ -183,11 +184,13 @@       Futhark.CLI.C       Futhark.CLI.CUDA       Futhark.CLI.Check+      Futhark.CLI.Benchcmp       Futhark.CLI.Datacmp       Futhark.CLI.Dataset       Futhark.CLI.Defs       Futhark.CLI.Dev       Futhark.CLI.Doc+      Futhark.CLI.Eval       Futhark.CLI.Literate       Futhark.CLI.LSP       Futhark.CLI.Main@@ -223,6 +226,7 @@       Futhark.CodeGen.Backends.GenericPython.Options       Futhark.CodeGen.Backends.GenericWASM       Futhark.CodeGen.Backends.MulticoreC+      Futhark.CodeGen.Backends.MulticoreC.Boilerplate       Futhark.CodeGen.Backends.MulticoreISPC       Futhark.CodeGen.Backends.MulticoreWASM       Futhark.CodeGen.Backends.PyOpenCL@@ -263,7 +267,6 @@       Futhark.CodeGen.ImpGen.Sequential       Futhark.CodeGen.ImpGen.Transpose       Futhark.CodeGen.OpenCL.Heuristics-      Futhark.CodeGen.SetDefaultSpace       Futhark.Compiler       Futhark.Compiler.CLI       Futhark.Compiler.Config@@ -299,6 +302,7 @@       Futhark.IR.Prop.TypeOf       Futhark.IR.Prop.Types       Futhark.IR.Rep+      Futhark.IR.Rephrase       Futhark.IR.RetType       Futhark.IR.SOACS       Futhark.IR.SOACS.SOAC@@ -348,7 +352,6 @@       Futhark.Optimise.ArrayShortCircuiting       Futhark.Optimise.ArrayShortCircuiting.ArrayCoalescing       Futhark.Optimise.ArrayShortCircuiting.DataStructs-      Futhark.Optimise.ArrayShortCircuiting.LastUse       Futhark.Optimise.ArrayShortCircuiting.MemRefAggreg       Futhark.Optimise.ArrayShortCircuiting.TopdownAnalysis       Futhark.Optimise.MergeGPUBodies@@ -507,7 +510,6 @@     , transformers >=0.3     , vector >=0.12     , versions >=5.0.0-    , zip-archive >=0.3.1.1     , zlib >=0.6.1.2     , statistics     , mwc-random@@ -517,7 +519,7 @@ executable futhark   import: common   main-is: src/main.hs-  ghc-options: -threaded -rtsopts "-with-rtsopts=-N -qg1 -A16M"+  ghc-options: -threaded -rtsopts "-with-rtsopts=-maxN16 -qg1 -A16M"   build-depends: base, futhark  test-suite unit

@@ -187,6 +187,8 @@   val log2: t -> t   -- | Base-10 logarithm.   val log10: t -> t+  -- | Compute `log (1 + x)` accurately even when `x` is very small.+  val log1p: t -> t    -- | Round towards infinity.   val ceil : t -> t@@ -927,6 +929,7 @@   def log (x: f64) = intrinsics.log64 x   def log2 (x: f64) = intrinsics.log2_64 x   def log10 (x: f64) = intrinsics.log10_64 x+  def log1p (x: f64) = intrinsics.log1p_64 x   def exp (x: f64) = intrinsics.exp64 x   def sin (x: f64) = intrinsics.sin64 x   def cos (x: f64) = intrinsics.cos64 x@@ -1041,6 +1044,7 @@   def log (x: f32) = intrinsics.log32 x   def log2 (x: f32) = intrinsics.log2_32 x   def log10 (x: f32) = intrinsics.log10_32 x+  def log1p (x: f32) = intrinsics.log1p_32 x   def exp (x: f32) = intrinsics.exp32 x   def sin (x: f32) = intrinsics.sin32 x   def cos (x: f32) = intrinsics.cos32 x@@ -1159,6 +1163,7 @@   def log (x: f16) = intrinsics.log16 x   def log2 (x: f16) = intrinsics.log2_16 x   def log10 (x: f16) = intrinsics.log10_16 x+  def log1p (x: f16) = intrinsics.log1p_16 x   def exp (x: f16) = intrinsics.exp16 x   def sin (x: f16) = intrinsics.sin16 x   def cos (x: f16) = intrinsics.cos16 x

@@ -211,7 +211,7 @@ def any [n] 'a (f: a -> bool) (as: [n]a): bool =   reduce (||) false (map f as) --- | `r = spread x k is vs` produces an array `r` such that `r[i] =+-- | `r = spread k x is vs` produces an array `r` such that `r[i] = -- vs[j]` where `is[j] == i`, or `x` if no such `j` exists. -- Intuitively, `is` is an array indicating where the corresponding -- elements of `vs` should be located in the result.  Out-of-bounds

@@ -0,0 +1,61 @@+// Start of backends/c.h++struct futhark_context_config {+  int in_use;+  int debugging;+  int profiling;+  int logging;+  const char *cache_fname;+  int num_tuning_params;+  int64_t *tuning_params;+  const char** tuning_param_names;+  const char** tuning_param_vars;+  const char** tuning_param_classes;+};++static void backend_context_config_setup(struct futhark_context_config* cfg) {+  (void)cfg;+}++static void backend_context_config_teardown(struct futhark_context_config* cfg) {+  (void)cfg;+}++int futhark_context_config_set_tuning_param(struct futhark_context_config* cfg, const char *param_name, size_t param_value) {+  (void)cfg; (void)param_name; (void)param_value;+  return 1;+}++struct futhark_context {+  struct futhark_context_config* cfg;+  int detail_memory;+  int debugging;+  int profiling;+  int profiling_paused;+  int logging;+  lock_t lock;+  char *error;+  lock_t error_lock;+  FILE *log;+  struct constants *constants;+  struct free_list free_list;+  int64_t peak_mem_usage_default;+  int64_t cur_mem_usage_default;+  struct program* program;+};++int backend_context_setup(struct futhark_context* ctx) {+  (void)ctx;+  return 0;+}++void backend_context_teardown(struct futhark_context* ctx) {+  (void)ctx;+}++int futhark_context_sync(struct futhark_context* ctx) {+  (void)ctx;+  return 0;+}++// End of backends/c.h

@@ -0,0 +1,947 @@+// Start of backends/cuda.h.++// Forward declarations.+// Invoked by setup_opencl() after the platform and device has been+// found, but before the program is loaded.  Its intended use is to+// tune constants based on the selected platform and device.+static void set_tuning_params(struct futhark_context* ctx);+static char* get_failure_msg(int failure_idx, int64_t args[]);++#define CUDA_SUCCEED_FATAL(x) cuda_api_succeed_fatal(x, #x, __FILE__, __LINE__)+#define CUDA_SUCCEED_NONFATAL(x) cuda_api_succeed_nonfatal(x, #x, __FILE__, __LINE__)+#define NVRTC_SUCCEED_FATAL(x) nvrtc_api_succeed_fatal(x, #x, __FILE__, __LINE__)+#define NVRTC_SUCCEED_NONFATAL(x) nvrtc_api_succeed_nonfatal(x, #x, __FILE__, __LINE__)+// Take care not to override an existing error.+#define CUDA_SUCCEED_OR_RETURN(e) {             \+    char *serror = CUDA_SUCCEED_NONFATAL(e);    \+    if (serror) {                               \+      if (!ctx->error) {                        \+        ctx->error = serror;                    \+        return bad;                             \+      } else {                                  \+        free(serror);                           \+      }                                         \+    }                                           \+  }++// CUDA_SUCCEED_OR_RETURN returns the value of the variable 'bad' in+// scope.  By default, it will be this one.  Create a local variable+// of some other type if needed.  This is a bit of a hack, but it+// saves effort in the code generator.+static const int bad = 1;++static inline void cuda_api_succeed_fatal(CUresult res, const char *call,+                                          const char *file, int line) {+  if (res != CUDA_SUCCESS) {+    const char *err_str;+    cuGetErrorString(res, &err_str);+    if (err_str == NULL) { err_str = "Unknown"; }+    futhark_panic(-1, "%s:%d: CUDA call\n  %s\nfailed with error code %d (%s)\n",+                  file, line, call, res, err_str);+  }+}++static char* cuda_api_succeed_nonfatal(CUresult res, const char *call,+                                       const char *file, int line) {+  if (res != CUDA_SUCCESS) {+    const char *err_str;+    cuGetErrorString(res, &err_str);+    if (err_str == NULL) { err_str = "Unknown"; }+    return msgprintf("%s:%d: CUDA call\n  %s\nfailed with error code %d (%s)\n",+                     file, line, call, res, err_str);+  } else {+    return NULL;+  }+}++static inline void nvrtc_api_succeed_fatal(nvrtcResult res, const char *call,+                                           const char *file, int line) {+  if (res != NVRTC_SUCCESS) {+    const char *err_str = nvrtcGetErrorString(res);+    futhark_panic(-1, "%s:%d: NVRTC call\n  %s\nfailed with error code %d (%s)\n",+                  file, line, call, res, err_str);+  }+}++static char* nvrtc_api_succeed_nonfatal(nvrtcResult res, const char *call,+                                        const char *file, int line) {+  if (res != NVRTC_SUCCESS) {+    const char *err_str = nvrtcGetErrorString(res);+    return msgprintf("%s:%d: NVRTC call\n  %s\nfailed with error code %d (%s)\n",+                     file, line, call, res, err_str);+  } else {+    return NULL;+  }+}++struct futhark_context_config {+  int in_use;+  int debugging;+  int profiling;+  int logging;+  const char *cache_fname;+  int num_tuning_params;+  int64_t *tuning_params;+  const char** tuning_param_names;+  const char** tuning_param_vars;+  const char** tuning_param_classes;+  // Uniform fields above.++  int num_nvrtc_opts;+  const char **nvrtc_opts;++  const char *preferred_device;+  int preferred_device_num;++  const char *dump_program_to;+  const char *load_program_from;++  const char *dump_ptx_to;+  const char *load_ptx_from;++  size_t default_block_size;+  size_t default_grid_size;+  size_t default_tile_size;+  size_t default_reg_tile_size;+  size_t default_threshold;++  int default_block_size_changed;+  int default_grid_size_changed;+  int default_tile_size_changed;+};++static void backend_context_config_setup(struct futhark_context_config *cfg) {+  cfg->num_nvrtc_opts = 0;+  cfg->nvrtc_opts = (const char**) malloc(sizeof(const char*));+  cfg->nvrtc_opts[0] = NULL;++  cfg->preferred_device_num = 0;+  cfg->preferred_device = "";+  cfg->dump_program_to = NULL;+  cfg->load_program_from = NULL;++  cfg->dump_ptx_to = NULL;+  cfg->load_ptx_from = NULL;++  cfg->default_block_size = 256;+  cfg->default_grid_size = 0; // Set properly later.+  cfg->default_tile_size = 32;+  cfg->default_reg_tile_size = 2;+  cfg->default_threshold = 32*1024;++  cfg->default_block_size_changed = 0;+  cfg->default_grid_size_changed = 0;+  cfg->default_tile_size_changed = 0;+}++static void backend_context_config_teardown(struct futhark_context_config* cfg) {+  free(cfg->nvrtc_opts);+}++void futhark_context_config_add_nvrtc_option(struct futhark_context_config *cfg, const char *opt) {+  cfg->nvrtc_opts[cfg->num_nvrtc_opts] = opt;+  cfg->num_nvrtc_opts++;+  cfg->nvrtc_opts = (const char **) realloc(cfg->nvrtc_opts, (cfg->num_nvrtc_opts + 1) * sizeof(const char *));+  cfg->nvrtc_opts[cfg->num_nvrtc_opts] = NULL;+}++void futhark_context_config_set_device(struct futhark_context_config *cfg, const char *s) {+  int x = 0;+  if (*s == '#') {+    s++;+    while (isdigit(*s)) {+      x = x * 10 + (*s++)-'0';+    }+    // Skip trailing spaces.+    while (isspace(*s)) {+      s++;+    }+  }+  cfg->preferred_device = s;+  cfg->preferred_device_num = x;+}++void futhark_context_config_dump_program_to(struct futhark_context_config *cfg, const char *path) {+  cfg->dump_program_to = path;+}++void futhark_context_config_load_program_from(struct futhark_context_config *cfg, const char *path) {+  cfg->load_program_from = path;+}++void futhark_context_config_dump_ptx_to(struct futhark_context_config *cfg, const char *path) {+  cfg->dump_ptx_to = path;+}++void futhark_context_config_load_ptx_from(struct futhark_context_config *cfg, const char *path) {+  cfg->load_ptx_from = path;+}++void futhark_context_config_set_default_group_size(struct futhark_context_config *cfg, int size) {+  cfg->default_block_size = size;+  cfg->default_block_size_changed = 1;+}++void futhark_context_config_set_default_num_groups(struct futhark_context_config *cfg, int num) {+  cfg->default_grid_size = num;+  cfg->default_grid_size_changed = 1;+}++void futhark_context_config_set_default_tile_size(struct futhark_context_config *cfg, int size) {+  cfg->default_tile_size = size;+  cfg->default_tile_size_changed = 1;+}++void futhark_context_config_set_default_reg_tile_size(struct futhark_context_config *cfg, int size) {+  cfg->default_reg_tile_size = size;+}++void futhark_context_config_set_default_threshold(struct futhark_context_config *cfg, int size) {+  cfg->default_threshold = size;+}++int futhark_context_config_set_tuning_param(struct futhark_context_config *cfg,+                                            const char *param_name,+                                            size_t new_value) {+  for (int i = 0; i < cfg->num_tuning_params; i++) {+    if (strcmp(param_name, cfg->tuning_param_names[i]) == 0) {+      cfg->tuning_params[i] = new_value;+      return 0;+    }+  }+  if (strcmp(param_name, "default_group_size") == 0) {+    cfg->default_block_size = new_value;+    return 0;+  }+  if (strcmp(param_name, "default_num_groups") == 0) {+    cfg->default_grid_size = new_value;+    return 0;+  }+  if (strcmp(param_name, "default_threshold") == 0) {+    cfg->default_threshold = new_value;+    return 0;+  }+  if (strcmp(param_name, "default_tile_size") == 0) {+    cfg->default_tile_size = new_value;+    return 0;+  }+  if (strcmp(param_name, "default_reg_tile_size") == 0) {+    cfg->default_reg_tile_size = new_value;+    return 0;+  }+  return 1;+}++// A record of something that happened.+struct profiling_record {+  cudaEvent_t *events; // Points to two events.+  int *runs;+  int64_t *runtime;+};++struct futhark_context {+  struct futhark_context_config* cfg;+  int detail_memory;+  int debugging;+  int profiling;+  int profiling_paused;+  int logging;+  lock_t lock;+  char *error;+  lock_t error_lock;+  FILE *log;+  struct constants *constants;+  struct free_list cu_free_list;+  int64_t peak_mem_usage_default;+  int64_t cur_mem_usage_default;+  // Uniform above++  CUdeviceptr global_failure;+  CUdeviceptr global_failure_args;+  struct tuning_params tuning_params;+  // True if a potentially failing kernel has been enqueued.+  int32_t failure_is_an_option;+  int total_runs;+  long int total_runtime;+  int64_t peak_mem_usage_device;+  int64_t cur_mem_usage_device;+  struct program* program;++  CUdevice dev;+  CUcontext cu_ctx;+  CUmodule module;++  struct free_list free_list;++  size_t max_block_size;+  size_t max_grid_size;+  size_t max_tile_size;+  size_t max_threshold;+  size_t max_shared_memory;+  size_t max_bespoke;++  size_t lockstep_width;++  struct profiling_record *profiling_records;+  int profiling_records_capacity;+  int profiling_records_used;+};++#define CU_DEV_ATTR(x) (CU_DEVICE_ATTRIBUTE_##x)+#define device_query(dev,attrib) _device_query(dev, CU_DEV_ATTR(attrib))+static int _device_query(CUdevice dev, CUdevice_attribute attrib) {+  int val;+  CUDA_SUCCEED_FATAL(cuDeviceGetAttribute(&val, attrib, dev));+  return val;+}++#define CU_FUN_ATTR(x) (CU_FUNC_ATTRIBUTE_##x)+#define function_query(fn,attrib) _function_query(dev, CU_FUN_ATTR(attrib))+static int _function_query(CUfunction dev, CUfunction_attribute attrib) {+  int val;+  CUDA_SUCCEED_FATAL(cuFuncGetAttribute(&val, attrib, dev));+  return val;+}++static int cuda_device_setup(struct futhark_context *ctx) {+  struct futhark_context_config *cfg = ctx->cfg;+  char name[256];+  int count, chosen = -1, best_cc = -1;+  int cc_major_best, cc_minor_best;+  int cc_major, cc_minor;+  CUdevice dev;++  CUDA_SUCCEED_FATAL(cuDeviceGetCount(&count));+  if (count == 0) { return 1; }++  int num_device_matches = 0;++  // XXX: Current device selection policy is to choose the device with the+  // highest compute capability (if no preferred device is set).+  // This should maybe be changed, since greater compute capability is not+  // necessarily an indicator of better performance.+  for (int i = 0; i < count; i++) {+    CUDA_SUCCEED_FATAL(cuDeviceGet(&dev, i));++    cc_major = device_query(dev, COMPUTE_CAPABILITY_MAJOR);+    cc_minor = device_query(dev, COMPUTE_CAPABILITY_MINOR);++    CUDA_SUCCEED_FATAL(cuDeviceGetName(name, sizeof(name) - 1, dev));+    name[sizeof(name) - 1] = 0;++    if (cfg->logging) {+      fprintf(stderr, "Device #%d: name=\"%s\", compute capability=%d.%d\n",+              i, name, cc_major, cc_minor);+    }++    if (device_query(dev, COMPUTE_MODE) == CU_COMPUTEMODE_PROHIBITED) {+      if (cfg->logging) {+        fprintf(stderr, "Device #%d is compute-prohibited, ignoring\n", i);+      }+      continue;+    }++    if (best_cc == -1 || cc_major > cc_major_best ||+        (cc_major == cc_major_best && cc_minor > cc_minor_best)) {+      best_cc = i;+      cc_major_best = cc_major;+      cc_minor_best = cc_minor;+    }++    if (strstr(name, cfg->preferred_device) != NULL &&+        num_device_matches++ == cfg->preferred_device_num) {+      chosen = i;+      break;+    }+  }++  if (chosen == -1) { chosen = best_cc; }+  if (chosen == -1) { return 1; }++  if (cfg->logging) {+    fprintf(stderr, "Using device #%d\n", chosen);+  }++  CUDA_SUCCEED_FATAL(cuDeviceGet(&ctx->dev, chosen));+  return 0;+}++static char *concat_fragments(const char *src_fragments[]) {+  size_t src_len = 0;+  const char **p;++  for (p = src_fragments; *p; p++) {+    src_len += strlen(*p);+  }++  char *src = (char*) malloc(src_len + 1);+  size_t n = 0;+  for (p = src_fragments; *p; p++) {+    strcpy(src + n, *p);+    n += strlen(*p);+  }++  return src;+}++static const char *cuda_nvrtc_get_arch(CUdevice dev) {+  struct {+    int major;+    int minor;+    const char *arch_str;+  } static const x[] = {+    { 3, 0, "compute_30" },+    { 3, 2, "compute_32" },+    { 3, 5, "compute_35" },+    { 3, 7, "compute_37" },+    { 5, 0, "compute_50" },+    { 5, 2, "compute_52" },+    { 5, 3, "compute_53" },+    { 6, 0, "compute_60" },+    { 6, 1, "compute_61" },+    { 6, 2, "compute_62" },+    { 7, 0, "compute_70" },+    { 7, 2, "compute_72" },+    { 7, 5, "compute_75" },+    { 8, 0, "compute_80" },+    { 8, 6, "compute_80" },+    { 8, 7, "compute_80" }+  };++  int major = device_query(dev, COMPUTE_CAPABILITY_MAJOR);+  int minor = device_query(dev, COMPUTE_CAPABILITY_MINOR);++  int chosen = -1;+  int num_archs = sizeof(x)/sizeof(x[0]);+  for (int i = 0; i < num_archs; i++) {+    if (x[i].major < major || (x[i].major == major && x[i].minor <= minor)) {+      chosen = i;+    } else {+      break;+    }+  }++  if (chosen == -1) {+    futhark_panic(-1, "Unsupported compute capability %d.%d\n", major, minor);+  }++  if (x[chosen].major != major || x[chosen].minor != minor) {+    fprintf(stderr,+            "Warning: device compute capability is %d.%d, but newest supported by Futhark is %d.%d.\n",+            major, minor, x[chosen].major, x[chosen].minor);+  }++  return x[chosen].arch_str;+}++static void cuda_nvrtc_mk_build_options(struct futhark_context *ctx, const char *extra_opts[],+                                        char*** opts_out, size_t *n_opts) {+  int arch_set = 0, num_extra_opts;+  struct futhark_context_config *cfg = ctx->cfg;++  // 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 i = 0, n_opts_alloc = 20 + num_extra_opts + cfg->num_tuning_params;+  char **opts = (char**) malloc(n_opts_alloc * sizeof(char *));+  if (!arch_set) {+    opts[i++] = strdup("-arch");+    opts[i++] = strdup(cuda_nvrtc_get_arch(ctx->dev));+  }+  opts[i++] = strdup("-default-device");+  if (cfg->debugging) {+    opts[i++] = strdup("-G");+    opts[i++] = strdup("-lineinfo");+  } else {+    opts[i++] = strdup("--disable-warnings");+  }+  opts[i++] = msgprintf("-D%s=%d",+                        "max_group_size",+                        (int)ctx->max_block_size);+  for (int j = 0; j < cfg->num_tuning_params; j++) {+    opts[i++] = msgprintf("-D%s=%zu", cfg->tuning_param_vars[j],+                          cfg->tuning_params[j]);+  }+  opts[i++] = msgprintf("-DLOCKSTEP_WIDTH=%zu", ctx->lockstep_width);+  opts[i++] = msgprintf("-DMAX_THREADS_PER_BLOCK=%zu", ctx->max_block_size);++  // Time for the best lines of the code in the entire compiler.+  if (getenv("CUDA_HOME") != NULL) {+    opts[i++] = msgprintf("-I%s/include", getenv("CUDA_HOME"));+  }+  if (getenv("CUDA_ROOT") != NULL) {+    opts[i++] = msgprintf("-I%s/include", getenv("CUDA_ROOT"));+  }+  if (getenv("CUDA_PATH") != NULL) {+    opts[i++] = msgprintf("-I%s/include", getenv("CUDA_PATH"));+  }+  opts[i++] = msgprintf("-I/usr/local/cuda/include");+  opts[i++] = msgprintf("-I/usr/include");++  for (int j = 0; extra_opts[j] != NULL; j++) {+    opts[i++] = strdup(extra_opts[j]);+  }++  *n_opts = i;+  *opts_out = opts;+}++static char* cuda_nvrtc_build(const char *src, const char *opts[], size_t n_opts,+                              char **ptx) {+  nvrtcProgram prog;+  char *problem = NULL;++  problem = NVRTC_SUCCEED_NONFATAL(nvrtcCreateProgram(&prog, src, "futhark-cuda", 0, NULL, NULL));++  if (problem) {+    return problem;+  }++  nvrtcResult res = nvrtcCompileProgram(prog, n_opts, opts);+  if (res != NVRTC_SUCCESS) {+    size_t log_size;+    if (nvrtcGetProgramLogSize(prog, &log_size) == NVRTC_SUCCESS) {+      char *log = (char*) malloc(log_size);+      if (nvrtcGetProgramLog(prog, log) == NVRTC_SUCCESS) {+        problem = msgprintf("NVRTC compilation failed.\n\n%s\n", log);+      } else {+        problem = msgprintf("Could not retrieve compilation log\n");+      }+      free(log);+    }+    return problem;+  }++  size_t ptx_size;+  NVRTC_SUCCEED_FATAL(nvrtcGetPTXSize(prog, &ptx_size));+  *ptx = (char*) malloc(ptx_size);+  NVRTC_SUCCEED_FATAL(nvrtcGetPTX(prog, *ptx));++  NVRTC_SUCCEED_FATAL(nvrtcDestroyProgram(&prog));++  return NULL;+}++static void cuda_load_ptx_from_cache(struct futhark_context_config *cfg,+                                     const char *src,+                                     const char *opts[], size_t n_opts,+                                     struct cache_hash *h, const char *cache_fname,+                                     char **ptx) {+  if (cfg->logging) {+    fprintf(stderr, "Restoring cache from from %s...\n", cache_fname);+  }+  cache_hash_init(h);+  for (size_t i = 0; i < n_opts; i++) {+    cache_hash(h, opts[i], strlen(opts[i]));+  }+  cache_hash(h, src, strlen(src));+  size_t ptxsize;+  errno = 0;+  if (cache_restore(cache_fname, h, (unsigned char**)ptx, &ptxsize) != 0) {+    if (cfg->logging) {+      fprintf(stderr, "Failed to restore cache (errno: %s)\n", strerror(errno));+    }+  }+}++static void cuda_size_setup(struct futhark_context *ctx)+{+  struct futhark_context_config *cfg = ctx->cfg;+  if (cfg->default_block_size > ctx->max_block_size) {+    if (cfg->default_block_size_changed) {+      fprintf(stderr,+              "Note: Device limits default block size to %zu (down from %zu).\n",+              ctx->max_block_size, cfg->default_block_size);+    }+    cfg->default_block_size = ctx->max_block_size;+  }+  if (cfg->default_grid_size > ctx->max_grid_size) {+    if (cfg->default_grid_size_changed) {+      fprintf(stderr,+              "Note: Device limits default grid size to %zu (down from %zu).\n",+              ctx->max_grid_size, cfg->default_grid_size);+    }+    cfg->default_grid_size = ctx->max_grid_size;+  }+  if (cfg->default_tile_size > ctx->max_tile_size) {+    if (cfg->default_tile_size_changed) {+      fprintf(stderr,+              "Note: Device limits default tile size to %zu (down from %zu).\n",+              ctx->max_tile_size, cfg->default_tile_size);+    }+    cfg->default_tile_size = ctx->max_tile_size;+  }++  if (!cfg->default_grid_size_changed) {+    cfg->default_grid_size =+      (device_query(ctx->dev, MULTIPROCESSOR_COUNT) *+       device_query(ctx->dev, MAX_THREADS_PER_MULTIPROCESSOR))+      / cfg->default_block_size;+  }++  for (int i = 0; i < cfg->num_tuning_params; i++) {+    const char *size_class = cfg->tuning_param_classes[i];+    int64_t *size_value = &cfg->tuning_params[i];+    const char* size_name = cfg->tuning_param_names[i];+    int64_t max_value = 0, default_value = 0;++    if (strstr(size_class, "group_size") == size_class) {+      max_value = ctx->max_block_size;+      default_value = cfg->default_block_size;+    } else if (strstr(size_class, "num_groups") == size_class) {+      max_value = ctx->max_grid_size;+      default_value = cfg->default_grid_size;+      // XXX: as a quick and dirty hack, use twice as many threads for+      // histograms by default.  We really should just be smarter+      // about sizes somehow.+      if (strstr(size_name, ".seghist_") != NULL) {+        default_value *= 2;+      }+    } else if (strstr(size_class, "tile_size") == size_class) {+      max_value = ctx->max_tile_size;+      default_value = cfg->default_tile_size;+    } else if (strstr(size_class, "reg_tile_size") == size_class) {+      max_value = 0; // No limit.+      default_value = cfg->default_reg_tile_size;+    } else if (strstr(size_class, "threshold") == size_class) {+      // Threshold can be as large as it takes.+      default_value = cfg->default_threshold;+    } else {+      // Bespoke sizes have no limit or default.+    }++    if (*size_value == 0) {+      *size_value = default_value;+    } else if (max_value > 0 && *size_value > max_value) {+      fprintf(stderr, "Note: Device limits %s to %zu (down from %zu)\n",+              size_name, max_value, *size_value);+      *size_value = max_value;+    }+  }+}++static char* cuda_module_setup(struct futhark_context *ctx,+                               const char *src_fragments[],+                               const char *extra_opts[],+                               const char* cache_fname) {+  char *ptx = NULL, *src = NULL;+  struct futhark_context_config *cfg = ctx->cfg;++  if (cfg->load_program_from == NULL) {+    src = concat_fragments(src_fragments);+  } else {+    src = slurp_file(cfg->load_program_from, NULL);+  }++  if (cfg->load_ptx_from) {+    if (cfg->load_program_from != NULL) {+      fprintf(stderr,+              "WARNING: Using PTX from %s instead of C code from %s\n",+              cfg->load_ptx_from, cfg->load_program_from);+    }+    ptx = slurp_file(cfg->load_ptx_from, NULL);+  }++  if (cfg->dump_program_to != NULL) {+    dump_file(cfg->dump_program_to, src, strlen(src));+  }++  char **opts;+  size_t n_opts;+  cuda_nvrtc_mk_build_options(ctx, extra_opts, &opts, &n_opts);++  if (cfg->logging) {+    fprintf(stderr, "NVRTC compile options:\n");+    for (size_t j = 0; j < n_opts; j++) {+      fprintf(stderr, "\t%s\n", opts[j]);+    }+    fprintf(stderr, "\n");+  }++  struct cache_hash h;+  int loaded_ptx_from_cache = 0;+  if (cache_fname != NULL) {+    cuda_load_ptx_from_cache(cfg, src, (const char**)opts, n_opts, &h, cache_fname, &ptx);++    if (ptx != NULL) {+      if (cfg->logging) {+        fprintf(stderr, "Restored PTX from cache; now loading module...\n");+      }+      if (cuModuleLoadData(&ctx->module, ptx) == CUDA_SUCCESS) {+        if (cfg->logging) {+          fprintf(stderr, "Success!\n");+        }+        loaded_ptx_from_cache = 1;+      } else {+        if (cfg->logging) {+          fprintf(stderr, "Failed!\n");+        }+        free(ptx);+        ptx = NULL;+      }+    }+  }++  if (ptx == NULL) {+    char* problem = cuda_nvrtc_build(src, (const char**)opts, n_opts, &ptx);+    if (problem != NULL) {+      free(src);+      return problem;+    }+  }++  if (cfg->dump_ptx_to != NULL) {+    dump_file(cfg->dump_ptx_to, ptx, strlen(ptx));+  }++  if (!loaded_ptx_from_cache) {+    CUDA_SUCCEED_FATAL(cuModuleLoadData(&ctx->module, ptx));+  }++  if (cache_fname != NULL && !loaded_ptx_from_cache) {+    if (cfg->logging) {+      fprintf(stderr, "Caching PTX in %s...\n", cache_fname);+    }+    errno = 0;+    if (cache_store(cache_fname, &h, (const unsigned char*)ptx, strlen(ptx)) != 0) {+      fprintf(stderr, "Failed to cache PTX: %s\n", strerror(errno));+    }+  }++  for (size_t i = 0; i < n_opts; i++) {+    free((char *)opts[i]);+  }+  free(opts);++  free(ptx);+  if (src != NULL) {+    free(src);+  }++  return NULL;+}++static char* cuda_setup(struct futhark_context *ctx, const char *src_fragments[],+                        const char *extra_opts[], const char* cache_fname) {+  CUDA_SUCCEED_FATAL(cuInit(0));++  if (cuda_device_setup(ctx) != 0) {+    futhark_panic(-1, "No suitable CUDA device found.\n");+  }+  CUDA_SUCCEED_FATAL(cuCtxCreate(&ctx->cu_ctx, 0, ctx->dev));++  free_list_init(&ctx->cu_free_list);++  ctx->max_shared_memory = device_query(ctx->dev, MAX_SHARED_MEMORY_PER_BLOCK);+  ctx->max_block_size = device_query(ctx->dev, MAX_THREADS_PER_BLOCK);+  ctx->max_grid_size = device_query(ctx->dev, MAX_GRID_DIM_X);+  ctx->max_tile_size = sqrt(ctx->max_block_size);+  ctx->max_threshold = 0;+  ctx->max_bespoke = 0;+  ctx->lockstep_width = device_query(ctx->dev, WARP_SIZE);++  cuda_size_setup(ctx);+  return cuda_module_setup(ctx, src_fragments, extra_opts, cache_fname);+}++// Count up the runtime all the profiling_records that occured during execution.+// Also clears the buffer of profiling_records.+static cudaError_t cuda_tally_profiling_records(struct futhark_context *ctx) {+  cudaError_t err;+  for (int i = 0; i < ctx->profiling_records_used; i++) {+    struct profiling_record record = ctx->profiling_records[i];++    float ms;+    if ((err = cudaEventElapsedTime(&ms, record.events[0], record.events[1])) != cudaSuccess) {+      return err;+    }++    // CUDA provides milisecond resolution, but we want microseconds.+    *record.runs += 1;+    *record.runtime += ms*1000;++    if ((err = cudaEventDestroy(record.events[0])) != cudaSuccess) {+      return err;+    }+    if ((err = cudaEventDestroy(record.events[1])) != cudaSuccess) {+      return err;+    }++    free(record.events);+  }++  ctx->profiling_records_used = 0;++  return cudaSuccess;+}++// Returns pointer to two events.+static cudaEvent_t* cuda_get_events(struct futhark_context *ctx, int *runs, int64_t *runtime) {+  if (ctx->profiling_records_used == ctx->profiling_records_capacity) {+    ctx->profiling_records_capacity *= 2;+    ctx->profiling_records =+      realloc(ctx->profiling_records,+              ctx->profiling_records_capacity *+              sizeof(struct profiling_record));+  }+  cudaEvent_t *events = calloc(2, sizeof(cudaEvent_t));+  cudaEventCreate(&events[0]);+  cudaEventCreate(&events[1]);+  ctx->profiling_records[ctx->profiling_records_used].events = events;+  ctx->profiling_records[ctx->profiling_records_used].runs = runs;+  ctx->profiling_records[ctx->profiling_records_used].runtime = runtime;+  ctx->profiling_records_used++;+  return events;+}++static CUresult cuda_alloc(struct futhark_context *ctx, FILE *log,+                           size_t min_size, const char *tag,+                           CUdeviceptr *mem_out, size_t *size_out) {+  if (min_size < sizeof(int)) {+    min_size = sizeof(int);+  }++  if (free_list_find(&ctx->cu_free_list, min_size, tag, size_out, (fl_mem*)mem_out) == 0) {+    if (*size_out >= min_size) {+      if (ctx->cfg->debugging) {+        fprintf(log, "No need to allocate: Found a block in the free list.\n");+      }+      return CUDA_SUCCESS;+    } else {+      if (ctx->cfg->debugging) {+        fprintf(log, "Found a free block, but it was too small.\n");+      }++      CUresult res = cuMemFree(*mem_out);+      if (res != CUDA_SUCCESS) {+        return res;+      }+    }+  }++  *size_out = min_size;++  if (ctx->cfg->debugging) {+    fprintf(log, "Actually allocating the desired block.\n");+  }++  CUresult res = cuMemAlloc(mem_out, min_size);+  while (res == CUDA_ERROR_OUT_OF_MEMORY) {+    CUdeviceptr mem;+    if (free_list_first(&ctx->cu_free_list, (fl_mem*)&mem) == 0) {+      res = cuMemFree(mem);+      if (res != CUDA_SUCCESS) {+        return res;+      }+    } else {+      break;+    }+    res = cuMemAlloc(mem_out, min_size);+  }++  return res;+}++static CUresult cuda_free(struct futhark_context *ctx,+                          CUdeviceptr mem, size_t size, const char *tag) {+  free_list_insert(&ctx->cu_free_list, size, (fl_mem)mem, tag);+  return CUDA_SUCCESS;+}++static CUresult cuda_free_all(struct futhark_context *ctx) {+  CUdeviceptr mem;+  free_list_pack(&ctx->cu_free_list);+  while (free_list_first(&ctx->cu_free_list, (fl_mem*)&mem) == 0) {+    CUresult res = cuMemFree(mem);+    if (res != CUDA_SUCCESS) {+      return res;+    }+  }++  return CUDA_SUCCESS;+}++int futhark_context_sync(struct futhark_context* ctx) {+  CUDA_SUCCEED_OR_RETURN(cuCtxPushCurrent(ctx->cu_ctx));+  CUDA_SUCCEED_OR_RETURN(cuCtxSynchronize());+  if (ctx->failure_is_an_option) {+    // Check for any delayed error.+    int32_t failure_idx;+    CUDA_SUCCEED_OR_RETURN(+                           cuMemcpyDtoH(&failure_idx,+                                        ctx->global_failure,+                                        sizeof(int32_t)));+    ctx->failure_is_an_option = 0;++    if (failure_idx >= 0) {+      // We have to clear global_failure so that the next entry point+      // is not considered a failure from the start.+      int32_t no_failure = -1;+      CUDA_SUCCEED_OR_RETURN(+                             cuMemcpyHtoD(ctx->global_failure,+                                          &no_failure,+                                          sizeof(int32_t)));++      int64_t args[max_failure_args+1];+      CUDA_SUCCEED_OR_RETURN(+                             cuMemcpyDtoH(&args,+                                          ctx->global_failure_args,+                                          sizeof(args)));++      ctx->error = get_failure_msg(failure_idx, args);++      return FUTHARK_PROGRAM_ERROR;+    }+  }+  CUDA_SUCCEED_OR_RETURN(cuCtxPopCurrent(&ctx->cu_ctx));+  return 0;+}++int backend_context_setup(struct futhark_context* ctx) {+  ctx->profiling_records_capacity = 200;+  ctx->profiling_records_used = 0;+  ctx->profiling_records =+    malloc(ctx->profiling_records_capacity *+           sizeof(struct profiling_record));+  ctx->failure_is_an_option = 0;+  ctx->total_runs = 0;+  ctx->total_runtime = 0;+  ctx->peak_mem_usage_device = 0;+  ctx->cur_mem_usage_device = 0;++  ctx->error = cuda_setup(ctx, cuda_program, ctx->cfg->nvrtc_opts, ctx->cfg->cache_fname);++  if (ctx->error != NULL) {+    futhark_panic(1, "%s\n", ctx->error);+  }++  int32_t no_error = -1;+  CUDA_SUCCEED_FATAL(cuMemAlloc(&ctx->global_failure, sizeof(no_error)));+  CUDA_SUCCEED_FATAL(cuMemcpyHtoD(ctx->global_failure, &no_error, sizeof(no_error)));+  // The +1 is to avoid zero-byte allocations.+  CUDA_SUCCEED_FATAL(cuMemAlloc(&ctx->global_failure_args, sizeof(int64_t)*(max_failure_args+1)));++  set_tuning_params(ctx);+  return 0;+}++void backend_context_teardown(struct futhark_context* ctx) {+  cuMemFree(ctx->global_failure);+  cuMemFree(ctx->global_failure_args);+  CUDA_SUCCEED_FATAL(cuda_free_all(ctx));+  (void)cuda_tally_profiling_records(ctx);+  free(ctx->profiling_records);+  CUDA_SUCCEED_FATAL(cuModuleUnload(ctx->module));+  CUDA_SUCCEED_FATAL(cuCtxDestroy(ctx->cu_ctx));+}++// End of backends/cuda.h.

@@ -0,0 +1,95 @@+// Start of backends/multicore.h++struct futhark_context_config {+  int in_use;+  int debugging;+  int profiling;+  int logging;+  const char *cache_fname;+  int num_tuning_params;+  int64_t *tuning_params;+  const char** tuning_param_names;+  const char** tuning_param_vars;+  const char** tuning_param_classes;+  // Uniform fields above.++  int num_threads;+};++static void backend_context_config_setup(struct futhark_context_config* cfg) {+  cfg->num_threads = 0;+}++static void backend_context_config_teardown(struct futhark_context_config* cfg) {+  (void)cfg;+}++void futhark_context_config_set_num_threads(struct futhark_context_config *cfg, int n) {+  cfg->num_threads = n;+}++int futhark_context_config_set_tuning_param(struct futhark_context_config* cfg, const char *param_name, size_t param_value) {+  (void)cfg; (void)param_name; (void)param_value;+  return 1;+}++struct futhark_context {+  struct futhark_context_config* cfg;+  int detail_memory;+  int debugging;+  int profiling;+  int profiling_paused;+  int logging;+  lock_t lock;+  char *error;+  lock_t error_lock;+  FILE *log;+  struct constants *constants;+  struct free_list free_list;+  int64_t peak_mem_usage_default;+  int64_t cur_mem_usage_default;+  struct program* program;+  // Uniform fields above.++  struct scheduler scheduler;+  int total_runs;+  long int total_runtime;+  int64_t tuning_timing;+  int64_t tuning_iter;+};++int backend_context_setup(struct futhark_context* ctx) {+  // Initialize rand()+  fast_srand(time(0));++  int tune_kappa = 0;+  double kappa = 5.1f * 1000;++  if (tune_kappa) {+    if (determine_kappa(&kappa) != 0) {+      ctx->error = strdup("Failed to determine kappa.");+      return 1;+    }+  }++  if (scheduler_init(&ctx->scheduler,+                     ctx->cfg->num_threads > 0 ?+                     ctx->cfg->num_threads : num_processors(),+                     kappa) != 0) {+    ctx->error = strdup("Failed to initialise scheduler.");+    return 1;+  }++  return 0;+}++void backend_context_teardown(struct futhark_context* ctx) {+  (void)scheduler_destroy(&ctx->scheduler);+}++int futhark_context_sync(struct futhark_context* ctx) {+  (void)ctx;+  return 0;+}++// End of backends/multicore.h

@@ -0,0 +1,1232 @@+// Start of backends/opencl.h++// Forward declarations.+struct opencl_device_option;+// Invoked by setup_opencl() after the platform and device has been+// found, but before the program is loaded.  Its intended use is to+// tune constants based on the selected platform and device.+static void post_opencl_setup(struct futhark_context*, struct opencl_device_option*);+static void set_tuning_params(struct futhark_context* ctx);+static char* get_failure_msg(int failure_idx, int64_t args[]);++#define OPENCL_SUCCEED_FATAL(e) opencl_succeed_fatal(e, #e, __FILE__, __LINE__)+#define OPENCL_SUCCEED_NONFATAL(e) opencl_succeed_nonfatal(e, #e, __FILE__, __LINE__)+// Take care not to override an existing error.+#define OPENCL_SUCCEED_OR_RETURN(e) {           \+    char *serror = OPENCL_SUCCEED_NONFATAL(e);  \+    if (serror) {                               \+      if (!ctx->error) {                        \+        ctx->error = serror;                    \+        return bad;                             \+      } else {                                  \+        free(serror);                           \+      }                                         \+    }                                           \+  }++// OPENCL_SUCCEED_OR_RETURN returns the value of the variable 'bad' in+// scope.  By default, it will be this one.  Create a local variable+// of some other type if needed.  This is a bit of a hack, but it+// saves effort in the code generator.+static const int bad = 1;++static const char* opencl_error_string(cl_int err) {+  switch (err) {+  case CL_SUCCESS:                            return "Success!";+  case CL_DEVICE_NOT_FOUND:                   return "Device not found.";+  case CL_DEVICE_NOT_AVAILABLE:               return "Device not available";+  case CL_COMPILER_NOT_AVAILABLE:             return "Compiler not available";+  case CL_MEM_OBJECT_ALLOCATION_FAILURE:      return "Memory object allocation failure";+  case CL_OUT_OF_RESOURCES:                   return "Out of resources";+  case CL_OUT_OF_HOST_MEMORY:                 return "Out of host memory";+  case CL_PROFILING_INFO_NOT_AVAILABLE:       return "Profiling information not available";+  case CL_MEM_COPY_OVERLAP:                   return "Memory copy overlap";+  case CL_IMAGE_FORMAT_MISMATCH:              return "Image format mismatch";+  case CL_IMAGE_FORMAT_NOT_SUPPORTED:         return "Image format not supported";+  case CL_BUILD_PROGRAM_FAILURE:              return "Program build failure";+  case CL_MAP_FAILURE:                        return "Map failure";+  case CL_INVALID_VALUE:                      return "Invalid value";+  case CL_INVALID_DEVICE_TYPE:                return "Invalid device type";+  case CL_INVALID_PLATFORM:                   return "Invalid platform";+  case CL_INVALID_DEVICE:                     return "Invalid device";+  case CL_INVALID_CONTEXT:                    return "Invalid context";+  case CL_INVALID_QUEUE_PROPERTIES:           return "Invalid queue properties";+  case CL_INVALID_COMMAND_QUEUE:              return "Invalid command queue";+  case CL_INVALID_HOST_PTR:                   return "Invalid host pointer";+  case CL_INVALID_MEM_OBJECT:                 return "Invalid memory object";+  case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR:    return "Invalid image format descriptor";+  case CL_INVALID_IMAGE_SIZE:                 return "Invalid image size";+  case CL_INVALID_SAMPLER:                    return "Invalid sampler";+  case CL_INVALID_BINARY:                     return "Invalid binary";+  case CL_INVALID_BUILD_OPTIONS:              return "Invalid build options";+  case CL_INVALID_PROGRAM:                    return "Invalid program";+  case CL_INVALID_PROGRAM_EXECUTABLE:         return "Invalid program executable";+  case CL_INVALID_KERNEL_NAME:                return "Invalid kernel name";+  case CL_INVALID_KERNEL_DEFINITION:          return "Invalid kernel definition";+  case CL_INVALID_KERNEL:                     return "Invalid kernel";+  case CL_INVALID_ARG_INDEX:                  return "Invalid argument index";+  case CL_INVALID_ARG_VALUE:                  return "Invalid argument value";+  case CL_INVALID_ARG_SIZE:                   return "Invalid argument size";+  case CL_INVALID_KERNEL_ARGS:                return "Invalid kernel arguments";+  case CL_INVALID_WORK_DIMENSION:             return "Invalid work dimension";+  case CL_INVALID_WORK_GROUP_SIZE:            return "Invalid work group size";+  case CL_INVALID_WORK_ITEM_SIZE:             return "Invalid work item size";+  case CL_INVALID_GLOBAL_OFFSET:              return "Invalid global offset";+  case CL_INVALID_EVENT_WAIT_LIST:            return "Invalid event wait list";+  case CL_INVALID_EVENT:                      return "Invalid event";+  case CL_INVALID_OPERATION:                  return "Invalid operation";+  case CL_INVALID_GL_OBJECT:                  return "Invalid OpenGL object";+  case CL_INVALID_BUFFER_SIZE:                return "Invalid buffer size";+  case CL_INVALID_MIP_LEVEL:                  return "Invalid mip-map level";+  default:                                    return "Unknown";+  }+}++static void opencl_succeed_fatal(cl_int ret,+                                 const char *call,+                                 const char *file,+                                 int line) {+  if (ret != CL_SUCCESS) {+    futhark_panic(-1, "%s:%d: OpenCL call\n  %s\nfailed with error code %d (%s)\n",+                  file, line, call, ret, opencl_error_string(ret));+  }+}++static char* opencl_succeed_nonfatal(cl_int ret,+                                     const char *call,+                                     const char *file,+                                     int line) {+  if (ret != CL_SUCCESS) {+    return msgprintf("%s:%d: OpenCL call\n  %s\nfailed with error code %d (%s)\n",+                     file, line, call, ret, opencl_error_string(ret));+  } else {+    return NULL;+  }+}++struct futhark_context_config {+  int in_use;+  int debugging;+  int profiling;+  int logging;+  const char *cache_fname;+  int num_tuning_params;+  int64_t *tuning_params;+  const char** tuning_param_names;+  const char** tuning_param_vars;+  const char** tuning_param_classes;+  // Uniform fields above.++  int preferred_device_num;+  const char *preferred_platform;+  const char *preferred_device;+  int ignore_blacklist;++  const char* dump_program_to;+  const char* load_program_from;+  const char* dump_binary_to;+  const char* load_binary_from;++  size_t default_group_size;+  size_t default_num_groups;+  size_t default_tile_size;+  size_t default_reg_tile_size;+  size_t default_threshold;++  int default_group_size_changed;+  int default_tile_size_changed;+  int num_build_opts;+  const char **build_opts;++  cl_command_queue queue;+  int queue_set;+};++static void backend_context_config_setup(struct futhark_context_config* cfg) {+  cfg->num_build_opts = 0;+  cfg->build_opts = (const char**) malloc(sizeof(const char*));+  cfg->build_opts[0] = NULL;+  cfg->preferred_device_num = 0;+  cfg->preferred_platform = "";+  cfg->preferred_device = "";+  cfg->ignore_blacklist = 0;+  cfg->dump_program_to = NULL;+  cfg->load_program_from = NULL;+  cfg->dump_binary_to = NULL;+  cfg->load_binary_from = NULL;++  // The following are dummy sizes that mean the concrete defaults+  // will be set during initialisation via hardware-inspection-based+  // heuristics.+  cfg->default_group_size = 0;+  cfg->default_num_groups = 0;+  cfg->default_tile_size = 0;+  cfg->default_reg_tile_size = 0;+  cfg->default_threshold = 0;++  cfg->default_group_size_changed = 0;+  cfg->default_tile_size_changed = 0;++  cfg->queue_set = 0;+}++static void backend_context_config_teardown(struct futhark_context_config* cfg) {+  free(cfg->build_opts);+}++void futhark_context_config_add_build_option(struct futhark_context_config* cfg, const char *opt) {+  cfg->build_opts[cfg->num_build_opts] = opt;+  cfg->num_build_opts++;+  cfg->build_opts = (const char**) realloc(cfg->build_opts, (cfg->num_build_opts+1) * sizeof(const char*));+  cfg->build_opts[cfg->num_build_opts] = NULL;+}++void futhark_context_config_set_device(struct futhark_context_config *cfg, const char* s) {+  int x = 0;+  if (*s == '#') {+    s++;+    while (isdigit(*s)) {+      x = x * 10 + (*s++)-'0';+    }+    // Skip trailing spaces.+    while (isspace(*s)) {+      s++;+    }+  }+  cfg->preferred_device = s;+  cfg->preferred_device_num = x;+  cfg->ignore_blacklist = 1;+}++void futhark_context_config_set_platform(struct futhark_context_config *cfg, const char *s) {+  cfg->preferred_platform = s;+  cfg->ignore_blacklist = 1;+}++void futhark_context_config_set_command_queue(struct futhark_context_config *cfg, cl_command_queue q) {+  cfg->queue = q;+  cfg->queue_set = 1;+}++struct opencl_device_option {+  cl_platform_id platform;+  cl_device_id device;+  cl_device_type device_type;+  char *platform_name;+  char *device_name;+};++static char* opencl_platform_info(cl_platform_id platform,+                                  cl_platform_info param) {+  size_t req_bytes;+  char *info;++  OPENCL_SUCCEED_FATAL(clGetPlatformInfo(platform, param, 0, NULL, &req_bytes));++  info = (char*) malloc(req_bytes);++  OPENCL_SUCCEED_FATAL(clGetPlatformInfo(platform, param, req_bytes, info, NULL));++  return info;+}++static char* opencl_device_info(cl_device_id device,+                                cl_device_info param) {+  size_t req_bytes;+  char *info;++  OPENCL_SUCCEED_FATAL(clGetDeviceInfo(device, param, 0, NULL, &req_bytes));++  info = (char*) malloc(req_bytes);++  OPENCL_SUCCEED_FATAL(clGetDeviceInfo(device, param, req_bytes, info, NULL));++  return info;+}++static int is_blacklisted(const char *platform_name, const char *device_name,+                          const struct futhark_context_config *cfg) {+  if (strcmp(cfg->preferred_platform, "") != 0 ||+      strcmp(cfg->preferred_device, "") != 0) {+    return 0;+  } else if (strstr(platform_name, "Apple") != NULL &&+             strstr(device_name, "Intel(R) Core(TM)") != NULL) {+    return 1;+  } else {+    return 0;+  }+}++static void opencl_all_device_options(struct opencl_device_option **devices_out,+                                      size_t *num_devices_out) {+  size_t num_devices = 0, num_devices_added = 0;++  cl_platform_id *all_platforms;+  cl_uint *platform_num_devices;++  cl_uint num_platforms;++  // Find the number of platforms.+  OPENCL_SUCCEED_FATAL(clGetPlatformIDs(0, NULL, &num_platforms));++  // Make room for them.+  all_platforms = calloc(num_platforms, sizeof(cl_platform_id));+  platform_num_devices = calloc(num_platforms, sizeof(cl_uint));++  // Fetch all the platforms.+  OPENCL_SUCCEED_FATAL(clGetPlatformIDs(num_platforms, all_platforms, NULL));++  // Count the number of devices for each platform, as well as the+  // total number of devices.+  for (cl_uint i = 0; i < num_platforms; i++) {+    if (clGetDeviceIDs(all_platforms[i], CL_DEVICE_TYPE_ALL,+                       0, NULL, &platform_num_devices[i]) == CL_SUCCESS) {+      num_devices += platform_num_devices[i];+    } else {+      platform_num_devices[i] = 0;+    }+  }++  // Make room for all the device options.+  struct opencl_device_option *devices =+    calloc(num_devices, sizeof(struct opencl_device_option));++  // Loop through the platforms, getting information about their devices.+  for (cl_uint i = 0; i < num_platforms; i++) {+    cl_platform_id platform = all_platforms[i];+    cl_uint num_platform_devices = platform_num_devices[i];++    if (num_platform_devices == 0) {+      continue;+    }++    char *platform_name = opencl_platform_info(platform, CL_PLATFORM_NAME);+    cl_device_id *platform_devices =+      calloc(num_platform_devices, sizeof(cl_device_id));++    // Fetch all the devices.+    OPENCL_SUCCEED_FATAL(clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL,+                                        num_platform_devices, platform_devices, NULL));++    // Loop through the devices, adding them to the devices array.+    for (cl_uint i = 0; i < num_platform_devices; i++) {+      char *device_name = opencl_device_info(platform_devices[i], CL_DEVICE_NAME);+      devices[num_devices_added].platform = platform;+      devices[num_devices_added].device = platform_devices[i];+      OPENCL_SUCCEED_FATAL(clGetDeviceInfo(platform_devices[i], CL_DEVICE_TYPE,+                                           sizeof(cl_device_type),+                                           &devices[num_devices_added].device_type,+                                           NULL));+      // We don't want the structs to share memory, so copy the platform name.+      // Each device name is already unique.+      devices[num_devices_added].platform_name = strclone(platform_name);+      devices[num_devices_added].device_name = device_name;+      num_devices_added++;+    }+    free(platform_devices);+    free(platform_name);+  }+  free(all_platforms);+  free(platform_num_devices);++  *devices_out = devices;+  *num_devices_out = num_devices;+}++void futhark_context_config_select_device_interactively(struct futhark_context_config *cfg) {+  struct opencl_device_option *devices;+  size_t num_devices;++  opencl_all_device_options(&devices, &num_devices);++  printf("Choose OpenCL device:\n");+  const char *cur_platform = "";+  for (size_t i = 0; i < num_devices; i++) {+    struct opencl_device_option device = devices[i];+    if (strcmp(cur_platform, device.platform_name) != 0) {+      printf("Platform: %s\n", device.platform_name);+      cur_platform = device.platform_name;+    }+    printf("[%d] %s\n", (int)i, device.device_name);+  }++  int selection;+  printf("Choice: ");+  if (scanf("%d", &selection) == 1) {+    cfg->preferred_platform = "";+    cfg->preferred_device = "";+    cfg->preferred_device_num = selection;+    cfg->ignore_blacklist = 1;+  }++  // Free all the platform and device names.+  for (size_t j = 0; j < num_devices; j++) {+    free(devices[j].platform_name);+    free(devices[j].device_name);+  }+  free(devices);+}++void futhark_context_config_list_devices(struct futhark_context_config *cfg) {+  (void)cfg;+  struct opencl_device_option *devices;+  size_t num_devices;++  opencl_all_device_options(&devices, &num_devices);++  const char *cur_platform = "";+  for (size_t i = 0; i < num_devices; i++) {+    struct opencl_device_option device = devices[i];+    if (strcmp(cur_platform, device.platform_name) != 0) {+      printf("Platform: %s\n", device.platform_name);+      cur_platform = device.platform_name;+    }+    printf("[%d]: %s\n", (int)i, device.device_name);+  }++  // Free all the platform and device names.+  for (size_t j = 0; j < num_devices; j++) {+    free(devices[j].platform_name);+    free(devices[j].device_name);+  }+  free(devices);+}++void futhark_context_config_dump_program_to(struct futhark_context_config *cfg, const char *path) {+  cfg->dump_program_to = path;+}++void futhark_context_config_load_program_from(struct futhark_context_config *cfg, const char *path) {+  cfg->load_program_from = path;+}++void futhark_context_config_dump_binary_to(struct futhark_context_config *cfg, const char *path) {+  cfg->dump_binary_to = path;+}++void futhark_context_config_load_binary_from(struct futhark_context_config *cfg, const char *path) {+  cfg->load_binary_from = path;+}++void futhark_context_config_set_default_group_size(struct futhark_context_config *cfg, int size) {+  cfg->default_group_size = size;+  cfg->default_group_size_changed = 1;+}++void futhark_context_config_set_default_num_groups(struct futhark_context_config *cfg, int num) {+  cfg->default_num_groups = num;+}++void futhark_context_config_set_default_tile_size(struct futhark_context_config *cfg, int size) {+  cfg->default_tile_size = size;+  cfg->default_tile_size_changed = 1;+}++void futhark_context_config_set_default_reg_tile_size(struct futhark_context_config *cfg, int size) {+  cfg->default_reg_tile_size = size;+}++void futhark_context_config_set_default_threshold(struct futhark_context_config *cfg, int size) {+  cfg->default_threshold = size;+}++int futhark_context_config_set_tuning_param(struct futhark_context_config *cfg,+                                            const char *param_name,+                                            size_t new_value) {+  for (int i = 0; i < cfg->num_tuning_params; i++) {+    if (strcmp(param_name, cfg->tuning_param_names[i]) == 0) {+      cfg->tuning_params[i] = new_value;+      return 0;+    }+  }+  if (strcmp(param_name, "default_group_size") == 0) {+    cfg->default_group_size = new_value;+    return 0;+  }+  if (strcmp(param_name, "default_num_groups") == 0) {+    cfg->default_num_groups = new_value;+    return 0;+  }+  if (strcmp(param_name, "default_threshold") == 0) {+    cfg->default_threshold = new_value;+    return 0;+  }+  if (strcmp(param_name, "default_tile_size") == 0) {+    cfg->default_tile_size = new_value;+    return 0;+  }+  if (strcmp(param_name, "default_reg_tile_size") == 0) {+    cfg->default_reg_tile_size = new_value;+    return 0;+  }+  return 1;+}++// A record of something that happened.+struct profiling_record {+  cl_event *event;+  int *runs;+  int64_t *runtime;+};++struct futhark_context {+  struct futhark_context_config* cfg;+  int detail_memory;+  int debugging;+  int profiling;+  int profiling_paused;+  int logging;+  lock_t lock;+  char *error;+  lock_t error_lock;+  FILE *log;+  struct constants *constants;+  struct free_list free_list;+  int64_t peak_mem_usage_default;+  int64_t cur_mem_usage_default;+  struct program* program;++  // Common fields above.++  cl_mem global_failure;+  cl_mem global_failure_args;+  struct tuning_params tuning_params;+  // True if a potentially failing kernel has been enqueued.+  cl_int failure_is_an_option;+  int total_runs;+  long int total_runtime;+  int64_t peak_mem_usage_device;+  int64_t cur_mem_usage_device;++  cl_device_id device;+  cl_context ctx;+  cl_command_queue queue;+  cl_program clprogram;++  struct free_list cl_free_list;++  size_t max_group_size;+  size_t max_num_groups;+  size_t max_tile_size;+  size_t max_threshold;+  size_t max_local_memory;++  size_t lockstep_width;++  struct profiling_record *profiling_records;+  int profiling_records_capacity;+  int profiling_records_used;++};++static cl_build_status build_opencl_program(cl_program program, cl_device_id device, const char* options) {+  cl_int clBuildProgram_error = clBuildProgram(program, 1, &device, options, NULL, NULL);++  // Avoid termination due to CL_BUILD_PROGRAM_FAILURE+  if (clBuildProgram_error != CL_SUCCESS &&+      clBuildProgram_error != CL_BUILD_PROGRAM_FAILURE) {+    OPENCL_SUCCEED_FATAL(clBuildProgram_error);+  }++  cl_build_status build_status;+  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;+    OPENCL_SUCCEED_FATAL(clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size));++    build_log = (char*) malloc(ret_val_size+1);+    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';++    fprintf(stderr, "Build log:\n%s\n", build_log);++    free(build_log);+  }++  return build_status;+}++static char* mk_compile_opts(struct futhark_context *ctx,+                             const char *extra_build_opts[],+                             struct opencl_device_option device_option) {+  int compile_opts_size = 1024;++  for (int i = 0; i < ctx->cfg->num_tuning_params; i++) {+    compile_opts_size += strlen(ctx->cfg->tuning_param_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 = (char*) malloc(compile_opts_size);++  int w = snprintf(compile_opts, compile_opts_size,+                   "-DLOCKSTEP_WIDTH=%d ",+                   (int)ctx->lockstep_width);++  w += snprintf(compile_opts+w, compile_opts_size-w,+                "-D%s=%d ",+                "max_group_size",+                (int)ctx->max_group_size);++  for (int i = 0; i < ctx->cfg->num_tuning_params; i++) {+    w += snprintf(compile_opts+w, compile_opts_size-w,+                  "-D%s=%d ",+                  ctx->cfg->tuning_param_vars[i],+                  (int)ctx->cfg->tuning_params[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]);+  }++  // Oclgrind claims to support cl_khr_fp16, but this is not actually+  // the case.+  if (strcmp(device_option.platform_name, "Oclgrind") == 0) {+    w += snprintf(compile_opts+w, compile_opts_size-w, "-DEMULATE_F16 ");+  }++  return compile_opts;+}++// Count up the runtime all the profiling_records that occured during execution.+// Also clears the buffer of profiling_records.+static cl_int opencl_tally_profiling_records(struct futhark_context *ctx) {+  cl_int err;+  for (int i = 0; i < ctx->profiling_records_used; i++) {+    struct profiling_record record = ctx->profiling_records[i];++    cl_ulong start_t, end_t;++    if ((err = clGetEventProfilingInfo(*record.event,+                                       CL_PROFILING_COMMAND_START,+                                       sizeof(start_t),+                                       &start_t,+                                       NULL)) != CL_SUCCESS) {+      return err;+    }++    if ((err = clGetEventProfilingInfo(*record.event,+                                       CL_PROFILING_COMMAND_END,+                                       sizeof(end_t),+                                       &end_t,+                                       NULL)) != CL_SUCCESS) {+      return err;+    }++    // OpenCL provides nanosecond resolution, but we want+    // microseconds.+    *record.runs += 1;+    *record.runtime += (end_t - start_t)/1000;++    if ((err = clReleaseEvent(*record.event)) != CL_SUCCESS) {+      return err;+    }+    free(record.event);+  }++  ctx->profiling_records_used = 0;++  return CL_SUCCESS;+}++// If profiling, produce an event associated with a profiling record.+static cl_event* opencl_get_event(struct futhark_context *ctx, int *runs, int64_t *runtime) {+  if (ctx->profiling_records_used == ctx->profiling_records_capacity) {+    ctx->profiling_records_capacity *= 2;+    ctx->profiling_records =+      realloc(ctx->profiling_records,+              ctx->profiling_records_capacity *+              sizeof(struct profiling_record));+  }+  cl_event *event = malloc(sizeof(cl_event));+  ctx->profiling_records[ctx->profiling_records_used].event = event;+  ctx->profiling_records[ctx->profiling_records_used].runs = runs;+  ctx->profiling_records[ctx->profiling_records_used].runtime = runtime;+  ctx->profiling_records_used++;+  return event;+}++// Allocate memory from driver. The problem is that OpenCL may perform+// lazy allocation, so we cannot know whether an allocation succeeded+// until the first time we try to use it.  Hence we immediately+// perform a write to see if the allocation succeeded.  This is slow,+// but the assumption is that this operation will be rare (most things+// will go through the free list).+static int opencl_alloc_actual(struct futhark_context *ctx, size_t size, cl_mem *mem_out) {+  int error;+  *mem_out = clCreateBuffer(ctx->ctx, CL_MEM_READ_WRITE, size, NULL, &error);++  if (error != CL_SUCCESS) {+    return error;+  }++  int x = 2;+  error = clEnqueueWriteBuffer(ctx->queue, *mem_out,+                               CL_TRUE,+                               0, sizeof(x), &x,+                               0, NULL, NULL);++  // No need to wait for completion here. clWaitForEvents() cannot+  // return mem object allocation failures. This implies that the+  // buffer is faulted onto the device on enqueue. (Observation by+  // Andreas Kloeckner.)++  return error;+}++static int opencl_alloc(struct futhark_context *ctx, FILE *log,+                        size_t min_size, const char *tag,+                        cl_mem *mem_out, size_t *size_out) {+  (void)tag;+  if (min_size < sizeof(int)) {+    min_size = sizeof(int);+  }++  cl_mem* memptr;+  if (free_list_find(&ctx->cl_free_list, min_size, tag, size_out, (fl_mem*)&memptr) == 0) {+    // Successfully found a free block.  Is it big enough?+    if (*size_out >= min_size) {+      if (ctx->cfg->debugging) {+        fprintf(log, "No need to allocate: Found a block in the free list.\n");+      }+      *mem_out = *memptr;+      free(memptr);+      return CL_SUCCESS;+    } else {+      if (ctx->cfg->debugging) {+        fprintf(log, "Found a free block, but it was too small.\n");+      }+      int error = clReleaseMemObject(*memptr);+      free(*memptr);+      if (error != CL_SUCCESS) {+        return error;+      }+    }+  }++  *size_out = min_size;++  // We have to allocate a new block from the driver.  If the+  // allocation does not succeed, then we might be in an out-of-memory+  // situation.  We now start freeing things from the free list until+  // we think we have freed enough that the allocation will succeed.+  // Since we don't know how far the allocation is from fitting, we+  // have to check after every deallocation.  This might be pretty+  // expensive.  Let's hope that this case is hit rarely.++  if (ctx->cfg->debugging) {+    fprintf(log, "Actually allocating the desired block.\n");+  }++  int error = opencl_alloc_actual(ctx, min_size, mem_out);++  while (error == CL_MEM_OBJECT_ALLOCATION_FAILURE) {+    if (ctx->cfg->debugging) {+      fprintf(log, "Out of OpenCL memory: releasing entry from the free list...\n");+    }+    cl_mem* memptr;+    if (free_list_first(&ctx->cl_free_list, (fl_mem*)&memptr) == 0) {+      cl_mem mem = *memptr;+      free(memptr);+      error = clReleaseMemObject(mem);+      if (error != CL_SUCCESS) {+        return error;+      }+    } else {+      break;+    }+    error = opencl_alloc_actual(ctx, min_size, mem_out);+  }++  return error;+}++static int opencl_free(struct futhark_context *ctx,+                       cl_mem mem, size_t size, const char *tag) {+  cl_mem* memptr = malloc(sizeof(cl_mem));+  *memptr = mem;+  free_list_insert(&ctx->cl_free_list, size, (fl_mem)memptr, tag);+  return CL_SUCCESS;+}++static int opencl_free_all(struct futhark_context *ctx) {+  free_list_pack(&ctx->cl_free_list);+  cl_mem* memptr;+  while (free_list_first(&ctx->cl_free_list, (fl_mem*)&memptr) == 0) {+    cl_mem mem = *memptr;+    free(memptr);+    int error = clReleaseMemObject(mem);+    if (error != CL_SUCCESS) {+      return error;+    }+  }++  return CL_SUCCESS;+}++int futhark_context_sync(struct futhark_context* ctx) {+  // Check for any delayed error.+  cl_int failure_idx = -1;+  if (ctx->failure_is_an_option) {+    OPENCL_SUCCEED_OR_RETURN(+                             clEnqueueReadBuffer(ctx->queue,+                                                 ctx->global_failure,+                                                 CL_FALSE,+                                                 0, sizeof(cl_int), &failure_idx,+                                                 0, NULL, NULL));+    ctx->failure_is_an_option = 0;+  }++  OPENCL_SUCCEED_OR_RETURN(clFinish(ctx->queue));++  if (failure_idx >= 0) {+    // We have to clear global_failure so that the next entry point+    // is not considered a failure from the start.+    cl_int no_failure = -1;+    OPENCL_SUCCEED_OR_RETURN(+                             clEnqueueWriteBuffer(ctx->queue, ctx->global_failure, CL_TRUE,+                                                  0, sizeof(cl_int), &no_failure,+                                                  0, NULL, NULL));++    int64_t args[max_failure_args+1];+    OPENCL_SUCCEED_OR_RETURN(+                             clEnqueueReadBuffer(ctx->queue,+                                                 ctx->global_failure_args,+                                                 CL_TRUE,+                                                 0, sizeof(args), &args,+                                                 0, NULL, NULL));++    ctx->error = get_failure_msg(failure_idx, args);++    return FUTHARK_PROGRAM_ERROR;+  }+  return 0;+}+++// We take as input several strings representing the program, because+// C does not guarantee that the compiler supports particularly large+// literals.  Notably, Visual C has a limit of 2048 characters.  The+// array must be NULL-terminated.+static void setup_opencl_with_command_queue(struct futhark_context *ctx,+                                            cl_command_queue queue,+                                            const char *srcs[],+                                            const char *extra_build_opts[],+                                            const char* cache_fname) {+  int error;++  free_list_init(&ctx->cl_free_list);+  ctx->queue = queue;++  OPENCL_SUCCEED_FATAL(clGetCommandQueueInfo(ctx->queue, CL_QUEUE_CONTEXT, sizeof(cl_context), &ctx->ctx, NULL));++  // Fill out the device info.  This is redundant work if we are+  // called from setup_opencl() (which is the common case), but I+  // doubt it matters much.+  struct opencl_device_option device_option;+  OPENCL_SUCCEED_FATAL(clGetCommandQueueInfo(ctx->queue, CL_QUEUE_DEVICE,+                                             sizeof(cl_device_id),+                                             &device_option.device,+                                             NULL));+  OPENCL_SUCCEED_FATAL(clGetDeviceInfo(device_option.device, CL_DEVICE_PLATFORM,+                                       sizeof(cl_platform_id),+                                       &device_option.platform,+                                       NULL));+  OPENCL_SUCCEED_FATAL(clGetDeviceInfo(device_option.device, CL_DEVICE_TYPE,+                                       sizeof(cl_device_type),+                                       &device_option.device_type,+                                       NULL));+  device_option.platform_name = opencl_platform_info(device_option.platform, CL_PLATFORM_NAME);+  device_option.device_name = opencl_device_info(device_option.device, CL_DEVICE_NAME);++  ctx->device = device_option.device;++  if (f64_required) {+    cl_uint supported;+    OPENCL_SUCCEED_FATAL(clGetDeviceInfo(device_option.device, CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE,+                                         sizeof(cl_uint), &supported, NULL));+    if (!supported) {+      futhark_panic(1, "Program uses double-precision floats, but this is not supported on the chosen device: %s\n",+                    device_option.device_name);+    }+  }++  size_t max_group_size;+  OPENCL_SUCCEED_FATAL(clGetDeviceInfo(device_option.device, CL_DEVICE_MAX_WORK_GROUP_SIZE,+                                       sizeof(size_t), &max_group_size, NULL));++  size_t max_tile_size = sqrt(max_group_size);++  cl_ulong max_local_memory;+  OPENCL_SUCCEED_FATAL(clGetDeviceInfo(device_option.device, CL_DEVICE_LOCAL_MEM_SIZE,+                                       sizeof(size_t), &max_local_memory, NULL));++  // Futhark reserves 4 bytes for bookkeeping information.+  max_local_memory -= 4;++  // The OpenCL implementation may reserve some local memory bytes for+  // various purposes.  In principle, we should use+  // clGetKernelWorkGroupInfo() to figure out for each kernel how much+  // is actually available, but our current code generator design+  // makes this infeasible.  Instead, we have this nasty hack where we+  // arbitrarily subtract some bytes, based on empirical measurements+  // (but which might be arbitrarily wrong).  Fortunately, we rarely+  // try to really push the local memory usage.+  if (strstr(device_option.platform_name, "NVIDIA CUDA") != NULL) {+    max_local_memory -= 12;+  } else if (strstr(device_option.platform_name, "AMD") != NULL) {+    max_local_memory -= 16;+  }++  // Make sure this function is defined.+  post_opencl_setup(ctx, &device_option);++  if (max_group_size < ctx->cfg->default_group_size) {+    if (ctx->cfg->default_group_size_changed) {+      fprintf(stderr, "Note: Device limits default group size to %zu (down from %zu).\n",+              max_group_size, ctx->cfg->default_group_size);+    }+    ctx->cfg->default_group_size = max_group_size;+  }++  if (max_tile_size < ctx->cfg->default_tile_size) {+    if (ctx->cfg->default_tile_size_changed) {+      fprintf(stderr, "Note: Device limits default tile size to %zu (down from %zu).\n",+              max_tile_size, ctx->cfg->default_tile_size);+    }+    ctx->cfg->default_tile_size = max_tile_size;+  }++  ctx->max_group_size = max_group_size;+  ctx->max_tile_size = max_tile_size; // No limit.+  ctx->max_threshold = ctx->max_num_groups = 0; // No limit.+  ctx->max_local_memory = max_local_memory;++  // Now we go through all the sizes, clamp them to the valid range,+  // or set them to the default.+  for (int i = 0; i < ctx->cfg->num_tuning_params; i++) {+    const char *size_class = ctx->cfg->tuning_param_classes[i];+    int64_t *size_value = &ctx->cfg->tuning_params[i];+    const char* size_name = ctx->cfg->tuning_param_names[i];+    int64_t max_value = 0, default_value = 0;++    if (strstr(size_class, "group_size") == size_class) {+      max_value = max_group_size;+      default_value = ctx->cfg->default_group_size;+    } else if (strstr(size_class, "num_groups") == size_class) {+      max_value = max_group_size; // Futhark assumes this constraint.+      default_value = ctx->cfg->default_num_groups;+      // XXX: as a quick and dirty hack, use twice as many threads for+      // histograms by default.  We really should just be smarter+      // about sizes somehow.+      if (strstr(size_name, ".seghist_") != NULL) {+        default_value *= 2;+      }+    } else if (strstr(size_class, "tile_size") == size_class) {+      max_value = sqrt(max_group_size);+      default_value = ctx->cfg->default_tile_size;+    } else if (strstr(size_class, "reg_tile_size") == size_class) {+      max_value = 0; // No limit.+      default_value = ctx->cfg->default_reg_tile_size;+    } else if (strstr(size_class, "threshold") == size_class) {+      // Threshold can be as large as it takes.+      default_value = ctx->cfg->default_threshold;+    } else {+      // Bespoke sizes have no limit or default.+    }+    if (*size_value == 0) {+      *size_value = default_value;+    } else if (max_value > 0 && *size_value > max_value) {+      fprintf(stderr, "Note: Device limits %s to %d (down from %d)\n",+              size_name, (int)max_value, (int)*size_value);+      *size_value = max_value;+    }+  }++  if (ctx->lockstep_width == 0) {+    ctx->lockstep_width = 1;+  }++  if (ctx->cfg->logging) {+    fprintf(stderr, "Lockstep width: %d\n", (int)ctx->lockstep_width);+    fprintf(stderr, "Default group size: %d\n", (int)ctx->cfg->default_group_size);+    fprintf(stderr, "Default number of groups: %d\n", (int)ctx->cfg->default_num_groups);+  }++  char *compile_opts = mk_compile_opts(ctx, extra_build_opts, device_option);++  if (ctx->cfg->logging) {+    fprintf(stderr, "OpenCL compiler options: %s\n", compile_opts);+  }++  char *fut_opencl_src = NULL;+  cl_program prog;+  error = CL_SUCCESS;++  struct cache_hash h;++  int loaded_from_cache = 0;+  if (ctx->cfg->load_binary_from == NULL) {+    size_t src_size = 0;++    // Maybe we have to read OpenCL source from somewhere else (used for debugging).+    if (ctx->cfg->load_program_from != NULL) {+      fut_opencl_src = slurp_file(ctx->cfg->load_program_from, NULL);+      assert(fut_opencl_src != NULL);+    } else {+      // Construct the OpenCL source concatenating all the fragments.+      for (const char **src = srcs; src && *src; src++) {+        src_size += strlen(*src);+      }++      fut_opencl_src = (char*) malloc(src_size + 1);++      size_t n, i;+      for (i = 0, n = 0; srcs && srcs[i]; i++) {+        strncpy(fut_opencl_src+n, srcs[i], src_size-n);+        n += strlen(srcs[i]);+      }+      fut_opencl_src[src_size] = 0;+    }++    if (ctx->cfg->dump_program_to != NULL) {+      if (ctx->cfg->logging) {+        fprintf(stderr, "Dumping OpenCL source to %s...\n", ctx->cfg->dump_program_to);+      }++      dump_file(ctx->cfg->dump_program_to, fut_opencl_src, strlen(fut_opencl_src));+    }++    if (cache_fname != NULL) {+      if (ctx->cfg->logging) {+        fprintf(stderr, "Restoring cache from from %s...\n", cache_fname);+      }+      cache_hash_init(&h);+      cache_hash(&h, fut_opencl_src, strlen(fut_opencl_src));+      cache_hash(&h, compile_opts, strlen(compile_opts));++      unsigned char *buf;+      size_t bufsize;+      errno = 0;+      if (cache_restore(cache_fname, &h, &buf, &bufsize) != 0) {+        if (ctx->cfg->logging) {+          fprintf(stderr, "Failed to restore cache (errno: %s)\n", strerror(errno));+        }+      } else {+        if (ctx->cfg->logging) {+          fprintf(stderr, "Cache restored; loading OpenCL binary...\n");+        }++        cl_int status = 0;+        prog = clCreateProgramWithBinary(ctx->ctx, 1, &device_option.device,+                                         &bufsize, (const unsigned char**)&buf,+                                         &status, &error);+        if (status == CL_SUCCESS) {+          loaded_from_cache = 1;+          if (ctx->cfg->logging) {+            fprintf(stderr, "Loading succeeded.\n");+          }+        } else {+          if (ctx->cfg->logging) {+            fprintf(stderr, "Loading failed.\n");+          }+        }+      }+    }++    if (!loaded_from_cache) {+      if (ctx->cfg->logging) {+        fprintf(stderr, "Creating OpenCL program...\n");+      }++      const char* src_ptr[] = {fut_opencl_src};+      prog = clCreateProgramWithSource(ctx->ctx, 1, src_ptr, &src_size, &error);+      OPENCL_SUCCEED_FATAL(error);+    }+  } else {+    if (ctx->cfg->logging) {+      fprintf(stderr, "Loading OpenCL binary from %s...\n", ctx->cfg->load_binary_from);+    }+    size_t binary_size;+    unsigned char *fut_opencl_bin =+      (unsigned char*) slurp_file(ctx->cfg->load_binary_from, &binary_size);+    assert(fut_opencl_bin != NULL);+    const unsigned char *binaries[1] = { fut_opencl_bin };+    cl_int status = 0;++    prog = clCreateProgramWithBinary(ctx->ctx, 1, &device_option.device,+                                     &binary_size, binaries,+                                     &status, &error);++    OPENCL_SUCCEED_FATAL(status);+    OPENCL_SUCCEED_FATAL(error);+  }++  if (ctx->cfg->logging) {+    fprintf(stderr, "Building OpenCL program...\n");+  }+  OPENCL_SUCCEED_FATAL(build_opencl_program(prog, device_option.device, compile_opts));++  free(compile_opts);+  free(fut_opencl_src);++  size_t binary_size = 0;+  unsigned char *binary = NULL;+  int store_in_cache = cache_fname != NULL && !loaded_from_cache;+  if (store_in_cache || ctx->cfg->dump_binary_to != NULL) {+    OPENCL_SUCCEED_FATAL(clGetProgramInfo(prog, CL_PROGRAM_BINARY_SIZES,+                                          sizeof(size_t), &binary_size, NULL));+    binary = (unsigned char*) malloc(binary_size);+    OPENCL_SUCCEED_FATAL(clGetProgramInfo(prog, CL_PROGRAM_BINARIES,+                                          sizeof(unsigned char*), &binary, NULL));+  }++  if (store_in_cache) {+    if (ctx->cfg->logging) {+      fprintf(stderr, "Caching OpenCL binary in %s...\n", cache_fname);+    }+    if (cache_store(cache_fname, &h, binary, binary_size) != 0) {+      printf("Failed to cache binary: %s\n", strerror(errno));+    }+  }++  if (ctx->cfg->dump_binary_to != NULL) {+    if (ctx->cfg->logging) {+      fprintf(stderr, "Dumping OpenCL binary to %s...\n", ctx->cfg->dump_binary_to);+    }+    dump_file(ctx->cfg->dump_binary_to, binary, binary_size);+  }++  ctx->clprogram = prog;+}++static struct opencl_device_option get_preferred_device(const struct futhark_context_config *cfg) {+  struct opencl_device_option *devices;+  size_t num_devices;++  opencl_all_device_options(&devices, &num_devices);++  int num_device_matches = 0;++  for (size_t i = 0; i < num_devices; i++) {+    struct opencl_device_option device = devices[i];+    if (strstr(device.platform_name, cfg->preferred_platform) != NULL &&+        strstr(device.device_name, cfg->preferred_device) != NULL &&+        (cfg->ignore_blacklist ||+         !is_blacklisted(device.platform_name, device.device_name, cfg)) &&+        num_device_matches++ == cfg->preferred_device_num) {+      // Free all the platform and device names, except the ones we have chosen.+      for (size_t j = 0; j < num_devices; j++) {+        if (j != i) {+          free(devices[j].platform_name);+          free(devices[j].device_name);+        }+      }+      free(devices);+      return device;+    }+  }++  futhark_panic(1, "Could not find acceptable OpenCL device.\n");+  exit(1); // Never reached+}++static void setup_opencl(struct futhark_context *ctx,+                         const char *srcs[],+                         const char *extra_build_opts[],+                         const char* cache_fname) {+  struct opencl_device_option device_option = get_preferred_device(ctx->cfg);++  if (ctx->cfg->logging) {+    fprintf(stderr, "Using platform: %s\n", device_option.platform_name);+    fprintf(stderr, "Using device: %s\n", device_option.device_name);+  }++  // Note that NVIDIA's OpenCL requires the platform property+  cl_context_properties properties[] = {+    CL_CONTEXT_PLATFORM,+    (cl_context_properties)device_option.platform,+    0+  };++  cl_int clCreateContext_error;+  ctx->ctx = clCreateContext(properties, 1, &device_option.device, NULL, NULL, &clCreateContext_error);+  OPENCL_SUCCEED_FATAL(clCreateContext_error);++  cl_int clCreateCommandQueue_error;+  cl_command_queue queue =+    clCreateCommandQueue(ctx->ctx,+                         device_option.device,+                         ctx->cfg->profiling ? CL_QUEUE_PROFILING_ENABLE : 0,+                         &clCreateCommandQueue_error);+  OPENCL_SUCCEED_FATAL(clCreateCommandQueue_error);++  setup_opencl_with_command_queue(ctx, queue, srcs, extra_build_opts, cache_fname);+}++int backend_context_setup(struct futhark_context* ctx) {+  ctx->lockstep_width = 0; // Real value set later.+  ctx->profiling_records_capacity = 200;+  ctx->profiling_records_used = 0;+  ctx->profiling_records =+    malloc(ctx->profiling_records_capacity *+           sizeof(struct profiling_record));+  ctx->failure_is_an_option = 0;+  ctx->total_runs = 0;+  ctx->total_runtime = 0;+  ctx->peak_mem_usage_device = 0;+  ctx->cur_mem_usage_device = 0;++  if (ctx->cfg->queue_set) {+    setup_opencl_with_command_queue(ctx, ctx->cfg->queue, opencl_program, ctx->cfg->build_opts, ctx->cfg->cache_fname);+  } else {+    setup_opencl(ctx, opencl_program, ctx->cfg->build_opts, ctx->cfg->cache_fname);+  }++  cl_int error;+  cl_int no_error = -1;+  ctx->global_failure =+    clCreateBuffer(ctx->ctx,+                   CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,+                   sizeof(cl_int), &no_error, &error);+  OPENCL_SUCCEED_OR_RETURN(error);++  // The +1 is to avoid zero-byte allocations.+  ctx->global_failure_args =+    clCreateBuffer(ctx->ctx,+                   CL_MEM_READ_WRITE,+                   sizeof(int64_t)*(max_failure_args+1), NULL, &error);+  OPENCL_SUCCEED_OR_RETURN(error);++  set_tuning_params(ctx);+  return 0;+}++void backend_context_teardown(struct futhark_context* ctx) {+  OPENCL_SUCCEED_FATAL(clReleaseMemObject(ctx->global_failure));+  OPENCL_SUCCEED_FATAL(clReleaseMemObject(ctx->global_failure_args));+  (void)opencl_tally_profiling_records(ctx);+  free(ctx->profiling_records);+  (void)opencl_free_all(ctx);+  (void)clReleaseProgram(ctx->clprogram);+  (void)clReleaseCommandQueue(ctx->queue);+  (void)clReleaseContext(ctx->ctx);+}++cl_command_queue futhark_context_get_command_queue(struct futhark_context* ctx) {+  return ctx->queue;+}++// End of backends/opencl.h

@@ -72,18 +72,75 @@   } } -static void context_setup(struct futhark_context *ctx) {+struct futhark_context_config* futhark_context_config_new(void) {+  struct futhark_context_config* cfg = malloc(sizeof(struct futhark_context_config));+  if (cfg == NULL) {+    return NULL;+  }+  cfg->in_use = 0;+  cfg->debugging = 0;+  cfg->profiling = 0;+  cfg->logging = 0;+  cfg->cache_fname = NULL;+  cfg->num_tuning_params = num_tuning_params;+  cfg->tuning_params = malloc(cfg->num_tuning_params * sizeof(int64_t));+  memcpy(cfg->tuning_params, tuning_param_defaults,+         cfg->num_tuning_params * sizeof(int64_t));+  cfg->tuning_param_names = tuning_param_names;+  cfg->tuning_param_vars = tuning_param_vars;+  cfg->tuning_param_classes = tuning_param_classes;+  backend_context_config_setup(cfg);+  return cfg;+}++void futhark_context_config_free(struct futhark_context_config* cfg) {+  assert(!cfg->in_use);+  backend_context_config_teardown(cfg);+  free(cfg->tuning_params);+  free(cfg);+}++struct futhark_context* futhark_context_new(struct futhark_context_config* cfg) {+  struct futhark_context* ctx = malloc(sizeof(struct futhark_context));+  if (ctx == NULL) {+    return NULL;+  }+  assert(!cfg->in_use);+  ctx->cfg = cfg;+  ctx->cfg->in_use = 1;   create_lock(&ctx->error_lock);   create_lock(&ctx->lock);   free_list_init(&ctx->free_list);+  ctx->peak_mem_usage_default = 0;+  ctx->cur_mem_usage_default = 0;+  ctx->constants = malloc(sizeof(struct constants));+  ctx->detail_memory = cfg->debugging;+  ctx->debugging = cfg->debugging;+  ctx->logging = cfg->logging;+  ctx->profiling = cfg->profiling;+  ctx->profiling_paused = 0;+  ctx->error = NULL;+  ctx->log = stderr;+  if (backend_context_setup(ctx) == 0) {+    setup_program(ctx);+    init_constants(ctx);+    (void)futhark_context_clear_caches(ctx);+    (void)futhark_context_sync(ctx);+  }+  return ctx; } -static void context_teardown(struct futhark_context *ctx) {+void futhark_context_free(struct futhark_context* ctx) {   free_constants(ctx);+  teardown_program(ctx);+  backend_context_teardown(ctx);   free_all_in_free_list(ctx);   free_list_destroy(&ctx->free_list);+  free(ctx->constants);   free_lock(&ctx->lock);   free_lock(&ctx->error_lock);+  ctx->cfg->in_use = 0;+  free(ctx); }  // End of context.h

@@ -1,20 +1,29 @@ // Start of context_prototypes.h //-// Prototypes for the functions in context.h that need to be-// available very early.+// Prototypes for the functions in context.h, or that will be called+// from those functions, that need to be available very early.  struct futhark_context_config; struct futhark_context;  static void set_error(struct futhark_context* ctx, char *error); -// These are called in context new/free functions and contain shared setup.-static void context_setup(struct futhark_context *ctx);-static void context_teardown(struct futhark_context *ctx);+// These are called in context/config new/free functions and contain+// shared setup.  They are generated by the compiler itself.+static int init_constants(struct futhark_context*);+static int free_constants(struct futhark_context*);+static void setup_program(struct futhark_context* ctx);+static void teardown_program(struct futhark_context *ctx);  // Allocate host memory.  Must be freed with host_free(). static void host_alloc(struct futhark_context* ctx, size_t size, const char* tag, size_t* size_out, void** mem_out); // Allocate memory allocated with host_alloc(). static void host_free(struct futhark_context* ctx, size_t size, const char* tag, void* mem);++// Functions that must be defined by the backend.+static void backend_context_config_setup(struct futhark_context_config* cfg);+static void backend_context_config_teardown(struct futhark_context_config* cfg);+static int backend_context_setup(struct futhark_context *ctx);+static void backend_context_teardown(struct futhark_context *ctx);  // End of of context_prototypes.h

@@ -1,766 +0,0 @@-// Start of cuda.h.--#define CUDA_SUCCEED_FATAL(x) cuda_api_succeed_fatal(x, #x, __FILE__, __LINE__)-#define CUDA_SUCCEED_NONFATAL(x) cuda_api_succeed_nonfatal(x, #x, __FILE__, __LINE__)-#define NVRTC_SUCCEED_FATAL(x) nvrtc_api_succeed_fatal(x, #x, __FILE__, __LINE__)-#define NVRTC_SUCCEED_NONFATAL(x) nvrtc_api_succeed_nonfatal(x, #x, __FILE__, __LINE__)-// Take care not to override an existing error.-#define CUDA_SUCCEED_OR_RETURN(e) {               \-    char *serror = CUDA_SUCCEED_NONFATAL(e);      \-    if (serror) {                                 \-      if (!ctx->error) {                          \-        ctx->error = serror;                      \-        return bad;                               \-      } else {                                    \-        free(serror);                             \-      }                                           \-    }                                             \-  }--// CUDA_SUCCEED_OR_RETURN returns the value of the variable 'bad' in-// scope.  By default, it will be this one.  Create a local variable-// of some other type if needed.  This is a bit of a hack, but it-// saves effort in the code generator.-static const int bad = 1;--static inline void cuda_api_succeed_fatal(CUresult res, const char *call,-    const char *file, int line) {-  if (res != CUDA_SUCCESS) {-    const char *err_str;-    cuGetErrorString(res, &err_str);-    if (err_str == NULL) { err_str = "Unknown"; }-    futhark_panic(-1, "%s:%d: CUDA call\n  %s\nfailed with error code %d (%s)\n",-        file, line, call, res, err_str);-  }-}--static char* cuda_api_succeed_nonfatal(CUresult res, const char *call,-    const char *file, int line) {-  if (res != CUDA_SUCCESS) {-    const char *err_str;-    cuGetErrorString(res, &err_str);-    if (err_str == NULL) { err_str = "Unknown"; }-    return msgprintf("%s:%d: CUDA call\n  %s\nfailed with error code %d (%s)\n",-        file, line, call, res, err_str);-  } else {-    return NULL;-  }-}--static inline void nvrtc_api_succeed_fatal(nvrtcResult res, const char *call,-                                           const char *file, int line) {-  if (res != NVRTC_SUCCESS) {-    const char *err_str = nvrtcGetErrorString(res);-    futhark_panic(-1, "%s:%d: NVRTC call\n  %s\nfailed with error code %d (%s)\n",-        file, line, call, res, err_str);-  }-}--static char* nvrtc_api_succeed_nonfatal(nvrtcResult res, const char *call,-                                        const char *file, int line) {-  if (res != NVRTC_SUCCESS) {-    const char *err_str = nvrtcGetErrorString(res);-    return msgprintf("%s:%d: NVRTC call\n  %s\nfailed with error code %d (%s)\n",-                     file, line, call, res, err_str);-  } else {-    return NULL;-  }-}--struct cuda_config {-  int debugging;-  int logging;-  const char *preferred_device;-  int preferred_device_num;--  const char *dump_program_to;-  const char *load_program_from;--  const char *dump_ptx_to;-  const char *load_ptx_from;--  size_t default_block_size;-  size_t default_grid_size;-  size_t default_tile_size;-  size_t default_reg_tile_size;-  size_t default_threshold;--  int default_block_size_changed;-  int default_grid_size_changed;-  int default_tile_size_changed;--  int num_sizes;-  const char **size_names;-  const char **size_vars;-  int64_t *size_values;-  const char **size_classes;-};--static void cuda_config_init(struct cuda_config *cfg,-                             int num_sizes,-                             const char *size_names[],-                             const char *size_vars[],-                             int64_t *size_values,-                             const char *size_classes[]) {-  cfg->debugging = 0;-  cfg->logging = 0;-  cfg->preferred_device_num = 0;-  cfg->preferred_device = "";-  cfg->dump_program_to = NULL;-  cfg->load_program_from = NULL;--  cfg->dump_ptx_to = NULL;-  cfg->load_ptx_from = NULL;--  cfg->default_block_size = 256;-  cfg->default_grid_size = 0; // Set properly later.-  cfg->default_tile_size = 32;-  cfg->default_reg_tile_size = 2;-  cfg->default_threshold = 32*1024;--  cfg->default_block_size_changed = 0;-  cfg->default_grid_size_changed = 0;-  cfg->default_tile_size_changed = 0;--  cfg->num_sizes = num_sizes;-  cfg->size_names = size_names;-  cfg->size_vars = size_vars;-  cfg->size_values = size_values;-  cfg->size_classes = size_classes;-}--// A record of something that happened.-struct profiling_record {-  cudaEvent_t *events; // Points to two events.-  int *runs;-  int64_t *runtime;-};--struct cuda_context {-  CUdevice dev;-  CUcontext cu_ctx;-  CUmodule module;--  struct cuda_config cfg;--  struct free_list free_list;--  size_t max_block_size;-  size_t max_grid_size;-  size_t max_tile_size;-  size_t max_threshold;-  size_t max_shared_memory;-  size_t max_bespoke;--  size_t lockstep_width;--  struct profiling_record *profiling_records;-  int profiling_records_capacity;-  int profiling_records_used;-};--#define CU_DEV_ATTR(x) (CU_DEVICE_ATTRIBUTE_##x)-#define device_query(dev,attrib) _device_query(dev, CU_DEV_ATTR(attrib))-static int _device_query(CUdevice dev, CUdevice_attribute attrib) {-  int val;-  CUDA_SUCCEED_FATAL(cuDeviceGetAttribute(&val, attrib, dev));-  return val;-}--#define CU_FUN_ATTR(x) (CU_FUNC_ATTRIBUTE_##x)-#define function_query(fn,attrib) _function_query(dev, CU_FUN_ATTR(attrib))-static int _function_query(CUfunction dev, CUfunction_attribute attrib) {-  int val;-  CUDA_SUCCEED_FATAL(cuFuncGetAttribute(&val, attrib, dev));-  return val;-}--static void set_preferred_device(struct cuda_config *cfg, const char *s) {-  int x = 0;-  if (*s == '#') {-    s++;-    while (isdigit(*s)) {-      x = x * 10 + (*s++)-'0';-    }-    // Skip trailing spaces.-    while (isspace(*s)) {-      s++;-    }-  }-  cfg->preferred_device = s;-  cfg->preferred_device_num = x;-}--static int cuda_device_setup(struct cuda_context *ctx) {-  char name[256];-  int count, chosen = -1, best_cc = -1;-  int cc_major_best, cc_minor_best;-  int cc_major, cc_minor;-  CUdevice dev;--  CUDA_SUCCEED_FATAL(cuDeviceGetCount(&count));-  if (count == 0) { return 1; }--  int num_device_matches = 0;--  // XXX: Current device selection policy is to choose the device with the-  // highest compute capability (if no preferred device is set).-  // This should maybe be changed, since greater compute capability is not-  // necessarily an indicator of better performance.-  for (int i = 0; i < count; i++) {-    CUDA_SUCCEED_FATAL(cuDeviceGet(&dev, i));--    cc_major = device_query(dev, COMPUTE_CAPABILITY_MAJOR);-    cc_minor = device_query(dev, COMPUTE_CAPABILITY_MINOR);--    CUDA_SUCCEED_FATAL(cuDeviceGetName(name, sizeof(name) - 1, dev));-    name[sizeof(name) - 1] = 0;--    if (ctx->cfg.logging) {-      fprintf(stderr, "Device #%d: name=\"%s\", compute capability=%d.%d\n",-          i, name, cc_major, cc_minor);-    }--    if (device_query(dev, COMPUTE_MODE) == CU_COMPUTEMODE_PROHIBITED) {-      if (ctx->cfg.logging) {-        fprintf(stderr, "Device #%d is compute-prohibited, ignoring\n", i);-      }-      continue;-    }--    if (best_cc == -1 || cc_major > cc_major_best ||-        (cc_major == cc_major_best && cc_minor > cc_minor_best)) {-      best_cc = i;-      cc_major_best = cc_major;-      cc_minor_best = cc_minor;-    }--    if (strstr(name, ctx->cfg.preferred_device) != NULL &&-        num_device_matches++ == ctx->cfg.preferred_device_num) {-      chosen = i;-      break;-    }-  }--  if (chosen == -1) { chosen = best_cc; }-  if (chosen == -1) { return 1; }--  if (ctx->cfg.logging) {-    fprintf(stderr, "Using device #%d\n", chosen);-  }--  CUDA_SUCCEED_FATAL(cuDeviceGet(&ctx->dev, chosen));-  return 0;-}--static char *concat_fragments(const char *src_fragments[]) {-  size_t src_len = 0;-  const char **p;--  for (p = src_fragments; *p; p++) {-    src_len += strlen(*p);-  }--  char *src = (char*) malloc(src_len + 1);-  size_t n = 0;-  for (p = src_fragments; *p; p++) {-    strcpy(src + n, *p);-    n += strlen(*p);-  }--  return src;-}--static const char *cuda_nvrtc_get_arch(CUdevice dev) {-  struct {-    int major;-    int minor;-    const char *arch_str;-  } static const x[] = {-    { 3, 0, "compute_30" },-    { 3, 2, "compute_32" },-    { 3, 5, "compute_35" },-    { 3, 7, "compute_37" },-    { 5, 0, "compute_50" },-    { 5, 2, "compute_52" },-    { 5, 3, "compute_53" },-    { 6, 0, "compute_60" },-    { 6, 1, "compute_61" },-    { 6, 2, "compute_62" },-    { 7, 0, "compute_70" },-    { 7, 2, "compute_72" },-    { 7, 5, "compute_75" },-    { 8, 0, "compute_80" },-    { 8, 6, "compute_80" },-    { 8, 7, "compute_80" }-  };--  int major = device_query(dev, COMPUTE_CAPABILITY_MAJOR);-  int minor = device_query(dev, COMPUTE_CAPABILITY_MINOR);--  int chosen = -1;-  for (int i = 0; i < sizeof(x)/sizeof(x[0]); i++) {-    if (x[i].major < major || (x[i].major == major && x[i].minor <= minor)) {-      chosen = i;-    } else {-      break;-    }-  }--  if (chosen == -1) {-    futhark_panic(-1, "Unsupported compute capability %d.%d\n", major, minor);-  }--  if (x[chosen].major != major || x[chosen].minor != minor) {-    fprintf(stderr,-            "Warning: device compute capability is %d.%d, but newest supported by Futhark is %d.%d.\n",-            major, minor, x[chosen].major, x[chosen].minor);-  }--  return x[chosen].arch_str;-}--static void cuda_nvrtc_mk_build_options(struct cuda_context *ctx, const char *extra_opts[],-                                        char*** opts_out, size_t *n_opts) {-  int arch_set = 0, num_extra_opts;--  // 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 i = 0, n_opts_alloc = 20 + num_extra_opts + ctx->cfg.num_sizes;-  char **opts = (char**) malloc(n_opts_alloc * sizeof(char *));-  if (!arch_set) {-    opts[i++] = strdup("-arch");-    opts[i++] = strdup(cuda_nvrtc_get_arch(ctx->dev));-  }-  opts[i++] = strdup("-default-device");-  if (ctx->cfg.debugging) {-    opts[i++] = strdup("-G");-    opts[i++] = strdup("-lineinfo");-  } else {-    opts[i++] = strdup("--disable-warnings");-  }-  opts[i++] = msgprintf("-D%s=%d",-                        "max_group_size",-                        (int)ctx->max_block_size);-  for (size_t j = 0; j < ctx->cfg.num_sizes; j++) {-    opts[i++] = msgprintf("-D%s=%zu", ctx->cfg.size_vars[j],-                          ctx->cfg.size_values[j]);-  }-  opts[i++] = msgprintf("-DLOCKSTEP_WIDTH=%zu", ctx->lockstep_width);-  opts[i++] = msgprintf("-DMAX_THREADS_PER_BLOCK=%zu", ctx->max_block_size);--  // Time for the best lines of the code in the entire compiler.-  if (getenv("CUDA_HOME") != NULL) {-    opts[i++] = msgprintf("-I%s/include", getenv("CUDA_HOME"));-  }-  if (getenv("CUDA_ROOT") != NULL) {-    opts[i++] = msgprintf("-I%s/include", getenv("CUDA_ROOT"));-  }-  if (getenv("CUDA_PATH") != NULL) {-    opts[i++] = msgprintf("-I%s/include", getenv("CUDA_PATH"));-  }-  opts[i++] = msgprintf("-I/usr/local/cuda/include");-  opts[i++] = msgprintf("-I/usr/include");--  for (int j = 0; extra_opts[j] != NULL; j++) {-    opts[i++] = strdup(extra_opts[j]);-  }--  *n_opts = i;-  *opts_out = opts;-}--static char* cuda_nvrtc_build(struct cuda_context *ctx, const char *src,-                              const char *opts[], size_t n_opts,-                              char **ptx) {-  nvrtcProgram prog;-  char *problem = NULL;--  problem = NVRTC_SUCCEED_NONFATAL(nvrtcCreateProgram(&prog, src, "futhark-cuda", 0, NULL, NULL));--  if (problem) {-    return problem;-  }--  nvrtcResult res = nvrtcCompileProgram(prog, n_opts, opts);-  if (res != NVRTC_SUCCESS) {-    size_t log_size;-    if (nvrtcGetProgramLogSize(prog, &log_size) == NVRTC_SUCCESS) {-      char *log = (char*) malloc(log_size);-      if (nvrtcGetProgramLog(prog, log) == NVRTC_SUCCESS) {-        problem = msgprintf("NVRTC compilation failed.\n\n%s\n", log);-      } else {-        problem = msgprintf("Could not retrieve compilation log\n");-      }-      free(log);-    }-    return problem;-  }--  size_t ptx_size;-  NVRTC_SUCCEED_FATAL(nvrtcGetPTXSize(prog, &ptx_size));-  *ptx = (char*) malloc(ptx_size);-  NVRTC_SUCCEED_FATAL(nvrtcGetPTX(prog, *ptx));--  NVRTC_SUCCEED_FATAL(nvrtcDestroyProgram(&prog));--  return NULL;-}--static void cuda_load_ptx_from_cache(struct cuda_context *ctx, const char *src,-                                     const char *opts[], size_t n_opts,-                                     struct cache_hash *h, const char *cache_fname,-                                     char **ptx) {-  if (ctx->cfg.logging) {-    fprintf(stderr, "Restoring cache from from %s...\n", cache_fname);-  }-  cache_hash_init(h);-  for (size_t i = 0; i < n_opts; i++) {-    cache_hash(h, opts[i], strlen(opts[i]));-  }-  cache_hash(h, src, strlen(src));-  size_t ptxsize;-  errno = 0;-  if (cache_restore(cache_fname, h, (unsigned char**)ptx, &ptxsize) != 0) {-    if (ctx->cfg.logging) {-      fprintf(stderr, "Failed to restore cache (errno: %s)\n", strerror(errno));-    }-  }-}--static void cuda_size_setup(struct cuda_context *ctx)-{-  if (ctx->cfg.default_block_size > ctx->max_block_size) {-    if (ctx->cfg.default_block_size_changed) {-      fprintf(stderr,-          "Note: Device limits default block size to %zu (down from %zu).\n",-          ctx->max_block_size, ctx->cfg.default_block_size);-    }-    ctx->cfg.default_block_size = ctx->max_block_size;-  }-  if (ctx->cfg.default_grid_size > ctx->max_grid_size) {-    if (ctx->cfg.default_grid_size_changed) {-      fprintf(stderr,-          "Note: Device limits default grid size to %zu (down from %zu).\n",-          ctx->max_grid_size, ctx->cfg.default_grid_size);-    }-    ctx->cfg.default_grid_size = ctx->max_grid_size;-  }-  if (ctx->cfg.default_tile_size > ctx->max_tile_size) {-    if (ctx->cfg.default_tile_size_changed) {-      fprintf(stderr,-          "Note: Device limits default tile size to %zu (down from %zu).\n",-          ctx->max_tile_size, ctx->cfg.default_tile_size);-    }-    ctx->cfg.default_tile_size = ctx->max_tile_size;-  }--  if (!ctx->cfg.default_grid_size_changed) {-    ctx->cfg.default_grid_size =-      (device_query(ctx->dev, MULTIPROCESSOR_COUNT) *-       device_query(ctx->dev, MAX_THREADS_PER_MULTIPROCESSOR))-      / ctx->cfg.default_block_size;-  }--  for (int i = 0; i < ctx->cfg.num_sizes; i++) {-    const char *size_class = ctx->cfg.size_classes[i];-    int64_t *size_value = &ctx->cfg.size_values[i];-    const char* size_name = ctx->cfg.size_names[i];-    int64_t max_value = 0, default_value = 0;--    if (strstr(size_class, "group_size") == size_class) {-      max_value = ctx->max_block_size;-      default_value = ctx->cfg.default_block_size;-    } else if (strstr(size_class, "num_groups") == size_class) {-      max_value = ctx->max_grid_size;-      default_value = ctx->cfg.default_grid_size;-      // XXX: as a quick and dirty hack, use twice as many threads for-      // histograms by default.  We really should just be smarter-      // about sizes somehow.-      if (strstr(size_name, ".seghist_") != NULL) {-        default_value *= 2;-      }-    } else if (strstr(size_class, "tile_size") == size_class) {-      max_value = ctx->max_tile_size;-      default_value = ctx->cfg.default_tile_size;-    } else if (strstr(size_class, "reg_tile_size") == size_class) {-      max_value = 0; // No limit.-      default_value = ctx->cfg.default_reg_tile_size;-    } else if (strstr(size_class, "threshold") == size_class) {-      // Threshold can be as large as it takes.-      default_value = ctx->cfg.default_threshold;-    } else {-      // Bespoke sizes have no limit or default.-    }--    if (*size_value == 0) {-      *size_value = default_value;-    } else if (max_value > 0 && *size_value > max_value) {-      fprintf(stderr, "Note: Device limits %s to %zu (down from %zu)\n",-              size_name, max_value, *size_value);-      *size_value = max_value;-    }-  }-}--static char* cuda_module_setup(struct cuda_context *ctx,-                               const char *src_fragments[],-                               const char *extra_opts[],-                               const char* cache_fname) {-  char *ptx = NULL, *src = NULL;--  if (ctx->cfg.load_program_from == NULL) {-    src = concat_fragments(src_fragments);-  } else {-    src = slurp_file(ctx->cfg.load_program_from, NULL);-  }--  if (ctx->cfg.load_ptx_from) {-    if (ctx->cfg.load_program_from != NULL) {-      fprintf(stderr,-              "WARNING: Using PTX from %s instead of C code from %s\n",-              ctx->cfg.load_ptx_from, ctx->cfg.load_program_from);-    }-    ptx = slurp_file(ctx->cfg.load_ptx_from, NULL);-  }--  if (ctx->cfg.dump_program_to != NULL) {-    dump_file(ctx->cfg.dump_program_to, src, strlen(src));-  }--  char **opts;-  size_t n_opts;-  cuda_nvrtc_mk_build_options(ctx, extra_opts, &opts, &n_opts);--  if (ctx->cfg.logging) {-    fprintf(stderr, "NVRTC compile options:\n");-    for (size_t j = 0; j < n_opts; j++) {-      fprintf(stderr, "\t%s\n", opts[j]);-    }-    fprintf(stderr, "\n");-  }--  struct cache_hash h;-  int loaded_ptx_from_cache = 0;-  if (cache_fname != NULL) {-    cuda_load_ptx_from_cache(ctx, src, (const char**)opts, n_opts, &h, cache_fname, &ptx);--    if (ptx != NULL) {-      if (ctx->cfg.logging) {-        fprintf(stderr, "Restored PTX from cache; now loading module...\n");-      }-      if (cuModuleLoadData(&ctx->module, ptx) == CUDA_SUCCESS) {-        if (ctx->cfg.logging) {-          fprintf(stderr, "Success!\n");-        }-        loaded_ptx_from_cache = 1;-      } else {-        if (ctx->cfg.logging) {-          fprintf(stderr, "Failed!\n");-        }-        free(ptx);-        ptx = NULL;-      }-    }-  }--  if (ptx == NULL) {-    char* problem = cuda_nvrtc_build(ctx, src, (const char**)opts, n_opts, &ptx);-    if (problem != NULL) {-      free(src);-      return problem;-    }-  }--  if (ctx->cfg.dump_ptx_to != NULL) {-    dump_file(ctx->cfg.dump_ptx_to, ptx, strlen(ptx));-  }--  if (!loaded_ptx_from_cache) {-    CUDA_SUCCEED_FATAL(cuModuleLoadData(&ctx->module, ptx));-  }--  if (cache_fname != NULL && !loaded_ptx_from_cache) {-    if (ctx->cfg.logging) {-      fprintf(stderr, "Caching PTX in %s...\n", cache_fname);-    }-    errno = 0;-    if (cache_store(cache_fname, &h, (const unsigned char*)ptx, strlen(ptx)) != 0) {-      fprintf(stderr, "Failed to cache PTX: %s\n", strerror(errno));-    }-  }--  for (size_t i = 0; i < n_opts; i++) {-    free((char *)opts[i]);-  }-  free(opts);--  free(ptx);-  if (src != NULL) {-    free(src);-  }--  return NULL;-}--static char* cuda_setup(struct cuda_context *ctx, const char *src_fragments[],-                        const char *extra_opts[], const char* cache_fname) {-  CUDA_SUCCEED_FATAL(cuInit(0));--  if (cuda_device_setup(ctx) != 0) {-    futhark_panic(-1, "No suitable CUDA device found.\n");-  }-  CUDA_SUCCEED_FATAL(cuCtxCreate(&ctx->cu_ctx, 0, ctx->dev));--  free_list_init(&ctx->free_list);--  ctx->max_shared_memory = device_query(ctx->dev, MAX_SHARED_MEMORY_PER_BLOCK);-  ctx->max_block_size = device_query(ctx->dev, MAX_THREADS_PER_BLOCK);-  ctx->max_grid_size = device_query(ctx->dev, MAX_GRID_DIM_X);-  ctx->max_tile_size = sqrt(ctx->max_block_size);-  ctx->max_threshold = 0;-  ctx->max_bespoke = 0;-  ctx->lockstep_width = device_query(ctx->dev, WARP_SIZE);--  cuda_size_setup(ctx);-  return cuda_module_setup(ctx, src_fragments, extra_opts, cache_fname);-}--// Count up the runtime all the profiling_records that occured during execution.-// Also clears the buffer of profiling_records.-static cudaError_t cuda_tally_profiling_records(struct cuda_context *ctx) {-  cudaError_t err;-  for (int i = 0; i < ctx->profiling_records_used; i++) {-    struct profiling_record record = ctx->profiling_records[i];--    float ms;-    if ((err = cudaEventElapsedTime(&ms, record.events[0], record.events[1])) != cudaSuccess) {-      return err;-    }--    // CUDA provides milisecond resolution, but we want microseconds.-    *record.runs += 1;-    *record.runtime += ms*1000;--    if ((err = cudaEventDestroy(record.events[0])) != cudaSuccess) {-      return err;-    }-    if ((err = cudaEventDestroy(record.events[1])) != cudaSuccess) {-      return err;-    }--    free(record.events);-  }--  ctx->profiling_records_used = 0;--  return cudaSuccess;-}--// Returns pointer to two events.-static cudaEvent_t* cuda_get_events(struct cuda_context *ctx, int *runs, int64_t *runtime) {-    if (ctx->profiling_records_used == ctx->profiling_records_capacity) {-      ctx->profiling_records_capacity *= 2;-      ctx->profiling_records =-        realloc(ctx->profiling_records,-                ctx->profiling_records_capacity *-                sizeof(struct profiling_record));-    }-    cudaEvent_t *events = calloc(2, sizeof(cudaEvent_t));-    cudaEventCreate(&events[0]);-    cudaEventCreate(&events[1]);-    ctx->profiling_records[ctx->profiling_records_used].events = events;-    ctx->profiling_records[ctx->profiling_records_used].runs = runs;-    ctx->profiling_records[ctx->profiling_records_used].runtime = runtime;-    ctx->profiling_records_used++;-    return events;-}--static CUresult cuda_free_all(struct cuda_context *ctx);--static void cuda_cleanup(struct cuda_context *ctx) {-  CUDA_SUCCEED_FATAL(cuda_free_all(ctx));-  (void)cuda_tally_profiling_records(ctx);-  free(ctx->profiling_records);-  CUDA_SUCCEED_FATAL(cuModuleUnload(ctx->module));-  CUDA_SUCCEED_FATAL(cuCtxDestroy(ctx->cu_ctx));-}--static CUresult cuda_alloc(struct cuda_context *ctx, FILE *log,-                           size_t min_size, const char *tag,-                           CUdeviceptr *mem_out, size_t *size_out) {-  if (min_size < sizeof(int)) {-    min_size = sizeof(int);-  }--  if (free_list_find(&ctx->free_list, min_size, tag, size_out, (fl_mem*)mem_out) == 0) {-    if (*size_out >= min_size) {-      if (ctx->cfg.debugging) {-        fprintf(log, "No need to allocate: Found a block in the free list.\n");-      }-      return CUDA_SUCCESS;-    } else {-      if (ctx->cfg.debugging) {-        fprintf(log, "Found a free block, but it was too small.\n");-      }--      CUresult res = cuMemFree(*mem_out);-      if (res != CUDA_SUCCESS) {-        return res;-      }-    }-  }--  *size_out = min_size;--  if (ctx->cfg.debugging) {-    fprintf(log, "Actually allocating the desired block.\n");-  }--  CUresult res = cuMemAlloc(mem_out, min_size);-  while (res == CUDA_ERROR_OUT_OF_MEMORY) {-    CUdeviceptr mem;-    if (free_list_first(&ctx->free_list, (fl_mem*)&mem) == 0) {-      res = cuMemFree(mem);-      if (res != CUDA_SUCCESS) {-        return res;-      }-    } else {-      break;-    }-    res = cuMemAlloc(mem_out, min_size);-  }--  return res;-}--static CUresult cuda_free(struct cuda_context *ctx,-                          CUdeviceptr mem, size_t size, const char *tag) {-  free_list_insert(&ctx->free_list, size, (fl_mem)mem, tag);-  return CUDA_SUCCESS;-}--static CUresult cuda_free_all(struct cuda_context *ctx) {-  CUdeviceptr mem;-  free_list_pack(&ctx->free_list);-  while (free_list_first(&ctx->free_list, (fl_mem*)&mem) == 0) {-    CUresult res = cuMemFree(mem);-    if (res != CUDA_SUCCESS) {-      return res;-    }-  }--  return CUDA_SUCCESS;-}--// End of cuda.h.

@@ -1,1066 +0,0 @@-// Start of opencl.h.--#define OPENCL_SUCCEED_FATAL(e) opencl_succeed_fatal(e, #e, __FILE__, __LINE__)-#define OPENCL_SUCCEED_NONFATAL(e) opencl_succeed_nonfatal(e, #e, __FILE__, __LINE__)-// Take care not to override an existing error.-#define OPENCL_SUCCEED_OR_RETURN(e) {             \-    char *serror = OPENCL_SUCCEED_NONFATAL(e);    \-    if (serror) {                                 \-      if (!ctx->error) {                          \-        ctx->error = serror;                      \-        return bad;                               \-      } else {                                    \-        free(serror);                             \-      }                                           \-    }                                             \-  }--// OPENCL_SUCCEED_OR_RETURN returns the value of the variable 'bad' in-// scope.  By default, it will be this one.  Create a local variable-// of some other type if needed.  This is a bit of a hack, but it-// saves effort in the code generator.-static const int bad = 1;--struct opencl_config {-  int debugging;-  int profiling;-  int logging;-  int preferred_device_num;-  const char *preferred_platform;-  const char *preferred_device;-  int ignore_blacklist;--  const char* dump_program_to;-  const char* load_program_from;-  const char* dump_binary_to;-  const char* load_binary_from;--  size_t default_group_size;-  size_t default_num_groups;-  size_t default_tile_size;-  size_t default_reg_tile_size;-  size_t default_threshold;--  int default_group_size_changed;-  int default_tile_size_changed;--  int num_sizes;-  const char **size_names;-  const char **size_vars;-  int64_t *size_values;-  const char **size_classes;-};--static void opencl_config_init(struct opencl_config *cfg,-                               int num_sizes,-                               const char *size_names[],-                               const char *size_vars[],-                               int64_t *size_values,-                               const char *size_classes[]) {-  cfg->debugging = 0;-  cfg->logging = 0;-  cfg->profiling = 0;-  cfg->preferred_device_num = 0;-  cfg->preferred_platform = "";-  cfg->preferred_device = "";-  cfg->ignore_blacklist = 0;-  cfg->dump_program_to = NULL;-  cfg->load_program_from = NULL;-  cfg->dump_binary_to = NULL;-  cfg->load_binary_from = NULL;--  // The following are dummy sizes that mean the concrete defaults-  // will be set during initialisation via hardware-inspection-based-  // heuristics.-  cfg->default_group_size = 0;-  cfg->default_num_groups = 0;-  cfg->default_tile_size = 0;-  cfg->default_reg_tile_size = 0;-  cfg->default_threshold = 0;--  cfg->default_group_size_changed = 0;-  cfg->default_tile_size_changed = 0;--  cfg->num_sizes = num_sizes;-  cfg->size_names = size_names;-  cfg->size_vars = size_vars;-  cfg->size_values = size_values;-  cfg->size_classes = size_classes;-}--// A record of something that happened.-struct profiling_record {-  cl_event *event;-  int *runs;-  int64_t *runtime;-};--struct opencl_context {-  cl_device_id device;-  cl_context ctx;-  cl_command_queue queue;--  struct opencl_config cfg;--  struct free_list free_list;--  size_t max_group_size;-  size_t max_num_groups;-  size_t max_tile_size;-  size_t max_threshold;-  size_t max_local_memory;--  size_t lockstep_width;--  struct profiling_record *profiling_records;-  int profiling_records_capacity;-  int profiling_records_used;-};--struct opencl_device_option {-  cl_platform_id platform;-  cl_device_id device;-  cl_device_type device_type;-  char *platform_name;-  char *device_name;-};--// This function must be defined by the user.  It is invoked by-// setup_opencl() after the platform and device has been found, but-// before the program is loaded.  Its intended use is to tune-// constants based on the selected platform and device.-static void post_opencl_setup(struct opencl_context*, struct opencl_device_option*);--static char *strclone(const char *str) {-  size_t size = strlen(str) + 1;-  char *copy = (char*) malloc(size);-  if (copy == NULL) {-    return NULL;-  }--  memcpy(copy, str, size);-  return copy;-}--static const char* opencl_error_string(cl_int err)-{-    switch (err) {-        case CL_SUCCESS:                            return "Success!";-        case CL_DEVICE_NOT_FOUND:                   return "Device not found.";-        case CL_DEVICE_NOT_AVAILABLE:               return "Device not available";-        case CL_COMPILER_NOT_AVAILABLE:             return "Compiler not available";-        case CL_MEM_OBJECT_ALLOCATION_FAILURE:      return "Memory object allocation failure";-        case CL_OUT_OF_RESOURCES:                   return "Out of resources";-        case CL_OUT_OF_HOST_MEMORY:                 return "Out of host memory";-        case CL_PROFILING_INFO_NOT_AVAILABLE:       return "Profiling information not available";-        case CL_MEM_COPY_OVERLAP:                   return "Memory copy overlap";-        case CL_IMAGE_FORMAT_MISMATCH:              return "Image format mismatch";-        case CL_IMAGE_FORMAT_NOT_SUPPORTED:         return "Image format not supported";-        case CL_BUILD_PROGRAM_FAILURE:              return "Program build failure";-        case CL_MAP_FAILURE:                        return "Map failure";-        case CL_INVALID_VALUE:                      return "Invalid value";-        case CL_INVALID_DEVICE_TYPE:                return "Invalid device type";-        case CL_INVALID_PLATFORM:                   return "Invalid platform";-        case CL_INVALID_DEVICE:                     return "Invalid device";-        case CL_INVALID_CONTEXT:                    return "Invalid context";-        case CL_INVALID_QUEUE_PROPERTIES:           return "Invalid queue properties";-        case CL_INVALID_COMMAND_QUEUE:              return "Invalid command queue";-        case CL_INVALID_HOST_PTR:                   return "Invalid host pointer";-        case CL_INVALID_MEM_OBJECT:                 return "Invalid memory object";-        case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR:    return "Invalid image format descriptor";-        case CL_INVALID_IMAGE_SIZE:                 return "Invalid image size";-        case CL_INVALID_SAMPLER:                    return "Invalid sampler";-        case CL_INVALID_BINARY:                     return "Invalid binary";-        case CL_INVALID_BUILD_OPTIONS:              return "Invalid build options";-        case CL_INVALID_PROGRAM:                    return "Invalid program";-        case CL_INVALID_PROGRAM_EXECUTABLE:         return "Invalid program executable";-        case CL_INVALID_KERNEL_NAME:                return "Invalid kernel name";-        case CL_INVALID_KERNEL_DEFINITION:          return "Invalid kernel definition";-        case CL_INVALID_KERNEL:                     return "Invalid kernel";-        case CL_INVALID_ARG_INDEX:                  return "Invalid argument index";-        case CL_INVALID_ARG_VALUE:                  return "Invalid argument value";-        case CL_INVALID_ARG_SIZE:                   return "Invalid argument size";-        case CL_INVALID_KERNEL_ARGS:                return "Invalid kernel arguments";-        case CL_INVALID_WORK_DIMENSION:             return "Invalid work dimension";-        case CL_INVALID_WORK_GROUP_SIZE:            return "Invalid work group size";-        case CL_INVALID_WORK_ITEM_SIZE:             return "Invalid work item size";-        case CL_INVALID_GLOBAL_OFFSET:              return "Invalid global offset";-        case CL_INVALID_EVENT_WAIT_LIST:            return "Invalid event wait list";-        case CL_INVALID_EVENT:                      return "Invalid event";-        case CL_INVALID_OPERATION:                  return "Invalid operation";-        case CL_INVALID_GL_OBJECT:                  return "Invalid OpenGL object";-        case CL_INVALID_BUFFER_SIZE:                return "Invalid buffer size";-        case CL_INVALID_MIP_LEVEL:                  return "Invalid mip-map level";-        default:                                    return "Unknown";-    }-}--static void opencl_succeed_fatal(cl_int ret,-                                 const char *call,-                                 const char *file,-                                 int line) {-  if (ret != CL_SUCCESS) {-    futhark_panic(-1, "%s:%d: OpenCL call\n  %s\nfailed with error code %d (%s)\n",-          file, line, call, ret, opencl_error_string(ret));-  }-}--static char* opencl_succeed_nonfatal(cl_int ret,-                                     const char *call,-                                     const char *file,-                                     int line) {-  if (ret != CL_SUCCESS) {-    return msgprintf("%s:%d: OpenCL call\n  %s\nfailed with error code %d (%s)\n",-                     file, line, call, ret, opencl_error_string(ret));-  } else {-    return NULL;-  }-}--static void set_preferred_platform(struct opencl_config *cfg, const char *s) {-  cfg->preferred_platform = s;-  cfg->ignore_blacklist = 1;-}--static void set_preferred_device(struct opencl_config *cfg, const char *s) {-  int x = 0;-  if (*s == '#') {-    s++;-    while (isdigit(*s)) {-      x = x * 10 + (*s++)-'0';-    }-    // Skip trailing spaces.-    while (isspace(*s)) {-      s++;-    }-  }-  cfg->preferred_device = s;-  cfg->preferred_device_num = x;-  cfg->ignore_blacklist = 1;-}--static char* opencl_platform_info(cl_platform_id platform,-                                  cl_platform_info param) {-  size_t req_bytes;-  char *info;--  OPENCL_SUCCEED_FATAL(clGetPlatformInfo(platform, param, 0, NULL, &req_bytes));--  info = (char*) malloc(req_bytes);--  OPENCL_SUCCEED_FATAL(clGetPlatformInfo(platform, param, req_bytes, info, NULL));--  return info;-}--static char* opencl_device_info(cl_device_id device,-                                cl_device_info param) {-  size_t req_bytes;-  char *info;--  OPENCL_SUCCEED_FATAL(clGetDeviceInfo(device, param, 0, NULL, &req_bytes));--  info = (char*) malloc(req_bytes);--  OPENCL_SUCCEED_FATAL(clGetDeviceInfo(device, param, req_bytes, info, NULL));--  return info;-}--static void opencl_all_device_options(struct opencl_device_option **devices_out,-                                      size_t *num_devices_out) {-  size_t num_devices = 0, num_devices_added = 0;--  cl_platform_id *all_platforms;-  cl_uint *platform_num_devices;--  cl_uint num_platforms;--  // Find the number of platforms.-  OPENCL_SUCCEED_FATAL(clGetPlatformIDs(0, NULL, &num_platforms));--  // Make room for them.-  all_platforms = calloc(num_platforms, sizeof(cl_platform_id));-  platform_num_devices = calloc(num_platforms, sizeof(cl_uint));--  // Fetch all the platforms.-  OPENCL_SUCCEED_FATAL(clGetPlatformIDs(num_platforms, all_platforms, NULL));--  // Count the number of devices for each platform, as well as the-  // total number of devices.-  for (cl_uint i = 0; i < num_platforms; i++) {-    if (clGetDeviceIDs(all_platforms[i], CL_DEVICE_TYPE_ALL,-                       0, NULL, &platform_num_devices[i]) == CL_SUCCESS) {-      num_devices += platform_num_devices[i];-    } else {-      platform_num_devices[i] = 0;-    }-  }--  // Make room for all the device options.-  struct opencl_device_option *devices =-    calloc(num_devices, sizeof(struct opencl_device_option));--  // Loop through the platforms, getting information about their devices.-  for (cl_uint i = 0; i < num_platforms; i++) {-    cl_platform_id platform = all_platforms[i];-    cl_uint num_platform_devices = platform_num_devices[i];--    if (num_platform_devices == 0) {-      continue;-    }--    char *platform_name = opencl_platform_info(platform, CL_PLATFORM_NAME);-    cl_device_id *platform_devices =-      calloc(num_platform_devices, sizeof(cl_device_id));--    // Fetch all the devices.-    OPENCL_SUCCEED_FATAL(clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL,-                                  num_platform_devices, platform_devices, NULL));--    // Loop through the devices, adding them to the devices array.-    for (cl_uint i = 0; i < num_platform_devices; i++) {-      char *device_name = opencl_device_info(platform_devices[i], CL_DEVICE_NAME);-      devices[num_devices_added].platform = platform;-      devices[num_devices_added].device = platform_devices[i];-      OPENCL_SUCCEED_FATAL(clGetDeviceInfo(platform_devices[i], CL_DEVICE_TYPE,-                                     sizeof(cl_device_type),-                                     &devices[num_devices_added].device_type,-                                     NULL));-      // We don't want the structs to share memory, so copy the platform name.-      // Each device name is already unique.-      devices[num_devices_added].platform_name = strclone(platform_name);-      devices[num_devices_added].device_name = device_name;-      num_devices_added++;-    }-    free(platform_devices);-    free(platform_name);-  }-  free(all_platforms);-  free(platform_num_devices);--  *devices_out = devices;-  *num_devices_out = num_devices;-}--// Returns 0 on success.-static int list_devices(void) {-  struct opencl_device_option *devices;-  size_t num_devices;--  opencl_all_device_options(&devices, &num_devices);--  const char *cur_platform = "";-  for (size_t i = 0; i < num_devices; i++) {-    struct opencl_device_option device = devices[i];-    if (strcmp(cur_platform, device.platform_name) != 0) {-      printf("Platform: %s\n", device.platform_name);-      cur_platform = device.platform_name;-    }-    printf("[%d]: %s\n", (int)i, device.device_name);-  }--  // Free all the platform and device names.-  for (size_t j = 0; j < num_devices; j++) {-    free(devices[j].platform_name);-    free(devices[j].device_name);-  }-  free(devices);--  return 0;-}--// Returns 0 on success.-static int select_device_interactively(struct opencl_config *cfg) {-  struct opencl_device_option *devices;-  size_t num_devices;-  int ret = 1;--  opencl_all_device_options(&devices, &num_devices);--  printf("Choose OpenCL device:\n");-  const char *cur_platform = "";-  for (size_t i = 0; i < num_devices; i++) {-    struct opencl_device_option device = devices[i];-    if (strcmp(cur_platform, device.platform_name) != 0) {-      printf("Platform: %s\n", device.platform_name);-      cur_platform = device.platform_name;-    }-    printf("[%d] %s\n", (int)i, device.device_name);-  }--  int selection;-  printf("Choice: ");-  if (scanf("%d", &selection) == 1) {-    ret = 0;-    cfg->preferred_platform = "";-    cfg->preferred_device = "";-    cfg->preferred_device_num = selection;-    cfg->ignore_blacklist = 1;-  }--  // Free all the platform and device names.-  for (size_t j = 0; j < num_devices; j++) {-    free(devices[j].platform_name);-    free(devices[j].device_name);-  }-  free(devices);--  return ret;-}--static int is_blacklisted(const char *platform_name, const char *device_name,-                          const struct opencl_config *cfg) {-  if (strcmp(cfg->preferred_platform, "") != 0 ||-      strcmp(cfg->preferred_device, "") != 0) {-    return 0;-  } else if (strstr(platform_name, "Apple") != NULL &&-             strstr(device_name, "Intel(R) Core(TM)") != NULL) {-    return 1;-  } else {-    return 0;-  }-}--static struct opencl_device_option get_preferred_device(const struct opencl_config *cfg) {-  struct opencl_device_option *devices;-  size_t num_devices;--  opencl_all_device_options(&devices, &num_devices);--  int num_device_matches = 0;--  for (size_t i = 0; i < num_devices; i++) {-    struct opencl_device_option device = devices[i];-    if (strstr(device.platform_name, cfg->preferred_platform) != NULL &&-        strstr(device.device_name, cfg->preferred_device) != NULL &&-        (cfg->ignore_blacklist ||-         !is_blacklisted(device.platform_name, device.device_name, cfg)) &&-        num_device_matches++ == cfg->preferred_device_num) {-      // Free all the platform and device names, except the ones we have chosen.-      for (size_t j = 0; j < num_devices; j++) {-        if (j != i) {-          free(devices[j].platform_name);-          free(devices[j].device_name);-        }-      }-      free(devices);-      return device;-    }-  }--  futhark_panic(1, "Could not find acceptable OpenCL device.\n");-  exit(1); // Never reached-}--static void describe_device_option(struct opencl_device_option device) {-  fprintf(stderr, "Using platform: %s\n", device.platform_name);-  fprintf(stderr, "Using device: %s\n", device.device_name);-}--static cl_build_status build_opencl_program(cl_program program, cl_device_id device, const char* options) {-  cl_int clBuildProgram_error = clBuildProgram(program, 1, &device, options, NULL, NULL);--  // Avoid termination due to CL_BUILD_PROGRAM_FAILURE-  if (clBuildProgram_error != CL_SUCCESS &&-      clBuildProgram_error != CL_BUILD_PROGRAM_FAILURE) {-    OPENCL_SUCCEED_FATAL(clBuildProgram_error);-  }--  cl_build_status build_status;-  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;-    OPENCL_SUCCEED_FATAL(clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size));--    build_log = (char*) malloc(ret_val_size+1);-    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';--    fprintf(stderr, "Build log:\n%s\n", build_log);--    free(build_log);-  }--  return build_status;-}--// Fields in a bitmask indicating which types we must be sure are-// available.-enum opencl_required_type { OPENCL_F64 = 1 };--static char* mk_compile_opts(struct opencl_context *ctx,-                             const char *extra_build_opts[],-                             struct opencl_device_option device_option) {-  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 = (char*) malloc(compile_opts_size);--  int w = snprintf(compile_opts, compile_opts_size,-                   "-DLOCKSTEP_WIDTH=%d ",-                   (int)ctx->lockstep_width);--  w += snprintf(compile_opts+w, compile_opts_size-w,-                "-D%s=%d ",-                "max_group_size",-                (int)ctx->max_group_size);--  for (int i = 0; i < ctx->cfg.num_sizes; i++) {-    w += snprintf(compile_opts+w, compile_opts_size-w,-                  "-D%s=%d ",-                  ctx->cfg.size_vars[i],-                  (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]);-  }--  // Oclgrind claims to support cl_khr_fp16, but this is not actually-  // the case.-  if (strcmp(device_option.platform_name, "Oclgrind") == 0) {-    w += snprintf(compile_opts+w, compile_opts_size-w, "-DEMULATE_F16 ");-  }--  return compile_opts;-}--// We take as input several strings representing the program, because-// C does not guarantee that the compiler supports particularly large-// literals.  Notably, Visual C has a limit of 2048 characters.  The-// array must be NULL-terminated.-static cl_program setup_opencl_with_command_queue(struct opencl_context *ctx,-                                                  cl_command_queue queue,-                                                  const char *srcs[],-                                                  int required_types,-                                                  const char *extra_build_opts[],-                                                  const char* cache_fname) {-  int error;--  free_list_init(&ctx->free_list);-  ctx->queue = queue;--  OPENCL_SUCCEED_FATAL(clGetCommandQueueInfo(ctx->queue, CL_QUEUE_CONTEXT, sizeof(cl_context), &ctx->ctx, NULL));--  // Fill out the device info.  This is redundant work if we are-  // called from setup_opencl() (which is the common case), but I-  // doubt it matters much.-  struct opencl_device_option device_option;-  OPENCL_SUCCEED_FATAL(clGetCommandQueueInfo(ctx->queue, CL_QUEUE_DEVICE,-                                       sizeof(cl_device_id),-                                       &device_option.device,-                                       NULL));-  OPENCL_SUCCEED_FATAL(clGetDeviceInfo(device_option.device, CL_DEVICE_PLATFORM,-                                 sizeof(cl_platform_id),-                                 &device_option.platform,-                                 NULL));-  OPENCL_SUCCEED_FATAL(clGetDeviceInfo(device_option.device, CL_DEVICE_TYPE,-                                 sizeof(cl_device_type),-                                 &device_option.device_type,-                                 NULL));-  device_option.platform_name = opencl_platform_info(device_option.platform, CL_PLATFORM_NAME);-  device_option.device_name = opencl_device_info(device_option.device, CL_DEVICE_NAME);--  ctx->device = device_option.device;--  if (required_types & OPENCL_F64) {-    cl_uint supported;-    OPENCL_SUCCEED_FATAL(clGetDeviceInfo(device_option.device, CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE,-                                   sizeof(cl_uint), &supported, NULL));-    if (!supported) {-      futhark_panic(1, "Program uses double-precision floats, but this is not supported on the chosen device: %s\n",-            device_option.device_name);-    }-  }--  size_t max_group_size;-  OPENCL_SUCCEED_FATAL(clGetDeviceInfo(device_option.device, CL_DEVICE_MAX_WORK_GROUP_SIZE,-                                 sizeof(size_t), &max_group_size, NULL));--  size_t max_tile_size = sqrt(max_group_size);--  cl_ulong max_local_memory;-  OPENCL_SUCCEED_FATAL(clGetDeviceInfo(device_option.device, CL_DEVICE_LOCAL_MEM_SIZE,-                                       sizeof(size_t), &max_local_memory, NULL));--  // Futhark reserves 4 bytes for bookkeeping information.-  max_local_memory -= 4;--  // The OpenCL implementation may reserve some local memory bytes for-  // various purposes.  In principle, we should use-  // clGetKernelWorkGroupInfo() to figure out for each kernel how much-  // is actually available, but our current code generator design-  // makes this infeasible.  Instead, we have this nasty hack where we-  // arbitrarily subtract some bytes, based on empirical measurements-  // (but which might be arbitrarily wrong).  Fortunately, we rarely-  // try to really push the local memory usage.-  if (strstr(device_option.platform_name, "NVIDIA CUDA") != NULL) {-    max_local_memory -= 12;-  } else if (strstr(device_option.platform_name, "AMD") != NULL) {-    max_local_memory -= 16;-  }--  // Make sure this function is defined.-  post_opencl_setup(ctx, &device_option);--  if (max_group_size < ctx->cfg.default_group_size) {-    if (ctx->cfg.default_group_size_changed) {-      fprintf(stderr, "Note: Device limits default group size to %zu (down from %zu).\n",-              max_group_size, ctx->cfg.default_group_size);-    }-    ctx->cfg.default_group_size = max_group_size;-  }--  if (max_tile_size < ctx->cfg.default_tile_size) {-    if (ctx->cfg.default_tile_size_changed) {-      fprintf(stderr, "Note: Device limits default tile size to %zu (down from %zu).\n",-              max_tile_size, ctx->cfg.default_tile_size);-    }-    ctx->cfg.default_tile_size = max_tile_size;-  }--  ctx->max_group_size = max_group_size;-  ctx->max_tile_size = max_tile_size; // No limit.-  ctx->max_threshold = ctx->max_num_groups = 0; // No limit.-  ctx->max_local_memory = max_local_memory;--  // Now we go through all the sizes, clamp them to the valid range,-  // or set them to the default.-  for (int i = 0; i < ctx->cfg.num_sizes; i++) {-    const char *size_class = ctx->cfg.size_classes[i];-    int64_t *size_value = &ctx->cfg.size_values[i];-    const char* size_name = ctx->cfg.size_names[i];-    int64_t max_value = 0, default_value = 0;--    if (strstr(size_class, "group_size") == size_class) {-      max_value = max_group_size;-      default_value = ctx->cfg.default_group_size;-    } else if (strstr(size_class, "num_groups") == size_class) {-      max_value = max_group_size; // Futhark assumes this constraint.-      default_value = ctx->cfg.default_num_groups;-      // XXX: as a quick and dirty hack, use twice as many threads for-      // histograms by default.  We really should just be smarter-      // about sizes somehow.-      if (strstr(size_name, ".seghist_") != NULL) {-        default_value *= 2;-      }-    } else if (strstr(size_class, "tile_size") == size_class) {-      max_value = sqrt(max_group_size);-      default_value = ctx->cfg.default_tile_size;-    } else if (strstr(size_class, "reg_tile_size") == size_class) {-      max_value = 0; // No limit.-      default_value = ctx->cfg.default_reg_tile_size;-    } else if (strstr(size_class, "threshold") == size_class) {-      // Threshold can be as large as it takes.-      default_value = ctx->cfg.default_threshold;-    } else {-      // Bespoke sizes have no limit or default.-    }-    if (*size_value == 0) {-      *size_value = default_value;-    } else if (max_value > 0 && *size_value > max_value) {-      fprintf(stderr, "Note: Device limits %s to %d (down from %d)\n",-              size_name, (int)max_value, (int)*size_value);-      *size_value = max_value;-    }-  }--  if (ctx->lockstep_width == 0) {-    ctx->lockstep_width = 1;-  }--  if (ctx->cfg.logging) {-    fprintf(stderr, "Lockstep width: %d\n", (int)ctx->lockstep_width);-    fprintf(stderr, "Default group size: %d\n", (int)ctx->cfg.default_group_size);-    fprintf(stderr, "Default number of groups: %d\n", (int)ctx->cfg.default_num_groups);-  }--  char *compile_opts = mk_compile_opts(ctx, extra_build_opts, device_option);--  if (ctx->cfg.logging) {-    fprintf(stderr, "OpenCL compiler options: %s\n", compile_opts);-  }--  char *fut_opencl_src = NULL;-  cl_program prog;-  error = CL_SUCCESS;--  struct cache_hash h;--  int loaded_from_cache = 0;-  if (ctx->cfg.load_binary_from == NULL) {-    size_t src_size = 0;--    // Maybe we have to read OpenCL source from somewhere else (used for debugging).-    if (ctx->cfg.load_program_from != NULL) {-      fut_opencl_src = slurp_file(ctx->cfg.load_program_from, NULL);-      assert(fut_opencl_src != NULL);-    } else {-      // Construct the OpenCL source concatenating all the fragments.-      for (const char **src = srcs; src && *src; src++) {-        src_size += strlen(*src);-      }--      fut_opencl_src = (char*) malloc(src_size + 1);--      size_t n, i;-      for (i = 0, n = 0; srcs && srcs[i]; i++) {-        strncpy(fut_opencl_src+n, srcs[i], src_size-n);-        n += strlen(srcs[i]);-      }-      fut_opencl_src[src_size] = 0;-    }--    if (ctx->cfg.dump_program_to != NULL) {-      if (ctx->cfg.logging) {-        fprintf(stderr, "Dumping OpenCL source to %s...\n", ctx->cfg.dump_program_to);-      }--      dump_file(ctx->cfg.dump_program_to, fut_opencl_src, strlen(fut_opencl_src));-    }--    if (cache_fname != NULL) {-      if (ctx->cfg.logging) {-        fprintf(stderr, "Restoring cache from from %s...\n", cache_fname);-      }-      cache_hash_init(&h);-      cache_hash(&h, fut_opencl_src, strlen(fut_opencl_src));-      cache_hash(&h, compile_opts, strlen(compile_opts));--      unsigned char *buf;-      size_t bufsize;-      errno = 0;-      if (cache_restore(cache_fname, &h, &buf, &bufsize) != 0) {-        if (ctx->cfg.logging) {-          fprintf(stderr, "Failed to restore cache (errno: %s)\n", strerror(errno));-        }-      } else {-        if (ctx->cfg.logging) {-          fprintf(stderr, "Cache restored; loading OpenCL binary...\n");-        }--        cl_int status = 0;-        prog = clCreateProgramWithBinary(ctx->ctx, 1, &device_option.device,-                                         &bufsize, (const unsigned char**)&buf,-                                         &status, &error);-        if (status == CL_SUCCESS) {-          loaded_from_cache = 1;-          if (ctx->cfg.logging) {-            fprintf(stderr, "Loading succeeded.\n");-          }-        } else {-          if (ctx->cfg.logging) {-            fprintf(stderr, "Loading failed.\n");-          }-        }-      }-    }--    if (!loaded_from_cache) {-      if (ctx->cfg.logging) {-        fprintf(stderr, "Creating OpenCL program...\n");-      }--      const char* src_ptr[] = {fut_opencl_src};-      prog = clCreateProgramWithSource(ctx->ctx, 1, src_ptr, &src_size, &error);-      OPENCL_SUCCEED_FATAL(error);-    }-  } else {-    if (ctx->cfg.logging) {-      fprintf(stderr, "Loading OpenCL binary from %s...\n", ctx->cfg.load_binary_from);-    }-    size_t binary_size;-    unsigned char *fut_opencl_bin =-      (unsigned char*) slurp_file(ctx->cfg.load_binary_from, &binary_size);-    assert(fut_opencl_bin != NULL);-    const unsigned char *binaries[1] = { fut_opencl_bin };-    cl_int status = 0;--    prog = clCreateProgramWithBinary(ctx->ctx, 1, &device_option.device,-                                     &binary_size, binaries,-                                     &status, &error);--    OPENCL_SUCCEED_FATAL(status);-    OPENCL_SUCCEED_FATAL(error);-  }--  if (ctx->cfg.logging) {-    fprintf(stderr, "Building OpenCL program...\n");-  }-  OPENCL_SUCCEED_FATAL(build_opencl_program(prog, device_option.device, compile_opts));--  free(compile_opts);-  free(fut_opencl_src);--  size_t binary_size = 0;-  unsigned char *binary = NULL;-  int store_in_cache = cache_fname != NULL && !loaded_from_cache;-  if (store_in_cache || ctx->cfg.dump_binary_to != NULL) {-    OPENCL_SUCCEED_FATAL(clGetProgramInfo(prog, CL_PROGRAM_BINARY_SIZES,-                                          sizeof(size_t), &binary_size, NULL));-    binary = (unsigned char*) malloc(binary_size);-    OPENCL_SUCCEED_FATAL(clGetProgramInfo(prog, CL_PROGRAM_BINARIES,-                                          sizeof(unsigned char*), &binary, NULL));-  }--  if (store_in_cache) {-    if (ctx->cfg.logging) {-      fprintf(stderr, "Caching OpenCL binary in %s...\n", cache_fname);-    }-    if (cache_store(cache_fname, &h, binary, binary_size) != 0) {-      printf("Failed to cache binary: %s\n", strerror(errno));-    }-  }--  if (ctx->cfg.dump_binary_to != NULL) {-    if (ctx->cfg.logging) {-      fprintf(stderr, "Dumping OpenCL binary to %s...\n", ctx->cfg.dump_binary_to);-    }-    dump_file(ctx->cfg.dump_binary_to, binary, binary_size);-  }--  return prog;-}--static cl_program setup_opencl(struct opencl_context *ctx,-                               const char *srcs[],-                               int required_types,-                               const char *extra_build_opts[],-                               const char* cache_fname) {--  ctx->lockstep_width = 0; // Real value set later.--  struct opencl_device_option device_option = get_preferred_device(&ctx->cfg);--  if (ctx->cfg.logging) {-    describe_device_option(device_option);-  }--  // Note that NVIDIA's OpenCL requires the platform property-  cl_context_properties properties[] = {-    CL_CONTEXT_PLATFORM,-    (cl_context_properties)device_option.platform,-    0-  };--  cl_int clCreateContext_error;-  ctx->ctx = clCreateContext(properties, 1, &device_option.device, NULL, NULL, &clCreateContext_error);-  OPENCL_SUCCEED_FATAL(clCreateContext_error);--  cl_int clCreateCommandQueue_error;-  cl_command_queue queue =-    clCreateCommandQueue(ctx->ctx,-                         device_option.device,-                         ctx->cfg.profiling ? CL_QUEUE_PROFILING_ENABLE : 0,-                         &clCreateCommandQueue_error);-  OPENCL_SUCCEED_FATAL(clCreateCommandQueue_error);--  return setup_opencl_with_command_queue(ctx, queue, srcs, required_types, extra_build_opts,-                                         cache_fname);-}--// Count up the runtime all the profiling_records that occured during execution.-// Also clears the buffer of profiling_records.-static cl_int opencl_tally_profiling_records(struct opencl_context *ctx) {-  cl_int err;-  for (int i = 0; i < ctx->profiling_records_used; i++) {-    struct profiling_record record = ctx->profiling_records[i];--    cl_ulong start_t, end_t;--    if ((err = clGetEventProfilingInfo(*record.event,-                                       CL_PROFILING_COMMAND_START,-                                       sizeof(start_t),-                                       &start_t,-                                       NULL)) != CL_SUCCESS) {-      return err;-    }--    if ((err = clGetEventProfilingInfo(*record.event,-                                       CL_PROFILING_COMMAND_END,-                                       sizeof(end_t),-                                       &end_t,-                                       NULL)) != CL_SUCCESS) {-      return err;-    }--    // OpenCL provides nanosecond resolution, but we want-    // microseconds.-    *record.runs += 1;-    *record.runtime += (end_t - start_t)/1000;--    if ((err = clReleaseEvent(*record.event)) != CL_SUCCESS) {-      return err;-    }-    free(record.event);-  }--  ctx->profiling_records_used = 0;--  return CL_SUCCESS;-}--// If profiling, produce an event associated with a profiling record.-static cl_event* opencl_get_event(struct opencl_context *ctx, int *runs, int64_t *runtime) {-    if (ctx->profiling_records_used == ctx->profiling_records_capacity) {-      ctx->profiling_records_capacity *= 2;-      ctx->profiling_records =-        realloc(ctx->profiling_records,-                ctx->profiling_records_capacity *-                sizeof(struct profiling_record));-    }-    cl_event *event = malloc(sizeof(cl_event));-    ctx->profiling_records[ctx->profiling_records_used].event = event;-    ctx->profiling_records[ctx->profiling_records_used].runs = runs;-    ctx->profiling_records[ctx->profiling_records_used].runtime = runtime;-    ctx->profiling_records_used++;-    return event;-}--// Allocate memory from driver. The problem is that OpenCL may perform-// lazy allocation, so we cannot know whether an allocation succeeded-// until the first time we try to use it.  Hence we immediately-// perform a write to see if the allocation succeeded.  This is slow,-// but the assumption is that this operation will be rare (most things-// will go through the free list).-static int opencl_alloc_actual(struct opencl_context *ctx, size_t size, cl_mem *mem_out) {-  int error;-  *mem_out = clCreateBuffer(ctx->ctx, CL_MEM_READ_WRITE, size, NULL, &error);--  if (error != CL_SUCCESS) {-    return error;-  }--  int x = 2;-  error = clEnqueueWriteBuffer(ctx->queue, *mem_out,-                               CL_TRUE,-                               0, sizeof(x), &x,-                               0, NULL, NULL);--  // No need to wait for completion here. clWaitForEvents() cannot-  // return mem object allocation failures. This implies that the-  // buffer is faulted onto the device on enqueue. (Observation by-  // Andreas Kloeckner.)--  return error;-}--static int opencl_alloc(struct opencl_context *ctx, FILE *log,-                        size_t min_size, const char *tag,-                        cl_mem *mem_out, size_t *size_out) {-  (void)tag;-  if (min_size < sizeof(int)) {-    min_size = sizeof(int);-  }--  cl_mem* memptr;-  if (free_list_find(&ctx->free_list, min_size, tag, size_out, (fl_mem*)&memptr) == 0) {-    // Successfully found a free block.  Is it big enough?-    if (*size_out >= min_size) {-      if (ctx->cfg.debugging) {-        fprintf(log, "No need to allocate: Found a block in the free list.\n");-      }-      *mem_out = *memptr;-      free(memptr);-      return CL_SUCCESS;-    } else {-      if (ctx->cfg.debugging) {-        fprintf(log, "Found a free block, but it was too small.\n");-      }-      int error = clReleaseMemObject(*memptr);-      free(*memptr);-      if (error != CL_SUCCESS) {-        return error;-      }-    }-  }--  *size_out = min_size;--  // We have to allocate a new block from the driver.  If the-  // allocation does not succeed, then we might be in an out-of-memory-  // situation.  We now start freeing things from the free list until-  // we think we have freed enough that the allocation will succeed.-  // Since we don't know how far the allocation is from fitting, we-  // have to check after every deallocation.  This might be pretty-  // expensive.  Let's hope that this case is hit rarely.--  if (ctx->cfg.debugging) {-    fprintf(log, "Actually allocating the desired block.\n");-  }--  int error = opencl_alloc_actual(ctx, min_size, mem_out);--  while (error == CL_MEM_OBJECT_ALLOCATION_FAILURE) {-    if (ctx->cfg.debugging) {-      fprintf(log, "Out of OpenCL memory: releasing entry from the free list...\n");-    }-    cl_mem* memptr;-    if (free_list_first(&ctx->free_list, (fl_mem*)&memptr) == 0) {-      cl_mem mem = *memptr;-      free(memptr);-      error = clReleaseMemObject(mem);-      if (error != CL_SUCCESS) {-        return error;-      }-    } else {-      break;-    }-    error = opencl_alloc_actual(ctx, min_size, mem_out);-  }--  return error;-}--static int opencl_free(struct opencl_context *ctx,-                       cl_mem mem, size_t size, const char *tag) {-  cl_mem* memptr = malloc(sizeof(cl_mem));-  *memptr = mem;-  free_list_insert(&ctx->free_list, size, (fl_mem)memptr, tag);-  return CL_SUCCESS;-}--static int opencl_free_all(struct opencl_context *ctx) {-  free_list_pack(&ctx->free_list);-  cl_mem* memptr;-  while (free_list_first(&ctx->free_list, (fl_mem*)&memptr) == 0) {-    cl_mem mem = *memptr;-    free(memptr);-    int error = clReleaseMemObject(mem);-    if (error != CL_SUCCESS) {-      return error;-    }-  }--  return CL_SUCCESS;-}--// Free everything that belongs to 'ctx', but do not free 'ctx'-// itself.-static void teardown_opencl(struct opencl_context *ctx) {-  (void)opencl_tally_profiling_records(ctx);-  free(ctx->profiling_records);-  (void)opencl_free_all(ctx);-  (void)clReleaseCommandQueue(ctx->queue);-  (void)clReleaseContext(ctx->ctx);-}--// End of opencl.h.

@@ -1774,6 +1774,10 @@   return log10(x); } +static inline float futrts_log1p_32(float x) {+  return log1p(x);+}+ static inline float futrts_sqrt32(float x) {   return sqrt(x); }@@ -1904,6 +1908,13 @@   return futrts_log32(x) / log(10.0f); } +static inline float futrts_log1p_32(float x) {+  if(x == -1.0f || (futrts_isinf32(x) && x > 0.0f)) return x / 0.0f;+  float y = 1.0f + x;+  float z = y - 1.0f;+  return log(y) - (z-x)/y;+}+ static inline float futrts_sqrt32(float x) {   return sqrt(x); }@@ -2106,6 +2117,10 @@   return log10f(x); } +static inline float futrts_log1p_32(float x) {+  return log1pf(x);+}+ static inline float futrts_sqrt32(float x) {   return sqrtf(x); }@@ -2357,6 +2372,13 @@   return futrts_log64(x)/log(10.0d); } +static inline double futrts_log1p_64(double x) {+  if(x == -1.0d || (futrts_isinf64(x) && x > 0.0d)) return x / 0.0d;+  double y = 1.0d + x;+  double z = y - 1.0d;+  return log(y) - (z-x)/y;+}+ static inline double futrts_sqrt64(double x) {   return sqrt(x); }@@ -2722,6 +2744,10 @@  static inline double futrts_log10_64(double x) {   return log10(x);+}++static inline double futrts_log1p_64(double x) {+  return log1p(x); }  static inline double futrts_sqrt64(double x) {

@@ -217,6 +217,10 @@   return log10(x); } +static inline f16 futrts_log1p_16(f16 x) {+  return log1p(x);+}+ static inline f16 futrts_sqrt16(f16 x) {   return sqrt(x); }@@ -346,6 +350,13 @@   return futrts_log16(x) / log(10.0f16); } +static inline f16 futrts_log1p_16(f16 x) {+  if(x == -1.0f16 || (futrts_isinf16(x) && x > 0.0f16)) return x / 0.0f16;+  f16 y = 1.0f16 + x;+  f16 z = y - 1.0f16;+  return log(y) - (z-x)/y;+}+ static inline f16 futrts_sqrt16(f16 x) {   return (float16)sqrt((float)x); }@@ -497,6 +508,10 @@   return hlog10(x); } +static inline f16 futrts_log1p_16(f16 x) {+  return (f16)log1pf((float)x);+}+ static inline f16 futrts_sqrt16(f16 x) {   return hsqrt(x); }@@ -693,6 +708,10 @@  static inline f16 futrts_log10_16(f16 x) {   return futrts_log10_32(x);+}++static inline f16 futrts_log1p_16(f16 x) {+  return futrts_log1p_32(x); }  static inline f16 futrts_sqrt16(f16 x) {

@@ -373,43 +373,49 @@   if (f == NULL) {     failure();     printf("Failed to open %s: %s\n", fname, strerror(errno));-  } else {-    int values = 0;-    for (int i = 1; arg_exists(args, i); i+=2, values++) {-      const char *vname = get_arg(args, i);-      const char *type = get_arg(args, i+1);+    return;+  } -      const struct type *t = get_type(s, type);-      struct variable *v = create_variable(s, vname, t);+  int bad = 0;+  int values = 0;+  for (int i = 1; arg_exists(args, i); i+=2, values++) {+    const char *vname = get_arg(args, i);+    const char *type = get_arg(args, i+1); -      if (v == NULL) {-        failure();-        printf("Variable already exists: %s\n", vname);-        return;-      }+    const struct type *t = get_type(s, type);+    struct variable *v = create_variable(s, vname, t); -      errno = 0;-      if (t->restore(t->aux, f, s->ctx, value_ptr(&v->value)) != 0) {-        failure();-        printf("Failed to restore variable %s.\n"-               "Possibly malformed data in %s (errno: %s)\n",-               vname, fname, strerror(errno));-        drop_variable(v);-        break;-      }+    if (v == NULL) {+      bad = 1;+      failure();+      printf("Variable already exists: %s\n", vname);+      break;     } -    if (end_of_input(f) != 0) {+    errno = 0;+    if (t->restore(t->aux, f, s->ctx, value_ptr(&v->value)) != 0) {+      bad = 1;       failure();-      printf("Expected EOF after reading %d values from %s\n",-             values, fname);+      printf("Failed to restore variable %s.\n"+             "Possibly malformed data in %s (errno: %s)\n",+             vname, fname, strerror(errno));+      drop_variable(v);+      break;     }+  } -    fclose(f);+  if (!bad && end_of_input(f) != 0) {+    failure();+    printf("Expected EOF after reading %d values from %s\n",+           values, fname);   } -  int err = futhark_context_sync(s->ctx);-  error_check(s, err);+  fclose(f);++  if (!bad) {+    int err = futhark_context_sync(s->ctx);+    error_check(s, err);+  } }  void cmd_store(struct server_state *s, const char *args[]) {

@@ -931,6 +931,13 @@   return futrts_log32(x) / log(10.0f); } +static inline uniform float futrts_log1p_32(uniform float x) {+  if(x == -1.0f || (futrts_isinf32(x) && x > 0.0f)) return x / 0.0f;+  uniform float y = 1.0f + x;+  uniform float z = y - 1.0f;+  return log(y) - (z-x)/y;+}+ static inline uniform float futrts_sqrt32(uniform float x) {   return sqrt(x); }@@ -1189,6 +1196,13 @@   return futrts_log64(x)/log(10.0d); } +static inline uniform double futrts_log1p_64(uniform double x) {+  if(x == -1.0d || (futrts_isinf64(x) && x > 0.0d)) return x / 0.0d;+  uniform double y = 1.0d + x;+  uniform double z = y - 1.0d;+  return log(y) - (z-x)/y;+}+ static inline uniform double futrts_sqrt64(uniform double x) {   return sqrt(x); }@@ -1559,6 +1573,13 @@  static inline uniform f16 futrts_log10_16(uniform f16 x) {   return futrts_log16(x) / log(10.0f16);+}++static inline uniform f16 futrts_log1p_16(uniform f16 x) {+  if(x == -1.0f16 || (futrts_isinf16(x) && x > 0.0f16)) return x / 0.0f16;+  uniform f16 y = 1.0f16 + x;+  uniform f16 z = y - 1.0f16;+  return log(y) - (z-x)/y; }  static inline uniform f16 futrts_sqrt16(uniform f16 x) {

@@ -14,6 +14,7 @@ struct str_builder; static void str_builder_init(struct str_builder *b); static void str_builder(struct str_builder *b, const char *s, ...);+static char *strclone(const char *str);  static void futhark_panic(int eval, const char *fmt, ...) {   va_list ap;@@ -133,6 +134,18 @@   va_start(vl, s); // Must re-init.   vsnprintf(b->str+b->used, b->capacity-b->used, s, vl);   b->used += needed;+}+++static char *strclone(const char *str) {+  size_t size = strlen(str) + 1;+  char *copy = (char*) malloc(size);+  if (copy == NULL) {+    return NULL;+  }++  memcpy(copy, str, size);+  return copy; }  // End of util.h.

@@ -59,7 +59,7 @@    _cmd_inputs(args) {     var entry = this._get_arg(args, 0);-    var inputs = this._get_entry_point(entry)[0];+    var inputs = this._get_entry_point(entry)[1];     for (var i = 0; i < inputs.length; i++) {       console.log(inputs[i]);     }@@ -67,7 +67,7 @@    _cmd_outputs(args) {     var entry = this._get_arg(args, 0);-    var outputs = this._get_entry_point(entry)[1];+    var outputs = this._get_entry_point(entry)[2];     for (var i = 0; i < outputs.length; i++) {       console.log(outputs[i]);     }@@ -99,8 +99,8 @@    _cmd_call(args) {     var entry = this._get_entry_point(this._get_arg(args, 0));-    var num_ins = entry[0].length;-    var num_outs = entry[1].length;+    var num_ins = entry[1].length;+    var num_outs = entry[2].length;     var expected_len = 1 + num_outs + num_ins      if (args.length != expected_len) {@@ -121,13 +121,13 @@     }     // Call entry point function from string name     var bef = performance.now()*1000;-    var vals = this.ctx[args[0]].apply(this.ctx, ins);+    var vals = this.ctx[entry[0]].apply(this.ctx, ins);     var aft = performance.now()*1000;     if (num_outs == 1) {-      this._set_var(out_vnames[0], vals, entry[1][0]);+      this._set_var(out_vnames[0], vals, entry[2][0]);     } else {       for (var i = 0; i < out_vnames.length; i++) {-        this._set_var(out_vnames[i], vals[i], entry[1][i]);+        this._set_var(out_vnames[i], vals[i], entry[2][i]);       }     }     console.log("runtime: " + Math.round(aft-bef));

@@ -436,6 +436,9 @@ def futhark_log10_64(x):   return np.float64(np.log10(x)) +def futhark_log1p_64(x):+  return np.float64(np.log1p(x))+ def futhark_sqrt64(x):   return np.sqrt(x) @@ -534,6 +537,9 @@ def futhark_log10_32(x):   return np.float32(np.log10(x)) +def futhark_log1p_32(x):+  return np.float32(np.log1p(x))+ def futhark_sqrt32(x):   return np.float32(np.sqrt(x)) @@ -631,6 +637,9 @@  def futhark_log10_16(x):   return np.float16(np.log10(x))++def futhark_log1p_16(x):+  return np.float16(np.log1p(x))  def futhark_sqrt16(x):   return np.float16(np.sqrt(x))

@@ -39,12 +39,12 @@      def _cmd_inputs(self, args):         entry = self._get_arg(args, 0)-        for t in self._get_entry_point(entry)[0]:+        for t in self._get_entry_point(entry)[1]:             print(t)      def _cmd_outputs(self, args):         entry = self._get_arg(args, 0)-        for t in self._get_entry_point(entry)[1]:+        for t in self._get_entry_point(entry)[2]:             print(t)      def _cmd_dummy(self, args):@@ -65,8 +65,9 @@      def _cmd_call(self, args):         entry = self._get_entry_point(self._get_arg(args, 0))-        num_ins = len(entry[0])-        num_outs = len(entry[1])+        entry_fname = entry[0]+        num_ins = len(entry[1])+        num_outs = len(entry[2])         exp_len = 1 + num_outs + num_ins          if len(args) != exp_len:@@ -81,7 +82,7 @@         ins = [ self._get_var(in_vname) for in_vname in in_vnames ]          try:-            (runtime, vals) = getattr(self._ctx, args[0])(*ins)+            (runtime, vals) = getattr(self._ctx, entry_fname)(*ins)         except Exception as e:             raise self.Failure(str(e)) @@ -175,7 +176,11 @@                  }      def _process_line(self, line):-        words = shlex.split(line)+        lex = shlex.shlex(line)+        lex.quotes = '"'+        lex.whitespace_split = True+        lex.commenters = ''+        words = list(lex)         if words == []:             raise self.Failure('Empty line')         else:

@@ -183,6 +183,12 @@   Just [untyped $ 1 / (isF32 x * log 2)] pdBuiltin "log2_64" [x] =   Just [untyped $ 1 / (isF64 x * log 2)]+pdBuiltin "log1p_16" [x] =+  Just [untyped $ 1 / (isF16 x + 1)]+pdBuiltin "log1p_32" [x] =+  Just [untyped $ 1 / (isF32 x + 1)]+pdBuiltin "log1p_64" [x] =+  Just [untyped $ 1 / (isF64 x + 1)] pdBuiltin "exp16" [x] =   Just [untyped $ exp (isF16 x)] pdBuiltin "exp32" [x] =

@@ -1,5 +1,4 @@ {-# LANGUAGE TypeFamilies #-}-{-# LANGUAGE UndecidableInstances #-}  module Futhark.AD.Fwd (fwdJVP) where @@ -7,7 +6,6 @@ import Control.Monad.RWS.Strict import Control.Monad.State.Strict import Data.Bifunctor (second)-import Data.Kind qualified import Data.List (transpose) import Data.List.NonEmpty (NonEmpty (..)) import Data.Map qualified as M@@ -90,15 +88,11 @@   modify $ \env -> env {stateTans = M.insert v v' (stateTans env)}  class TanBuilder a where-  type Bundled a :: Data.Kind.Type-  type Bundled a = [a]   newTan :: a -> ADM a-  bundleNew :: a -> ADM (Bundled a)+  bundleNew :: a -> ADM [a] -instance (Monoid (Bundled a), TanBuilder a) => TanBuilder [a] where-  type Bundled [a] = Bundled a-  newTan = mapM newTan-  bundleNew = fmap mconcat . mapM bundleNew+bundleNewList :: TanBuilder a => [a] -> ADM [a]+bundleNewList = fmap mconcat . mapM bundleNew  instance TanBuilder (PatElem (TypeBase s u)) where   newTan (PatElem p t)@@ -117,11 +111,12 @@       then pure [pe']       else pure [pe, pe'] -instance TanBuilder (Pat (TypeBase s u)) where-  type Bundled (Pat (TypeBase s u)) = Pat (TypeBase s u)-  newTan (Pat pes) = Pat <$> newTan pes-  bundleNew (Pat pes) = Pat <$> bundleNew pes+newTanPat :: TanBuilder (PatElem t) => Pat t -> ADM (Pat t)+newTanPat (Pat pes) = Pat <$> mapM newTan pes +bundleNewPat :: TanBuilder (PatElem t) => Pat t -> ADM (Pat t)+bundleNewPat (Pat pes) = Pat <$> bundleNewList pes+ instance TanBuilder (Param (TypeBase s u)) where   newTan (Param _ p t) = do     PatElem p' t' <- newTan $ PatElem p t@@ -142,10 +137,8 @@     pure $ zip b [x, x_tan]  class Tangent a where-  type BundledTan a :: Data.Kind.Type-  type BundledTan a = [a]   tangent :: a -> ADM a-  bundleTan :: a -> ADM (BundledTan a)+  bundleTan :: a -> ADM [a]  instance Tangent (TypeBase s u) where   tangent = tanType@@ -157,10 +150,8 @@         t' <- tangent t         pure [t, t'] -instance (Monoid (BundledTan a), Tangent a) => Tangent [a] where-  type BundledTan [a] = BundledTan a-  tangent = mapM tangent-  bundleTan = (mconcat <$>) . mapM bundleTan+bundleTangents :: Tangent a => [a] -> ADM [a]+bundleTangents = (mconcat <$>) . mapM bundleTan  instance Tangent VName where   tangent v = do@@ -192,7 +183,7 @@  basicFwd :: Pat Type -> StmAux () -> BasicOp -> ADM () basicFwd pat aux op = do-  pat_tan <- newTan pat+  pat_tan <- newTanPat pat   case op of     SubExp se -> do       se_tan <- tangent se@@ -201,7 +192,7 @@       se_tan <- tangent se       addStm $ Let pat_tan aux $ BasicOp $ Opaque opaqueop se_tan     ArrayLit ses t -> do-      ses_tan <- tangent ses+      ses_tan <- mapM tangent ses       addStm $ Let pat_tan aux $ BasicOp $ ArrayLit ses_tan t     UnOp unop x -> do       let t = unOpType unop@@ -233,7 +224,7 @@       addStm $ Let pat_tan aux $ BasicOp $ Update safety arr_tan slice se_tan     Concat d (arr :| arrs) w -> do       arr_tan <- tangent arr-      arrs_tans <- tangent arrs+      arrs_tans <- mapM tangent arrs       addStm $ Let pat_tan aux $ BasicOp $ Concat d (arr_tan :| arrs_tans) w     Copy arr -> do       arr_tan <- tangent arr@@ -261,11 +252,11 @@  fwdLambda :: Lambda SOACS -> ADM (Lambda SOACS) fwdLambda l@(Lambda params body ret) =-  Lambda <$> bundleNew params <*> inScopeOf l (fwdBody body) <*> bundleTan ret+  Lambda <$> bundleNewList params <*> inScopeOf l (fwdBody body) <*> bundleTangents ret  fwdStreamLambda :: Lambda SOACS -> ADM (Lambda SOACS) fwdStreamLambda l@(Lambda params body ret) =-  Lambda <$> ((take 1 params ++) <$> bundleNew (drop 1 params)) <*> inScopeOf l (fwdBody body) <*> bundleTan ret+  Lambda <$> ((take 1 params ++) <$> bundleNewList (drop 1 params)) <*> inScopeOf l (fwdBody body) <*> bundleTangents ret  interleave :: [a] -> [a] -> [a] interleave xs ys = concat $ transpose [xs, ys]@@ -280,8 +271,8 @@  fwdSOAC :: Pat Type -> StmAux () -> SOAC SOACS -> ADM () fwdSOAC pat aux (Screma size xs (ScremaForm scs reds f)) = do-  pat' <- bundleNew pat-  xs' <- bundleTan xs+  pat' <- bundleNewPat pat+  xs' <- bundleTangents xs   scs' <- mapM fwdScan scs   reds' <- mapM fwdRed reds   f' <- fwdLambda f@@ -307,34 +298,34 @@             redNeutral = redNeutral red `interleave` map Var neutral_tans           } fwdSOAC pat aux (Stream size xs nes lam) = do-  pat' <- bundleNew pat+  pat' <- bundleNewPat pat   lam' <- fwdStreamLambda lam-  xs' <- bundleTan xs+  xs' <- bundleTangents xs   nes_tan <- mapM (fmap Var . zeroFromSubExp) nes   let nes' = interleave nes nes_tan   addStm $ Let pat' aux $ Op $ Stream size xs' nes' lam' fwdSOAC pat aux (Hist w arrs ops bucket_fun) = do-  pat' <- bundleNew pat+  pat' <- bundleNewPat pat   ops' <- mapM fwdHist ops   bucket_fun' <- fwdHistBucket bucket_fun-  arrs' <- bundleTan arrs+  arrs' <- bundleTangents arrs   addStm $ Let pat' aux $ Op $ Hist w arrs' ops' bucket_fun'   where     n_indices = sum $ map (shapeRank . histShape) ops     fwdBodyHist (Body _ stms res) = buildBody_ $ do       mapM_ fwdStm stms       let (res_is, res_vs) = splitAt n_indices res-      (res_is ++) <$> bundleTan res_vs+      (res_is ++) <$> bundleTangents res_vs     fwdHistBucket l@(Lambda params body ret) =       let (r_is, r_vs) = splitAt n_indices ret        in Lambda-            <$> bundleNew params+            <$> bundleNewList params             <*> inScopeOf l (fwdBodyHist body)-            <*> ((r_is ++) <$> bundleTan r_vs)+            <*> ((r_is ++) <$> bundleTangents r_vs)      fwdHist :: HistOp SOACS -> ADM (HistOp SOACS)     fwdHist (HistOp shape rf dest nes op) = do-      dest' <- bundleTan dest+      dest' <- bundleTangents dest       nes_tan <- mapM (fmap Var . zeroFromSubExp) nes       op' <- fwdLambda op       pure $@@ -347,8 +338,8 @@           } fwdSOAC (Pat pes) aux (Scatter w ivs lam as) = do   as_tan <- mapM (\(s, n, a) -> do a_tan <- tangent a; pure (s, n, a_tan)) as-  pes_tan <- newTan pes-  ivs' <- bundleTan ivs+  pes_tan <- mapM newTan pes+  ivs' <- bundleTangents ivs   let (as_ws, as_ns, _as_vs) = unzip3 as       n_indices = sum $ zipWith (*) as_ns $ map length as_ws   lam' <- fwdScatterLambda n_indices lam@@ -357,8 +348,8 @@   where     fwdScatterLambda :: Int -> Lambda SOACS -> ADM (Lambda SOACS)     fwdScatterLambda n_indices (Lambda params body ret) = do-      params' <- bundleNew params-      ret_tan <- tangent $ drop n_indices ret+      params' <- bundleNewList params+      ret_tan <- mapM tangent $ drop n_indices ret       body' <- fwdBodyScatter n_indices body       let indices = concat $ replicate 2 $ take n_indices ret           ret' = indices ++ drop n_indices ret ++ ret_tan@@ -367,7 +358,7 @@     fwdBodyScatter n_indices (Body _ stms res) = do       (res_tan, stms') <- collectStms $ do         mapM_ fwdStm stms-        tangent $ drop n_indices res+        mapM tangent $ drop n_indices res       let indices = concat $ replicate 2 $ take n_indices res           res' = indices ++ drop n_indices res ++ res_tan       pure $ mkBody stms' res'@@ -378,8 +369,8 @@  fwdStm :: Stm SOACS -> ADM () fwdStm (Let pat aux (BasicOp (UpdateAcc acc i x))) = do-  pat' <- bundleNew pat-  x' <- bundleTan x+  pat' <- bundleNewPat pat+  x' <- bundleTangents x   acc_tan <- tangent acc   addStm $ Let pat' aux $ BasicOp $ UpdateAcc acc_tan i x' fwdStm stm@(Let pat aux (BasicOp e)) = do@@ -392,7 +383,7 @@       addStm stm       arg_tans <-         zipWith primExpFromSubExp argts <$> mapM (tangent . fst) args-      pat_tan <- newTan pat+      pat_tan <- newTanPat pat       let arg_pes = zipWith primExpFromSubExp argts (map fst args)       case pdBuiltin f arg_pes of         Nothing ->@@ -414,27 +405,27 @@ fwdStm (Let pat aux (Match ses cases defbody (MatchDec ret ifsort))) = do   cases' <- slocal' $ mapM (traverse fwdBody) cases   defbody' <- slocal' $ fwdBody defbody-  pat' <- bundleNew pat-  ret' <- bundleTan ret+  pat' <- bundleNewPat pat+  ret' <- bundleTangents ret   addStm $ Let pat' aux $ Match ses cases' defbody' $ MatchDec ret' ifsort fwdStm (Let pat aux (DoLoop val_pats loop@(WhileLoop v) body)) = do-  val_pats' <- bundleNew val_pats-  pat' <- bundleNew pat+  val_pats' <- bundleNewList val_pats+  pat' <- bundleNewPat pat   body' <-     localScope (scopeOfFParams (map fst val_pats) <> scopeOf loop) . slocal' $       fwdBody body   addStm $ Let pat' aux $ DoLoop val_pats' (WhileLoop v) body' fwdStm (Let pat aux (DoLoop val_pats loop@(ForLoop i it bound loop_vars) body)) = do-  pat' <- bundleNew pat-  val_pats' <- bundleNew val_pats-  loop_vars' <- bundleNew loop_vars+  pat' <- bundleNewPat pat+  val_pats' <- bundleNewList val_pats+  loop_vars' <- bundleNewList loop_vars   body' <-     localScope (scopeOfFParams (map fst val_pats) <> scopeOf loop) . slocal' $       fwdBody body   addStm $ Let pat' aux $ DoLoop val_pats' (ForLoop i it bound loop_vars') body' fwdStm (Let pat aux (WithAcc inputs lam)) = do   inputs' <- forM inputs $ \(shape, arrs, op) -> do-    arrs_tan <- tangent arrs+    arrs_tan <- mapM tangent arrs     op' <- case op of       Nothing -> pure Nothing       Just (op_lam, nes) -> do@@ -445,7 +436,7 @@             let op_lam'' = Lambda (removeIndexTans (shapeRank shape) ps) body ret             pure $ Just (op_lam'', interleave nes nes_tan)     pure (shape, arrs <> arrs_tan, op')-  pat' <- bundleNew pat+  pat' <- bundleNewPat pat   lam' <- fwdLambda lam   addStm $ Let pat' aux $ WithAcc inputs' lam'   where@@ -459,17 +450,17 @@ fwdBody :: Body SOACS -> ADM (Body SOACS) fwdBody (Body _ stms res) = buildBody_ $ do   mapM_ fwdStm stms-  bundleTan res+  bundleTangents res  fwdBodyTansLast :: Body SOACS -> ADM (Body SOACS) fwdBodyTansLast (Body _ stms res) = buildBody_ $ do   mapM_ fwdStm stms-  (res <>) <$> tangent res+  (res <>) <$> mapM tangent res  fwdJVP :: MonadFreshNames m => Scope SOACS -> Lambda SOACS -> m (Lambda SOACS) fwdJVP scope l@(Lambda params body ret) =   runADM . localScope scope . inScopeOf l $ do-    params_tan <- newTan params+    params_tan <- mapM newTan params     body_tan <- fwdBodyTansLast body-    ret_tan <- tangent ret+    ret_tan <- mapM tangent ret     pure $ Lambda (params ++ params_tan) body_tan (ret <> ret_tan)

@@ -4,7 +4,6 @@   ( printAction,     printAliasesAction,     printLastUseGPU,-    printLastUseGPUSS,     printFusionGraph,     printInterferenceGPU,     printMemAliasGPU,@@ -55,10 +54,8 @@ import Futhark.IR import Futhark.IR.GPUMem (GPUMem) import Futhark.IR.MCMem (MCMem)-import Futhark.IR.Prop.Aliases import Futhark.IR.SOACS (SOACS) import Futhark.IR.SeqMem (SeqMem)-import Futhark.Optimise.ArrayShortCircuiting.LastUse qualified as LastUseSS import Futhark.Optimise.Fusion.GraphRep qualified import Futhark.Util (runProgramWithExitCode, unixEnvironment) import Futhark.Version (versionString)@@ -77,7 +74,7 @@     }  -- | Print the result to stdout, alias annotations.-printAliasesAction :: (ASTRep rep, CanBeAliased (Op rep)) => Action rep+printAliasesAction :: AliasableRep rep => Action rep printAliasesAction =   Action     { actionName = "Prettyprint",@@ -91,21 +88,12 @@   Action     { actionName = "print last use gpu",       actionDescription = "Print last use information on gpu.",-      actionProcedure = liftIO . putStrLn . prettyString . M.toList . LastUse.analyseGPUMem-    }---- | Print last use information to stdout.-printLastUseGPUSS :: Action GPUMem-printLastUseGPUSS =-  Action-    { actionName = "print last use ss gpu",-      actionDescription = "Print last use ss information on gpu.",       actionProcedure =         liftIO           . putStrLn           . prettyString           . bimap M.toList (M.toList . fmap M.toList)-          . LastUseSS.lastUseGPUMem+          . LastUse.lastUseGPUMem           . aliasAnalysis     } 

@@ -10,10 +10,12 @@ -- the building blocks do). module Futhark.Analysis.Alias   ( aliasAnalysis,+    AliasableRep,      -- * Ad-hoc utilities     analyseFun,     analyseStms,+    analyseStm,     analyseExp,     analyseBody,     analyseLambda,@@ -26,7 +28,7 @@  -- | Perform alias analysis on a Futhark program. aliasAnalysis ::-  (ASTRep rep, CanBeAliased (Op rep)) =>+  AliasableRep rep =>   Prog rep ->   Prog (Aliases rep) aliasAnalysis prog =@@ -37,7 +39,7 @@  -- | Perform alias analysis on function. analyseFun ::-  (ASTRep rep, CanBeAliased (Op rep)) =>+  AliasableRep rep =>   FunDef rep ->   FunDef (Aliases rep) analyseFun (FunDef entry attrs fname restype params body) =@@ -47,9 +49,7 @@  -- | Perform alias analysis on Body. analyseBody ::-  ( ASTRep rep,-    CanBeAliased (Op rep)-  ) =>+  AliasableRep rep =>   AliasTable ->   Body rep ->   Body (Aliases rep)@@ -59,20 +59,26 @@  -- | Perform alias analysis on statements. analyseStms ::-  (ASTRep rep, CanBeAliased (Op rep)) =>+  AliasableRep rep =>   AliasTable ->   Stms rep ->   (Stms (Aliases rep), AliasesAndConsumed) analyseStms orig_aliases =-  foldl' f (mempty, (orig_aliases, mempty)) . stmsToList+  withoutBound . foldl' f (mempty, (orig_aliases, mempty)) . stmsToList   where+    withoutBound (stms, (aliases, consumed)) =+      let bound = foldMap (namesFromList . patNames . stmPat) stms+          consumed' = consumed `namesSubtract` bound+       in (stms, (aliases, consumed'))+     f (stms, aliases) stm =       let stm' = analyseStm (fst aliases) stm           atable' = trackAliases aliases stm'        in (stms <> oneStm stm', atable') +-- | Perform alias analysis on statement. analyseStm ::-  (ASTRep rep, CanBeAliased (Op rep)) =>+  AliasableRep rep =>   AliasTable ->   Stm rep ->   Stm (Aliases rep)@@ -84,7 +90,7 @@  -- | Perform alias analysis on expression. analyseExp ::-  (ASTRep rep, CanBeAliased (Op rep)) =>+  AliasableRep rep =>   AliasTable ->   Exp rep ->   Exp (Aliases rep)@@ -124,7 +130,7 @@  -- | Perform alias analysis on lambda. analyseLambda ::-  (ASTRep rep, CanBeAliased (Op rep)) =>+  AliasableRep rep =>   AliasTable ->   Lambda rep ->   Lambda (Aliases rep)

@@ -66,6 +66,7 @@     ViewL (..),     ArrayTransform (..),     transformFromExp,+    transformToExp,     soacToStream,   ) where@@ -233,6 +234,22 @@   Just (v, Replicate cs shape) transformFromExp _ _ = Nothing +-- | Turn an array transform on an array back into an expression.+transformToExp :: (Monad m, HasScope rep m) => ArrayTransform -> VName -> m (Certs, Exp rep)+transformToExp (Replicate cs n) ia =+  pure (cs, BasicOp $ Futhark.Replicate n (Var ia))+transformToExp (Rearrange cs perm) ia = do+  r <- arrayRank <$> lookupType ia+  pure (cs, BasicOp $ Futhark.Rearrange (perm ++ [length perm .. r - 1]) ia)+transformToExp (Reshape cs k shape) ia = do+  pure (cs, BasicOp $ Futhark.Reshape k shape ia)+transformToExp (ReshapeOuter cs k shape) ia = do+  shape' <- reshapeOuter shape 1 . arrayShape <$> lookupType ia+  pure (cs, BasicOp $ Futhark.Reshape k shape' ia)+transformToExp (ReshapeInner cs k shape) ia = do+  shape' <- reshapeInner shape 1 . arrayShape <$> lookupType ia+  pure (cs, BasicOp $ Futhark.Reshape k shape' ia)+ -- | One array input to a SOAC - a SOAC may have multiple inputs, but -- all are of this form.  Only the array inputs are expressed with -- this type; other arguments, such as initial accumulator values, are@@ -285,24 +302,16 @@ addInitialTransforms ts (Input ots a t) = Input (ts <> ots) a t  applyTransform :: MonadBuilder m => ArrayTransform -> VName -> m VName-applyTransform (Replicate cs n) ia =-  certifying cs . letExp "repeat" . BasicOp $-    Futhark.Replicate n (Var ia)-applyTransform (Rearrange cs perm) ia = do-  r <- arrayRank <$> lookupType ia-  certifying cs . letExp "rearrange" . BasicOp $-    Futhark.Rearrange (perm ++ [length perm .. r - 1]) ia-applyTransform (Reshape cs k shape) ia =-  certifying cs . letExp "reshape" . BasicOp $-    Futhark.Reshape k shape ia-applyTransform (ReshapeOuter cs k shape) ia = do-  shape' <- reshapeOuter shape 1 . arrayShape <$> lookupType ia-  certifying cs . letExp "reshape_outer" . BasicOp $-    Futhark.Reshape k shape' ia-applyTransform (ReshapeInner cs k shape) ia = do-  shape' <- reshapeInner shape 1 . arrayShape <$> lookupType ia-  certifying cs . letExp "reshape_inner" . BasicOp $-    Futhark.Reshape k shape' ia+applyTransform tr ia = do+  (cs, e) <- transformToExp tr ia+  certifying cs $ letExp s e+  where+    s = case tr of+      Replicate {} -> "replicate"+      Rearrange {} -> "rearrange"+      Reshape {} -> "reshape"+      ReshapeOuter {} -> "reshape_outer"+      ReshapeInner {} -> "reshape_inner"  applyTransforms :: MonadBuilder m => ArrayTransforms -> VName -> m VName applyTransforms (ArrayTransforms ts) a = foldlM (flip applyTransform) a ts

@@ -12,7 +12,8 @@ import Data.Maybe (catMaybes, fromMaybe, mapMaybe) import Data.Set (Set) import Data.Set qualified as S-import Futhark.Analysis.LastUse (LastUseMap)+import Futhark.Analysis.Alias qualified as AnlAls+import Futhark.Analysis.LastUse (LUTabFun) import Futhark.Analysis.LastUse qualified as LastUse import Futhark.Analysis.MemAlias qualified as MemAlias import Futhark.IR.GPUMem@@ -39,7 +40,7 @@  analyseStm ::   LocalScope GPUMem m =>-  LastUseMap ->+  LUTabFun ->   InUse ->   Stm GPUMem ->   m (InUse, LastUsed, Graph VName)@@ -103,7 +104,7 @@  analyseExp ::   LocalScope GPUMem m =>-  LastUseMap ->+  LUTabFun ->   InUse ->   Exp GPUMem ->   m (InUse, LastUsed, Graph VName)@@ -122,7 +123,7 @@  analyseKernelBody ::   LocalScope GPUMem m =>-  LastUseMap ->+  LUTabFun ->   InUse ->   KernelBody GPUMem ->   m (InUse, LastUsed, Graph VName)@@ -130,7 +131,7 @@  analyseBody ::   LocalScope GPUMem m =>-  LastUseMap ->+  LUTabFun ->   InUse ->   Body GPUMem ->   m (InUse, LastUsed, Graph VName)@@ -138,7 +139,7 @@  analyseStms ::   LocalScope GPUMem m =>-  LastUseMap ->+  LUTabFun ->   InUse ->   Stms GPUMem ->   m (InUse, LastUsed, Graph VName)@@ -151,7 +152,7 @@  analyseSegOp ::   LocalScope GPUMem m =>-  LastUseMap ->+  LUTabFun ->   InUse ->   SegOp lvl GPUMem ->   m (InUse, LastUsed, Graph VName)@@ -168,7 +169,7 @@  segWithBinOps ::   LocalScope GPUMem m =>-  LastUseMap ->+  LUTabFun ->   InUse ->   [SegBinOp GPUMem] ->   KernelBody GPUMem ->@@ -184,7 +185,7 @@  analyseSegBinOp ::   LocalScope GPUMem m =>-  LastUseMap ->+  LUTabFun ->   InUse ->   SegBinOp GPUMem ->   m (InUse, LastUsed, Graph VName)@@ -193,7 +194,7 @@  analyseHistOp ::   LocalScope GPUMem m =>-  LastUseMap ->+  LUTabFun ->   InUse ->   HistOp GPUMem ->   m (InUse, LastUsed, Graph VName)@@ -202,7 +203,7 @@  analyseLambda ::   LocalScope GPUMem m =>-  LastUseMap ->+  LUTabFun ->   InUse ->   Lambda GPUMem ->   m (InUse, LastUsed, Graph VName)@@ -213,14 +214,14 @@ analyseProgGPU prog = onConsts (progConsts prog) <> foldMap onFun (progFuns prog)   where     (consts_aliases, funs_aliases) = MemAlias.analyzeGPUMem prog-    lumap = LastUse.analyseGPUMem prog+    (lumap_consts, lumap) = LastUse.lastUseGPUMem $ AnlAls.aliasAnalysis prog     onFun f =       applyAliases (fromMaybe mempty $ M.lookup (funDefName f) funs_aliases) $-        runReader (analyseGPU lumap $ bodyStms $ funDefBody f) $+        runReader (analyseGPU (lumap M.! funDefName f) $ bodyStms $ funDefBody f) $           scopeOf f     onConsts stms =       applyAliases consts_aliases $-        runReader (analyseGPU lumap stms) (mempty :: Scope GPUMem)+        runReader (analyseGPU lumap_consts stms) (mempty :: Scope GPUMem)  applyAliases :: MemAlias.MemAliases -> Graph VName -> Graph VName applyAliases aliases =@@ -237,7 +238,7 @@ -- within, and the resulting graph. analyseGPU ::   LocalScope GPUMem m =>-  LastUseMap ->+  LUTabFun ->   Stms GPUMem ->   m (Graph VName) analyseGPU lumap stms = do@@ -304,7 +305,7 @@  analyseGPU' ::   LocalScope GPUMem m =>-  LastUseMap ->+  LUTabFun ->   Stms GPUMem ->   m (InUse, LastUsed, Graph VName) analyseGPU' lumap stms =

@@ -1,240 +1,445 @@ {-# LANGUAGE TypeFamilies #-} --- | Provides last-use analysis for Futhark programs.+-- | Last use analysis for array short circuiting+--+-- Last-Use analysis of a Futhark program in aliased explicit-memory lore form.+-- Takes as input such a program or a function and produces a 'M.Map VName+-- Names', in which the key identified the let stmt, and the list argument+-- identifies the variables that were lastly used in that stmt.  Note that the+-- results of a body do not have a last use, and neither do a function+-- parameters if it happens to not be used inside function's body.  Such cases+-- are supposed to be treated separately. module Futhark.Analysis.LastUse-  ( LastUseMap,-    LastUse,-    Used,-    analyseGPUMem,-    analyseSeqMem,+  ( lastUseSeqMem,+    lastUseGPUMem,+    lastUseMCMem,+    LUTabFun,+    LUTabProg,   ) where  import Control.Monad.Reader-import Data.Bifunctor (bimap, first)+import Control.Monad.State.Strict+import Data.Bifunctor (bimap) import Data.Function ((&))-import Data.Map (Map)-import Data.Map qualified as M-import Data.Tuple-import Futhark.Analysis.Alias (aliasAnalysis)+import Data.Map.Strict qualified as M+import Data.Maybe+import Data.Sequence (Seq (..)) import Futhark.IR.Aliases import Futhark.IR.GPUMem+import Futhark.IR.GPUMem qualified as GPU+import Futhark.IR.MCMem+import Futhark.IR.MCMem qualified as MC import Futhark.IR.SeqMem+import Futhark.Optimise.ArrayShortCircuiting.DataStructs+import Futhark.Util --- | 'LastUseMap' tells which names were last used in a given statement.--- Statements are uniquely identified by the 'VName' of the first value--- parameter in the statement pattern. 'Names' is the set of names last used.-type LastUseMap = Map VName Names+-- | Maps a name indentifying a Stm to the last uses in that Stm.+type LUTabFun = M.Map VName Names --- | 'LastUse' is a mapping from a 'VName' to the statement identifying it's--- last use. 'LastUseMap' is the inverse of 'LastUse'.-type LastUse = Map VName VName+-- | LU-table for the constants, and for each function.+type LUTabProg = (LUTabFun, M.Map Name LUTabFun) --- | 'Used' is the set of 'VName' that were used somewhere in a--- statement, body or otherwise.-type Used = Names+type LastUseOp rep = Op (Aliases rep) -> Names -> LastUseM rep (LUTabFun, Names, Names) -type LastUseOp rep =-  VName -> (LastUse, Used) -> Op (Aliases rep) -> LastUseM rep+-- | 'LastUseReader' allows us to abstract over representations by supplying the+-- 'onOp' function.+data LastUseReader rep = LastUseReader+  { onOp :: LastUseOp rep,+    scope :: Scope (Aliases rep)+  } -newtype Env rep = Env {envLastUseOp :: LastUseOp rep}+-- | Maps a variable or memory block to its aliases.+type AliasTab = M.Map VName Names -type LastUseM rep = Reader (Env rep) (LastUse, Used)+newtype LastUseM rep a = LastUseM (StateT AliasTab (Reader (LastUseReader rep)) a)+  deriving+    ( Monad,+      Functor,+      Applicative,+      MonadReader (LastUseReader rep),+      MonadState AliasTab+    ) -analyseGPUMem :: Prog GPUMem -> LastUseMap-analyseGPUMem = analyseProg analyseGPUOp+instance RepTypes (Aliases rep) => HasScope (Aliases rep) (LastUseM rep) where+  askScope = asks scope -analyseSeqMem :: Prog SeqMem -> LastUseMap-analyseSeqMem = analyseProg analyseSeqOp+instance RepTypes (Aliases rep) => LocalScope (Aliases rep) (LastUseM rep) where+  localScope sc (LastUseM m) = LastUseM $ do+    local (\rd -> rd {scope = scope rd <> sc}) m -analyseGPUOp :: LastUseOp GPUMem-analyseGPUOp pat_name (lumap, used) (Alloc se sp) = do-  let nms = (freeIn se <> freeIn sp) `namesSubtract` used-  pure (insertNames pat_name nms lumap, used <> nms)-analyseGPUOp pat_name (lumap, used) (Inner (SizeOp sop)) = do-  let nms = freeIn sop `namesSubtract` used-  pure (insertNames pat_name nms lumap, used <> nms)-analyseGPUOp _ (lumap, used) (Inner (OtherOp ())) =-  pure (lumap, used)-analyseGPUOp pat_name (lumap, used) (Inner (SegOp (SegMap lvl _ tps body))) = do-  (lumap', used') <- analyseKernelBody (lumap, used) body-  let nms = (freeIn lvl <> freeIn tps) `namesSubtract` used'-  pure (insertNames pat_name nms lumap', used' <> nms)-analyseGPUOp pat_name (lumap, used) (Inner (SegOp (SegRed lvl _ binops tps body))) =-  segOpHelper pat_name lumap used lvl binops tps body-analyseGPUOp pat_name (lumap, used) (Inner (SegOp (SegScan lvl _ binops tps body))) =-  segOpHelper pat_name lumap used lvl binops tps body-analyseGPUOp pat_name (lumap, used) (Inner (SegOp (SegHist lvl _ binops tps body))) = do-  (lumap', used') <- foldM analyseHistOp (lumap, used) $ reverse binops-  (lumap'', used'') <- analyseKernelBody (lumap', used') body-  let nms = (freeIn lvl <> freeIn tps) `namesSubtract` used''-  pure (insertNames pat_name nms lumap'', used'' <> nms)-analyseGPUOp pat_name (lumap, used) (Inner (GPUBody ts body)) = do-  (lumap', used') <- analyseBody lumap used body-  let nms = freeIn ts-  pure (insertNames pat_name nms lumap', used' <> nms)+type Constraints rep =+  ( LocalScope (Aliases rep) (LastUseM rep),+    HasMemBlock (Aliases rep),+    AliasableRep rep+  ) -segOpHelper ::-  (FreeIn (OpWithAliases (Op rep)), ASTRep rep) =>-  VName ->-  LastUse ->-  Used ->-  SegLevel ->-  [SegBinOp (Aliases rep)] ->-  [Type] ->+runLastUseM :: LastUseOp rep -> LastUseM rep a -> a+runLastUseM onOp (LastUseM m) =+  runReader (evalStateT m mempty) (LastUseReader onOp mempty)++aliasLookup :: VName -> LastUseM rep Names+aliasLookup vname =+  gets $ fromMaybe mempty . M.lookup vname++lastUseProg ::+  Constraints rep =>+  Prog (Aliases rep) ->+  LastUseM rep LUTabProg+lastUseProg prog =+  let bound_in_consts =+        progConsts prog+          & concatMap (patNames . stmPat)+          & namesFromList+      consts = progConsts prog+      funs = progFuns prog+   in inScopeOf consts $ do+        (consts_lu, _) <- lastUseStms consts mempty mempty+        lus <- mapM (lastUseFun bound_in_consts) funs+        pure (consts_lu, M.fromList $ zip (map funDefName funs) lus)++lastUseFun ::+  Constraints rep =>+  Names ->+  FunDef (Aliases rep) ->+  LastUseM rep LUTabFun+lastUseFun bound_in_consts f =+  inScopeOf f $ fst <$> lastUseBody (funDefBody f) (mempty, bound_in_consts)++-- | Perform last-use analysis on a 'Prog' in 'SeqMem'+lastUseSeqMem :: Prog (Aliases SeqMem) -> LUTabProg+lastUseSeqMem = runLastUseM lastUseSeqOp . lastUseProg++-- | Perform last-use analysis on a 'Prog' in 'GPUMem'+lastUseGPUMem :: Prog (Aliases GPUMem) -> LUTabProg+lastUseGPUMem = runLastUseM (lastUseMemOp lastUseGPUOp) . lastUseProg++-- | Perform last-use analysis on a 'Prog' in 'MCMem'+lastUseMCMem :: Prog (Aliases MCMem) -> LUTabProg+lastUseMCMem = runLastUseM (lastUseMemOp lastUseMCOp) . lastUseProg++-- | Performing the last-use analysis on a body.+--+-- The implementation consists of a bottom-up traversal of the body's statements+-- in which the the variables lastly used in a statement are computed as the+-- difference between the free-variables in that stmt and the set of variables+-- known to be used after that statement.+lastUseBody ::+  Constraints rep =>+  -- | The body of statements+  Body (Aliases rep) ->+  -- | The current last-use table, tupled with the known set of already used names+  (LUTabFun, Names) ->+  -- | The result is:+  --      (i) an updated last-use table,+  --     (ii) an updated set of used names (including the binding).+  LastUseM rep (LUTabFun, Names)+lastUseBody bdy@(Body _ stms result) (lutab, used_nms) =+  -- perform analysis bottom-up in bindings: results are known to be used,+  -- hence they are added to the used_nms set.+  inScopeOf stms $ do+    (lutab', _) <-+      lastUseStms stms (lutab, used_nms) $+        namesToList $+          freeIn $+            map resSubExp result+    -- Clean up the used names by recomputing the aliasing transitive-closure+    -- of the free names in body based on the current alias table @alstab@.+    used_in_body <- aliasTransitiveClosure $ freeIn bdy+    pure (lutab', used_nms <> used_in_body)++-- | Performing the last-use analysis on a body.+--+-- The implementation consists of a bottom-up traversal of the body's statements+-- in which the the variables lastly used in a statement are computed as the+-- difference between the free-variables in that stmt and the set of variables+-- known to be used after that statement.+lastUseKernelBody ::+  Constraints rep =>+  -- | The body of statements   KernelBody (Aliases rep) ->-  LastUseM rep-segOpHelper pat_name lumap used lvl binops tps body = do-  (lumap', used') <- foldM analyseSegBinOp (lumap, used) $ reverse binops-  (lumap'', used'') <- analyseKernelBody (lumap', used') body-  let nms = (freeIn lvl <> freeIn tps) `namesSubtract` used''-  pure (insertNames pat_name nms lumap'', used'' <> nms)+  -- | The current last-use table, tupled with the known set of already used names+  (LUTabFun, Names) ->+  -- | The result is:+  --      (i) an updated last-use table,+  --     (ii) an updated set of used names (including the binding).+  LastUseM rep (LUTabFun, Names)+lastUseKernelBody bdy@(KernelBody _ stms result) (lutab, used_nms) =+  inScopeOf stms $ do+    -- perform analysis bottom-up in bindings: results are known to be used,+    -- hence they are added to the used_nms set.+    (lutab', _) <-+      lastUseStms stms (lutab, used_nms) $ namesToList $ freeIn result+    -- Clean up the used names by recomputing the aliasing transitive-closure+    -- of the free names in body based on the current alias table @alstab@.+    used_in_body <- aliasTransitiveClosure $ freeIn bdy+    pure (lutab', used_nms <> used_in_body) -analyseSeqOp :: LastUseOp SeqMem-analyseSeqOp pat_name (lumap, used) (Alloc se sp) = do-  let nms = (freeIn se <> freeIn sp) `namesSubtract` used-  pure (insertNames pat_name nms lumap, used <> nms)-analyseSeqOp _ (lumap, used) (Inner ()) =-  pure (lumap, used)+lastUseStms ::+  Constraints rep =>+  Stms (Aliases rep) ->+  (LUTabFun, Names) ->+  [VName] ->+  LastUseM rep (LUTabFun, Names)+lastUseStms Empty (lutab, nms) res_nms = do+  aliases <- concatMapM aliasLookup res_nms+  pure (lutab, nms <> aliases <> namesFromList res_nms)+lastUseStms (stm@(Let pat _ e) :<| stms) (lutab, nms) res_nms =+  inScopeOf stm $ do+    let extra_alias = case e of+          BasicOp (Update _ old _ _) -> oneName old+          BasicOp (FlatUpdate old _ _) -> oneName old+          _ -> mempty+    -- We build up aliases top-down+    updateAliasing extra_alias pat+    -- But compute last use bottom-up+    (lutab', nms') <- lastUseStms stms (lutab, nms) res_nms+    (lutab'', nms'') <- lastUseStm stm (lutab', nms')+    pure (lutab'', nms'') --- | Analyses a program to return a last-use map, mapping each simple statement--- in the program to the values that were last used within that statement, and--- the set of all `VName` that were used inside.-analyseProg :: (CanBeAliased (Op rep), Mem rep inner) => LastUseOp rep -> Prog rep -> LastUseMap-analyseProg onOp prog =-  runReader helper (Env onOp)-  where-    helper = do-      let bound_in_consts =-            progConsts prog-              & concatMap (patNames . stmPat)-              & namesFromList-          prog_alias = aliasAnalysis prog-          consts = progConsts prog_alias-          funs = progFuns prog_alias-      (consts_lu, _) <- analyseStms mempty mempty consts-      (lus, _) <- mconcat <$> mapM (analyseFun mempty bound_in_consts) funs-      pure $ flipMap $ consts_lu <> lus+lastUseStm ::+  Constraints rep =>+  Stm (Aliases rep) ->+  (LUTabFun, Names) ->+  LastUseM rep (LUTabFun, Names)+lastUseStm (Let pat _ e) (lutab, used_nms) = do+  -- analyse the expression and get the+  --  (i)  a new last-use table (in case the @e@ contains bodies of stmts)+  -- (ii) the set of variables lastly used in the current binding.+  -- (iii)  aliased transitive-closure of used names, and+  (lutab', last_uses, used_nms') <- lastUseExp e used_nms+  sc <- asks scope+  let lu_mems =+        namesToList last_uses+          & mapMaybe (`getScopeMemInfo` sc)+          & map memName+          & namesFromList+          & flip namesSubtract used_nms -analyseFun :: (FreeIn (OpWithAliases (Op rep)), ASTRep rep) => LastUse -> Used -> FunDef (Aliases rep) -> LastUseM rep-analyseFun lumap used fun = do-  (lumap', used') <- analyseBody lumap used $ funDefBody fun-  pure (lumap', used' <> freeIn (funDefParams fun))+  -- filter-out the binded names from the set of used variables,+  -- since they go out of scope, and update the last-use table.+  let patnms = patNames pat+      used_nms'' = used_nms' `namesSubtract` namesFromList patnms+      lutab'' =+        M.union lutab' $ M.insert (head patnms) (last_uses <> lu_mems) lutab+  pure (lutab'', used_nms'') -analyseStms :: (FreeIn (OpWithAliases (Op rep)), ASTRep rep) => LastUse -> Used -> Stms (Aliases rep) -> LastUseM rep-analyseStms lumap used stms = foldM analyseStm (lumap, used) $ reverse $ stmsToList stms+-------------------------------- -analyseStm :: (FreeIn (OpWithAliases (Op rep)), ASTRep rep) => (LastUse, Used) -> Stm (Aliases rep) -> LastUseM rep-analyseStm (lumap0, used0) (Let pat _ e) = do-  let (lumap', used') = patElems pat & foldl helper (lumap0, used0)-  analyseExp (lumap', used') e+-- | Last-Use Analysis for an expression.+lastUseExp ::+  Constraints rep =>+  -- | The expression to analyse+  Exp (Aliases rep) ->+  -- | The set of used names "after" this expression+  Names ->+  -- | Result:+  --    1. an extra LUTab recording the last use for expression's inner bodies,+  --    2. the set of last-used vars in the expression at this level,+  --    3. the updated used names, now including expression's free vars.+  LastUseM rep (LUTabFun, Names, Names)+lastUseExp (Match _ cases body _) used_nms = do+  -- For an if-then-else, we duplicate the last use at each body level, meaning+  -- we record the last use of the outer statement, and also the last use in the+  -- statement in the inner bodies. We can safely ignore the if-condition as it is+  -- a boolean scalar.+  (lutab_cases, used_cases) <-+    bimap mconcat mconcat . unzip+      <$> mapM (flip lastUseBody (M.empty, used_nms) . caseBody) cases+  (lutab', body_used_nms) <- lastUseBody body (M.empty, used_nms)+  let free_in_body = freeIn body+  let free_in_cases = freeIn cases+  let used_nms' = used_cases <> body_used_nms+  (_, last_used_arrs) <- lastUsedInNames used_nms $ free_in_body <> free_in_cases+  pure (lutab_cases <> lutab', last_used_arrs, used_nms')+lastUseExp (DoLoop var_ses lf body) used_nms0 = inScopeOf lf $ do+  free_in_body <- aliasTransitiveClosure $ freeIn body+  -- compute the aliasing transitive closure of initializers that are not last-uses+  var_inis <- catMaybes <$> mapM (initHelper (free_in_body <> used_nms0)) var_ses+  let -- To record last-uses inside the loop body, we call 'lastUseBody' with used-names+      -- being:  (free_in_body - loop-variants-a) + used_nms0. As such we disable cases b)+      -- and c) to produce loop-variant last uses inside the loop, and also we prevent+      -- the free-loop-variables to having last uses inside the loop.+      free_in_body' = free_in_body `namesSubtract` namesFromList (map fst var_inis)+      used_nms = used_nms0 <> free_in_body' <> freeIn (bodyResult body)+  (body_lutab, _) <- lastUseBody body (mempty, used_nms)++  -- add var_inis_a to the body_lutab, i.e., record the last-use of+  -- initializer in the corresponding loop variant.+  let lutab_res = body_lutab <> M.fromList var_inis++      -- the result used names are:+      fpar_nms = namesFromList $ map (identName . paramIdent . fst) var_ses+      used_nms' = (free_in_body <> freeIn (map snd var_ses)) `namesSubtract` fpar_nms+      used_nms_res = used_nms0 <> used_nms' <> freeIn (bodyResult body)++      -- the last-uses at loop-statement level are the loop free variables that+      -- do not belong to @used_nms0@; this includes the initializers of b), @lu_ini_b@+      lu_arrs = used_nms' `namesSubtract` used_nms0+  pure (lutab_res, lu_arrs, used_nms_res)   where-    helper (lumap_acc, used_acc) (PatElem name (aliases, _)) =-      -- Any aliases of `name` should have the same last-use as `name`-      ( case M.lookup name lumap_acc of-          Just name' ->-            insertNames name' (unAliases aliases) lumap_acc-          Nothing -> lumap_acc,-        used_acc <> unAliases aliases-      )+    initHelper free_and_used (fp, se) = do+      names <- aliasTransitiveClosure $ maybe mempty oneName $ subExpVar se+      if names `namesIntersect` free_and_used+        then pure Nothing+        else pure $ Just (identName $ paramIdent fp, names)+lastUseExp (Op op) used_nms = do+  on_op <- reader onOp+  on_op op used_nms+lastUseExp e used_nms = do+  let free_in_e = freeIn e+  (used_nms', lu_vars) <- lastUsedInNames used_nms free_in_e+  pure (M.empty, lu_vars, used_nms') -    addAliasesFromBodyRes (lumap_acc, used_acc) (PatElem {}, Constant _) = (lumap_acc, used_acc)-    addAliasesFromBodyRes (lumap_acc, used_acc) (PatElem name _, Var body_res) =-      -- Any aliases of `name` should have the same last-use as `name`-      ( case M.lookup name lumap_acc of-          Just name' ->-            insertNames name' (oneName body_res) lumap_acc-          Nothing -> lumap_acc,-        used_acc <> oneName body_res-      )+lastUseMemOp ::+  (inner (Aliases rep) -> Names -> LastUseM rep (LUTabFun, Names, Names)) ->+  MemOp inner (Aliases rep) ->+  Names ->+  LastUseM rep (LUTabFun, Names, Names)+lastUseMemOp _ (Alloc se sp) used_nms = do+  let free_in_e = freeIn se <> freeIn sp+  (used_nms', lu_vars) <- lastUsedInNames used_nms free_in_e+  pure (M.empty, lu_vars, used_nms')+lastUseMemOp onInner (Inner op) used_nms = onInner op used_nms -    pat_name = patElemName $ head $ patElems pat+lastUseSegOp ::+  Constraints rep =>+  SegOp lvl (Aliases rep) ->+  Names ->+  LastUseM rep (LUTabFun, Names, Names)+lastUseSegOp (SegMap _ _ tps kbody) used_nms = do+  (used_nms', lu_vars) <- lastUsedInNames used_nms $ freeIn tps+  (body_lutab, used_nms'') <- lastUseKernelBody kbody (mempty, used_nms')+  pure (body_lutab, lu_vars, used_nms' <> used_nms'')+lastUseSegOp (SegRed _ _ sbos tps kbody) used_nms = do+  (lutab_sbo, lu_vars_sbo, used_nms_sbo) <- lastUseSegBinOp sbos used_nms+  (used_nms', lu_vars) <- lastUsedInNames used_nms_sbo $ freeIn tps+  (body_lutab, used_nms'') <- lastUseKernelBody kbody (mempty, used_nms')+  pure (M.union lutab_sbo body_lutab, lu_vars <> lu_vars_sbo, used_nms_sbo <> used_nms' <> used_nms'')+lastUseSegOp (SegScan _ _ sbos tps kbody) used_nms = do+  (lutab_sbo, lu_vars_sbo, used_nms_sbo) <- lastUseSegBinOp sbos used_nms+  (used_nms', lu_vars) <- lastUsedInNames used_nms_sbo $ freeIn tps+  (body_lutab, used_nms'') <- lastUseKernelBody kbody (mempty, used_nms')+  pure (M.union lutab_sbo body_lutab, lu_vars <> lu_vars_sbo, used_nms_sbo <> used_nms' <> used_nms'')+lastUseSegOp (SegHist _ _ hos tps kbody) used_nms = do+  (lutab_sbo, lu_vars_sbo, used_nms_sbo) <- lastUseHistOp hos used_nms+  (used_nms', lu_vars) <- lastUsedInNames used_nms_sbo $ freeIn tps+  (body_lutab, used_nms'') <- lastUseKernelBody kbody (mempty, used_nms')+  pure (M.union lutab_sbo body_lutab, lu_vars <> lu_vars_sbo, used_nms_sbo <> used_nms' <> used_nms'') -    analyseExp (lumap, used) (BasicOp _) = do-      let nms = freeIn e `namesSubtract` used-      pure (insertNames pat_name nms lumap, used <> nms)-    analyseExp (lumap, used) (Apply _ args _ _) = do-      let nms = freeIn $ map fst args-      pure (insertNames pat_name nms lumap, used <> nms)-    analyseExp (lumap, used) (Match ses cases defbody dec) = do-      (lumap_cases, used_cases) <--        bimap mconcat mconcat . unzip-          <$> mapM (analyseBody lumap used . caseBody) cases-      (lumap_defbody, used_defbody) <- analyseBody lumap used defbody-      let (lumap', used') =-            (lumap_defbody <> lumap_cases, used_cases <> used_defbody)-              & flip (foldl addAliasesFromBodyRes) (zip (patElems pat) (map resSubExp $ bodyResult defbody))-          nms = (freeIn ses <> freeIn dec) `namesSubtract` used'-      pure-        ( insertNames pat_name nms lumap',-          used' <> nms-        )-    analyseExp (lumap, used) (DoLoop merge form body) = do-      let (lumap', used') =-            zip (patElems pat) (map resSubExp $ bodyResult body)-              & foldl addAliasesFromBodyRes (lumap, used)-      (lumap'', used'') <- analyseBody lumap' used' body-      let nms = (freeIn merge <> freeIn form) `namesSubtract` used''-      pure (insertNames pat_name nms lumap'', used'' <> nms)-    analyseExp (lumap, used) (Op op) = do-      onOp <- asks envLastUseOp-      onOp pat_name (lumap, used) op-    analyseExp (lumap, used) (WithAcc _ l) =-      analyseLambda (lumap, used) l+lastUseGPUOp :: HostOp NoOp (Aliases GPUMem) -> Names -> LastUseM GPUMem (LUTabFun, Names, Names)+lastUseGPUOp (GPU.OtherOp NoOp) used_nms =+  pure (mempty, mempty, used_nms)+lastUseGPUOp (SizeOp sop) used_nms = do+  (used_nms', lu_vars) <- lastUsedInNames used_nms $ freeIn sop+  pure (mempty, lu_vars, used_nms')+lastUseGPUOp (GPUBody tps body) used_nms = do+  (used_nms', lu_vars) <- lastUsedInNames used_nms $ freeIn tps+  (body_lutab, used_nms'') <- lastUseBody body (mempty, used_nms')+  pure (body_lutab, lu_vars, used_nms' <> used_nms'')+lastUseGPUOp (SegOp op) used_nms =+  lastUseSegOp op used_nms -analyseBody :: (FreeIn (OpWithAliases (Op rep)), ASTRep rep) => LastUse -> Used -> Body (Aliases rep) -> LastUseM rep-analyseBody lumap used (Body _ stms result) = do-  let used' = used <> freeIn result-  analyseStms lumap used' stms+lastUseMCOp :: MCOp NoOp (Aliases MCMem) -> Names -> LastUseM MCMem (LUTabFun, Names, Names)+lastUseMCOp (MC.OtherOp NoOp) used_nms =+  pure (mempty, mempty, used_nms)+lastUseMCOp (MC.ParOp par_op op) used_nms = do+  (lutab_par_op, lu_vars_par_op, used_names_par_op) <-+    maybe (pure mempty) (`lastUseSegOp` used_nms) par_op+  (lutab_op, lu_vars_op, used_names_op) <-+    lastUseSegOp op used_nms+  pure+    ( lutab_par_op <> lutab_op,+      lu_vars_par_op <> lu_vars_op,+      used_names_par_op <> used_names_op+    ) -analyseKernelBody ::-  (FreeIn (OpWithAliases (Op rep)), ASTRep rep) =>-  (LastUse, Used) ->-  KernelBody (Aliases rep) ->-  LastUseM rep-analyseKernelBody (lumap, used) (KernelBody _ stms result) =-  let used' = used <> freeIn result-   in analyseStms lumap used' stms+lastUseSegBinOp ::+  Constraints rep =>+  [SegBinOp (Aliases rep)] ->+  Names ->+  LastUseM rep (LUTabFun, Names, Names)+lastUseSegBinOp sbos used_nms = do+  (lutab, lu_vars, used_nms') <- unzip3 <$> mapM helper sbos+  pure (mconcat lutab, mconcat lu_vars, mconcat used_nms')+  where+    helper (SegBinOp _ l@(Lambda _ body _) neutral shp) = inScopeOf l $ do+      (used_nms', lu_vars) <- lastUsedInNames used_nms $ freeIn neutral <> freeIn shp+      (body_lutab, used_nms'') <- lastUseBody body (mempty, used_nms')+      pure (body_lutab, lu_vars, used_nms'') -analyseSegBinOp ::-  (FreeIn (OpWithAliases (Op rep)), ASTRep rep) =>-  (LastUse, Used) ->-  SegBinOp (Aliases rep) ->-  LastUseM rep-analyseSegBinOp (lumap, used) (SegBinOp _ lambda neutral shp) = do-  (lumap', used') <- analyseLambda (lumap, used) lambda-  let nms = (freeIn neutral <> freeIn shp) `namesSubtract` used'-  pure (lumap', used' <> nms)+lastUseHistOp ::+  Constraints rep =>+  [HistOp (Aliases rep)] ->+  Names ->+  LastUseM rep (LUTabFun, Names, Names)+lastUseHistOp hos used_nms = do+  (lutab, lu_vars, used_nms') <- unzip3 <$> mapM helper hos+  pure (mconcat lutab, mconcat lu_vars, mconcat used_nms')+  where+    helper (HistOp shp rf dest neutral shp' l@(Lambda _ body _)) = inScopeOf l $ do+      (used_nms', lu_vars) <- lastUsedInNames used_nms $ freeIn shp <> freeIn rf <> freeIn dest <> freeIn neutral <> freeIn shp'+      (body_lutab, used_nms'') <- lastUseBody body (mempty, used_nms')+      pure (body_lutab, lu_vars, used_nms'') -analyseHistOp ::-  (FreeIn (OpWithAliases (Op rep)), ASTRep rep) =>-  (LastUse, Used) ->-  HistOp (Aliases rep) ->-  LastUseM rep-analyseHistOp (lumap, used) (HistOp width race dest neutral shp lambda) = do-  (lumap', used') <- analyseLambda (lumap, used) lambda-  let nms =-        ( freeIn width-            <> freeIn race-            <> freeIn dest-            <> freeIn neutral-            <> freeIn shp-        )-          `namesSubtract` used'-  pure (lumap', used' <> nms)+lastUseSeqOp :: Op (Aliases SeqMem) -> Names -> LastUseM SeqMem (LUTabFun, Names, Names)+lastUseSeqOp (Alloc se sp) used_nms = do+  let free_in_e = freeIn se <> freeIn sp+  (used_nms', lu_vars) <- lastUsedInNames used_nms free_in_e+  pure (mempty, lu_vars, used_nms')+lastUseSeqOp (Inner NoOp) used_nms = do+  pure (mempty, mempty, used_nms) -analyseLambda :: (FreeIn (OpWithAliases (Op rep)), ASTRep rep) => (LastUse, Used) -> Lambda (Aliases rep) -> LastUseM rep-analyseLambda (lumap, used) (Lambda params body ret) = do-  (lumap', used') <- analyseBody lumap used body-  let used'' = used' <> freeIn params <> freeIn ret-  pure (lumap', used'')+------------------------------------------------------ -flipMap :: Map VName VName -> Map VName Names-flipMap m =-  M.toList m-    & fmap (swap . first oneName)-    & foldr (uncurry $ M.insertWith (<>)) mempty+-- | Given already used names and newly encountered 'Names', return an updated+-- set used names and the set of names that were last used here.+--+-- For a given name @x@ in the new uses, if neither @x@ nor any of its aliases+-- are present in the set of used names, this is a last use of @x@.+lastUsedInNames ::+  -- | Used names+  Names ->+  -- | New uses+  Names ->+  LastUseM rep (Names, Names)+lastUsedInNames used_nms new_uses = do+  -- a use of an argument x is also a use of any variable in x alias set+  -- so we update the alias-based transitive-closure of used names.+  new_uses_with_aliases <- aliasTransitiveClosure new_uses+  -- if neither a variable x, nor any of its alias set have been used before (in+  -- the backward traversal), then it is a last use of both that variable and+  -- all other variables in its alias set+  last_uses <- filterM isLastUse $ namesToList new_uses+  last_uses' <- aliasTransitiveClosure $ namesFromList last_uses+  pure (used_nms <> new_uses_with_aliases, last_uses')+  where+    isLastUse x = do+      with_aliases <- aliasTransitiveClosure $ oneName x+      pure $ not $ with_aliases `namesIntersect` used_nms -insertNames :: VName -> Names -> LastUse -> LastUse-insertNames name names lumap =-  foldr (flip (M.insertWith $ \_ x -> x) name) lumap $ namesToList names+-- | Compute the transitive closure of the aliases of a set of 'Names'.+aliasTransitiveClosure :: Names -> LastUseM rep Names+aliasTransitiveClosure args = do+  res <- foldl (<>) args <$> mapM aliasLookup (namesToList args)+  if res == args+    then pure res+    else aliasTransitiveClosure res++-- | For each 'PatElem' in the 'Pat', add its aliases to the 'AliasTab' in+-- 'LastUseM'. Additionally, 'Names' are added as aliases of all the 'PatElemT'.+updateAliasing ::+  AliasesOf dec =>+  -- | Extra names that all 'PatElem' should alias.+  Names ->+  -- | Pattern to process+  Pat dec ->+  LastUseM rep ()+updateAliasing extra_aliases =+  mapM_ update . patElems+  where+    update :: AliasesOf dec => PatElem dec -> LastUseM rep ()+    update (PatElem name dec) = do+      let aliases = aliasesOf dec+      aliases' <- aliasTransitiveClosure $ extra_aliases <> aliases+      modify $ M.insert name aliases'