futhark 0.25.5 → 0.25.6

raw patch · 120 files changed

+1697/−1550 lines, 120 filesPVP: major bump suggested

API removals or changes: PVP suggests a major version bump

API changes (from Hackage documentation)

Files

• docs/c-api.rst +7/−2
• docs/conf.py +12/−3
• docs/index.rst +1/−0
• docs/man/futhark-bench.rst +7/−1
• docs/man/futhark-c.rst +19/−0
• docs/man/futhark-cuda.rst +0/−19
• docs/man/futhark-hip.rst +0/−23
• docs/man/futhark-multicore.rst +0/−4
• docs/man/futhark-opencl.rst +0/−29
• docs/man/futhark-profile.rst +95/−0
• futhark.cabal +5/−1
• rts/c/backends/c.h +2/−1
• rts/c/backends/cuda.h +128/−115
• rts/c/backends/hip.h +119/−113
• rts/c/backends/multicore.h +28/−1
• rts/c/backends/opencl.h +109/−113
• rts/c/context.h +15/−4
• rts/c/context_prototypes.h +6/−0
• rts/c/event_list.h +62/−0
• rts/c/util.h +30/−63
• rts/python/memory.py +12/−27
• rts/python/opencl.py +2/−2
• rts/python/scalar.py +3/−3
• rts/python/values.py +3/−3
• src/Futhark/AD/Fwd.hs +16/−17
• src/Futhark/AD/Rev.hs +6/−6
• src/Futhark/AD/Rev/Hist.hs +1/−1
• src/Futhark/AD/Rev/Loop.hs +50/−100
• src/Futhark/AD/Rev/Monad.hs +1/−1
• src/Futhark/AD/Rev/SOAC.hs +1/−1
• src/Futhark/Analysis/DataDependencies.hs +71/−2
• src/Futhark/Analysis/HORep/SOAC.hs +49/−33
• src/Futhark/Analysis/Interference.hs +2/−4
• src/Futhark/Analysis/LastUse.hs +3/−3
• src/Futhark/Bench.hs +57/−15
• src/Futhark/Builder.hs +1/−1
• src/Futhark/CLI/Autotune.hs +3/−3
• src/Futhark/CLI/Bench.hs +28/−19
• src/Futhark/CLI/Main.hs +3/−1
• src/Futhark/CLI/Profile.hs +136/−0
• src/Futhark/CodeGen/Backends/CCUDA.hs +0/−7
• src/Futhark/CodeGen/Backends/COpenCL.hs +0/−7
• src/Futhark/CodeGen/Backends/GPU.hs +0/−22
• src/Futhark/CodeGen/Backends/GenericC.hs +14/−15
• src/Futhark/CodeGen/Backends/GenericC/CLI.hs +9/−13
• src/Futhark/CodeGen/Backends/GenericC/Monad.hs +0/−7
• src/Futhark/CodeGen/Backends/GenericC/Server.hs +7/−0
• src/Futhark/CodeGen/Backends/GenericPython.hs +16/−32
• src/Futhark/CodeGen/Backends/GenericPython/AST.hs +0/−7
• src/Futhark/CodeGen/Backends/HIP.hs +0/−7
• src/Futhark/CodeGen/Backends/MulticoreC.hs +20/−122
• src/Futhark/CodeGen/Backends/MulticoreISPC.hs +2/−7
• src/Futhark/CodeGen/Backends/PyOpenCL.hs +0/−50
• src/Futhark/CodeGen/Backends/SequentialPython.hs +1/−8
• src/Futhark/CodeGen/ImpCode.hs +2/−2
• src/Futhark/CodeGen/ImpGen.hs +2/−11
• src/Futhark/CodeGen/ImpGen/GPU/SegHist.hs +8/−4
• src/Futhark/CodeGen/RTS/C.hs +6/−0
• src/Futhark/Construct.hs +1/−1
• src/Futhark/IR/GPU/Op.hs +7/−0
• src/Futhark/IR/MC/Op.hs +3/−0
• src/Futhark/IR/Mem.hs +2/−0
• src/Futhark/IR/Parse.hs +3/−6
• src/Futhark/IR/Pretty.hs +3/−14
• src/Futhark/IR/Prop.hs +4/−0
• src/Futhark/IR/Prop/Aliases.hs +1/−9
• src/Futhark/IR/Prop/Names.hs +5/−4
• src/Futhark/IR/Prop/Reshape.hs +3/−2
• src/Futhark/IR/Prop/Scope.hs +5/−4
• src/Futhark/IR/SOACS/SOAC.hs +69/−5
• src/Futhark/IR/SOACS/Simplify.hs +41/−29
• src/Futhark/IR/SegOp.hs +2/−0
• src/Futhark/IR/Syntax.hs +8/−14
• src/Futhark/IR/Syntax/Core.hs +5/−0
• src/Futhark/IR/Traversals.hs +15/−25
• src/Futhark/IR/TypeCheck.hs +4/−27
• src/Futhark/Internalise/Exps.hs +10/−10
• src/Futhark/Internalise/Lambdas.hs +1/−1
• src/Futhark/Optimise/ArrayShortCircuiting/ArrayCoalescing.hs +3/−3
• src/Futhark/Optimise/ArrayShortCircuiting/TopdownAnalysis.hs +4/−4
• src/Futhark/Optimise/DoubleBuffer.hs +1/−1
• src/Futhark/Optimise/Fusion.hs +27/−1
• src/Futhark/Optimise/Fusion/GraphRep.hs +6/−6
• src/Futhark/Optimise/Fusion/TryFusion.hs +52/−4
• src/Futhark/Optimise/GenRedOpt.hs +1/−1
• src/Futhark/Optimise/InPlaceLowering.hs +1/−1
• src/Futhark/Optimise/InPlaceLowering/LowerIntoStm.hs +1/−1
• src/Futhark/Optimise/InliningDeadFun.hs +2/−2
• src/Futhark/Optimise/MergeGPUBodies.hs +8/−4
• src/Futhark/Optimise/ReduceDeviceSyncs.hs +2/−28
• src/Futhark/Optimise/ReduceDeviceSyncs/MigrationTable.hs +8/−17
• src/Futhark/Optimise/Simplify/Engine.hs +8/−21
• src/Futhark/Optimise/Simplify/Rep.hs +2/−5
• src/Futhark/Optimise/Simplify/Rule.hs +1/−1
• src/Futhark/Optimise/Simplify/Rules.hs +4/−4
• src/Futhark/Optimise/Simplify/Rules/BasicOp.hs +1/−2
• src/Futhark/Optimise/Simplify/Rules/Index.hs +4/−2
• src/Futhark/Optimise/Simplify/Rules/Loop.hs +8/−82
• src/Futhark/Optimise/Sink.hs +7/−29
• src/Futhark/Optimise/TileLoops.hs +13/−14
• src/Futhark/Optimise/TileLoops/Shared.hs +2/−2
• src/Futhark/Pass/AD.hs +1/−1
• src/Futhark/Pass/ExpandAllocations.hs +6/−7
• src/Futhark/Pass/ExplicitAllocations.hs +2/−25
• src/Futhark/Pass/ExtractKernels.hs +4/−10
• src/Futhark/Pass/ExtractKernels/ISRWIM.hs +4/−4
• src/Futhark/Pass/ExtractKernels/Interchange.hs +7/−7
• src/Futhark/Pass/ExtractKernels/Intragroup.hs +2/−6
• src/Futhark/Pass/ExtractMulticore.hs +4/−10
• src/Futhark/Profile.hs +65/−0
• src/Futhark/Transform/FirstOrderTransform.hs +9/−9
• src/Futhark/Transform/Rename.hs +6/−12
• src/Futhark/Transform/Substitute.hs +2/−2
• src/Futhark/Util/Pretty.hs +1/−1
• src/Language/Futhark/Parser/Lexer.x +2/−0
• src/Language/Futhark/Parser/Lexer/Tokens.hs +1/−0
• src/Language/Futhark/Parser/Monad.hs +3/−0
• src/Language/Futhark/Primitive.hs +14/−14
• unittests/Futhark/BenchTests.hs +3/−1
• unittests/Futhark/ProfileTests.hs +23/−0

@@ -22,6 +22,11 @@ ``NULL`` pointer.  Use :c:func:`futhark_context_get_error` to get a (possibly) more precise error message. +Some functions take a C string (``const char*``) as argument.  Unless+otherwise indicated, the string will be copied if necessary, meaning+the argument string can always be modified (or freed) after the+function returns.+ .. c:macro:: FUTHARK_BACKEND_foo     A preprocessor macro identifying that the backend *foo* was used to@@ -32,7 +37,7 @@ Error codes ----------- -Most errors are result in a not otherwise specified nonzero return+Most errors result in a not otherwise specified nonzero return code, but a few classes of errors have distinct error codes.  .. c:macro:: FUTHARK_SUCCESS@@ -202,7 +207,7 @@  .. c:function:: char *futhark_context_report(struct futhark_context *ctx) -   Produce a human-readable C string with debug and profiling+   Produce a C string encoding a JSON object with debug and profiling    information collected during program runtime.  It is the caller's    responsibility to free the returned string.  It is likely to only    contain interesting information if

@@ -31,6 +31,7 @@     Whitespace, ) from sphinx.highlighting import lexers+from typing import Dict, Any, List, Tuple  # If extensions (or modules to document with autodoc) are in another directory, # add these directories to sys.path here. If the directory is relative to the@@ -292,7 +293,7 @@ # Theme options are theme-specific and customize the look and feel of a theme # further.  For a list of options available for each theme, see the # documentation.-html_theme_options = {}+html_theme_options: Dict[Any, Any] = {}  # Add any paths that contain custom themes here, relative to this directory. # html_theme_path = []@@ -377,7 +378,7 @@  # -- Options for LaTeX output --------------------------------------------- -latex_elements = {+latex_elements: Dict[str, str] = {     # The paper size ('letterpaper' or 'a4paper').     #'papersize': 'letterpaper',     # The font size ('10pt', '11pt' or '12pt').@@ -418,7 +419,7 @@  # One entry per manual page. List of tuples # (source start file, name, description, authors, manual section).-man_pages = [+man_pages: List[Tuple[str, str, str, List[Any], int]] = [     ("man/futhark", "futhark", "a parallel functional array language", [], 1),     (         "man/futhark-autotune",@@ -450,6 +451,7 @@         1,     ),     ("man/futhark-cuda", "futhark-cuda", "compile Futhark to CUDA", [], 1),+    ("man/futhark-hip", "futhark-hip", "compile Futhark to HIP", [], 1),     (         "man/futhark-python",         "futhark-python",@@ -513,6 +515,13 @@         "man/futhark-literate",         "futhark-literate",         "execute literate Futhark program",+        [],+        1,+    ),+    (+        "man/futhark-profile",+        "futhark-profile",+        "profile Futhark programs",         [],         1,     ),

@@ -57,6 +57,7 @@    man/futhark-multicore.rst    man/futhark-opencl.rst    man/futhark-pkg.rst+   man/futhark-profile.rst    man/futhark-pyopencl.rst    man/futhark-python.rst    man/futhark-repl.rst

@@ -120,6 +120,12 @@    Pass an extra option to the compiler when compiling the programs. +--profile++  Enable profiling for the binary (by passing ``--profiling`` and+  ``--logging``) and store the recorded information in the file+  indicated by ``--json``, along with the other benchmarking results .+ --runner=program    If set to a non-empty string, compiled programs are not run@@ -185,7 +191,7 @@  -- output { 49995000i64 }  -- compiled input { 1000000i64 }  -- output { 499999500000i64 }- +  let main(n: i64): i64 =    reduce (+) 0 (iota n) 

@@ -116,6 +116,22 @@   Do not print the program result.  Not accepted by server-mode   executables. +-P, --profile++  Gather profiling data during execution.  Mostly interesting in+  ``--server`` mode.  Implied by ``-D``.++--param=ASSIGNMENT++  Set a tuning parameter to the given+  value. ``ASSIGNMENT`` must be of the form ``NAME=INT`` Use+  ``--print-params`` to see which names are available.++--print-params++  Print all tuning parameters that can be set with ``--param`` or+  ``--tuning``.+ -r, --runs=NUM    Perform NUM runs of the program.  With ``-t``, the runtime for each@@ -130,6 +146,9 @@   an integral number of microseconds.  Not accepted by server-mode   executables. +--tuning=FILE++  Read size=value assignments from the given file.  SEE ALSO ========

@@ -101,30 +101,11 @@    Load PTX code from the indicated file. --n, --no-print-result--  Do not print the program result.- --nvrtc-option=OPT    Add an additional build option to the string passed to NVRTC.  Refer   to the CUDA documentation for which options are supported.  Be   careful - some options can easily result in invalid results.----param=ASSIGNMENT--  Set a tuning parameter to the given-  value. ``ASSIGNMENT`` must be of the form ``NAME=INT`` Use-  ``--print-params`` to see which names are available.----print-params--  Print all tuning parameters that can be set with ``--param`` or-  ``--tuning``.----tuning=FILE--  Read size=value assignments from the given file.  ENVIRONMENT ===========

@@ -54,10 +54,6 @@  The following additional options are accepted. --h, --help--  Print help text to standard output and exit.- --default-group-size=INT    The default size of thread blocks that are launched.  Capped to the@@ -89,31 +85,12 @@   Instead of using the embedded HIP kernels, load them from the   indicated file. --n, --no-print-result--  Do not print the program result.- --build-option=OPT    Add an additional build option to the string passed to the kernel   compiler (HIPRTC).  Refer to the HIP documentation for which options   are supported.  Be careful - some options can easily result in   invalid results.----param=ASSIGNMENT--  Set a tuning parameter to the given-  value. ``ASSIGNMENT`` must be of the form ``NAME=INT`` Use-  ``--print-params`` to see which names are available.----print-params--  Print all tuning parameters that can be set with ``--param`` or-  ``--tuning``.----tuning=FILE--  Read size=value assignments from the given file.  ENVIRONMENT ===========

@@ -49,10 +49,6 @@ Generated executables accept the same options as those generated by :ref:`futhark-c(1)`.  The following additional options are accepted. --P, --profile--  Collect and report various bits of profiling information.- --num-threads=INT    Use this many physical threads.

@@ -57,10 +57,6 @@  The following additional options are accepted. --h, --help--  Print help text to standard output and exit.- --build-option=OPT    Add an additional build option to the string passed to@@ -117,36 +113,11 @@    Load an OpenCL binary from the indicated file. --n, --no-print-result--  Do not print the program result.- -p, --platform=NAME    Use the first OpenCL platform whose name contains the given string.   The special string ``#k``, where ``k`` is an integer, can be used to   pick the *k*-th platform, numbered from zero.---P, --profile--  Gather profiling data while executing and print out a summary at the-  end.  When ``-r`` is used, only the last run will be profiled.-  Implied by ``-D``.----param=ASSIGNMENT--  Set a tuning parameter to the given-  value. ``ASSIGNMENT`` must be of the form ``NAME=INT`` Use-  ``--print-params`` to see which names are available.----print-params--  Print all tuning parameters that can be set with ``--param`` or-  ``--tuning``.----tuning=FILE--  Read size=value assignments from the given file.  --list-devices 

@@ -0,0 +1,95 @@+.. role:: ref(emphasis)++.. _futhark-profile(1):++===============+futhark-profile+===============++SYNOPSIS+========++futhark profile JSONFILE++DESCRIPTION+===========++This tool produces human-readable profiling information based on+information collected with :ref:`futhark bench<futhark-bench(1)>`.+Futhark has only rudimentary support for profiling.  While the system+can collect information about the run-time behaviour of the program,+there is currently no automatic way to connect the information to the+program source code.  However, the collected information can still be+useful for estimating the source of inefficiencies.++USAGE+=====++The first step is to run :ref:`futhark bench<futhark-bench(1)>` on+your program, while passing ``--profile`` and ``--json``.  This will+produce a JSON file containing runtime measurements, as well as+collected profiling information.  If you neglect to pass+``--profile``, the latter will be missing.  If you neglect to pass+``--json``, no file will be created.  The information in the JSON file+is complete, but it is difficult for humans to read.++The next step is to run ``futhark profile`` on the JSON file.  For a+JSON file ``prog.json``, this will create a *top level directory*+``prog.prof`` that contains files with human-readable profiling+information.  A set of files will be created for each benchmark+dataset.  If the original invocation of ``futhark bench`` included+multiple programs, then ``futhark profile`` will create subdirectories+for each program (although all inside the same top level directory).++Files produced+--------------++Supposing a dataset ``foo``, ``futhark profile`` will produce the+following files in the top level directory.++* ``foo.log``: the running log produced during execution.  Contains+  many details on dynamic behaviour, depending on the exact backend.++* ``foo.summary``: a summary of memory usage and cost centres.  For+  the GPU backends, the cost centres are kernel executions and memory+  copies.++Technicalities+--------------++The profiling information, including the log, is collected from a+*final* run performed after all the measured runs.  Profiling+information is not collected during the runs that contribute to the+runtime measurement reported by ``futhark bench``.  However, enabling+profiling may still affect performance, as it changes the+behaviour of the run time system.++EXAMPLES+========++This shows the sequence of commands one might use to profile the+program ``LocVolCalib.fut``, which has three datasets associated with+it, using the ``hip`` backend::++ $ futhark bench --backend=hip --profile --json result.json LocVolCalib.fut+ $ futhark profile result.json+ $ tree result.prof/+ result.prof/+ ├── LocVolCalib-data_large.in.log+ ├── LocVolCalib-data_large.in.summary+ ├── LocVolCalib-data_medium.in.log+ ├── LocVolCalib-data_medium.in.summary+ ├── LocVolCalib-data_small.in.log+ └── LocVolCalib-data_small.in.summary++BUGS+====++Only the C-based backends currently support profiling.++The ``c`` backend does not actually record useful profiling information.++SEE ALSO+========++:ref:`futhark-bench(1)`

@@ -1,6 +1,6 @@ cabal-version: 2.4 name:           futhark-version:        0.25.5+version:        0.25.6 synopsis:       An optimising compiler for a functional, array-oriented language.  description:    Futhark is a small programming language designed to be compiled to@@ -49,6 +49,7 @@     rts/c/timing.h     rts/c/errors.h     rts/c/free_list.h+    rts/c/event_list.h     rts/c/gpu.h     rts/c/gpu_prototypes.h     rts/c/tuning.h@@ -209,6 +210,7 @@       Futhark.CLI.MulticoreWASM       Futhark.CLI.OpenCL       Futhark.CLI.Pkg+      Futhark.CLI.Profile       Futhark.CLI.PyOpenCL       Futhark.CLI.Python       Futhark.CLI.Query@@ -415,6 +417,7 @@       Futhark.Pkg.Info       Futhark.Pkg.Solve       Futhark.Pkg.Types+      Futhark.Profile       Futhark.Script       Futhark.Test       Futhark.Test.Spec@@ -545,6 +548,7 @@   other-modules:       Futhark.AD.DerivativesTests       Futhark.BenchTests+      Futhark.ProfileTests       Futhark.Pkg.SolveTests       Futhark.Analysis.AlgSimplifyTests       Futhark.Internalise.TypesValuesTests

@@ -5,7 +5,7 @@   int debugging;   int profiling;   int logging;-  const char *cache_fname;+  char *cache_fname;   int num_tuning_params;   int64_t *tuning_params;   const char** tuning_param_names;@@ -39,6 +39,7 @@   FILE *log;   struct constants *constants;   struct free_list free_list;+  struct event_list event_list;   int64_t peak_mem_usage_default;   int64_t cur_mem_usage_default;   struct program* program;

@@ -79,7 +79,7 @@   int debugging;   int profiling;   int logging;-  const char *cache_fname;+  char* cache_fname;   int num_tuning_params;   int64_t *tuning_params;   const char** tuning_param_names;@@ -89,13 +89,13 @@    char* program;   int num_nvrtc_opts;-  const char **nvrtc_opts;+  char* *nvrtc_opts; -  const char *preferred_device;+  char* preferred_device;   int preferred_device_num; -  const char *dump_ptx_to;-  const char *load_ptx_from;+  char* dump_ptx_to;+  char* load_ptx_from;    size_t default_block_size;   size_t default_grid_size;@@ -110,13 +110,13 @@  static void backend_context_config_setup(struct futhark_context_config *cfg) {   cfg->num_nvrtc_opts = 0;-  cfg->nvrtc_opts = (const char**) malloc(sizeof(const char*));+  cfg->nvrtc_opts = (char**) malloc(sizeof(char*));   cfg->nvrtc_opts[0] = NULL;    cfg->program = strconcat(gpu_program);    cfg->preferred_device_num = 0;-  cfg->preferred_device = "";+  cfg->preferred_device = strdup("");    cfg->dump_ptx_to = NULL;   cfg->load_ptx_from = NULL;@@ -133,13 +133,20 @@ }  static void backend_context_config_teardown(struct futhark_context_config* cfg) {+  for (int i = 0; i < cfg->num_nvrtc_opts; i++) {+    free(cfg->nvrtc_opts[i]);+  }   free(cfg->nvrtc_opts);+  free(cfg->dump_ptx_to);+  free(cfg->load_ptx_from);+  free(cfg->preferred_device);+  free(cfg->program); }  void futhark_context_config_add_nvrtc_option(struct futhark_context_config *cfg, const char *opt) {-  cfg->nvrtc_opts[cfg->num_nvrtc_opts] = opt;+  cfg->nvrtc_opts[cfg->num_nvrtc_opts] = strdup(opt);   cfg->num_nvrtc_opts++;-  cfg->nvrtc_opts = (const char **) realloc(cfg->nvrtc_opts, (cfg->num_nvrtc_opts + 1) * sizeof(const char *));+  cfg->nvrtc_opts = (char **) realloc(cfg->nvrtc_opts, (cfg->num_nvrtc_opts + 1) * sizeof(char *));   cfg->nvrtc_opts[cfg->num_nvrtc_opts] = NULL; } @@ -155,7 +162,8 @@       s++;     }   }-  cfg->preferred_device = s;+  free(cfg->preferred_device);+  cfg->preferred_device = strdup(s);   cfg->preferred_device_num = x; } @@ -164,15 +172,18 @@ }  void futhark_context_config_set_program(struct futhark_context_config *cfg, const char *s) {+  free(cfg->program);   cfg->program = strdup(s); }  void futhark_context_config_dump_ptx_to(struct futhark_context_config *cfg, const char *path) {-  cfg->dump_ptx_to = path;+  free(cfg->dump_ptx_to);+  cfg->dump_ptx_to = strdup(path); }  void futhark_context_config_load_ptx_from(struct futhark_context_config *cfg, const char *path) {-  cfg->load_ptx_from = path;+  free(cfg->load_ptx_from);+  cfg->load_ptx_from = strdup(path); }  void futhark_context_config_set_default_group_size(struct futhark_context_config *cfg, int size) {@@ -249,6 +260,7 @@   FILE *log;   struct constants *constants;   struct free_list free_list;+  struct event_list event_list;   int64_t peak_mem_usage_default;   int64_t cur_mem_usage_default;   // Uniform fields above.@@ -280,10 +292,6 @@    size_t lockstep_width; -  struct profiling_record *profiling_records;-  int profiling_records_capacity;-  int profiling_records_used;-   struct builtin_kernels* kernels; }; @@ -696,60 +704,43 @@   return NULL; } -// Count up the runtime all the profiling_records that occured during execution.-// Also clears the buffer of profiling_records.-static CUresult tally_profiling_records(struct futhark_context *ctx,-                                        struct cost_centres* ccs) {-  CUresult err;-  for (int i = 0; i < ctx->profiling_records_used; i++) {-    struct profiling_record record = ctx->profiling_records[i];+struct cuda_event {+  cudaEvent_t start;+  cudaEvent_t end;+}; -    float ms;-    if ((err = cuEventElapsedTime(&ms, record.events[0], record.events[1])) != CUDA_SUCCESS) {-      return err;-    }+static struct cuda_event* cuda_event_new(struct futhark_context* ctx) {+  if (ctx->profiling && !ctx->profiling_paused) {+    struct cuda_event* e = malloc(sizeof(struct cuda_event));+    cudaEventCreate(&e->start);+    cudaEventCreate(&e->end);+    return e;+  } else {+    return NULL;+  }+} -    if (ccs) {-      // CUDA provides milisecond resolution, but we want microseconds.-      struct cost_centre c = {-        .name = record.name,-        .runs = 1,-        .runtime = ms*1000-      };-      cost_centres_add(ccs, c);-    }+static int cuda_event_report(struct str_builder* sb, struct cuda_event* e) {+  float ms;+  CUresult err;+  if ((err = cuEventElapsedTime(&ms, e->start, e->end)) != CUDA_SUCCESS) {+    return err;+  } -    if ((err = cuEventDestroy(record.events[0])) != CUDA_SUCCESS) {-      return err;-    }-    if ((err = cuEventDestroy(record.events[1])) != CUDA_SUCCESS) {-      return err;-    }+  // CUDA provides milisecond resolution, but we want microseconds.+  str_builder(sb, ",\"duration\":%f", ms*1000); -    free(record.events);+  if ((err = cuEventDestroy(e->start)) != CUDA_SUCCESS) {+    return 1;   } -  ctx->profiling_records_used = 0;+  if ((err = cuEventDestroy(e->end)) != CUDA_SUCCESS) {+    return 1;+  } -  return CUDA_SUCCESS;-}+  free(e); -// Returns pointer to two events.-static cudaEvent_t* cuda_get_events(struct futhark_context *ctx, const char* name) {-  if (ctx->profiling_records_used == ctx->profiling_records_capacity) {-    ctx->profiling_records_capacity *= 2;-    ctx->profiling_records =-      realloc(ctx->profiling_records,-              ctx->profiling_records_capacity *-              sizeof(struct profiling_record));-  }-  cudaEvent_t *events = calloc(2, sizeof(cudaEvent_t));-  cudaEventCreate(&events[0]);-  cudaEventCreate(&events[1]);-  ctx->profiling_records[ctx->profiling_records_used].events = events;-  ctx->profiling_records[ctx->profiling_records_used].name = name;-  ctx->profiling_records_used++;-  return events;+  return 0; }  int futhark_context_sync(struct futhark_context* ctx) {@@ -793,11 +784,6 @@ void free_builtin_kernels(struct futhark_context* ctx, struct builtin_kernels* kernels);  int backend_context_setup(struct futhark_context* ctx) {-  ctx->profiling_records_capacity = 200;-  ctx->profiling_records_used = 0;-  ctx->profiling_records =-    malloc(ctx->profiling_records_capacity *-           sizeof(struct profiling_record));   ctx->failure_is_an_option = 0;   ctx->total_runs = 0;   ctx->total_runtime = 0;@@ -821,8 +807,10 @@   ctx->lockstep_width = device_query(ctx->dev, WARP_SIZE);   CUDA_SUCCEED_FATAL(cuStreamCreate(&ctx->stream, CU_STREAM_DEFAULT));   cuda_size_setup(ctx);-  ctx->error = cuda_module_setup(ctx, ctx->cfg->program,-                                 ctx->cfg->nvrtc_opts, ctx->cfg->cache_fname);+  ctx->error = cuda_module_setup(ctx,+                                 ctx->cfg->program,+                                 (const char**)ctx->cfg->nvrtc_opts,+                                 ctx->cfg->cache_fname);    if (ctx->error != NULL) {     futhark_panic(1, "During CUDA initialisation:\n%s\n", ctx->error);@@ -846,8 +834,6 @@   cuMemFree(ctx->global_failure);   cuMemFree(ctx->global_failure_args);   CUDA_SUCCEED_FATAL(gpu_free_all(ctx));-  (void)tally_profiling_records(ctx, NULL);-  free(ctx->profiling_records);   CUDA_SUCCEED_FATAL(cuStreamDestroy(ctx->stream));   CUDA_SUCCEED_FATAL(cuModuleUnload(ctx->module));   CUDA_SUCCEED_FATAL(cuCtxDestroy(ctx->cu_ctx));@@ -878,14 +864,18 @@ static int gpu_scalar_to_device(struct futhark_context* ctx,                                 gpu_mem dst, size_t offset, size_t size,                                 void *src) {-  CUevent *pevents = NULL;-  if (ctx->profiling && !ctx->profiling_paused) {-    pevents = cuda_get_events(ctx, "copy_scalar_to_dev");-    CUDA_SUCCEED_FATAL(cuEventRecord(pevents[0], ctx->stream));+  struct cuda_event *event = cuda_event_new(ctx);+  if (event != NULL) {+    add_event(ctx,+              "copy_scalar_to_dev",+              strdup(""),+              event,+              (event_report_fn)cuda_event_report);+    CUDA_SUCCEED_FATAL(cuEventRecord(event->start, ctx->stream));   }   CUDA_SUCCEED_OR_RETURN(cuMemcpyHtoD(dst + offset, src, size));-  if (pevents != NULL) {-    CUDA_SUCCEED_FATAL(cuEventRecord(pevents[1], ctx->stream));+  if (event != NULL) {+    CUDA_SUCCEED_FATAL(cuEventRecord(event->start, ctx->stream));   }   return FUTHARK_SUCCESS; }@@ -893,14 +883,18 @@ static int gpu_scalar_from_device(struct futhark_context* ctx,                                   void *dst,                                   gpu_mem src, size_t offset, size_t size) {-  CUevent *pevents = NULL;-  if (ctx->profiling && !ctx->profiling_paused) {-    pevents = cuda_get_events(ctx, "copy_scalar_from_dev");-    CUDA_SUCCEED_FATAL(cuEventRecord(pevents[0], ctx->stream));+  struct cuda_event *event = cuda_event_new(ctx);+  if (event != NULL) {+    add_event(ctx,+              "copy_scalar_from_dev",+              strdup(""),+              event,+              (event_report_fn)cuda_event_report);+    CUDA_SUCCEED_FATAL(cuEventRecord(event->start, ctx->stream));   }   CUDA_SUCCEED_OR_RETURN(cuMemcpyDtoH(dst, src + offset, size));-  if (pevents != NULL) {-    CUDA_SUCCEED_FATAL(cuEventRecord(pevents[1], ctx->stream));+  if (event != NULL) {+    CUDA_SUCCEED_FATAL(cuEventRecord(event->end, ctx->stream));   }   return FUTHARK_SUCCESS; }@@ -909,14 +903,18 @@                       gpu_mem dst, int64_t dst_offset,                       gpu_mem src, int64_t src_offset,                       int64_t nbytes) {-  CUevent *pevents = NULL;-  if (ctx->profiling && !ctx->profiling_paused) {-    pevents = cuda_get_events(ctx, "copy_dev_to_dev");-    CUDA_SUCCEED_FATAL(cuEventRecord(pevents[0], ctx->stream));+  struct cuda_event *event = cuda_event_new(ctx);+  if (event != NULL) {+    add_event(ctx,+              "copy_dev_to_dev",+              strdup(""),+              event,+              (event_report_fn)cuda_event_report);+    CUDA_SUCCEED_FATAL(cuEventRecord(event->start, ctx->stream));   }   CUDA_SUCCEED_OR_RETURN(cuMemcpy(dst+dst_offset, src+src_offset, nbytes));-  if (pevents != NULL) {-    CUDA_SUCCEED_FATAL(cuEventRecord(pevents[1], ctx->stream));+  if (event != NULL) {+    CUDA_SUCCEED_FATAL(cuEventRecord(event->end, ctx->stream));   }   return FUTHARK_SUCCESS; }@@ -926,10 +924,14 @@                            const unsigned char* src, int64_t src_offset,                            int64_t nbytes) {   if (nbytes > 0) {-    CUevent* pevents = NULL;-    if (ctx->profiling && !ctx->profiling_paused) {-      pevents = cuda_get_events(ctx, "copy_host_to_dev");-      CUDA_SUCCEED_FATAL(cuEventRecord(pevents[0], ctx->stream));+    struct cuda_event *event = cuda_event_new(ctx);+    if (event != NULL) {+      add_event(ctx,+                "copy_host_to_dev",+                strdup(""),+                event,+                (event_report_fn)cuda_event_report);+      CUDA_SUCCEED_FATAL(cuEventRecord(event->start, ctx->stream));     }     if (sync) {       CUDA_SUCCEED_OR_RETURN@@ -938,8 +940,8 @@       CUDA_SUCCEED_OR_RETURN         (cuMemcpyHtoDAsync(dst + dst_offset, src + src_offset, nbytes, ctx->stream));     }-    if (pevents != NULL) {-      CUDA_SUCCEED_FATAL(cuEventRecord(pevents[1], ctx->stream));+    if (event != NULL) {+      CUDA_SUCCEED_FATAL(cuEventRecord(event->end, ctx->stream));     }   }   return FUTHARK_SUCCESS;@@ -950,10 +952,14 @@                            gpu_mem src, int64_t src_offset,                            int64_t nbytes) {   if (nbytes > 0) {-    CUevent* pevents = NULL;-    if (ctx->profiling && !ctx->profiling_paused) {-      pevents = cuda_get_events(ctx, "copy_dev_to_host");-      CUDA_SUCCEED_FATAL(cuEventRecord(pevents[0], ctx->stream));+    struct cuda_event *event = cuda_event_new(ctx);+    if (event != NULL) {+      add_event(ctx,+                "copy_dev_to_host",+                strdup(""),+                event,+                (event_report_fn)cuda_event_report);+      CUDA_SUCCEED_FATAL(cuEventRecord(event->start, ctx->stream));     }     if (sync) {       CUDA_SUCCEED_OR_RETURN@@ -962,6 +968,9 @@       CUDA_SUCCEED_OR_RETURN         (cuMemcpyDtoHAsync(dst + dst_offset, src + src_offset, nbytes, ctx->stream));     }+    if (event != NULL) {+      CUDA_SUCCEED_FATAL(cuEventRecord(event->end, ctx->stream));+    }     if (sync &&         ctx->failure_is_an_option &&         futhark_context_sync(ctx) != 0) {@@ -981,23 +990,27 @@                              size_t args_sizes[num_args]) {   (void) args_sizes;   int64_t time_start = 0, time_end = 0;-  if (ctx->logging) {-    fprintf(ctx->log,-            "Launching kernel %s with\n"-            "  grid=(%d,%d,%d)\n"-            "  block=(%d,%d,%d)\n"-            "  local memory=%d\n",-            name,-            grid[0], grid[1], grid[2],-            block[0], block[1], block[2],-            local_mem_bytes);++  if (ctx->debugging) {     time_start = get_wall_time();   } -  CUevent *pevents = NULL;-  if (ctx->profiling && !ctx->profiling_paused) {-    pevents = cuda_get_events(ctx, name);-    CUDA_SUCCEED_FATAL(cuEventRecord(pevents[0], ctx->stream));+  struct cuda_event *event = cuda_event_new(ctx);++  if (event != NULL) {+    CUDA_SUCCEED_FATAL(cuEventRecord(event->start, ctx->stream));+    add_event(ctx,+              name,+              msgprintf("Kernel %s with\n"+                        "  grid=(%d,%d,%d)\n"+                        "  block=(%d,%d,%d)\n"+                        "  local memory=%d",+                        name,+                        grid[0], grid[1], grid[2],+                        block[0], block[1], block[2],+                        local_mem_bytes),+              event,+              (event_report_fn)cuda_event_report);   }    CUDA_SUCCEED_OR_RETURN@@ -1007,8 +1020,8 @@                     local_mem_bytes, ctx->stream,                     args, NULL)); -  if (pevents != NULL) {-    CUDA_SUCCEED_FATAL(cuEventRecord(pevents[1], ctx->stream));+  if (event != NULL) {+    CUDA_SUCCEED_FATAL(cuEventRecord(event->end, ctx->stream));   }    if (ctx->debugging) {

@@ -77,7 +77,7 @@   int debugging;   int profiling;   int logging;-  const char *cache_fname;+  char* cache_fname;   int num_tuning_params;   int64_t *tuning_params;   const char** tuning_param_names;@@ -87,9 +87,9 @@    char* program;   int num_build_opts;-  const char **build_opts;+  char* *build_opts; -  const char *preferred_device;+  char* preferred_device;   int preferred_device_num;    size_t default_block_size;@@ -105,10 +105,10 @@  static void backend_context_config_setup(struct futhark_context_config *cfg) {   cfg->num_build_opts = 0;-  cfg->build_opts = (const char**) malloc(sizeof(const char*));+  cfg->build_opts = (char**) malloc(sizeof(char*));   cfg->build_opts[0] = NULL;   cfg->preferred_device_num = 0;-  cfg->preferred_device = "";+  cfg->preferred_device = strdup("");   cfg->program = strconcat(gpu_program);    cfg->default_block_size = 256;@@ -123,14 +123,19 @@ }  static void backend_context_config_teardown(struct futhark_context_config* cfg) {+  for (int i = 0; i < cfg->num_build_opts; i++) {+    free(cfg->build_opts[i]);+  }+   free(cfg->build_opts);+  free(cfg->preferred_device);   free(cfg->program); }  void futhark_context_config_add_build_option(struct futhark_context_config *cfg, const char *opt) {-  cfg->build_opts[cfg->num_build_opts] = opt;+  cfg->build_opts[cfg->num_build_opts] = strdup(opt);   cfg->num_build_opts++;-  cfg->build_opts = (const char **) realloc(cfg->build_opts, (cfg->num_build_opts + 1) * sizeof(const char *));+  cfg->build_opts = (char **) realloc(cfg->build_opts, (cfg->num_build_opts + 1) * sizeof(char *));   cfg->build_opts[cfg->num_build_opts] = NULL; } @@ -146,7 +151,8 @@       s++;     }   }-  cfg->preferred_device = s;+  free(cfg->preferred_device);+  cfg->preferred_device = strdup(s);   cfg->preferred_device_num = x; } @@ -156,6 +162,7 @@ }  void futhark_context_config_set_program(struct futhark_context_config *cfg, const char *s) {+  free(cfg->program);   cfg->program = strdup(s); } @@ -214,12 +221,6 @@   return 1; } -// A record of something that happened.-struct profiling_record {-  hipEvent_t *events; // Points to two events.-  const char *name;-};- struct futhark_context {   struct futhark_context_config* cfg;   int detail_memory;@@ -233,6 +234,7 @@   FILE *log;   struct constants *constants;   struct free_list free_list;+  struct event_list event_list;   int64_t peak_mem_usage_default;   int64_t cur_mem_usage_default;   // Uniform fields above.@@ -264,10 +266,6 @@    size_t lockstep_width; -  struct profiling_record *profiling_records;-  int profiling_records_capacity;-  int profiling_records_used;-   struct builtin_kernels* kernels; }; @@ -568,56 +566,43 @@   return NULL; } -static int tally_profiling_records(struct futhark_context *ctx,-                                   struct cost_centres* ccs) {-  hipError_t err;-  for (int i = 0; i < ctx->profiling_records_used; i++) {-    struct profiling_record record = ctx->profiling_records[i];+struct hip_event {+  hipEvent_t start;+  hipEvent_t end;+}; -    float ms;-    if ((err = hipEventElapsedTime(&ms, record.events[0], record.events[1])) != hipSuccess) {-      return err;-    }+static struct hip_event* hip_event_new(struct futhark_context* ctx) {+  if (ctx->profiling && !ctx->profiling_paused) {+    struct hip_event* e = malloc(sizeof(struct hip_event));+    hipEventCreate(&e->start);+    hipEventCreate(&e->end);+    return e;+  } else {+    return NULL;+  }+} -    if (ccs) {-      struct cost_centre c = {-        .name = record.name,-        .runs = 1,-        .runtime = ms*1000-      };-      cost_centres_add(ccs, c);-    }+static int hip_event_report(struct str_builder* sb, struct hip_event* e) {+  float ms;+  hipError_t err;+  if ((err = hipEventElapsedTime(&ms, e->start, e->end)) != hipSuccess) {+    return err;+  } -    if ((err = hipEventDestroy(record.events[0])) != hipSuccess) {-      return 1;-    }-    if ((err = hipEventDestroy(record.events[1])) != hipSuccess) {-      return 1;-    }+  // HIP provides milisecond resolution, but we want microseconds.+  str_builder(sb, ",\"duration\":%f", ms*1000); -    free(record.events);+  if ((err = hipEventDestroy(e->start)) != hipSuccess) {+    return 1;   } -  ctx->profiling_records_used = 0;+  if ((err = hipEventDestroy(e->end)) != hipSuccess) {+    return 1;+  } -  return 0;-}+  free(e); -static hipEvent_t* hip_get_events(struct futhark_context *ctx, const char* name) {-  if (ctx->profiling_records_used == ctx->profiling_records_capacity) {-    ctx->profiling_records_capacity *= 2;-    ctx->profiling_records =-      realloc(ctx->profiling_records,-              ctx->profiling_records_capacity *-              sizeof(struct profiling_record));-  }-  hipEvent_t *events = calloc(2, sizeof(hipEvent_t));-  hipEventCreate(&events[0]);-  hipEventCreate(&events[1]);-  ctx->profiling_records[ctx->profiling_records_used].events = events;-  ctx->profiling_records[ctx->profiling_records_used].name = name;-  ctx->profiling_records_used++;-  return events;+  return 0; }  int futhark_context_sync(struct futhark_context* ctx) {@@ -655,11 +640,6 @@ void free_builtin_kernels(struct futhark_context* ctx, struct builtin_kernels* kernels);  int backend_context_setup(struct futhark_context* ctx) {-  ctx->profiling_records_capacity = 200;-  ctx->profiling_records_used = 0;-  ctx->profiling_records =-    malloc(ctx->profiling_records_capacity *-           sizeof(struct profiling_record));   ctx->failure_is_an_option = 0;   ctx->total_runs = 0;   ctx->total_runtime = 0;@@ -688,8 +668,10 @@   ctx->lockstep_width = 32;   HIP_SUCCEED_FATAL(hipStreamCreate(&ctx->stream));   hip_size_setup(ctx);-  ctx->error = hip_module_setup(ctx, ctx->cfg->program,-                                ctx->cfg->build_opts, ctx->cfg->cache_fname);+  ctx->error = hip_module_setup(ctx,+                                ctx->cfg->program,+                                (const char**)ctx->cfg->build_opts,+                                ctx->cfg->cache_fname);    if (ctx->error != NULL) {     futhark_panic(1, "During HIP initialisation:\n%s\n", ctx->error);@@ -713,8 +695,6 @@   hipFree(ctx->global_failure);   hipFree(ctx->global_failure_args);   HIP_SUCCEED_FATAL(gpu_free_all(ctx));-  (void)tally_profiling_records(ctx, NULL);-  free(ctx->profiling_records);   HIP_SUCCEED_FATAL(hipStreamDestroy(ctx->stream));   HIP_SUCCEED_FATAL(hipModuleUnload(ctx->module)); }@@ -742,14 +722,18 @@ static int gpu_scalar_to_device(struct futhark_context* ctx,                                 gpu_mem dst, size_t offset, size_t size,                                 void *src) {-  hipEvent_t *pevents = NULL;-  if (ctx->profiling && !ctx->profiling_paused) {-    pevents = hip_get_events(ctx, "copy_scalar_to_dev");-    HIP_SUCCEED_FATAL(hipEventRecord(pevents[0], ctx->stream));+  struct hip_event *event = hip_event_new(ctx);+  if (event != NULL) {+    add_event(ctx,+              "copy_scalar_to_dev",+              strdup(""),+              event,+              (event_report_fn)hip_event_report);+    HIP_SUCCEED_FATAL(hipEventRecord(event->start, ctx->stream));   }   HIP_SUCCEED_OR_RETURN(hipMemcpyHtoD((unsigned char*)dst + offset, src, size));-  if (pevents != NULL) {-    HIP_SUCCEED_FATAL(hipEventRecord(pevents[1], ctx->stream));+  if (event != NULL) {+    HIP_SUCCEED_FATAL(hipEventRecord(event->end, ctx->stream));   }   return FUTHARK_SUCCESS; }@@ -757,14 +741,18 @@ static int gpu_scalar_from_device(struct futhark_context* ctx,                                   void *dst,                                   gpu_mem src, size_t offset, size_t size) {-  hipEvent_t *pevents = NULL;-  if (ctx->profiling && !ctx->profiling_paused) {-    pevents = hip_get_events(ctx, "copy_scalar_from_dev");-    HIP_SUCCEED_FATAL(hipEventRecord(pevents[0], ctx->stream));+  struct hip_event *event = hip_event_new(ctx);+  if (event != NULL) {+    add_event(ctx,+              "copy_scalar_from_dev",+              strdup(""),+              event,+              (event_report_fn)hip_event_report);+    HIP_SUCCEED_FATAL(hipEventRecord(event->start, ctx->stream));   }   HIP_SUCCEED_OR_RETURN(hipMemcpyDtoH(dst, (unsigned char*)src + offset, size));-  if (pevents != NULL) {-    HIP_SUCCEED_FATAL(hipEventRecord(pevents[1], ctx->stream));+  if (event != NULL) {+    HIP_SUCCEED_FATAL(hipEventRecord(event->end, ctx->stream));   }   return FUTHARK_SUCCESS; }@@ -773,15 +761,19 @@                       gpu_mem dst, int64_t dst_offset,                       gpu_mem src, int64_t src_offset,                       int64_t nbytes) {-  hipEvent_t *pevents = NULL;-  if (ctx->profiling && !ctx->profiling_paused) {-    pevents = hip_get_events(ctx, "copy_dev_to_dev");-    HIP_SUCCEED_FATAL(hipEventRecord(pevents[0], ctx->stream));+  struct hip_event *event = hip_event_new(ctx);+  if (event != NULL) {+    add_event(ctx,+              "copy_dev_to_dev",+              strdup(""),+              event,+              (event_report_fn)hip_event_report);+    HIP_SUCCEED_FATAL(hipEventRecord(event->start, ctx->stream));   }   HIP_SUCCEED_OR_RETURN(hipMemcpyWithStream((unsigned char*)dst+dst_offset, (unsigned char*)src+src_offset,                                             nbytes, hipMemcpyDeviceToDevice ,ctx->stream));-  if (pevents != NULL) {-    HIP_SUCCEED_FATAL(hipEventRecord(pevents[1], ctx->stream));+  if (event != NULL) {+    HIP_SUCCEED_FATAL(hipEventRecord(event->end, ctx->stream));   }   return FUTHARK_SUCCESS; }@@ -791,10 +783,14 @@                            const unsigned char* src, int64_t src_offset,                            int64_t nbytes) {   if (nbytes > 0) {-    hipEvent_t* pevents = NULL;-    if (ctx->profiling && !ctx->profiling_paused) {-      pevents = hip_get_events(ctx, "copy_host_to_dev");-      HIP_SUCCEED_FATAL(hipEventRecord(pevents[0], ctx->stream));+    struct hip_event *event = hip_event_new(ctx);+    if (event != NULL) {+      add_event(ctx,+                "copy_host_to_dev",+                strdup(""),+                event,+                (event_report_fn)hip_event_report);+      HIP_SUCCEED_FATAL(hipEventRecord(event->start, ctx->stream));     }     if (sync) {       HIP_SUCCEED_OR_RETURN@@ -806,8 +802,8 @@                             (unsigned char*)src + src_offset,                             nbytes, ctx->stream));     }-    if (pevents != NULL) {-      HIP_SUCCEED_FATAL(hipEventRecord(pevents[1], ctx->stream));+    if (event != NULL) {+      HIP_SUCCEED_FATAL(hipEventRecord(event->end, ctx->stream));     }   }   return FUTHARK_SUCCESS;@@ -818,10 +814,14 @@                            gpu_mem src, int64_t src_offset,                            int64_t nbytes) {   if (nbytes > 0) {-    hipEvent_t* pevents = NULL;-    if (ctx->profiling && !ctx->profiling_paused) {-      pevents = hip_get_events(ctx, "copy_dev_to_host");-      HIP_SUCCEED_FATAL(hipEventRecord(pevents[0], ctx->stream));+    struct hip_event *event = hip_event_new(ctx);+    if (event != NULL) {+      add_event(ctx,+                "copy_dev_to_host",+                strdup(""),+                event,+                (event_report_fn)hip_event_report);+      HIP_SUCCEED_FATAL(hipEventRecord(event->start, ctx->stream));     }     if (sync) {       HIP_SUCCEED_OR_RETURN@@ -834,6 +834,9 @@                             (unsigned char*)src + src_offset,                             nbytes, ctx->stream));     }+    if (event != NULL) {+      HIP_SUCCEED_FATAL(hipEventRecord(event->end, ctx->stream));+    }     if (sync &&         ctx->failure_is_an_option &&         futhark_context_sync(ctx) != 0) {@@ -853,23 +856,26 @@                              size_t args_sizes[num_args]) {   (void) args_sizes;   int64_t time_start = 0, time_end = 0;-  if (ctx->logging) {-    fprintf(ctx->log,-            "Launching kernel %s with\n"-            "  grid=(%d,%d,%d)\n"-            "  block=(%d,%d,%d)\n"-            "  local memory=%d\n",-            name,-            grid[0], grid[1], grid[2],-            block[0], block[1], block[2],-            local_mem_bytes);+  if (ctx->debugging) {     time_start = get_wall_time();   } -  hipEvent_t *pevents = NULL;-  if (ctx->profiling && !ctx->profiling_paused) {-    pevents = hip_get_events(ctx, name);-    HIP_SUCCEED_FATAL(hipEventRecord(pevents[0], ctx->stream));+  struct hip_event *event = hip_event_new(ctx);++  if (event != NULL) {+    HIP_SUCCEED_FATAL(hipEventRecord(event->start, ctx->stream));+    add_event(ctx,+              name,+              msgprintf("Kernel %s with\n"+                        "  grid=(%d,%d,%d)\n"+                        "  block=(%d,%d,%d)\n"+                        "  local memory=%d",+                        name,+                        grid[0], grid[1], grid[2],+                        block[0], block[1], block[2],+                        local_mem_bytes),+              event,+              (event_report_fn)hip_event_report);   }    HIP_SUCCEED_OR_RETURN@@ -879,8 +885,8 @@                            local_mem_bytes, ctx->stream,                            args, NULL)); -  if (pevents != NULL) {-    HIP_SUCCEED_FATAL(hipEventRecord(pevents[1], ctx->stream));+  if (event != NULL) {+    HIP_SUCCEED_FATAL(hipEventRecord(event->end, ctx->stream));   }    if (ctx->debugging) {

@@ -5,7 +5,7 @@   int debugging;   int profiling;   int logging;-  const char *cache_fname;+  char *cache_fname;   int num_tuning_params;   int64_t *tuning_params;   const char** tuning_param_names;@@ -46,11 +46,13 @@   FILE *log;   struct constants *constants;   struct free_list free_list;+  struct event_list event_list;   int64_t peak_mem_usage_default;   int64_t cur_mem_usage_default;   struct program* program;   // Uniform fields above. +  lock_t event_list_lock;   struct scheduler scheduler;   int total_runs;   long int total_runtime;@@ -80,16 +82,41 @@     return 1;   } +  create_lock(&ctx->event_list_lock);+   return 0; }  void backend_context_teardown(struct futhark_context* ctx) {   (void)scheduler_destroy(&ctx->scheduler);+  free_lock(&ctx->event_list_lock); }  int futhark_context_sync(struct futhark_context* ctx) {   (void)ctx;   return 0; }++struct mc_event {+  // Time in microseconds.+  uint64_t bef, aft;+};++static struct mc_event* mc_event_new(struct futhark_context* ctx) {+  if (ctx->profiling && !ctx->profiling_paused) {+    struct mc_event* e = malloc(sizeof(struct mc_event));+    return e;+  } else {+    return NULL;+  }+}++static int mc_event_report(struct str_builder* sb, struct mc_event* e) {+  float ms = e->aft - e->bef;+  str_builder(sb, ",\"duration\":%f", ms);+  free(e);+  return 0;+}+  // End of backends/multicore.h

@@ -122,7 +122,7 @@   int debugging;   int profiling;   int logging;-  const char *cache_fname;+  char *cache_fname;   int num_tuning_params;   int64_t *tuning_params;   const char** tuning_param_names;@@ -132,12 +132,12 @@    char* program;   int preferred_device_num;-  const char *preferred_platform;-  const char *preferred_device;+  char* preferred_platform;+  char* preferred_device;   int ignore_blacklist; -  const char* dump_binary_to;-  const char* load_binary_from;+  char* dump_binary_to;+  char* load_binary_from;    size_t default_group_size;   size_t default_num_groups;@@ -148,7 +148,7 @@   int default_group_size_changed;   int default_tile_size_changed;   int num_build_opts;-  const char **build_opts;+  char* *build_opts;    cl_command_queue queue;   int queue_set;@@ -156,11 +156,11 @@  static void backend_context_config_setup(struct futhark_context_config* cfg) {   cfg->num_build_opts = 0;-  cfg->build_opts = (const char**) malloc(sizeof(const char*));+  cfg->build_opts = (char**) malloc(sizeof(const char*));   cfg->build_opts[0] = NULL;   cfg->preferred_device_num = 0;-  cfg->preferred_platform = "";-  cfg->preferred_device = "";+  cfg->preferred_platform = strdup("");+  cfg->preferred_device = strdup("");   cfg->ignore_blacklist = 0;   cfg->dump_binary_to = NULL;   cfg->load_binary_from = NULL;@@ -182,14 +182,21 @@ }  static void backend_context_config_teardown(struct futhark_context_config* cfg) {+  for (int i = 0; i < cfg->num_build_opts; i++) {+    free(cfg->build_opts[i]);+  }   free(cfg->build_opts);+  free(cfg->dump_binary_to);+  free(cfg->load_binary_from);+  free(cfg->preferred_device);+  free(cfg->preferred_platform);   free(cfg->program); }  void futhark_context_config_add_build_option(struct futhark_context_config* cfg, const char *opt) {-  cfg->build_opts[cfg->num_build_opts] = opt;+  cfg->build_opts[cfg->num_build_opts] = strdup(opt);   cfg->num_build_opts++;-  cfg->build_opts = (const char**) realloc(cfg->build_opts, (cfg->num_build_opts+1) * sizeof(const char*));+  cfg->build_opts = (char**) realloc(cfg->build_opts, (cfg->num_build_opts+1) * sizeof(char*));   cfg->build_opts[cfg->num_build_opts] = NULL; } @@ -205,13 +212,15 @@       s++;     }   }-  cfg->preferred_device = s;+  free(cfg->preferred_device);+  cfg->preferred_device = strdup(s);   cfg->preferred_device_num = x;   cfg->ignore_blacklist = 1; }  void futhark_context_config_set_platform(struct futhark_context_config *cfg, const char *s) {-  cfg->preferred_platform = s;+  free(cfg->preferred_platform);+  cfg->preferred_platform = strdup(s);   cfg->ignore_blacklist = 1; } @@ -409,15 +418,18 @@ }  void futhark_context_config_set_program(struct futhark_context_config *cfg, const char *s) {+  free(cfg->program);   cfg->program = strdup(s); }  void futhark_context_config_dump_binary_to(struct futhark_context_config *cfg, const char *path) {-  cfg->dump_binary_to = path;+  free(cfg->dump_binary_to);+  cfg->dump_binary_to = strdup(path); }  void futhark_context_config_load_binary_from(struct futhark_context_config *cfg, const char *path) {-  cfg->load_binary_from = path;+  free(cfg->load_binary_from);+  cfg->load_binary_from = strdup(path); }  void futhark_context_config_set_default_group_size(struct futhark_context_config *cfg, int size) {@@ -474,12 +486,6 @@   return 1; } -// A record of something that happened.-struct profiling_record {-  cl_event *event;-  const char* name;-};- struct futhark_context {   struct futhark_context_config* cfg;   int detail_memory;@@ -493,6 +499,7 @@   FILE *log;   struct constants *constants;   struct free_list free_list;+  struct event_list event_list;   int64_t peak_mem_usage_default;   int64_t cur_mem_usage_default;   struct program* program;@@ -524,10 +531,6 @@    size_t lockstep_width; -  struct profiling_record *profiling_records;-  int profiling_records_capacity;-  int profiling_records_used;-   struct builtin_kernels* kernels; }; @@ -616,61 +619,38 @@   return compile_opts; } +static cl_event* opencl_event_new(struct futhark_context* ctx) {+  if (ctx->profiling && !ctx->profiling_paused) {+    return malloc(sizeof(cl_event));+  } else {+    return NULL;+  }+} -// Count up the runtime all the profiling_records that occured during execution.-// Also clears the buffer of profiling_records.-static void tally_profiling_records(struct futhark_context *ctx,-                                    struct cost_centres* ccs) {+static int opencl_event_report(struct str_builder* sb, cl_event* e) {   cl_int err;-  for (int i = 0; i < ctx->profiling_records_used; i++) {-    struct profiling_record record = ctx->profiling_records[i];--    cl_ulong start_t, end_t;--    OPENCL_SUCCEED_FATAL(clGetEventProfilingInfo(*record.event,-                                                 CL_PROFILING_COMMAND_START,-                                                 sizeof(start_t),-                                                 &start_t,-                                                 NULL));+  cl_ulong start_t, end_t; -    OPENCL_SUCCEED_FATAL(clGetEventProfilingInfo(*record.event,-                                                 CL_PROFILING_COMMAND_END,-                                                 sizeof(end_t),-                                                 &end_t,-                                                 NULL));+  assert(e != NULL);+  OPENCL_SUCCEED_FATAL(clGetEventProfilingInfo(*e,+                                               CL_PROFILING_COMMAND_START,+                                               sizeof(start_t),+                                               &start_t,+                                               NULL));+  OPENCL_SUCCEED_FATAL(clGetEventProfilingInfo(*e,+                                               CL_PROFILING_COMMAND_END,+                                               sizeof(end_t),+                                               &end_t,+                                               NULL)); -    if (ccs) {-      // Note that OpenCL provides nanosecond resolution, but we want-      // microseconds.-      struct cost_centre c = {-        .name = record.name,-        .runs = 1,-        .runtime = (end_t - start_t)/1000-      };-      cost_centres_add(ccs, c);-    }+  // OpenCL provides nanosecond resolution, but we want microseconds.+  str_builder(sb, ",\"duration\":%f", (end_t - start_t)/1000.0); -    OPENCL_SUCCEED_FATAL(clReleaseEvent(*record.event));-    free(record.event);-  }+  OPENCL_SUCCEED_FATAL(clReleaseEvent(*e)); -  ctx->profiling_records_used = 0;-}+  free(e); -// If profiling, produce an event associated with a profiling record.-static cl_event* opencl_get_event(struct futhark_context *ctx, const char *name) {-  if (ctx->profiling_records_used == ctx->profiling_records_capacity) {-    ctx->profiling_records_capacity *= 2;-    ctx->profiling_records =-      realloc(ctx->profiling_records,-              ctx->profiling_records_capacity *-              sizeof(struct profiling_record));-  }-  cl_event *event = malloc(sizeof(cl_event));-  ctx->profiling_records[ctx->profiling_records_used].event = event;-  ctx->profiling_records[ctx->profiling_records_used].name = name;-  ctx->profiling_records_used++;-  return event;+  return 0; }  int futhark_context_sync(struct futhark_context* ctx) {@@ -719,7 +699,7 @@ // array must be NULL-terminated. static void setup_opencl_with_command_queue(struct futhark_context *ctx,                                             cl_command_queue queue,-                                            const char *extra_build_opts[],+                                            const char* extra_build_opts[],                                             const char* cache_fname) {   int error; @@ -1047,11 +1027,6 @@  int backend_context_setup(struct futhark_context* ctx) {   ctx->lockstep_width = 0; // Real value set later.-  ctx->profiling_records_capacity = 200;-  ctx->profiling_records_used = 0;-  ctx->profiling_records =-    malloc(ctx->profiling_records_capacity *-           sizeof(struct profiling_record));   ctx->failure_is_an_option = 0;   ctx->total_runs = 0;   ctx->total_runtime = 0;@@ -1059,9 +1034,9 @@   ctx->cur_mem_usage_device = 0;    if (ctx->cfg->queue_set) {-    setup_opencl_with_command_queue(ctx, ctx->cfg->queue, ctx->cfg->build_opts, ctx->cfg->cache_fname);+    setup_opencl_with_command_queue(ctx, ctx->cfg->queue, (const char**)ctx->cfg->build_opts, ctx->cfg->cache_fname);   } else {-    setup_opencl(ctx, ctx->cfg->build_opts, ctx->cfg->cache_fname);+    setup_opencl(ctx, (const char**)ctx->cfg->build_opts, ctx->cfg->cache_fname);   }    cl_int error;@@ -1092,8 +1067,6 @@   free_builtin_kernels(ctx, ctx->kernels);   OPENCL_SUCCEED_FATAL(clReleaseMemObject(ctx->global_failure));   OPENCL_SUCCEED_FATAL(clReleaseMemObject(ctx->global_failure_args));-  (void)tally_profiling_records(ctx, NULL);-  free(ctx->profiling_records);   (void)gpu_free_all(ctx);   (void)clReleaseProgram(ctx->clprogram);   (void)clReleaseCommandQueue(ctx->queue);@@ -1131,9 +1104,13 @@ static int gpu_scalar_to_device(struct futhark_context* ctx,                                 gpu_mem dst, size_t offset, size_t size,                                 void *src) {-  cl_event* event = NULL;-  if (ctx->profiling && !ctx->profiling_paused) {-    event = opencl_get_event(ctx, "copy_scalar_to_dev");+  cl_event* event = opencl_event_new(ctx);+  if (event != NULL) {+    add_event(ctx,+              "copy_scalar_to_dev",+              strdup(""),+              event,+              (event_report_fn)opencl_event_report);   }   OPENCL_SUCCEED_OR_RETURN     (clEnqueueWriteBuffer@@ -1145,9 +1122,13 @@ static int gpu_scalar_from_device(struct futhark_context* ctx,                                   void *dst,                                   gpu_mem src, size_t offset, size_t size) {-  cl_event* event = NULL;-  if (ctx->profiling && !ctx->profiling_paused) {-    event = opencl_get_event(ctx, "copy_scalar_from_dev");+  cl_event* event = opencl_event_new(ctx);+  if (event != NULL) {+    add_event(ctx,+              "copy_scalar_from_dev",+              strdup(""),+              event,+              (event_report_fn)opencl_event_report);   }   OPENCL_SUCCEED_OR_RETURN     (clEnqueueReadBuffer@@ -1161,9 +1142,13 @@                       gpu_mem src, int64_t src_offset,                       int64_t nbytes) {   if (nbytes > 0) {-    cl_event* event = NULL;-    if (ctx->profiling && !ctx->profiling_paused) {-      event = opencl_get_event(ctx, "copy_dev_to_dev");+    cl_event* event = opencl_event_new(ctx);+    if (event != NULL) {+      add_event(ctx,+                "copy_dev_to_dev",+                strdup(""),+                event,+                (event_report_fn)opencl_event_report);     }     // OpenCL swaps the usual order of operands for memcpy()-like     // functions.  The order below is not a typo.@@ -1183,9 +1168,13 @@                            const unsigned char* src, int64_t src_offset,                            int64_t nbytes) {   if (nbytes > 0) {-    cl_event* event = NULL;-    if (ctx->profiling && !ctx->profiling_paused) {-      event = opencl_get_event(ctx, "copy_host_to_dev");+    cl_event* event = opencl_event_new(ctx);+    if (event != NULL) {+      add_event(ctx,+                "copy_host_to_dev",+                strdup(""),+                event,+                (event_report_fn)opencl_event_report);     }     OPENCL_SUCCEED_OR_RETURN       (clEnqueueWriteBuffer(ctx->queue,@@ -1206,9 +1195,13 @@                            gpu_mem src, int64_t src_offset,                            int64_t nbytes) {   if (nbytes > 0) {-    cl_event* event = NULL;-    if (ctx->profiling && !ctx->profiling_paused) {-      event = opencl_get_event(ctx, "copy_dev_to_host");+    cl_event* event = opencl_event_new(ctx);+    if (event != NULL) {+      add_event(ctx,+                "copy_dev_to_host",+                strdup(""),+                event,+                (event_report_fn)opencl_event_report);     }     OPENCL_SUCCEED_OR_RETURN       (clEnqueueReadBuffer(ctx->queue, src,@@ -1235,16 +1228,24 @@                              void* args[num_args],                              size_t args_sizes[num_args]) {   int64_t time_start = 0, time_end = 0;-  if (ctx->logging) {-    fprintf(ctx->log,-            "Launching kernel %s with\n"-            "  grid=(%d,%d,%d)\n"-            "  block=(%d,%d,%d)\n"-            "  local memory=%d\n",-            name,-            grid[0], grid[1], grid[2],-            block[0], block[1], block[2],-            local_mem_bytes);++  cl_event* event = opencl_event_new(ctx);+  if (event != NULL) {+    add_event(ctx,+              name,+              msgprintf("Kernel %s with\n"+                        "  grid=(%d,%d,%d)\n"+                        "  block=(%d,%d,%d)\n"+                        "  local memory=%d",+                        name,+                        grid[0], grid[1], grid[2],+                        block[0], block[1], block[2],+                        local_mem_bytes),+              event,+              (event_report_fn)opencl_event_report);+  }++  if (ctx->debugging) {     time_start = get_wall_time();   } @@ -1268,11 +1269,6 @@     {block[0],      block[1],      block[2]};--  cl_event* event = NULL;-  if (ctx->profiling && !ctx->profiling_paused) {-    event = opencl_get_event(ctx, name);-  }    OPENCL_SUCCEED_OR_RETURN     (clEnqueueNDRangeKernel(ctx->queue,

@@ -1,9 +1,5 @@ // Start of context.h -// Eventually it would be nice to move the context definition in here-// instead of generating it in the compiler.  For now it defines-// various helper functions that must be available.- // Internal functions.  static void set_error(struct futhark_context* ctx, char *error) {@@ -72,6 +68,17 @@   } } +static void add_event(struct futhark_context* ctx,+                      const char* name,+                      char* description,+                      void* data,+                      event_report_fn f) {+  if (ctx->logging) {+    fprintf(ctx->log, "Event: %s\n%s\n", name, description);+  }+  add_event_to_list(&ctx->event_list, name, description, data, f);+}+ struct futhark_context_config* futhark_context_config_new(void) {   struct futhark_context_config* cfg = malloc(sizeof(struct futhark_context_config));   if (cfg == NULL) {@@ -96,6 +103,7 @@ void futhark_context_config_free(struct futhark_context_config* cfg) {   assert(!cfg->in_use);   backend_context_config_teardown(cfg);+  free(cfg->cache_fname);   free(cfg->tuning_params);   free(cfg); }@@ -111,6 +119,7 @@   create_lock(&ctx->error_lock);   create_lock(&ctx->lock);   free_list_init(&ctx->free_list);+  event_list_init(&ctx->event_list);   ctx->peak_mem_usage_default = 0;   ctx->cur_mem_usage_default = 0;   ctx->constants = malloc(sizeof(struct constants));@@ -137,7 +146,9 @@   backend_context_teardown(ctx);   free_all_in_free_list(ctx);   free_list_destroy(&ctx->free_list);+  event_list_free(&ctx->event_list);   free(ctx->constants);+  free(ctx->error);   free_lock(&ctx->lock);   free_lock(&ctx->error_lock);   ctx->cfg->in_use = 0;

@@ -41,6 +41,12 @@ static bool lmad_memcpyable(int r,                             int64_t dst_strides[r], int64_t src_strides[r], int64_t shape[r]); +static void add_event(struct futhark_context* ctx,+                      const char* name,+                      char* description,+                      void* data,+                      event_report_fn f);+ // Functions that must be defined by the backend. static void backend_context_config_setup(struct futhark_context_config* cfg); static void backend_context_config_teardown(struct futhark_context_config* cfg);

@@ -0,0 +1,62 @@+// Start of event_list.h++typedef int (*event_report_fn)(struct str_builder*, void*);++struct event {+  void* data;+  event_report_fn f;+  const char* name;+  char *description;+};++struct event_list {+  struct event *events;+  int num_events;+  int capacity;+};++static void event_list_init(struct event_list *l) {+  l->capacity = 100;+  l->num_events = 0;+  l->events = calloc(l->capacity, sizeof(struct event));+}++static void event_list_free(struct event_list *l) {+  free(l->events);+}++static void add_event_to_list(struct event_list *l,+                              const char* name,+                              char* description,+                              void* data,+                              event_report_fn f) {+  if (l->num_events == l->capacity) {+    l->capacity *= 2;+    l->events = realloc(l->events, l->capacity * sizeof(struct event));+  }+  l->events[l->num_events].name = name;+  l->events[l->num_events].description = description;+  l->events[l->num_events].data = data;+  l->events[l->num_events].f = f;+  l->num_events++;+}++static void report_events_in_list(struct event_list *l,+                                  struct str_builder* sb) {+  for (int i = 0; i < l->num_events; i++) {+    if (i != 0) {+      str_builder_str(sb, ",");+    }+    str_builder_str(sb, "{\"name\":");+    str_builder_json_str(sb, l->events[i].name);+    str_builder_str(sb, ",\"description\":");+    str_builder_json_str(sb, l->events[i].description);+    free(l->events[i].description);+    l->events[i].f(sb, l->events[i].data);+    str_builder(sb, "}");+  }+  event_list_free(l);+  event_list_init(l);+}++// End of event_list.h

@@ -136,76 +136,43 @@   b->used += needed; } -struct cost_centre {-  const char *name;-  int64_t runs;-  int64_t runtime;-};--// Dynamic dictionary for tallying cost centres when aggregating-// profiling information.  Not performance-critical.-struct cost_centres {-  size_t capacity;-  size_t used;-  struct cost_centre* centres;-};--static struct cost_centres *cost_centres_new() {-  struct cost_centres *ccs = malloc(sizeof(struct cost_centres));-  ccs->capacity = 100;-  ccs->used = 0;-  ccs->centres = calloc(ccs->capacity, sizeof(struct cost_centre));-  return ccs;-}--static void cost_centres_free(struct cost_centres* ccs) {-  free(ccs->centres);-  free(ccs);-}--static void cost_centres_init(struct cost_centres* ccs, const char *name) {-  if (ccs->used == ccs->capacity) {-    ccs->capacity *= 2;-    ccs->centres = realloc(ccs->centres, ccs->capacity*sizeof(struct cost_centre));+static void str_builder_str(struct str_builder *b, const char *s) {+  size_t needed = strlen(s);+  if (b->capacity < b->used + needed + 1) {+    b->capacity *= 2;+    b->str = realloc(b->str, b->capacity);   }-  ccs->centres[ccs->used].name = name;-  ccs->centres[ccs->used].runs = 0;-  ccs->centres[ccs->used].runtime = 0;-  ccs->used++;+  strcpy(b->str+b->used, s);+  b->used += needed; } -static void cost_centres_add(struct cost_centres* ccs, struct cost_centre c) {-  size_t i = 0;-  for (i = 0; i < ccs->used; i++) {-    if (strcmp(c.name, ccs->centres[i].name) == 0) {-      ccs->centres[i].runs += c.runs;-      ccs->centres[i].runtime += c.runtime;-      return;-    }-  }-  if (i == ccs->capacity) {-    ccs->capacity *= 2;-    ccs->centres = realloc(ccs->centres, ccs->capacity*sizeof(struct cost_centre));+static void str_builder_char(struct str_builder *b, char c) {+  size_t needed = 1;+  if (b->capacity < b->used + needed + 1) {+    b->capacity *= 2;+    b->str = realloc(b->str, b->capacity);   }-  ccs->centres[i] = c;-  ccs->used++;+  b->str[b->used] = c;+  b->str[b->used+1] = 0;+  b->used += needed; } -static void cost_centre_report(struct cost_centres* ccs, struct str_builder *b) {-  int64_t total_runs = 0;-  int64_t total_runtime = 0;-  for (size_t i = 0; i < ccs->used; i++) {-    struct cost_centre c = ccs->centres[i];-    str_builder(b,-                "%-40s ran %5d times; avg %8ldus; total: %8ldus\n",-                c.name,-                c.runs, c.runs == 0 ? 0 : c.runtime/c.runs, c.runtime);-    total_runs += c.runs;-    total_runtime += c.runtime;+static void str_builder_json_str(struct str_builder* sb, const char* s) {+  str_builder_char(sb, '"');+  for (int j = 0; s[j]; j++) {+    char c = s[j];+    switch (c) {+    case '\n':+      str_builder_str(sb, "\\n");+      break;+    case '"':+      str_builder_str(sb, "\\\"");+      break;+    default:+      str_builder_char(sb, c);+    }   }-  str_builder(b,-              "%d operations with cumulative runtime: %6ldus\n",-              total_runs, total_runtime);+  str_builder_char(sb, '"'); }  static char *strclone(const char *str) {

@@ -3,12 +3,8 @@ import ctypes as ct  -def addressOffset(x, offset, bt):-    return ct.cast(ct.addressof(x.contents) + int(offset), ct.POINTER(bt))-- def allocateMem(size):-    return ct.cast((ct.c_byte * max(0, size))(), ct.POINTER(ct.c_byte))+    return np.empty(size, dtype=np.byte)   # Copy an array if its is not-None.  This is important for treating@@ -21,28 +17,15 @@   def unwrapArray(x):-    return normaliseArray(x).ctypes.data_as(ct.POINTER(ct.c_byte))---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=t)-    else:-        return np.ctypeslib.as_array(x, shape=shape).view(t)+    return x.ravel().view(np.byte)   def indexArray(x, offset, bt):-    return addressOffset(x, offset * ct.sizeof(bt), bt)[0]+    return x.view(bt)[offset]   def writeScalarArray(x, offset, v):-    ct.memmove(-        ct.addressof(x.contents) + int(offset) * ct.sizeof(v),-        ct.addressof(v),-        ct.sizeof(v),-    )+    x.view(type(v))[offset] = v   # An opaque Futhark value.@@ -149,7 +132,7 @@             writeScalarArray(                 dst,                 dst_offset + i * dst_strides[0],-                pt(indexArray(src, src_offset + i * src_strides[0], pt)),+                indexArray(src, src_offset + i * src_strides[0], pt),             )     else:         for i in range(shape[0]):@@ -169,11 +152,13 @@     pt, dst, dst_offset, dst_strides, src, src_offset, src_strides, shape ):     if lmad_memcpyable(dst_strides, src_strides, shape):-        ct.memmove(-            addressOffset(dst, dst_offset * ct.sizeof(pt), ct.c_byte),-            addressOffset(src, src_offset * ct.sizeof(pt), ct.c_byte),-            np.prod(shape) * ct.sizeof(pt),-        )+        dst[+            dst_offset * ct.sizeof(pt) : dst_offset * ct.sizeof(pt)+            + np.prod(shape) * ct.sizeof(pt)+        ] = src[+            src_offset * ct.sizeof(pt) : src_offset * ct.sizeof(pt)+            + np.prod(shape) * ct.sizeof(pt)+        ]     else:         lmad_copy_elements(             pt,

@@ -430,7 +430,7 @@             "Futhark runtime limitation:\nCannot copy array of greater than rank 8.\n"         ) -    n = np.product(shape)+    n = np.prod(shape)     zero = np.int64(0)     layout_args = [None] * (8 * 3)     for i in range(8):@@ -464,7 +464,7 @@     self, pt, dst, dst_offset, dst_strides, src, src_offset, src_strides, shape ):     elem_size = ct.sizeof(pt)-    nbytes = np.product(shape) * elem_size+    nbytes = np.prod(shape) * elem_size     if nbytes == 0:         return None     if lmad_memcpyable(dst_strides, src_strides, shape):

@@ -334,7 +334,7 @@     for i in range(bits):         if x < 0:             break-        n += 1+        n += np.int32(1)         x <<= np.int8(1)     return n @@ -345,7 +345,7 @@     for i in range(bits):         if (x & 1) == 1:             break-        n += 1+        n += np.int32(1)         x >>= np.int8(1)     return n @@ -354,7 +354,7 @@     c = np.int32(0)     while x != 0:         x &= x - np.int8(1)-        c += np.int8(1)+        c += np.int32(1)     return c  

@@ -357,7 +357,7 @@         parse_specific_string(f, "[")         dims += [int(parse_int(f))]         parse_specific_string(f, "]")-    if np.product(dims) != 0:+    if np.prod(dims) != 0:         raise ValueError     parse_specific_string(f, type_name)     parse_specific_char(f, b")")@@ -672,7 +672,7 @@     )     arr.shape = shape -    return arr+    return arr.copy()  # To ensure it is writeable.   if sys.version_info >= (3, 0):@@ -758,7 +758,7 @@         else:             out.write("%.6ff64" % v)     elif type(v) == np.ndarray:-        if np.product(v.shape) == 0:+        if np.prod(v.shape) == 0:             tname = numpy_type_to_type_name(v.dtype)             out.write(                 "empty({}{})".format(

@@ -245,12 +245,12 @@     _ -> error $ "basicFwd: Unsupported op " ++ prettyString op  fwdLambda :: Lambda SOACS -> ADM (Lambda SOACS)-fwdLambda l@(Lambda params body ret) =-  Lambda <$> bundleNewList params <*> inScopeOf l (fwdBody body) <*> bundleTangents ret+fwdLambda l@(Lambda params ret body) =+  Lambda <$> bundleNewList params <*> bundleTangents ret <*> inScopeOf l (fwdBody body)  fwdStreamLambda :: Lambda SOACS -> ADM (Lambda SOACS)-fwdStreamLambda l@(Lambda params body ret) =-  Lambda <$> ((take 1 params ++) <$> bundleNewList (drop 1 params)) <*> inScopeOf l (fwdBody body) <*> bundleTangents ret+fwdStreamLambda l@(Lambda params ret body) =+  Lambda <$> ((take 1 params ++) <$> bundleNewList (drop 1 params)) <*> bundleTangents ret <*> inScopeOf l (fwdBody body)  interleave :: [a] -> [a] -> [a] interleave xs ys = concat $ transpose [xs, ys]@@ -310,12 +310,12 @@       mapM_ fwdStm stms       let (res_is, res_vs) = splitAt n_indices res       (res_is ++) <$> bundleTangents res_vs-    fwdHistBucket l@(Lambda params body ret) =+    fwdHistBucket l@(Lambda params ret body) =       let (r_is, r_vs) = splitAt n_indices ret        in Lambda             <$> bundleNewList params-            <*> inScopeOf l (fwdBodyHist body)             <*> ((r_is ++) <$> bundleTangents r_vs)+            <*> inScopeOf l (fwdBodyHist body)      fwdHist :: HistOp SOACS -> ADM (HistOp SOACS)     fwdHist (HistOp shape rf dest nes op) = do@@ -341,13 +341,13 @@   addStm s   where     fwdScatterLambda :: Int -> Lambda SOACS -> ADM (Lambda SOACS)-    fwdScatterLambda n_indices (Lambda params body ret) = do+    fwdScatterLambda n_indices (Lambda params ret body) = do       params' <- bundleNewList params       ret_tan <- mapM tangent $ drop n_indices ret       body' <- fwdBodyScatter n_indices body       let indices = concat $ replicate 2 $ take n_indices ret           ret' = indices ++ drop n_indices ret ++ ret_tan-      pure $ Lambda params' body' ret'+      pure $ Lambda params' ret' body'     fwdBodyScatter :: Int -> Body SOACS -> ADM (Body SOACS)     fwdBodyScatter n_indices (Body _ stms res) = do       (res_tan, stms') <- collectStms $ do@@ -406,17 +406,16 @@   val_pats' <- bundleNewList val_pats   pat' <- bundleNewPat pat   body' <--    localScope (scopeOfFParams (map fst val_pats) <> scopeOf loop) . slocal' $+    localScope (scopeOfFParams (map fst val_pats) <> scopeOfLoopForm loop) . slocal' $       fwdBody body   addStm $ Let pat' aux $ Loop val_pats' (WhileLoop v) body'-fwdStm (Let pat aux (Loop val_pats loop@(ForLoop i it bound loop_vars) body)) = do+fwdStm (Let pat aux (Loop val_pats loop@(ForLoop i it bound) body)) = do   pat' <- bundleNewPat pat   val_pats' <- bundleNewList val_pats-  loop_vars' <- bundleNewList loop_vars   body' <--    localScope (scopeOfFParams (map fst val_pats) <> scopeOf loop) . slocal' $+    localScope (scopeOfFParams (map fst val_pats) <> scopeOfLoopForm loop) . slocal' $       fwdBody body-  addStm $ Let pat' aux $ Loop val_pats' (ForLoop i it bound loop_vars') body'+  addStm $ Let pat' aux $ Loop val_pats' (ForLoop i it bound) body' fwdStm (Let pat aux (WithAcc inputs lam)) = do   inputs' <- forM inputs $ \(shape, arrs, op) -> do     arrs_tan <- mapM tangent arrs@@ -426,8 +425,8 @@         nes_tan <- mapM (fmap Var . zeroFromSubExp) nes         op_lam' <- fwdLambda op_lam         case op_lam' of-          Lambda ps body ret -> do-            let op_lam'' = Lambda (removeIndexTans (shapeRank shape) ps) body ret+          Lambda ps ret body -> do+            let op_lam'' = Lambda (removeIndexTans (shapeRank shape) ps) ret body             pure $ Just (op_lam'', interleave nes nes_tan)     pure (shape, arrs <> arrs_tan, op')   pat' <- bundleNewPat pat@@ -452,9 +451,9 @@   (res <>) <$> mapM tangent res  fwdJVP :: (MonadFreshNames m) => Scope SOACS -> Lambda SOACS -> m (Lambda SOACS)-fwdJVP scope l@(Lambda params body ret) =+fwdJVP scope l@(Lambda params ret body) =   runADM . localScope scope . inScopeOf l $ do     params_tan <- mapM newTan params     body_tan <- fwdBodyTansLast body     ret_tan <- mapM tangent ret-    pure $ Lambda (params ++ params_tan) body_tan (ret <> ret_tan)+    pure $ Lambda (params ++ params_tan) (ret <> ret_tan) body_tan

@@ -300,12 +300,12 @@       f' <- renameLambda f       pure (shape, as, Just (f', nes))     renameInputLambda input = pure input-    diffLambda' res_adjs get_adjs_for (Lambda params body ts) =+    diffLambda' res_adjs get_adjs_for (Lambda params ts body) =       localScope (scopeOfLParams params) $ do         Body () stms res <- diffBody res_adjs get_adjs_for body         let body' = Body () stms $ take (length inputs) res <> takeLast (length get_adjs_for) res         ts' <- mapM lookupType get_adjs_for-        pure $ Lambda params body' $ take (length inputs) ts <> ts'+        pure $ Lambda params (take (length inputs) ts <> ts') body' diffStm stm _ = error $ "diffStm unhandled:\n" ++ prettyString stm  diffStms :: Stms SOACS -> ADM ()@@ -336,15 +336,15 @@     pure $ Body () stms' $ res <> varsRes adjs  diffLambda :: [Adj] -> [VName] -> Lambda SOACS -> ADM (Lambda SOACS)-diffLambda res_adjs get_adjs_for (Lambda params body _) =+diffLambda res_adjs get_adjs_for (Lambda params _ body) =   localScope (scopeOfLParams params) $ do     Body () stms res <- diffBody res_adjs get_adjs_for body     let body' = Body () stms $ takeLast (length get_adjs_for) res     ts' <- mapM lookupType get_adjs_for-    pure $ Lambda params body' ts'+    pure $ Lambda params ts' body'  revVJP :: (MonadFreshNames m) => Scope SOACS -> Lambda SOACS -> m (Lambda SOACS)-revVJP scope (Lambda params body ts) =+revVJP scope (Lambda params ts body) =   runADM . localScope (scope <> scopeOfLParams params) $ do     params_adj <- forM (zip (map resSubExp (bodyResult body)) ts) $ \(se, t) ->       Param mempty <$> maybe (newVName "const_adj") adjVName (subExpVar se) <*> pure t@@ -356,7 +356,7 @@           (map paramName params)           body -    pure $ Lambda (params ++ params_adj) body' (ts <> map paramType params)+    pure $ Lambda (params ++ params_adj) (ts <> map paramType params) body'  -- Note [Adjoints of accumulators] --

@@ -646,7 +646,7 @@   letTupExp "sorted" $     Loop       (zip params $ map Var xs)-      (ForLoop i Int64 iters [])+      (ForLoop i Int64 iters)       loopbody   where     log2 :: SubExp -> ADM SubExp

@@ -24,17 +24,16 @@   (PrettyRep rep) =>   Exp rep ->   ( [(Param (FParamInfo rep), SubExp)] ->-    LoopForm rep ->+    LoopForm ->     VName ->     IntType ->     SubExp ->-    [(Param (LParamInfo rep), VName)] ->     Body rep ->     a   ) ->   a-bindForLoop (Loop val_pats form@(ForLoop i it bound loop_vars) body) f =-  f val_pats form i it bound loop_vars body+bindForLoop (Loop val_pats form@(ForLoop i it bound) body) f =+  f val_pats form i it bound body bindForLoop e _ = error $ "bindForLoop: not a for-loop:\n" <> prettyString e  -- | A convenience function to rename a for-loop and then bind the@@ -44,11 +43,10 @@   Exp rep ->   ( Exp rep ->     [(Param (FParamInfo rep), SubExp)] ->-    LoopForm rep ->+    LoopForm ->     VName ->     IntType ->     SubExp ->-    [(Param (LParamInfo rep), VName)] ->     Body rep ->     m a   ) ->@@ -60,21 +58,6 @@ isWhileLoop (Loop _ WhileLoop {} _) = True isWhileLoop _ = False --- | Transforms a 'ForLoop' into a 'ForLoop' with an empty list of--- loop variables.-removeLoopVars :: (MonadBuilder m) => Exp (Rep m) -> m (Exp (Rep m))-removeLoopVars loop =-  bindForLoop loop $ \val_pats form i _it _bound loop_vars body -> do-    let indexify (x_param, xs) = do-          xs_t <- lookupType xs-          x' <--            letExp (baseString $ paramName x_param) . BasicOp . Index xs $-              fullSlice xs_t [DimFix (Var i)]-          pure (paramName x_param, x')-    (substs_list, subst_stms) <- collectStms $ mapM indexify loop_vars-    let Body aux' stms' res' = substituteNames (M.fromList substs_list) body-    pure $ Loop val_pats form $ Body aux' (subst_stms <> stms') res'- -- | Augments a while-loop to also compute the number of iterations. computeWhileIters :: Exp SOACS -> ADM SubExp computeWhileIters (Loop val_pats (WhileLoop b) body) = do@@ -109,7 +92,7 @@             (pure body)             (resultBodyM $ map (Var . paramName . fst) val_pats)         ]-    pure $ Loop val_pats (ForLoop i Int64 bound_se mempty) body'+    pure $ Loop val_pats (ForLoop i Int64 bound_se) body' convertWhileLoop _ e = error $ "convertWhileLoopBound: not a while-loop:\n" <> prettyString e  -- | @nestifyLoop n bound loop@ transforms a loop into a depth-@n@ loop nest@@ -125,9 +108,9 @@   where     nestifyLoop' offset n loop = bindForLoop loop nestify       where-        nestify val_pats _form i it _bound loop_vars body+        nestify val_pats _form i it _bound body           | n > 1 = do-              renameForLoop loop $ \_loop' val_pats' _form' i' it' _bound' loop_vars' body' -> do+              renameForLoop loop $ \_loop' val_pats' _form' i' it' _bound' body' -> do                 let loop_params = map fst val_pats                     loop_params' = map fst val_pats'                     loop_inits' = map (Var . paramName) loop_params@@ -149,12 +132,11 @@                         =<< nestifyLoop'                           offset'                           (n - 1)-                          (Loop val_pats'' (ForLoop i' it' bound_se loop_vars') inner_body)+                          (Loop val_pats'' (ForLoop i' it' bound_se) inner_body)                     pure $ varsRes inner_loop-                pure $-                  Loop val_pats (ForLoop i it bound_se loop_vars) outer_body+                pure $ Loop val_pats (ForLoop i it bound_se) outer_body           | n == 1 =-              pure $ Loop val_pats (ForLoop i it bound_se loop_vars) body+              pure $ Loop val_pats (ForLoop i it bound_se) body           | otherwise = pure loop  -- | @stripmine n pat loop@ stripmines a loop into a depth-@n@ loop nest.@@ -163,8 +145,7 @@ -- so that the stripmined loop is semantically equivalent to the original loop. stripmine :: Integer -> Pat Type -> Exp SOACS -> ADM (Stms SOACS) stripmine n pat loop = do-  loop' <- removeLoopVars loop-  bindForLoop loop' $ \_val_pats _form _i it bound _loop_vars _body -> do+  bindForLoop loop $ \_val_pats _form _i it bound _body -> do     let n_root = Constant $ FloatValue $ floatValue Float64 (1 / fromIntegral n :: Double)     bound_float <- letSubExp "bound_f64" $ BasicOp $ ConvOp (UIToFP it Float64) bound     bound' <- letSubExp "bound" $ BasicOp $ BinOp (FPow Float64) bound_float n_root@@ -176,7 +157,7 @@       letSubExp "remain_iters" $ BasicOp $ BinOp (Sub it OverflowUndef) bound total_iters     mined_loop <- nestifyLoop bound_int n loop     pat' <- renamePat pat-    renameForLoop loop $ \_loop' val_pats' _form' i' it' _bound' loop_vars' body' -> do+    renameForLoop loop $ \_loop val_pats' _form' i' it' _bound' body' -> do       remain_body <- insertStmsM $ do         i_remain <-           letExp "i_remain" . BasicOp $@@ -185,7 +166,7 @@       let loop_params_rem = map fst val_pats'           loop_inits_rem = map (Var . patElemName) $ patElems pat'           val_pats_rem = zip loop_params_rem loop_inits_rem-          remain_loop = Loop val_pats_rem (ForLoop i' it' remain_iters loop_vars') remain_body+          remain_loop = Loop val_pats_rem (ForLoop i' it' remain_iters) remain_body       collectStms_ $ do         letBind pat' mined_loop         letBind pat remain_loop@@ -213,7 +194,7 @@ -- the originals (which will be consumed later in the reverse pass). fwdLoop :: Pat Type -> StmAux () -> Exp SOACS -> ADM () fwdLoop pat aux loop =-  bindForLoop loop $ \val_pats form i _it bound _loop_vars body -> do+  bindForLoop loop $ \val_pats form i _it bound body -> do     bound64 <- asIntS Int64 bound     let loop_params = map fst val_pats         is_true_dep = inAttrs (AttrName "true_dep") . paramAttrs@@ -228,7 +209,7 @@      (body', (saved_pats, saved_params)) <- buildBody $       localScope (scopeOfFParams loop_params) $-        inScopeOf form $ do+        localScope (scopeOfLoopForm form) $ do           copy_substs <- copyConsumedArrsInBody dont_copy body           addStms $ bodyStms body           i_i64 <- asIntS Int64 $ Var i@@ -270,34 +251,22 @@ valPatAdjs :: LoopInfo [VName] -> ADM (LoopInfo [(Param DeclType, SubExp)]) valPatAdjs = (mapM . mapM) valPatAdj --- | Reverses a loop by substituting the loop index as well as reversing--- the arrays that loop variables are bound to.-reverseIndices :: Exp SOACS -> ADM (Substitutions, Substitutions, Stms SOACS)+-- | Reverses a loop by substituting the loop index.+reverseIndices :: Exp SOACS -> ADM (Substitutions, Stms SOACS) reverseIndices loop = do-  bindForLoop loop $ \_val_pats form i it bound loop_vars _body -> do+  bindForLoop loop $ \_val_pats form i it bound _body -> do     bound_minus_one <--      inScopeOf form $+      localScope (scopeOfLoopForm form) $         let one = Constant $ IntValue $ intValue it (1 :: Int)          in letSubExp "bound-1" $ BasicOp $ BinOp (Sub it OverflowUndef) bound one -    var_arrays_substs <- fmap M.fromList $-      inScopeOf form $ do-        forM (map snd loop_vars) $ \xs -> do-          xs_t <- lookupType xs-          xs_rev <--            letExp "reverse" . BasicOp . Index xs $-              fullSlice-                xs_t-                [DimSlice bound_minus_one bound (Constant (IntValue (Int64Value (-1))))]-          pure (xs, xs_rev)-     (i_rev, i_stms) <- collectStms $-      inScopeOf form $ do+      localScope (scopeOfLoopForm form) $ do         letExp (baseString i <> "_rev") $           BasicOp $             BinOp (Sub it OverflowWrap) bound_minus_one (Var i) -    pure (var_arrays_substs, M.singleton i i_rev, i_stms)+    pure (M.singleton i i_rev, i_stms)  -- | Pures a substitution which substitutes values in the reverse -- loop body with values from the tape.@@ -332,7 +301,6 @@ data LoopInfo a = LoopInfo   { loopRes :: a,     loopFree :: a,-    loopVars :: a,     loopVals :: a   }   deriving (Functor, Foldable, Traversable, Show)@@ -340,20 +308,18 @@ -- | Transforms a for-loop into its reverse-mode derivative. revLoop :: (Stms SOACS -> ADM ()) -> Pat Type -> Exp SOACS -> ADM () revLoop diffStms pat loop =-  bindForLoop loop $ \val_pats _form _i _it _bound _loop_vars _body ->+  bindForLoop loop $ \val_pats _form _i _it _bound _body ->     renameForLoop loop $-      \loop' val_pats' form' i' _it' _bound' loop_vars' body' -> do+      \loop' val_pats' form' i' _it' _bound' body' -> do         let loop_params = map fst val_pats             (loop_params', loop_vals') = unzip val_pats'-            loop_var_arrays' = map snd loop_vars'             getVName Constant {} = Nothing             getVName (Var v) = Just v             loop_vnames =               LoopInfo                 { loopRes = mapMaybe subExpResVName $ bodyResult body',                   loopFree =-                    (namesToList (freeIn loop') \\ loop_var_arrays') \\ mapMaybe getVName loop_vals',-                  loopVars = loop_var_arrays',+                    namesToList (freeIn loop') \\ mapMaybe getVName loop_vals',                   loopVals = nubOrd $ mapMaybe getVName loop_vals'                 } @@ -364,50 +330,37 @@             Just v -> setAdj v =<< lookupAdj (patElemName pe)             Nothing -> pure () -        (var_array_substs, i_subst, i_stms) <--          reverseIndices loop'+        (i_subst, i_stms) <- reverseIndices loop'          val_pat_adjs <- valPatAdjs loop_vnames         let val_pat_adjs_list = concat $ toList val_pat_adjs          (loop_adjs, stms_adj) <- collectStms $-          inScopeOf form' $-            localScope (scopeOfFParams (map fst val_pat_adjs_list <> loop_params')) $ do-              addStms i_stms-              (loop_adjs, stms_adj) <- collectStms $-                subAD $ do-                  zipWithM_-                    (\val_pat v -> insAdj v (paramName $ fst val_pat))-                    val_pat_adjs_list-                    (concat $ toList loop_vnames)-                  diffStms $ bodyStms body'--                  let update_var_arrays v vs = do-                        vs_t <- lookupType vs-                        v_adj <- lookupAdjVal v-                        updateAdjSlice (fullSlice vs_t [DimFix $ Var i']) vs v_adj-                  zipWithM_-                    update_var_arrays-                    (map (paramName . fst) loop_vars')-                    (loopVars loop_vnames)+          localScope (scopeOfLoopForm form' <> scopeOfFParams (map fst val_pat_adjs_list <> loop_params')) $ do+            addStms i_stms+            (loop_adjs, stms_adj) <- collectStms $+              subAD $ do+                zipWithM_+                  (\val_pat v -> insAdj v (paramName $ fst val_pat))+                  val_pat_adjs_list+                  (concat $ toList loop_vnames)+                diffStms $ bodyStms body' -                  loop_res_adjs <- mapM (lookupAdjVal . paramName) loop_params'-                  loop_free_adjs <- mapM lookupAdjVal $ loopFree loop_vnames-                  loop_vars_adjs <- mapM lookupAdjVal $ loopVars loop_vnames-                  loop_vals_adjs <- mapM lookupAdjVal $ loopVals loop_vnames+                loop_res_adjs <- mapM (lookupAdjVal . paramName) loop_params'+                loop_free_adjs <- mapM lookupAdjVal $ loopFree loop_vnames+                loop_vals_adjs <- mapM lookupAdjVal $ loopVals loop_vnames -                  pure $-                    LoopInfo-                      { loopRes = loop_res_adjs,-                        loopFree = loop_free_adjs,-                        loopVars = loop_vars_adjs,-                        loopVals = loop_vals_adjs-                      }-              (substs, restore_stms) <--                collectStms $ restore stms_adj loop_params' i'-              addStms $ substituteNames i_subst restore_stms-              addStms $ substituteNames i_subst $ substituteNames substs stms_adj-              pure loop_adjs+                pure $+                  LoopInfo+                    { loopRes = loop_res_adjs,+                      loopFree = loop_free_adjs,+                      loopVals = loop_vals_adjs+                    }+            (substs, restore_stms) <-+              collectStms $ restore stms_adj loop_params' i'+            addStms $ substituteNames i_subst restore_stms+            addStms $ substituteNames i_subst $ substituteNames substs stms_adj+            pure loop_adjs          inScopeOf stms_adj $           localScope (scopeOfFParams $ map fst val_pat_adjs_list) $ do@@ -419,18 +372,15 @@                       else Nothing             adjs' <-               letTupExp "loop_adj" $-                substituteNames (restore_true_deps <> var_array_substs) $+                substituteNames restore_true_deps $                   Loop val_pat_adjs_list form' body_adj             let (loop_res_adjs, loop_free_var_val_adjs) =                   splitAt (length $ loopRes loop_adjs) adjs'-                (loop_free_adjs, loop_var_val_adjs) =+                (loop_free_adjs, loop_val_adjs) =                   splitAt (length $ loopFree loop_adjs) loop_free_var_val_adjs-                (loop_var_adjs, loop_val_adjs) =-                  splitAt (length $ loopVars loop_adjs) loop_var_val_adjs             returnSweepCode $ do               zipWithM_ updateSubExpAdj loop_vals' loop_res_adjs               zipWithM_ insAdj (loopFree loop_vnames) loop_free_adjs-              zipWithM_ insAdj (loopVars loop_vnames) loop_var_adjs               zipWithM_ updateAdj (loopVals loop_vnames) loop_val_adjs  -- | Transforms a loop into its reverse-mode derivative.

@@ -328,7 +328,7 @@         res <- tabNest' (paramName iparam : is) (n - 1) (map paramName params) f         ret <- mapM lookupType res         pure (ret, varsRes res)-      let lam = Lambda (iparam : params) (Body () stms res) ret+      let lam = Lambda (iparam : params) ret (Body () stms res)       letTupExp "tab" $ Op $ Screma w (iota : vs) (mapSOAC lam)  -- | Construct a lambda for adding two values of the given type.

@@ -216,7 +216,7 @@ --   if the lambda argument is a map, then returns --   just the map's lambda; otherwise nothing. mapOp :: Lambda SOACS -> Maybe (Lambda SOACS)-mapOp (Lambda [pa1, pa2] lam_body _)+mapOp (Lambda [pa1, pa2] _ lam_body)   | [SubExpRes cs r] <- bodyResult lam_body,     cs == mempty,     [map_stm] <- stmsToList (bodyStms lam_body),

@@ -2,6 +2,12 @@ module Futhark.Analysis.DataDependencies   ( Dependencies,     dataDependencies,+    depsOf,+    depsOf',+    depsOfArrays,+    depsOfShape,+    lambdaDependencies,+    reductionDependencies,     findNecessaryForReturned,   ) where@@ -27,6 +33,23 @@   Dependencies dataDependencies' startdeps = foldl grow startdeps . bodyStms   where+    grow deps (Let pat _ (WithAcc inputs lam)) =+      let input_deps = foldMap depsOfWithAccInput inputs+          -- Dependencies of each input reduction are concatenated.+          -- Input to lam is cert_1, ..., cert_n, acc_1, ..., acc_n.+          lam_deps = lambdaDependencies deps lam (input_deps <> input_deps)+          transitive = map (depsOfNames deps) lam_deps+       in M.fromList (zip (patNames pat) transitive) `M.union` deps+      where+        depsOfArrays' shape =+          map (\arr -> oneName arr <> depsOfShape shape)+        depsOfWithAccInput (shape, arrs, Nothing) =+          depsOfArrays' shape arrs+        depsOfWithAccInput (shape, arrs, Just (lam', nes)) =+          reductionDependencies deps lam' nes (depsOfArrays' shape arrs)+    grow deps (Let pat _ (Op op)) =+      let op_deps = map (depsOfNames deps) (opDependencies op)+       in M.fromList (zip (patNames pat) op_deps) `M.union` deps     grow deps (Let pat _ (Match c cases defbody _)) =       let cases_deps = map (dataDependencies' deps . caseBody) cases           defbody_deps = dataDependencies' deps defbody@@ -50,15 +73,61 @@        in M.unions $ [branchdeps, deps, defbody_deps] ++ cases_deps     grow deps (Let pat _ e) =       let free = freeIn pat <> freeIn e-          freeDeps = mconcat $ map (depsOfVar deps) $ namesToList free-       in M.fromList [(name, freeDeps) | name <- patNames pat] `M.union` deps+          free_deps = depsOfNames deps free+       in M.fromList [(name, free_deps) | name <- patNames pat] `M.union` deps  depsOf :: Dependencies -> SubExp -> Names depsOf _ (Constant _) = mempty depsOf deps (Var v) = depsOfVar deps v +depsOf' :: SubExp -> Names+depsOf' (Constant _) = mempty+depsOf' (Var v) = depsOfVar mempty v+ depsOfVar :: Dependencies -> VName -> Names depsOfVar deps name = oneName name <> M.findWithDefault mempty name deps++depsOfRes :: Dependencies -> SubExpRes -> Names+depsOfRes deps (SubExpRes _ se) = depsOf deps se++-- | Extend @names@ with direct dependencies in @deps@.+depsOfNames :: Dependencies -> Names -> Names+depsOfNames deps names = mconcat $ map (depsOfVar deps) $ namesToList names++depsOfArrays :: SubExp -> [VName] -> [Names]+depsOfArrays size = map (\arr -> oneName arr <> depsOf mempty size)++depsOfShape :: Shape -> Names+depsOfShape shape = mconcat $ map (depsOf mempty) (shapeDims shape)++-- | Determine the variables on which the results of applying+-- anonymous function @lam@ to @inputs@ depend.+lambdaDependencies ::+  (ASTRep rep) =>+  Dependencies ->+  Lambda rep ->+  [Names] ->+  [Names]+lambdaDependencies deps lam inputs =+  let names_in_scope = freeIn lam <> mconcat inputs+      deps_in = M.fromList $ zip (boundByLambda lam) inputs+      deps' = dataDependencies' (deps_in <> deps) (lambdaBody lam)+   in map+        (namesIntersection names_in_scope . depsOfRes deps')+        (bodyResult $ lambdaBody lam)++-- | Like 'lambdaDependencies', but @lam@ is a binary operation+-- with a neutral element.+reductionDependencies ::+  (ASTRep rep) =>+  Dependencies ->+  Lambda rep ->+  [SubExp] ->+  [Names] ->+  [Names]+reductionDependencies deps lam nes inputs =+  let nes' = map (depsOf deps) nes+   in lambdaDependencies deps lam (zipWith (<>) nes' inputs)  -- | @findNecessaryForReturned p merge deps@ computes which of the -- loop parameters (@merge@) are necessary for the result of the loop,

@@ -76,7 +76,8 @@ import Data.Sequence qualified as Seq import Futhark.Construct hiding (toExp) import Futhark.IR hiding-  ( Iota,+  ( Index,+    Iota,     Rearrange,     Replicate,     Reshape,@@ -107,6 +108,8 @@     ReshapeInner Certs ReshapeKind Shape   | -- | Replicate the rows of the array a number of times.     Replicate Certs Shape+  | -- | An array indexing operation.+    Index Certs (Slice SubExp)   deriving (Show, Eq, Ord)  instance Substitute ArrayTransform where@@ -120,6 +123,8 @@     ReshapeInner (substituteNames substs cs) k (substituteNames substs ses)   substituteNames substs (Replicate cs se) =     Replicate (substituteNames substs cs) (substituteNames substs se)+  substituteNames substs (Index cs slice) =+    Index (substituteNames substs cs) (substituteNames substs slice)  -- | A sequence of array transformations, heavily inspired by -- "Data.Seq".  You can decompose it using 'viewf' and 'viewl', and@@ -232,6 +237,8 @@   Just (v, Reshape cs k shape) transformFromExp cs (BasicOp (Futhark.Replicate shape (Var v))) =   Just (v, Replicate cs shape)+transformFromExp cs (BasicOp (Futhark.Index v slice)) =+  Just (v, Index cs slice) transformFromExp _ _ = Nothing  -- | Turn an array transform on an array back into an expression.@@ -249,6 +256,8 @@ transformToExp (ReshapeInner cs k shape) ia = do   shape' <- reshapeInner shape 1 . arrayShape <$> lookupType ia   pure (cs, BasicOp $ Futhark.Reshape k shape' ia)+transformToExp (Index cs slice) ia = do+  pure (cs, BasicOp $ Futhark.Index ia slice)  -- | One array input to a SOAC - a SOAC may have multiple inputs, but -- all are of this form.  Only the array inputs are expressed with@@ -312,6 +321,7 @@       Reshape {} -> "reshape"       ReshapeOuter {} -> "reshape_outer"       ReshapeInner {} -> "reshape_inner"+      Index {} -> "index"  applyTransforms :: (MonadBuilder m) => ArrayTransforms -> VName -> m VName applyTransforms (ArrayTransforms ts) a = foldlM (flip applyTransform) a ts@@ -350,6 +360,8 @@     transformType t (ReshapeInner _ _ shape) =       let Shape oldshape = arrayShape t        in t `setArrayShape` Shape (take 1 oldshape ++ shapeDims shape)+    transformType t (Index _ slice) =+      t `setArrayShape` sliceShape slice  -- | Return the row type of an input.  Just a convenient alias. inputRowType :: Input -> Type@@ -396,32 +408,6 @@   | Hist SubExp [HistOp rep] (Lambda rep) [Input]   deriving (Eq, Show) -instance PP.Pretty Input where-  pretty (Input (ArrayTransforms ts) arr _) = foldl f (pretty arr) ts-    where-      f e (Rearrange cs perm) =-        "rearrange" <> pretty cs <> PP.apply [PP.apply (map pretty perm), e]-      f e (Reshape cs ReshapeArbitrary shape) =-        "reshape" <> pretty cs <> PP.apply [pretty shape, e]-      f e (ReshapeOuter cs ReshapeArbitrary shape) =-        "reshape_outer" <> pretty cs <> PP.apply [pretty shape, e]-      f e (ReshapeInner cs ReshapeArbitrary shape) =-        "reshape_inner" <> pretty cs <> PP.apply [pretty shape, e]-      f e (Reshape cs ReshapeCoerce shape) =-        "coerce" <> pretty cs <> PP.apply [pretty shape, e]-      f e (ReshapeOuter cs ReshapeCoerce shape) =-        "coerce_outer" <> pretty cs <> PP.apply [pretty shape, e]-      f e (ReshapeInner cs ReshapeCoerce shape) =-        "coerce_inner" <> pretty cs <> PP.apply [pretty shape, e]-      f e (Replicate cs ne) =-        "replicate" <> pretty cs <> PP.apply [pretty ne, e]--instance (PrettyRep rep) => PP.Pretty (SOAC rep) where-  pretty (Screma w form arrs) = Futhark.ppScrema w arrs form-  pretty (Hist len ops bucket_fun imgs) = Futhark.ppHist len imgs ops bucket_fun-  pretty (Stream w lam nes arrs) = Futhark.ppStream w arrs nes lam-  pretty (Scatter w lam arrs dests) = Futhark.ppScatter w arrs lam dests- -- | Returns the inputs used in a SOAC. inputs :: SOAC rep -> [Input] inputs (Stream _ _ _ arrs) = arrs@@ -572,7 +558,7 @@               insstm = mkLet strm_resids $ Op insoac               strmbdy = mkBody (oneStm insstm) $ map (subExpRes . Var . identName) strm_resids               strmpar = chunk_param : strm_inpids-              strmlam = Lambda strmpar strmbdy loutps+              strmlam = Lambda strmpar loutps strmbdy           -- map(f,a) creates a stream with NO accumulators           pure (Stream w strmlam [] inps, [])       | Just (scans, _) <- Futhark.isScanomapSOAC form,@@ -614,9 +600,7 @@                   outszm1id                   (constant (0 :: Int64))             -- 3. let lasteel_ids = ...-            let indexLast arr = do-                  arr_t <- lookupType arr-                  pure $ BasicOp . Index arr $ fullSlice arr_t [DimFix outszm1id]+            let indexLast arr = eIndex arr [eSubExp outszm1id]             lastel_ids <-               letTupExp "lastel"                 =<< eIf@@ -671,7 +655,7 @@                 mkBody (oneStm insstm <> addaccstm) $                   addaccres ++ map (subExpRes . Var . identName) strm_resids               strmpar = chunk_param : inpacc_ids ++ strm_inpids-              strmlam = Lambda strmpar strmbdy (accrtps ++ loutps')+              strmlam = Lambda strmpar (accrtps ++ loutps') strmbdy           pure (Stream w strmlam nes inps, [])      -- Otherwise it cannot become a stream.@@ -695,7 +679,7 @@               (bodyDec plus_bdy)               (stmsFromList parstms <> bodyStms plus_bdy)               (bodyResult plus_bdy)-      renameLambda $ Lambda rempars newlambdy $ lambdaReturnType plus+      renameLambda $ Lambda rempars (lambdaReturnType plus) newlambdy      mkPlusBnds ::       (MonadFreshNames m, Buildable rep) =>@@ -711,3 +695,35 @@               accels           body = lambdaBody plus'       pure $ body {bodyStms = stmsFromList parstms <> bodyStms body}++ppArrayTransform :: PP.Doc a -> ArrayTransform -> PP.Doc a+ppArrayTransform e (Rearrange cs perm) =+  "rearrange" <> pretty cs <> PP.apply [PP.apply (map pretty perm), e]+ppArrayTransform e (Reshape cs ReshapeArbitrary shape) =+  "reshape" <> pretty cs <> PP.apply [pretty shape, e]+ppArrayTransform e (ReshapeOuter cs ReshapeArbitrary shape) =+  "reshape_outer" <> pretty cs <> PP.apply [pretty shape, e]+ppArrayTransform e (ReshapeInner cs ReshapeArbitrary shape) =+  "reshape_inner" <> pretty cs <> PP.apply [pretty shape, e]+ppArrayTransform e (Reshape cs ReshapeCoerce shape) =+  "coerce" <> pretty cs <> PP.apply [pretty shape, e]+ppArrayTransform e (ReshapeOuter cs ReshapeCoerce shape) =+  "coerce_outer" <> pretty cs <> PP.apply [pretty shape, e]+ppArrayTransform e (ReshapeInner cs ReshapeCoerce shape) =+  "coerce_inner" <> pretty cs <> PP.apply [pretty shape, e]+ppArrayTransform e (Replicate cs ne) =+  "replicate" <> pretty cs <> PP.apply [pretty ne, e]+ppArrayTransform e (Index cs slice) =+  e <> pretty cs <> pretty slice++instance PP.Pretty Input where+  pretty (Input (ArrayTransforms ts) arr _) = foldl ppArrayTransform (pretty arr) ts++instance PP.Pretty ArrayTransform where+  pretty = ppArrayTransform "INPUT"++instance (PrettyRep rep) => PP.Pretty (SOAC rep) where+  pretty (Screma w form arrs) = Futhark.ppScrema w arrs form+  pretty (Hist len ops bucket_fun imgs) = Futhark.ppHist len imgs ops bucket_fun+  pretty (Stream w lam nes arrs) = Futhark.ppStream w arrs nes lam+  pretty (Scatter w lam arrs dests) = Futhark.ppScatter w arrs lam dests

@@ -199,8 +199,7 @@   InUse ->   HistOp GPUMem ->   m (InUse, LastUsed, Graph VName)-analyseHistOp lumap inuse histop =-  analyseLambda lumap inuse (histOp histop)+analyseHistOp lumap inuse = analyseLambda lumap inuse . histOp  analyseLambda ::   (LocalScope GPUMem m) =>@@ -208,8 +207,7 @@   InUse ->   Lambda GPUMem ->   m (InUse, LastUsed, Graph VName)-analyseLambda lumap inuse (Lambda _ body _) =-  analyseBody lumap inuse body+analyseLambda lumap inuse = analyseBody lumap inuse . lambdaBody  analyseProgGPU :: Prog GPUMem -> Graph VName analyseProgGPU prog = onConsts (progConsts prog) <> foldMap onFun (progFuns prog)

@@ -253,7 +253,7 @@   let used_nms' = used_cases <> body_used_nms   (_, last_used_arrs) <- lastUsedInNames used_nms $ free_in_body <> free_in_cases   pure (lutab_cases <> lutab', last_used_arrs, used_nms')-lastUseExp (Loop var_ses lf body) used_nms0 = inScopeOf lf $ do+lastUseExp (Loop var_ses form body) used_nms0 = localScope (scopeOfLoopForm form) $ do   free_in_body <- aliasTransitiveClosure $ freeIn body   -- compute the aliasing transitive closure of initializers that are not last-uses   var_inis <- catMaybes <$> mapM (initHelper (free_in_body <> used_nms0)) var_ses@@ -364,7 +364,7 @@   (lutab, lu_vars, used_nms') <- unzip3 <$> mapM helper sbos   pure (mconcat lutab, mconcat lu_vars, mconcat used_nms')   where-    helper (SegBinOp _ l@(Lambda _ body _) neutral shp) = inScopeOf l $ do+    helper (SegBinOp _ l@(Lambda _ _ body) neutral shp) = inScopeOf l $ do       (used_nms', lu_vars) <- lastUsedInNames used_nms $ freeIn neutral <> freeIn shp       (body_lutab, used_nms'') <- lastUseBody body (mempty, used_nms')       pure (body_lutab, lu_vars, used_nms'')@@ -378,7 +378,7 @@   (lutab, lu_vars, used_nms') <- unzip3 <$> mapM helper hos   pure (mconcat lutab, mconcat lu_vars, mconcat used_nms')   where-    helper (HistOp shp rf dest neutral shp' l@(Lambda _ body _)) = inScopeOf l $ do+    helper (HistOp shp rf dest neutral shp' l@(Lambda _ _ body)) = inScopeOf l $ do       (used_nms', lu_vars) <- lastUsedInNames used_nms $ freeIn shp <> freeIn rf <> freeIn dest <> freeIn neutral <> freeIn shp'       (body_lutab, used_nms'') <- lastUseBody body (mempty, used_nms')       pure (body_lutab, lu_vars, used_nms'')

@@ -12,6 +12,7 @@     RunOptions (..),     prepareBenchmarkProgram,     CompileOptions (..),+    module Futhark.Profile,   ) where @@ -30,6 +31,7 @@ import Data.Text qualified as T import Data.Time.Clock import Data.Vector.Unboxed qualified as U+import Futhark.Profile import Futhark.Server import Futhark.Test import Futhark.Util (showText)@@ -47,20 +49,30 @@ -- | The measurements resulting from various successful runs of a -- benchmark (same dataset). data Result = Result-  { runResults :: [RunResult],+  { -- | Runtime of every run.+    runResults :: [RunResult],+    -- | Memory usage.     memoryMap :: M.Map T.Text Int,-    stdErr :: Maybe T.Text+    -- | The error output produced during execution.  Often 'Nothing'+    -- for space reasons, and otherwise only the output from the last+    -- run.+    stdErr :: Maybe T.Text,+    -- | Profiling report.  This will have been measured based on the+    -- last run.+    report :: Maybe ProfilingReport   }   deriving (Eq, Show)  -- | The results for a single named dataset is either an error message, or--- runtime measurements, the number of bytes used, and the stderr that was--- produced.+-- a result. data DataResult = DataResult T.Text (Either T.Text Result)   deriving (Eq, Show)  -- | The results for all datasets for some benchmark program.-data BenchResult = BenchResult FilePath [DataResult]+data BenchResult = BenchResult+  { benchResultProg :: FilePath,+    benchResultResults :: [DataResult]+  }   deriving (Eq, Show)  newtype DataResults = DataResults {unDataResults :: [DataResult]}@@ -68,10 +80,13 @@ newtype BenchResults = BenchResults {unBenchResults :: [BenchResult]}  instance JSON.ToJSON Result where-  toJSON (Result runres memmap err) = JSON.toJSON (runres, memmap, err)+  toJSON (Result runres memmap err profiling) =+    JSON.toJSON (runres, memmap, err, profiling)  instance JSON.FromJSON Result where-  parseJSON = fmap (\(runres, memmap, err) -> Result runres memmap err) . JSON.parseJSON+  parseJSON = fmap f . JSON.parseJSON+    where+      f (runres, memmap, err, profiling) = Result runres memmap err profiling  instance JSON.ToJSON RunResult where   toJSON = JSON.toJSON . runMicroseconds@@ -93,20 +108,27 @@         DataResult (JSON.toText k)           <$> ((Right <$> success v) <|> (Left <$> JSON.parseJSON v))       success = JSON.withObject "result" $ \o ->-        Result <$> o JSON..: "runtimes" <*> o JSON..: "bytes" <*> o JSON..:? "stderr"+        Result+          <$> o JSON..: "runtimes"+          <*> o JSON..: "bytes"+          <*> o JSON..:? "stderr"+          <*> o JSON..:? "profiling"  dataResultJSON :: DataResult -> (JSON.Key, JSON.Value) dataResultJSON (DataResult desc (Left err)) =   (JSON.fromText desc, JSON.toJSON err)-dataResultJSON (DataResult desc (Right (Result runtimes bytes progerr_opt))) =+dataResultJSON (DataResult desc (Right (Result runtimes bytes progerr_opt profiling_opt))) =   ( JSON.fromText desc,     JSON.object $       [ ("runtimes", JSON.toJSON $ map runMicroseconds runtimes),         ("bytes", JSON.toJSON bytes)       ]-        ++ case progerr_opt of+        <> case progerr_opt of           Nothing -> []           Just progerr -> [("stderr", JSON.toJSON progerr)]+        <> case profiling_opt of+          Nothing -> []+          Just profiling -> [("profiling", JSON.toJSON profiling)]   )  benchResultJSON :: BenchResult -> (JSON.Key, JSON.Value)@@ -152,7 +174,10 @@     runConvergenceMaxTime :: NominalDiffTime,     -- | Invoked for every runtime measured during the run.  Can be     -- used to provide a progress bar.-    runResultAction :: (Int, Maybe Double) -> IO ()+    runResultAction :: (Int, Maybe Double) -> IO (),+    -- | Perform a final run at the end with profiling information+    -- enabled.+    runProfile :: Bool   }  -- | A list of @(autocorrelation,rse)@ pairs.  When the@@ -275,7 +300,7 @@   Values ->   Maybe Success ->   FilePath ->-  IO (Either T.Text ([RunResult], T.Text))+  IO (Either T.Text ([RunResult], T.Text, ProfilingReport)) benchmarkDataset server opts futhark program entry input_spec expected_spec ref_out = runExceptT $ do   output_types <- cmdEither $ cmdOutputs server entry   input_types <- cmdEither $ cmdInputs server entry@@ -284,6 +309,8 @@    cmdMaybe . liftIO $ cmdClear server +  cmdMaybe . liftIO $ cmdPauseProfiling server+   let freeOuts = cmdMaybe (cmdFree server outs)       freeIns = cmdMaybe (cmdFree server ins)       loadInput = valuesAsVars server (zip ins $ map inputType input_types) futhark dir input_spec@@ -314,11 +341,21 @@      xs <- runConvergence (freeOuts *> doRun) opts ys +    -- Possibly a profiled run at the end.+    profile_log <-+      if not (runProfile opts)+        then pure Nothing+        else do+          cmdMaybe . liftIO $ cmdUnpauseProfiling server+          profile_log <- freeOuts *> doRun+          cmdMaybe . liftIO $ cmdPauseProfiling server+          pure $ Just profile_log+     vs <- readResults server outs <* freeOuts -    pure (vs, DL.toList xs)+    pure (vs, DL.toList xs, profile_log) -  (vs, call_logs) <- case maybe_call_logs of+  (vs, call_logs, profile_log) <- case maybe_call_logs of     Nothing ->       throwError . T.pack $         "Execution exceeded " ++ show (runTimeout opts) ++ " seconds."@@ -328,6 +365,10 @@    report <- cmdEither $ cmdReport server +  report' <-+    maybe (throwError "Program produced invalid profiling report.") pure $+      profilingInfoFromText (T.unlines report)+   maybe_expected <-     liftIO $ maybe (pure Nothing) (fmap Just . getExpectedValues) expected_spec @@ -337,7 +378,8 @@    pure     ( map fst call_logs,-      T.unlines $ map (T.unlines . snd) call_logs <> report+      T.unlines $ snd $ fromMaybe (last call_logs) profile_log,+      report'     )   where     getExpectedValues (SuccessValues vs) =

@@ -222,7 +222,7 @@     res <- m     ret <- mapM subExpResType res     pure (res, ret)-  pure $ Lambda params (mkBody stms res) ret+  pure $ Lambda params ret $ mkBody stms res  -- Utility instance defintions for MTL classes.  These require -- UndecidableInstances, but save on typing elsewhere.

@@ -70,7 +70,8 @@       runVerbose = optVerbose opts,       runConvergencePhase = True,       runConvergenceMaxTime = fromIntegral timeout_s,-      runResultAction = const $ pure ()+      runResultAction = const $ pure (),+      runProfile = False     }  type Path = [(T.Text, Int)]@@ -167,8 +168,7 @@       pure (dataset, do_run, iosEntryPoint ios)   where     run server entry_point trun expected timeout path = do-      let bestRuntime :: ([RunResult], T.Text) -> ([(T.Text, Int)], Int)-          bestRuntime (runres, errout) =+      let bestRuntime (runres, errout, _) =             ( comparisons errout,               minimum $ map runMicroseconds runres             )

@@ -1,10 +1,10 @@ -- | @futhark bench@ module Futhark.CLI.Bench (main) where -import Control.Arrow (first) import Control.Exception import Control.Monad import Control.Monad.IO.Class (liftIO)+import Data.Bifunctor (first) import Data.ByteString.Char8 qualified as SBS import Data.ByteString.Lazy.Char8 qualified as LBS import Data.Either@@ -65,6 +65,7 @@     optConvergencePhase :: Bool,     optConvergenceMaxTime :: NominalDiffTime,     optConcurrency :: Maybe Int,+    optProfile :: Bool,     optVerbose :: Int,     optTestSpec :: Maybe FilePath   }@@ -90,6 +91,7 @@       optConvergencePhase = True,       optConvergenceMaxTime = 5 * 60,       optConcurrency = Nothing,+      optProfile = False,       optVerbose = 0,       optTestSpec = Nothing     }@@ -226,7 +228,11 @@ runBenchmark :: BenchOptions -> FutharkExe -> (FilePath, [InputOutputs]) -> IO (Maybe [BenchResult]) runBenchmark opts futhark (program, cases) = do   (tuning_opts, tuning_desc) <- determineTuning (optTuning opts) program-  let runopts = optExtraOptions opts ++ tuning_opts ++ determineCache (optCacheExt opts) program+  let runopts =+        optExtraOptions opts+          ++ tuning_opts+          ++ determineCache (optCacheExt opts) program+          ++ if optProfile opts then ["--profile", "--log"] else []   withProgramServer program (optRunner opts) runopts $ \server ->     mapM (forInputOutputs server tuning_desc) $ filter relevant cases   where@@ -254,7 +260,8 @@       runVerbose = optVerbose opts,       runConvergencePhase = optConvergencePhase opts,       runConvergenceMaxTime = optConvergenceMaxTime opts,-      runResultAction = f+      runResultAction = f,+      runProfile = optProfile opts     }  descText :: T.Text -> Int -> T.Text@@ -370,7 +377,8 @@ runBenchmarkCase _ opts _ _ _ _ (TestRun tags _ _ _ _)   | any (`elem` tags) $ optExcludeCase opts =       pure Nothing-runBenchmarkCase server opts futhark program entry pad_to tr@(TestRun _ input_spec (Succeeds expected_spec) _ dataset_desc) = do+runBenchmarkCase server opts futhark program entry pad_to tr = do+  let (TestRun _ input_spec (Succeeds expected_spec) _ dataset_desc) = tr   start_time <- liftIO getCurrentTime   prompt <- mkProgressPrompt opts pad_to dataset_desc start_time @@ -399,7 +407,7 @@       putStrLn ""       putRedLn err       pure $ Just $ DataResult dataset_desc $ Left err-    Right (runtimes, errout) -> do+    Right (runtimes, errout, report) -> do       let vec_runtimes = U.fromList $ map (fromIntegral . runMicroseconds) runtimes       g <- create       resampled <- liftIO $ resample g [Mean] 70000 vec_runtimes@@ -412,24 +420,20 @@               _ -> (0, 0)        reportResult runtimes bootstrapCI-      -- We throw away the 'errout' because it is almost always-      -- useless and adds too much to the .json file size.  This-      -- behaviour could be moved into a command line option if we-      -- wish.-      Result runtimes (getMemoryUsage errout) Nothing++      -- We throw away the 'errout' unless profiling is enabled,+      -- because it is otherwise useless and adds too much to the+      -- .json file size.+      let errout' = guard (optProfile opts) >> Just errout+          report' = guard (optProfile opts) >> Just report+      Result runtimes (getMemoryUsage report) errout' report'         & Right         & DataResult dataset_desc         & Just         & pure -getMemoryUsage :: T.Text -> M.Map T.Text Int-getMemoryUsage t =-  foldMap matchMap $ T.lines t-  where-    mem_regex = "Peak memory usage for space '([^']+)': ([0-9]+) bytes." :: T.Text-    matchMap l = case (l =~ mem_regex :: (T.Text, T.Text, T.Text, [T.Text])) of-      (_, _, _, [device, bytes]) -> M.singleton device (read $ T.unpack bytes)-      _ -> mempty+getMemoryUsage :: ProfilingReport -> M.Map T.Text Int+getMemoryUsage = fmap fromInteger . profilingMemory  commandLineOptions :: [FunOptDescr BenchOptions] commandLineOptions =@@ -619,7 +623,12 @@       "v"       ["verbose"]       (NoArg $ Right $ \config -> config {optVerbose = optVerbose config + 1})-      "Enable logging.  Pass multiple times for more."+      "Enable logging.  Pass multiple times for more.",+    Option+      "P"+      ["profile"]+      (NoArg $ Right $ \config -> config {optProfile = True})+      "Collect profiling information."   ]   where     max_timeout :: Int

@@ -26,6 +26,7 @@ import Futhark.CLI.MulticoreWASM qualified as MulticoreWASM import Futhark.CLI.OpenCL qualified as OpenCL import Futhark.CLI.Pkg qualified as Pkg+import Futhark.CLI.Profile qualified as Profile import Futhark.CLI.PyOpenCL qualified as PyOpenCL import Futhark.CLI.Python qualified as Python import Futhark.CLI.Query qualified as Query@@ -81,7 +82,8 @@       ("lsp", (LSP.main, "Run LSP server.")),       ("thanks", (Misc.mainThanks, "Express gratitude.")),       ("tokens", (Misc.mainTokens, "Print tokens from Futhark file.")),-      ("benchcmp", (Benchcmp.main, "Compare two benchmark results."))+      ("benchcmp", (Benchcmp.main, "Compare two benchmark results.")),+      ("profile", (Profile.main, "Analyse profiling data."))     ]  msg :: String

@@ -0,0 +1,136 @@+module Futhark.CLI.Profile (main) where++import Control.Exception (catch)+import Data.ByteString.Lazy.Char8 qualified as BS+import Data.Map qualified as M+import Data.Text qualified as T+import Data.Text.IO qualified as T+import Futhark.Bench+import Futhark.Util (showText)+import Futhark.Util.Options+import System.Directory (createDirectoryIfMissing, removePathForcibly)+import System.FilePath+import System.IO+import Text.Printf++commonPrefix :: (Eq e) => [e] -> [e] -> [e]+commonPrefix _ [] = []+commonPrefix [] _ = []+commonPrefix (x : xs) (y : ys)+  | x == y = x : commonPrefix xs ys+  | otherwise = []++longestCommonPrefix :: [FilePath] -> FilePath+longestCommonPrefix [] = ""+longestCommonPrefix (x : xs) = foldr commonPrefix x xs++memoryReport :: M.Map T.Text Integer -> T.Text+memoryReport = T.unlines . ("Peak memory usage in bytes" :) . map f . M.toList+  where+    f (space, bytes) = space <> ": " <> showText bytes++padRight :: Int -> T.Text -> T.Text+padRight k s = s <> T.replicate (k - T.length s) " "++padLeft :: Int -> T.Text -> T.Text+padLeft k s = T.replicate (k - T.length s) " " <> s++tabulateEvents :: [ProfilingEvent] -> T.Text+tabulateEvents = mkRows . M.toList . M.fromListWith comb . map pair+  where+    pair (ProfilingEvent name dur _) = (name, (1, dur))+    comb (xn, xdur) (yn, ydur) = (xn + yn, xdur + ydur)+    numpad = 15+    mkRows rows =+      let longest = foldl max numpad $ map (T.length . fst) rows+          header = headerRow longest+          splitter = T.map (const '-') header+          bottom =+            T.unwords+              [ showText (sum (map (fst . snd) rows)),+                "events with a total runtime of ",+                T.pack $ printf "%.2fμs" $ sum $ map (snd . snd) rows+              ]+       in T.unlines $+            header+              : splitter+              : map (mkRow longest) rows+                <> [splitter, bottom]+    headerRow longest =+      T.unwords+        [ padLeft longest "Cost centre",+          padLeft numpad "count",+          padLeft numpad "total",+          padLeft numpad "mean"+        ]+    mkRow longest (name, (n, dur)) =+      T.unwords+        [ padRight longest name,+          padLeft numpad (showText n),+          padLeft numpad $ T.pack $ printf "%.2fμs" dur,+          padLeft numpad $ T.pack $ printf "%.2fμs" $ dur / fromInteger n+        ]++analyseProfileReport :: FilePath -> [BenchResult] -> IO ()+analyseProfileReport json_path bench_results = do+  let top_dir = takeFileName json_path -<.> "prof"+  T.hPutStrLn stderr $ "Writing results to " <> T.pack top_dir <> "/"+  T.hPutStrLn stderr $ "Stripping '" <> T.pack prefix <> "' from program paths."+  removePathForcibly top_dir+  mapM_ (onBenchResult top_dir) bench_results+  where+    prefix = longestCommonPrefix $ map benchResultProg bench_results++    -- Eliminate characters that are filesystem-meaningful.+    escape '/' = '_'+    escape c = c++    problem prog_name name what =+      T.hPutStrLn stderr $ prog_name <> " dataset " <> name <> ": " <> what++    onBenchResult top_dir (BenchResult prog_path data_results) = do+      let (prog_path', entry) = span (/= ':') prog_path+          prog_name = drop (length prefix) prog_path'+          prog_dir = top_dir </> dropExtension prog_name </> drop 1 entry+      createDirectoryIfMissing True prog_dir+      mapM_ (onDataResult prog_dir (T.pack prog_name)) data_results++    onDataResult _ prog_name (DataResult name (Left _)) =+      problem prog_name name "execution failed"+    onDataResult prog_dir prog_name (DataResult name (Right res)) = do+      let name' = prog_dir </> T.unpack (T.map escape name)+      case stdErr res of+        Nothing -> problem prog_name name "no log recorded"+        Just text -> T.writeFile (name' <.> ".log") text+      case report res of+        Nothing -> problem prog_name name "no profiling information"+        Just r ->+          T.writeFile (name' <> ".summary") $+            memoryReport (profilingMemory r)+              <> "\n\n"+              <> tabulateEvents (profilingEvents r)++readFileSafely :: FilePath -> IO (Either String BS.ByteString)+readFileSafely filepath =+  (Right <$> BS.readFile filepath) `catch` couldNotRead+  where+    couldNotRead e = pure $ Left $ show (e :: IOError)++decodeFileBenchResults ::+  FilePath ->+  IO (Either String [BenchResult])+decodeFileBenchResults path = do+  file <- readFileSafely path+  pure $ file >>= decodeBenchResults++-- | Run @futhark profile@.+main :: String -> [String] -> IO ()+main = mainWithOptions () [] "<file>" f+  where+    f [json_path] () = Just $ do+      res_either <- decodeFileBenchResults json_path++      case res_either of+        Left a -> hPutStrLn stderr a+        Right a -> analyseProfileReport json_path a+    f _ _ = Nothing

@@ -102,13 +102,6 @@              optionArgument = RequiredArgument "OPT",              optionDescription = "Add an additional build option to the string passed to NVRTC.",              optionAction = [C.cstm|futhark_context_config_add_nvrtc_option(cfg, optarg);|]-           },-         Option-           { optionLongName = "profile",-             optionShortName = Just 'P',-             optionArgument = NoArgument,-             optionDescription = "Gather profiling data while executing and print out a summary at the end.",-             optionAction = [C.cstm|futhark_context_config_set_profiling(cfg, 1);|]            }        ] 

@@ -169,13 +169,6 @@              optionAction = [C.cstm|futhark_context_config_add_build_option(cfg, optarg);|]            },          Option-           { optionLongName = "profile",-             optionShortName = Just 'P',-             optionArgument = NoArgument,-             optionDescription = "Gather profiling data while executing and print out a summary at the end.",-             optionAction = [C.cstm|futhark_context_config_set_profiling(cfg, 1);|]-           },-         Option            { optionLongName = "list-devices",              optionShortName = Nothing,              optionArgument = NoArgument,

@@ -381,26 +381,6 @@                   return strdup("Unknown error.  This is a compiler bug.");                 }|] -genProfileReport :: [Name] -> GC.CompilerM op s ()-genProfileReport cost_centres =-  GC.profileReport-    [C.citem|{struct cost_centres* ccs = cost_centres_new();-              $stms:(map initCostCentre (def_cost_centres<>cost_centres))-              tally_profiling_records(ctx, ccs);-              cost_centre_report(ccs, &builder);-              cost_centres_free(ccs);-              }|]-  where-    def_cost_centres =-      [ "copy_dev_to_dev",-        "copy_dev_to_host",-        "copy_host_to_dev",-        "copy_scalar_to_dev",-        "copy_scalar_from_dev"-      ]-    initCostCentre v =-      [C.cstm|cost_centres_init(ccs, $string:(nameToString v));|]- -- | Called after most code has been generated to generate the bulk of -- the boilerplate. generateGPUBoilerplate ::@@ -436,5 +416,3 @@   GC.generateProgramStruct    GC.onClear [C.citem|if (ctx->error == NULL) { gpu_free_all(ctx); }|]--  genProfileReport kernels

@@ -22,6 +22,7 @@ import Control.Monad.State import Data.Bifunctor (second) import Data.DList qualified as DL+import Data.List qualified as L import Data.Loc import Data.Map.Strict qualified as M import Data.Maybe@@ -37,7 +38,7 @@ import Futhark.CodeGen.Backends.GenericC.Server (serverDefs) import Futhark.CodeGen.Backends.GenericC.Types import Futhark.CodeGen.ImpCode-import Futhark.CodeGen.RTS.C (cacheH, contextH, contextPrototypesH, copyH, errorsH, freeListH, halfH, lockH, timingH, utilH)+import Futhark.CodeGen.RTS.C (cacheH, contextH, contextPrototypesH, copyH, errorsH, eventListH, freeListH, halfH, lockH, timingH, utilH) import Futhark.IR.GPU.Sizes import Futhark.Manifest qualified as Manifest import Futhark.MonadFreshNames@@ -70,7 +71,7 @@           2 -> ("lmad_copy_2b", [C.cty|typename uint16_t|])           4 -> ("lmad_copy_4b", [C.cty|typename uint32_t|])           8 -> ("lmad_copy_8b", [C.cty|typename uint64_t|])-          k -> error $ "lmadcopyCPU: " <> error (show k)+          k -> error $ "lmadcopyCPU: " <> show k       r = length shape       dststride_inits = [[C.cinit|$exp:e|] | Count e <- dststride]       srcstride_inits = [[C.cinit|$exp:e|] | Count e <- srcstride]@@ -261,15 +262,10 @@    onClear [C.citem|ctx->$id:peakname = 0;|] -  let peakmsg = "Peak memory usage for " ++ spacedesc ++ ": %lld bytes.\n"+  let peakmsg = "\"" <> spacedesc <> "\": %lld"   pure     ( [unrefdef, allocdef, setdef],-      -- Do not report memory usage for DefaultSpace (CPU memory),-      -- because it would not be accurate anyway.  This whole-      -- tracking probably needs to be rethought.-      if space == DefaultSpace-        then [C.citem|{}|]-        else [C.citem|str_builder(&builder, $string:peakmsg, (long long) ctx->$id:peakname);|]+      [C.citem|str_builder(&builder, $string:peakmsg, (long long) ctx->$id:peakname);|]     )   where     mty = fatMemType space@@ -442,6 +438,7 @@ $timingH $lockH $freeListH+$eventListH |]    let early_decls = definitionsText $ DL.toList $ compEarlyDecls endstate@@ -588,7 +585,6 @@   ctx <- contextType   cfg <- configType   ops <- asks envOperations-  profilereport <- gets $ DL.toList . compProfileItems    publicDef_ "context_config_set_debugging" InitDecl $ \s ->     ( [C.cedecl|void $id:s($ty:cfg* cfg, int flag);|],@@ -614,7 +610,7 @@   publicDef_ "context_config_set_cache_file" MiscDecl $ \s ->     ( [C.cedecl|void $id:s($ty:cfg* cfg, const char *f);|],       [C.cedecl|void $id:s($ty:cfg* cfg, const char *f) {-                 cfg->cache_fname = f;+                 cfg->cache_fname = strdup(f);                }|]     ) @@ -640,6 +636,7 @@     )    sync <- publicName "context_sync"+  let comma = [C.citem|str_builder_char(&builder, ',');|]   publicDef_ "context_report" MiscDecl $ \s ->     ( [C.cedecl|char* $id:s($ty:ctx *ctx);|],       [C.cedecl|char* $id:s($ty:ctx *ctx) {@@ -649,10 +646,12 @@                   struct str_builder builder;                  str_builder_init(&builder);-                 $items:memreport-                 if (ctx->profiling) {-                   $items:profilereport-                 }+                 str_builder_char(&builder, '{');+                 str_builder_str(&builder, "\"memory\":{");+                 $items:(L.intersperse comma memreport)+                 str_builder_str(&builder, "},\"events\":[");+                 report_events_in_list(&ctx->event_list, &builder);+                 str_builder_str(&builder, "]}");                  return builder.str;                }|]     )

@@ -46,9 +46,7 @@         optionArgument = NoArgument,         optionDescription = "Perform possibly expensive internal correctness checks and verbose logging.",         optionAction =-          [C.cstm|{futhark_context_config_set_debugging(cfg, 1);-                                print_report = 1;-                               }|]+          [C.cstm|{futhark_context_config_set_debugging(cfg, 1);}|]       },     Option       { optionLongName = "log",@@ -56,11 +54,16 @@         optionArgument = NoArgument,         optionDescription = "Print various low-overhead logging information to stderr while running.",         optionAction =-          [C.cstm|{futhark_context_config_set_logging(cfg, 1);-                   print_report = 1;-                  }|]+          [C.cstm|{futhark_context_config_set_logging(cfg, 1);}|]       },     Option+      { optionLongName = "profile",+        optionShortName = Just 'P',+        optionArgument = NoArgument,+        optionDescription = "Enable the collection of profiling information.",+        optionAction = [C.cstm|futhark_context_config_set_profiling(cfg, 1);|]+      },+    Option       { optionLongName = "entry-point",         optionShortName = Just 'e',         optionArgument = RequiredArgument "NAME",@@ -423,7 +426,6 @@  static int binary_output = 0; static int print_result = 1;-static int print_report = 0; static typename FILE *runtime_file; static int perform_warmup = 0; static int num_runs = 1;@@ -498,12 +500,6 @@      if (runtime_file != NULL) {       fclose(runtime_file);-    }--    if (print_report) {-      char *report = futhark_context_report(ctx);-      fputs(report, stderr);-      free(report);     }   } 

@@ -40,7 +40,6 @@     headerDecl,     publicDef,     publicDef_,-    profileReport,     onClear,     HeaderSection (..),     libDecl,@@ -111,7 +110,6 @@     compHeaderDecls :: M.Map HeaderSection (DL.DList C.Definition),     compLibDecls :: DL.DList C.Definition,     compCtxFields :: DL.DList (C.Id, C.Type, Maybe C.Exp, Maybe (C.Stm, C.Stm)),-    compProfileItems :: DL.DList C.BlockItem,     compClearItems :: DL.DList C.BlockItem,     compDeclaredMem :: [(VName, Space)],     compItems :: DL.DList C.BlockItem@@ -127,7 +125,6 @@       compHeaderDecls = mempty,       compLibDecls = mempty,       compCtxFields = mempty,-      compProfileItems = mempty,       compClearItems = mempty,       compDeclaredMem = mempty,       compItems = mempty@@ -422,10 +419,6 @@ contextFieldDyn :: C.Id -> C.Type -> C.Stm -> C.Stm -> CompilerM op s () contextFieldDyn name ty create free = modify $ \s ->   s {compCtxFields = compCtxFields s <> DL.singleton (name, ty, Nothing, Just (create, free))}--profileReport :: C.BlockItem -> CompilerM op s ()-profileReport x = modify $ \s ->-  s {compProfileItems = compProfileItems s <> DL.singleton x}  onClear :: C.BlockItem -> CompilerM op s () onClear x = modify $ \s ->

@@ -36,6 +36,13 @@         optionAction = [C.cstm|futhark_context_config_set_logging(cfg, 1);|]       },     Option+      { optionLongName = "profile",+        optionShortName = Just 'P',+        optionArgument = NoArgument,+        optionDescription = "Enable the collection of profiling information.",+        optionAction = [C.cstm|futhark_context_config_set_profiling(cfg, 1);|]+      },+    Option       { optionLongName = "help",         optionShortName = Just 'h',         optionArgument = NoArgument,

@@ -38,7 +38,6 @@     collect',     collect,     simpleCall,-    copyMemoryDefaultSpace,   ) where @@ -143,7 +142,6 @@   { opsWriteScalar :: WriteScalar op s,     opsReadScalar :: ReadScalar op s,     opsAllocate :: Allocate op s,-    opsCopy :: Copy op s,     -- | @(dst,src)@-space mapping to copy functions.     opsCopies :: M.Map (Space, Space) (DoLMADCopy op s),     opsCompiler :: OpCompiler op s,@@ -160,7 +158,6 @@     { opsWriteScalar = defWriteScalar,       opsReadScalar = defReadScalar,       opsAllocate = defAllocate,-      opsCopy = defCopy,       opsCopies = M.singleton (DefaultSpace, DefaultSpace) lmadcopyCPU,       opsCompiler = defCompiler,       opsEntryOutput = defEntryOutput,@@ -173,8 +170,6 @@       error "Cannot read from non-default memory space"     defAllocate _ _ _ =       error "Cannot allocate in non-default memory space"-    defCopy _ _ _ _ _ _ _ _ =-      error "Cannot copy to or from non-default memory space"     defCompiler _ =       error "The default compiler cannot compile extended operations"     defEntryOutput _ _ _ _ =@@ -582,9 +577,16 @@   pack_output <- asks envEntryOutput   pack_output mem sid bt ept dims entryPointOutput (Imp.TransparentValue (Imp.ArrayValue mem _ bt ept dims)) = do-  mem' <- Cast <$> compileVar mem <*> pure (compilePrimTypeExt bt ept)+  mem' <- compileVar mem   dims' <- mapM compileDim dims-  pure $ simpleCall "createArray" [mem', Tuple dims', Var $ compilePrimToExtNp bt ept]+  pure $+    simpleCall+      "np.reshape"+      [ Index+          (Call (Field mem' "view") [Arg $ Var $ compilePrimToExtNp bt ept])+          (IdxRange (Integer 0) (foldl1 (BinOp "*") dims')),+        Tuple dims'+      ]  badInput :: Int -> PyExp -> T.Text -> PyStmt badInput i e t =@@ -853,24 +855,6 @@       zip (map snd results) res     ) -copyMemoryDefaultSpace ::-  PyExp ->-  PyExp ->-  PyExp ->-  PyExp ->-  PyExp ->-  CompilerM op s ()-copyMemoryDefaultSpace destmem destidx srcmem srcidx nbytes = do-  let offset_call1 =-        simpleCall-          "addressOffset"-          [destmem, destidx, Var "ct.c_byte"]-  let offset_call2 =-        simpleCall-          "addressOffset"-          [srcmem, srcidx, Var "ct.c_byte"]-  stm $ Exp $ simpleCall "ct.memmove" [offset_call1, offset_call2, nbytes]- data ReturnTiming = ReturnTiming | DoNotReturnTiming  compileEntryFun ::@@ -1116,21 +1100,21 @@   simpleCall "np.int64" [Integer $ toInteger v] compilePrimValue (FloatValue (Float16Value v))   | isInfinite v =-      if v > 0 then Var "np.inf" else Var "-np.inf"+      if v > 0 then Var "np.float16(np.inf)" else Var "np.float16(-np.inf)"   | isNaN v =-      Var "np.nan"+      Var "np.float16(np.nan)"   | otherwise = simpleCall "np.float16" [Float $ fromRational $ toRational v] compilePrimValue (FloatValue (Float32Value v))   | isInfinite v =-      if v > 0 then Var "np.inf" else Var "-np.inf"+      if v > 0 then Var "np.float32(np.inf)" else Var "np.float32(-np.inf)"   | isNaN v =-      Var "np.nan"+      Var "np.float32(np.nan)"   | otherwise = simpleCall "np.float32" [Float $ fromRational $ toRational v] compilePrimValue (FloatValue (Float64Value v))   | isInfinite v =       if v > 0 then Var "np.inf" else Var "-np.inf"   | isNaN v =-      Var "np.nan"+      Var "np.float64(np.nan)"   | otherwise = simpleCall "np.float64" [Float $ fromRational $ toRational v] compilePrimValue (BoolValue v) = Bool v compilePrimValue UnitValue = Var "np.byte(0)"@@ -1222,8 +1206,8 @@ generateWrite dst iexp pt (Imp.Space space) elemexp = do   writer <- asks envWriteScalar   writer dst iexp pt space elemexp-generateWrite dst iexp pt DefaultSpace elemexp =-  stm $ Exp $ simpleCall "writeScalarArray" [dst, iexp, toStorage pt elemexp]+generateWrite dst iexp _ DefaultSpace elemexp =+  stm $ Exp $ simpleCall "writeScalarArray" [dst, iexp, elemexp]  -- | Compile an 'LMADCopy' using sequential nested loops, but -- parameterised over how to do the reads and writes.

@@ -38,7 +38,6 @@   | Cond PyExp PyExp PyExp   | Index PyExp PyIdx   | Call PyExp [PyArg]-  | Cast PyExp String   | Tuple [PyExp]   | List [PyExp]   | Field PyExp String@@ -112,12 +111,6 @@   pretty (BinOp s e1 e2) = parens (pretty e1 <+> pretty s <+> pretty e2)   pretty (UnOp s e) = pretty s <> parens (pretty e)   pretty (Cond e1 e2 e3) = pretty e2 <+> "if" <+> pretty e1 <+> "else" <+> pretty e3-  pretty (Cast src bt) =-    "ct.cast"-      <> parens-        ( pretty src <> ","-            <+> "ct.POINTER" <> parens (pretty bt)-        )   pretty (Index src idx) = pretty src <> brackets (pretty idx)   pretty (Call fun exps) = pretty fun <> parens (commasep $ map pretty exps)   pretty (Tuple [dim]) = parens (pretty dim <> ",")

@@ -84,13 +84,6 @@              optionArgument = RequiredArgument "OPT",              optionDescription = "Add an additional build option to the string passed to NVRTC.",              optionAction = [C.cstm|futhark_context_config_add_build_option(cfg, optarg);|]-           },-         Option-           { optionLongName = "profile",-             optionShortName = Just 'P',-             optionArgument = NoArgument,-             optionDescription = "Gather profiling data while executing and print out a summary at the end.",-             optionAction = [C.cstm|futhark_context_config_set_profiling(cfg, 1);|]            }        ] 

@@ -19,7 +19,6 @@     addTimingFields,     functionTiming,     functionIterations,-    multiCoreReport,     multicoreDef,     multicoreName,     DefSpecifier,@@ -64,13 +63,6 @@ cliOptions :: [Option] cliOptions =   [ Option-      { optionLongName = "profile",-        optionShortName = Just 'P',-        optionArgument = NoArgument,-        optionAction = [C.cstm|futhark_context_config_set_profiling(cfg, 1);|],-        optionDescription = "Gather profiling information."-      },-    Option       { optionLongName = "num-threads",         optionShortName = Nothing,         optionArgument = RequiredArgument "INT",@@ -208,113 +200,26 @@         else RawMem     ) -functionRuntime :: Name -> C.Id-functionRuntime = (`C.toIdent` mempty) . (<> "_total_runtime")--functionRuns :: Name -> C.Id-functionRuns = (`C.toIdent` mempty) . (<> "_runs")--functionIter :: Name -> C.Id-functionIter = (`C.toIdent` mempty) . (<> "_iter")--multiCoreReport :: [(Name, Bool)] -> [C.BlockItem]-multiCoreReport names = report_kernels-  where-    report_kernels = concatMap reportKernel names-    max_name_len_pad = 40-    format_string name True =-      let name_s = nameToString name-          padding = replicate (max_name_len_pad - length name_s) ' '-       in unwords ["tid %2d -", name_s ++ padding, "ran %10d times; avg: %10ldus; total: %10ldus; time pr. iter %9.6f; iters %9ld; avg %ld\n"]-    format_string name False =-      let name_s = nameToString name-          padding = replicate (max_name_len_pad - length name_s) ' '-       in unwords ["        ", name_s ++ padding, "ran %10d times; avg: %10ldus; total: %10ldus; time pr. iter %9.6f; iters %9ld; avg %ld\n"]-    reportKernel (name, is_array) =-      let runs = functionRuns name-          total_runtime = functionRuntime name-          iters = functionIter name-       in if is_array-            then-              [ [C.citem|-                     for (int i = 0; i < ctx->scheduler.num_threads; i++) {-                       fprintf(ctx->log,-                         $string:(format_string name is_array),-                         i,-                         ctx->program->$id:runs[i],-                         (long int) ctx->program->$id:total_runtime[i] / (ctx->program->$id:runs[i] != 0 ? ctx->program->$id:runs[i] : 1),-                         (long int) ctx->program->$id:total_runtime[i],-                         (double) ctx->program->$id:total_runtime[i] /  (ctx->program->$id:iters[i] == 0 ? 1 : (double)ctx->program->$id:iters[i]),-                         (long int) (ctx->program->$id:iters[i]),-                         (long int) (ctx->program->$id:iters[i]) / (ctx->program->$id:runs[i] != 0 ? ctx->program->$id:runs[i] : 1)-                         );-                     }-                   |]-              ]-            else-              [ [C.citem|-                    fprintf(ctx->log,-                       $string:(format_string name is_array),-                       ctx->program->$id:runs,-                       (long int) ctx->program->$id:total_runtime / (ctx->program->$id:runs != 0 ? ctx->program->$id:runs : 1),-                       (long int) ctx->program->$id:total_runtime,-                       (double) ctx->program->$id:total_runtime /  (ctx->program->$id:iters == 0 ? 1 : (double)ctx->program->$id:iters),-                       (long int) (ctx->program->$id:iters),-                       (long int) (ctx->program->$id:iters) / (ctx->program->$id:runs != 0 ? ctx->program->$id:runs : 1));-                   |],-                [C.citem|ctx->total_runtime += ctx->program->$id:total_runtime;|],-                [C.citem|ctx->total_runs += ctx->program->$id:runs;|]-              ]--addBenchmarkFields :: Name -> Maybe C.Id -> GC.CompilerM op s ()-addBenchmarkFields name (Just _) = do-  GC.contextFieldDyn-    (functionRuntime name)-    [C.cty|typename int64_t*|]-    [C.cstm|ctx->program->$id:(functionRuntime name) = calloc(sizeof(typename int64_t), ctx->scheduler.num_threads);|]-    [C.cstm|free(ctx->program->$id:(functionRuntime name));|]-  GC.contextFieldDyn-    (functionRuns name)-    [C.cty|int*|]-    [C.cstm|ctx->program->$id:(functionRuns name) = calloc(sizeof(int), ctx->scheduler.num_threads);|]-    [C.cstm|free(ctx->program->$id:(functionRuns name));|]-  GC.contextFieldDyn-    (functionIter name)-    [C.cty|typename int64_t*|]-    [C.cstm|ctx->program->$id:(functionIter name) = calloc(sizeof(sizeof(typename int64_t)), ctx->scheduler.num_threads);|]-    [C.cstm|free(ctx->program->$id:(functionIter name));|]-addBenchmarkFields name Nothing = do-  GC.contextField (functionRuntime name) [C.cty|typename int64_t|] $ Just [C.cexp|0|]-  GC.contextField (functionRuns name) [C.cty|int|] $ Just [C.cexp|0|]-  GC.contextField (functionIter name) [C.cty|typename int64_t|] $ Just [C.cexp|0|]--benchmarkCode :: Name -> Maybe C.Id -> [C.BlockItem] -> GC.CompilerM op s [C.BlockItem]-benchmarkCode name tid code = do-  addBenchmarkFields name tid+benchmarkCode :: Name -> [C.BlockItem] -> GC.CompilerM op s [C.BlockItem]+benchmarkCode name code = do+  event <- newVName "event"   pure     [C.citems|-     typename uint64_t $id:start = 0;-     if (ctx->profiling && !ctx->profiling_paused) {-       $id:start = get_wall_time();+     struct mc_event* $id:event = mc_event_new(ctx);+     if ($id:event != NULL) {+       $id:event->bef = get_wall_time();      }      $items:code-     if (ctx->profiling && !ctx->profiling_paused) {-       typename uint64_t $id:end = get_wall_time();-       typename uint64_t elapsed = $id:end - $id:start;-       $items:(updateFields tid)-     }-     |]-  where-    start = name <> "_start"-    end = name <> "_end"-    updateFields Nothing =-      [C.citems|__atomic_fetch_add(&ctx->program->$id:(functionRuns name), 1, __ATOMIC_RELAXED);-                                            __atomic_fetch_add(&ctx->program->$id:(functionRuntime name), elapsed, __ATOMIC_RELAXED);-                                            __atomic_fetch_add(&ctx->program->$id:(functionIter name), iterations, __ATOMIC_RELAXED);|]-    updateFields (Just _tid') =-      [C.citems|ctx->program->$id:(functionRuns name)[tid]++;-                                            ctx->program->$id:(functionRuntime name)[tid] += elapsed;-                                            ctx->program->$id:(functionIter name)[tid] += iterations;|]+     if ($id:event != NULL) {+       $id:event->aft = get_wall_time();+       lock_lock(&ctx->event_list_lock);+       add_event(ctx,+                 $string:(nameToString name),+                 strdup("nothing further"),+                 $id:event,+                 (typename event_report_fn)mc_event_report);+       lock_unlock(&ctx->event_list_lock);+     }|]  functionTiming :: Name -> C.Id functionTiming = (`C.toIdent` mempty) . (<> "_total_time")@@ -353,7 +258,7 @@   let (fargs, fctypes) = unzip free   let (retval_args, retval_ctypes) = unzip retval   multicoreDef basename $ \s -> do-    fbody <- benchmarkCode s (Just "tid") <=< GC.inNewFunction $+    fbody <- benchmarkCode s <=< GC.inNewFunction $       GC.cachingMemory lexical $ \decl_cached free_cached -> GC.collect $ do         mapM_ GC.item [C.citems|$decls:(compileGetStructVals fstruct fargs fctypes)|]         mapM_ GC.item [C.citems|$decls:(compileGetRetvalStructVals fstruct retval_args retval_ctypes)|]@@ -443,15 +348,13 @@     Static -> GC.stm [C.cstm|$id:ftask_name.sched = STATIC;|]    -- Generate the nested segop function if available-  fnpar_task <- case par_task of+  case par_task of     Just (ParallelTask nested_code) -> do       let lexical_nested = lexicalMemoryUsageMC TraverseKernels $ Function Nothing [] params nested_code       fnpar_task <- generateParLoopFn lexical_nested (name ++ "_nested_task") nested_code fstruct free retval       GC.stm [C.cstm|$id:ftask_name.nested_fn = $id:fnpar_task;|]-      pure $ zip [fnpar_task] [True]-    Nothing -> do+    Nothing ->       GC.stm [C.cstm|$id:ftask_name.nested_fn=NULL;|]-      pure mempty    let ftask_err = fpar_task <> "_err"       code =@@ -462,9 +365,6 @@                   }|]    mapM_ GC.item code--  -- Add profile fields for -P option-  mapM_ GC.profileReport $ multiCoreReport $ (fpar_task, True) : fnpar_task compileOp (ParLoop s' body free) = do   free_ctypes <- mapM paramToCType free   let free_args = map paramName free@@ -475,7 +375,7 @@     prepareTaskStruct multicoreDef (s' ++ "_parloop_struct") free_args free_ctypes mempty mempty    ftask <- multicoreDef (s' ++ "_parloop") $ \s -> do-    fbody <- benchmarkCode s (Just "tid") <=< GC.inNewFunction $+    fbody <- benchmarkCode s <=< GC.inNewFunction $       GC.cachingMemory lexical $ \decl_cached free_cached -> GC.collect $ do         GC.items [C.citems|$decls:(compileGetStructVals fstruct free_args free_ctypes)|] @@ -515,7 +415,6 @@   code' <-     benchmarkCode       ftask_total-      Nothing       [C.citems|int $id:ftask_err = scheduler_execute_task(&ctx->scheduler,                                                            &$id:ftask_name);                if ($id:ftask_err != 0) {@@ -524,7 +423,6 @@                }|]    mapM_ GC.item code'-  mapM_ GC.profileReport $ multiCoreReport $ zip [ftask, ftask_total] [True, False] compileOp (Atomic aop) =   atomicOps aop (\ty _ -> pure [C.cty|$ty:ty*|]) compileOp (ISPCKernel body _) =

@@ -776,20 +776,15 @@       Static -> GC.stm [C.cstm|$id:ftask_name.sched = STATIC;|]      -- Generate the nested segop function if available-    fnpar_task <- case par_task of+    case par_task of       Just (ParallelTask nested_code) -> do         let lexical_nested = lexicalMemoryUsageMC OpaqueKernels $ Function Nothing [] params nested_code         fnpar_task <- MC.generateParLoopFn lexical_nested (name ++ "_nested_task") nested_code fstruct free retval         GC.stm [C.cstm|$id:ftask_name.nested_fn = $id:fnpar_task;|]-        pure $ zip [fnpar_task] [True]-      Nothing -> do+      Nothing ->         GC.stm [C.cstm|$id:ftask_name.nested_fn=NULL;|]-        pure mempty      GC.stm [C.cstm|return scheduler_prepare_task(&ctx->scheduler, &$id:ftask_name);|]--    -- Add profile fields for -P option-    mapM_ GC.profileReport $ MC.multiCoreReport $ (fpar_task, True) : fnpar_task    schedn <- MC.multicoreDef "schedule_shim" $ \s ->     pure

@@ -190,7 +190,6 @@           opsWriteScalar = writeOpenCLScalar,           opsReadScalar = readOpenCLScalar,           opsAllocate = allocateOpenCLBuffer,-          opsCopy = copyOpenCLMemory,           opsCopies =             M.insert (Imp.Space "device", Imp.Space "device") copygpu2gpu $               opsCopies defaultOperations,@@ -330,55 +329,6 @@       simpleCall "opencl_alloc" [Var "self", size, String $ prettyText mem] allocateOpenCLBuffer _ _ space =   error $ "Cannot allocate in '" ++ space ++ "' space"--copyOpenCLMemory :: Copy Imp.OpenCL ()-copyOpenCLMemory destmem destidx Imp.DefaultSpace srcmem srcidx (Imp.Space "device") nbytes bt = do-  let divide = BinOp "//" nbytes (Integer $ Imp.primByteSize bt)-      end = BinOp "+" destidx divide-      dest = Index destmem (IdxRange destidx end)-  stm $-    ifNotZeroSize nbytes $-      Exp $-        Call-          (Var "cl.enqueue_copy")-          [ Arg $ Var "self.queue",-            Arg dest,-            Arg srcmem,-            ArgKeyword "device_offset" $ asLong srcidx,-            ArgKeyword "is_blocking" $ Var "synchronous"-          ]-copyOpenCLMemory destmem destidx (Imp.Space "device") srcmem srcidx Imp.DefaultSpace nbytes _ = do-  let end = BinOp "+" srcidx nbytes-      src = Index (simpleCall "createArray" [srcmem, List [nbytes], Var "np.byte"]) (IdxRange srcidx end)-  stm $-    ifNotZeroSize nbytes $-      Exp $-        Call-          (Var "cl.enqueue_copy")-          [ Arg $ Var "self.queue",-            Arg destmem,-            Arg src,-            ArgKeyword "device_offset" $ asLong destidx,-            ArgKeyword "is_blocking" $ Var "synchronous"-          ]-copyOpenCLMemory destmem destidx (Imp.Space "device") srcmem srcidx (Imp.Space "device") nbytes _ = do-  stm $-    ifNotZeroSize nbytes $-      Exp $-        Call-          (Var "cl.enqueue_copy")-          [ Arg $ Var "self.queue",-            Arg destmem,-            Arg srcmem,-            ArgKeyword "dst_offset" $ asLong destidx,-            ArgKeyword "src_offset" $ asLong srcidx,-            ArgKeyword "byte_count" $ asLong nbytes-          ]-  finishIfSynchronous-copyOpenCLMemory destmem destidx Imp.DefaultSpace srcmem srcidx Imp.DefaultSpace nbytes _ =-  copyMemoryDefaultSpace destmem destidx srcmem srcidx nbytes-copyOpenCLMemory _ _ destspace _ _ srcspace _ _ =-  error $ "Cannot copy to " ++ show destspace ++ " from " ++ show srcspace  packArrayOutput :: EntryOutput Imp.OpenCL () packArrayOutput mem "device" bt ept dims = do

@@ -45,12 +45,5 @@     operations :: GenericPython.Operations Imp.Sequential ()     operations =       GenericPython.defaultOperations-        { GenericPython.opsCompiler = const $ pure (),-          GenericPython.opsCopy = copySequentialMemory+        { GenericPython.opsCompiler = const $ pure ()         }--copySequentialMemory :: GenericPython.Copy Imp.Sequential ()-copySequentialMemory destmem destidx DefaultSpace srcmem srcidx DefaultSpace nbytes _bt =-  GenericPython.copyMemoryDefaultSpace destmem destidx srcmem srcidx nbytes-copySequentialMemory _ _ destspace _ _ srcspace _ _ =-  error $ "Cannot copy to " ++ show destspace ++ " from " ++ show srcspace

@@ -518,9 +518,9 @@         Unsigned -> " (unsigned)"         Signed -> mempty   pretty (ArrayValue mem space et ept shape) =-    foldr f (pretty et) shape <+> "at" <+> pretty mem <> pretty space <+> ept'+    foldMap (brackets . pretty) shape+      <> (pretty et <+> "at" <+> pretty mem <> pretty space <+> ept')     where-      f e s = brackets $ s <> comma <> pretty e       ept' = case ept of         Unsigned -> " (unsigned)"         Signed -> mempty

@@ -818,18 +818,9 @@   let doBody = compileLoopBody params body    case form of-    ForLoop i _ bound loopvars -> do-      let setLoopParam (p, a)-            | Prim _ <- paramType p =-                copyDWIM (paramName p) [] (Var a) [DimFix $ Imp.le64 i]-            | otherwise =-                pure ()-+    ForLoop i _ bound -> do       bound' <- toExp bound--      dLParams $ map fst loopvars-      sFor' i bound' $-        mapM_ setLoopParam loopvars >> doBody+      sFor' i bound' doBody     WhileLoop cond ->       sWhile (TPrimExp $ Imp.var cond Bool) doBody 

@@ -157,13 +157,14 @@  prepareAtomicUpdateGlobal ::   Maybe Locking ->+  Shape ->   [VName] ->   SegHistSlug ->   CallKernelGen     ( Maybe Locking,       [Imp.TExp Int64] -> InKernelGen ()     )-prepareAtomicUpdateGlobal l dests slug =+prepareAtomicUpdateGlobal l segments dests slug =   -- We need a separate lock array if the operators are not all of a   -- particularly simple form that permits pure atomic operations.   case (l, slugAtomicUpdate slug) of@@ -181,7 +182,8 @@       let num_locks = 100151           dims =             map pe64 $-              shapeDims (histOpShape (slugOp slug))+              shapeDims segments+                ++ shapeDims (histOpShape (slugOp slug))                 ++ [tvSize (slugNumSubhistos slug)]                 ++ shapeDims (histShape (slugOp slug)) @@ -202,6 +204,7 @@  prepareIntermediateArraysGlobal ::   Passage ->+  Shape ->   Imp.TExp Int32 ->   Imp.TExp Int64 ->   [SegHistSlug] ->@@ -209,7 +212,7 @@     ( Imp.TExp Int32,       [[Imp.TExp Int64] -> InKernelGen ()]     )-prepareIntermediateArraysGlobal passage hist_T hist_N slugs = do+prepareIntermediateArraysGlobal passage segments hist_T hist_N slugs = do   -- The paper formulae assume there is only one histogram, but in our   -- implementation there can be multiple that have been horisontally   -- fused.  We do a bit of trickery with summings and averages to@@ -379,7 +382,7 @@          pure $ subhistosArray info -      (l', do_op') <- prepareAtomicUpdateGlobal l dests slug+      (l', do_op') <- prepareAtomicUpdateGlobal l segments dests slug        pure (l', do_op') @@ -492,6 +495,7 @@   (hist_S, histograms) <-     prepareIntermediateArraysGlobal       (bodyPassage kbody)+      (Shape (init space_sizes))       num_threads       (pe64 $ last space_sizes)       slugs

@@ -7,6 +7,7 @@     contextPrototypesH,     copyH,     freeListH,+    eventListH,     gpuH,     gpuPrototypesH,     halfH,@@ -51,6 +52,11 @@ freeListH :: T.Text freeListH = $(embedStringFile "rts/c/free_list.h") {-# NOINLINE freeListH #-}++-- | @rts/c/event_list.h@+eventListH :: T.Text+eventListH = $(embedStringFile "rts/c/event_list.h")+{-# NOINLINE eventListH #-}  -- | @rts/c/gpu.h@ gpuH :: T.Text

@@ -546,7 +546,7 @@     res <- m     ret <- mapM subExpResType res     pure (res, ret)-  pure $ Lambda params body ret+  pure $ Lambda params ret body  -- | Slice a full dimension of the given size. sliceDim :: SubExp -> DimIndex SubExp

@@ -196,6 +196,7 @@ instance IsOp SizeOp where   safeOp _ = True   cheapOp _ = True+  opDependencies op = [freeIn op]  instance TypedOp SizeOp where   opType (GetSize _ _) = pure [Prim int64]@@ -290,6 +291,12 @@     -- Current GPUBody usage only benefits from hoisting kernels that     -- transfer scalars to device.     SQ.null (bodyStms body) && all ((== 0) . arrayRank) types++  opDependencies (SegOp op) = opDependencies op+  opDependencies (OtherOp op) = opDependencies op+  opDependencies op@(SizeOp {}) = [freeIn op]+  opDependencies (GPUBody _ body) =+    replicate (length . bodyResult $ body) (freeIn body)  instance (TypedOp (op rep)) => TypedOp (HostOp op rep) where   opType (SegOp op) = opType op

@@ -78,6 +78,9 @@   cheapOp (ParOp _ op) = cheapOp op   cheapOp (OtherOp op) = cheapOp op +  opDependencies (ParOp _ op) = opDependencies op+  opDependencies (OtherOp op) = opDependencies op+ instance (TypedOp (op rep)) => TypedOp (MCOp op rep) where   opType (ParOp _ op) = opType op   opType (OtherOp op) = opType op

@@ -239,6 +239,8 @@   safeOp (Inner k) = safeOp k   cheapOp (Inner k) = cheapOp k   cheapOp Alloc {} = True+  opDependencies op@(Alloc {}) = [freeIn op]+  opDependencies (Inner op) = opDependencies op  instance (CanBeWise inner) => CanBeWise (MemOp inner) where   addOpWisdom (Alloc size space) = Alloc size space

@@ -499,8 +499,7 @@             <* lexeme ":"             <*> pIntType             <* lexeme "<"-            <*> pSubExp-            <*> many ((,) <$> pLParam pr <* keyword "in" <*> pVName),+            <*> pSubExp,           keyword "while" $> WhileLoop <*> pVName         ] @@ -508,16 +507,14 @@ pLambda pr =   choice     [ lexeme "\\"-        $> lam+        $> Lambda         <*> pLParams pr         <* pColon         <*> pTypes         <* pArrow         <*> pBody pr,-      keyword "nilFn" $> Lambda mempty (Body (pBodyDec pr) mempty []) []+      keyword "nilFn" $> Lambda mempty [] (Body (pBodyDec pr) mempty [])     ]-  where-    lam params ret body = Lambda params body ret  pReduce :: PR rep -> Parser (SOAC.Reduce rep) pReduce pr =

@@ -315,23 +315,13 @@       <+> equals       <+> ppTuple' (map pretty args)       </> ( case form of-              ForLoop i it bound [] ->+              ForLoop i it bound ->                 "for"                   <+> align                     ( pretty i <> ":" <> pretty it                         <+> "<"                         <+> align (pretty bound)                     )-              ForLoop i it bound loop_vars ->-                "for"-                  <+> align-                    ( pretty i-                        <> ":"-                        <> pretty it-                        <+> "<"-                        <+> align (pretty bound)-                        </> stack (map prettyLoopVar loop_vars)-                    )               WhileLoop cond ->                 "while" <+> pretty cond           )@@ -339,7 +329,6 @@       <+> nestedBlock "{" "}" (pretty loopbody)     where       (params, args) = unzip merge-      prettyLoopVar (p, a) = pretty p <+> "in" <+> pretty a   pretty (WithAcc inputs lam) =     "with_acc"       <> parens (braces (commastack $ map ppInput inputs) <> comma </> pretty lam)@@ -355,8 +344,8 @@           )  instance (PrettyRep rep) => Pretty (Lambda rep) where-  pretty (Lambda [] (Body _ stms []) []) | stms == mempty = "nilFn"-  pretty (Lambda params body rettype) =+  pretty (Lambda [] [] (Body _ stms [])) | stms == mempty = "nilFn"+  pretty (Lambda params rettype body) =     "\\"       <+> braces (commastack $ map pretty params)       </> indent 2 (colon <+> ppTupleLines' (map pretty rettype) <+> "->")

@@ -191,9 +191,13 @@   -- | Should we try to hoist this out of branches?   cheapOp :: op -> Bool +  -- | Compute the data dependencies of an operation.+  opDependencies :: op -> [Names]+ instance IsOp (NoOp rep) where   safeOp NoOp = True   cheapOp NoOp = True+  opDependencies NoOp = []  -- | Representation-specific attributes; also means the rep supports -- some basic facilities.

@@ -166,19 +166,11 @@     consumeArg _ = mempty consumedInExp (Match _ cases defbody _) =   foldMap (consumedInBody . caseBody) cases <> consumedInBody defbody-consumedInExp (Loop merge form body) =+consumedInExp (Loop merge _ _) =   mconcat     ( map (subExpAliases . snd) $         filter (unique . paramDeclType . fst) merge     )-    <> consumedInForm form-  where-    body_consumed = consumedInBody body-    varConsumed = (`nameIn` body_consumed) . paramName . fst-    consumedInForm (ForLoop _ _ _ loopvars) =-      namesFromList $ map snd $ filter varConsumed loopvars-    consumedInForm WhileLoop {} =-      mempty consumedInExp (WithAcc inputs lam) =   mconcat (map inputConsumed inputs)     <> ( consumedByLambda lam

@@ -49,7 +49,6 @@ import Data.Map.Strict qualified as M import Data.Set qualified as S import Futhark.IR.Prop.Patterns-import Futhark.IR.Prop.Scope import Futhark.IR.Syntax import Futhark.IR.Traversals import Futhark.Util.Pretty@@ -287,7 +286,9 @@   freeIn' (Loop merge form loopbody) =     let (params, args) = unzip merge         bound_here =-          namesFromList $ M.keys $ scopeOf form <> scopeOfFParams params+          case form of+            WhileLoop {} -> namesFromList $ map paramName params+            ForLoop i _ _ -> namesFromList $ i : map paramName params      in fvBind bound_here $           freeIn' args <> freeIn' form <> freeIn' params <> freeIn' loopbody   freeIn' (WithAcc inputs lam) =@@ -363,8 +364,8 @@ instance (FreeIn dec) => FreeIn (PatElem dec) where   freeIn' (PatElem _ dec) = freeIn' dec -instance (FreeIn (LParamInfo rep)) => FreeIn (LoopForm rep) where-  freeIn' (ForLoop _ _ bound loop_vars) = freeIn' bound <> freeIn' loop_vars+instance FreeIn LoopForm where+  freeIn' (ForLoop _ _ bound) = freeIn' bound   freeIn' (WhileLoop cond) = freeIn' cond  instance (FreeIn d) => FreeIn (DimIndex d) where

@@ -82,8 +82,9 @@   [num] ->   [num] ->   num-flattenIndex dims is =-  sum $ zipWith (*) is slicesizes+flattenIndex dims is+  | length is /= length slicesizes = error "flattenIndex: length mismatch"+  | otherwise = sum $ zipWith (*) is slicesizes   where     slicesizes = drop 1 $ sliceSizes dims 

@@ -23,6 +23,7 @@     inScopeOf,     scopeOfLParams,     scopeOfFParams,+    scopeOfLoopForm,     scopeOfPat,     scopeOfPatElem,     SameScope,@@ -171,10 +172,10 @@ instance Scoped rep (VName, NameInfo rep) where   scopeOf = uncurry M.singleton -instance Scoped rep (LoopForm rep) where-  scopeOf (WhileLoop _) = mempty-  scopeOf (ForLoop i it _ xs) =-    M.insert i (IndexName it) $ scopeOfLParams (map fst xs)+-- | The scope of a loop form.+scopeOfLoopForm :: LoopForm -> Scope rep+scopeOfLoopForm (WhileLoop _) = mempty+scopeOfLoopForm (ForLoop i it _) = M.singleton i $ IndexName it  -- | The scope of a pattern. scopeOfPat :: (LetDec rep ~ dec) => Pat dec -> Scope rep

@@ -58,6 +58,7 @@ import Data.Map.Strict qualified as M import Data.Maybe import Futhark.Analysis.Alias qualified as Alias+import Futhark.Analysis.DataDependencies import Futhark.Analysis.Metrics import Futhark.Analysis.PrimExp.Convert import Futhark.Analysis.SymbolTable qualified as ST@@ -69,7 +70,7 @@ import Futhark.Optimise.Simplify.Rep import Futhark.Transform.Rename import Futhark.Transform.Substitute-import Futhark.Util (chunks, maybeNth)+import Futhark.Util (chunks, maybeNth, splitAt3) import Futhark.Util.Pretty (Doc, align, comma, commasep, docText, parens, ppTuple', pretty, (<+>), (</>)) import Futhark.Util.Pretty qualified as PP import Prelude hiding (id, (.))@@ -170,9 +171,13 @@   }   deriving (Eq, Ord, Show) +-- | What are the sizes of reduction results produced by these 'Scan's?+scanSizes :: [Scan rep] -> [Int]+scanSizes = map (length . scanNeutral)+ -- | How many reduction results are produced by these 'Scan's? scanResults :: [Scan rep] -> Int-scanResults = sum . map (length . scanNeutral)+scanResults = sum . scanSizes  -- | Combine multiple scan operators to a single operator. singleScan :: (Buildable rep) => [Scan rep] -> Scan rep@@ -189,9 +194,13 @@   }   deriving (Eq, Ord, Show) +-- | What are the sizes of reduction results produced by these 'Reduce's?+redSizes :: [Reduce rep] -> [Int]+redSizes = map (length . redNeutral)+ -- | How many reduction results are produced by these 'Reduce's? redResults :: [Reduce rep] -> Int-redResults = sum . map (length . redNeutral)+redResults = sum . redSizes  -- | Combine multiple reduction operators to a single operator. singleReduce :: (Buildable rep) => [Reduce rep] -> Reduce rep@@ -235,7 +244,7 @@  -- | A lambda with no parameters that returns no values. nilFn :: (Buildable rep) => Lambda rep-nilFn = Lambda mempty (mkBody mempty mempty) mempty+nilFn = Lambda mempty mempty (mkBody mempty mempty)  -- | Construct a Screma with possibly multiple scans, and -- the given map function.@@ -514,7 +523,7 @@   where     nms = map paramName $ take (1 + length accs) params     substs = M.fromList $ zip nms (outersize : accs)-    Lambda params _ rtp = lam+    Lambda params rtp _ = lam soacType (Scatter _w _ivs lam dests) =   zipWith arrayOfShape (map (snd . head) rets) shapes   where@@ -588,6 +597,61 @@ instance (ASTRep rep) => IsOp (SOAC rep) where   safeOp _ = False   cheapOp _ = False+  opDependencies (Stream w arrs accs lam) =+    let accs_deps = map depsOf' accs+        arrs_deps = depsOfArrays w arrs+     in lambdaDependencies mempty lam (arrs_deps <> accs_deps)+  opDependencies (Hist w arrs ops lam) =+    let bucket_fun_deps' = lambdaDependencies mempty lam (depsOfArrays w arrs)+        -- Bucket function results are indices followed by values.+        -- Reshape this to align with list of histogram operations.+        ranks = [length (histShape op) | op <- ops]+        value_lengths = [length (histNeutral op) | op <- ops]+        (indices, values) = splitAt (sum ranks) bucket_fun_deps'+        bucket_fun_deps =+          zipWith+            concatIndicesToEachValue+            (chunks ranks indices)+            (chunks value_lengths values)+     in mconcat $ zipWith (<>) bucket_fun_deps (map depsOfHistOp ops)+    where+      depsOfHistOp (HistOp dest_shape rf dests nes op) =+        let shape_deps = depsOfShape dest_shape+            in_deps = map (\vn -> oneName vn <> shape_deps <> depsOf' rf) dests+         in reductionDependencies mempty op nes in_deps+      -- A histogram operation may use the same index for multiple values.+      concatIndicesToEachValue is vs =+        let is_flat = mconcat is+         in map (is_flat <>) vs+  opDependencies (Scatter w arrs lam outputs) =+    let deps = lambdaDependencies mempty lam (depsOfArrays w arrs)+     in map flattenGroups (groupScatterResults' outputs deps)+    where+      flattenGroups (indicess, values) = mconcat indicess <> values+  opDependencies (JVP lam args vec) =+    mconcat $+      replicate 2 $+        lambdaDependencies mempty lam $+          zipWith (<>) (map depsOf' args) (map depsOf' vec)+  opDependencies (VJP lam args vec) =+    mconcat $+      replicate 2 $+        lambdaDependencies mempty lam $+          zipWith (<>) (map depsOf' args) (map depsOf' vec)+  opDependencies (Screma w arrs (ScremaForm scans reds map_lam)) =+    let (scans_in, reds_in, map_deps) =+          splitAt3 (scanResults scans) (redResults reds) $+            lambdaDependencies mempty map_lam (depsOfArrays w arrs)+        scans_deps =+          concatMap depsOfScan (zip scans $ chunks (scanSizes scans) scans_in)+        reds_deps =+          concatMap depsOfRed (zip reds $ chunks (redSizes reds) reds_in)+     in scans_deps <> reds_deps <> map_deps+    where+      depsOfScan (Scan lam nes, deps_in) =+        reductionDependencies mempty lam nes deps_in+      depsOfRed (Reduce _ lam nes, deps_in) =+        reductionDependencies mempty lam nes deps_in  substNamesInType :: M.Map VName SubExp -> Type -> Type substNamesInType _ t@Prim {} = t

@@ -521,37 +521,49 @@ -- the data dependencies to see that the "dead" result is not -- actually used for computing one of the live ones. removeDeadReduction :: BottomUpRuleOp (Wise SOACS)-removeDeadReduction (_, used) pat aux (Screma w arrs form)-  | Just ([Reduce comm redlam nes], maplam) <- isRedomapSOAC form,-    not $ all (`UT.used` used) $ patNames pat, -- Quick/cheap check-    let (red_pes, map_pes) = splitAt (length nes) $ patElems pat,-    let redlam_deps = dataDependencies $ lambdaBody redlam,-    let redlam_res = bodyResult $ lambdaBody redlam,-    let redlam_params = lambdaParams redlam,-    let used_after =-          map snd . filter ((`UT.used` used) . patElemName . fst) $-            zip red_pes redlam_params,-    let necessary =-          findNecessaryForReturned-            (`elem` used_after)-            (zip redlam_params $ map resSubExp $ redlam_res <> redlam_res)-            redlam_deps,-    let alive_mask = map ((`nameIn` necessary) . paramName) redlam_params,-    not $ and (take (length nes) alive_mask) = Simplify $ do-      let fixDeadToNeutral lives ne = if lives then Nothing else Just ne-          dead_fix = zipWith fixDeadToNeutral alive_mask nes-          (used_red_pes, _, used_nes) =-            unzip3 . filter (\(_, x, _) -> paramName x `nameIn` necessary) $-              zip3 red_pes redlam_params nes+removeDeadReduction (_, used) pat aux (Screma w arrs form) =+  case isRedomapSOAC form of+    Just ([Reduce comm redlam rednes], maplam) ->+      let mkOp lam nes' = redomapSOAC [Reduce comm lam nes']+       in removeDeadReduction' redlam rednes maplam mkOp+    _ ->+      case isScanomapSOAC form of+        Just ([Scan scanlam nes], maplam) ->+          let mkOp lam nes' = scanomapSOAC [Scan lam nes']+           in removeDeadReduction' scanlam nes maplam mkOp+        _ -> Skip+  where+    removeDeadReduction' redlam nes maplam mkOp+      | not $ all (`UT.used` used) $ patNames pat, -- Quick/cheap check+        let (red_pes, map_pes) = splitAt (length nes) $ patElems pat,+        let redlam_deps = dataDependencies $ lambdaBody redlam,+        let redlam_res = bodyResult $ lambdaBody redlam,+        let redlam_params = lambdaParams redlam,+        let used_after =+              map snd . filter ((`UT.used` used) . patElemName . fst) $+                zip red_pes redlam_params,+        let necessary =+              findNecessaryForReturned+                (`elem` used_after)+                (zip redlam_params $ map resSubExp $ redlam_res <> redlam_res)+                redlam_deps,+        let alive_mask = map ((`nameIn` necessary) . paramName) redlam_params,+        not $ and (take (length nes) alive_mask) = Simplify $ do+          let fixDeadToNeutral lives ne = if lives then Nothing else Just ne+              dead_fix = zipWith fixDeadToNeutral alive_mask nes+              (used_red_pes, _, used_nes) =+                unzip3 . filter (\(_, x, _) -> paramName x `nameIn` necessary) $+                  zip3 red_pes redlam_params nes -      let maplam' = removeLambdaResults (take (length nes) alive_mask) maplam-      redlam' <- removeLambdaResults (take (length nes) alive_mask) <$> fixLambdaParams redlam (dead_fix ++ dead_fix)+          let maplam' = removeLambdaResults (take (length nes) alive_mask) maplam+          redlam' <- removeLambdaResults (take (length nes) alive_mask) <$> fixLambdaParams redlam (dead_fix ++ dead_fix) -      auxing aux $-        letBind (Pat $ used_red_pes ++ map_pes) $-          Op $-            Screma w arrs $-              redomapSOAC [Reduce comm redlam' used_nes] maplam'+          auxing aux $+            letBind (Pat $ used_red_pes ++ map_pes) $+              Op $+                Screma w arrs $+                  mkOp redlam' used_nes maplam'+    removeDeadReduction' _ _ _ _ = Skip removeDeadReduction _ _ _ _ = Skip  -- | If we are writing to an array that is never used, get rid of it.

@@ -1005,6 +1005,8 @@   where   cheapOp _ = False   safeOp _ = True+  opDependencies op =+    replicate (length . kernelBodyResult $ segBody op) (freeIn op)  --- Simplification 

@@ -283,8 +283,8 @@ -- | The result of a body is a sequence of subexpressions. type Result = [SubExpRes] --- | A body consists of a number of bindings, terminating in a result--- (essentially a tuple literal).+-- | A body consists of a sequence of statements, terminating in a+-- list of result values. data Body rep = Body   { bodyDec :: BodyDec rep,     bodyStms :: Stms rep,@@ -435,7 +435,7 @@     -- body/ is picked.     Match [SubExp] [Case (Body rep)] (Body rep) (MatchDec (BranchType rep))   | -- | @loop {a} = {v} (for i < n|while b) do b@.-    Loop [(FParam rep, SubExp)] (LoopForm rep) (Body rep)+    Loop [(FParam rep, SubExp)] LoopForm (Body rep)   | -- | Create accumulators backed by the given arrays (which are     -- consumed) and pass them to the lambda, which must return the     -- updated accumulators and possibly some extra values.  The@@ -453,7 +453,7 @@ deriving instance (RepTypes rep) => Ord (Exp rep)  -- | For-loop or while-loop?-data LoopForm rep+data LoopForm   = ForLoop       VName       -- ^ The loop iterator var@@ -461,14 +461,8 @@       -- ^ The type of the loop iterator var       SubExp       -- ^ The number of iterations.-      [(LParam rep, VName)]   | WhileLoop VName--deriving instance (RepTypes rep) => Eq (LoopForm rep)--deriving instance (RepTypes rep) => Show (LoopForm rep)--deriving instance (RepTypes rep) => Ord (LoopForm rep)+  deriving (Eq, Ord, Show)  -- | Data associated with a branch. data MatchDec rt = MatchDec@@ -499,8 +493,8 @@ -- | Anonymous function for use in a SOAC. data Lambda rep = Lambda   { lambdaParams :: [LParam rep],-    lambdaBody :: Body rep,-    lambdaReturnType :: [Type]+    lambdaReturnType :: [Type],+    lambdaBody :: Body rep   }  deriving instance (RepTypes rep) => Eq (Lambda rep)@@ -515,7 +509,7 @@ -- | A lambda parameter. type LParam rep = Param (LParamInfo rep) --- | Function Declarations+-- | Function definitions. data FunDef rep = FunDef   { -- | Contains a value if this function is     -- an entry point.

@@ -60,6 +60,7 @@     dimFix,     sliceIndices,     sliceDims,+    sliceShape,     unitSlice,     fixSlice,     sliceSlice,@@ -388,6 +389,10 @@   where     dimSlice (DimSlice _ d _) = Just d     dimSlice DimFix {} = Nothing++-- | The shape of the array produced by this slice.+sliceShape :: Slice d -> ShapeBase d+sliceShape = Shape . sliceDims  -- | A slice with a stride of one. unitSlice :: (Num d) => d -> d -> DimIndex d

@@ -184,7 +184,7 @@ mapExpM tv (Loop merge form loopbody) = do   params' <- mapM (mapOnFParam tv) params   form' <- mapOnLoopForm tv form-  let scope = scopeOf form' <> scopeOfFParams params'+  let scope = scopeOfLoopForm form' <> scopeOfFParams params'   Loop     <$> (zip params' <$> mapM (mapOnSubExp tv) args)     <*> pure form'@@ -200,16 +200,10 @@ mapOnLoopForm ::   (Monad m) =>   Mapper frep trep m ->-  LoopForm frep ->-  m (LoopForm trep)-mapOnLoopForm tv (ForLoop i it bound loop_vars) =-  ForLoop-    <$> mapOnVName tv i-    <*> pure it-    <*> mapOnSubExp tv bound-    <*> (zip <$> mapM (mapOnLParam tv) loop_lparams <*> mapM (mapOnVName tv) loop_arrs)-  where-    (loop_lparams, loop_arrs) = unzip loop_vars+  LoopForm ->+  m LoopForm+mapOnLoopForm tv (ForLoop i it bound) =+  ForLoop <$> mapOnVName tv i <*> pure it <*> mapOnSubExp tv bound mapOnLoopForm tv (WhileLoop cond) =   WhileLoop <$> mapOnVName tv cond @@ -218,11 +212,11 @@   Mapper frep trep m ->   Lambda frep ->   m (Lambda trep)-mapOnLambda tv (Lambda params body ret) = do+mapOnLambda tv (Lambda params ret body) = do   params' <- mapM (mapOnLParam tv) params   Lambda params'-    <$> mapOnBody tv (scopeOfLParams params') body-    <*> mapM (mapOnType (mapOnSubExp tv)) ret+    <$> mapM (mapOnType (mapOnSubExp tv)) ret+    <*> mapOnBody tv (scopeOfLParams params') body  -- | Like 'mapExpM', but in the 'Identity' monad. mapExp :: Mapper frep trep Identity -> Exp frep -> Exp trep@@ -268,19 +262,15 @@ walkOnType _ Mem {} = pure () walkOnType tv (Array _ shape _) = walkOnShape tv shape -walkOnLoopForm :: (Monad m) => Walker rep m -> LoopForm rep -> m ()-walkOnLoopForm tv (ForLoop i _ bound loop_vars) =+walkOnLoopForm :: (Monad m) => Walker rep m -> LoopForm -> m ()+walkOnLoopForm tv (ForLoop i _ bound) = do   walkOnVName tv i-    >> walkOnSubExp tv bound-    >> mapM_ (walkOnLParam tv) loop_lparams-    >> mapM_ (walkOnVName tv) loop_arrs-  where-    (loop_lparams, loop_arrs) = unzip loop_vars+  walkOnSubExp tv bound walkOnLoopForm tv (WhileLoop cond) =   walkOnVName tv cond  walkOnLambda :: (Monad m) => Walker rep m -> Lambda rep -> m ()-walkOnLambda tv (Lambda params body ret) = do+walkOnLambda tv (Lambda params ret body) = do   mapM_ (walkOnLParam tv) params   walkOnBody tv (scopeOfLParams params) body   mapM_ (walkOnType tv) ret@@ -351,7 +341,7 @@   mapM_ (walkOnFParam tv) params   walkOnLoopForm tv form   mapM_ (walkOnSubExp tv) args-  let scope = scopeOfFParams params <> scopeOf form+  let scope = scopeOfFParams params <> scopeOfLoopForm form   walkOnBody tv scope loopbody   where     (params, args) = unzip merge@@ -370,5 +360,5 @@  -- | A helper for defining 'traverseOpStms'. traverseLambdaStms :: (Monad m) => OpStmsTraverser m (Lambda rep) rep-traverseLambdaStms f (Lambda ps (Body dec stms res) ret) =-  Lambda ps <$> (Body dec <$> f (scopeOfLParams ps) stms <*> pure res) <*> pure ret+traverseLambdaStms f (Lambda ps ret (Body dec stms res)) =+  Lambda ps ret <$> (Body dec <$> f (scopeOfLParams ps) stms <*> pure res)

@@ -1023,7 +1023,7 @@    checkLoopArgs -  binding (scopeOf form) $ do+  binding (scopeOfLoopForm form) $ do     form_consumable <- checkForm mergeargs form      let rettype = map paramDeclType mergepat@@ -1059,38 +1059,15 @@           map (`namesSubtract` bound_here)             <$> mapM (subExpAliasesM . resSubExp) (bodyResult loopbody)   where-    checkLoopVar (p, a) = do-      a_t <- lookupType a-      observe a-      case peelArray 1 a_t of-        Just a_t_r -> do-          checkLParamDec (paramName p) $ paramDec p-          unless (a_t_r `subtypeOf` typeOf (paramDec p)) $-            bad . TypeError $-              "Loop parameter "-                <> prettyText p-                <> " not valid for element of "-                <> prettyText a-                <> ", which has row type "-                <> prettyText a_t_r-          als <- lookupAliases a-          pure (paramName p, als)-        _ ->-          bad . TypeError $-            "Cannot loop over "-              <> prettyText a-              <> " of type "-              <> prettyText a_t-    checkForm mergeargs (ForLoop loopvar it boundexp loopvars) = do+    checkForm mergeargs (ForLoop loopvar it boundexp) = do       iparam <- primFParam loopvar $ IntType it       let mergepat = map fst merge           funparams = iparam : mergepat           paramts = map paramDeclType funparams -      consumable <- mapM checkLoopVar loopvars       boundarg <- checkArg boundexp       checkFuncall Nothing paramts $ boundarg : mergeargs-      pure consumable+      pure mempty     checkForm mergeargs (WhileLoop cond) = do       case find ((== cond) . paramName . fst) merge of         Just (condparam, _) ->@@ -1396,7 +1373,7 @@ -- parameters. checkAnyLambda ::   (Checkable rep) => Bool -> Lambda (Aliases rep) -> [Arg] -> TypeM rep ()-checkAnyLambda soac (Lambda params body rettype) args = do+checkAnyLambda soac (Lambda params rettype body) args = do   let fname = nameFromString "<anonymous>"   if length params == length args     then do

@@ -424,7 +424,7 @@   let merge = zip (shapepat ++ mergepat') args'       merge_ts = map (I.paramType . fst) merge   loopbody'' <--    localScope (scopeOfFParams $ map fst merge) . inScopeOf form' . buildBody_ $+    localScope (scopeOfFParams (map fst merge) <> scopeOfLoopForm form') . buildBody_ $       fmap subExpsRes         . ensureArgShapes           "shape of loop result does not match shapes in loop parameter"@@ -442,8 +442,10 @@   where     sparams' = map (`TypeParamDim` mempty) sparams -    forLoop mergepat' shapepat mergeinit form' =-      bodyFromStms . inScopeOf form' $ do+    forLoop mergepat' shapepat mergeinit i loopvars form' =+      bodyFromStms . localScope (scopeOfLoopForm form') $ do+        forM_ loopvars $ \(p, arr) ->+          letBindNames [I.paramName p] =<< eIndex arr [eSubExp (I.Var i)]         ses <- internaliseExp "loopres" loopbody         sets <- mapM subExpType ses         shapeargs <- argShapes (map I.paramName shapepat) mergepat' sets@@ -467,8 +469,7 @@       bindingLoopParams sparams' mergepat ts $ \shapepat mergepat' ->         bindingLambdaParams [toParam E.Observe <$> x] (map rowType arr_ts) $ \x_params -> do           let loopvars = zip x_params arr'-          forLoop mergepat' shapepat mergeinit $-            I.ForLoop i Int64 w loopvars+          forLoop mergepat' shapepat mergeinit i loopvars $ I.ForLoop i Int64 w     handleForm mergeinit (E.For i num_iterations) = do       num_iterations' <- internaliseExp1 "upper_bound" num_iterations       num_iterations_t <- I.subExpType num_iterations'@@ -477,10 +478,9 @@         _ -> error "internaliseExp Loop: invalid type"        ts <- mapM subExpType mergeinit-      bindingLoopParams sparams' mergepat ts $-        \shapepat mergepat' ->-          forLoop mergepat' shapepat mergeinit $-            I.ForLoop (E.identName i) it num_iterations' []+      bindingLoopParams sparams' mergepat ts $ \shapepat mergepat' ->+        forLoop mergepat' shapepat mergeinit (E.identName i) [] $+          I.ForLoop (E.identName i) it num_iterations'     handleForm mergeinit (E.While cond) = do       ts <- mapM subExpType mergeinit       bindingLoopParams sparams' mergepat ts $ \shapepat mergepat' -> do@@ -1237,7 +1237,7 @@         (lam_params, lam_body, lam_rettype) <-           internaliseLambda op_lam $ ne_ts ++ ne_ts         idxp <- newParam "idx" $ I.Prim int64-        let op_lam' = I.Lambda (idxp : lam_params) lam_body lam_rettype+        let op_lam' = I.Lambda (idxp : lam_params) lam_rettype lam_body         pure $ Just (op_lam', ne')       Nothing ->         pure Nothing

@@ -54,7 +54,7 @@   body' <-     localScope (scopeOfLParams params) $       lambdaWithIncrement body-  pure $ I.Lambda params body' rettype+  pure $ I.Lambda params rettype body'   where     rettype = replicate (k + 2) $ I.Prim int64     result i =

@@ -1072,10 +1072,10 @@     translateIxFnInScope _ = Nothing     se0 = intConst Int64 0     mbLoopIndexRange ::-      LoopForm (Aliases rep) ->+      LoopForm ->       Maybe (VName, (TPrimExp Int64 VName, TPrimExp Int64 VName))     mbLoopIndexRange (WhileLoop _) = Nothing-    mbLoopIndexRange (ForLoop inm _inttp seN _) = Just (inm, (pe64 se0, pe64 seN))+    mbLoopIndexRange (ForLoop inm _inttp seN) = Just (inm, (pe64 se0, pe64 seN))     addBeforeLoop actv_bef m_b etry =       case M.lookup m_b actv_bef of         Nothing -> etry@@ -1704,7 +1704,7 @@     ( computeScalarTable $         scope_table           <> scopeOfFParams (map fst loop_inits)-          <> scopeOf loop_form+          <> scopeOfLoopForm loop_form           <> scopeOf (bodyStms body)     )     (stmsToList $ bodyStms body)

@@ -202,19 +202,19 @@   foldMap (namesFromList . mapMaybe subExpVar . arrayDims . typeOf) elems  -- | The topdown handler for loops.-updateTopdownEnvLoop :: TopdownEnv rep -> [(FParam rep, SubExp)] -> LoopForm (Aliases rep) -> TopdownEnv rep+updateTopdownEnvLoop :: TopdownEnv rep -> [(FParam rep, SubExp)] -> LoopForm -> TopdownEnv rep updateTopdownEnvLoop td_env arginis lform =   let scopetab =         scope td_env           <> scopeOfFParams (map fst arginis)-          <> scopeOf lform+          <> scopeOfLoopForm lform       non_negatives =         nonNegatives td_env <> case lform of-          ForLoop v _ _ _ -> oneName v+          ForLoop v _ _ -> oneName v           _ -> mempty       less_than =         case lform of-          ForLoop v _ b _ -> [(v, primExpFromSubExp (IntType Int64) b)]+          ForLoop v _ b -> [(v, primExpFromSubExp (IntType Int64) b)]           _ -> mempty    in td_env         { scope = scopetab,

@@ -161,7 +161,7 @@ optimiseStm :: forall rep. (ASTRep rep) => Stm rep -> DoubleBufferM rep (Stms rep) optimiseStm (Let pat aux (Loop merge form body)) = do   body' <--    localScope (scopeOf form <> scopeOfFParams (map fst merge)) $+    localScope (scopeOfLoopForm form <> scopeOfFParams (map fst merge)) $       optimiseBody body   opt_loop <- asks envOptimiseLoop   (stms, pat', merge', body'') <- opt_loop pat merge body'

@@ -276,6 +276,32 @@             (TF.fsSOAC ker')             (aux1 <> aux2)       Nothing -> pure Nothing+vFuseNodeT+  _+  infusible+  (SoacNode ots1 pat1 (H.Screma w form inps) aux1, _, _)+  (TransNode stm2_out (H.Index cs slice@(Slice (ds@(DimSlice _ w' _) : ds_rest))) _, _)+    | null infusible,+      w /= w',+      ots1 == mempty,+      Just _ <- isMapSOAC form,+      [pe] <- patElems pat1 = do+        let out_t = patElemType pe `setArrayShape` sliceShape slice+            inps' = map sliceInput inps+            -- Even if we move the slice of the outermost dimension, there+            -- might still be some slicing of the inner ones.+            ots1' = ots1 H.|> H.Index cs (Slice (sliceDim w' : ds_rest))+        fusedSomething $+          SoacNode+            ots1'+            (Pat [PatElem stm2_out out_t])+            (H.Screma w' form inps')+            aux1+    where+      sliceInput inp =+        H.addTransform+          (H.Index cs (fullSlice (H.inputType inp) [ds]))+          inp vFuseNodeT _ _ _ _ = pure Nothing  resFromLambda :: Lambda rep -> Result@@ -402,7 +428,7 @@ runInnerFusionOnContext :: DepContext -> FusionM DepContext runInnerFusionOnContext c@(incoming, node, nodeT, outgoing) = case nodeT of   DoNode (Let pat aux (Loop params form body)) to_fuse ->-    doFuseScans . localScope (scopeOfFParams (map fst params) <> scopeOf form) $ do+    doFuseScans . localScope (scopeOfFParams (map fst params) <> scopeOfLoopForm form) $ do       b <- doFusionWithDelayed body to_fuse       pure (incoming, node, DoNode (Let pat aux (Loop params form b)) [], outgoing)   MatchNode (Let pat aux (Match cond cases defbody dec)) to_fuse -> doFuseScans $ do

@@ -8,11 +8,11 @@ -- by the @fgl@ package ("Data.Graph.Inductive").  The graph provided -- by this package allows nodes and edges to have arbitrarily-typed -- "labels".  It is these labels ('EdgeT', 'NodeT') that we use to--- contain Futhark-specific information.  An edge goes *from* uses of--- variables to the node that produces that variable.  There are also--- edges that do not represent normal data dependencies, but other--- things.  This means that a node can have multiple edges for the--- same name, indicating different kinds of dependencies.+-- contain Futhark-specific information.  An edge goes *from*+-- consumers to producers.  There are also edges that do not represent+-- normal data dependencies, but other things.  This means that a node+-- can have multiple edges for the same name, indicating different+-- kinds of dependencies. module Futhark.Optimise.Fusion.GraphRep   ( -- * Data structure     EdgeT (..),@@ -102,7 +102,7 @@ instance Show NodeT where   show (StmNode (Let pat _ _)) = L.intercalate ", " $ map prettyString $ patNames pat   show (SoacNode _ pat _ _) = prettyString pat-  show (TransNode _ tr _) = prettyString (show tr)+  show (TransNode _ tr _) = prettyString tr   show (ResNode name) = prettyString $ "Res: " ++ prettyString name   show (FreeNode name) = prettyString $ "Input: " ++ prettyString name   show (MatchNode stm _) = "Match: " ++ L.intercalate ", " (map prettyString $ stmNames stm)

@@ -571,7 +571,7 @@           scan_params = lambdaParams map_fun           scan_body = lambdaBody map_fun           scan_rettype = lambdaReturnType map_fun-          scan_fun' = Lambda scan_params scan_body scan_rettype+          scan_fun' = Lambda scan_params scan_rettype scan_body           nes' = map Var $ take (length map_nes) $ map paramName map_params           arrs' = drop (length map_nes) $ map paramName map_params @@ -586,7 +586,7 @@               )               $ varsRes               $ patNames map_pat-          map_fun' = Lambda map_params map_body map_rettype+          map_fun' = Lambda map_params map_rettype map_body           perm = case lambdaReturnType scan_fun of -- instead of map_fun             [] -> []             t : _ -> 1 : 0 : [2 .. arrayRank t]@@ -669,6 +669,39 @@       pure (soac', ots')     else fail "Cannot pull transpose" +pullIndex ::+  SOAC ->+  SOAC.ArrayTransforms ->+  TryFusion (SOAC, SOAC.ArrayTransforms)+pullIndex (SOAC.Screma w form inps) ots+  | SOAC.Index cs slice@(Slice (ds@(DimSlice _ w' _) : inner_ds))+      SOAC.:< ots' <-+      SOAC.viewf ots,+    w /= w',+    Just lam <- isMapSOAC form = do+      let sliceInput inp =+            SOAC.addTransform+              (SOAC.Index cs (fullSlice (SOAC.inputType inp) [ds]))+              inp+          sliceRes (SubExpRes rcs (Var v)) =+            certifying rcs+              . fmap subExpRes+              . letSubExp (baseString v <> "_sliced")+              $ BasicOp (Index v (Slice inner_ds))+          sliceRes r = pure r+          inner_changed =+            any+              ((/= stripDims 1 (sliceShape slice)) . arrayShape)+              (lambdaReturnType lam)+      lam' <-+        if not inner_changed+          then pure lam+          else+            runLambdaBuilder (lambdaParams lam) $+              mapM sliceRes =<< bodyBind (lambdaBody lam)+      pure (SOAC.Screma w' (mapSOAC lam') (map sliceInput inps), ots')+pullIndex _ _ = fail "Cannot pull index"+ pushRearrange ::   [VName] ->   SOAC ->@@ -790,6 +823,7 @@ exposeInputs inpIds ker =   (exposeInputs' =<< pushRearrange')     <|> (exposeInputs' =<< pullRearrange')+    <|> (exposeInputs' =<< pullIndex')     <|> exposeInputs' ker   where     ot = fsOutputTransform ker@@ -812,6 +846,16 @@             fsOutputTransform = SOAC.noTransforms           } +    pullIndex' = do+      (soac', ot') <- pullIndex (fsSOAC ker) ot+      unless (SOAC.nullTransforms ot') $+        fail "pullIndex was not enough"+      pure+        ker+          { fsSOAC = soac',+            fsOutputTransform = SOAC.noTransforms+          }+     exposeInputs' ker' =       case commonTransforms inpIds $ inputs ker' of         (ot', inps')@@ -823,8 +867,12 @@       | SOAC.nullTransforms ts = True     exposed inp = SOAC.inputArray inp `notElem` inpIds -outputTransformPullers :: [SOAC -> SOAC.ArrayTransforms -> TryFusion (SOAC, SOAC.ArrayTransforms)]-outputTransformPullers = [pullRearrange, pullReshape]+outputTransformPullers ::+  [ SOAC ->+    SOAC.ArrayTransforms ->+    TryFusion (SOAC, SOAC.ArrayTransforms)+  ]+outputTransformPullers = [pullRearrange, pullReshape, pullIndex]  pullOutputTransforms ::   SOAC ->

@@ -146,7 +146,7 @@                 stmsToList code1       (code1'', code1_tr_host) <- transposeFVs (freeIn kerstm) variance invar_gid code1'       let map_lam_body = mkBody code1'' $ map (SubExpRes (Certs [])) acc_vals-          map_lam0 = Lambda [Param mempty invar_gid (Prim int64)] map_lam_body el_tps+          map_lam0 = Lambda [Param mempty invar_gid (Prim int64)] el_tps map_lam_body       map_lam <- renameLambda map_lam0       (k1_res, ker1_stms) <- runBuilderT' $ do         iota <- letExp "iota" $ BasicOp $ Iota inv_dim_len (intConst Int64 0) (intConst Int64 1) Int64

@@ -186,7 +186,7 @@  optimiseExp :: (Constraints rep) => Exp (Aliases rep) -> ForwardingM rep (Exp (Aliases rep)) optimiseExp (Loop merge form body) =-  bindingScope (scopeOf form) . bindingFParams (map fst merge) $+  bindingScope (scopeOfLoopForm form) . bindingFParams (map fst merge) $     Loop merge form <$> optimiseBody body optimiseExp (Op op) = do   f <- asks topOnOp

@@ -181,7 +181,7 @@   [DesiredUpdate (LetDec rep)] ->   Pat (LetDec rep) ->   [(FParam rep, SubExp)] ->-  LoopForm rep ->+  LoopForm ->   Body rep ->   Maybe     ( m

@@ -219,8 +219,8 @@     onSOAC =       mapSOACM identitySOACMapper {mapOnSOACLambda = onLambda} -    onLambda (Lambda params body ret) =-      Lambda params <$> onBody body <*> pure ret+    onLambda (Lambda params ret body) =+      Lambda params ret <$> onBody body  -- Propagate source locations and attributes to the inlined -- statements.  Attributes are propagated only when applicable (this

@@ -80,9 +80,9 @@   AliasTable ->   Lambda (Aliases GPU) ->   PassM (Lambda GPU, Dependencies)-transformLambda aliases (Lambda params body types) = do+transformLambda aliases (Lambda params types body) = do   (body', deps) <- transformBody aliases body-  pure (Lambda params body' types, deps)+  pure (Lambda params types body', deps)  -- | Optimize a body and determine its dependencies. transformBody ::@@ -188,11 +188,15 @@       let (params, args) = unzip merge       let deps = body_deps <> depsOf params <> depsOf args <> depsOf lform -      let scope = scopeOf lform <> scopeOfFParams params+      let scope =+            scopeOfLoopForm lform <> scopeOfFParams params ::+              Scope (Aliases GPU)       let bound = IS.fromList $ map baseTag (M.keys scope)       let deps' = deps \\ bound -      let dummy = Loop merge lform (Body (bodyDec body) SQ.empty [])+      let dummy =+            Loop merge lform (Body (bodyDec body) SQ.empty []) ::+              Exp (Aliases GPU)       let Loop merge' lform' _ = removeExpAliases dummy        pure (Loop merge' lform' body', deps')

@@ -15,7 +15,7 @@ import Data.IntMap.Strict qualified as IM import Data.List (transpose, zip4) import Data.Map.Strict qualified as M-import Data.Sequence ((<|), (><), (|>))+import Data.Sequence ((><), (|>)) import Data.Text qualified as T import Futhark.Construct (fullSlice, mkBody, sliceDim) import Futhark.Error@@ -195,10 +195,7 @@          -- Read migrated scalars that are used on host.         foldM addRead (out |> stm') (zip pes pes')-      Loop ps lf b -> do-        -- Enable the migration of for-in loop variables.-        (params, lform, body) <- rewriteForIn (ps, lf, b)-+      Loop params lform body -> do         -- Update statement bound variables and parameters if their values         -- have been migrated to device.         let lmerge (res, stms, rebinds) (pe, param, StayOnHost) =@@ -306,29 +303,6 @@     addRead stms (pe@(PatElem n _), PatElem dev _)       | n == dev = pure stms       | otherwise = pe `migratedTo` (dev, stms)---- | Rewrite a for-in loop such that relevant source array reads can be delayed.-rewriteForIn ::-  ([(FParam GPU, SubExp)], LoopForm GPU, Body GPU) ->-  ReduceM ([(FParam GPU, SubExp)], LoopForm GPU, Body GPU)-rewriteForIn loop@(_, WhileLoop {}, _) =-  pure loop-rewriteForIn (params, ForLoop i t n elems, body) = do-  mt <- ask-  let (elems', stms') = foldr (inline mt) ([], bodyStms body) elems-  pure (params, ForLoop i t n elems', body {bodyStms = stms'})-  where-    inline mt (x, arr) (arrs, stms)-      | pn <- paramName x,-        not (usedOnHost pn mt) =-          let pt = typeOf x-              stm = bind (PatElem pn pt) (BasicOp $ index arr pt)-           in (arrs, stm <| stms)-      | otherwise =-          ((x, arr) : arrs, stms)--    index arr of_type =-      Index arr $ Slice $ DimFix (Var i) : map sliceDim (arrayDims of_type)  -- | Optimize an accumulator input. The 'VName' is the accumulator token. optimizeWithAccInput :: VName -> WithAccInput GPU -> ReduceM (WithAccInput GPU)

@@ -127,7 +127,7 @@   statusOf n mt /= StayOnHost shouldMoveStm (Let _ _ (Match cond _ _ _)) mt =   all ((== MoveToDevice) . (`statusOf` mt)) $ subExpVars cond-shouldMoveStm (Let _ _ (Loop _ (ForLoop _ _ (Var n) _) _)) mt =+shouldMoveStm (Let _ _ (Loop _ (ForLoop _ _ (Var n)) _)) mt =   statusOf n mt == MoveToDevice shouldMoveStm (Let _ _ (Loop _ (WhileLoop n) _)) mt =   statusOf n mt == MoveToDevice@@ -200,9 +200,6 @@      checkPats = check isArray -    checkLoopForm (ForLoop _ _ _ (_ : _)) = hostOnly-    checkLoopForm _ = ok-     checkBody = checkStms . bodyStms      checkStms stms = S.unions <$> mapM checkStm stms@@ -216,9 +213,8 @@     checkExp (Apply fn _ _ _) = Just (S.singleton fn)     checkExp (Match _ cases defbody _) =       mconcat <$> mapM checkBody (defbody : map caseBody cases)-    checkExp (Loop params lform body) = do+    checkExp (Loop params _ body) = do       checkLParams params-      checkLoopForm lform       checkBody body     checkExp BasicOp {} = Just S.empty @@ -675,7 +671,7 @@ graphLoop ::   [Binding] ->   [(FParam GPU, SubExp)] ->-  LoopForm GPU ->+  LoopForm ->   Body GPU ->   Grapher () graphLoop [] _ _ _ =@@ -698,7 +694,7 @@   -- as a whole or not. See 'shouldMoveStm'.   let may_migrate = not (bodyHostOnly stats) && may_copy_results   unless may_migrate $ case lform of-    ForLoop _ _ (Var n) _ -> connectToSink (nameToId n)+    ForLoop _ _ (Var n) -> connectToSink (nameToId n)     WhileLoop n       | Just (_, p, _, res) <- loopValueFor n -> do           connectToSink p@@ -738,7 +734,7 @@   --   -- For more details see the similar description for if statements.   when may_migrate $ case lform of-    ForLoop _ _ n _ ->+    ForLoop _ _ n ->       onlyGraphedScalarSubExp n >>= addEdges (ToNodes bindings Nothing)     WhileLoop n       | Just (_, _, arg, _) <- loopValueFor n ->@@ -784,16 +780,11 @@       -- TODO: Track memory reuse through merge parameters.        case lform of-        ForLoop _ _ n elems -> do+        ForLoop _ _ n ->           onlyGraphedScalarSubExp n >>= tellOperands-          mapM_ graphForInElem elems         WhileLoop _ -> pure ()       graphBody body       where-        graphForInElem (p, arr) = do-          when (isScalar p) $ addSource (nameToId $ paramName p, typeOf p)-          when (isArray p) $ (nameToId (paramName p), typeOf p) `reuses` arr-         graphParam ((_, t), p, arg, _) =           do             -- It is unknown whether a read can be delayed via the parameter@@ -949,7 +940,7 @@   st <- get    -- Collect statistics about the operator statements.-  let lambda = fromMaybe (Lambda [] (Body () SQ.empty []) []) op+  let lambda = fromMaybe (Lambda [] [] (Body () SQ.empty [])) op   let m = graphBody (lambdaBody lambda)   let stats = R.runReader (evalStateT (captureBodyStats m) st) env   -- We treat GPUBody kernels as host-only to not bother rewriting them inside@@ -1050,7 +1041,7 @@           captureAcc a >> collectBasic ua     collectStm stm = collect (stmExp stm) -    collectLForm (ForLoop _ _ b _) = collectSubExp b+    collectLForm (ForLoop _ _ b) = collectSubExp b     -- WhileLoop condition is declared as a loop parameter.     collectLForm (WhileLoop _) = pure () 

@@ -239,14 +239,6 @@ bindLParams params =   localVtable $ \vtable -> foldr ST.insertLParam vtable params -bindArrayLParams ::-  (SimplifiableRep rep) =>-  [LParam (Wise rep)] ->-  SimpleM rep a ->-  SimpleM rep a-bindArrayLParams params =-  localVtable $ \vtable -> foldl' (flip ST.insertLParam) vtable params- bindMerge ::   (SimplifiableRep rep) =>   [(FParam (Wise rep), SubExp, SubExpRes)] ->@@ -336,7 +328,7 @@ protectLoopHoisted ::   (SimplifiableRep rep) =>   [(FParam (Wise rep), SubExp)] ->-  LoopForm (Wise rep) ->+  LoopForm ->   SimpleM rep (a, b, Stms (Wise rep)) ->   SimpleM rep (a, b, Stms (Wise rep)) protectLoopHoisted merge form m = do@@ -357,7 +349,7 @@               find ((== cond) . paramName . fst) merge ->               pure cond_init           | otherwise -> pure $ constant True -- infinite loop-        ForLoop _ it bound _ ->+        ForLoop _ it bound ->           letSubExp "loop_nonempty" $             BasicOp $               CmpOp (CmpSlt it) (intConst it 0) bound@@ -838,18 +830,13 @@   args' <- mapM simplify args   let merge' = zip params' args'   (form', boundnames, wrapbody) <- case form of-    ForLoop loopvar it boundexp loopvars -> do+    ForLoop loopvar it boundexp -> do       boundexp' <- simplify boundexp-      let (loop_params, loop_arrs) = unzip loopvars-      loop_params' <- mapM (traverse simplify) loop_params-      loop_arrs' <- mapM simplify loop_arrs-      let form' = ForLoop loopvar it boundexp' (zip loop_params' loop_arrs')+      let form' = ForLoop loopvar it boundexp'       pure         ( form',-          namesFromList (loopvar : map paramName loop_params') <> fparamnames,-          bindLoopVar loopvar it boundexp'-            . protectLoopHoisted merge' form'-            . bindArrayLParams loop_params'+          oneName loopvar <> fparamnames,+          bindLoopVar loopvar it boundexp' . protectLoopHoisted merge' form'         )     WhileLoop cond -> do       cond' <- simplify cond@@ -1093,7 +1080,7 @@   UT.UsageTable ->   Lambda (Wise rep) ->   SimpleM rep (Lambda (Wise rep), Stms (Wise rep))-simplifyLambdaWith f blocked usage lam@(Lambda params body rettype) = do+simplifyLambdaWith f blocked usage lam@(Lambda params rettype body) = do   params' <- mapM (traverse simplify) params   let paramnames = namesFromList $ boundByLambda lam   (hoisted, body') <-@@ -1104,7 +1091,7 @@         (map (const mempty) rettype)         body   rettype' <- simplify rettype-  pure (Lambda params' body' rettype', hoisted)+  pure (Lambda params' rettype' body', hoisted)  instance Simplifiable Certs where   simplify (Certs ocs) = Certs . nubOrd . concat <$> mapM check ocs

@@ -318,17 +318,14 @@  -- | Construct a 'Wise' lambda. informLambda :: (Informing rep) => Lambda rep -> Lambda (Wise rep)-informLambda (Lambda ps body ret) = Lambda ps (informBody body) ret+informLambda (Lambda ps ret body) = Lambda ps ret (informBody body)  -- | Construct a 'Wise' expression. informExp :: (Informing rep) => Exp rep -> Exp (Wise rep) informExp (Match cond cases defbody (MatchDec ts ifsort)) =   Match cond (map (fmap informBody) cases) (informBody defbody) (MatchDec ts ifsort) informExp (Loop merge form loopbody) =-  let form' = case form of-        ForLoop i it bound params -> ForLoop i it bound params-        WhileLoop cond -> WhileLoop cond-   in Loop merge form' $ informBody loopbody+  Loop merge form $ informBody loopbody informExp e = runIdentity $ mapExpM mapper e   where     mapper =

@@ -130,7 +130,7 @@   Pat (LetDec rep) ->   StmAux (ExpDec rep) ->   ( [(FParam rep, SubExp)],-    LoopForm rep,+    LoopForm,     Body rep   ) ->   Rule rep

@@ -27,7 +27,6 @@ import Futhark.Analysis.UsageTable qualified as UT import Futhark.Construct import Futhark.IR-import Futhark.IR.Prop.Aliases import Futhark.Optimise.Simplify.Rule import Futhark.Optimise.Simplify.Rules.BasicOp import Futhark.Optimise.Simplify.Rules.Index@@ -51,7 +50,7 @@ -- | A set of standard simplification rules.  These assume pure -- functional semantics, and so probably should not be applied after -- memory block merging.-standardRules :: (BuilderOps rep, TraverseOpStms rep, Aliased rep) => RuleBook rep+standardRules :: (BuilderOps rep, TraverseOpStms rep) => RuleBook rep standardRules =   ruleBook topDownRules bottomUpRules     <> loopRules@@ -126,7 +125,7 @@  simplifyIndex :: (BuilderOps rep) => BottomUpRuleBasicOp rep simplifyIndex (vtable, used) pat@(Pat [pe]) (StmAux cs attrs _) (Index idd inds)-  | Just m <- simplifyIndexing vtable seType idd inds consumed =+  | Just m <- simplifyIndexing vtable seType idd inds consumed consuming =       Simplify $ certifying cs $ do         res <- m         attributing attrs $ case res of@@ -135,7 +134,8 @@           IndexResult extra_cs idd' inds' ->             certifying extra_cs $ letBindNames (patNames pat) $ BasicOp $ Index idd' inds'   where-    consumed = patElemName pe `UT.isConsumed` used+    consuming = (`UT.isConsumed` used)+    consumed = consuming $ patElemName pe     seType (Var v) = ST.lookupType v vtable     seType (Constant v) = Just $ Prim $ primValueType v simplifyIndex _ _ _ _ = Skip

@@ -15,7 +15,6 @@ import Futhark.Analysis.SymbolTable qualified as ST import Futhark.Construct import Futhark.IR-import Futhark.IR.Prop.Aliases import Futhark.Optimise.Simplify.Rule import Futhark.Optimise.Simplify.Rules.Loop import Futhark.Optimise.Simplify.Rules.Simple@@ -380,5 +379,5 @@  -- | A set of simplification rules for t'BasicOp's.  Includes rules -- from "Futhark.Optimise.Simplify.Rules.Simple".-basicOpRules :: (BuilderOps rep, Aliased rep) => RuleBook rep+basicOpRules :: (BuilderOps rep) => RuleBook rep basicOpRules = ruleBook topDownRules bottomUpRules <> loopRules

@@ -39,8 +39,9 @@   VName ->   Slice SubExp ->   Bool ->+  (VName -> Bool) ->   Maybe (m IndexResult)-simplifyIndexing vtable seType idd (Slice inds) consuming =+simplifyIndexing vtable seType idd (Slice inds) consuming consumed =   case defOf idd of     _       | Just t <- seType (Var idd),@@ -100,11 +101,11 @@           <$> subExpSlice (sliceSlice (primExpSlice ais) (primExpSlice (Slice inds)))     Just (Replicate (Shape [_]) (Var vv), cs)       | [DimFix {}] <- inds,-        not consuming,         ST.available vv vtable ->           Just $ pure $ SubExpResult cs $ Var vv       | DimFix {} : is' <- inds,         not consuming,+        not $ consumed vv,         ST.available vv vtable ->           Just $ pure $ IndexResult cs vv $ Slice is'     Just (Replicate (Shape [_]) val@(Constant _), cs)@@ -130,6 +131,7 @@     Just (Replicate (Shape []) (Var src), cs)       | Just dims <- arrayDims <$> seType (Var src),         length inds == length dims,+        not $ consumed src,         -- It is generally not safe to simplify a slice of a copy,         -- because the result may be used in an in-place update of the         -- original.  But we know this can only happen if the original

@@ -11,10 +11,8 @@ import Futhark.Analysis.UsageTable qualified as UT import Futhark.Construct import Futhark.IR-import Futhark.IR.Prop.Aliases import Futhark.Optimise.Simplify.Rule import Futhark.Optimise.Simplify.Rules.ClosedForm-import Futhark.Optimise.Simplify.Rules.Index import Futhark.Transform.Rename  -- This next one is tricky - it's easy enough to determine that some@@ -173,85 +171,18 @@         || (name `nameIn` namesOfInvariant)  simplifyClosedFormLoop :: (BuilderOps rep) => TopDownRuleLoop rep-simplifyClosedFormLoop _ pat _ (val, ForLoop i it bound [], body) =+simplifyClosedFormLoop _ pat _ (val, ForLoop i it bound, body) =   Simplify $ loopClosedForm pat val (oneName i) it bound body simplifyClosedFormLoop _ _ _ _ = Skip -simplifyLoopVariables :: (BuilderOps rep, Aliased rep) => TopDownRuleLoop rep-simplifyLoopVariables vtable pat aux (merge, form@(ForLoop i it num_iters loop_vars), body)-  | simplifiable <- map checkIfSimplifiable loop_vars,-    not $ all isNothing simplifiable = Simplify $ do-      -- Check if the simplifications throw away more information than-      -- we are comfortable with at this stage.-      (maybe_loop_vars, body_prefix_stms) <--        localScope (scopeOf form) $-          unzip <$> zipWithM onLoopVar loop_vars simplifiable-      if maybe_loop_vars == map Just loop_vars-        then cannotSimplify-        else do-          body' <- buildBody_ $ do-            addStms $ mconcat body_prefix_stms-            bodyBind body-          let form' = ForLoop i it num_iters $ catMaybes maybe_loop_vars-          auxing aux $ letBind pat $ Loop merge form' body'-  where-    seType (Var v)-      | v == i = Just $ Prim $ IntType it-      | otherwise = ST.lookupType v vtable-    seType (Constant v) = Just $ Prim $ primValueType v-    consumed_in_body = consumedInBody body--    vtable' = ST.fromScope (scopeOf form) <> vtable--    checkIfSimplifiable (p, arr) =-      simplifyIndexing-        vtable'-        seType-        arr-        (Slice (DimFix (Var i) : unSlice (fullSlice (paramType p) [])))-        $ paramName p `nameIn` consumed_in_body--    -- We only want this simplification if the result does not refer-    -- to 'i' at all, or does not contain accesses.-    onLoopVar (p, arr) Nothing =-      pure (Just (p, arr), mempty)-    onLoopVar (p, arr) (Just m) = do-      (x, x_stms) <- collectStms m-      case x of-        IndexResult cs arr' (Slice slice)-          | not $ any ((i `nameIn`) . freeIn) x_stms,-            DimFix (Var j) : slice' <- slice,-            j == i,-            i `notNameIn` freeIn slice -> do-              addStms x_stms-              w <- arraySize 0 <$> lookupType arr'-              for_in_partial <--                certifying cs $-                  letExp "for_in_partial" . BasicOp . Index arr' . Slice $-                    DimSlice (intConst Int64 0) w (intConst Int64 1) : slice'-              pure (Just (p, for_in_partial), mempty)-        SubExpResult cs se-          | all (notIndex . stmExp) x_stms -> do-              x_stms' <- collectStms_ $-                certifying cs $ do-                  addStms x_stms-                  letBindNames [paramName p] $ BasicOp $ SubExp se-              pure (Nothing, x_stms')-        _ -> pure (Just (p, arr), mempty)--    notIndex (BasicOp Index {}) = False-    notIndex _ = True-simplifyLoopVariables _ _ _ _ = Skip- unroll ::   (BuilderOps rep) =>   Integer ->   [(FParam rep, SubExpRes)] ->   (VName, IntType, Integer) ->-  [(LParam rep, VName)] ->   Body rep ->   RuleM rep [SubExpRes]-unroll n merge (iv, it, i) loop_vars body+unroll n merge (iv, it, i) body   | i >= n =       pure $ map snd merge   | otherwise = do@@ -261,10 +192,6 @@          letBindNames [iv] $ BasicOp $ SubExp $ intConst it i -        forM_ loop_vars $ \(p, arr) ->-          letBindNames [paramName p] . BasicOp . Index arr . Slice $-            DimFix (intConst Int64 i) : unSlice (fullSlice (paramType p) [])-         -- Some of the sizes in the types here might be temporarily wrong         -- until copy propagation fixes it up.         pure body@@ -273,24 +200,23 @@       addStms $ bodyStms iter_body'        let merge' = zip (map fst merge) $ bodyResult iter_body'-      unroll n merge' (iv, it, i + 1) loop_vars body+      unroll n merge' (iv, it, i + 1) body  simplifyKnownIterationLoop :: (BuilderOps rep) => TopDownRuleLoop rep-simplifyKnownIterationLoop _ pat aux (merge, ForLoop i it (Constant iters) loop_vars, body)+simplifyKnownIterationLoop _ pat aux (merge, ForLoop i it (Constant iters), body)   | IntValue n <- iters,     zeroIshInt n || oneIshInt n || "unroll" `inAttrs` stmAuxAttrs aux = Simplify $ do-      res <- unroll (valueIntegral n) (map (second subExpRes) merge) (i, it, 0) loop_vars body+      res <- unroll (valueIntegral n) (map (second subExpRes) merge) (i, it, 0) body       forM_ (zip (patNames pat) res) $ \(v, SubExpRes cs se) ->         certifying cs $ letBindNames [v] $ BasicOp $ SubExp se simplifyKnownIterationLoop _ _ _ _ =   Skip -topDownRules :: (BuilderOps rep, Aliased rep) => [TopDownRule rep]+topDownRules :: (BuilderOps rep) => [TopDownRule rep] topDownRules =   [ RuleLoop hoistLoopInvariantMergeVariables,     RuleLoop simplifyClosedFormLoop,-    RuleLoop simplifyKnownIterationLoop,-    RuleLoop simplifyLoopVariables+    RuleLoop simplifyKnownIterationLoop   ]  bottomUpRules :: (BuilderOps rep) => [BottomUpRule rep]@@ -299,5 +225,5 @@   ]  -- | Standard loop simplification rules.-loopRules :: (BuilderOps rep, Aliased rep) => RuleBook rep+loopRules :: (BuilderOps rep) => RuleBook rep loopRules = ruleBook topDownRules bottomUpRules

@@ -47,12 +47,9 @@ import Data.Bifunctor import Data.List (foldl') import Data.Map qualified as M-import Data.Sequence ((<|))-import Data.Sequence qualified as SQ import Futhark.Analysis.Alias qualified as Alias import Futhark.Analysis.SymbolTable qualified as ST import Futhark.Builder.Class-import Futhark.Construct (sliceDim) import Futhark.IR.Aliases import Futhark.IR.GPU import Futhark.IR.MC@@ -113,36 +110,17 @@ optimiseLoop ::   (Constraints rep) =>   Sinker rep (Op rep) ->-  Sinker rep ([(FParam rep, SubExp)], LoopForm rep, Body rep)-optimiseLoop onOp vtable sinking (merge, form, body0)-  | WhileLoop {} <- form =-      let (body1, sunk) = optimiseBody onOp vtable' sinking body0-       in ((merge, form, body1), sunk)-  | ForLoop i it bound loop_vars <- form =-      let stms' = foldr (inline i) (bodyStms body0) loop_vars-          body1 = body0 {bodyStms = stms'}-          (body2, sunk) = optimiseBody onOp vtable' sinking body1-          notSunk (x, _) = paramName x `notNameIn` sunk-          loop_vars' = filter notSunk loop_vars-          form' = ForLoop i it bound loop_vars'-          body3 = body2 {bodyStms = SQ.drop (length loop_vars') (bodyStms body2)}-       in ((merge, form', body3), sunk)+  Sinker rep ([(FParam rep, SubExp)], LoopForm, Body rep)+optimiseLoop onOp vtable sinking (merge, form, body0) =+  let (body1, sunk) = optimiseBody onOp vtable' sinking body0+   in ((merge, form, body1), sunk)   where     (params, _) = unzip merge     scope = case form of       WhileLoop {} -> scopeOfFParams params-      ForLoop i it _ _ -> M.insert i (IndexName it) $ scopeOfFParams params+      ForLoop i it _ -> M.insert i (IndexName it) $ scopeOfFParams params     vtable' = ST.fromScope scope <> vtable -    inline i (x, arr) stms =-      let pt = typeOf x-          slice = Slice $ DimFix (Var i) : map sliceDim (arrayDims pt)-          e = BasicOp (Index arr slice)-          pat = mkExpPat [Ident (paramName x) pt] e-          aux = StmAux mempty mempty (mkExpDec pat e)-          stm = Let pat aux e-       in stm <| stms- optimiseStms ::   (Constraints rep) =>   Sinker rep (Op rep) ->@@ -186,10 +164,10 @@       | Loop merge lform body <- stmExp stm =           let comps = (merge, lform, body)               (comps', loop_sunk) = optimiseLoop onOp vtable sinking comps-              (merge', lform', body') = comps'+              (merge', _, body') = comps'                (stms', stms_sunk) = optimiseStms' vtable' sinking stms-           in ( stm {stmExp = Loop merge' lform' body'} : stms',+           in ( stm {stmExp = Loop merge' lform body'} : stms',                 stms_sunk <> loop_sunk               )       | Op op <- stmExp stm =

@@ -164,7 +164,7 @@               poststms'               stms_res       -- Tiling inside for-loop.-      | Loop merge (ForLoop i it bound []) loopbody <- stmExp stm_to_tile,+      | Loop merge (ForLoop i it bound) loopbody <- stmExp stm_to_tile,         not $ any ((`nameIn` freeIn merge) . paramName . fst) merge,         Just (prestms', poststms') <-           preludeToPostlude variance prestms stm_to_tile (stmsFromList poststms) = do@@ -432,7 +432,7 @@             resultBody . map Var <$> tiledBody private' privstms'         accs' <-           letTupExp "tiled_inside_loop" $-            Loop merge' (ForLoop i it bound []) loopbody'+            Loop merge' (ForLoop i it bound) loopbody'          postludeGeneric tiling (privstms <> inloop_privstms) pat accs' poststms poststms_res res_ts @@ -699,20 +699,19 @@           <*> pure (Var mergeinit)        tile_id <- newVName "tile_id"-      let loopform = ForLoop tile_id Int64 num_whole_tiles []+      let loopform = ForLoop tile_id Int64 num_whole_tiles       loopbody <- renameBody <=< runBodyBuilder $-        inScopeOf loopform $-          localScope (scopeOfFParams $ map fst merge) $ do-            -- Collectively read a tile.-            tile <- tilingReadTile tiling TilePartial privstms (Var tile_id) inputs+        localScope (scopeOfLoopForm loopform <> scopeOfFParams (map fst merge)) $ do+          -- Collectively read a tile.+          tile <- tilingReadTile tiling TilePartial privstms (Var tile_id) inputs -            -- Now each thread performs a traversal of the tile and-            -- updates its accumulator.-            let accs =-                  map (paramName . fst) merge-                tile_args =-                  ProcessTileArgs privstms red_comm red_lam map_lam tile accs (Var tile_id)-            resultBody . map Var <$> tilingProcessTile tiling tile_args+          -- Now each thread performs a traversal of the tile and+          -- updates its accumulator.+          let accs =+                map (paramName . fst) merge+              tile_args =+                ProcessTileArgs privstms red_comm red_lam map_lam tile accs (Var tile_id)+          resultBody . map Var <$> tilingProcessTile tiling tile_args        accs <- letTupExp "accs" $ Loop merge loopform loopbody 

@@ -56,13 +56,13 @@   Builder GPU [VName] forLoop' i_bound merge body = do   i <- newVName "i" -- could give this as arg to the function-  let loop_form = ForLoop i Int64 i_bound []+  let loop_form = ForLoop i Int64 i_bound    merge_ts <- mapM lookupType merge   loop_inits <- mapM (\merge_t -> newParam "merge" $ toDecl merge_t Unique) merge_ts    loop_body <--    runBodyBuilder . inScopeOf loop_form . localScope (scopeOfFParams loop_inits) $+    runBodyBuilder . localScope (scopeOfLoopForm loop_form <> scopeOfFParams loop_inits) $       body i $         map paramName loop_inits 

@@ -25,7 +25,7 @@   Lambda SOACS ->   [SubExp] ->   m ()-bindLambda pat aux (Lambda params body _) args = do+bindLambda pat aux (Lambda params _ body) args = do   auxing aux . forM_ (zip params args) $ \(param, arg) ->     letBindNames [paramName param] $       BasicOp $ case paramType param of

@@ -74,10 +74,9 @@ transformBody (Body () stms res) = Body () <$> transformStms stms <*> pure res  transformLambda :: Lambda GPUMem -> ExpandM (Lambda GPUMem)-transformLambda (Lambda params body ret) =-  Lambda params+transformLambda (Lambda params ret body) =+  Lambda params ret     <$> localScope (scopeOfLParams params) (transformBody body)-    <*> pure ret  transformStms :: Stms GPUMem -> ExpandM (Stms GPUMem) transformStms stms =@@ -733,7 +732,7 @@   offsetMemoryInLoopParams merge $ \merge' -> do     body' <-       localScope-        (scopeOfFParams (map fst merge') <> scopeOf form)+        (scopeOfFParams (map fst merge') <> scopeOfLoopForm form)         (offsetMemoryInBody body)     pure $ Loop merge' form body' offsetMemoryInExp e = mapExpM recurse e@@ -779,8 +778,8 @@     unAllocStm _ (Let pat dec e) =       Let <$> unAllocPat pat <*> pure dec <*> mapExpM unAlloc' e -    unAllocLambda (Lambda params body ret) =-      Lambda (map unParam params) <$> unAllocBody body <*> pure ret+    unAllocLambda (Lambda params ret body) =+      Lambda (map unParam params) ret <$> unAllocBody body      unAllocPat (Pat pes) =       Pat <$> mapM (rephrasePatElem (Right . unMem)) pes@@ -865,7 +864,7 @@         forM (zip xs ys) $ \(x, y) ->           fmap subExpRes . letSubExp "z" . BasicOp $             BinOp (SMax Int64) (Var $ paramName x) (Var $ paramName y)-    pure $ Lambda (xs ++ ys) (mkBody stms zs) i64s+    pure $ Lambda (xs ++ ys) i64s (mkBody stms zs)    flat_gtid_lparam <- newParam "flat_gtid" (Prim (IntType Int64)) 

@@ -929,13 +929,12 @@   AllocM fromrep torep (Exp torep) allocInExp (Loop merge form (Body () bodystms bodyres)) =   allocInMergeParams merge $ \merge' mk_loop_val -> do-    form' <- allocInLoopForm form-    localScope (scopeOf form') $ do+    localScope (scopeOfLoopForm form) $ do       body' <-         buildBody_ . allocInStms bodystms $ do           (valctx, valres') <- mk_loop_val $ map resSubExp bodyres           pure $ subExpsRes valctx <> zipWith SubExpRes (map resCerts bodyres) valres'-      pure $ Loop merge' form' body'+      pure $ Loop merge' form body' allocInExp (Apply fname args rettype loc) = do   args' <- funcallArgs args   space <- askDefaultSpace@@ -1028,28 +1027,6 @@             handle <- asks allocInOp             handle op         }--allocInLoopForm ::-  (Allocable fromrep torep inner) =>-  LoopForm fromrep ->-  AllocM fromrep torep (LoopForm torep)-allocInLoopForm (WhileLoop v) = pure $ WhileLoop v-allocInLoopForm (ForLoop i it n loopvars) =-  ForLoop i it n <$> mapM allocInLoopVar loopvars-  where-    allocInLoopVar (p, a) = do-      (mem, ixfun) <- lookupArraySummary a-      case paramType p of-        Array pt shape u -> do-          dims <- map pe64 . arrayDims <$> lookupType a-          let ixfun' = IxFun.slice ixfun $ fullSliceNum dims [DimFix $ le64 i]-          pure (p {paramDec = MemArray pt shape u $ ArrayIn mem ixfun'}, a)-        Prim bt ->-          pure (p {paramDec = MemPrim bt}, a)-        Mem space ->-          pure (p {paramDec = MemMem space}, a)-        Acc acc ispace ts u ->-          pure (p {paramDec = MemAcc acc ispace ts u}, a)  class SizeSubst op where   opIsConst :: op -> Bool

@@ -340,10 +340,9 @@     MatchDec (staticShapes (patTypes pat)) MatchEquiv  transformLambda :: KernelPath -> Lambda SOACS -> DistribM (Lambda GPU)-transformLambda path (Lambda params body ret) =-  Lambda params+transformLambda path (Lambda params ret body) =+  Lambda params ret     <$> localScope (scopeOfLParams params) (transformBody path body)-    <*> pure ret  transformStm :: KernelPath -> Stm SOACS -> DistribM GPUStms transformStm _ stm@@ -364,15 +363,10 @@     transformInput (shape, arrs, op) =       (shape, arrs, fmap (first soacsLambdaToGPU) op) transformStm path (Let pat aux (Loop merge form body)) =-  localScope (castScope (scopeOf form) <> scopeOfFParams params) $-    oneStm . Let pat aux . Loop merge form' <$> transformBody path body+  localScope (scopeOfLoopForm form <> scopeOfFParams params) $+    oneStm . Let pat aux . Loop merge form <$> transformBody path body   where     params = map fst merge-    form' = case form of-      WhileLoop cond ->-        WhileLoop cond-      ForLoop i it bound ps ->-        ForLoop i it bound ps transformStm path (Let pat aux (Op (Screma w arrs form)))   | Just lam <- isMapSOAC form =       onMap path $ MapLoop pat aux w lam arrs

@@ -40,7 +40,7 @@           scan_params = lambdaParams map_fun           scan_body = lambdaBody map_fun           scan_rettype = lambdaReturnType map_fun-          scan_fun' = Lambda scan_params scan_body scan_rettype+          scan_fun' = Lambda scan_params scan_rettype scan_body           scan_input' =             map (first Var) $               uncurry zip $@@ -58,7 +58,7 @@               )               $ varsRes               $ patNames map_pat-          map_fun' = Lambda map_params map_body map_rettype+          map_fun' = Lambda map_params map_rettype map_body        res_pat' <-         fmap basicPat $@@ -113,7 +113,7 @@           red_params = lambdaParams map_fun           red_body = lambdaBody map_fun           red_rettype = lambdaReturnType map_fun-          red_fun' = Lambda red_params red_body red_rettype+          red_fun' = Lambda red_params red_rettype red_body           red_input' = zip accs' $ map paramName map_params           red_pat = stripPatOuterDim map_pat @@ -134,7 +134,7 @@             map_body_stms <- collectStms_ m             pure $ mkBody map_body_stms $ varsRes $ patNames map_pat -      let map_fun' = Lambda map_params map_body map_rettype+      let map_fun' = Lambda map_params map_rettype map_body        addStm $         Let res_pat (StmAux map_cs mempty ()) $

@@ -34,7 +34,7 @@ -- | An encoding of a sequential do-loop with no existential context, -- alongside its result pattern. data SeqLoop-  = SeqLoop [Int] (Pat Type) [(FParam SOACS, SubExp)] (LoopForm SOACS) (Body SOACS)+  = SeqLoop [Int] (Pat Type) [(FParam SOACS, SubExp)] LoopForm (Body SOACS)  loopPerm :: SeqLoop -> [Int] loopPerm (SeqLoop perm _ _ _ _) = perm@@ -74,7 +74,7 @@         localScope (scopeOfLParams new_params) $           unzip . catMaybes <$> mapM copyOrRemoveParam params_and_arrs -    let lam = Lambda (params' <> new_params) body rettype+    let lam = Lambda (params' <> new_params) rettype body         map_stm =           Let loop_pat_expanded aux $             Op $@@ -146,8 +146,8 @@       lam =         Lambda           params-          (mkBody stms' $ varsRes res')           (map rowType $ patTypes (loopNestingPat n))+          (mkBody stms' $ varsRes res')    in ( patNames $ loopNestingPat n,         oneStm $           Let (loopNestingPat n) (loopNestingAux n) $@@ -214,7 +214,7 @@           Pat $ map (fmap (`arrayOfRow` w)) $ patElems branch_pat          mkBranch branch = (renameBody =<<) $ do-          let lam = Lambda params branch lam_ret+          let lam = Lambda params lam_ret branch               res = varsRes $ patNames branch_pat'               map_stm = Let branch_pat' aux $ Op $ Screma w arrs $ mapSOAC lam           pure $ mkBody (oneStm map_stm) res@@ -262,7 +262,7 @@           Iota w (intConst Int64 0) (intConst Int64 1) Int64       let (params, arrs) = unzip params_and_arrs           maplam_ret = lambdaReturnType acc_lam-          maplam = Lambda (iota_p : orig_acc_params ++ params) (lambdaBody acc_lam) maplam_ret+          maplam = Lambda (iota_p : orig_acc_params ++ params) maplam_ret (lambdaBody acc_lam)       auxing map_aux . fmap subExpsRes . letTupExp' "withacc_inter" $         Op $           Screma w (iota_w : map paramName acc_params ++ arrs) (mapSOAC maplam)@@ -299,9 +299,9 @@       trParam :: Param (TypeBase shape u) -> Param (TypeBase shape u)       trParam = fmap trType -      trLam i (Lambda params body ret) =+      trLam i (Lambda params ret body) =         localScope (scopeOfLParams params) $-          Lambda (map trParam params) <$> trBody i body <*> pure (map trType ret)+          Lambda (map trParam params) (map trType ret) <$> trBody i body        trBody i (Body dec stms res) =         inScopeOf stms $ Body dec <$> traverse (trStm i) stms <*> pure res

@@ -203,14 +203,10 @@    case e of     Loop merge form loopbody ->-      localScope (scopeOf form' <> scopeOfFParams (map fst merge)) $ do+      localScope (scopeOfLoopForm form <> scopeOfFParams (map fst merge)) $ do         loopbody' <- intraGroupBody loopbody         certifying (stmAuxCerts aux) . letBind pat $-          Loop merge form' loopbody'-      where-        form' = case form of-          ForLoop i it bound inps -> ForLoop i it bound inps-          WhileLoop cond -> WhileLoop cond+          Loop merge form loopbody'     Match cond cases defbody ifdec -> do       cases' <- mapM (traverse intraGroupBody) cases       defbody' <- intraGroupBody defbody

@@ -101,21 +101,16 @@   let space = SegSpace flat [(gtid, w)]   pure (gtid, space) -transformLoopForm :: LoopForm SOACS -> LoopForm MC-transformLoopForm (WhileLoop cond) = WhileLoop cond-transformLoopForm (ForLoop i it bound params) = ForLoop i it bound params- transformStm :: Stm SOACS -> ExtractM (Stms MC) transformStm (Let pat aux (BasicOp op)) =   pure $ oneStm $ Let pat aux $ BasicOp op transformStm (Let pat aux (Apply f args ret info)) =   pure $ oneStm $ Let pat aux $ Apply f args ret info transformStm (Let pat aux (Loop merge form body)) = do-  let form' = transformLoopForm form   body' <--    localScope (scopeOfFParams (map fst merge) <> scopeOf form') $+    localScope (scopeOfFParams (map fst merge) <> scopeOfLoopForm form) $       transformBody body-  pure $ oneStm $ Let pat aux $ Loop merge form' body'+  pure $ oneStm $ Let pat aux $ Loop merge form body' transformStm (Let pat aux (Match ses cases defbody ret)) =   oneStm . Let pat aux     <$> (Match ses <$> mapM transformCase cases <*> transformBody defbody <*> pure ret)@@ -131,10 +126,9 @@   fmap (certify (stmAuxCerts aux)) <$> transformSOAC pat (stmAuxAttrs aux) op  transformLambda :: Lambda SOACS -> ExtractM (Lambda MC)-transformLambda (Lambda params body ret) =-  Lambda params+transformLambda (Lambda params ret body) =+  Lambda params ret     <$> localScope (scopeOfLParams params) (transformBody body)-    <*> pure ret  transformStms :: Stms SOACS -> ExtractM (Stms MC) transformStms stms =

@@ -0,0 +1,65 @@+-- | Profiling information emitted by a running Futhark program.+module Futhark.Profile+  ( ProfilingEvent (..),+    ProfilingReport (..),+    profilingInfoFromText,+  )+where++import Data.Aeson qualified as JSON+import Data.Aeson.Key qualified as JSON+import Data.Aeson.KeyMap qualified as JSON+import Data.Bifunctor+import Data.ByteString.Builder (toLazyByteString)+import Data.Map qualified as M+import Data.Text qualified as T+import Data.Text.Encoding (encodeUtf8Builder)++-- | A thing that has occurred during execution.+data ProfilingEvent = ProfilingEvent+  { -- | Short, single line.+    eventName :: T.Text,+    -- | In microseconds.+    eventDuration :: Double,+    -- | Long, may be multiple lines.+    eventDescription :: T.Text+  }+  deriving (Eq, Ord, Show)++instance JSON.ToJSON ProfilingEvent where+  toJSON (ProfilingEvent name duration description) =+    JSON.object+      [ ("name", JSON.toJSON name),+        ("duration", JSON.toJSON duration),+        ("description", JSON.toJSON description)+      ]++instance JSON.FromJSON ProfilingEvent where+  parseJSON = JSON.withObject "event" $ \o ->+    ProfilingEvent+      <$> o JSON..: "name"+      <*> o JSON..: "duration"+      <*> o JSON..: "description"++data ProfilingReport = ProfilingReport+  { profilingEvents :: [ProfilingEvent],+    -- | Mapping memory spaces to bytes.+    profilingMemory :: M.Map T.Text Integer+  }+  deriving (Eq, Ord, Show)++instance JSON.ToJSON ProfilingReport where+  toJSON (ProfilingReport events memory) =+    JSON.object+      [ ("events", JSON.toJSON events),+        ("memory", JSON.object $ map (bimap JSON.fromText JSON.toJSON) $ M.toList memory)+      ]++instance JSON.FromJSON ProfilingReport where+  parseJSON = JSON.withObject "profiling-info" $ \o ->+    ProfilingReport+      <$> o JSON..: "events"+      <*> (JSON.toMapText <$> o JSON..: "memory")++profilingInfoFromText :: T.Text -> Maybe ProfilingReport+profilingInfoFromText = JSON.decode . toLazyByteString . encodeUtf8Builder

@@ -158,11 +158,11 @@             zip mapout_params $ map Var map_arrs           ]   i <- newVName "i"-  let loopform = ForLoop i Int64 w []+  let loopform = ForLoop i Int64 w       lam_cons = consumedByLambda $ Alias.analyseLambda mempty map_lam    loop_body <- runBodyBuilder-    . localScope (scopeOfFParams (map fst merge) <> scopeOf loopform)+    . localScope (scopeOfFParams (map fst merge) <> scopeOfLoopForm loopform)     $ do       -- Bind the parameters to the lambda.       forM_ (zip3 (lambdaParams map_lam) arrs arr_ts) $ \(p, arr, arr_t) ->@@ -258,13 +258,13 @@    i <- newVName "i" -  let loop_form = ForLoop i Int64 w []+  let loop_form = ForLoop i Int64 w    letBindNames [paramName chunk_size_param] . BasicOp . SubExp $     intConst Int64 1    loop_body <- runBodyBuilder $-    localScope (scopeOf loop_form <> scopeOfFParams merge_params) $ do+    localScope (scopeOfLoopForm loop_form <> scopeOfFParams merge_params) $ do       let slice = [DimSlice (Var i) (Var (paramName chunk_size_param)) (intConst Int64 1)]       forM_ (zip chunk_params arrs) $ \(p, arr) ->         letBindNames [paramName p] . BasicOp . Index arr $@@ -308,7 +308,7 @@          foldM saveInArray arr indexes'       pure $ resultBody (map Var ress)-  letBind pat $ Loop merge (ForLoop iter Int64 len []) loopBody+  letBind pat $ Loop merge (ForLoop iter Int64 len) loopBody transformSOAC pat (Hist len imgs ops bucket_fun) = do   iter <- newVName "iter" @@ -364,7 +364,7 @@     pure $ resultBody $ map Var $ concat hists_out''    -- Wrap up the above into a for-loop.-  letBind pat $ Loop merge (ForLoop iter Int64 len []) loopBody+  letBind pat $ Loop merge (ForLoop iter Int64 len) loopBody  -- | Recursively first-order-transform a lambda. transformLambda ::@@ -378,12 +378,12 @@   ) =>   Lambda SOACS ->   m (AST.Lambda rep)-transformLambda (Lambda params body rettype) = do+transformLambda (Lambda params rettype body) = do   body' <-     runBodyBuilder $       localScope (scopeOfLParams params) $         transformBody body-  pure $ Lambda params body' rettype+  pure $ Lambda params rettype body'  letwith :: (Transformer m) => [VName] -> SubExp -> [SubExp] -> m [VName] letwith ks i vs = do@@ -401,7 +401,7 @@   AST.Lambda (Rep m) ->   [AST.Exp (Rep m)] ->   m Result-bindLambda (Lambda params body _) args = do+bindLambda (Lambda params _ body) args = do   forM_ (zip params args) $ \(param, arg) ->     if primType $ paramType param       then letBindNames [paramName param] arg

@@ -270,19 +270,16 @@       -- It is important that 'i' is renamed before the loop_vars, as       -- 'i' may be used in the annotations for loop_vars (e.g. index       -- functions).-      ForLoop i it boundexp loop_vars -> renameBound [i] $ do-        let (arr_params, loop_arrs) = unzip loop_vars+      ForLoop i it boundexp -> renameBound [i] $ do         boundexp' <- rename boundexp-        loop_arrs' <- rename loop_arrs-        renameBound (map paramName params ++ map paramName arr_params) $ do+        renameBound (map paramName params) $ do           params' <- mapM rename params-          arr_params' <- mapM rename arr_params           i' <- rename i           loopbody' <- rename loopbody           pure $             Loop               (zip params' args')-              (ForLoop i' it boundexp' $ zip arr_params' loop_arrs')+              (ForLoop i' it boundexp')               loopbody'       WhileLoop cond ->         renameBound (map paramName params) $ do@@ -315,12 +312,9 @@     Acc <$> rename acc <*> rename ispace <*> rename ts <*> pure u  instance (Renameable rep) => Rename (Lambda rep) where-  rename (Lambda params body ret) =-    renameBound (map paramName params) $ do-      params' <- mapM rename params-      body' <- rename body-      ret' <- mapM rename ret-      pure $ Lambda params' body' ret'+  rename (Lambda params ret body) =+    renameBound (map paramName params) $+      Lambda <$> mapM rename params <*> mapM rename ret <*> rename body  instance Rename Names where   rename = fmap namesFromList . mapM rename . namesToList

@@ -170,11 +170,11 @@     Mem space  instance (Substitutable rep) => Substitute (Lambda rep) where-  substituteNames substs (Lambda params body rettype) =+  substituteNames substs (Lambda params rettype body) =     Lambda       (substituteNames substs params)-      (substituteNames substs body)       (map (substituteNames substs) rettype)+      (substituteNames substs body)  instance Substitute Ident where   substituteNames substs v =

@@ -104,7 +104,7 @@ prettyText :: (Pretty a) => a -> Text prettyText = docText . pretty --- | Convert a 'Doc' to text.  Thsi ignores any annotations (i.e. it+-- | Convert a 'Doc' to text.  This ignores any annotations (i.e. it -- will be non-coloured output). docText :: Doc a -> T.Text docText = Prettyprinter.Render.Text.renderStrict . layouter

@@ -179,6 +179,8 @@    @binop                   { tokenS $ symbol [] . nameFromText }   @qualbinop               { tokenS $ uncurry symbol . mkQualId }++  .                        { tokenS ERROR } {  nameFromBS :: BS.ByteString -> Name

@@ -129,6 +129,7 @@   | DOC T.Text   | EOF   | HOLE+  | ERROR T.Text   deriving (Show, Eq, Ord)  mkQualId :: T.Text -> ([Name], Name)

@@ -194,6 +194,9 @@ parseError (L loc DOC {}, _) =   parseErrorAt (locOf loc) $     Just "Documentation comments ('-- |') are only permitted when preceding declarations."+parseError (L loc (ERROR "\""), _) =+  parseErrorAt (locOf loc) $+    Just "Unclosed string literal." parseError (L loc _, expected) = do   input <- lift $ gets parserInput   let ~(Loc (Pos _ _ _ beg) (Pos _ _ _ end)) = locOf loc

@@ -1680,20 +1680,6 @@ negativeIshInt (Int32Value k) = k < 0 negativeIshInt (Int64Value k) = k < 0 --- | The size of a value of a given primitive type in bits.-primBitSize :: PrimType -> Int-primBitSize = (* 8) . primByteSize---- | The size of a value of a given primitive type in eight-bit bytes.------ Warning: note that this is 0 for 'Unit', but a 'Unit' takes up a--- byte in the binary data format.-primByteSize :: (Num a) => PrimType -> a-primByteSize (IntType t) = intByteSize t-primByteSize (FloatType t) = floatByteSize t-primByteSize Bool = 1-primByteSize Unit = 0- -- | The size of a value of a given integer type in eight-bit bytes. intByteSize :: (Num a) => IntType -> a intByteSize Int8 = 1@@ -1706,6 +1692,20 @@ floatByteSize Float16 = 2 floatByteSize Float32 = 4 floatByteSize Float64 = 8++-- | The size of a value of a given primitive type in eight-bit bytes.+--+-- Warning: note that this is 0 for 'Unit', but a 'Unit' takes up a+-- byte in the binary data format.+primByteSize :: (Num a) => PrimType -> a+primByteSize (IntType t) = intByteSize t+primByteSize (FloatType t) = floatByteSize t+primByteSize Bool = 1+primByteSize Unit = 0++-- | The size of a value of a given primitive type in bits.+primBitSize :: PrimType -> Int+primBitSize = (* 8) . primByteSize  -- | True if the given binary operator is commutative. commutativeBinOp :: BinOp -> Bool

@@ -5,6 +5,7 @@ import Data.Map qualified as M import Data.Text qualified as T import Futhark.Bench+import Futhark.ProfileTests () import Test.Tasty import Test.Tasty.QuickCheck @@ -12,7 +13,7 @@   arbitrary = RunResult . getPositive <$> arbitrary  printable :: Gen String-printable = getPrintableString <$> arbitrary+printable = getASCIIString <$> arbitrary  instance Arbitrary DataResult where   arbitrary =@@ -25,6 +26,7 @@                     <$> arbitrary                     <*> arbMap                     <*> oneof [pure Nothing, Just <$> arbText]+                    <*> arbitrary                 )         ]     where

@@ -0,0 +1,23 @@+{-# OPTIONS_GHC -fno-warn-orphans #-}++module Futhark.ProfileTests () where++import Data.Map qualified as M+import Data.Text qualified as T+import Futhark.Profile+import Test.Tasty.QuickCheck++printable :: Gen String+printable = getPrintableString <$> arbitrary++arbText :: Gen T.Text+arbText = T.pack <$> printable++instance Arbitrary ProfilingEvent where+  arbitrary = ProfilingEvent <$> arbText <*> arbitrary <*> arbText++instance Arbitrary ProfilingReport where+  arbitrary =+    ProfilingReport+      <$> arbitrary+      <*> (M.fromList <$> listOf ((,) <$> arbText <*> arbitrary))