futhark 0.19.7 → 0.20.1

raw patch · 227 files changed

+12429/−7920 lines, 227 filesdep +base16-bytestringdep +cryptohash-md5dep +githashdep −gitrevdep −hashabledep −utf8-stringdep ~file-embeddep ~futhark-datadep ~mainland-prettyPVP ok

version bump matches the API change (PVP)

Dependencies added: base16-bytestring, cryptohash-md5, githash, half

Dependencies removed: gitrev, hashable, utf8-string

Dependency ranges changed: file-embed, futhark-data, mainland-pretty

API changes (from Hackage documentation)

Files

• docs/binary-data-format.rst +3/−4
• docs/c-api.rst +11/−5
• docs/conf.py +2/−0
• docs/index.rst +3/−0
• docs/js-api.rst +142/−0
• docs/language-reference.rst +41/−5
• docs/man/futhark-cuda.rst +7/−1
• docs/man/futhark-repl.rst +2/−1
• docs/man/futhark-run.rst +4/−12
• docs/man/futhark-wasm-multicore.rst +99/−0
• docs/man/futhark-wasm.rst +98/−0
• docs/man/futhark.rst +1/−1
• docs/server-protocol.rst +2/−2
• futhark.cabal +20/−11
• prelude/math.fut +157/−29
• prelude/prelude.fut +7/−6
• rts/c/cuda.h +12/−0
• rts/c/half.h +238/−0
• rts/c/opencl.h +8/−2
• rts/c/scalar.h +1654/−0
• rts/c/scalar_f16.h +627/−0
• rts/c/scheduler.h +8/−1
• rts/c/server.h +8/−2
• rts/c/tuning.h +1/−1
• rts/c/util.h +15/−8
• rts/c/values.h +58/−18
• rts/javascript/server.js +251/−0
• rts/javascript/values.js +165/−0
• rts/javascript/wrapperclasses.js +83/−0
• rts/python/memory.py +5/−5
• rts/python/opencl.py +10/−3
• rts/python/scalar.py +118/−11
• rts/python/values.py +37/−0
• src/Futhark/Actions.hs +157/−29
• src/Futhark/Analysis/Alias.hs +1/−1
• src/Futhark/Analysis/CallGraph.hs +70/−19
• src/Futhark/Analysis/DataDependencies.hs +3/−3
• src/Futhark/Analysis/HORep/MapNest.hs +7/−7
• src/Futhark/Analysis/HORep/SOAC.hs +42/−60
• src/Futhark/Analysis/Interference.hs +38/−21
• src/Futhark/Analysis/LastUse.hs +5/−5
• src/Futhark/Analysis/Metrics.hs +26/−24
• src/Futhark/Analysis/PrimExp.hs +30/−0
• src/Futhark/Analysis/PrimExp/Convert.hs +12/−12
• src/Futhark/Analysis/PrimExp/Parse.hs +3/−0
• src/Futhark/Analysis/Rephrase.hs +6/−6
• src/Futhark/Analysis/SymbolTable.hs +15/−16
• src/Futhark/Analysis/UsageTable.hs +6/−4
• src/Futhark/Bench.hs +2/−0
• src/Futhark/Binder.hs +0/−259
• src/Futhark/Binder/Class.hs +0/−173
• src/Futhark/Builder.hs +246/−0
• src/Futhark/Builder/Class.hs +173/−0
• src/Futhark/CLI/Bench.hs +1/−6
• src/Futhark/CLI/Dataset.hs +19/−1
• src/Futhark/CLI/Dev.hs +83/−56
• src/Futhark/CLI/Literate.hs +72/−61
• src/Futhark/CLI/Misc.hs +4/−3
• src/Futhark/CLI/MulticoreWASM.hs +19/−0
• src/Futhark/CLI/REPL.hs +6/−25
• src/Futhark/CLI/Run.hs +2/−2
• src/Futhark/CLI/Test.hs +8/−7
• src/Futhark/CLI/WASM.hs +19/−0
• src/Futhark/CodeGen/Backends/CCUDA.hs +15/−19
• src/Futhark/CodeGen/Backends/CCUDA/Boilerplate.hs +9/−9
• src/Futhark/CodeGen/Backends/COpenCL.hs +38/−29
• src/Futhark/CodeGen/Backends/COpenCL/Boilerplate.hs +14/−15
• src/Futhark/CodeGen/Backends/GenericC.hs +203/−196
• src/Futhark/CodeGen/Backends/GenericC/CLI.hs +20/−21
• src/Futhark/CodeGen/Backends/GenericC/Server.hs +10/−10
• src/Futhark/CodeGen/Backends/GenericPython.hs +70/−37
• src/Futhark/CodeGen/Backends/GenericPython/AST.hs +5/−4
• src/Futhark/CodeGen/Backends/GenericPython/Definitions.hs +0/−31
• src/Futhark/CodeGen/Backends/GenericWASM.hs +338/−0
• src/Futhark/CodeGen/Backends/MulticoreC.hs +142/−138
• src/Futhark/CodeGen/Backends/MulticoreWASM.hs +61/−0
• src/Futhark/CodeGen/Backends/PyOpenCL.hs +6/−4
• src/Futhark/CodeGen/Backends/PyOpenCL/Boilerplate.hs +3/−11
• src/Futhark/CodeGen/Backends/SequentialC.hs +3/−1
• src/Futhark/CodeGen/Backends/SequentialC/Boilerplate.hs +1/−1
• src/Futhark/CodeGen/Backends/SequentialPython.hs +2/−1
• src/Futhark/CodeGen/Backends/SequentialWASM.hs +61/−0
• src/Futhark/CodeGen/Backends/SimpleRep.hs +69/−928
• src/Futhark/CodeGen/ImpCode.hs +16/−1
• src/Futhark/CodeGen/ImpCode/OpenCL.hs +3/−2
• src/Futhark/CodeGen/ImpGen.hs +289/−274
• src/Futhark/CodeGen/ImpGen/GPU.hs +54/−56
• src/Futhark/CodeGen/ImpGen/GPU/Base.hs +103/−98
• src/Futhark/CodeGen/ImpGen/GPU/SegHist.hs +11/−11
• src/Futhark/CodeGen/ImpGen/GPU/SegMap.hs +3/−3
• src/Futhark/CodeGen/ImpGen/GPU/SegRed.hs +11/−11
• src/Futhark/CodeGen/ImpGen/GPU/SegScan.hs +5/−6
• src/Futhark/CodeGen/ImpGen/GPU/SegScan/SinglePass.hs +8/−8
• src/Futhark/CodeGen/ImpGen/GPU/SegScan/TwoPass.hs +8/−8
• src/Futhark/CodeGen/ImpGen/GPU/ToOpenCL.hs +73/−90
• src/Futhark/CodeGen/ImpGen/Multicore.hs +4/−4
• src/Futhark/CodeGen/ImpGen/Multicore/Base.hs +4/−4
• src/Futhark/CodeGen/ImpGen/Multicore/SegHist.hs +14/−14
• src/Futhark/CodeGen/ImpGen/Multicore/SegMap.hs +8/−13
• src/Futhark/CodeGen/ImpGen/Multicore/SegRed.hs +12/−12
• src/Futhark/CodeGen/ImpGen/Multicore/SegScan.hs +17/−17
• src/Futhark/CodeGen/RTS/C.hs +102/−0
• src/Futhark/CodeGen/RTS/JavaScript.hs +24/−0
• src/Futhark/CodeGen/RTS/Python.hs +44/−0
• src/Futhark/CodeGen/SetDefaultSpace.hs +35/−28
• src/Futhark/Construct.hs +90/−120
• src/Futhark/Doc/Generator.hs +2/−2
• src/Futhark/IR/Aliases.hs +43/−81
• src/Futhark/IR/GPU.hs +7/−7
• src/Futhark/IR/GPU/Kernel.hs +0/−346
• src/Futhark/IR/GPU/Op.hs +346/−0
• src/Futhark/IR/GPU/Simplify.hs +1/−1
• src/Futhark/IR/GPUMem.hs +13/−11
• src/Futhark/IR/MC.hs +6/−6
• src/Futhark/IR/MCMem.hs +12/−9
• src/Futhark/IR/Mem.hs +154/−208
• src/Futhark/IR/Mem/IxFun.hs +46/−10
• src/Futhark/IR/Mem/Simplify.hs +30/−38
• src/Futhark/IR/Parse.hs +81/−66
• src/Futhark/IR/Pretty.hs +41/−27
• src/Futhark/IR/Primitive.hs +213/−24
• src/Futhark/IR/Primitive/Parse.hs +3/−0
• src/Futhark/IR/Prop.hs +25/−22
• src/Futhark/IR/Prop/Aliases.hs +14/−12
• src/Futhark/IR/Prop/Names.hs +23/−17
• src/Futhark/IR/Prop/Patterns.hs +21/−51
• src/Futhark/IR/Prop/Scope.hs +5/−6
• src/Futhark/IR/Prop/TypeOf.hs +46/−38
• src/Futhark/IR/Prop/Types.hs +3/−10
• src/Futhark/IR/RetType.hs +2/−2
• src/Futhark/IR/SOACS.hs +9/−9
• src/Futhark/IR/SOACS/SOAC.hs +13/−13
• src/Futhark/IR/SOACS/Simplify.hs +78/−81
• src/Futhark/IR/SegOp.hs +145/−107
• src/Futhark/IR/Seq.hs +6/−6
• src/Futhark/IR/SeqMem.hs +5/−9
• src/Futhark/IR/Syntax.hs +71/−37
• src/Futhark/IR/Syntax/Core.hs +90/−38
• src/Futhark/IR/Traversals.hs +37/−26
• src/Futhark/Internalise.hs +1/−1
• src/Futhark/Internalise/AccurateSizes.hs +5/−5
• src/Futhark/Internalise/Bindings.hs +45/−45
• src/Futhark/Internalise/Defunctionalise.hs +90/−87
• src/Futhark/Internalise/Defunctorise.hs +3/−5
• src/Futhark/Internalise/Exps.hs +214/−96
• src/Futhark/Internalise/FreeVars.hs +10/−14
• src/Futhark/Internalise/Lambdas.hs +4/−4
• src/Futhark/Internalise/LiftLambdas.hs +3/−4
• src/Futhark/Internalise/Monad.hs +14/−14
• src/Futhark/Internalise/Monomorphise.hs +52/−50
• src/Futhark/Optimise/BlkRegTiling.hs +52/−54
• src/Futhark/Optimise/CSE.hs +36/−30
• src/Futhark/Optimise/DoubleBuffer.hs +216/−81
• src/Futhark/Optimise/Fusion.hs +59/−72
• src/Futhark/Optimise/Fusion/Composing.hs +9/−14
• src/Futhark/Optimise/Fusion/LoopKernel.hs +9/−9
• src/Futhark/Optimise/InPlaceLowering.hs +13/−17
• src/Futhark/Optimise/InPlaceLowering/LowerIntoStm.hs +58/−68
• src/Futhark/Optimise/InPlaceLowering/SubstituteIndices.hs +29/−28
• src/Futhark/Optimise/InliningDeadFun.hs +110/−91
• src/Futhark/Optimise/ReuseAllocations.hs +16/−23
• src/Futhark/Optimise/ReuseAllocations/GreedyColoring.hs +11/−13
• src/Futhark/Optimise/Simplify/Engine.hs +78/−116
• src/Futhark/Optimise/Simplify/Rep.hs +19/−25
• src/Futhark/Optimise/Simplify/Rule.hs +10/−11
• src/Futhark/Optimise/Simplify/Rules.hs +80/−91
• src/Futhark/Optimise/Simplify/Rules/BasicOp.hs +53/−46
• src/Futhark/Optimise/Simplify/Rules/ClosedForm.hs +15/−20
• src/Futhark/Optimise/Simplify/Rules/Index.hs +41/−31
• src/Futhark/Optimise/Simplify/Rules/Loop.hs +54/−102
• src/Futhark/Optimise/Simplify/Rules/Simple.hs +55/−52
• src/Futhark/Optimise/Sink.hs +2/−2
• src/Futhark/Optimise/TileLoops.hs +77/−82
• src/Futhark/Optimise/TileLoops/Shared.hs +26/−26
• src/Futhark/Optimise/Unstream.hs +3/−3
• src/Futhark/Pass/ExpandAllocations.hs +80/−73
• src/Futhark/Pass/ExplicitAllocations.hs +397/−517
• src/Futhark/Pass/ExplicitAllocations/GPU.hs +19/−18
• src/Futhark/Pass/ExplicitAllocations/MC.hs +2/−0
• src/Futhark/Pass/ExplicitAllocations/SegOp.hs +4/−4
• src/Futhark/Pass/ExtractKernels.hs +52/−58
• src/Futhark/Pass/ExtractKernels/BlockedKernel.hs +31/−28
• src/Futhark/Pass/ExtractKernels/DistributeNests.hs +109/−125
• src/Futhark/Pass/ExtractKernels/Distribution.hs +39/−41
• src/Futhark/Pass/ExtractKernels/ISRWIM.hs +26/−31
• src/Futhark/Pass/ExtractKernels/Interchange.hs +24/−26
• src/Futhark/Pass/ExtractKernels/Intragroup.hs +25/−19
• src/Futhark/Pass/ExtractKernels/StreamKernel.hs +26/−38
• src/Futhark/Pass/ExtractKernels/ToGPU.hs +2/−2
• src/Futhark/Pass/ExtractMulticore.hs +31/−28
• src/Futhark/Pass/KernelBabysitting.hs +21/−21
• src/Futhark/Passes.hs +3/−1
• src/Futhark/Script.hs +3/−3
• src/Futhark/Test.hs +6/−2
• src/Futhark/Tools.hs +29/−37
• src/Futhark/Transform/FirstOrderTransform.hs +46/−44
• src/Futhark/Transform/Rename.hs +30/−43
• src/Futhark/Transform/Substitute.hs +19/−6
• src/Futhark/TypeCheck.hs +94/−86
• src/Futhark/Util.hs +24/−14
• src/Futhark/Util/Console.hs +0/−5
• src/Futhark/Util/Options.hs +8/−1
• src/Futhark/Util/Pretty.hs +5/−1
• src/Futhark/Version.hs +27/−16
• src/Language/Futhark.hs +2/−2
• src/Language/Futhark/Core.hs +3/−1
• src/Language/Futhark/Interpreter.hs +199/−46
• src/Language/Futhark/Parser/Lexer.x +8/−11
• src/Language/Futhark/Parser/Parser.y +79/−95
• src/Language/Futhark/Pretty.hs +10/−9
• src/Language/Futhark/Prop.hs +139/−67
• src/Language/Futhark/Query.hs +26/−26
• src/Language/Futhark/Syntax.hs +56/−59
• src/Language/Futhark/Traversals.hs +47/−44
• src/Language/Futhark/TypeChecker.hs +86/−58
• src/Language/Futhark/TypeChecker/Match.hs +8/−8
• src/Language/Futhark/TypeChecker/Modules.hs +1/−1
• src/Language/Futhark/TypeChecker/Monad.hs +4/−7
• src/Language/Futhark/TypeChecker/Terms.hs +164/−159
• src/Language/Futhark/TypeChecker/Types.hs +10/−15
• src/Language/Futhark/TypeChecker/Unify.hs +12/−12
• src/futhark.hs +4/−0
• unittests/Futhark/IR/Mem/IxFun/Alg.hs +21/−4
• unittests/Futhark/IR/Mem/IxFunTests.hs +194/−76
• unittests/Futhark/IR/Mem/IxFunWrapper.hs +9/−1
• unittests/Futhark/IR/PrimitiveTests.hs +7/−1
• unittests/Futhark/Optimise/ReuseAllocations/GreedyColoringTests.hs +15/−15

@@ -12,10 +12,8 @@ not used for Futhark programs compiled to libraries, which instead use whichever format is supported by their host language. -Currently reading binary input is only supported for programs-generated by ``futhark c``/``futhark opencl``, and-``futhark py``/``futhark pyopencl``.  It is *not* supported for-``futhark run``.+Currently reading binary input is only supported for compiled programs.+It is *not* supported for ``futhark run``.  You can generate random data in the binary format with ``futhark dataset`` (:ref:`futhark-dataset(1)`).  This tool can also be used to@@ -72,6 +70,7 @@   " u16"   " u32"   " u64"+  " f16"   " f32"   " f64"   "bool"

@@ -148,11 +148,17 @@ ------  Primitive types (``i32``, ``bool``, etc) are mapped directly to their-corresponding C type.  For each distinct array type of primitives-(ignoring sizes), an opaque C struct is defined.  For types that do-not map cleanly to C, including records, sum types, and arrays of-tuples, see :ref:`opaques`.+corresponding C type.  The ``f16`` type is mapped to ``uint16_t``,+because C does not have a standard ``half`` type.  This integer+contains the bitwise representation of the ``f16`` value in the IEEE+754 binary16 format. +For each distinct array type of primitives (ignoring sizes), an opaque+C struct is defined.  Arrays of ``f16`` are presented as containing+``uint16_t`` elements.  For types that do not map cleanly to C,+including records, sum types, and arrays of tuples, see+:ref:`opaques`.+ All array values share a similar API, which is illustrated here for the case of the type ``[]i32``.  The creation/retrieval functions are all asynchronous, so make sure to call :c:func:`futhark_context_sync`@@ -281,7 +287,7 @@  Results in the following C function: -.. c:function:: int futhark_entry_main(struct futhark_context *ctx, int32_t *out0, const struct futhark_i32_1d *in0)+.. c:function:: int futhark_entry_sum(struct futhark_context *ctx, int32_t *out0, const struct futhark_i32_1d *in0)     Asynchronously call the entry point with the given arguments.  Make    sure to call :c:func:`futhark_context_sync` before using the value

@@ -268,6 +268,8 @@     ('man/futhark-cuda', 'futhark-cuda', 'compile Futhark to CUDA', [], 1),     ('man/futhark-python', 'futhark-python', 'compile Futhark to sequential Python', [], 1),     ('man/futhark-pyopencl', 'futhark-pyopencl', 'compile Futhark to Python and OpenCL', [], 1),+    ('man/futhark-wasm', 'futhark-wasm', 'compile Futhark to WebAssembly', [], 1),+    ('man/futhark-wasm-multicore', 'futhark-wasm-multicore', 'compile Futhark to parallel WebAssembly', [], 1),     ('man/futhark-run', 'futhark-run', 'interpret Futhark program', [], 1),     ('man/futhark-repl', 'futhark-repl', 'interactive Futhark read-eval-print-loop', [], 1),     ('man/futhark-test', 'futhark-test', 'test Futhark programs', [], 1),

@@ -31,6 +31,7 @@    usage.rst    language-reference.rst    c-api.rst+   js-api.rst    package-management.rst    server-protocol.rst    c-porting-guide.rst@@ -54,6 +55,8 @@    man/futhark-pkg.rst    man/futhark-pyopencl.rst    man/futhark-python.rst+   man/futhark-wasm.rst+   man/futhark-wasm-multicore.rst    man/futhark-repl.rst    man/futhark-run.rst    man/futhark-literate.rst

@@ -0,0 +1,142 @@+.. _js-api:++JavaScript API Reference+========================++The :ref:`futhark-wasm(1)` and :ref:`futhark-wasm-multicore(1)`+compilers produce JavaScript wrapper code to allow JavaScript programs+to invoke the generated WebAssembly code.  This chapter describes the+API exposed by the wrapper.++First a warning: **the JavaScript API is experimental**.  It may+change incompatibly even in minor versions of the compiler.++A Futhark program ``futlib.fut`` compiled with a WASM backend as a library+with the ``--library`` command line option produces four files:++* ``futlib.c``, ``futlib.h``: Implementation and header C files+  generated by the compiler, similar to ``futhark c``.  You can delete+  these - they are not needed at run-time.+* ``futlib.class.js``: An intermediate build artifact.  Feel free to+  delete it.+* ``futlib.wasm``: A compiled WebAssembly module, which must be+   present at runtime.+* ``futlib.mjs``: An ES6 module that can can be imported by other+  JavaScript code, and implements the API given in the following.++The module exports a function, ``newFutharkContext``, which is a factory+function that returns a Promise producing a ``FutharkContext``+instance (see below).  A simple usage example:++.. code-block:: javascript++   import { newFutharkContext } from './futlib.mjs';+   var fc;+   newFutharkContext().then(x => fc = x);++General concerns+----------------++Memory management is completely manual, as JavaScript does not support+finalizers that could let Futhark hook into the garbage collector.+You are responsible for eventually freeing all objects produced by the+API, using the appropriate methods.++FutharkContext+--------------++FutharkContext is a class that contains information about the context+and configuration from the C API. It has methods for invoking the+Futhark entry points and creating FutharkArrays on the WebAssembly+heap.++.. js:function:: newFutharkContext()++   Asynchronously create a new ``FutharkContext`` object.++.. js:class:: FutharkContext()++   A bookkeeping class representing an instance of a Futhark program.+   Do *not* directly invoke its constructor - always use the+   ``newFutharkContext()`` factory function.++.. js:function:: FutharkContext.free()++   Frees all memory created by the ``FutharkContext`` object. Should+   be called when the ``FutharkContext`` is done being used. It is an+   error use a ``FutharkArray`` or ``FutharkOpaque`` after the+   ``FutharkContext`` on which they were defined has been freed.++Values+------++Numeric types ``u8``, ``u16``, ``u32``, ``i8``, ``i16``, ``i32``, ``f32``,+and ``f64`` are mapped to JavaScript's standard number type. 64-bit integers+``u64``, and ``i64`` are mapped to  ``BigInt``. ``bool`` is mapped to+JavaScript's ``boolean`` type. Arrays are represented by the ``FutharkArray``.+complex types (records, nested tuples, etc) are represented by the+``FutharkOpaque`` class.++FutharkArray+------------++``FutharkArray`` has the following API++.. js:function:: FutharkArray.toArray()++   Returns a nested JavaScript array++.. js:function:: FutharkArray.toTypedArray()++   Returns a flat typed array of the underlying data.++.. js:function:: FutharkArray.shape()++   Returns the shape of the FutharkArray as an array of BigInts.++.. js:function:: FutharkArray.free()++   Frees the memory used by the FutharkArray class++``FutharkContext`` also contains two functions for creating+``FutharkArrays`` from JavaScript arrays, and typed arrays for each+array type that appears in an entry point.  All array types share+similar API methods on the ``FutharkContext``, which is illustrated+here for the case of the type ``[]i32``.++.. js:function:: FutharkContext.new_i32_1d_from_jsarray(jsarray)++  Creates and returns a one-dimensional ``i32`` ``FutharkArray`` representing+  the JavaScript array jsarray++.. js:function:: FutharkContext.new_i32_1d(array, dim1)++  Creates and returns a one-dimensional ``i32`` ``FutharkArray`` representing+  the typed array of array, with the size given by dim1.+++FutharkOpaque+-------------++Complex types (records, nested tuples, etc) are represented by+``FutharkOpaque``.  It has no use outside of being accepted and+returned by entry point functions. For this reason the method only has+one function for freeing the memory when ``FutharkOpaque`` is no+longer used.++.. js:function:: FutharkOpaque.free()++   Frees  memory used by FutharkOpaque. Should be called when Futhark+   Opaque is no longer used.++Entry Points+------------++Each entry point in the compiled futhark program has an entry point method on+the FutharkContext++.. js:function:: FutharkContext.<entry_point_name>(in1, ..., inN)++  The entry point function taking the N arguments of the Futhark entry point+  function, and returns the result. If the result is a tuple the return value+  is an array.

@@ -51,13 +51,12 @@ Boolean literals are written ``true`` and ``false``.  The primitive types in Futhark are the signed integer types ``i8``, ``i16``, ``i32``, ``i64``, the unsigned integer types ``u8``, ``u16``, ``u32``,-``u64``, the floating-point types ``f32``, ``f64``, as well as-``bool``.  An ``f32`` is always a single-precision float and a ``f64``-is a double-precision float.+``u64``, the floating-point types ``f16``, ``f32``, ``f64``, as well+as ``bool``.  .. productionlist::    int_type: "i8" | "i16" | "i32" | "i64" | "u8" | "u16" | "u32" | "u64"-   float_type: "f32" | "f64"+   float_type: "f16" | "f32" | "f64"  Numeric literals can be suffixed with their intended type.  For example ``42i8`` is of type ``i8``, and ``1337e2f64`` is of type@@ -429,6 +428,8 @@    exp:   `atom`       : | `exp` `qualbinop` `exp`       : | `exp` `exp`+      : | "!" `exp`+      : | "-" `exp`       : | `constructor` `exp`*       : | `exp` ":" `type`       : | `exp` ":>" `type`@@ -503,6 +504,10 @@   enclosed in parentheses, rather than an operator section partially   applying the infix operator ``-``. +* Function application and prefix operators bind more tightly than any+  infix operator.  Note that the only prefix operators are ``!`` and+  ``-``, and more cannot be defined.+ * The following table describes the precedence and associativity of   infix operators.  All operators in the same row have the same   precedence.  The rows are listed in increasing order of precedence.@@ -1401,7 +1406,7 @@ .. productionlist::    mod_bind: "module" `id` `mod_param`* "=" [":" mod_type_exp] "=" `mod_exp`    mod_param: "(" `id` ":" `mod_type_exp` ")"-   mod_type_bind: "module" "type" `id` `type_param`* "=" `mod_type_exp`+   mod_type_bind: "module" "type" `id` "=" `mod_type_exp`  Futhark supports an ML-style higher-order module system.  *Modules* can contain types, functions, and other modules and module types.@@ -1657,6 +1662,32 @@  The following expression attributes are supported. +``trace``+.........++Print the value produced by the attributed expression.  Used for+debugging.  Somewhat unreliable outside of the interpreter, and in+particular does not work for GPU device code.++``trace(tag)``+..............++Like ``trace``, but prefix output with *tag*, which must lexically be+an identifier.++``break``+.........++In the interpreter, pause execution *before* evaluating the expression.+No effect for compiled code.++``opaque``+..........++The compiler will treat the attributed expression as a black box.+This is used to work around optimisation deficiencies (or bugs),+although it should hopefully rarely be necessary.+ ``incremental_flattening(no_outer)`` .................................... @@ -1744,6 +1775,11 @@ Do not inline any calls to this function.  If the function is then used within a parallel construct (e.g. ``map``), this will likely prevent the GPU backends from generating working code.++``inline``+..........++Always inline calls to this function.  Spec attributes ~~~~~~~~~~~~~~~

@@ -161,13 +161,19 @@ compiler to compile the generated C program into a binary.  This only works if the C compiler can find the necessary CUDA libraries.  On most systems, CUDA is installed in ``/usr/local/cuda``, which is-usually not part of the default compiler search path.  You may need to+usually not part of the default compiler search path. You may need to set the following environment variables before running ``futhark cuda``::    LIBRARY_PATH=/usr/local/cuda/lib64   LD_LIBRARY_PATH=/usr/local/cuda/lib64/   CPATH=/usr/local/cuda/include++At runtime the generated program must be able to find the CUDA+installation directory, which is normally located at+``/usr/local/cuda``.  If you have CUDA installed elsewhere, set any of+the ``CUDA_HOME``, ``CUDA_ROOT``, or ``CUDA_PATH`` environment+variables to the proper directory.  SEE ALSO ========

@@ -25,7 +25,8 @@ colon.  ``futhark repl`` uses the Futhark interpreter, which grants access to-certain special functions.  See :ref:`futhark-run(1)` for a description.+the ``#[trace]`` and ``#[break]`` attributes.  See+:ref:`futhark-run(1)` for a description.  OPTIONS =======

@@ -19,18 +19,10 @@ output.  ``futhark run`` is very slow, and in practice only useful for testing,-teaching, and experimenting with the language.  Certain special-debugging functions are available in ``futhark run``:--``trace 'a : a -> a``-  Semantically identity, but prints the value on standard output.--``break 'a : a -> a``-  Semantically identity, but interrupts execution at the calling-  point, such that the environment can be inspected.  Continue-  execution by entering an empty input line.  Breakpoints are only-  respected when starting a program from the prompt, not when-  passing a program on the command line.+teaching, and experimenting with the language.  The ``#[trace]`` and+``#[break]`` attributes are fully supported in the interpreter.+Tracing prints values to stdout in contrast to compiled code, which+prints to stderr.  OPTIONS =======

@@ -0,0 +1,99 @@+.. role:: ref(emphasis)++.. _futhark-wasm-multicore(1):++======================+futhark-wasm-multicore+======================++SYNOPSIS+========++futhark wasm-multicore [options...] <program.fut>++DESCRIPTION+===========++``futhark wasm-multicore`` translates a Futhark program to+multi-threaded WebAssembly code by first generating C as ``futhark+c``, and then using Emscripten (``emcc``).  This produces a ``.js``+file that allows the compiled code to be invoked from JavaScript.+Executables implement the Futhark server protocol and can be run with+Node.js.++OPTIONS+=======++-h+  Print help text to standard output and exit.++--entry-point NAME+  Treat this top-level function as an entry point.++--library+  Generate a library instead of an executable.  Appends ``.js``+  to the name indicated by the ``-o`` option to determine output+  file names.++-o outfile+  Where to write the result.  If the source program is named+  ``foo.fut``, this defaults to ``foo``.++--safe+  Ignore ``unsafe`` in program and perform safety checks unconditionally.++--server+  Generate a server-mode executable that reads commands from stdin.++-v verbose+  Enable debugging output.  If compilation fails due to a compiler+  error, the result of the last successful compiler step will be+  printed to standard error.++-V+  Print version information on standard output and exit.++-W+  Do not print any warnings.++--Werror+  Treat warnings as errors.++++ENVIRONMENT VARIABLES+=====================++``CFLAGS``++  Space-separated list of options passed to ``emcc``.  Defaults+  to ``-O3 -std=c99`` if unset.++``EMCFLAGS``++  Space-separated list of options passed to ``emcc``.++EXECUTABLE OPTIONS+==================++The following options are accepted by executables generated by+``futhark wasm-multicore``.++-h, --help++  Print help text to standard output and exit.++-D, --debugging++  Perform possibly expensive internal correctness checks and verbose+  logging.  Implies ``-L``.++-L, --log++  Print various low-overhead logging information to stderr while+  running.++SEE ALSO+========++:ref:`futhark-c(1)`, :ref:`futhark-wasm(1)`

@@ -0,0 +1,98 @@+.. role:: ref(emphasis)++.. _futhark-wasm(1):++============+futhark-wasm+============++SYNOPSIS+========++futhark wasm [options...] <program.fut>++DESCRIPTION+===========++``futhark wasm`` translates a Futhark program to sequential+WebAssembly code by first generating C as ``futhark c``, and then+using Emscripten (``emcc``).  This produces a ``.js`` file that allows+the compiled code to be invoked from JavaScript.  Executables+implement the Futhark server protocol and can be run with Node.js.++OPTIONS+=======++-h+  Print help text to standard output and exit.++--entry-point NAME+  Treat this top-level function as an entry point.++--library+  Generate a library instead of an executable.  Appends ``.js``+  to the name indicated by the ``-o`` option to determine output+  file names.++-o outfile+  Where to write the result.  If the source program is named+  ``foo.fut``, this defaults to ``foo``.++--safe+  Ignore ``unsafe`` in program and perform safety checks unconditionally.++--server+  Generate a server-mode executable that reads commands from stdin.+  This is the default.++-v verbose+  Enable debugging output.  If compilation fails due to a compiler+  error, the result of the last successful compiler step will be+  printed to standard error.++-V+  Print version information on standard output and exit.++-W+  Do not print any warnings.++--Werror+  Treat warnings as errors.++ENVIRONMENT VARIABLES+=====================++``CFLAGS``++  Space-separated list of options passed to ``emcc``.  Defaults+  to ``-O3 -std=c99`` if unset.++``EMCFLAGS``++  Space-separated list of options passed to ``emcc``.++EXECUTABLE OPTIONS+==================++The following options are accepted by executables generated by+``futhark wasm``.++-h, --help++  Print help text to standard output and exit.++-D, --debugging++  Perform possibly expensive internal correctness checks and verbose+  logging.  Implies ``-L``.++-L, --log++  Print various low-overhead logging information to stderr while+  running.+++SEE ALSO+========++:ref:`futhark-c(1)`, :ref:`futhark-wasm-multicore(1)`

@@ -74,4 +74,4 @@ SEE ALSO ======== -:ref:`futhark-opencl(1)`, :ref:`futhark-c(1)`, :ref:`futhark-py(1)`, :ref:`futhark-pyopencl(1)`, :ref:`futhark-dataset(1)`, :ref:`futhark-doc(1)`, :ref:`futhark-test(1)`, :ref:`futhark-bench(1)`, :ref:`futhark-run(1)`, :ref:`futhark-repl(1)`, :ref:`futhark-literate(1)`+:ref:`futhark-opencl(1)`, :ref:`futhark-c(1)`, :ref:`futhark-py(1)`, :ref:`futhark-pyopencl(1)`, :ref:`futhark-wasm(1)`, :ref:`futhark-wasm-multicore(1)`, :ref:`futhark-dataset(1)`, :ref:`futhark-doc(1)`, :ref:`futhark-test(1)`, :ref:`futhark-bench(1)`, :ref:`futhark-run(1)`, :ref:`futhark-repl(1)`, :ref:`futhark-literate(1)`

@@ -71,10 +71,10 @@  The following commands are supported. -``call`` *entry* *o1* ... *oN* *i1* ... *oM*+``call`` *entry* *o1* ... *oN* *i1* ... *iM* ............................................ -Call the given entry point with input from the variables *i1* to *oM*.+Call the given entry point with input from the variables *i1* to *iM*. The results are stored in *o1* to *oN*, which must not already exist.  ``restore`` *file* *v1* *t1* ... *vN* *tN*

@@ -1,6 +1,6 @@ cabal-version: 2.4 name:           futhark-version:        0.19.7+version:        0.20.1 synopsis:       An optimising compiler for a functional, array-oriented language.  description:    Futhark is a small programming language designed to be compiled to@@ -35,6 +35,7 @@ extra-source-files:     rts/c/*.h     rts/futhark-doc/*.css+    rts/javascript/*.js     rts/python/*.py     prelude/*.fut -- Just enough of the docs to build the manpages.@@ -70,8 +71,8 @@       Futhark.Analysis.SymbolTable       Futhark.Analysis.UsageTable       Futhark.Bench-      Futhark.Binder-      Futhark.Binder.Class+      Futhark.Builder+      Futhark.Builder.Class       Futhark.CLI.Autotune       Futhark.CLI.Bench       Futhark.CLI.C@@ -84,6 +85,7 @@       Futhark.CLI.Literate       Futhark.CLI.Misc       Futhark.CLI.Multicore+      Futhark.CLI.MulticoreWASM       Futhark.CLI.OpenCL       Futhark.CLI.Pkg       Futhark.CLI.PyOpenCL@@ -92,6 +94,7 @@       Futhark.CLI.REPL       Futhark.CLI.Run       Futhark.CLI.Test+      Futhark.CLI.WASM       Futhark.CodeGen.Backends.CCUDA       Futhark.CodeGen.Backends.CCUDA.Boilerplate       Futhark.CodeGen.Backends.COpenCL@@ -102,15 +105,20 @@       Futhark.CodeGen.Backends.GenericC.Server       Futhark.CodeGen.Backends.GenericPython       Futhark.CodeGen.Backends.GenericPython.AST-      Futhark.CodeGen.Backends.GenericPython.Definitions       Futhark.CodeGen.Backends.GenericPython.Options+      Futhark.CodeGen.Backends.GenericWASM       Futhark.CodeGen.Backends.MulticoreC+      Futhark.CodeGen.Backends.MulticoreWASM       Futhark.CodeGen.Backends.PyOpenCL       Futhark.CodeGen.Backends.PyOpenCL.Boilerplate       Futhark.CodeGen.Backends.SequentialC       Futhark.CodeGen.Backends.SequentialC.Boilerplate       Futhark.CodeGen.Backends.SequentialPython+      Futhark.CodeGen.Backends.SequentialWASM       Futhark.CodeGen.Backends.SimpleRep+      Futhark.CodeGen.RTS.C+      Futhark.CodeGen.RTS.Python+      Futhark.CodeGen.RTS.JavaScript       Futhark.CodeGen.ImpCode       Futhark.CodeGen.ImpCode.GPU       Futhark.CodeGen.ImpCode.Multicore@@ -150,7 +158,7 @@       Futhark.IR       Futhark.IR.Aliases       Futhark.IR.GPU-      Futhark.IR.GPU.Kernel+      Futhark.IR.GPU.Op       Futhark.IR.GPU.Simplify       Futhark.IR.GPU.Sizes       Futhark.IR.GPUMem@@ -303,25 +311,27 @@     , ansi-terminal >=0.6.3.1     , array >=0.4     , base >=4.13 && <5+    , base16-bytestring     , binary >=0.8.3     , blaze-html >=0.9.0.1     , bytestring >=0.10.8     , bytestring-to-vector >=0.3.0.1     , bmp >=1.2.6.3     , containers >=0.6.2.1+    , cryptohash-md5     , directory >=1.3.0.0     , directory-tree >=0.12.1     , dlist >=0.6.0.1-    , file-embed >=0.0.9+    , file-embed >=0.0.14.0     , filepath >=1.4.1.1     , free >=4.12.4-    , futhark-data >= 1.0.0.1+    , futhark-data >= 1.0.2.0     , futhark-server >= 1.1.0.0-    , gitrev >=1.2.0-    , hashable+    , githash >=0.1.6.1+    , half >= 0.3     , haskeline     , language-c-quote >=0.12-    , mainland-pretty >=0.6.1+    , mainland-pretty >=0.7.1     , cmark-gfm >=0.2.1     , megaparsec >=9.0.0     , mtl >=2.2.1@@ -339,7 +349,6 @@     , time >=1.6.0.1     , transformers >=0.3     , unordered-containers >=0.2.7-    , utf8-string >=1     , vector >=0.12     , vector-binary-instances >=0.2.2.0     , versions >=5.0.0

@@ -17,6 +17,7 @@   val u32: u32 -> t   val u64: u64 -> t +  val f16: f16 -> t   val f32: f32 -> t   val f64: f64 -> t @@ -89,8 +90,10 @@   val &: t -> t -> t   val |: t -> t -> t   val ^: t -> t -> t-  val !: t -> t +  -- | Bitwise negation.+  val not: t -> t+   val <<: t -> t -> t   val >>: t -> t -> t   val >>>: t -> t -> t@@ -111,12 +114,13 @@   val mad_hi: (a: t) -> (b: t) -> (c: t) -> t    -- | Count number of zero bits preceding the most significant set-  -- bit.+  -- bit.  Returns the number of bits in the type if the argument is+  -- zero.   val clz: t -> i32    -- | Count number of trailing zero bits following the least   -- significant set bit.  Returns the number of bits in the type if-  -- the argument is all-zero.+  -- the argument is zero.   val ctz: t -> i32 } @@ -233,6 +237,7 @@   let u32 (x: u32) = intrinsics.itob_i32_bool (intrinsics.sign_i32 x)   let u64 (x: u64) = intrinsics.itob_i64_bool (intrinsics.sign_i64 x) +  let f16 (x: f16) = x != 0f16   let f32 (x: f32) = x != 0f32   let f64 (x: f64) = x != 0f64 @@ -254,7 +259,7 @@   let (x: i8) & (y: i8) = intrinsics.and8 (x, y)   let (x: i8) | (y: i8) = intrinsics.or8 (x, y)   let (x: i8) ^ (y: i8) = intrinsics.xor8 (x, y)-  let ! (x: i8) = intrinsics.complement8 x+  let not (x: i8) = intrinsics.complement8 x    let (x: i8) << (y: i8) = intrinsics.shl8 (x, y)   let (x: i8) >> (y: i8) = intrinsics.ashr8 (x, y)@@ -270,6 +275,7 @@   let u32 (x: u32) = intrinsics.zext_i32_i8 (intrinsics.sign_i32 x)   let u64 (x: u64) = intrinsics.zext_i64_i8 (intrinsics.sign_i64 x) +  let f16 (x: f16) = intrinsics.fptosi_f16_i8 x   let f32 (x: f32) = intrinsics.fptosi_f32_i8 x   let f64 (x: f64) = intrinsics.fptosi_f64_i8 x @@ -283,7 +289,7 @@   let (x: i8) > (y: i8) = intrinsics.slt8 (y, x)   let (x: i8) <= (y: i8) = intrinsics.sle8 (x, y)   let (x: i8) >= (y: i8) = intrinsics.sle8 (y, x)-  let (x: i8) != (y: i8) = intrinsics.! (x == y)+  let (x: i8) != (y: i8) = !(x == y)    let sgn (x: i8) = intrinsics.ssignum8 x   let abs (x: i8) = intrinsics.abs8 x@@ -298,7 +304,7 @@   let num_bits = 8i32   let get_bit (bit: i32) (x: t) = to_i32 ((x >> i32 bit) & i32 1)   let set_bit (bit: i32) (x: t) (b: i32) =-    ((x & i32 (intrinsics.!(1 intrinsics.<< bit))) | i32 (b intrinsics.<< bit))+    ((x & i32 (!(1 intrinsics.<< bit))) | i32 (b intrinsics.<< bit))   let popc = intrinsics.popc8   let mul_hi a b = intrinsics.mul_hi8 (i8 a, i8 b)   let mad_hi a b c = intrinsics.mad_hi8 (i8 a, i8 b, i8 c)@@ -326,7 +332,7 @@   let (x: i16) & (y: i16) = intrinsics.and16 (x, y)   let (x: i16) | (y: i16) = intrinsics.or16 (x, y)   let (x: i16) ^ (y: i16) = intrinsics.xor16 (x, y)-  let ! (x: i16) = intrinsics.complement16 x+  let not (x: i16) = intrinsics.complement16 x    let (x: i16) << (y: i16) = intrinsics.shl16 (x, y)   let (x: i16) >> (y: i16) = intrinsics.ashr16 (x, y)@@ -342,6 +348,7 @@   let u32 (x: u32) = intrinsics.zext_i32_i16 (intrinsics.sign_i32 x)   let u64 (x: u64) = intrinsics.zext_i64_i16 (intrinsics.sign_i64 x) +  let f16 (x: f16) = intrinsics.fptosi_f16_i16 x   let f32 (x: f32) = intrinsics.fptosi_f32_i16 x   let f64 (x: f64) = intrinsics.fptosi_f64_i16 x @@ -355,7 +362,7 @@   let (x: i16) > (y: i16) = intrinsics.slt16 (y, x)   let (x: i16) <= (y: i16) = intrinsics.sle16 (x, y)   let (x: i16) >= (y: i16) = intrinsics.sle16 (y, x)-  let (x: i16) != (y: i16) = intrinsics.! (x == y)+  let (x: i16) != (y: i16) = !(x == y)    let sgn (x: i16) = intrinsics.ssignum16 x   let abs (x: i16) = intrinsics.abs16 x@@ -370,7 +377,7 @@   let num_bits = 16i32   let get_bit (bit: i32) (x: t) = to_i32 ((x >> i32 bit) & i32 1)   let set_bit (bit: i32) (x: t) (b: i32) =-    ((x & i32 (intrinsics.!(1 intrinsics.<< bit))) | i32 (b intrinsics.<< bit))+    ((x & i32 (!(1 intrinsics.<< bit))) | i32 (b intrinsics.<< bit))   let popc = intrinsics.popc16   let mul_hi a b = intrinsics.mul_hi16 (i16 a, i16 b)   let mad_hi a b c = intrinsics.mad_hi16 (i16 a, i16 b, i16 c)@@ -401,7 +408,7 @@   let (x: i32) & (y: i32) = intrinsics.and32 (x, y)   let (x: i32) | (y: i32) = intrinsics.or32 (x, y)   let (x: i32) ^ (y: i32) = intrinsics.xor32 (x, y)-  let ! (x: i32) = intrinsics.complement32 x+  let not (x: i32) = intrinsics.complement32 x    let (x: i32) << (y: i32) = intrinsics.shl32 (x, y)   let (x: i32) >> (y: i32) = intrinsics.ashr32 (x, y)@@ -417,6 +424,7 @@   let u32 (x: u32) = intrinsics.zext_i32_i32 (intrinsics.sign_i32 x)   let u64 (x: u64) = intrinsics.zext_i64_i32 (intrinsics.sign_i64 x) +  let f16 (x: f16) = intrinsics.fptosi_f16_i32 x   let f32 (x: f32) = intrinsics.fptosi_f32_i32 x   let f64 (x: f64) = intrinsics.fptosi_f64_i32 x @@ -430,7 +438,7 @@   let (x: i32) > (y: i32) = intrinsics.slt32 (y, x)   let (x: i32) <= (y: i32) = intrinsics.sle32 (x, y)   let (x: i32) >= (y: i32) = intrinsics.sle32 (y, x)-  let (x: i32) != (y: i32) = intrinsics.! (x == y)+  let (x: i32) != (y: i32) = !(x == y)    let sgn (x: i32) = intrinsics.ssignum32 x   let abs (x: i32) = intrinsics.abs32 x@@ -445,7 +453,7 @@   let num_bits = 32i32   let get_bit (bit: i32) (x: t) = to_i32 ((x >> i32 bit) & i32 1)   let set_bit (bit: i32) (x: t) (b: i32) =-    ((x & i32 (intrinsics.!(1 intrinsics.<< bit))) | i32 (b intrinsics.<< bit))+    ((x & i32 (!(1 intrinsics.<< bit))) | i32 (b intrinsics.<< bit))   let popc = intrinsics.popc32   let mul_hi a b = intrinsics.mul_hi32 (i32 a, i32 b)   let mad_hi a b c = intrinsics.mad_hi32 (i32 a, i32 b, i32 c)@@ -476,7 +484,7 @@   let (x: i64) & (y: i64) = intrinsics.and64 (x, y)   let (x: i64) | (y: i64) = intrinsics.or64 (x, y)   let (x: i64) ^ (y: i64) = intrinsics.xor64 (x, y)-  let ! (x: i64) = intrinsics.complement64 x+  let not (x: i64) = intrinsics.complement64 x    let (x: i64) << (y: i64) = intrinsics.shl64 (x, y)   let (x: i64) >> (y: i64) = intrinsics.ashr64 (x, y)@@ -492,6 +500,7 @@   let u32 (x: u32) = intrinsics.zext_i32_i64 (intrinsics.sign_i32 x)   let u64 (x: u64) = intrinsics.zext_i64_i64 (intrinsics.sign_i64 x) +  let f16 (x: f16) = intrinsics.fptosi_f16_i64 x   let f32 (x: f32) = intrinsics.fptosi_f32_i64 x   let f64 (x: f64) = intrinsics.fptosi_f64_i64 x @@ -505,7 +514,7 @@   let (x: i64) > (y: i64) = intrinsics.slt64 (y, x)   let (x: i64) <= (y: i64) = intrinsics.sle64 (x, y)   let (x: i64) >= (y: i64) = intrinsics.sle64 (y, x)-  let (x: i64) != (y: i64) = intrinsics.! (x == y)+  let (x: i64) != (y: i64) = !(x == y)    let sgn (x: i64) = intrinsics.ssignum64 x   let abs (x: i64) = intrinsics.abs64 x@@ -520,7 +529,7 @@   let num_bits = 64i32   let get_bit (bit: i32) (x: t) = to_i32 ((x >> i32 bit) & i32 1)   let set_bit (bit: i32) (x: t) (b: i32) =-    ((x & i32 (intrinsics.!(1 intrinsics.<< bit))) | intrinsics.zext_i32_i64 (b intrinsics.<< bit))+    ((x & i32 (!(1 intrinsics.<< bit))) | intrinsics.zext_i32_i64 (b intrinsics.<< bit))   let popc = intrinsics.popc64   let mul_hi a b = intrinsics.mul_hi64 (i64 a, i64 b)   let mad_hi a b c = intrinsics.mad_hi64 (i64 a, i64 b, i64 c)@@ -551,7 +560,7 @@   let (x: u8) & (y: u8) = unsign (intrinsics.and8 (sign x, sign y))   let (x: u8) | (y: u8) = unsign (intrinsics.or8 (sign x, sign y))   let (x: u8) ^ (y: u8) = unsign (intrinsics.xor8 (sign x, sign y))-  let ! (x: u8) = unsign (intrinsics.complement8 (sign x))+  let not (x: u8) = unsign (intrinsics.complement8 (sign x))    let (x: u8) << (y: u8) = unsign (intrinsics.shl8 (sign x, sign y))   let (x: u8) >> (y: u8) = unsign (intrinsics.ashr8 (sign x, sign y))@@ -567,6 +576,7 @@   let i32 (x: i32) = unsign (intrinsics.zext_i32_i8 x)   let i64 (x: i64) = unsign (intrinsics.zext_i64_i8 x) +  let f16 (x: f16) = unsign (intrinsics.fptoui_f16_i8 x)   let f32 (x: f32) = unsign (intrinsics.fptoui_f32_i8 x)   let f64 (x: f64) = unsign (intrinsics.fptoui_f64_i8 x) @@ -580,7 +590,7 @@   let (x: u8) > (y: u8) = intrinsics.ult8 (sign y, sign x)   let (x: u8) <= (y: u8) = intrinsics.ule8 (sign x, sign y)   let (x: u8) >= (y: u8) = intrinsics.ule8 (sign y, sign x)-  let (x: u8) != (y: u8) = intrinsics.! (x == y)+  let (x: u8) != (y: u8) = !(x == y)    let sgn (x: u8) = unsign (intrinsics.usignum8 (sign x))   let abs (x: u8) = x@@ -595,7 +605,7 @@   let num_bits = 8i32   let get_bit (bit: i32) (x: t) = to_i32 ((x >> i32 bit) & i32 1)   let set_bit (bit: i32) (x: t) (b: i32) =-    ((x & i32 (intrinsics.!(1 intrinsics.<< bit))) | i32 (b intrinsics.<< bit))+    ((x & i32 (!(1 intrinsics.<< bit))) | i32 (b intrinsics.<< bit))   let popc x = intrinsics.popc8 (sign x)   let mul_hi a b = unsign (intrinsics.mul_hi8 (sign a, sign b))   let mad_hi a b c = unsign (intrinsics.mad_hi8 (sign a, sign b, sign c))@@ -626,7 +636,7 @@   let (x: u16) & (y: u16) = unsign (intrinsics.and16 (sign x, sign y))   let (x: u16) | (y: u16) = unsign (intrinsics.or16 (sign x, sign y))   let (x: u16) ^ (y: u16) = unsign (intrinsics.xor16 (sign x, sign y))-  let ! (x: u16) = unsign (intrinsics.complement16 (sign x))+  let not (x: u16) = unsign (intrinsics.complement16 (sign x))    let (x: u16) << (y: u16) = unsign (intrinsics.shl16 (sign x, sign y))   let (x: u16) >> (y: u16) = unsign (intrinsics.ashr16 (sign x, sign y))@@ -642,6 +652,7 @@   let i32 (x: i32) = unsign (intrinsics.zext_i32_i16 x)   let i64 (x: i64) = unsign (intrinsics.zext_i64_i16 x) +  let f16 (x: f16) = unsign (intrinsics.fptoui_f16_i16 x)   let f32 (x: f32) = unsign (intrinsics.fptoui_f32_i16 x)   let f64 (x: f64) = unsign (intrinsics.fptoui_f64_i16 x) @@ -655,7 +666,7 @@   let (x: u16) > (y: u16) = intrinsics.ult16 (sign y, sign x)   let (x: u16) <= (y: u16) = intrinsics.ule16 (sign x, sign y)   let (x: u16) >= (y: u16) = intrinsics.ule16 (sign y, sign x)-  let (x: u16) != (y: u16) = intrinsics.! (x == y)+  let (x: u16) != (y: u16) = !(x == y)    let sgn (x: u16) = unsign (intrinsics.usignum16 (sign x))   let abs (x: u16) = x@@ -670,7 +681,7 @@   let num_bits = 16i32   let get_bit (bit: i32) (x: t) = to_i32 ((x >> i32 bit) & i32 1)   let set_bit (bit: i32) (x: t) (b: i32) =-    ((x & i32 (intrinsics.!(1 intrinsics.<< bit))) | i32 (b intrinsics.<< bit))+    ((x & i32 (!(1 intrinsics.<< bit))) | i32 (b intrinsics.<< bit))   let popc x = intrinsics.popc16 (sign x)   let mul_hi a b = unsign (intrinsics.mul_hi16 (sign a, sign b))   let mad_hi a b c = unsign (intrinsics.mad_hi16 (sign a, sign b, sign c))@@ -701,7 +712,7 @@   let (x: u32) & (y: u32) = unsign (intrinsics.and32 (sign x, sign y))   let (x: u32) | (y: u32) = unsign (intrinsics.or32 (sign x, sign y))   let (x: u32) ^ (y: u32) = unsign (intrinsics.xor32 (sign x, sign y))-  let ! (x: u32) = unsign (intrinsics.complement32 (sign x))+  let not (x: u32) = unsign (intrinsics.complement32 (sign x))    let (x: u32) << (y: u32) = unsign (intrinsics.shl32 (sign x, sign y))   let (x: u32) >> (y: u32) = unsign (intrinsics.ashr32 (sign x, sign y))@@ -717,6 +728,7 @@   let i32 (x: i32) = unsign (intrinsics.zext_i32_i32 x)   let i64 (x: i64) = unsign (intrinsics.zext_i64_i32 x) +  let f16 (x: f16) = unsign (intrinsics.fptoui_f16_i32 x)   let f32 (x: f32) = unsign (intrinsics.fptoui_f32_i32 x)   let f64 (x: f64) = unsign (intrinsics.fptoui_f64_i32 x) @@ -730,7 +742,7 @@   let (x: u32) > (y: u32) = intrinsics.ult32 (sign y, sign x)   let (x: u32) <= (y: u32) = intrinsics.ule32 (sign x, sign y)   let (x: u32) >= (y: u32) = intrinsics.ule32 (sign y, sign x)-  let (x: u32) != (y: u32) = intrinsics.! (x == y)+  let (x: u32) != (y: u32) = !(x == y)    let sgn (x: u32) = unsign (intrinsics.usignum32 (sign x))   let abs (x: u32) = x@@ -745,7 +757,7 @@   let num_bits = 32i32   let get_bit (bit: i32) (x: t) = to_i32 ((x >> i32 bit) & i32 1)   let set_bit (bit: i32) (x: t) (b: i32) =-    ((x & i32 (intrinsics.!(1 intrinsics.<< bit))) | i32 (b intrinsics.<< bit))+    ((x & i32 (!(1 intrinsics.<< bit))) | i32 (b intrinsics.<< bit))   let popc x = intrinsics.popc32 (sign x)   let mul_hi a b = unsign (intrinsics.mul_hi32 (sign a, sign b))   let mad_hi a b c = unsign (intrinsics.mad_hi32 (sign a, sign b, sign c))@@ -776,7 +788,7 @@   let (x: u64) & (y: u64) = unsign (intrinsics.and64 (sign x, sign y))   let (x: u64) | (y: u64) = unsign (intrinsics.or64 (sign x, sign y))   let (x: u64) ^ (y: u64) = unsign (intrinsics.xor64 (sign x, sign y))-  let ! (x: u64) = unsign (intrinsics.complement64 (sign x))+  let not (x: u64) = unsign (intrinsics.complement64 (sign x))    let (x: u64) << (y: u64) = unsign (intrinsics.shl64 (sign x, sign y))   let (x: u64) >> (y: u64) = unsign (intrinsics.ashr64 (sign x, sign y))@@ -792,6 +804,7 @@   let i32 (x: i32) = unsign (intrinsics.zext_i32_i64 x)   let i64 (x: i64) = unsign (intrinsics.zext_i64_i64 x) +  let f16 (x: f16) = unsign (intrinsics.fptoui_f16_i64 x)   let f32 (x: f32) = unsign (intrinsics.fptoui_f32_i64 x)   let f64 (x: f64) = unsign (intrinsics.fptoui_f64_i64 x) @@ -805,7 +818,7 @@   let (x: u64) > (y: u64) = intrinsics.ult64 (sign y, sign x)   let (x: u64) <= (y: u64) = intrinsics.ule64 (sign x, sign y)   let (x: u64) >= (y: u64) = intrinsics.ule64 (sign y, sign x)-  let (x: u64) != (y: u64) = intrinsics.! (x == y)+  let (x: u64) != (y: u64) = !(x == y)    let sgn (x: u64) = unsign (intrinsics.usignum64 (sign x))   let abs (x: u64) = x@@ -820,7 +833,7 @@   let num_bits = 64i32   let get_bit (bit: i32) (x: t) = to_i32 ((x >> i32 bit) & i32 1)   let set_bit (bit: i32) (x: t) (b: i32) =-    ((x & i32 (intrinsics.!(1 intrinsics.<< bit))) | i32 (b intrinsics.<< bit))+    ((x & i32 (!(1 intrinsics.<< bit))) | i32 (b intrinsics.<< bit))   let popc x = intrinsics.popc64 (sign x)   let mul_hi a b = unsign (intrinsics.mul_hi64 (sign a, sign b))   let mad_hi a b c = unsign (intrinsics.mad_hi64 (sign a, sign b, sign c))@@ -857,6 +870,7 @@   let i32 (x: i32) = intrinsics.sitofp_i32_f64 x   let i64 (x: i64) = intrinsics.sitofp_i64_f64 x +  let f16 (x: f16) = intrinsics.fpconv_f16_f64 x   let f32 (x: f32) = intrinsics.fpconv_f32_f64 x   let f64 (x: f64) = intrinsics.fpconv_f64_f64 x @@ -871,7 +885,7 @@   let (x: f64) > (y: f64) = intrinsics.lt64 (y, x)   let (x: f64) <= (y: f64) = intrinsics.le64 (x, y)   let (x: f64) >= (y: f64) = intrinsics.le64 (y, x)-  let (x: f64) != (y: f64) = intrinsics.! (x == y)+  let (x: f64) != (y: f64) = !(x == y)    let neg (x: t) = -x   let recip (x: t) = 1/x@@ -965,6 +979,7 @@   let i32 (x: i32) = intrinsics.sitofp_i32_f32 x   let i64 (x: i64) = intrinsics.sitofp_i64_f32 x +  let f16 (x: f16) = intrinsics.fpconv_f16_f32 x   let f32 (x: f32) = intrinsics.fpconv_f32_f32 x   let f64 (x: f64) = intrinsics.fpconv_f64_f32 x @@ -979,7 +994,7 @@   let (x: f32) > (y: f32) = intrinsics.lt32 (y, x)   let (x: f32) <= (y: f32) = intrinsics.le32 (x, y)   let (x: f32) >= (y: f32) = intrinsics.le32 (y, x)-  let (x: f32) != (y: f32) = intrinsics.! (x == y)+  let (x: f32) != (y: f32) = !(x == y)    let neg (x: t) = -x   let recip (x: t) = 1/x@@ -1038,6 +1053,119 @@   let highest = inf   let lowest = -inf   let epsilon = 1.1920929e-7f32++  let pi = f64 f64m.pi+  let e = f64 f64m.e++  let sum = reduce (+) (i32 0)+  let product = reduce (*) (i32 1)+  let maximum = reduce max lowest+  let minimum = reduce min highest+}++-- | Emulated with single precision on systems that do not natively+-- support half precision.  This means you might get more accurate+-- results than on real systems, but it is also likely to be+-- significantly slower than just using `f32` in the first place.+module f16: (float with t = f16 with int_t = u16) = {+  type t = f16+  type int_t = u16++  module i16m = i16+  module u16m = u16+  module f64m = f64++  let (x: f16) + (y: f16) = intrinsics.fadd16 (x, y)+  let (x: f16) - (y: f16) = intrinsics.fsub16 (x, y)+  let (x: f16) * (y: f16) = intrinsics.fmul16 (x, y)+  let (x: f16) / (y: f16) = intrinsics.fdiv16 (x, y)+  let (x: f16) % (y: f16) = intrinsics.fmod16 (x, y)+  let (x: f16) ** (y: f16) = intrinsics.fpow16 (x, y)++  let u8  (x: u8)  = intrinsics.uitofp_i8_f16  (i8.u8 x)+  let u16 (x: u16) = intrinsics.uitofp_i16_f16 (i16.u16 x)+  let u32 (x: u32) = intrinsics.uitofp_i32_f16 (i32.u32 x)+  let u64 (x: u64) = intrinsics.uitofp_i64_f16 (i64.u64 x)++  let i8 (x: i8) = intrinsics.sitofp_i8_f16 x+  let i16 (x: i16) = intrinsics.sitofp_i16_f16 x+  let i32 (x: i32) = intrinsics.sitofp_i32_f16 x+  let i64 (x: i64) = intrinsics.sitofp_i64_f16 x++  let f16 (x: f16) = intrinsics.fpconv_f16_f16 x+  let f32 (x: f32) = intrinsics.fpconv_f32_f16 x+  let f64 (x: f64) = intrinsics.fpconv_f64_f16 x++  let bool (x: bool) = if x then 1f16 else 0f16++  let from_fraction (x: i64) (y: i64) = i64 x / i64 y+  let to_i64 (x: f16) = intrinsics.fptosi_f16_i64 x+  let to_f64 (x: f16) = intrinsics.fpconv_f16_f64 x++  let (x: f16) == (y: f16) = intrinsics.eq_f16 (x, y)+  let (x: f16) < (y: f16) = intrinsics.lt16 (x, y)+  let (x: f16) > (y: f16) = intrinsics.lt16 (y, x)+  let (x: f16) <= (y: f16) = intrinsics.le16 (x, y)+  let (x: f16) >= (y: f16) = intrinsics.le16 (y, x)+  let (x: f16) != (y: f16) = !(x == y)++  let neg (x: t) = -x+  let recip (x: t) = 1/x+  let max (x: t) (y: t) = intrinsics.fmax16 (x, y)+  let min (x: t) (y: t) = intrinsics.fmin16 (x, y)++  let sgn (x: f16) = intrinsics.fsignum16 x+  let abs (x: f16) = intrinsics.fabs16 x++  let sqrt (x: f16) = intrinsics.sqrt16 x++  let log (x: f16) = intrinsics.log16 x+  let log2 (x: f16) = intrinsics.log2_16 x+  let log10 (x: f16) = intrinsics.log10_16 x+  let exp (x: f16) = intrinsics.exp16 x+  let sin (x: f16) = intrinsics.sin16 x+  let cos (x: f16) = intrinsics.cos16 x+  let tan (x: f16) = intrinsics.tan16 x+  let acos (x: f16) = intrinsics.acos16 x+  let asin (x: f16) = intrinsics.asin16 x+  let atan (x: f16) = intrinsics.atan16 x+  let sinh (x: f16) = intrinsics.sinh16 x+  let cosh (x: f16) = intrinsics.cosh16 x+  let tanh (x: f16) = intrinsics.tanh16 x+  let acosh (x: f16) = intrinsics.acosh16 x+  let asinh (x: f16) = intrinsics.asinh16 x+  let atanh (x: f16) = intrinsics.atanh16 x+  let atan2 (x: f16) (y: f16) = intrinsics.atan2_16 (x, y)+  let hypot (x: f16) (y: f16) = intrinsics.hypot16 (x, y)+  let gamma = intrinsics.gamma16+  let lgamma = intrinsics.lgamma16++  let lerp v0 v1 t = intrinsics.lerp16 (v0,v1,t)+  let fma a b c = intrinsics.fma16 (a,b,c)+  let mad a b c = intrinsics.mad16 (a,b,c)++  let ceil = intrinsics.ceil16+  let floor = intrinsics.floor16+  let trunc (x: f16) : f16 = i16 (i16m.f16 x)++  let round = intrinsics.round16++  let to_bits (x: f16): u16 = u16m.i16 (intrinsics.to_bits16 x)+  let from_bits (x: u16): f16 = intrinsics.from_bits16 (intrinsics.sign_i16 x)++  let num_bits = 16i32+  let get_bit (bit: i32) (x: t) = u16m.get_bit bit (to_bits x)+  let set_bit (bit: i32) (x: t) (b: i32) = from_bits (u16m.set_bit bit (to_bits x) b)++  let isinf (x: f16) = intrinsics.isinf16 x+  let isnan (x: f16) = intrinsics.isnan16 x++  let inf = 1f16 / 0f16+  let nan = 0f16 / 0f16++  let highest = inf+  let lowest = -inf+  let epsilon = 1.1920929e-7f16    let pi = f64 f64m.pi   let e = f64 f64m.e

@@ -20,15 +20,16 @@ -- inhibitor.  The compiler will treat this function as a black box. -- You can use this to work around optimisation deficiencies (or -- bugs), although it should hopefully rarely be necessary.+-- Deprecated: use `#[opaque]` attribute instead. let opaque 't (x: t): t =-  intrinsics.opaque x+  #[opaque] x --- | Semantically just identity, but when run in the interpreter, the--- argument value will be printed.+-- | Semantically just identity, but at runtime, the argument value+-- will be printed.  Deprecated: use `#[trace]` attribute instead. let trace 't (x: t): t =-  intrinsics.trace x+  #[trace(trace)] x  -- | Semantically just identity, but acts as a break point in--- `futhark repl`.+-- `futhark repl`.  Deprecated: use `#[break]` attribute instead. let break 't (x: t): t =-  intrinsics.break x+  #[break] x

@@ -345,6 +345,18 @@   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");+   // It is crucial that the extra_opts are last, so that the free()   // logic below does not cause problems.   for (int j = 0; extra_opts[j] != NULL; j++) {

@@ -0,0 +1,238 @@+// Start of half.h.++// Conversion functions are from http://half.sourceforge.net/, but+// translated to C.+//+// Copyright (c) 2012-2021 Christian Rau+//+// Permission is hereby granted, free of charge, to any person obtaining a copy+// of this software and associated documentation files (the "Software"), to deal+// in the Software without restriction, including without limitation the rights+// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell+// copies of the Software, and to permit persons to whom the Software is+// furnished to do so, subject to the following conditions:+//+// The above copyright notice and this permission notice shall be included in+// all copies or substantial portions of the Software.+//+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR+// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,+// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE+// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER+// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,+// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN+// THE SOFTWARE.++#ifndef __OPENCL_VERSION__+#define __constant+#endif++__constant static const uint16_t base_table[512] = {+  0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000,+  0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000,+  0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000,+  0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000,+  0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000,+  0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000,+  0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0001, 0x0002, 0x0004, 0x0008, 0x0010, 0x0020, 0x0040, 0x0080, 0x0100,+  0x0200, 0x0400, 0x0800, 0x0C00, 0x1000, 0x1400, 0x1800, 0x1C00, 0x2000, 0x2400, 0x2800, 0x2C00, 0x3000, 0x3400, 0x3800, 0x3C00,+  0x4000, 0x4400, 0x4800, 0x4C00, 0x5000, 0x5400, 0x5800, 0x5C00, 0x6000, 0x6400, 0x6800, 0x6C00, 0x7000, 0x7400, 0x7800, 0x7C00,+  0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00,+  0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00,+  0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00,+  0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00,+  0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00,+  0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00,+  0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00, 0x7C00,+  0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000,+  0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000,+  0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000,+  0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000,+  0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000,+  0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000,+  0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8001, 0x8002, 0x8004, 0x8008, 0x8010, 0x8020, 0x8040, 0x8080, 0x8100,+  0x8200, 0x8400, 0x8800, 0x8C00, 0x9000, 0x9400, 0x9800, 0x9C00, 0xA000, 0xA400, 0xA800, 0xAC00, 0xB000, 0xB400, 0xB800, 0xBC00,+  0xC000, 0xC400, 0xC800, 0xCC00, 0xD000, 0xD400, 0xD800, 0xDC00, 0xE000, 0xE400, 0xE800, 0xEC00, 0xF000, 0xF400, 0xF800, 0xFC00,+  0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00,+  0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00,+  0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00,+  0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00,+  0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00,+  0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00,+  0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00, 0xFC00 };++__constant static const unsigned char shift_table[512] = {+  24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24,+  24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24,+  24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24,+  24, 24, 24, 24, 24, 24, 24, 23, 22, 21, 20, 19, 18, 17, 16, 15, 14, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13,+  13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24,+  24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24,+  24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24,+  24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 13,+  24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24,+  24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24,+  24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24,+  24, 24, 24, 24, 24, 24, 24, 23, 22, 21, 20, 19, 18, 17, 16, 15, 14, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13,+  13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24,+  24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24,+  24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24,+  24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 13 };++__constant static const uint32_t mantissa_table[2048] = {+  0x00000000, 0x33800000, 0x34000000, 0x34400000, 0x34800000, 0x34A00000, 0x34C00000, 0x34E00000, 0x35000000, 0x35100000, 0x35200000, 0x35300000, 0x35400000, 0x35500000, 0x35600000, 0x35700000,+  0x35800000, 0x35880000, 0x35900000, 0x35980000, 0x35A00000, 0x35A80000, 0x35B00000, 0x35B80000, 0x35C00000, 0x35C80000, 0x35D00000, 0x35D80000, 0x35E00000, 0x35E80000, 0x35F00000, 0x35F80000,+  0x36000000, 0x36040000, 0x36080000, 0x360C0000, 0x36100000, 0x36140000, 0x36180000, 0x361C0000, 0x36200000, 0x36240000, 0x36280000, 0x362C0000, 0x36300000, 0x36340000, 0x36380000, 0x363C0000,+  0x36400000, 0x36440000, 0x36480000, 0x364C0000, 0x36500000, 0x36540000, 0x36580000, 0x365C0000, 0x36600000, 0x36640000, 0x36680000, 0x366C0000, 0x36700000, 0x36740000, 0x36780000, 0x367C0000,+  0x36800000, 0x36820000, 0x36840000, 0x36860000, 0x36880000, 0x368A0000, 0x368C0000, 0x368E0000, 0x36900000, 0x36920000, 0x36940000, 0x36960000, 0x36980000, 0x369A0000, 0x369C0000, 0x369E0000,+  0x36A00000, 0x36A20000, 0x36A40000, 0x36A60000, 0x36A80000, 0x36AA0000, 0x36AC0000, 0x36AE0000, 0x36B00000, 0x36B20000, 0x36B40000, 0x36B60000, 0x36B80000, 0x36BA0000, 0x36BC0000, 0x36BE0000,+  0x36C00000, 0x36C20000, 0x36C40000, 0x36C60000, 0x36C80000, 0x36CA0000, 0x36CC0000, 0x36CE0000, 0x36D00000, 0x36D20000, 0x36D40000, 0x36D60000, 0x36D80000, 0x36DA0000, 0x36DC0000, 0x36DE0000,+  0x36E00000, 0x36E20000, 0x36E40000, 0x36E60000, 0x36E80000, 0x36EA0000, 0x36EC0000, 0x36EE0000, 0x36F00000, 0x36F20000, 0x36F40000, 0x36F60000, 0x36F80000, 0x36FA0000, 0x36FC0000, 0x36FE0000,+  0x37000000, 0x37010000, 0x37020000, 0x37030000, 0x37040000, 0x37050000, 0x37060000, 0x37070000, 0x37080000, 0x37090000, 0x370A0000, 0x370B0000, 0x370C0000, 0x370D0000, 0x370E0000, 0x370F0000,+  0x37100000, 0x37110000, 0x37120000, 0x37130000, 0x37140000, 0x37150000, 0x37160000, 0x37170000, 0x37180000, 0x37190000, 0x371A0000, 0x371B0000, 0x371C0000, 0x371D0000, 0x371E0000, 0x371F0000,+  0x37200000, 0x37210000, 0x37220000, 0x37230000, 0x37240000, 0x37250000, 0x37260000, 0x37270000, 0x37280000, 0x37290000, 0x372A0000, 0x372B0000, 0x372C0000, 0x372D0000, 0x372E0000, 0x372F0000,+  0x37300000, 0x37310000, 0x37320000, 0x37330000, 0x37340000, 0x37350000, 0x37360000, 0x37370000, 0x37380000, 0x37390000, 0x373A0000, 0x373B0000, 0x373C0000, 0x373D0000, 0x373E0000, 0x373F0000,+  0x37400000, 0x37410000, 0x37420000, 0x37430000, 0x37440000, 0x37450000, 0x37460000, 0x37470000, 0x37480000, 0x37490000, 0x374A0000, 0x374B0000, 0x374C0000, 0x374D0000, 0x374E0000, 0x374F0000,+  0x37500000, 0x37510000, 0x37520000, 0x37530000, 0x37540000, 0x37550000, 0x37560000, 0x37570000, 0x37580000, 0x37590000, 0x375A0000, 0x375B0000, 0x375C0000, 0x375D0000, 0x375E0000, 0x375F0000,+  0x37600000, 0x37610000, 0x37620000, 0x37630000, 0x37640000, 0x37650000, 0x37660000, 0x37670000, 0x37680000, 0x37690000, 0x376A0000, 0x376B0000, 0x376C0000, 0x376D0000, 0x376E0000, 0x376F0000,+  0x37700000, 0x37710000, 0x37720000, 0x37730000, 0x37740000, 0x37750000, 0x37760000, 0x37770000, 0x37780000, 0x37790000, 0x377A0000, 0x377B0000, 0x377C0000, 0x377D0000, 0x377E0000, 0x377F0000,+  0x37800000, 0x37808000, 0x37810000, 0x37818000, 0x37820000, 0x37828000, 0x37830000, 0x37838000, 0x37840000, 0x37848000, 0x37850000, 0x37858000, 0x37860000, 0x37868000, 0x37870000, 0x37878000,+  0x37880000, 0x37888000, 0x37890000, 0x37898000, 0x378A0000, 0x378A8000, 0x378B0000, 0x378B8000, 0x378C0000, 0x378C8000, 0x378D0000, 0x378D8000, 0x378E0000, 0x378E8000, 0x378F0000, 0x378F8000,+  0x37900000, 0x37908000, 0x37910000, 0x37918000, 0x37920000, 0x37928000, 0x37930000, 0x37938000, 0x37940000, 0x37948000, 0x37950000, 0x37958000, 0x37960000, 0x37968000, 0x37970000, 0x37978000,+  0x37980000, 0x37988000, 0x37990000, 0x37998000, 0x379A0000, 0x379A8000, 0x379B0000, 0x379B8000, 0x379C0000, 0x379C8000, 0x379D0000, 0x379D8000, 0x379E0000, 0x379E8000, 0x379F0000, 0x379F8000,+  0x37A00000, 0x37A08000, 0x37A10000, 0x37A18000, 0x37A20000, 0x37A28000, 0x37A30000, 0x37A38000, 0x37A40000, 0x37A48000, 0x37A50000, 0x37A58000, 0x37A60000, 0x37A68000, 0x37A70000, 0x37A78000,+  0x37A80000, 0x37A88000, 0x37A90000, 0x37A98000, 0x37AA0000, 0x37AA8000, 0x37AB0000, 0x37AB8000, 0x37AC0000, 0x37AC8000, 0x37AD0000, 0x37AD8000, 0x37AE0000, 0x37AE8000, 0x37AF0000, 0x37AF8000,+  0x37B00000, 0x37B08000, 0x37B10000, 0x37B18000, 0x37B20000, 0x37B28000, 0x37B30000, 0x37B38000, 0x37B40000, 0x37B48000, 0x37B50000, 0x37B58000, 0x37B60000, 0x37B68000, 0x37B70000, 0x37B78000,+  0x37B80000, 0x37B88000, 0x37B90000, 0x37B98000, 0x37BA0000, 0x37BA8000, 0x37BB0000, 0x37BB8000, 0x37BC0000, 0x37BC8000, 0x37BD0000, 0x37BD8000, 0x37BE0000, 0x37BE8000, 0x37BF0000, 0x37BF8000,+  0x37C00000, 0x37C08000, 0x37C10000, 0x37C18000, 0x37C20000, 0x37C28000, 0x37C30000, 0x37C38000, 0x37C40000, 0x37C48000, 0x37C50000, 0x37C58000, 0x37C60000, 0x37C68000, 0x37C70000, 0x37C78000,+  0x37C80000, 0x37C88000, 0x37C90000, 0x37C98000, 0x37CA0000, 0x37CA8000, 0x37CB0000, 0x37CB8000, 0x37CC0000, 0x37CC8000, 0x37CD0000, 0x37CD8000, 0x37CE0000, 0x37CE8000, 0x37CF0000, 0x37CF8000,+  0x37D00000, 0x37D08000, 0x37D10000, 0x37D18000, 0x37D20000, 0x37D28000, 0x37D30000, 0x37D38000, 0x37D40000, 0x37D48000, 0x37D50000, 0x37D58000, 0x37D60000, 0x37D68000, 0x37D70000, 0x37D78000,+  0x37D80000, 0x37D88000, 0x37D90000, 0x37D98000, 0x37DA0000, 0x37DA8000, 0x37DB0000, 0x37DB8000, 0x37DC0000, 0x37DC8000, 0x37DD0000, 0x37DD8000, 0x37DE0000, 0x37DE8000, 0x37DF0000, 0x37DF8000,+  0x37E00000, 0x37E08000, 0x37E10000, 0x37E18000, 0x37E20000, 0x37E28000, 0x37E30000, 0x37E38000, 0x37E40000, 0x37E48000, 0x37E50000, 0x37E58000, 0x37E60000, 0x37E68000, 0x37E70000, 0x37E78000,+  0x37E80000, 0x37E88000, 0x37E90000, 0x37E98000, 0x37EA0000, 0x37EA8000, 0x37EB0000, 0x37EB8000, 0x37EC0000, 0x37EC8000, 0x37ED0000, 0x37ED8000, 0x37EE0000, 0x37EE8000, 0x37EF0000, 0x37EF8000,+  0x37F00000, 0x37F08000, 0x37F10000, 0x37F18000, 0x37F20000, 0x37F28000, 0x37F30000, 0x37F38000, 0x37F40000, 0x37F48000, 0x37F50000, 0x37F58000, 0x37F60000, 0x37F68000, 0x37F70000, 0x37F78000,+  0x37F80000, 0x37F88000, 0x37F90000, 0x37F98000, 0x37FA0000, 0x37FA8000, 0x37FB0000, 0x37FB8000, 0x37FC0000, 0x37FC8000, 0x37FD0000, 0x37FD8000, 0x37FE0000, 0x37FE8000, 0x37FF0000, 0x37FF8000,+  0x38000000, 0x38004000, 0x38008000, 0x3800C000, 0x38010000, 0x38014000, 0x38018000, 0x3801C000, 0x38020000, 0x38024000, 0x38028000, 0x3802C000, 0x38030000, 0x38034000, 0x38038000, 0x3803C000,+  0x38040000, 0x38044000, 0x38048000, 0x3804C000, 0x38050000, 0x38054000, 0x38058000, 0x3805C000, 0x38060000, 0x38064000, 0x38068000, 0x3806C000, 0x38070000, 0x38074000, 0x38078000, 0x3807C000,+  0x38080000, 0x38084000, 0x38088000, 0x3808C000, 0x38090000, 0x38094000, 0x38098000, 0x3809C000, 0x380A0000, 0x380A4000, 0x380A8000, 0x380AC000, 0x380B0000, 0x380B4000, 0x380B8000, 0x380BC000,+  0x380C0000, 0x380C4000, 0x380C8000, 0x380CC000, 0x380D0000, 0x380D4000, 0x380D8000, 0x380DC000, 0x380E0000, 0x380E4000, 0x380E8000, 0x380EC000, 0x380F0000, 0x380F4000, 0x380F8000, 0x380FC000,+  0x38100000, 0x38104000, 0x38108000, 0x3810C000, 0x38110000, 0x38114000, 0x38118000, 0x3811C000, 0x38120000, 0x38124000, 0x38128000, 0x3812C000, 0x38130000, 0x38134000, 0x38138000, 0x3813C000,+  0x38140000, 0x38144000, 0x38148000, 0x3814C000, 0x38150000, 0x38154000, 0x38158000, 0x3815C000, 0x38160000, 0x38164000, 0x38168000, 0x3816C000, 0x38170000, 0x38174000, 0x38178000, 0x3817C000,+  0x38180000, 0x38184000, 0x38188000, 0x3818C000, 0x38190000, 0x38194000, 0x38198000, 0x3819C000, 0x381A0000, 0x381A4000, 0x381A8000, 0x381AC000, 0x381B0000, 0x381B4000, 0x381B8000, 0x381BC000,+  0x381C0000, 0x381C4000, 0x381C8000, 0x381CC000, 0x381D0000, 0x381D4000, 0x381D8000, 0x381DC000, 0x381E0000, 0x381E4000, 0x381E8000, 0x381EC000, 0x381F0000, 0x381F4000, 0x381F8000, 0x381FC000,+  0x38200000, 0x38204000, 0x38208000, 0x3820C000, 0x38210000, 0x38214000, 0x38218000, 0x3821C000, 0x38220000, 0x38224000, 0x38228000, 0x3822C000, 0x38230000, 0x38234000, 0x38238000, 0x3823C000,+  0x38240000, 0x38244000, 0x38248000, 0x3824C000, 0x38250000, 0x38254000, 0x38258000, 0x3825C000, 0x38260000, 0x38264000, 0x38268000, 0x3826C000, 0x38270000, 0x38274000, 0x38278000, 0x3827C000,+  0x38280000, 0x38284000, 0x38288000, 0x3828C000, 0x38290000, 0x38294000, 0x38298000, 0x3829C000, 0x382A0000, 0x382A4000, 0x382A8000, 0x382AC000, 0x382B0000, 0x382B4000, 0x382B8000, 0x382BC000,+  0x382C0000, 0x382C4000, 0x382C8000, 0x382CC000, 0x382D0000, 0x382D4000, 0x382D8000, 0x382DC000, 0x382E0000, 0x382E4000, 0x382E8000, 0x382EC000, 0x382F0000, 0x382F4000, 0x382F8000, 0x382FC000,+  0x38300000, 0x38304000, 0x38308000, 0x3830C000, 0x38310000, 0x38314000, 0x38318000, 0x3831C000, 0x38320000, 0x38324000, 0x38328000, 0x3832C000, 0x38330000, 0x38334000, 0x38338000, 0x3833C000,+  0x38340000, 0x38344000, 0x38348000, 0x3834C000, 0x38350000, 0x38354000, 0x38358000, 0x3835C000, 0x38360000, 0x38364000, 0x38368000, 0x3836C000, 0x38370000, 0x38374000, 0x38378000, 0x3837C000,+  0x38380000, 0x38384000, 0x38388000, 0x3838C000, 0x38390000, 0x38394000, 0x38398000, 0x3839C000, 0x383A0000, 0x383A4000, 0x383A8000, 0x383AC000, 0x383B0000, 0x383B4000, 0x383B8000, 0x383BC000,+  0x383C0000, 0x383C4000, 0x383C8000, 0x383CC000, 0x383D0000, 0x383D4000, 0x383D8000, 0x383DC000, 0x383E0000, 0x383E4000, 0x383E8000, 0x383EC000, 0x383F0000, 0x383F4000, 0x383F8000, 0x383FC000,+  0x38400000, 0x38404000, 0x38408000, 0x3840C000, 0x38410000, 0x38414000, 0x38418000, 0x3841C000, 0x38420000, 0x38424000, 0x38428000, 0x3842C000, 0x38430000, 0x38434000, 0x38438000, 0x3843C000,+  0x38440000, 0x38444000, 0x38448000, 0x3844C000, 0x38450000, 0x38454000, 0x38458000, 0x3845C000, 0x38460000, 0x38464000, 0x38468000, 0x3846C000, 0x38470000, 0x38474000, 0x38478000, 0x3847C000,+  0x38480000, 0x38484000, 0x38488000, 0x3848C000, 0x38490000, 0x38494000, 0x38498000, 0x3849C000, 0x384A0000, 0x384A4000, 0x384A8000, 0x384AC000, 0x384B0000, 0x384B4000, 0x384B8000, 0x384BC000,+  0x384C0000, 0x384C4000, 0x384C8000, 0x384CC000, 0x384D0000, 0x384D4000, 0x384D8000, 0x384DC000, 0x384E0000, 0x384E4000, 0x384E8000, 0x384EC000, 0x384F0000, 0x384F4000, 0x384F8000, 0x384FC000,+  0x38500000, 0x38504000, 0x38508000, 0x3850C000, 0x38510000, 0x38514000, 0x38518000, 0x3851C000, 0x38520000, 0x38524000, 0x38528000, 0x3852C000, 0x38530000, 0x38534000, 0x38538000, 0x3853C000,+  0x38540000, 0x38544000, 0x38548000, 0x3854C000, 0x38550000, 0x38554000, 0x38558000, 0x3855C000, 0x38560000, 0x38564000, 0x38568000, 0x3856C000, 0x38570000, 0x38574000, 0x38578000, 0x3857C000,+  0x38580000, 0x38584000, 0x38588000, 0x3858C000, 0x38590000, 0x38594000, 0x38598000, 0x3859C000, 0x385A0000, 0x385A4000, 0x385A8000, 0x385AC000, 0x385B0000, 0x385B4000, 0x385B8000, 0x385BC000,+  0x385C0000, 0x385C4000, 0x385C8000, 0x385CC000, 0x385D0000, 0x385D4000, 0x385D8000, 0x385DC000, 0x385E0000, 0x385E4000, 0x385E8000, 0x385EC000, 0x385F0000, 0x385F4000, 0x385F8000, 0x385FC000,+  0x38600000, 0x38604000, 0x38608000, 0x3860C000, 0x38610000, 0x38614000, 0x38618000, 0x3861C000, 0x38620000, 0x38624000, 0x38628000, 0x3862C000, 0x38630000, 0x38634000, 0x38638000, 0x3863C000,+  0x38640000, 0x38644000, 0x38648000, 0x3864C000, 0x38650000, 0x38654000, 0x38658000, 0x3865C000, 0x38660000, 0x38664000, 0x38668000, 0x3866C000, 0x38670000, 0x38674000, 0x38678000, 0x3867C000,+  0x38680000, 0x38684000, 0x38688000, 0x3868C000, 0x38690000, 0x38694000, 0x38698000, 0x3869C000, 0x386A0000, 0x386A4000, 0x386A8000, 0x386AC000, 0x386B0000, 0x386B4000, 0x386B8000, 0x386BC000,+  0x386C0000, 0x386C4000, 0x386C8000, 0x386CC000, 0x386D0000, 0x386D4000, 0x386D8000, 0x386DC000, 0x386E0000, 0x386E4000, 0x386E8000, 0x386EC000, 0x386F0000, 0x386F4000, 0x386F8000, 0x386FC000,+  0x38700000, 0x38704000, 0x38708000, 0x3870C000, 0x38710000, 0x38714000, 0x38718000, 0x3871C000, 0x38720000, 0x38724000, 0x38728000, 0x3872C000, 0x38730000, 0x38734000, 0x38738000, 0x3873C000,+  0x38740000, 0x38744000, 0x38748000, 0x3874C000, 0x38750000, 0x38754000, 0x38758000, 0x3875C000, 0x38760000, 0x38764000, 0x38768000, 0x3876C000, 0x38770000, 0x38774000, 0x38778000, 0x3877C000,+  0x38780000, 0x38784000, 0x38788000, 0x3878C000, 0x38790000, 0x38794000, 0x38798000, 0x3879C000, 0x387A0000, 0x387A4000, 0x387A8000, 0x387AC000, 0x387B0000, 0x387B4000, 0x387B8000, 0x387BC000,+  0x387C0000, 0x387C4000, 0x387C8000, 0x387CC000, 0x387D0000, 0x387D4000, 0x387D8000, 0x387DC000, 0x387E0000, 0x387E4000, 0x387E8000, 0x387EC000, 0x387F0000, 0x387F4000, 0x387F8000, 0x387FC000,+  0x38000000, 0x38002000, 0x38004000, 0x38006000, 0x38008000, 0x3800A000, 0x3800C000, 0x3800E000, 0x38010000, 0x38012000, 0x38014000, 0x38016000, 0x38018000, 0x3801A000, 0x3801C000, 0x3801E000,+  0x38020000, 0x38022000, 0x38024000, 0x38026000, 0x38028000, 0x3802A000, 0x3802C000, 0x3802E000, 0x38030000, 0x38032000, 0x38034000, 0x38036000, 0x38038000, 0x3803A000, 0x3803C000, 0x3803E000,+  0x38040000, 0x38042000, 0x38044000, 0x38046000, 0x38048000, 0x3804A000, 0x3804C000, 0x3804E000, 0x38050000, 0x38052000, 0x38054000, 0x38056000, 0x38058000, 0x3805A000, 0x3805C000, 0x3805E000,+  0x38060000, 0x38062000, 0x38064000, 0x38066000, 0x38068000, 0x3806A000, 0x3806C000, 0x3806E000, 0x38070000, 0x38072000, 0x38074000, 0x38076000, 0x38078000, 0x3807A000, 0x3807C000, 0x3807E000,+  0x38080000, 0x38082000, 0x38084000, 0x38086000, 0x38088000, 0x3808A000, 0x3808C000, 0x3808E000, 0x38090000, 0x38092000, 0x38094000, 0x38096000, 0x38098000, 0x3809A000, 0x3809C000, 0x3809E000,+  0x380A0000, 0x380A2000, 0x380A4000, 0x380A6000, 0x380A8000, 0x380AA000, 0x380AC000, 0x380AE000, 0x380B0000, 0x380B2000, 0x380B4000, 0x380B6000, 0x380B8000, 0x380BA000, 0x380BC000, 0x380BE000,+  0x380C0000, 0x380C2000, 0x380C4000, 0x380C6000, 0x380C8000, 0x380CA000, 0x380CC000, 0x380CE000, 0x380D0000, 0x380D2000, 0x380D4000, 0x380D6000, 0x380D8000, 0x380DA000, 0x380DC000, 0x380DE000,+  0x380E0000, 0x380E2000, 0x380E4000, 0x380E6000, 0x380E8000, 0x380EA000, 0x380EC000, 0x380EE000, 0x380F0000, 0x380F2000, 0x380F4000, 0x380F6000, 0x380F8000, 0x380FA000, 0x380FC000, 0x380FE000,+  0x38100000, 0x38102000, 0x38104000, 0x38106000, 0x38108000, 0x3810A000, 0x3810C000, 0x3810E000, 0x38110000, 0x38112000, 0x38114000, 0x38116000, 0x38118000, 0x3811A000, 0x3811C000, 0x3811E000,+  0x38120000, 0x38122000, 0x38124000, 0x38126000, 0x38128000, 0x3812A000, 0x3812C000, 0x3812E000, 0x38130000, 0x38132000, 0x38134000, 0x38136000, 0x38138000, 0x3813A000, 0x3813C000, 0x3813E000,+  0x38140000, 0x38142000, 0x38144000, 0x38146000, 0x38148000, 0x3814A000, 0x3814C000, 0x3814E000, 0x38150000, 0x38152000, 0x38154000, 0x38156000, 0x38158000, 0x3815A000, 0x3815C000, 0x3815E000,+  0x38160000, 0x38162000, 0x38164000, 0x38166000, 0x38168000, 0x3816A000, 0x3816C000, 0x3816E000, 0x38170000, 0x38172000, 0x38174000, 0x38176000, 0x38178000, 0x3817A000, 0x3817C000, 0x3817E000,+  0x38180000, 0x38182000, 0x38184000, 0x38186000, 0x38188000, 0x3818A000, 0x3818C000, 0x3818E000, 0x38190000, 0x38192000, 0x38194000, 0x38196000, 0x38198000, 0x3819A000, 0x3819C000, 0x3819E000,+  0x381A0000, 0x381A2000, 0x381A4000, 0x381A6000, 0x381A8000, 0x381AA000, 0x381AC000, 0x381AE000, 0x381B0000, 0x381B2000, 0x381B4000, 0x381B6000, 0x381B8000, 0x381BA000, 0x381BC000, 0x381BE000,+  0x381C0000, 0x381C2000, 0x381C4000, 0x381C6000, 0x381C8000, 0x381CA000, 0x381CC000, 0x381CE000, 0x381D0000, 0x381D2000, 0x381D4000, 0x381D6000, 0x381D8000, 0x381DA000, 0x381DC000, 0x381DE000,+  0x381E0000, 0x381E2000, 0x381E4000, 0x381E6000, 0x381E8000, 0x381EA000, 0x381EC000, 0x381EE000, 0x381F0000, 0x381F2000, 0x381F4000, 0x381F6000, 0x381F8000, 0x381FA000, 0x381FC000, 0x381FE000,+  0x38200000, 0x38202000, 0x38204000, 0x38206000, 0x38208000, 0x3820A000, 0x3820C000, 0x3820E000, 0x38210000, 0x38212000, 0x38214000, 0x38216000, 0x38218000, 0x3821A000, 0x3821C000, 0x3821E000,+  0x38220000, 0x38222000, 0x38224000, 0x38226000, 0x38228000, 0x3822A000, 0x3822C000, 0x3822E000, 0x38230000, 0x38232000, 0x38234000, 0x38236000, 0x38238000, 0x3823A000, 0x3823C000, 0x3823E000,+  0x38240000, 0x38242000, 0x38244000, 0x38246000, 0x38248000, 0x3824A000, 0x3824C000, 0x3824E000, 0x38250000, 0x38252000, 0x38254000, 0x38256000, 0x38258000, 0x3825A000, 0x3825C000, 0x3825E000,+  0x38260000, 0x38262000, 0x38264000, 0x38266000, 0x38268000, 0x3826A000, 0x3826C000, 0x3826E000, 0x38270000, 0x38272000, 0x38274000, 0x38276000, 0x38278000, 0x3827A000, 0x3827C000, 0x3827E000,+  0x38280000, 0x38282000, 0x38284000, 0x38286000, 0x38288000, 0x3828A000, 0x3828C000, 0x3828E000, 0x38290000, 0x38292000, 0x38294000, 0x38296000, 0x38298000, 0x3829A000, 0x3829C000, 0x3829E000,+  0x382A0000, 0x382A2000, 0x382A4000, 0x382A6000, 0x382A8000, 0x382AA000, 0x382AC000, 0x382AE000, 0x382B0000, 0x382B2000, 0x382B4000, 0x382B6000, 0x382B8000, 0x382BA000, 0x382BC000, 0x382BE000,+  0x382C0000, 0x382C2000, 0x382C4000, 0x382C6000, 0x382C8000, 0x382CA000, 0x382CC000, 0x382CE000, 0x382D0000, 0x382D2000, 0x382D4000, 0x382D6000, 0x382D8000, 0x382DA000, 0x382DC000, 0x382DE000,+  0x382E0000, 0x382E2000, 0x382E4000, 0x382E6000, 0x382E8000, 0x382EA000, 0x382EC000, 0x382EE000, 0x382F0000, 0x382F2000, 0x382F4000, 0x382F6000, 0x382F8000, 0x382FA000, 0x382FC000, 0x382FE000,+  0x38300000, 0x38302000, 0x38304000, 0x38306000, 0x38308000, 0x3830A000, 0x3830C000, 0x3830E000, 0x38310000, 0x38312000, 0x38314000, 0x38316000, 0x38318000, 0x3831A000, 0x3831C000, 0x3831E000,+  0x38320000, 0x38322000, 0x38324000, 0x38326000, 0x38328000, 0x3832A000, 0x3832C000, 0x3832E000, 0x38330000, 0x38332000, 0x38334000, 0x38336000, 0x38338000, 0x3833A000, 0x3833C000, 0x3833E000,+  0x38340000, 0x38342000, 0x38344000, 0x38346000, 0x38348000, 0x3834A000, 0x3834C000, 0x3834E000, 0x38350000, 0x38352000, 0x38354000, 0x38356000, 0x38358000, 0x3835A000, 0x3835C000, 0x3835E000,+  0x38360000, 0x38362000, 0x38364000, 0x38366000, 0x38368000, 0x3836A000, 0x3836C000, 0x3836E000, 0x38370000, 0x38372000, 0x38374000, 0x38376000, 0x38378000, 0x3837A000, 0x3837C000, 0x3837E000,+  0x38380000, 0x38382000, 0x38384000, 0x38386000, 0x38388000, 0x3838A000, 0x3838C000, 0x3838E000, 0x38390000, 0x38392000, 0x38394000, 0x38396000, 0x38398000, 0x3839A000, 0x3839C000, 0x3839E000,+  0x383A0000, 0x383A2000, 0x383A4000, 0x383A6000, 0x383A8000, 0x383AA000, 0x383AC000, 0x383AE000, 0x383B0000, 0x383B2000, 0x383B4000, 0x383B6000, 0x383B8000, 0x383BA000, 0x383BC000, 0x383BE000,+  0x383C0000, 0x383C2000, 0x383C4000, 0x383C6000, 0x383C8000, 0x383CA000, 0x383CC000, 0x383CE000, 0x383D0000, 0x383D2000, 0x383D4000, 0x383D6000, 0x383D8000, 0x383DA000, 0x383DC000, 0x383DE000,+  0x383E0000, 0x383E2000, 0x383E4000, 0x383E6000, 0x383E8000, 0x383EA000, 0x383EC000, 0x383EE000, 0x383F0000, 0x383F2000, 0x383F4000, 0x383F6000, 0x383F8000, 0x383FA000, 0x383FC000, 0x383FE000,+  0x38400000, 0x38402000, 0x38404000, 0x38406000, 0x38408000, 0x3840A000, 0x3840C000, 0x3840E000, 0x38410000, 0x38412000, 0x38414000, 0x38416000, 0x38418000, 0x3841A000, 0x3841C000, 0x3841E000,+  0x38420000, 0x38422000, 0x38424000, 0x38426000, 0x38428000, 0x3842A000, 0x3842C000, 0x3842E000, 0x38430000, 0x38432000, 0x38434000, 0x38436000, 0x38438000, 0x3843A000, 0x3843C000, 0x3843E000,+  0x38440000, 0x38442000, 0x38444000, 0x38446000, 0x38448000, 0x3844A000, 0x3844C000, 0x3844E000, 0x38450000, 0x38452000, 0x38454000, 0x38456000, 0x38458000, 0x3845A000, 0x3845C000, 0x3845E000,+  0x38460000, 0x38462000, 0x38464000, 0x38466000, 0x38468000, 0x3846A000, 0x3846C000, 0x3846E000, 0x38470000, 0x38472000, 0x38474000, 0x38476000, 0x38478000, 0x3847A000, 0x3847C000, 0x3847E000,+  0x38480000, 0x38482000, 0x38484000, 0x38486000, 0x38488000, 0x3848A000, 0x3848C000, 0x3848E000, 0x38490000, 0x38492000, 0x38494000, 0x38496000, 0x38498000, 0x3849A000, 0x3849C000, 0x3849E000,+  0x384A0000, 0x384A2000, 0x384A4000, 0x384A6000, 0x384A8000, 0x384AA000, 0x384AC000, 0x384AE000, 0x384B0000, 0x384B2000, 0x384B4000, 0x384B6000, 0x384B8000, 0x384BA000, 0x384BC000, 0x384BE000,+  0x384C0000, 0x384C2000, 0x384C4000, 0x384C6000, 0x384C8000, 0x384CA000, 0x384CC000, 0x384CE000, 0x384D0000, 0x384D2000, 0x384D4000, 0x384D6000, 0x384D8000, 0x384DA000, 0x384DC000, 0x384DE000,+  0x384E0000, 0x384E2000, 0x384E4000, 0x384E6000, 0x384E8000, 0x384EA000, 0x384EC000, 0x384EE000, 0x384F0000, 0x384F2000, 0x384F4000, 0x384F6000, 0x384F8000, 0x384FA000, 0x384FC000, 0x384FE000,+  0x38500000, 0x38502000, 0x38504000, 0x38506000, 0x38508000, 0x3850A000, 0x3850C000, 0x3850E000, 0x38510000, 0x38512000, 0x38514000, 0x38516000, 0x38518000, 0x3851A000, 0x3851C000, 0x3851E000,+  0x38520000, 0x38522000, 0x38524000, 0x38526000, 0x38528000, 0x3852A000, 0x3852C000, 0x3852E000, 0x38530000, 0x38532000, 0x38534000, 0x38536000, 0x38538000, 0x3853A000, 0x3853C000, 0x3853E000,+  0x38540000, 0x38542000, 0x38544000, 0x38546000, 0x38548000, 0x3854A000, 0x3854C000, 0x3854E000, 0x38550000, 0x38552000, 0x38554000, 0x38556000, 0x38558000, 0x3855A000, 0x3855C000, 0x3855E000,+  0x38560000, 0x38562000, 0x38564000, 0x38566000, 0x38568000, 0x3856A000, 0x3856C000, 0x3856E000, 0x38570000, 0x38572000, 0x38574000, 0x38576000, 0x38578000, 0x3857A000, 0x3857C000, 0x3857E000,+  0x38580000, 0x38582000, 0x38584000, 0x38586000, 0x38588000, 0x3858A000, 0x3858C000, 0x3858E000, 0x38590000, 0x38592000, 0x38594000, 0x38596000, 0x38598000, 0x3859A000, 0x3859C000, 0x3859E000,+  0x385A0000, 0x385A2000, 0x385A4000, 0x385A6000, 0x385A8000, 0x385AA000, 0x385AC000, 0x385AE000, 0x385B0000, 0x385B2000, 0x385B4000, 0x385B6000, 0x385B8000, 0x385BA000, 0x385BC000, 0x385BE000,+  0x385C0000, 0x385C2000, 0x385C4000, 0x385C6000, 0x385C8000, 0x385CA000, 0x385CC000, 0x385CE000, 0x385D0000, 0x385D2000, 0x385D4000, 0x385D6000, 0x385D8000, 0x385DA000, 0x385DC000, 0x385DE000,+  0x385E0000, 0x385E2000, 0x385E4000, 0x385E6000, 0x385E8000, 0x385EA000, 0x385EC000, 0x385EE000, 0x385F0000, 0x385F2000, 0x385F4000, 0x385F6000, 0x385F8000, 0x385FA000, 0x385FC000, 0x385FE000,+  0x38600000, 0x38602000, 0x38604000, 0x38606000, 0x38608000, 0x3860A000, 0x3860C000, 0x3860E000, 0x38610000, 0x38612000, 0x38614000, 0x38616000, 0x38618000, 0x3861A000, 0x3861C000, 0x3861E000,+  0x38620000, 0x38622000, 0x38624000, 0x38626000, 0x38628000, 0x3862A000, 0x3862C000, 0x3862E000, 0x38630000, 0x38632000, 0x38634000, 0x38636000, 0x38638000, 0x3863A000, 0x3863C000, 0x3863E000,+  0x38640000, 0x38642000, 0x38644000, 0x38646000, 0x38648000, 0x3864A000, 0x3864C000, 0x3864E000, 0x38650000, 0x38652000, 0x38654000, 0x38656000, 0x38658000, 0x3865A000, 0x3865C000, 0x3865E000,+  0x38660000, 0x38662000, 0x38664000, 0x38666000, 0x38668000, 0x3866A000, 0x3866C000, 0x3866E000, 0x38670000, 0x38672000, 0x38674000, 0x38676000, 0x38678000, 0x3867A000, 0x3867C000, 0x3867E000,+  0x38680000, 0x38682000, 0x38684000, 0x38686000, 0x38688000, 0x3868A000, 0x3868C000, 0x3868E000, 0x38690000, 0x38692000, 0x38694000, 0x38696000, 0x38698000, 0x3869A000, 0x3869C000, 0x3869E000,+  0x386A0000, 0x386A2000, 0x386A4000, 0x386A6000, 0x386A8000, 0x386AA000, 0x386AC000, 0x386AE000, 0x386B0000, 0x386B2000, 0x386B4000, 0x386B6000, 0x386B8000, 0x386BA000, 0x386BC000, 0x386BE000,+  0x386C0000, 0x386C2000, 0x386C4000, 0x386C6000, 0x386C8000, 0x386CA000, 0x386CC000, 0x386CE000, 0x386D0000, 0x386D2000, 0x386D4000, 0x386D6000, 0x386D8000, 0x386DA000, 0x386DC000, 0x386DE000,+  0x386E0000, 0x386E2000, 0x386E4000, 0x386E6000, 0x386E8000, 0x386EA000, 0x386EC000, 0x386EE000, 0x386F0000, 0x386F2000, 0x386F4000, 0x386F6000, 0x386F8000, 0x386FA000, 0x386FC000, 0x386FE000,+  0x38700000, 0x38702000, 0x38704000, 0x38706000, 0x38708000, 0x3870A000, 0x3870C000, 0x3870E000, 0x38710000, 0x38712000, 0x38714000, 0x38716000, 0x38718000, 0x3871A000, 0x3871C000, 0x3871E000,+  0x38720000, 0x38722000, 0x38724000, 0x38726000, 0x38728000, 0x3872A000, 0x3872C000, 0x3872E000, 0x38730000, 0x38732000, 0x38734000, 0x38736000, 0x38738000, 0x3873A000, 0x3873C000, 0x3873E000,+  0x38740000, 0x38742000, 0x38744000, 0x38746000, 0x38748000, 0x3874A000, 0x3874C000, 0x3874E000, 0x38750000, 0x38752000, 0x38754000, 0x38756000, 0x38758000, 0x3875A000, 0x3875C000, 0x3875E000,+  0x38760000, 0x38762000, 0x38764000, 0x38766000, 0x38768000, 0x3876A000, 0x3876C000, 0x3876E000, 0x38770000, 0x38772000, 0x38774000, 0x38776000, 0x38778000, 0x3877A000, 0x3877C000, 0x3877E000,+  0x38780000, 0x38782000, 0x38784000, 0x38786000, 0x38788000, 0x3878A000, 0x3878C000, 0x3878E000, 0x38790000, 0x38792000, 0x38794000, 0x38796000, 0x38798000, 0x3879A000, 0x3879C000, 0x3879E000,+  0x387A0000, 0x387A2000, 0x387A4000, 0x387A6000, 0x387A8000, 0x387AA000, 0x387AC000, 0x387AE000, 0x387B0000, 0x387B2000, 0x387B4000, 0x387B6000, 0x387B8000, 0x387BA000, 0x387BC000, 0x387BE000,+  0x387C0000, 0x387C2000, 0x387C4000, 0x387C6000, 0x387C8000, 0x387CA000, 0x387CC000, 0x387CE000, 0x387D0000, 0x387D2000, 0x387D4000, 0x387D6000, 0x387D8000, 0x387DA000, 0x387DC000, 0x387DE000,+  0x387E0000, 0x387E2000, 0x387E4000, 0x387E6000, 0x387E8000, 0x387EA000, 0x387EC000, 0x387EE000, 0x387F0000, 0x387F2000, 0x387F4000, 0x387F6000, 0x387F8000, 0x387FA000, 0x387FC000, 0x387FE000 };+__constant static const uint32_t exponent_table[64] = {+  0x00000000, 0x00800000, 0x01000000, 0x01800000, 0x02000000, 0x02800000, 0x03000000, 0x03800000, 0x04000000, 0x04800000, 0x05000000, 0x05800000, 0x06000000, 0x06800000, 0x07000000, 0x07800000,+  0x08000000, 0x08800000, 0x09000000, 0x09800000, 0x0A000000, 0x0A800000, 0x0B000000, 0x0B800000, 0x0C000000, 0x0C800000, 0x0D000000, 0x0D800000, 0x0E000000, 0x0E800000, 0x0F000000, 0x47800000,+  0x80000000, 0x80800000, 0x81000000, 0x81800000, 0x82000000, 0x82800000, 0x83000000, 0x83800000, 0x84000000, 0x84800000, 0x85000000, 0x85800000, 0x86000000, 0x86800000, 0x87000000, 0x87800000,+  0x88000000, 0x88800000, 0x89000000, 0x89800000, 0x8A000000, 0x8A800000, 0x8B000000, 0x8B800000, 0x8C000000, 0x8C800000, 0x8D000000, 0x8D800000, 0x8E000000, 0x8E800000, 0x8F000000, 0xC7800000 };+__constant static const unsigned short offset_table[64] = {+  0, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024,+  0, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024 };++static uint16_t float2halfbits(float value) {+  union { float x; uint32_t y; } u;+  u.x = value;+  uint32_t bits = u.y;++  uint16_t hbits = base_table[bits>>23] + (uint16_t)((bits&0x7FFFFF)>>shift_table[bits>>23]);;++  return hbits;+}++static float halfbits2float(uint16_t value) {+  uint32_t bits = mantissa_table[offset_table[value>>10]+(value&0x3FF)] + exponent_table[value>>10];++  union { uint32_t x; float y; } u;+  u.x = bits;+  return u.y;+}++// End of half.h.

@@ -195,7 +195,7 @@     } } -static void opencl_succeed_fatal(unsigned int ret,+static void opencl_succeed_fatal(cl_int ret,                                  const char *call,                                  const char *file,                                  int line) {@@ -205,7 +205,7 @@   } } -static char* opencl_succeed_nonfatal(unsigned int ret,+static char* opencl_succeed_nonfatal(cl_int ret,                                      const char *call,                                      const char *file,                                      int line) {@@ -734,6 +734,12 @@   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 ");   }    if (ctx->cfg.debugging) {

@@ -0,0 +1,1654 @@+// Start of scalar.h.++// Implementation of the primitive scalar operations.  Very+// repetitive.  This code is inserted directly into both CUDA and+// OpenCL programs, as well as the CPU code, so it has some #ifdefs to+// work everywhere.  Some operations are defined as macros because+// this allows us to use them as constant expressions in things like+// array sizes and static initialisers.++// Some of the #ifdefs are because OpenCL uses type-generic functions+// for some operations (e.g. sqrt), while C and CUDA sensibly use+// distinct functions for different precisions (e.g. sqrtf() and+// sqrt()).  This is quite annoying.  Due to C's unfortunate casting+// rules, it is also really easy to accidentally implement+// floating-point functions in the wrong precision, so be careful.++// Double-precision definitions are only included if the preprocessor+// macro FUTHARK_F64_ENABLED is set.++static inline uint8_t add8(uint8_t x, uint8_t y) {+  return x + y;+}++static inline uint16_t add16(uint16_t x, uint16_t y) {+  return x + y;+}++static inline uint32_t add32(uint32_t x, uint32_t y) {+  return x + y;+}++static inline uint64_t add64(uint64_t x, uint64_t y) {+  return x + y;+}++static inline uint8_t sub8(uint8_t x, uint8_t y) {+  return x - y;+}++static inline uint16_t sub16(uint16_t x, uint16_t y) {+  return x - y;+}++static inline uint32_t sub32(uint32_t x, uint32_t y) {+  return x - y;+}++static inline uint64_t sub64(uint64_t x, uint64_t y) {+  return x - y;+}++static inline uint8_t mul8(uint8_t x, uint8_t y) {+  return x * y;+}++static inline uint16_t mul16(uint16_t x, uint16_t y) {+  return x * y;+}++static inline uint32_t mul32(uint32_t x, uint32_t y) {+  return x * y;+}++static inline uint64_t mul64(uint64_t x, uint64_t y) {+  return x * y;+}++static inline uint8_t udiv8(uint8_t x, uint8_t y) {+  return x / y;+}++static inline uint16_t udiv16(uint16_t x, uint16_t y) {+  return x / y;+}++static inline uint32_t udiv32(uint32_t x, uint32_t y) {+  return x / y;+}++static inline uint64_t udiv64(uint64_t x, uint64_t y) {+  return x / y;+}++static inline uint8_t udiv_up8(uint8_t x, uint8_t y) {+  return (x + y - 1) / y;+}++static inline uint16_t udiv_up16(uint16_t x, uint16_t y) {+  return (x + y - 1) / y;+}++static inline uint32_t udiv_up32(uint32_t x, uint32_t y) {+  return (x + y - 1) / y;+}++static inline uint64_t udiv_up64(uint64_t x, uint64_t y) {+  return (x + y - 1) / y;+}++static inline uint8_t umod8(uint8_t x, uint8_t y) {+  return x % y;+}++static inline uint16_t umod16(uint16_t x, uint16_t y) {+  return x % y;+}++static inline uint32_t umod32(uint32_t x, uint32_t y) {+  return x % y;+}++static inline uint64_t umod64(uint64_t x, uint64_t y) {+  return x % y;+}++static inline uint8_t udiv_safe8(uint8_t x, uint8_t y) {+  return y == 0 ? 0 : x / y;+}++static inline uint16_t udiv_safe16(uint16_t x, uint16_t y) {+  return y == 0 ? 0 : x / y;+}++static inline uint32_t udiv_safe32(uint32_t x, uint32_t y) {+  return y == 0 ? 0 : x / y;+}++static inline uint64_t udiv_safe64(uint64_t x, uint64_t y) {+  return y == 0 ? 0 : x / y;+}++static inline uint8_t udiv_up_safe8(uint8_t x, uint8_t y) {+  return y == 0 ? 0 : (x + y - 1) / y;+}++static inline uint16_t udiv_up_safe16(uint16_t x, uint16_t y) {+  return y == 0 ? 0 : (x + y - 1) / y;+}++static inline uint32_t udiv_up_safe32(uint32_t x, uint32_t y) {+  return y == 0 ? 0 : (x + y - 1) / y;+}++static inline uint64_t udiv_up_safe64(uint64_t x, uint64_t y) {+  return y == 0 ? 0 : (x + y - 1) / y;+}++static inline uint8_t umod_safe8(uint8_t x, uint8_t y) {+  return y == 0 ? 0 : x % y;+}++static inline uint16_t umod_safe16(uint16_t x, uint16_t y) {+  return y == 0 ? 0 : x % y;+}++static inline uint32_t umod_safe32(uint32_t x, uint32_t y) {+  return y == 0 ? 0 : x % y;+}++static inline uint64_t umod_safe64(uint64_t x, uint64_t y) {+  return y == 0 ? 0 : x % y;+}++static inline int8_t sdiv8(int8_t x, int8_t y) {+  int8_t q = x / y;+  int8_t r = x % y;++  return q - ((r != 0 && r < 0 != y < 0) ? 1 : 0);+}++static inline int16_t sdiv16(int16_t x, int16_t y) {+  int16_t q = x / y;+  int16_t r = x % y;++  return q - ((r != 0 && r < 0 != y < 0) ? 1 : 0);+}++static inline int32_t sdiv32(int32_t x, int32_t y) {+  int32_t q = x / y;+  int32_t r = x % y;++  return q - ((r != 0 && r < 0 != y < 0) ? 1 : 0);+}++static inline int64_t sdiv64(int64_t x, int64_t y) {+  int64_t q = x / y;+  int64_t r = x % y;++  return q - ((r != 0 && r < 0 != y < 0) ? 1 : 0);+}++static inline int8_t sdiv_up8(int8_t x, int8_t y) {+  return sdiv8(x + y - 1, y);+}++static inline int16_t sdiv_up16(int16_t x, int16_t y) {+  return sdiv16(x + y - 1, y);+}++static inline int32_t sdiv_up32(int32_t x, int32_t y) {+  return sdiv32(x + y - 1, y);+}++static inline int64_t sdiv_up64(int64_t x, int64_t y) {+  return sdiv64(x + y - 1, y);+}++static inline int8_t smod8(int8_t x, int8_t y) {+  int8_t r = x % y;++  return r + (r == 0 || (x > 0 && y > 0) || (x < 0 && y < 0) ? 0 : y);+}++static inline int16_t smod16(int16_t x, int16_t y) {+  int16_t r = x % y;++  return r + (r == 0 || (x > 0 && y > 0) || (x < 0 && y < 0) ? 0 : y);+}++static inline int32_t smod32(int32_t x, int32_t y) {+  int32_t r = x % y;++  return r + (r == 0 || (x > 0 && y > 0) || (x < 0 && y < 0) ? 0 : y);+}++static inline int64_t smod64(int64_t x, int64_t y) {+  int64_t r = x % y;++  return r + (r == 0 || (x > 0 && y > 0) || (x < 0 && y < 0) ? 0 : y);+}++static inline int8_t sdiv_safe8(int8_t x, int8_t y) {+  return y == 0 ? 0 : sdiv8(x, y);+}++static inline int16_t sdiv_safe16(int16_t x, int16_t y) {+  return y == 0 ? 0 : sdiv16(x, y);+}++static inline int32_t sdiv_safe32(int32_t x, int32_t y) {+  return y == 0 ? 0 : sdiv32(x, y);+}++static inline int64_t sdiv_safe64(int64_t x, int64_t y) {+  return y == 0 ? 0 : sdiv64(x, y);+}++static inline int8_t sdiv_up_safe8(int8_t x, int8_t y) {+  return sdiv_safe8(x + y - 1, y);+}++static inline int16_t sdiv_up_safe16(int16_t x, int16_t y) {+  return sdiv_safe16(x + y - 1, y);+}++static inline int32_t sdiv_up_safe32(int32_t x, int32_t y) {+  return sdiv_safe32(x + y - 1, y);+}++static inline int64_t sdiv_up_safe64(int64_t x, int64_t y) {+  return sdiv_safe64(x + y - 1, y);+}++static inline int8_t smod_safe8(int8_t x, int8_t y) {+  return y == 0 ? 0 : smod8(x, y);+}++static inline int16_t smod_safe16(int16_t x, int16_t y) {+  return y == 0 ? 0 : smod16(x, y);+}++static inline int32_t smod_safe32(int32_t x, int32_t y) {+  return y == 0 ? 0 : smod32(x, y);+}++static inline int64_t smod_safe64(int64_t x, int64_t y) {+  return y == 0 ? 0 : smod64(x, y);+}++static inline int8_t squot8(int8_t x, int8_t y) {+  return x / y;+}++static inline int16_t squot16(int16_t x, int16_t y) {+  return x / y;+}++static inline int32_t squot32(int32_t x, int32_t y) {+  return x / y;+}++static inline int64_t squot64(int64_t x, int64_t y) {+  return x / y;+}++static inline int8_t srem8(int8_t x, int8_t y) {+  return x % y;+}++static inline int16_t srem16(int16_t x, int16_t y) {+  return x % y;+}++static inline int32_t srem32(int32_t x, int32_t y) {+  return x % y;+}++static inline int64_t srem64(int64_t x, int64_t y) {+  return x % y;+}++static inline int8_t squot_safe8(int8_t x, int8_t y) {+  return y == 0 ? 0 : x / y;+}++static inline int16_t squot_safe16(int16_t x, int16_t y) {+  return y == 0 ? 0 : x / y;+}++static inline int32_t squot_safe32(int32_t x, int32_t y) {+  return y == 0 ? 0 : x / y;+}++static inline int64_t squot_safe64(int64_t x, int64_t y) {+  return y == 0 ? 0 : x / y;+}++static inline int8_t srem_safe8(int8_t x, int8_t y) {+  return y == 0 ? 0 : x % y;+}++static inline int16_t srem_safe16(int16_t x, int16_t y) {+  return y == 0 ? 0 : x % y;+}++static inline int32_t srem_safe32(int32_t x, int32_t y) {+  return y == 0 ? 0 : x % y;+}++static inline int64_t srem_safe64(int64_t x, int64_t y) {+  return y == 0 ? 0 : x % y;+}++static inline int8_t smin8(int8_t x, int8_t y) {+  return x < y ? x : y;+}++static inline int16_t smin16(int16_t x, int16_t y) {+  return x < y ? x : y;+}++static inline int32_t smin32(int32_t x, int32_t y) {+  return x < y ? x : y;+}++static inline int64_t smin64(int64_t x, int64_t y) {+  return x < y ? x : y;+}++static inline uint8_t umin8(uint8_t x, uint8_t y) {+  return x < y ? x : y;+}++static inline uint16_t umin16(uint16_t x, uint16_t y) {+  return x < y ? x : y;+}++static inline uint32_t umin32(uint32_t x, uint32_t y) {+  return x < y ? x : y;+}++static inline uint64_t umin64(uint64_t x, uint64_t y) {+  return x < y ? x : y;+}++static inline int8_t smax8(int8_t x, int8_t y) {+  return x < y ? y : x;+}++static inline int16_t smax16(int16_t x, int16_t y) {+  return x < y ? y : x;+}++static inline int32_t smax32(int32_t x, int32_t y) {+  return x < y ? y : x;+}++static inline int64_t smax64(int64_t x, int64_t y) {+  return x < y ? y : x;+}++static inline uint8_t umax8(uint8_t x, uint8_t y) {+  return x < y ? y : x;+}++static inline uint16_t umax16(uint16_t x, uint16_t y) {+  return x < y ? y : x;+}++static inline uint32_t umax32(uint32_t x, uint32_t y) {+  return x < y ? y : x;+}++static inline uint64_t umax64(uint64_t x, uint64_t y) {+  return x < y ? y : x;+}++static inline uint8_t shl8(uint8_t x, uint8_t y) {+  return (uint8_t)(x << y);+}++static inline uint16_t shl16(uint16_t x, uint16_t y) {+  return (uint16_t)(x << y);+}++static inline uint32_t shl32(uint32_t x, uint32_t y) {+  return x << y;+}++static inline uint64_t shl64(uint64_t x, uint64_t y) {+  return x << y;+}++static inline uint8_t lshr8(uint8_t x, uint8_t y) {+  return x >> y;+}++static inline uint16_t lshr16(uint16_t x, uint16_t y) {+  return x >> y;+}++static inline uint32_t lshr32(uint32_t x, uint32_t y) {+  return x >> y;+}++static inline uint64_t lshr64(uint64_t x, uint64_t y) {+  return x >> y;+}++static inline int8_t ashr8(int8_t x, int8_t y) {+  return x >> y;+}++static inline int16_t ashr16(int16_t x, int16_t y) {+  return x >> y;+}++static inline int32_t ashr32(int32_t x, int32_t y) {+  return x >> y;+}++static inline int64_t ashr64(int64_t x, int64_t y) {+  return x >> y;+}++static inline uint8_t and8(uint8_t x, uint8_t y) {+  return x & y;+}++static inline uint16_t and16(uint16_t x, uint16_t y) {+  return x & y;+}++static inline uint32_t and32(uint32_t x, uint32_t y) {+  return x & y;+}++static inline uint64_t and64(uint64_t x, uint64_t y) {+  return x & y;+}++static inline uint8_t or8(uint8_t x, uint8_t y) {+  return x | y;+}++static inline uint16_t or16(uint16_t x, uint16_t y) {+  return x | y;+}++static inline uint32_t or32(uint32_t x, uint32_t y) {+  return x | y;+}++static inline uint64_t or64(uint64_t x, uint64_t y) {+  return x | y;+}++static inline uint8_t xor8(uint8_t x, uint8_t y) {+  return x ^ y;+}++static inline uint16_t xor16(uint16_t x, uint16_t y) {+  return x ^ y;+}++static inline uint32_t xor32(uint32_t x, uint32_t y) {+  return x ^ y;+}++static inline uint64_t xor64(uint64_t x, uint64_t y) {+  return x ^ y;+}++static inline bool ult8(uint8_t x, uint8_t y) {+  return x < y;+}++static inline bool ult16(uint16_t x, uint16_t y) {+  return x < y;+}++static inline bool ult32(uint32_t x, uint32_t y) {+  return x < y;+}++static inline bool ult64(uint64_t x, uint64_t y) {+  return x < y;+}++static inline bool ule8(uint8_t x, uint8_t y) {+  return x <= y;+}++static inline bool ule16(uint16_t x, uint16_t y) {+  return x <= y;+}++static inline bool ule32(uint32_t x, uint32_t y) {+  return x <= y;+}++static inline bool ule64(uint64_t x, uint64_t y) {+  return x <= y;+}++static inline bool slt8(int8_t x, int8_t y) {+  return x < y;+}++static inline bool slt16(int16_t x, int16_t y) {+  return x < y;+}++static inline bool slt32(int32_t x, int32_t y) {+  return x < y;+}++static inline bool slt64(int64_t x, int64_t y) {+  return x < y;+}++static inline bool sle8(int8_t x, int8_t y) {+  return x <= y;+}++static inline bool sle16(int16_t x, int16_t y) {+  return x <= y;+}++static inline bool sle32(int32_t x, int32_t y) {+  return x <= y;+}++static inline bool sle64(int64_t x, int64_t y) {+  return x <= y;+}++static inline uint8_t pow8(uint8_t x, uint8_t y) {+  uint8_t res = 1, rem = y;++  while (rem != 0) {+    if (rem & 1)+      res *= x;+    rem >>= 1;+    x *= x;+  }+  return res;+}++static inline uint16_t pow16(uint16_t x, uint16_t y) {+  uint16_t res = 1, rem = y;++  while (rem != 0) {+    if (rem & 1)+      res *= x;+    rem >>= 1;+    x *= x;+  }+  return res;+}++static inline uint32_t pow32(uint32_t x, uint32_t y) {+  uint32_t res = 1, rem = y;++  while (rem != 0) {+    if (rem & 1)+      res *= x;+    rem >>= 1;+    x *= x;+  }+  return res;+}++static inline uint64_t pow64(uint64_t x, uint64_t y) {+  uint64_t res = 1, rem = y;++  while (rem != 0) {+    if (rem & 1)+      res *= x;+    rem >>= 1;+    x *= x;+  }+  return res;+}++static inline bool itob_i8_bool(int8_t x) {+  return x;+}++static inline bool itob_i16_bool(int16_t x) {+  return x;+}++static inline bool itob_i32_bool(int32_t x) {+  return x;+}++static inline bool itob_i64_bool(int64_t x) {+  return x;+}++static inline int8_t btoi_bool_i8(bool x) {+  return x;+}++static inline int16_t btoi_bool_i16(bool x) {+  return x;+}++static inline int32_t btoi_bool_i32(bool x) {+  return x;+}++static inline int64_t btoi_bool_i64(bool x) {+  return x;+}++#define sext_i8_i8(x) ((int8_t) (int8_t) (x))+#define sext_i8_i16(x) ((int16_t) (int8_t) (x))+#define sext_i8_i32(x) ((int32_t) (int8_t) (x))+#define sext_i8_i64(x) ((int64_t) (int8_t) (x))+#define sext_i16_i8(x) ((int8_t) (int16_t) (x))+#define sext_i16_i16(x) ((int16_t) (int16_t) (x))+#define sext_i16_i32(x) ((int32_t) (int16_t) (x))+#define sext_i16_i64(x) ((int64_t) (int16_t) (x))+#define sext_i32_i8(x) ((int8_t) (int32_t) (x))+#define sext_i32_i16(x) ((int16_t) (int32_t) (x))+#define sext_i32_i32(x) ((int32_t) (int32_t) (x))+#define sext_i32_i64(x) ((int64_t) (int32_t) (x))+#define sext_i64_i8(x) ((int8_t) (int64_t) (x))+#define sext_i64_i16(x) ((int16_t) (int64_t) (x))+#define sext_i64_i32(x) ((int32_t) (int64_t) (x))+#define sext_i64_i64(x) ((int64_t) (int64_t) (x))+#define zext_i8_i8(x) ((int8_t) (uint8_t) (x))+#define zext_i8_i16(x) ((int16_t) (uint8_t) (x))+#define zext_i8_i32(x) ((int32_t) (uint8_t) (x))+#define zext_i8_i64(x) ((int64_t) (uint8_t) (x))+#define zext_i16_i8(x) ((int8_t) (uint16_t) (x))+#define zext_i16_i16(x) ((int16_t) (uint16_t) (x))+#define zext_i16_i32(x) ((int32_t) (uint16_t) (x))+#define zext_i16_i64(x) ((int64_t) (uint16_t) (x))+#define zext_i32_i8(x) ((int8_t) (uint32_t) (x))+#define zext_i32_i16(x) ((int16_t) (uint32_t) (x))+#define zext_i32_i32(x) ((int32_t) (uint32_t) (x))+#define zext_i32_i64(x) ((int64_t) (uint32_t) (x))+#define zext_i64_i8(x) ((int8_t) (uint64_t) (x))+#define zext_i64_i16(x) ((int16_t) (uint64_t) (x))+#define zext_i64_i32(x) ((int32_t) (uint64_t) (x))+#define zext_i64_i64(x) ((int64_t) (uint64_t) (x))++static int8_t abs8(int8_t x) {+  return (int8_t)abs(x);+}++static int16_t abs16(int16_t x) {+  return (int16_t)abs(x);+}++static int32_t abs32(int32_t x) {+  return abs(x);+}++static int64_t abs64(int64_t x) {+#if defined(__OPENCL_VERSION__)+  return abs(x);+#else+  return llabs(x);+#endif+}++#if defined(__OPENCL_VERSION__)+static int32_t futrts_popc8(int8_t x) {+  return popcount(x);+}++static int32_t futrts_popc16(int16_t x) {+  return popcount(x);+}++static int32_t futrts_popc32(int32_t x) {+  return popcount(x);+}++static int32_t futrts_popc64(int64_t x) {+  return popcount(x);+}+#elif defined(__CUDA_ARCH__)++static int32_t futrts_popc8(int8_t x) {+  return __popc(zext_i8_i32(x));+}++static int32_t futrts_popc16(int16_t x) {+  return __popc(zext_i16_i32(x));+}++static int32_t futrts_popc32(int32_t x) {+  return __popc(x);+}++static int32_t futrts_popc64(int64_t x) {+  return __popcll(x);+}++#else // Not OpenCL or CUDA, but plain C.++static int32_t futrts_popc8(uint8_t x) {+  int c = 0;+  for (; x; ++c) { x &= x - 1; }+  return c;+}++static int32_t futrts_popc16(uint16_t x) {+  int c = 0;+  for (; x; ++c) { x &= x - 1; }+  return c;+}++static int32_t futrts_popc32(uint32_t x) {+  int c = 0;+  for (; x; ++c) { x &= x - 1; }+  return c;+}++static int32_t futrts_popc64(uint64_t x) {+  int c = 0;+  for (; x; ++c) { x &= x - 1; }+  return c;+}+#endif++#if defined(__OPENCL_VERSION__)+static uint8_t futrts_mul_hi8(uint8_t a, uint8_t b) {+  return mul_hi(a, b);+}++static uint16_t futrts_mul_hi16(uint16_t a, uint16_t b) {+  return mul_hi(a, b);+}++static uint32_t futrts_mul_hi32(uint32_t a, uint32_t b) {+  return mul_hi(a, b);+}++static uint64_t futrts_mul_hi64(uint64_t a, uint64_t b) {+  return mul_hi(a, b);+}++#elif defined(__CUDA_ARCH__)++static uint8_t futrts_mul_hi8(uint8_t a, uint8_t b) {+  uint16_t aa = a;+  uint16_t bb = b;++  return aa * bb >> 8;+}++static uint16_t futrts_mul_hi16(uint16_t a, uint16_t b) {+  uint32_t aa = a;+  uint32_t bb = b;++  return aa * bb >> 16;+}++static uint32_t futrts_mul_hi32(uint32_t a, uint32_t b) {+  return mulhi(a, b);+}++static uint64_t futrts_mul_hi64(uint64_t a, uint64_t b) {+  return mul64hi(a, b);+}++#else // Not OpenCL or CUDA, but plain C.++static uint8_t futrts_mul_hi8(uint8_t a, uint8_t b) {+  uint16_t aa = a;+  uint16_t bb = b;++  return aa * bb >> 8;+}++static uint16_t futrts_mul_hi16(uint16_t a, uint16_t b) {+  uint32_t aa = a;+  uint32_t bb = b;++  return aa * bb >> 16;+}++static uint32_t futrts_mul_hi32(uint32_t a, uint32_t b) {+  uint64_t aa = a;+  uint64_t bb = b;++  return aa * bb >> 32;+}++static uint64_t futrts_mul_hi64(uint64_t a, uint64_t b) {+  __uint128_t aa = a;+  __uint128_t bb = b;++  return aa * bb >> 64;+}+#endif++#if defined(__OPENCL_VERSION__)+static uint8_t futrts_mad_hi8(uint8_t a, uint8_t b, uint8_t c) {+  return mad_hi(a, b, c);+}++static uint16_t futrts_mad_hi16(uint16_t a, uint16_t b, uint16_t c) {+  return mad_hi(a, b, c);+}++static uint32_t futrts_mad_hi32(uint32_t a, uint32_t b, uint32_t c) {+  return mad_hi(a, b, c);+}++static uint64_t futrts_mad_hi64(uint64_t a, uint64_t b, uint64_t c) {+  return mad_hi(a, b, c);+}++#else // Not OpenCL++static uint8_t futrts_mad_hi8(uint8_t a, uint8_t b, uint8_t c) {+  return futrts_mul_hi8(a, b) + c;+}++static uint16_t futrts_mad_hi16(uint16_t a, uint16_t b, uint16_t c) {+  return futrts_mul_hi16(a, b) + c;+}++static uint32_t futrts_mad_hi32(uint32_t a, uint32_t b, uint32_t c) {+  return futrts_mul_hi32(a, b) + c;+}++static uint64_t futrts_mad_hi64(uint64_t a, uint64_t b, uint64_t c) {+  return futrts_mul_hi64(a, b) + c;+}+#endif++#if defined(__OPENCL_VERSION__)+static int32_t futrts_clzz8(int8_t x) {+  return clz(x);+}++static int32_t futrts_clzz16(int16_t x) {+  return clz(x);+}++static int32_t futrts_clzz32(int32_t x) {+  return clz(x);+}++static int32_t futrts_clzz64(int64_t x) {+  return clz(x);+}++#elif defined(__CUDA_ARCH__)++static int32_t futrts_clzz8(int8_t x) {+  return __clz(zext_i8_i32(x)) - 24;+}++static int32_t futrts_clzz16(int16_t x) {+  return __clz(zext_i16_i32(x)) - 16;+}++static int32_t futrts_clzz32(int32_t x) {+  return __clz(x);+}++static int32_t futrts_clzz64(int64_t x) {+  return __clzll(x);+}++#else // Not OpenCL or CUDA, but plain C.++static int32_t futrts_clzz8(int8_t x) {+  return x == 0 ? 8 : __builtin_clz((uint32_t)zext_i8_i32(x)) - 24;+}++static int32_t futrts_clzz16(int16_t x) {+  return x == 0 ? 16 : __builtin_clz((uint32_t)zext_i16_i32(x)) - 16;+}++static int32_t futrts_clzz32(int32_t x) {+  return x == 0 ? 32 : __builtin_clz((uint32_t)x);+}++static int32_t futrts_clzz64(int64_t x) {+  return x == 0 ? 64 : __builtin_clzll((uint64_t)x);+}+#endif++#if defined(__OPENCL_VERSION__)+static int32_t futrts_ctzz8(int8_t x) {+  int i = 0;+  for (; i < 8 && (x & 1) == 0; i++, x >>= 1)+    ;+  return i;+}++static int32_t futrts_ctzz16(int16_t x) {+  int i = 0;+  for (; i < 16 && (x & 1) == 0; i++, x >>= 1)+    ;+  return i;+}++static int32_t futrts_ctzz32(int32_t x) {+  int i = 0;+  for (; i < 32 && (x & 1) == 0; i++, x >>= 1)+    ;+  return i;+}++static int32_t futrts_ctzz64(int64_t x) {+  int i = 0;+  for (; i < 64 && (x & 1) == 0; i++, x >>= 1)+    ;+  return i;+}++#elif defined(__CUDA_ARCH__)++static int32_t futrts_ctzz8(int8_t x) {+  int y = __ffs(x);+  return y == 0 ? 8 : y - 1;+}++static int32_t futrts_ctzz16(int16_t x) {+  int y = __ffs(x);+  return y == 0 ? 16 : y - 1;+}++static int32_t futrts_ctzz32(int32_t x) {+  int y = __ffs(x);+  return y == 0 ? 32 : y - 1;+}++static int32_t futrts_ctzz64(int64_t x) {+  int y = __ffsll(x);+  return y == 0 ? 64 : y - 1;+}++#else // Not OpenCL or CUDA, but plain C.++static int32_t futrts_ctzz8(int8_t x) {+  return x == 0 ? 8 : __builtin_ctz((uint32_t)x);+}++static int32_t futrts_ctzz16(int16_t x) {+  return x == 0 ? 16 : __builtin_ctz((uint32_t)x);+}++static int32_t futrts_ctzz32(int32_t x) {+  return x == 0 ? 32 : __builtin_ctz((uint32_t)x);+}++static int32_t futrts_ctzz64(int64_t x) {+  return x == 0 ? 64 : __builtin_ctzll((uint64_t)x);+}+#endif++static inline float fdiv32(float x, float y) {+  return x / y;+}++static inline float fadd32(float x, float y) {+  return x + y;+}++static inline float fsub32(float x, float y) {+  return x - y;+}++static inline float fmul32(float x, float y) {+  return x * y;+}++static inline bool cmplt32(float x, float y) {+  return x < y;+}++static inline bool cmple32(float x, float y) {+  return x <= y;+}++static inline float sitofp_i8_f32(int8_t x) {+  return (float) x;+}++static inline float sitofp_i16_f32(int16_t x) {+  return (float) x;+}++static inline float sitofp_i32_f32(int32_t x) {+  return (float) x;+}++static inline float sitofp_i64_f32(int64_t x) {+  return (float) x;+}++static inline float uitofp_i8_f32(uint8_t x) {+  return (float) x;+}++static inline float uitofp_i16_f32(uint16_t x) {+  return (float) x;+}++static inline float uitofp_i32_f32(uint32_t x) {+  return (float) x;+}++static inline float uitofp_i64_f32(uint64_t x) {+  return (float) x;+}++static inline int8_t fptosi_f32_i8(float x) {+  return (int8_t) x;+}++static inline int16_t fptosi_f32_i16(float x) {+  return (int16_t) x;+}++static inline int32_t fptosi_f32_i32(float x) {+  return (int32_t) x;+}++static inline int64_t fptosi_f32_i64(float x) {+  return (int64_t) x;+}++static inline uint8_t fptoui_f32_i8(float x) {+  return (uint8_t) x;+}++static inline uint16_t fptoui_f32_i16(float x) {+  return (uint16_t) x;+}++static inline uint32_t fptoui_f32_i32(float x) {+  return (uint32_t) x;+}++static inline uint64_t fptoui_f32_i64(float x) {+  return (uint64_t) x;+}++#ifdef __OPENCL_VERSION__+static inline float fabs32(float x) {+  return fabs(x);+}++static inline float fmax32(float x, float y) {+  return fmax(x, y);+}++static inline float fmin32(float x, float y) {+  return fmin(x, y);+}++static inline float fpow32(float x, float y) {+  return pow(x, y);+}++#else // Not OpenCL, but CUDA or plain C.++static inline float fabs32(float x) {+  return fabsf(x);+}++static inline float fmax32(float x, float y) {+  return fmaxf(x, y);+}++static inline float fmin32(float x, float y) {+  return fminf(x, y);+}++static inline float fpow32(float x, float y) {+  return powf(x, y);+}+#endif++static inline bool futrts_isnan32(float x) {+  return isnan(x);+}++static inline bool futrts_isinf32(float x) {+  return isinf(x);+}++#ifdef __OPENCL_VERSION__+static inline float futrts_log32(float x) {+  return log(x);+}++static inline float futrts_log2_32(float x) {+  return log2(x);+}++static inline float futrts_log10_32(float x) {+  return log10(x);+}++static inline float futrts_sqrt32(float x) {+  return sqrt(x);+}++static inline float futrts_exp32(float x) {+  return exp(x);+}++static inline float futrts_cos32(float x) {+  return cos(x);+}++static inline float futrts_sin32(float x) {+  return sin(x);+}++static inline float futrts_tan32(float x) {+  return tan(x);+}++static inline float futrts_acos32(float x) {+  return acos(x);+}++static inline float futrts_asin32(float x) {+  return asin(x);+}++static inline float futrts_atan32(float x) {+  return atan(x);+}++static inline float futrts_cosh32(float x) {+  return cosh(x);+}++static inline float futrts_sinh32(float x) {+  return sinh(x);+}++static inline float futrts_tanh32(float x) {+  return tanh(x);+}++static inline float futrts_acosh32(float x) {+  return acosh(x);+}++static inline float futrts_asinh32(float x) {+  return asinh(x);+}++static inline float futrts_atanh32(float x) {+  return atanh(x);+}++static inline float futrts_atan2_32(float x, float y) {+  return atan2(x, y);+}++static inline float futrts_hypot32(float x, float y) {+  return hypot(x, y);+}++static inline float futrts_gamma32(float x) {+  return tgamma(x);+}++static inline float futrts_lgamma32(float x) {+  return lgamma(x);+}++static inline float fmod32(float x, float y) {+  return fmod(x, y);+}++static inline float futrts_round32(float x) {+  return rint(x);+}++static inline float futrts_floor32(float x) {+  return floor(x);+}++static inline float futrts_ceil32(float x) {+  return ceil(x);+}++static inline float futrts_lerp32(float v0, float v1, float t) {+  return mix(v0, v1, t);+}++static inline float futrts_mad32(float a, float b, float c) {+  return mad(a, b, c);+}++static inline float futrts_fma32(float a, float b, float c) {+  return fma(a, b, c);+}++#else // Not OpenCL, but CUDA or plain C.++static inline float futrts_log32(float x) {+  return logf(x);+}++static inline float futrts_log2_32(float x) {+  return log2f(x);+}++static inline float futrts_log10_32(float x) {+  return log10f(x);+}++static inline float futrts_sqrt32(float x) {+  return sqrtf(x);+}++static inline float futrts_exp32(float x) {+  return expf(x);+}++static inline float futrts_cos32(float x) {+  return cosf(x);+}++static inline float futrts_sin32(float x) {+  return sinf(x);+}++static inline float futrts_tan32(float x) {+  return tanf(x);+}++static inline float futrts_acos32(float x) {+  return acosf(x);+}++static inline float futrts_asin32(float x) {+  return asinf(x);+}++static inline float futrts_atan32(float x) {+  return atanf(x);+}++static inline float futrts_cosh32(float x) {+  return coshf(x);+}++static inline float futrts_sinh32(float x) {+  return sinhf(x);+}++static inline float futrts_tanh32(float x) {+  return tanhf(x);+}++static inline float futrts_acosh32(float x) {+  return acoshf(x);+}++static inline float futrts_asinh32(float x) {+  return asinhf(x);+}++static inline float futrts_atanh32(float x) {+  return atanhf(x);+}++static inline float futrts_atan2_32(float x, float y) {+  return atan2f(x, y);+}++static inline float futrts_hypot32(float x, float y) {+  return hypotf(x, y);+}++static inline float futrts_gamma32(float x) {+  return tgammaf(x);+}++static inline float futrts_lgamma32(float x) {+  return lgammaf(x);+}++static inline float fmod32(float x, float y) {+  return fmodf(x, y);+}++static inline float futrts_round32(float x) {+  return rintf(x);+}++static inline float futrts_floor32(float x) {+  return floorf(x);+}++static inline float futrts_ceil32(float x) {+  return ceilf(x);+}++static inline float futrts_lerp32(float v0, float v1, float t) {+  return v0 + (v1 - v0) * t;+}++static inline float futrts_mad32(float a, float b, float c) {+  return a * b + c;+}++static inline float futrts_fma32(float a, float b, float c) {+  return fmaf(a, b, c);+}+#endif++static inline int32_t futrts_to_bits32(float x) {+  union {+    float f;+    int32_t t;+  } p;++  p.f = x;+  return p.t;+}++static inline float futrts_from_bits32(int32_t x) {+  union {+    int32_t f;+    float t;+  } p;++  p.f = x;+  return p.t;+}++static inline float fsignum32(float x) {+  return futrts_isnan32(x) ? x : (x > 0) - (x < 0);+}++#ifdef FUTHARK_F64_ENABLED++static inline double fdiv64(double x, double y) {+  return x / y;+}++static inline double fadd64(double x, double y) {+  return x + y;+}++static inline double fsub64(double x, double y) {+  return x - y;+}++static inline double fmul64(double x, double y) {+  return x * y;+}++static inline bool cmplt64(double x, double y) {+  return x < y;+}++static inline bool cmple64(double x, double y) {+  return x <= y;+}++static inline double sitofp_i8_f64(int8_t x) {+  return (double) x;+}++static inline double sitofp_i16_f64(int16_t x) {+  return (double) x;+}++static inline double sitofp_i32_f64(int32_t x) {+  return (double) x;+}++static inline double sitofp_i64_f64(int64_t x) {+  return (double) x;+}++static inline double uitofp_i8_f64(uint8_t x) {+  return (double) x;+}++static inline double uitofp_i16_f64(uint16_t x) {+  return (double) x;+}++static inline double uitofp_i32_f64(uint32_t x) {+  return (double) x;+}++static inline double uitofp_i64_f64(uint64_t x) {+  return (double) x;+}++static inline int8_t fptosi_f64_i8(double x) {+  return (int8_t) x;+}++static inline int16_t fptosi_f64_i16(double x) {+  return (int16_t) x;+}++static inline int32_t fptosi_f64_i32(double x) {+  return (int32_t) x;+}++static inline int64_t fptosi_f64_i64(double x) {+  return (int64_t) x;+}++static inline uint8_t fptoui_f64_i8(double x) {+  return (uint8_t) x;+}++static inline uint16_t fptoui_f64_i16(double x) {+  return (uint16_t) x;+}++static inline uint32_t fptoui_f64_i32(double x) {+  return (uint32_t) x;+}++static inline uint64_t fptoui_f64_i64(double x) {+  return (uint64_t) x;+}++static inline double fabs64(double x) {+  return fabs(x);+}++static inline double fmax64(double x, double y) {+  return fmax(x, y);+}++static inline double fmin64(double x, double y) {+  return fmin(x, y);+}++static inline double fpow64(double x, double y) {+  return pow(x, y);+}++static inline double futrts_log64(double x) {+  return log(x);+}++static inline double futrts_log2_64(double x) {+  return log2(x);+}++static inline double futrts_log10_64(double x) {+  return log10(x);+}++static inline double futrts_sqrt64(double x) {+  return sqrt(x);+}++static inline double futrts_exp64(double x) {+  return exp(x);+}++static inline double futrts_cos64(double x) {+  return cos(x);+}++static inline double futrts_sin64(double x) {+  return sin(x);+}++static inline double futrts_tan64(double x) {+  return tan(x);+}++static inline double futrts_acos64(double x) {+  return acos(x);+}++static inline double futrts_asin64(double x) {+  return asin(x);+}++static inline double futrts_atan64(double x) {+  return atan(x);+}++static inline double futrts_cosh64(double x) {+  return cosh(x);+}++static inline double futrts_sinh64(double x) {+  return sinh(x);+}++static inline double futrts_tanh64(double x) {+  return tanh(x);+}++static inline double futrts_acosh64(double x) {+  return acosh(x);+}++static inline double futrts_asinh64(double x) {+  return asinh(x);+}++static inline double futrts_atanh64(double x) {+  return atanh(x);+}++static inline double futrts_atan2_64(double x, double y) {+  return atan2(x, y);+}++static inline double futrts_hypot64(double x, double y) {+  return hypot(x, y);+}++static inline double futrts_gamma64(double x) {+  return tgamma(x);+}++static inline double futrts_lgamma64(double x) {+  return lgamma(x);+}++static inline double futrts_fma64(double a, double b, double c) {+  return fma(a, b, c);+}++static inline double futrts_round64(double x) {+  return rint(x);+}++static inline double futrts_ceil64(double x) {+  return ceil(x);+}++static inline double futrts_floor64(double x) {+  return floor(x);+}++static inline bool futrts_isnan64(double x) {+  return isnan(x);+}++static inline bool futrts_isinf64(double x) {+  return isinf(x);+}++static inline int64_t futrts_to_bits64(double x) {+  union {+    double f;+    int64_t t;+  } p;++  p.f = x;+  return p.t;+}++static inline double futrts_from_bits64(int64_t x) {+  union {+    int64_t f;+    double t;+  } p;++  p.f = x;+  return p.t;+}++static inline double fmod64(double x, double y) {+  return fmod(x, y);+}++static inline double fsignum64(double x) {+  return futrts_isnan64(x) ? x : (x > 0) - (x < 0);+}++static inline double futrts_lerp64(double v0, double v1, double t) {+#ifdef __OPENCL_VERSION__+  return mix(v0, v1, t);+#else+  return v0 + (v1 - v0) * t;+#endif+}++static inline double futrts_mad64(double a, double b, double c) {+#ifdef __OPENCL_VERSION__+  return mad(a, b, c);+#else+  return a * b + c;+#endif+}++static inline float fpconv_f32_f32(float x) {+  return (float) x;+}++static inline double fpconv_f32_f64(float x) {+  return (double) x;+}++static inline float fpconv_f64_f32(double x) {+  return (float) x;+}++static inline double fpconv_f64_f64(double x) {+  return (double) x;+}++#endif++// End of scalar.h.

@@ -0,0 +1,627 @@+// Start of scalar_f16.h.++// Half-precision is emulated if needed (e.g. in straight C) with the+// native type used if possible.  The emulation works by typedef'ing+// 'float' to 'f16', and then implementing all operations on single+// precision.  To cut down on duplication, we use the same code for+// those Futhark functions that require just operators or casts.  The+// in-memory representation for arrays will still be 16 bits even+// under emulation, so the compiler will have to be careful when+// generating reads or writes.++#if !defined(cl_khr_fp16) && !(defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 600)+#define EMULATE_F16+#endif++#if !defined(EMULATE_F16) && defined(__OPENCL_VERSION__)+#pragma OPENCL EXTENSION cl_khr_fp16 : enable+#endif++#ifdef EMULATE_F16++// Note that the half-precision storage format is still 16 bits - the+// compiler will have to be real careful!+typedef float f16;++#else++#ifdef __CUDA_ARCH__+#include <cuda_fp16.h>+#endif++typedef half f16;++#endif++// Some of these functions convert to single precision because half+// precision versions are not available.++static inline f16 fadd16(f16 x, f16 y) {+  return x + y;+}++static inline f16 fsub16(f16 x, f16 y) {+  return x - y;+}++static inline f16 fmul16(f16 x, f16 y) {+  return x * y;+}++static inline bool cmplt16(f16 x, f16 y) {+  return x < y;+}++static inline bool cmple16(f16 x, f16 y) {+  return x <= y;+}++static inline f16 sitofp_i8_f16(int8_t x) {+  return (f16) x;+}++static inline f16 sitofp_i16_f16(int16_t x) {+  return (f16) x;+}++static inline f16 sitofp_i32_f16(int32_t x) {+  return (f16) x;+}++static inline f16 sitofp_i64_f16(int64_t x) {+  return (f16) x;+}++static inline f16 uitofp_i8_f16(uint8_t x) {+  return (f16) x;+}++static inline f16 uitofp_i16_f16(uint16_t x) {+  return (f16) x;+}++static inline f16 uitofp_i32_f16(uint32_t x) {+  return (f16) x;+}++static inline f16 uitofp_i64_f16(uint64_t x) {+  return (f16) x;+}++static inline int8_t fptosi_f16_i8(f16 x) {+  return (int8_t) (float) x;+}++static inline int16_t fptosi_f16_i16(f16 x) {+  return (int16_t) x;+}++static inline int32_t fptosi_f16_i32(f16 x) {+  return (int32_t) x;+}++static inline int64_t fptosi_f16_i64(f16 x) {+  return (int64_t) x;+}++static inline uint8_t fptoui_f16_i8(f16 x) {+  return (uint8_t) (float) x;+}++static inline uint16_t fptoui_f16_i16(f16 x) {+  return (uint16_t) x;+}++static inline uint32_t fptoui_f16_i32(f16 x) {+  return (uint32_t) x;+}++static inline uint64_t fptoui_f16_i64(f16 x) {+  return (uint64_t) x;+}++#ifndef EMULATE_F16++#ifdef __OPENCL_VERSION__+static inline f16 fabs16(f16 x) {+  return fabs(x);+}++static inline f16 fmax16(f16 x, f16 y) {+  return fmax(x, y);+}++static inline f16 fmin16(f16 x, f16 y) {+  return fmin(x, y);+}++static inline f16 fpow16(f16 x, f16 y) {+  return pow(x, y);+}++#else // Assuming CUDA.++static inline f16 fabs16(f16 x) {+  return fabsf(x);+}++static inline f16 fmax16(f16 x, f16 y) {+  return fmaxf(x, y);+}++static inline f16 fmin16(f16 x, f16 y) {+  return fminf(x, y);+}++static inline f16 fpow16(f16 x, f16 y) {+  return powf(x, y);+}+#endif++static inline bool futrts_isnan16(f16 x) {+  return isnan((float)x);+}++static inline bool futrts_isinf16(f16 x) {+  return isinf((float)x);+}++#ifdef __OPENCL_VERSION__+static inline f16 futrts_log16(f16 x) {+  return log(x);+}++static inline f16 futrts_log2_16(f16 x) {+  return log2(x);+}++static inline f16 futrts_log10_16(f16 x) {+  return log10(x);+}++static inline f16 futrts_sqrt16(f16 x) {+  return sqrt(x);+}++static inline f16 futrts_exp16(f16 x) {+  return exp(x);+}++static inline f16 futrts_cos16(f16 x) {+  return cos(x);+}++static inline f16 futrts_sin16(f16 x) {+  return sin(x);+}++static inline f16 futrts_tan16(f16 x) {+  return tan(x);+}++static inline f16 futrts_acos16(f16 x) {+  return acos(x);+}++static inline f16 futrts_asin16(f16 x) {+  return asin(x);+}++static inline f16 futrts_atan16(f16 x) {+  return atan(x);+}++static inline f16 futrts_cosh16(f16 x) {+  return cosh(x);+}++static inline f16 futrts_sinh16(f16 x) {+  return sinh(x);+}++static inline f16 futrts_tanh16(f16 x) {+  return tanh(x);+}++static inline f16 futrts_acosh16(f16 x) {+  return acosh(x);+}++static inline f16 futrts_asinh16(f16 x) {+  return asinh(x);+}++static inline f16 futrts_atanh16(f16 x) {+  return atanh(x);+}++static inline f16 futrts_atan2_16(f16 x, f16 y) {+  return atan2(x, y);+}++static inline f16 futrts_hypot16(f16 x, f16 y) {+  return hypot(x, y);+}++static inline f16 futrts_gamma16(f16 x) {+  return tgamma(x);+}++static inline f16 futrts_lgamma16(f16 x) {+  return lgamma(x);+}++static inline f16 fmod16(f16 x, f16 y) {+  return fmod(x, y);+}++static inline f16 futrts_round16(f16 x) {+  return rint(x);+}++static inline f16 futrts_floor16(f16 x) {+  return floor(x);+}++static inline f16 futrts_ceil16(f16 x) {+  return ceil(x);+}++static inline f16 futrts_lerp16(f16 v0, f16 v1, f16 t) {+  return mix(v0, v1, t);+}++static inline f16 futrts_mad16(f16 a, f16 b, f16 c) {+  return mad(a, b, c);+}++static inline f16 futrts_fma16(f16 a, f16 b, f16 c) {+  return fma(a, b, c);+}++#else // Assume CUDA.++static inline f16 futrts_log16(f16 x) {+  return hlog(x);+}++static inline f16 futrts_log2_16(f16 x) {+  return hlog2(x);+}++static inline f16 futrts_log10_16(f16 x) {+  return hlog10(x);+}++static inline f16 futrts_sqrt16(f16 x) {+  return hsqrt(x);+}++static inline f16 futrts_exp16(f16 x) {+  return hexp(x);+}++static inline f16 futrts_cos16(f16 x) {+  return hcos(x);+}++static inline f16 futrts_sin16(f16 x) {+  return hsin(x);+}++static inline f16 futrts_tan16(f16 x) {+  return tanf(x);+}++static inline f16 futrts_acos16(f16 x) {+  return acosf(x);+}++static inline f16 futrts_asin16(f16 x) {+  return asinf(x);+}++static inline f16 futrts_atan16(f16 x) {+  return atanf(x);+}++static inline f16 futrts_cosh16(f16 x) {+  return coshf(x);+}++static inline f16 futrts_sinh16(f16 x) {+  return sinhf(x);+}++static inline f16 futrts_tanh16(f16 x) {+  return tanhf(x);+}++static inline f16 futrts_acosh16(f16 x) {+  return acoshf(x);+}++static inline f16 futrts_asinh16(f16 x) {+  return asinhf(x);+}++static inline f16 futrts_atanh16(f16 x) {+  return atanhf(x);+}++static inline f16 futrts_atan2_16(f16 x, f16 y) {+  return atan2f(x, y);+}++static inline f16 futrts_hypot16(f16 x, f16 y) {+  return hypotf(x, y);+}++static inline f16 futrts_gamma16(f16 x) {+  return tgammaf(x);+}++static inline f16 futrts_lgamma16(f16 x) {+  return lgammaf(x);+}++static inline f16 fmod16(f16 x, f16 y) {+  return fmodf(x, y);+}++static inline f16 futrts_round16(f16 x) {+  return rintf(x);+}++static inline f16 futrts_floor16(f16 x) {+  return hfloor(x);+}++static inline f16 futrts_ceil16(f16 x) {+  return hceil(x);+}++static inline f16 futrts_lerp16(f16 v0, f16 v1, f16 t) {+  return v0 + (v1 - v0) * t;+}++static inline f16 futrts_mad16(f16 a, f16 b, f16 c) {+  return a * b + c;+}++static inline f16 futrts_fma16(f16 a, f16 b, f16 c) {+  return fmaf(a, b, c);+}++#endif++// The CUDA __half type cannot be put in unions for some reason, so we+// use bespoke conversion functions instead.+#ifdef __CUDA_ARCH__+static inline int16_t futrts_to_bits16(f16 x) {+  return __half_as_ushort(x);+}+static inline f16 futrts_from_bits16(int16_t x) {+  return __ushort_as_half(x);+}+#else+static inline int16_t futrts_to_bits16(f16 x) {+  union {+    f16 f;+    int16_t t;+  } p;++  p.f = x;+  return p.t;+}++static inline f16 futrts_from_bits16(int16_t x) {+  union {+    int16_t f;+    f16 t;+  } p;++  p.f = x;+  return p.t;+}+#endif++#else // No native f16 - emulate.++static inline f16 fabs16(f16 x) {+  return fabs32(x);+}++static inline f16 fmax16(f16 x, f16 y) {+  return fmax32(x, y);+}++static inline f16 fmin16(f16 x, f16 y) {+  return fmin32(x, y);+}++static inline f16 fpow16(f16 x, f16 y) {+  return fpow32(x, y);+}++static inline bool futrts_isnan16(f16 x) {+  return futrts_isnan32(x);+}++static inline bool futrts_isinf16(f16 x) {+  return futrts_isinf32(x);+}++static inline f16 futrts_log16(f16 x) {+  return futrts_log32(x);+}++static inline f16 futrts_log2_16(f16 x) {+  return futrts_log2_32(x);+}++static inline f16 futrts_log10_16(f16 x) {+  return futrts_log10_32(x);+}++static inline f16 futrts_sqrt16(f16 x) {+  return futrts_sqrt32(x);+}++static inline f16 futrts_exp16(f16 x) {+  return futrts_exp32(x);+}++static inline f16 futrts_cos16(f16 x) {+  return futrts_cos32(x);+}++static inline f16 futrts_sin16(f16 x) {+  return futrts_sin32(x);+}++static inline f16 futrts_tan16(f16 x) {+  return futrts_tan32(x);+}++static inline f16 futrts_acos16(f16 x) {+  return futrts_acos32(x);+}++static inline f16 futrts_asin16(f16 x) {+  return futrts_asin32(x);+}++static inline f16 futrts_atan16(f16 x) {+  return futrts_atan32(x);+}++static inline f16 futrts_cosh16(f16 x) {+  return futrts_cosh32(x);+}++static inline f16 futrts_sinh16(f16 x) {+  return futrts_sinh32(x);+}++static inline f16 futrts_tanh16(f16 x) {+  return futrts_tanh32(x);+}++static inline f16 futrts_acosh16(f16 x) {+  return futrts_acosh32(x);+}++static inline f16 futrts_asinh16(f16 x) {+  return futrts_asinh32(x);+}++static inline f16 futrts_atanh16(f16 x) {+  return futrts_atanh32(x);+}++static inline f16 futrts_atan2_16(f16 x, f16 y) {+  return futrts_atan2_32(x, y);+}++static inline f16 futrts_hypot16(f16 x, f16 y) {+  return futrts_hypot32(x, y);+}++static inline f16 futrts_gamma16(f16 x) {+  return futrts_gamma32(x);+}++static inline f16 futrts_lgamma16(f16 x) {+  return futrts_lgamma32(x);+}++static inline f16 fmod16(f16 x, f16 y) {+  return fmod32(x, y);+}++static inline f16 futrts_round16(f16 x) {+  return futrts_round32(x);+}++static inline f16 futrts_floor16(f16 x) {+  return futrts_floor32(x);+}++static inline f16 futrts_ceil16(f16 x) {+  return futrts_ceil32(x);+}++static inline f16 futrts_lerp16(f16 v0, f16 v1, f16 t) {+  return futrts_lerp32(v0, v1, t);+}++static inline f16 futrts_mad16(f16 a, f16 b, f16 c) {+  return futrts_mad32(a, b, c);+}++static inline f16 futrts_fma16(f16 a, f16 b, f16 c) {+  return futrts_fma32(a, b, c);+}++// Even when we are using an OpenCL that does not support cl_khr_fp16,+// it must still support vload_half for actually creating a+// half-precision number, which can then be efficiently converted to a+// float.  Similarly for vstore_half.+#ifdef __OPENCL_VERSION__++static inline int16_t futrts_to_bits16(f16 x) {+  int16_t y;+  // Violating strict aliasing here.+  vstore_half((float)x, 0, (half*)&y);+  return y;+}++static inline f16 futrts_from_bits16(int16_t x) {+  return (f16)vload_half(0, (half*)&x);+}++#else++static inline int16_t futrts_to_bits16(f16 x) {+  return (int16_t)float2halfbits(x);+}++static inline f16 futrts_from_bits16(int16_t x) {+  return halfbits2float((uint16_t)x);+}++static inline f16 fsignum16(f16 x) {+  return futrts_isnan16(x) ? x : (x > 0) - (x < 0);+}++#endif++#endif++static inline float fpconv_f16_f16(f16 x) {+  return x;+}++static inline float fpconv_f16_f32(f16 x) {+  return x;+}++static inline f16 fpconv_f32_f16(float x) {+  return x;+}++#ifdef FUTHARK_F64_ENABLED++static inline double fpconv_f16_f64(f16 x) {+  return (double) x;+}++static inline f16 fpconv_f64_f16(double x) {+  return (f16) x;+}++#endif+++// End of scalar_f16.h.

@@ -100,6 +100,11 @@ #include <sys/sysinfo.h> #include <sys/resource.h> #include <signal.h>+#elif defined(__EMSCRIPTEN__)+#include <emscripten/threading.h>+#include <sys/sysinfo.h>+#include <sys/resource.h>+#include <signal.h> #endif  /* Multicore Utility functions */@@ -128,7 +133,7 @@     } else {         errno = EINVAL;     }-#elif defined(__linux__)+#elif defined(__linux__) || __EMSCRIPTEN__     err = getrusage(RUSAGE_THREAD, rusage); #endif     return err;@@ -152,6 +157,8 @@     return ncores; #elif defined(__linux__)   return get_nprocs();+#elif __EMSCRIPTEN__+  return emscripten_num_logical_cores(); #else   fprintf(stderr, "operating system not recognised\n");   return -1;

@@ -57,6 +57,7 @@ DEF_SCALAR_TYPE(u16); DEF_SCALAR_TYPE(u32); DEF_SCALAR_TYPE(u64);+DEF_SCALAR_TYPE(f16); DEF_SCALAR_TYPE(f32); DEF_SCALAR_TYPE(f64); DEF_SCALAR_TYPE(bool);@@ -75,6 +76,7 @@     uint32_t v_u32;     uint64_t v_u64; +    uint16_t v_f16;     float v_f32;     double v_f64; @@ -107,6 +109,9 @@   if (v->type == &type_u64) {     return &v->value.v_u64;   }+  if (v->type == &type_f16) {+    return &v->value.v_f16;+  }   if (v->type == &type_f32) {     return &v->value.v_f32;   }@@ -286,8 +291,9 @@    int num_outs = entry_num_outs(e);   int num_ins = entry_num_ins(e);-  void* outs[num_outs];-  void* ins[num_ins];+  // +1 to avoid zero-size arrays, which is UB.+  void* outs[num_outs+1];+  void* ins[num_ins+1];    for (int i = 0; i < num_ins; i++) {     const char *in_name = get_arg(args, 1+num_outs+i);

@@ -20,7 +20,7 @@     if (eql) {       *eql = 0;       int value = atoi(eql+1);-      if (set_size(cfg, line, value) != 0) {+      if (set_size(cfg, line, (size_t)value) != 0) {         char* err = (char*) malloc(max_line_len + 50);         snprintf(err, max_line_len + 50, "Unknown name '%s' on line %d.", line, lineno);         free(line);

@@ -7,6 +7,14 @@  static const char *fut_progname = "(embedded Futhark)"; +static void futhark_panic(int eval, const char *fmt, ...) __attribute__((noreturn));+static char* msgprintf(const char *s, ...);+static void* slurp_file(const char *filename, size_t *size);+static int dump_file(const char *file, const void *buf, size_t n);+struct str_builder;+static void str_builder_init(struct str_builder *b);+static void str_builder(struct str_builder *b, const char *s, ...);+ static void futhark_panic(int eval, const char *fmt, ...) {   va_list ap;   va_start(ap, fmt);@@ -28,7 +36,6 @@   return buffer; } - static inline void check_err(int errval, int sets_errno, const char *fun, int line,                             const char *msg, ...) {   if (errval) {@@ -46,18 +53,18 @@   } } -#define CHECK_ERR(err, msg...) check_err(err, 0, __func__, __LINE__, msg)-#define CHECK_ERRNO(err, msg...) check_err(err, 1, __func__, __LINE__, msg)+#define CHECK_ERR(err, ...) check_err(err, 0, __func__, __LINE__, __VA_ARGS__)+#define CHECK_ERRNO(err, ...) check_err(err, 1, __func__, __LINE__, __VA_ARGS__)  // Read the rest of an open file into a NUL-terminated string; returns // NULL on error. static void* fslurp_file(FILE *f, size_t *size) {-  size_t start = ftell(f);+  long start = ftell(f);   fseek(f, 0, SEEK_END);-  size_t src_size = ftell(f)-start;+  long src_size = ftell(f)-start;   fseek(f, start, SEEK_SET);-  unsigned char *s = (unsigned char*) malloc(src_size + 1);-  if (fread(s, 1, src_size, f) != src_size) {+  unsigned char *s = (unsigned char*) malloc((size_t)src_size + 1);+  if (fread(s, 1, (size_t)src_size, f) != (size_t)src_size) {     free(s);     s = NULL;   } else {@@ -65,7 +72,7 @@   }    if (size) {-    *size = src_size;+    *size = (size_t)src_size;   }    return s;

@@ -107,7 +107,7 @@   int ret;   int first = 1;   char *knows_dimsize = (char*) calloc((size_t)dims, sizeof(char));-  int cur_dim = dims-1;+  int cur_dim = (int)dims-1;   int64_t *elems_read_in_dim = (int64_t*) calloc((size_t)dims, sizeof(int64_t));    while (1) {@@ -337,16 +337,40 @@   READ_STR(SCNu64, uint64_t, "u64"); } +static int read_str_f16(char *buf, void* dest) {+  remove_underscores(buf);+  if (strcmp(buf, "f16.nan") == 0) {+    *(uint16_t*)dest = float2halfbits(NAN);+    return 0;+  } else if (strcmp(buf, "f16.inf") == 0) {+    *(uint16_t*)dest = float2halfbits(INFINITY);+    return 0;+  } else if (strcmp(buf, "-f16.inf") == 0) {+    *(uint16_t*)dest = float2halfbits(-INFINITY);+    return 0;+  } else {+    int j;+    float x;+    if (sscanf(buf, "%f%n", &x, &j) == 1) {+      if (strcmp(buf+j, "") == 0 || strcmp(buf+j, "f16") == 0) {+        *(uint16_t*)dest = float2halfbits(x);+        return 0;+      }+    }+    return 1;+  }+}+ static int read_str_f32(char *buf, void* dest) {   remove_underscores(buf);   if (strcmp(buf, "f32.nan") == 0) {-    *(float*)dest = NAN;+    *(float*)dest = (float)NAN;     return 0;   } else if (strcmp(buf, "f32.inf") == 0) {-    *(float*)dest = INFINITY;+    *(float*)dest = (float)INFINITY;     return 0;   } else if (strcmp(buf, "-f32.inf") == 0) {-    *(float*)dest = -INFINITY;+    *(float*)dest = (float)-INFINITY;     return 0;   } else {     READ_STR("f", float, "f32");@@ -356,13 +380,13 @@ static int read_str_f64(char *buf, void* dest) {   remove_underscores(buf);   if (strcmp(buf, "f64.nan") == 0) {-    *(double*)dest = NAN;+    *(double*)dest = (double)NAN;     return 0;   } else if (strcmp(buf, "f64.inf") == 0) {-    *(double*)dest = INFINITY;+    *(double*)dest = (double)INFINITY;     return 0;   } else if (strcmp(buf, "-f64.inf") == 0) {-    *(double*)dest = -INFINITY;+    *(double*)dest = (double)-INFINITY;     return 0;   } else {     READ_STR("lf", double, "f64");@@ -413,6 +437,19 @@   return fprintf(out, "%"PRIu64"u64", *src); } +static int write_str_f16(FILE *out, uint16_t *src) {+  float x = halfbits2float(*src);+  if (isnan(x)) {+    return fprintf(out, "f16.nan");+  } else if (isinf(x) && x >= 0) {+    return fprintf(out, "f16.inf");+  } else if (isinf(x)) {+    return fprintf(out, "-f16.inf");+  } else {+    return fprintf(out, "%.6ff16", x);+  }+}+ static int write_str_f32(FILE *out, float *src) {   float x = *src;   if (isnan(x)) {@@ -448,10 +485,10 @@ #define BINARY_FORMAT_VERSION 2 #define IS_BIG_ENDIAN (!*(unsigned char *)&(uint16_t){1}) -static void flip_bytes(int elem_size, unsigned char *elem) {-  for (int j=0; j<elem_size/2; j++) {+static void flip_bytes(size_t elem_size, unsigned char *elem) {+  for (size_t j=0; j<elem_size/2; j++) {     unsigned char head = elem[j];-    int tail_index = elem_size-1-j;+    size_t tail_index = elem_size-1-j;     elem[j] = elem[tail_index];     elem[tail_index] = head;   }@@ -472,7 +509,7 @@ #endif  static int read_byte(FILE *f, void* dest) {-  int num_elems_read = fread(dest, 1, 1, f);+  size_t num_elems_read = fread(dest, 1, 1, f);   return num_elems_read == 1 ? 0 : 1; } @@ -510,6 +547,9 @@ static const struct primtype_info_t u64_info =   {.binname = " u64", .type_name = "u64",  .size = 8,    .write_str = (writer)write_str_u64, .read_str = (str_reader)read_str_u64};+static const struct primtype_info_t f16_info =+  {.binname = " f16", .type_name = "f16",  .size = 2,+   .write_str = (writer)write_str_f16, .read_str = (str_reader)read_str_f16}; static const struct primtype_info_t f32_info =   {.binname = " f32", .type_name = "f32",  .size = 4,    .write_str = (writer)write_str_f32, .read_str = (str_reader)read_str_f32};@@ -523,7 +563,7 @@ static const struct primtype_info_t* primtypes[] = {   &i8_info, &i16_info, &i32_info, &i64_info,   &u8_info, &u16_info, &u32_info, &u64_info,-  &f32_info, &f64_info,+  &f16_info, &f32_info, &f64_info,   &bool_info,   NULL // NULL-terminated };@@ -612,7 +652,7 @@   int64_t elem_count = 1;   for (int i=0; i<dims; i++) {     int64_t bin_shape;-    ret = fread(&bin_shape, sizeof(bin_shape), 1, f);+    ret = (int)fread(&bin_shape, sizeof(bin_shape), 1, f);     if (ret != 1) {       futhark_panic(1, "binary-input: Couldn't read size for dimension %i of array.\n", i);     }@@ -640,7 +680,7 @@   // If we're on big endian platform we must change all multibyte elements   // from using little endian to big endian   if (IS_BIG_ENDIAN && elem_size != 1) {-    flip_bytes(elem_size, (unsigned char*) *data);+    flip_bytes((size_t)elem_size, (unsigned char*) *data);   }    return 0;@@ -671,7 +711,7 @@                            const int64_t *shape,                            int8_t rank) {   if (rank==0) {-    elem_type->write_str(out, (void*)data);+    elem_type->write_str(out, (const void*)data);   } else {     int64_t len = (int64_t)shape[0];     int64_t slice_size = 1;@@ -691,7 +731,7 @@     } else if (rank==1) {       fputc('[', out);       for (int64_t i = 0; i < len; i++) {-        elem_type->write_str(out, (void*) (data + i * elem_size));+        elem_type->write_str(out, (const void*) (data + i * elem_size));         if (i != len-1) {           fprintf(out, ", ");         }@@ -763,8 +803,8 @@     return expected_type->read_str(buf, dest);   } else {     read_bin_ensure_scalar(f, expected_type);-    int64_t elem_size = expected_type->size;-    int num_elems_read = fread(dest, (size_t)elem_size, 1, f);+    size_t elem_size = (size_t)expected_type->size;+    size_t num_elems_read = fread(dest, elem_size, 1, f);     if (IS_BIG_ENDIAN) {       flip_bytes(elem_size, (unsigned char*) dest);     }

@@ -0,0 +1,251 @@+// Start of server.js++class Server {++  constructor(ctx) {+    this.ctx = ctx;+    this._vars = {};+    this._types = {};+    this._commands = [ 'inputs',+                       'outputs',+                       'call',+                       'restore',+                       'store',+                       'free',+                       'clear',+                       'pause_profiling',+                       'unpause_profiling',+                       'report',+                       'rename'+                     ];+  }++  _get_arg(args, i) {+    if (i < args.length) {+      return args[i];+    } else {+      throw 'Insufficient command args';+    }+  }++  _get_entry_point(entry) {+    if (entry in this.ctx.get_entry_points()) {+      return this.ctx.get_entry_points()[entry];+    } else {+      throw "Unkown entry point: " + entry;+    }+  }++  _check_var(vname) {+    if (!(vname in this._vars)) {+      throw 'Unknown variable: ' + vname;+    }+  }++  _set_var(vname, v, t) {+    this._vars[vname] = v;+    this._types[vname] = t;+  }++  _get_type(vname) {+    this._check_var(vname);+    return this._types[vname];+  }++  _get_var(vname) {+    this._check_var(vname);+    return this._vars[vname];+  }++  _cmd_inputs(args) {+    var entry = this._get_arg(args, 0);+    var inputs = this._get_entry_point(entry)[0];+    for (var i = 0; i < inputs.length; i++) {+      console.log(inputs[i]);+    }+  }++  _cmd_outputs(args) {+    var entry = this._get_arg(args, 0);+    var outputs = this._get_entry_point(entry)[1];+    for (var i = 0; i < outputs.length; i++) {+      console.log(outputs[i]);+    }+  }++  _cmd_dummy(args) {+    // pass+  }++  _cmd_free(args) {+    for (var i = 0; i < args.length; i++) {+      var vname = args[i];+      this._check_var(vname);+      delete this._vars[vname];+    }+  }++  _cmd_rename(args) {+    var oldname = this._get_arg(args, 0)+    var newname = this._get_arg(args, 1)+    if (newname in this._vars) {+      throw "Variable already exists: " + newname;+    }+    this._vars[newname] = this._vars[oldname];+    this._types[newname] = this._types[oldname];+    delete this._vars[oldname];+    delete this._types[oldname];+  }++  _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 expected_len = 1 + num_outs + num_ins++    if (args.length != expected_len) {+      throw "Invalid argument count, expected " + expected_len+    }++    var out_vnames = args.slice(1, num_outs+1)+    for (var i = 0; i < out_vnames.length; i++) {+      var out_vname = out_vnames[i];+      if (out_vname in this._vars) {+        throw "Variable already exists: " + out_vname;+      }+    }+    var in_vnames = args.slice(1+num_outs);+    var ins = [];+    for (var i = 0; i < in_vnames.length; i++) {+      ins.push(this._get_var(in_vnames[i]));+    }+    // Call entry point function from string name+    var bef = performance.now()*1000;+    var vals = this.ctx[args[0]].apply(this.ctx, ins);+    var aft = performance.now()*1000;+    if (num_outs == 1) {+      this._set_var(out_vnames[0], vals, entry[1][0]);+    } else {+      for (var i = 0; i < out_vnames.length; i++) {+        this._set_var(out_vnames[i], vals[i], entry[1][i]);+      }+    }+    console.log("runtime: " + Math.round(aft-bef));+  }++  _cmd_store(args) {+    var fname = this._get_arg(args, 0);+    for (var i = 1; i < args.length; i++) {+      var vname = args[i];+      var value = this._get_var(vname);+      var typ = this._get_type(vname);+      var fs = require("fs");+      var bin_val = construct_binary_value(value, typ);+      fs.appendFileSync(fname, bin_val, 'binary')+    }+  }++  fut_to_dim_typ(typ) {+    var type = typ;+    var count = 0;+    while (type.substr(0, 2) == '[]') {+      count = count + 1;+      type = type.slice(2);+    }+    return [count, type];+  }++  _cmd_restore(args) {+    if (args.length % 2 == 0) {+      throw "Invalid argument count";+    }++    var fname = args[0];+    var args = args.slice(1);++    var as = args;+    var reader = new Reader(fname);+    while (as.length != 0) {+      var vname = as[0];+      var typename = as[1];+      as = as.slice(2);++      if (vname in this._vars) {+        throw "Variable already exists: " + vname;+      }+      try {+        var value = read_value(typename, reader);+        if (typeof value == 'number' || typeof value == 'bigint') {+          this._set_var(vname, value, typename);+        } else {+          // We are working with an array and need to create to convert [shape, arr] to futhark ptr+          var shape= value[0];+          var arr = value[1];+          var dimtyp = this.fut_to_dim_typ(typename);+          var dim = dimtyp[0];+          var typ = dimtyp[1];+          var arg_list = [arr, ...shape];+          var fnam = "new_" + typ + "_" + dim + "d";+          var ptr = this.ctx[fnam].apply(this.ctx, arg_list);+          this._set_var(vname, ptr, typename);+        }+      } catch (err) {+        var err_msg = "Failed to restore variable " + vname + ".\nPossibly malformed data in " + fname + ".\n" + err.toString();+        throw err_msg;+      }+    }+    skip_spaces(reader);+    if (reader.get_buff().length != 0) {+      throw "Expected EOF after reading values";+    }+  }++  _process_line(line) {+    // TODO make sure it splits on anywhite space+    var words = line.split(" ");+    if (words.length == 0) {+      throw "Empty line";+    } else {+      var cmd = words[0];+      var args = words.splice(1);+      if (this._commands.includes(cmd)) {+        switch (cmd) {+        case 'inputs': this._cmd_inputs(args); break;+        case 'outputs': this._cmd_outputs(args); break+        case 'call': this._cmd_call(args); break+        case 'restore': this._cmd_restore(args); break+        case 'store': this._cmd_store(args); break+        case 'free': this._cmd_free(args); break+        case 'clear': this._cmd_dummy(args); break+        case 'pause_profiling': this._cmd_dummy(args); break+        case 'unpause_profiling': this._cmd_dummy(args); break+        case 'report': this._cmd_dummy(args); break+        case 'rename': this._cmd_rename(args); break+        }+      } else {+        throw "Unknown command: " + cmd;+      }+    }+  }++  run() {+    console.log('%%% OK'); // TODO figure out if flushing is neccesary for JS+    const readline = require('readline');+    const rl = readline.createInterface(process.stdin);+    rl.on('line', (line) => {+      if (line == "") {+        rl.close();+      }+      try {+        this._process_line(line);+        console.log('%%% OK');+      } catch (err) {+        console.log('%%% FAILURE');+        console.log(err);+        console.log('%%% OK');+      }+    }).on('close', () => { process.exit(0); });+  }+}++// End of server.js

@@ -0,0 +1,165 @@+// Start of values.js++var typToType = { '  i8' : Int8Array ,+                  ' i16' : Int16Array ,+                  ' i32' : Int32Array ,+                  ' i64' : BigInt64Array ,+                  '  u8' : Uint8Array ,+                  ' u16' : Uint16Array ,+                  ' u32' : Uint32Array ,+                  ' u64' : BigUint64Array ,+                  ' f16' : Uint16Array ,+                  ' f32' : Float32Array ,+                  ' f64' : Float64Array ,+                  'bool' : Uint8Array+                };++function binToStringArray(buff, array) {+  for (var i = 0; i < array.length; i++) {+    array[i] = buff[i];+  }+}++function fileToBuff(fname) {+  var readline = require('readline');+  var fs = require('fs');+  var buff =  fs.readFileSync(fname);+  return buff;+}++var typToSize = {+  "bool" : 1,+  "  u8" : 1,+  "  i8" : 1,+  " u16" : 2,+  " i16" : 2,+  " u32" : 4,+  " i32" : 4,+  " f16" : 2,+  " f32" : 4,+  " u64" : 8,+  " i64" : 8,+  " f64" : 8,+}++function toU8(ta) {+  return new Uint8Array(ta.buffer, ta.byteOffset, ta.byteLength);+}++function construct_binary_value(v, typ) {+  var dims;+  var payload_bytes;+  var filler;+  if (v instanceof FutharkOpaque) {+    throw "Opaques are not supported";+  } else if (v instanceof FutharkArray) {+    var t = v.futharkType();+    var ftype = "    ".slice(t.length) + t;+    var shape = v.shape();+    var ta = v.toTypedArray(shape);+    var da = new BigInt64Array(shape);+    dims = shape.length;+    payload_bytes = da.byteLength + ta.byteLength;+    filler = (bytes) => {+      bytes.set(toU8(da), 7);+      bytes.set(toU8(ta), 7 + da.byteLength);+    }+  } else {+    var ftype = "    ".slice(typ.length) + typ;+    dims = 0;+    payload_bytes = typToSize[ftype];+    filler = (bytes) => {+      var scalar = new (typToType[ftype])([v]);+      bytes.set(toU8(scalar), 7);+    }+  }+  var total_bytes = 7 + payload_bytes;+  var bytes = new Uint8Array(total_bytes);+  bytes[0] = Buffer.from('b').readUInt8();+  bytes[1] = 2;+  bytes[2] = dims;+  for (var i = 0; i < 4; i++) {+    bytes[3+i] = ftype.charCodeAt(i);+  }+  filler(bytes);+  return Buffer.from(bytes);+}++class Reader {+  constructor(f) {+    this.f = f;+    this.buff = fileToBuff(f);+  }++  read_bin_array(num_dim, typ) {+    var u8_array = new Uint8Array(num_dim * 8);+    binToStringArray(this.buff.slice(0, num_dim * 8), u8_array);+    var shape = new BigInt64Array(u8_array.buffer);+    var length = shape[0];+    for (var i = 1; i < shape.length; i++) {+      length = length * shape[i];+    }+    length = Number(length);+    var dbytes = typToSize[typ];+    var u8_data = new Uint8Array(length * dbytes);+    binToStringArray(this.buff.slice(num_dim * 8, num_dim * 8 + dbytes * length), u8_data);+    var data  = new (typToType[typ])(u8_data.buffer);+    var tmp_buff = this.buff.slice(num_dim * 8, num_dim * 8 + dbytes * length);+    this.buff = this.buff.slice(num_dim * 8 + dbytes * length);+   return [shape, data];+  }++  read_bin_scalar(typ) {+    var size = typToSize[typ];+    var u8_array = new Uint8Array(size);+    binToStringArray(this.buff, u8_array);+    var array = new (typToType[typ])(u8_array.buffer);+    this.buff = this.buff.slice(u8_array.length); // Update buff to be unread part of the string+    return array[0];+  }++  skip_spaces() {+    while (this.buff.length > 0 && this.buff.slice(0, 1).toString().trim() == "") {+      this.buff = this.buff.slice(1);+    }+  }++  read_binary() {+    // Skip leading white space+    while (this.buff.slice(0, 1).toString().trim() == "") {+      this.buff = this.buff.slice(1);+    }+    if (this.buff[0] != 'b'.charCodeAt(0)) {+      throw "Not in binary format"+    }+    var version = this.buff[1];+    if (version != 2) {+      throw "Not version 2";+    }+    var num_dim = this.buff[2];+    var typ = this.buff.slice(3, 7);+    this.buff = this.buff.slice(7);+    if (num_dim == 0) {+      return this.read_bin_scalar(typ);+    } else {+      return this.read_bin_array(num_dim, typ);+    }+  }++  get_buff() {+    return this.buff;+  }+}++// Function is redudant but is helpful for keeping consistent with python implementation+function skip_spaces(reader) {+  reader.skip_spaces();+}++function read_value(typename, reader) {+  // TODO include typename in implementation+  var val = reader.read_binary();+  return val;+}++// End of values.js

@@ -0,0 +1,83 @@+// Start of wrapperclasses.js++class FutharkArray {+  constructor(ctx, ptr, type_name, dim, heap, fshape, fvalues, ffree) {+    this.ctx = ctx;+    this.ptr = ptr;+    this.type_name = type_name;+    this.dim = dim;+    this.heap = heap;+    this.fshape = fshape;+    this.fvalues = fvalues;+    this.ffree = ffree;+    this.valid = true;+  }++  validCheck() {+    if (!this.valid) {+      throw "Using freed memory"+    }+  }++  futharkType() {+    return this.type_name;+  }++  free() {+    this.validCheck();+    this.ffree(this.ctx.ctx, this.ptr);+    this.valid = false;+  }++  shape() {+    this.validCheck();+    var s = this.fshape(this.ctx.ctx, this.ptr) >> 3;+    return Array.from(this.ctx.wasm.HEAP64.subarray(s, s + this.dim));+  }++  toTypedArray(dims = this.shape()) {+    this.validCheck();+    console.assert(dims.length === this.dim, "dim=%s,dims=%s", this.dim, dims.toString());+    var length = Number(dims.reduce((a, b) => a * b));+    var v = this.fvalues(this.ctx.ctx, this.ptr) / this.heap.BYTES_PER_ELEMENT;+    return this.heap.subarray(v, v + length);+  }++  toArray() {+    this.validCheck();+    var dims = this.shape();+    var ta = this.toTypedArray(dims);+    return (function nest(offs, ds) {+      var d0 = Number(ds[0]);+      if (ds.length === 1) {+        return Array.from(ta.subarray(offs, offs + d0));+      } else {+        var d1 = Number(ds[1]);+        return Array.from(Array(d0), (x,i) => nest(offs + i * d1, ds.slice(1)));+      }+    })(0, dims);+  }+}++class FutharkOpaque {+  constructor(ctx, ptr, ffree) {+    this.ctx = ctx;+    this.ptr = ptr;+    this.ffree = ffree;+    this.valid = true;+  }++  validCheck() {+    if (!this.valid) {+      throw "Using freed memory"+    }+  }++  free() {+    this.validCheck();+    this.ffree(this.ctx.ctx, this.ptr);+    this.valid = false;+  }+}++// End of wrapperclasses.js

@@ -19,16 +19,16 @@ def unwrapArray(x):   return normaliseArray(x).ctypes.data_as(ct.POINTER(ct.c_byte)) -def createArray(x, shape):+def createArray(x, shape, t):   # HACK: np.ctypeslib.as_array may fail if the shape contains zeroes,   # for some reason.   if any(map(lambda x: x == 0, shape)):-      return np.ndarray(shape, dtype=x._type_)+      return np.ndarray(shape, dtype=t)   else:-      return np.ctypeslib.as_array(x, shape=shape)+      return np.ctypeslib.as_array(x, shape=shape).view(t) -def indexArray(x, offset, bt, nptype):-  return nptype(addressOffset(x, offset*ct.sizeof(bt), bt)[0])+def indexArray(x, offset, bt):+  return addressOffset(x, offset*ct.sizeof(bt), bt)[0]  def writeScalarArray(x, offset, v):   ct.memmove(ct.addressof(x.contents)+int(offset)*ct.sizeof(v), ct.addressof(v), ct.sizeof(v))

@@ -214,9 +214,16 @@     # compiler should provide us with the variables to which     # parameters are mapped.     if (len(program_src) >= 0):-        return cl.Program(self.ctx, program_src).build(-            ["-DLOCKSTEP_WIDTH={}".format(lockstep_width)]-            + ["-D{}={}".format(s.replace('z', 'zz').replace('.', 'zi').replace('#', 'zh'),v) for (s,v) in self.sizes.items()])+        build_options = []++        build_options += ["-DLOCKSTEP_WIDTH={}".format(lockstep_width)]++        build_options += ["-D{}={}".format(s.replace('z', 'zz').replace('.', 'zi').replace('#', 'zh'),v) for (s,v) in self.sizes.items()]++        if (self.platform.name == 'Oclgrind'):+            build_options += ['-DEMULATE_F16']++        return cl.Program(self.ctx, program_src).build(build_options)  def opencl_alloc(self, min_size, tag):     min_size = 1 if min_size == 0 else min_size

@@ -238,9 +238,9 @@ umax8 = umax16 = umax32 = umax64 = umaxN umin8 = umin16 = umin32 = umin64 = uminN pow8 = pow16 = pow32 = pow64 = powN-fpow32 = fpow64 = fpowN-fmax32 = fmax64 = fmaxN-fmin32 = fmin64 = fminN+fpow16 = fpow32 = fpow64 = fpowN+fmax16 = fmax32 = fmax64 = fmaxN+fmin16 = fmin32 = fmin64 = fminN sle8 = sle16 = sle32 = sle64 = sleN slt8 = slt16 = slt32 = slt64 = sltN ule8 = ule16 = ule32 = ule64 = uleN@@ -309,39 +309,51 @@  def fptosi_T_i8(x):   return np.int8(np.trunc(x))-fptosi_f32_i8 = fptosi_f64_i8 = fptosi_T_i8+fptosi_f16_i8 = fptosi_f32_i8 = fptosi_f64_i8 = fptosi_T_i8  def fptosi_T_i16(x):   return np.int16(np.trunc(x))-fptosi_f32_i16 = fptosi_f64_i16 = fptosi_T_i16+fptosi_f16_i16 = fptosi_f32_i16 = fptosi_f64_i16 = fptosi_T_i16  def fptosi_T_i32(x):   return np.int32(np.trunc(x))-fptosi_f32_i32 = fptosi_f64_i32 = fptosi_T_i32+fptosi_f16_i32 = fptosi_f32_i32 = fptosi_f64_i32 = fptosi_T_i32  def fptosi_T_i64(x):   return np.int64(np.trunc(x))-fptosi_f32_i64 = fptosi_f64_i64 = fptosi_T_i64+fptosi_f16_i64 = fptosi_f32_i64 = fptosi_f64_i64 = fptosi_T_i64  def fptoui_T_i8(x):   return np.uint8(np.trunc(x))-fptoui_f32_i8 = fptoui_f64_i8 = fptoui_T_i8+fptoui_f16_i8 = fptoui_f32_i8 = fptoui_f64_i8 = fptoui_T_i8  def fptoui_T_i16(x):   return np.uint16(np.trunc(x))-fptoui_f32_i16 = fptoui_f64_i16 = fptoui_T_i16+fptoui_f16_i16 = fptoui_f32_i16 = fptoui_f64_i16 = fptoui_T_i16  def fptoui_T_i32(x):   return np.uint32(np.trunc(x))-fptoui_f32_i32 = fptoui_f64_i32 = fptoui_T_i32+fptoui_f16_i32 = fptoui_f32_i32 = fptoui_f64_i32 = fptoui_T_i32  def fptoui_T_i64(x):   return np.uint64(np.trunc(x))-fptoui_f32_i64 = fptoui_f64_i64 = fptoui_T_i64+fptoui_f16_i64 = fptoui_f32_i64 = fptoui_f64_i64 = fptoui_T_i64 +def fpconv_f16_f32(x):+  return np.float32(x)++def fpconv_f16_f64(x):+  return np.float64(x)++def fpconv_f32_f16(x):+  return np.float16(x)+ def fpconv_f32_f64(x):   return np.float64(x) +def fpconv_f64_f16(x):+  return np.float16(x)+ def fpconv_f64_f32(x):   return np.float32(x) @@ -550,16 +562,111 @@   s = struct.pack('>l', x)   return np.float32(struct.unpack('>f', s)[0]) +def futhark_log16(x):+  return np.float16(np.log(x))++def futhark_log2_16(x):+  return np.float16(np.log2(x))++def futhark_log10_16(x):+  return np.float16(np.log10(x))++def futhark_sqrt16(x):+  return np.float16(np.sqrt(x))++def futhark_exp16(x):+  return np.exp(x)++def futhark_cos16(x):+  return np.cos(x)++def futhark_sin16(x):+  return np.sin(x)++def futhark_tan16(x):+  return np.tan(x)++def futhark_acos16(x):+  return np.arccos(x)++def futhark_asin16(x):+  return np.arcsin(x)++def futhark_atan16(x):+  return np.arctan(x)++def futhark_cosh16(x):+  return np.cosh(x)++def futhark_sinh16(x):+  return np.sinh(x)++def futhark_tanh16(x):+  return np.tanh(x)++def futhark_acosh16(x):+  return np.arccosh(x)++def futhark_asinh16(x):+  return np.arcsinh(x)++def futhark_atanh16(x):+  return np.arctanh(x)++def futhark_atan2_16(x, y):+  return np.arctan2(x, y)++def futhark_hypot16(x, y):+  return np.hypot(x, y)++def futhark_gamma16(x):+  return np.float16(math.gamma(x))++def futhark_lgamma16(x):+  return np.float16(math.lgamma(x))++def futhark_round16(x):+  return np.round(x)++def futhark_ceil16(x):+  return np.ceil(x)++def futhark_floor16(x):+  return np.floor(x)++def futhark_isnan16(x):+  return np.isnan(x)++def futhark_isinf16(x):+  return np.isinf(x)++def futhark_to_bits16(x):+  s = struct.pack('>e', x)+  return np.int16(struct.unpack('>H', s)[0])++def futhark_from_bits16(x):+  s = struct.pack('>H', np.uint16(x))+  return np.float16(struct.unpack('>e', s)[0])++def futhark_lerp16(v0, v1, t):+  return v0 + (v1-v0)*t+ def futhark_lerp32(v0, v1, t):   return v0 + (v1-v0)*t  def futhark_lerp64(v0, v1, t):   return v0 + (v1-v0)*t +def futhark_mad16(a, b, c):+  return a * b + c+ def futhark_mad32(a, b, c):   return a * b + c  def futhark_mad64(a, b, c):+  return a * b + c++def futhark_fma16(a, b, c):   return a * b + c  def futhark_fma32(a, b, c):

@@ -249,6 +249,24 @@         expt = b'0'     return float(sign + bef + b'.' + aft + b'E' + expt) +def read_str_f16(f):+    skip_spaces(f)+    try:+        parse_specific_string(f, 'f16.nan')+        return np.float32(np.nan)+    except ValueError:+        try:+            parse_specific_string(f, 'f16.inf')+            return np.float32(np.inf)+        except ValueError:+            try:+               parse_specific_string(f, '-f16.inf')+               return np.float32(-np.inf)+            except ValueError:+               x = read_str_decimal(f)+               optional_specific_string(f, 'f16')+               return x+ def read_str_f32(f):     skip_spaces(f)     try:@@ -385,6 +403,7 @@ read_bin_u32 = mk_bin_scalar_reader('u32') read_bin_u64 = mk_bin_scalar_reader('u64') +read_bin_f16 = mk_bin_scalar_reader('f16') read_bin_f32 = mk_bin_scalar_reader('f32') read_bin_f64 = mk_bin_scalar_reader('f64') @@ -460,6 +479,13 @@             'bin_format': 'Q',             'numpy_type': np.uint64 }, +    'f16': {'binname' : b" f16",+            'size' : 2,+            'bin_reader': read_bin_f16,+            'str_reader': read_str_f16,+            'bin_format': 'e',+            'numpy_type': np.float16 },+     'f32': {'binname' : b" f32",             'size' : 4,             'bin_reader': read_bin_f32,@@ -600,6 +626,16 @@             out.write("true")         else:             out.write("false")+    elif type(v) == np.float16:+        if np.isnan(v):+            out.write('f16.nan')+        elif np.isinf(v):+            if v >= 0:+                out.write('f16.inf')+            else:+                out.write('-f16.inf')+        else:+            out.write("%.6ff16" % v)     elif type(v) == np.float32:         if np.isnan(v):             out.write('f32.nan')@@ -644,6 +680,7 @@               np.dtype('uint16'): b' u16',               np.dtype('uint32'): b' u32',               np.dtype('uint64'): b' u64',+              np.dtype('float16'): b' f16',               np.dtype('float32'): b' f32',               np.dtype('float64'): b' f64',               np.dtype('bool'): b'bool'}

@@ -1,18 +1,22 @@ {-# LANGUAGE FlexibleContexts #-}+{-# LANGUAGE OverloadedStrings #-}  -- | All (almost) compiler pipelines end with an 'Action', which does -- something with the result of the pipeline. module Futhark.Actions   ( printAction,     printAliasesAction,+    callGraphAction,     impCodeGenAction,     kernelImpCodeGenAction,     multicoreImpCodeGenAction,     metricsAction,     compileCAction,+    compileCtoWASMAction,     compileOpenCLAction,     compileCUDAAction,     compileMulticoreAction,+    compileMulticoreToWASMAction,     compilePythonAction,     compilePyOpenCLAction,   )@@ -20,15 +24,21 @@  import Control.Monad import Control.Monad.IO.Class+import Data.List (intercalate) import Data.Maybe (fromMaybe)+import qualified Data.Text as T+import qualified Data.Text.IO as T import Futhark.Analysis.Alias+import Futhark.Analysis.CallGraph (buildCallGraph) import Futhark.Analysis.Metrics import qualified Futhark.CodeGen.Backends.CCUDA as CCUDA import qualified Futhark.CodeGen.Backends.COpenCL as COpenCL import qualified Futhark.CodeGen.Backends.MulticoreC as MulticoreC+import qualified Futhark.CodeGen.Backends.MulticoreWASM as MulticoreWASM import qualified Futhark.CodeGen.Backends.PyOpenCL as PyOpenCL import qualified Futhark.CodeGen.Backends.SequentialC as SequentialC import qualified Futhark.CodeGen.Backends.SequentialPython as SequentialPy+import qualified Futhark.CodeGen.Backends.SequentialWASM as SequentialWASM import qualified Futhark.CodeGen.ImpGen.GPU as ImpGenGPU import qualified Futhark.CodeGen.ImpGen.Multicore as ImpGenMulticore import qualified Futhark.CodeGen.ImpGen.Sequential as ImpGenSequential@@ -37,6 +47,7 @@ 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.Util (runProgramWithExitCode, unixEnvironment) import Futhark.Version (versionString)@@ -63,6 +74,15 @@       actionProcedure = liftIO . putStrLn . pretty . aliasAnalysis     } +-- | Print call graph to stdout.+callGraphAction :: Action SOACS+callGraphAction =+  Action+    { actionName = "call-graph",+      actionDescription = "Prettyprint the callgraph of the result to standard output.",+      actionProcedure = liftIO . putStrLn . pretty . buildCallGraph+    }+ -- | Print metrics about AST node counts to stdout. metricsAction :: OpMetrics (Op rep) => Action rep metricsAction =@@ -100,20 +120,20 @@     }  -- Lines that we prepend (in comments) to generated code.-headerLines :: [String]-headerLines = lines $ "Generated by Futhark " ++ versionString+headerLines :: [T.Text]+headerLines = T.lines $ "Generated by Futhark " <> T.pack versionString -cHeaderLines :: [String]+cHeaderLines :: [T.Text] cHeaderLines = map ("// " <>) headerLines -pyHeaderLines :: [String]+pyHeaderLines :: [T.Text] pyHeaderLines = map ("# " <>) headerLines -cPrependHeader :: String -> String-cPrependHeader = (unlines cHeaderLines ++)+cPrependHeader :: T.Text -> T.Text+cPrependHeader = (T.unlines cHeaderLines <>) -pyPrependHeader :: String -> String-pyPrependHeader = (unlines pyHeaderLines ++)+pyPrependHeader :: T.Text -> T.Text+pyPrependHeader = (T.unlines pyHeaderLines <>)  cmdCC :: String cmdCC = fromMaybe "cc" $ lookup "CC" unixEnvironment@@ -163,13 +183,13 @@       case mode of         ToLibrary -> do           let (header, impl) = SequentialC.asLibrary cprog-          liftIO $ writeFile hpath $ cPrependHeader header-          liftIO $ writeFile cpath $ cPrependHeader impl+          liftIO $ T.writeFile hpath $ cPrependHeader header+          liftIO $ T.writeFile cpath $ cPrependHeader impl         ToExecutable -> do-          liftIO $ writeFile cpath $ SequentialC.asExecutable cprog+          liftIO $ T.writeFile cpath $ SequentialC.asExecutable cprog           runCC cpath outpath ["-O3", "-std=c99"] ["-lm"]         ToServer -> do-          liftIO $ writeFile cpath $ SequentialC.asServer cprog+          liftIO $ T.writeFile cpath $ SequentialC.asServer cprog           runCC cpath outpath ["-O3", "-std=c99"] ["-lm"]  -- | The @futhark opencl@ action.@@ -196,13 +216,13 @@       case mode of         ToLibrary -> do           let (header, impl) = COpenCL.asLibrary cprog-          liftIO $ writeFile hpath $ cPrependHeader header-          liftIO $ writeFile cpath $ cPrependHeader impl+          liftIO $ T.writeFile hpath $ cPrependHeader header+          liftIO $ T.writeFile cpath $ cPrependHeader impl         ToExecutable -> do-          liftIO $ writeFile cpath $ cPrependHeader $ COpenCL.asExecutable cprog+          liftIO $ T.writeFile cpath $ cPrependHeader $ COpenCL.asExecutable cprog           runCC cpath outpath ["-O", "-std=c99"] ("-lm" : extra_options)         ToServer -> do-          liftIO $ writeFile cpath $ cPrependHeader $ COpenCL.asServer cprog+          liftIO $ T.writeFile cpath $ cPrependHeader $ COpenCL.asServer cprog           runCC cpath outpath ["-O", "-std=c99"] ("-lm" : extra_options)  -- | The @futhark cuda@ action.@@ -226,13 +246,13 @@       case mode of         ToLibrary -> do           let (header, impl) = CCUDA.asLibrary cprog-          liftIO $ writeFile hpath $ cPrependHeader header-          liftIO $ writeFile cpath $ cPrependHeader impl+          liftIO $ T.writeFile hpath $ cPrependHeader header+          liftIO $ T.writeFile cpath $ cPrependHeader impl         ToExecutable -> do-          liftIO $ writeFile cpath $ cPrependHeader $ CCUDA.asExecutable cprog+          liftIO $ T.writeFile cpath $ cPrependHeader $ CCUDA.asExecutable cprog           runCC cpath outpath ["-O", "-std=c99"] ("-lm" : extra_options)         ToServer -> do-          liftIO $ writeFile cpath $ cPrependHeader $ CCUDA.asServer cprog+          liftIO $ T.writeFile cpath $ cPrependHeader $ CCUDA.asServer cprog           runCC cpath outpath ["-O", "-std=c99"] ("-lm" : extra_options)  -- | The @futhark multicore@ action.@@ -252,17 +272,17 @@       case mode of         ToLibrary -> do           let (header, impl) = MulticoreC.asLibrary cprog-          liftIO $ writeFile hpath $ cPrependHeader header-          liftIO $ writeFile cpath $ cPrependHeader impl+          liftIO $ T.writeFile hpath $ cPrependHeader header+          liftIO $ T.writeFile cpath $ cPrependHeader impl         ToExecutable -> do-          liftIO $ writeFile cpath $ cPrependHeader $ MulticoreC.asExecutable cprog-          runCC cpath outpath ["-O", "-std=c99"] ["-lm", "-pthread"]+          liftIO $ T.writeFile cpath $ cPrependHeader $ MulticoreC.asExecutable cprog+          runCC cpath outpath ["-O3", "-std=c99"] ["-lm", "-pthread"]         ToServer -> do-          liftIO $ writeFile cpath $ cPrependHeader $ MulticoreC.asServer cprog-          runCC cpath outpath ["-O", "-std=c99"] ["-lm", "-pthread"]+          liftIO $ T.writeFile cpath $ cPrependHeader $ MulticoreC.asServer cprog+          runCC cpath outpath ["-O3", "-std=c99"] ["-lm", "-pthread"]  pythonCommon ::-  (CompilerMode -> String -> prog -> FutharkM (Warnings, String)) ->+  (CompilerMode -> String -> prog -> FutharkM (Warnings, T.Text)) ->   FutharkConfig ->   CompilerMode ->   FilePath ->@@ -277,12 +297,13 @@    case mode of     ToLibrary ->-      liftIO $ writeFile (outpath `addExtension` "py") $ pyPrependHeader pyprog+      liftIO $ T.writeFile (outpath `addExtension` "py") $ pyPrependHeader pyprog     _ -> liftIO $ do-      writeFile outpath $ "#!/usr/bin/env python3\n" ++ pyPrependHeader pyprog+      T.writeFile outpath $ "#!/usr/bin/env python3\n" <> pyPrependHeader pyprog       perms <- liftIO $ getPermissions outpath       setPermissions outpath $ setOwnerExecutable True perms +-- | The @futhark python@ action. compilePythonAction :: FutharkConfig -> CompilerMode -> FilePath -> Action SeqMem compilePythonAction fcfg mode outpath =   Action@@ -291,6 +312,7 @@       actionProcedure = pythonCommon SequentialPy.compileProg fcfg mode outpath     } +-- | The @futhark pyopencl@ action. compilePyOpenCLAction :: FutharkConfig -> CompilerMode -> FilePath -> Action GPUMem compilePyOpenCLAction fcfg mode outpath =   Action@@ -298,3 +320,109 @@       actionDescription = "Compile to Python with OpenCL",       actionProcedure = pythonCommon PyOpenCL.compileProg fcfg mode outpath     }++cmdEMCFLAGS :: [String] -> [String]+cmdEMCFLAGS def = maybe def words $ lookup "EMCFLAGS" unixEnvironment++runEMCC :: String -> String -> FilePath -> [String] -> [String] -> [String] -> Bool -> FutharkM ()+runEMCC cpath outpath classpath cflags_def ldflags expfuns lib = do+  ret <-+    liftIO $+      runProgramWithExitCode+        "emcc"+        ( [cpath, "-o", outpath]+            ++ ["-lnodefs.js"]+            ++ ["-s", "--extern-post-js", classpath]+            ++ ( if lib+                   then ["-s", "EXPORT_NAME=loadWASM"]+                   else []+               )+            ++ ["-s", "WASM_BIGINT"]+            ++ cmdCFLAGS cflags_def+            ++ cmdEMCFLAGS [""]+            ++ [ "-s",+                 "EXPORTED_FUNCTIONS=["+                   ++ intercalate "," ("'_malloc'" : "'_free'" : expfuns)+                   ++ "]"+               ]+            -- The default LDFLAGS are always added.+            ++ ldflags+        )+        mempty+  case ret of+    Left err ->+      externalErrorS $ "Failed to run emcc: " ++ show err+    Right (ExitFailure code, _, emccerr) ->+      externalErrorS $+        "emcc failed with code "+          ++ show code+          ++ ":\n"+          ++ emccerr+    Right (ExitSuccess, _, _) ->+      return ()++-- | The @futhark wasm@ action.+compileCtoWASMAction :: FutharkConfig -> CompilerMode -> FilePath -> Action SeqMem+compileCtoWASMAction fcfg mode outpath =+  Action+    { actionName = "Compile to sequential C",+      actionDescription = "Compile to sequential C",+      actionProcedure = helper+    }+  where+    helper prog = do+      (cprog, jsprog, exps) <- handleWarnings fcfg $ SequentialWASM.compileProg prog+      case mode of+        ToLibrary -> do+          writeLibs cprog jsprog+          liftIO $ T.appendFile classpath SequentialWASM.libraryExports+          runEMCC cpath mjspath classpath ["-O3", "-msimd128"] ["-lm"] exps True+        _ -> do+          -- Non-server executables are not supported.+          writeLibs cprog jsprog+          liftIO $ T.appendFile classpath SequentialWASM.runServer+          runEMCC cpath outpath classpath ["-O3", "-msimd128"] ["-lm"] exps False+    writeLibs cprog jsprog = do+      let (h, imp) = SequentialC.asLibrary cprog+      liftIO $ T.writeFile hpath h+      liftIO $ T.writeFile cpath imp+      liftIO $ T.writeFile classpath jsprog++    cpath = outpath `addExtension` "c"+    hpath = outpath `addExtension` "h"+    mjspath = outpath `addExtension` "mjs"+    classpath = outpath `addExtension` ".class.js"++-- | The @futhark wasm-multicore@ action.+compileMulticoreToWASMAction :: FutharkConfig -> CompilerMode -> FilePath -> Action MCMem+compileMulticoreToWASMAction fcfg mode outpath =+  Action+    { actionName = "Compile to sequential C",+      actionDescription = "Compile to sequential C",+      actionProcedure = helper+    }+  where+    helper prog = do+      (cprog, jsprog, exps) <- handleWarnings fcfg $ MulticoreWASM.compileProg prog++      case mode of+        ToLibrary -> do+          writeLibs cprog jsprog+          liftIO $ T.appendFile classpath MulticoreWASM.libraryExports+          runEMCC cpath mjspath classpath ["-O3", "-msimd128"] ["-lm", "-pthread"] exps True+        _ -> do+          -- Non-server executables are not supported.+          writeLibs cprog jsprog+          liftIO $ T.appendFile classpath MulticoreWASM.runServer+          runEMCC cpath outpath classpath ["-O3", "-msimd128"] ["-lm", "-pthread"] exps False++    writeLibs cprog jsprog = do+      let (h, imp) = MulticoreC.asLibrary cprog+      liftIO $ T.writeFile hpath h+      liftIO $ T.writeFile cpath imp+      liftIO $ T.writeFile classpath jsprog++    cpath = outpath `addExtension` "c"+    hpath = outpath `addExtension` "h"+    mjspath = outpath `addExtension` "mjs"+    classpath = outpath `addExtension` ".class.js"

@@ -72,7 +72,7 @@   Stm (Aliases rep) analyseStm aliases (Let pat (StmAux cs attrs dec) e) =   let e' = analyseExp aliases e-      pat' = addAliasesToPattern pat e'+      pat' = addAliasesToPat pat e'       rep' = (AliasDec $ consumedInExp e', dec)    in Let pat' (StmAux cs attrs rep') e' 

@@ -9,6 +9,7 @@     calls,     calledByConsts,     allCalledBy,+    numOccurences,     findNoninlined,   ) where@@ -16,9 +17,10 @@ import Control.Monad.Writer.Strict import Data.List (foldl') import qualified Data.Map.Strict as M-import Data.Maybe (fromMaybe)+import Data.Maybe (isJust) import qualified Data.Set as S import Futhark.IR.SOACS+import Futhark.Util.Pretty  type FunctionTable = M.Map Name (FunDef SOACS) @@ -27,45 +29,78 @@   where     expand ftab f = M.insert (funDefName f) f ftab -type FunGraph = M.Map Name (S.Set Name)+-- | A unique (at least within a function) name identifying a function+-- call.  In practice the first element of the corresponding pattern.+type CallId = VName +data FunCalls = FunCalls+  { fcMap :: M.Map CallId (Attrs, Name),+    fcAllCalled :: S.Set Name+  }+  deriving (Eq, Ord, Show)++instance Monoid FunCalls where+  mempty = FunCalls mempty mempty++instance Semigroup FunCalls where+  FunCalls x1 y1 <> FunCalls x2 y2 = FunCalls (x1 <> x2) (y1 <> y2)++fcCalled :: Name -> FunCalls -> Bool+fcCalled f fcs = f `S.member` fcAllCalled fcs++type FunGraph = M.Map Name FunCalls+ -- | The call graph is a mapping from a function name, i.e., the--- caller, to a set of the names of functions called *directly* (not+-- caller, to a record of the names of functions called *directly* (not -- transitively!) by the function. -- -- We keep track separately of the functions called by constants. data CallGraph = CallGraph-  { calledByFuns :: M.Map Name (S.Set Name),-    calledInConsts :: S.Set Name+  { cgCalledByFuns :: FunGraph,+    cgCalledByConsts :: FunCalls   }+  deriving (Eq, Ord, Show)  -- | Is the given function known to the call graph? isFunInCallGraph :: Name -> CallGraph -> Bool-isFunInCallGraph f = M.member f . calledByFuns+isFunInCallGraph f = M.member f . cgCalledByFuns  -- | Does the first function call the second? calls :: Name -> Name -> CallGraph -> Bool calls caller callee =-  maybe False (S.member callee) . M.lookup caller . calledByFuns+  maybe False (fcCalled callee) . M.lookup caller . cgCalledByFuns  -- | Is the function called in any of the constants? calledByConsts :: Name -> CallGraph -> Bool-calledByConsts f = S.member f . calledInConsts+calledByConsts callee = fcCalled callee . cgCalledByConsts  -- | All functions called by this function. allCalledBy :: Name -> CallGraph -> S.Set Name-allCalledBy f = fromMaybe mempty . M.lookup f . calledByFuns+allCalledBy f = maybe mempty fcAllCalled . M.lookup f . cgCalledByFuns  -- | @buildCallGraph prog@ build the program's call graph. buildCallGraph :: Prog SOACS -> CallGraph buildCallGraph prog =-  CallGraph fg $ buildFGStms $ progConsts prog+  CallGraph fg cg   where     fg = foldl' (buildFGfun ftable) M.empty entry_points+    cg = buildFGStms $ progConsts prog -    entry_points = map funDefName $ progFuns prog+    entry_points =+      S.fromList (map funDefName (filter (isJust . funDefEntryPoint) $ progFuns prog))+        <> fcAllCalled cg     ftable = buildFunctionTable prog +count :: Ord k => [k] -> M.Map k Int+count ks = M.fromListWith (+) $ zip ks $ repeat 1++-- | Produce a mapping of the number of occurences in the call graph+-- of each function.  Only counts functions that are called at least+-- once.+numOccurences :: CallGraph -> M.Map Name Int+numOccurences (CallGraph funs consts) =+  count $ map snd $ M.elems (fcMap consts <> foldMap fcMap (M.elems funs))+ -- | @buildCallGraph ftable fg fname@ updates @fg@ with the -- contributions of function @fname@. buildFGfun :: FunctionTable -> FunGraph -> Name -> FunGraph@@ -77,18 +112,19 @@       let callees = buildFGBody $ funDefBody f           fg' = M.insert fname callees fg       -- recursively build the callees-      foldl' (buildFGfun ftable) fg' callees+      foldl' (buildFGfun ftable) fg' $ fcAllCalled callees     _ -> fg -buildFGStms :: Stms SOACS -> S.Set Name-buildFGStms = mconcat . map (buildFGexp . stmExp) . stmsToList+buildFGStms :: Stms SOACS -> FunCalls+buildFGStms = mconcat . map buildFGstm . stmsToList -buildFGBody :: Body -> S.Set Name+buildFGBody :: Body -> FunCalls buildFGBody = buildFGStms . bodyStms -buildFGexp :: Exp -> S.Set Name-buildFGexp (Apply fname _ _ _) = S.singleton fname-buildFGexp (Op op) = execWriter $ mapSOACM folder op+buildFGstm :: Stm -> FunCalls+buildFGstm (Let (Pat (p : _)) aux (Apply fname _ _ _)) =+  FunCalls (M.singleton (patElemName p) (stmAuxAttrs aux, fname)) (S.singleton fname)+buildFGstm (Let _ _ (Op op)) = execWriter $ mapSOACM folder op   where     folder =       identitySOACMapper@@ -96,7 +132,7 @@             tell $ buildFGBody $ lambdaBody lam             return lam         }-buildFGexp e = execWriter $ mapExpM folder e+buildFGstm (Let _ _ e) = execWriter $ mapExpM folder e   where     folder =       identityMapper@@ -137,3 +173,18 @@         S.singleton $ funDefName fd       | otherwise =         mempty++instance Pretty FunCalls where+  ppr = stack . map f . M.toList . fcMap+    where+      f (x, (attrs, y)) = "=>" <+> ppr y <+> parens ("at" <+> ppr x <+> ppr attrs)++instance Pretty CallGraph where+  ppr (CallGraph fg cg) =+    stack $+      punctuate line $+        ppFunCalls ("called at top level", cg) : map ppFunCalls (M.toList fg)+    where+      ppFunCalls (f, fcalls) =+        ppr f </> text (map (const '=') (nameToString f))+          </> indent 2 (ppr fcalls)

@@ -32,7 +32,7 @@       let tdeps = dataDependencies' deps tb           fdeps = dataDependencies' deps fb           cdeps = depsOf deps c-          comb (pe, tres, fres) =+          comb (pe, SubExpRes _ tres, SubExpRes _ fres) =             ( patElemName pe,               mconcat $                 [freeIn pe, cdeps, depsOf tdeps tres, depsOf fdeps fres]@@ -42,14 +42,14 @@             M.fromList $               map comb $                 zip3-                  (patternElements pat)+                  (patElems pat)                   (bodyResult tb)                   (bodyResult fb)        in M.unions [branchdeps, deps, tdeps, fdeps]     grow deps (Let pat _ e) =       let free = freeIn pat <> freeIn e           freeDeps = mconcat $ map (depsOfVar deps) $ namesToList free-       in M.fromList [(name, freeDeps) | name <- patternNames pat] `M.union` deps+       in M.fromList [(name, freeDeps) | name <- patNames pat] `M.union` deps  depsOf :: Dependencies -> SubExp -> Names depsOf _ (Constant _) = mempty