sync 2e7068faf55e

Browse files

Files changed (16) hide show

README.md +88 -0
build/webgpu/bench.json +127 -0
build/webgpu/manifest.json +1061 -0
build/webgpu/metadata.json +28 -0
build/webgpu/reduce-axis-split-reduce.wgsl.jinja +71 -0
build/webgpu/reduce-axis0-splitk-combine.wgsl.jinja +54 -0
build/webgpu/reduce-axis0-splitk-reduce.wgsl.jinja +68 -0
build/webgpu/reduce-axis0-tilecols.wgsl.jinja +174 -0
build/webgpu/reduce-flat-combine-logsumexp.wgsl.jinja +88 -0
build/webgpu/reduce-flat-partial-logsumexp.wgsl.jinja +101 -0
build/webgpu/reduce-i32-axes02.wgsl.jinja +33 -0
build/webgpu/reduce-noop-empty-axes.wgsl.jinja +13 -0
build/webgpu/reduce-row-subgroup.wgsl.jinja +110 -0
build/webgpu/reduce-row-tree.wgsl.jinja +161 -0
build/webgpu/reduce-serial-axis.wgsl.jinja +196 -0
build/webgpu/test.json +0 -0

README.md CHANGED Viewed

@@ -1,3 +1,91 @@
 ---
 license: apache-2.0
 ---

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.ReduceLogSumExp
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 18
+## Description
+Computes `log(sum(exp(x)))` over the specified axes of the input tensor. The output rank matches the input when `keepdims` is 1; reduced dimensions are pruned when `keepdims` is 0. Reduction over an empty set of values yields negative infinity.
+See the [ONNX `ReduceLogSumExp` spec](https://onnx.ai/onnx/operators/onnx__ReduceLogSumExp.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `data` | `x` | `T` | — | — | The input tensor to reduce. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `reduced` | `y` | `T` | derived | — | The reduced output tensor. | required |
+## Attributes
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `keepdims` | `1` | If 1, retains the reduced dimension with size 1 in the output; if 0, the reduced dimension is removed. |
+| `noop_with_empty_axes` | `0` | When 1 and `axes` is empty, acts as an identity (no reduction); when 0 and `axes` is empty, reduces over all axes. |
+| `axes` | `[]` | Values of the optional ONNX `axes` tensor input, supplied through this request attribute; an empty list follows `noop_with_empty_axes`. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16`, `int32` |
+## Device requirements
+Some implementation variants require `subgroups`. These are route-specific capabilities, not package-wide requirements; availability also depends on the request shape and dtype.
+## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
+- [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
+- [`test.json`](build/webgpu/test.json) — correctness cases
+- [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
+- [`reduce-axis-split-reduce.wgsl.jinja`](build/webgpu/reduce-axis-split-reduce.wgsl.jinja)
+- [`reduce-axis0-splitk-combine.wgsl.jinja`](build/webgpu/reduce-axis0-splitk-combine.wgsl.jinja)
+- [`reduce-axis0-splitk-reduce.wgsl.jinja`](build/webgpu/reduce-axis0-splitk-reduce.wgsl.jinja)
+- [`reduce-axis0-tilecols.wgsl.jinja`](build/webgpu/reduce-axis0-tilecols.wgsl.jinja)
+- [`reduce-flat-combine-logsumexp.wgsl.jinja`](build/webgpu/reduce-flat-combine-logsumexp.wgsl.jinja)
+- [`reduce-flat-partial-logsumexp.wgsl.jinja`](build/webgpu/reduce-flat-partial-logsumexp.wgsl.jinja)
+- [`reduce-i32-axes02.wgsl.jinja`](build/webgpu/reduce-i32-axes02.wgsl.jinja)
+- [`reduce-noop-empty-axes.wgsl.jinja`](build/webgpu/reduce-noop-empty-axes.wgsl.jinja)
+- [`reduce-row-subgroup.wgsl.jinja`](build/webgpu/reduce-row-subgroup.wgsl.jinja)
+- [`reduce-row-tree.wgsl.jinja`](build/webgpu/reduce-row-tree.wgsl.jinja)
+- [`reduce-serial-axis.wgsl.jinja`](build/webgpu/reduce-serial-axis.wgsl.jinja)
+## Use with `@huggingface/kernels`
+The loader automatically allocates outputs whose metadata it can derive from the manifest contract and this call.
+The explicit `outputs` entries provide shape and logical dtype metadata for the results listed below:
+- `y`
+Each entry either requests an optional result or supplies metadata that cannot be inferred from the inputs.
+The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+Replace each `*Data` placeholder with a typed array containing the corresponding input data.
+```js
+import { getKernel } from "@huggingface/kernels";
+const kernel = await getKernel("webgpu-kernels/ai.onnx.ReduceLogSumExp", { version: 1 });
+// Explicit destinations request optional results or supply metadata that cannot be inferred.
+const { y } = await kernel({ x: { data: xData, shape: [] } }, {
+  outputs: { y: { shape: [], dtype: "float32" } },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,127 @@

+{
+  "op": "ai.onnx.ReduceLogSumExp",
+  "cases": [
+    {
+      "name": "1024x1024_axis1",
+      "attrs": { "axes": [1], "keepdims": 0 },
+      "inputs": { "x": { "dtype": "float32", "shape": [1024, 1024] } },
+      "outputs": { "y": { "dtype": "float32", "shape": [1024] } }
+    },
+    {
+      "name": "reducelogsumexp-fullreduce-r3-256x256x256-keepdims",
+      "preset": "smoke",
+      "vars": { "rows": 65536, "cols": 256 },
+      "attrs": { "keepdims": 1 },
+      "inputs": { "x": { "shape": [256, 256, 256], "dtype": "float32", "dist": "normal", "seed": 811, "scale": 0.5 } },
+      "outputs": { "y": { "shape": [1, 1, 1], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "name": "stable max-plus-exp traffic", "value": "2 * args.rows * args.cols * 4" }
+        ]
+      }
+    },
+    {
+      "name": "reducelogsumexp-axis0-f32-1024x512",
+      "preset": "smoke",
+      "vars": { "rows": 1024, "cols": 512 },
+      "attrs": { "axes": [0], "keepdims": 0 },
+      "inputs": { "x": { "shape": [1024, 512], "dtype": "float32", "dist": "normal", "seed": 118, "scale": 0.2 } },
+      "outputs": { "y": { "shape": [512], "dtype": "float32" } },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          { "type": "bandwidth", "name": "stable max-plus-exp traffic", "value": "2 * args.rows * args.cols * 4" }
+        ]
+      }
+    },
+    {
+      "name": "reducelogsumexp-lastaxis-f32-4096x4096",
+      "preset": "smoke",
+      "vars": { "rows": 4096, "cols": 4096 },
+      "attrs": { "axes": [1], "keepdims": 0 },
+      "inputs": { "x": { "shape": [4096, 4096], "dtype": "float32", "dist": "normal", "seed": 117, "scale": 0.2 } },
+      "outputs": { "y": { "shape": [4096], "dtype": "float32" } },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          { "type": "bandwidth", "name": "stable max-plus-exp traffic", "value": "2 * args.rows * args.cols * 4" }
+        ]
+      }
+    },
+    {
+      "name": "axis0-large-65536x512",
+      "preset": "smoke",
+      "vars": { "rows": 65536, "cols": 512 },
+      "attrs": { "axes": [0], "keepdims": 0 },
+      "inputs": { "x": { "shape": [65536, 512], "dtype": "float32", "dist": "normal", "scale": 0.2, "seed": 311 } },
+      "outputs": { "y": { "shape": [512], "dtype": "float32" } },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          { "type": "bandwidth", "name": "stable max-plus-exp traffic", "value": "2 * args.rows * args.cols * 4" }
+        ]
+      }
+    },
+    {
+      "name": "axis0-tall-16384x1024",
+      "preset": "smoke",
+      "vars": { "rows": 16384, "cols": 1024 },
+      "attrs": { "axes": [0], "keepdims": 0 },
+      "inputs": { "x": { "shape": [16384, 1024], "dtype": "float32", "dist": "normal", "scale": 0.2, "seed": 312 } },
+      "outputs": { "y": { "shape": [1024], "dtype": "float32" } },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          { "type": "bandwidth", "name": "stable max-plus-exp traffic", "value": "2 * args.rows * args.cols * 4" }
+        ]
+      }
+    },
+    {
+      "name": "fullreduce-single-lane-r3-101x101x101-numel-not-mult4",
+      "preset": "stress",
+      "vars": { "numel": 1030301 },
+      "attrs": { "keepdims": 1 },
+      "inputs": { "x": { "shape": [101, 101, 101], "dtype": "float32", "dist": "normal", "seed": 733, "scale": 0.3 } },
+      "outputs": { "y": { "shape": [1, 1, 1], "dtype": "float32", "dist": "empty" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "name": "stable max-plus-exp traffic", "value": "2 * numel(shapes.x) * 4" }]
+      }
+    },
+    {
+      "name": "axis0-few-lane-starvation-131072x8",
+      "preset": "stress",
+      "vars": { "rows": 131072, "cols": 8 },
+      "attrs": { "axes": [0], "keepdims": 0 },
+      "inputs": { "x": { "shape": [131072, 8], "dtype": "float32", "dist": "normal", "seed": 941, "scale": 0.2 } },
+      "outputs": { "y": { "shape": [8], "dtype": "float32", "dist": "empty" } },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          { "type": "bandwidth", "name": "stable max-plus-exp traffic", "value": "2 * args.rows * args.cols * 4" }
+        ]
+      }
+    },
+    {
+      "name": "reducelogsumexp-spatial-axes12-f32-16x256x256-low-lane",
+      "preset": "stress",
+      "provenance": {
+        "source": "authored for variant coverage",
+        "notes": "Low-lane rank-3 contiguous-suffix reduction that verifies the vec4 subgroup reducer and no-subgroup workgroup-tree fallback."
+      },
+      "vars": { "batch": 16, "height": 256, "width": 256 },
+      "attrs": { "axes": [1, 2], "keepdims": 1 },
+      "inputs": { "x": { "shape": [16, 256, 256], "dtype": "float32", "dist": "normal", "seed": 910, "scale": 0.2 } },
+      "outputs": { "y": { "shape": [16, 1, 1], "dtype": "float32", "dist": "empty" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "name": "stable max-plus-exp traffic",
+            "value": "2 * args.batch * args.height * args.width * 4"
+          }
+        ]
+      }
+    }
+  ],
+  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] }
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,1061 @@

+{
+  "domain": "ai.onnx",
+  "name": "ReduceLogSumExp",
+  "sinceVersion": 18,
+  "description": "Computes `log(sum(exp(x)))` over the specified axes of the input tensor. The output rank matches the input when `keepdims` is 1; reduced dimensions are pruned when `keepdims` is 0. Reduction over an empty set of values yields negative infinity.",
+  "inputs": [{ "role": "data", "dtype": "T", "description": "The input tensor to reduce." }],
+  "outputs": [
+    {
+      "role": "reduced",
+      "dtype": "T",
+      "rank": "ranks.data if attrs.keepdims == 1 or ((attrs.axes | length) == 0 and attrs.noop_with_empty_axes == 1) else (ranks.data - (attrs.axes | length) if (attrs.axes | length) > 0 else 0)",
+      "description": "The reduced output tensor."
+    }
+  ],
+  "attributes": { "keepdims": 1, "noop_with_empty_axes": 0, "axes": [] },
+  "attributeDescriptions": {
+    "keepdims": "If 1, retains the reduced dimension with size 1 in the output; if 0, the reduced dimension is removed.",
+    "noop_with_empty_axes": "When 1 and `axes` is empty, acts as an identity (no reduction); when 0 and `axes` is empty, reduces over all axes.",
+    "axes": "Values of the optional ONNX `axes` tensor input, supplied through this request attribute; an empty list follows `noop_with_empty_axes`."
+  },
+  "attributeConstraints": { "keepdims": { "values": [0, 1] }, "noop_with_empty_axes": { "values": [0, 1] } },
+  "typeConstraints": { "T": ["float32", "float16", "int32"] },
+  "args": {
+    "x": { "kind": "tensor", "semantic": "data", "role": "input" },
+    "y": { "kind": "tensor", "semantic": "reduced", "role": "output" }
+  },
+  "derive": {
+    "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
+    "reduceWorkgroupSize": "min(tunables.WORKGROUP_SIZE, deviceWorkgroupCap)",
+    "treeWorkgroupOk": "reduceWorkgroupSize > 0 and pow2ceil(reduceWorkgroupSize) == reduceWorkgroupSize and reduceWorkgroupSize * dtypeBytes(\"float32\") <= device.limits.maxComputeWorkgroupStorageSize",
+    "subgroupWorkgroupFloor": "min(reduceWorkgroupSize, max(1, device.adapterInfo.subgroupMaxSize))",
+    "lastAxisRows": "rows(shapes.data, ranks.data - 1) if ranks.data > 0 else 1",
+    "lastAxisCols": "dim(shapes.data, ranks.data - 1) if ranks.data > 0 else 1",
+    "rowSerialPreferred": "lastAxisRows >= tunables.ROW_SERIAL_MIN_ROWS and lastAxisCols <= tunables.ROW_SERIAL_MAX_COLS",
+    "axis0Rows": "dim(shapes.data, 0) if ranks.data >= 2 else 0",
+    "axis0Cols": "dim(shapes.data, 1) if ranks.data >= 2 else 0",
+    "axis0SplitCount": "min(tunables.AXIS0_MAX_SPLITS, pow2ceil(ceilDiv(axis0Rows, tunables.AXIS0_SPLIT_TARGET_ROWS)))",
+    "axis0SplitScratchBytes": "3 * axis0SplitCount * axis0Cols * dtypeBytes(\"float32\")",
+    "axis0SplitPathFits": "axis0SplitCount <= device.limits.maxComputeWorkgroupsPerDimension and ceilDiv(ceilDiv(axis0Cols, reduceWorkgroupSize), device.limits.maxComputeWorkgroupsPerDimension) <= device.limits.maxComputeWorkgroupsPerDimension and axis0SplitScratchBytes <= device.limits.maxStorageBufferBindingSize and axis0SplitScratchBytes <= device.limits.maxBufferSize",
+    "reduceAxis": "(attrs.axes[0] + ranks.data if attrs.axes[0] < 0 else attrs.axes[0]) if ((attrs.axes | length) == 1 and isUniqueIntList(attrs.axes, 0 - ranks.data, ranks.data, 1)) else ranks.data",
+    "axisSplitDim": "dim(shapes.data, reduceAxis) if ranks.data >= 2 and reduceAxis < ranks.data else 0",
+    "axisSplitInner": "inner(shapes.data, reduceAxis) if ranks.data >= 2 and reduceAxis < ranks.data else 1",
+    "axisSplitOutputs": "numel(shapes.reduced)",
+    "axisSplitCount": "min(tunables.AXIS0_MAX_SPLITS, pow2ceil(ceilDiv(axisSplitDim, tunables.AXIS0_SPLIT_TARGET_ROWS)))",
+    "axisSplitScratchBytes": "3 * axisSplitCount * axisSplitOutputs * 4",
+    "axisSplitPathFits": "axisSplitCount <= device.limits.maxComputeWorkgroupsPerDimension and ceilDiv(ceilDiv(axisSplitOutputs, reduceWorkgroupSize), device.limits.maxComputeWorkgroupsPerDimension) <= device.limits.maxComputeWorkgroupsPerDimension and axisSplitScratchBytes <= device.limits.maxStorageBufferBindingSize and axisSplitScratchBytes <= device.limits.maxBufferSize",
+    "axis0TilePathFits": "treeWorkgroupOk and tunables.AXIS0_TILE_COLS > 0 and tunables.AXIS0_TILE_COLS <= reduceWorkgroupSize and reduceWorkgroupSize % tunables.AXIS0_TILE_COLS == 0",
+    "flatItems": "numel(shapes.data) / tunables.VECTOR_WIDTH if numel(shapes.data) % tunables.VECTOR_WIDTH == 0 else numel(shapes.data)",
+    "flatSplitCount": "max(1, min(tunables.FULL_REDUCE_MAX_SPLITS, ceilDiv(flatItems, reduceWorkgroupSize)))",
+    "flatScratchBytes": "3 * flatSplitCount * dtypeBytes(\"float32\")",
+    "flatPathFits": "treeWorkgroupOk and flatSplitCount <= device.limits.maxComputeWorkgroupsPerDimension and flatScratchBytes <= device.limits.maxStorageBufferBindingSize and flatScratchBytes <= device.limits.maxBufferSize",
+    "flatParallelCovered": "(dtypes.T == \"f32\" or dtypes.T == \"f16\") and f16Ok(dtypes.T) and numel(shapes.reduced) == 1 and numel(shapes.data) >= tunables.FULL_REDUCE_MIN_ELEMENTS and flatPathFits",
+    "contiguousSuffixParallelCovered": "(dtypes.T == \"f32\" or dtypes.T == \"f16\") and f16Ok(dtypes.T) and numel(shapes.reduced) > 0 and numel(shapes.data) % numel(shapes.reduced) == 0 and numel(shapes.data) / numel(shapes.reduced) >= tunables.CONTIGUOUS_SUFFIX_MIN_COLS and ((ranks.data == 3 and hasAxis(attrs.axes, 0, 3) == false and hasAxis(attrs.axes, 1, 3) and hasAxis(attrs.axes, 2, 3) and numel(shapes.reduced) == dim(shapes.data, 0)) or (ranks.data == 4 and hasAxis(attrs.axes, 0, 4) == false and hasAxis(attrs.axes, 1, 4) == false and hasAxis(attrs.axes, 2, 4) and hasAxis(attrs.axes, 3, 4) and numel(shapes.reduced) == dim(shapes.data, 0) * dim(shapes.data, 1)) or (ranks.data == 4 and hasAxis(attrs.axes, 0, 4) == false and hasAxis(attrs.axes, 1, 4) and hasAxis(attrs.axes, 2, 4) and hasAxis(attrs.axes, 3, 4) and numel(shapes.reduced) == dim(shapes.data, 0)))"
+  },
+  "tunables": {
+    "WORKGROUP_SIZE": 256,
+    "VECTOR_WIDTH": 4,
+    "ROW_PARALLEL_MIN_COLS": 64,
+    "SUBGROUP_MIN_COLS": 256,
+    "SUBGROUP_SMALL_ROW_LIMIT": 32768,
+    "AXIS0_SPLIT_MIN_ROWS": 8192,
+    "AXIS0_SPLIT_TARGET_ROWS": 256,
+    "AXIS0_MAX_SPLITS": 128,
+    "AXIS0_TILE_MIN_ROWS": 64,
+    "AXIS0_TILE_MIN_COLS": 16,
+    "AXIS0_TILE_COLS": 16,
+    "AXIS_SPLIT_TILE_COLS": 8,
+    "FULL_REDUCE_MIN_ELEMENTS": 8192,
+    "FULL_REDUCE_MAX_SPLITS": 256,
+    "CONTIGUOUS_SUFFIX_MIN_COLS": 256,
+    "AXES02_WORKGROUP_SIZE": 256,
+    "ROW_SERIAL_MIN_ROWS": 8192,
+    "ROW_SERIAL_MAX_COLS": 1024
+  },
+  "bindingSets": {
+    "axes02": [
+      { "name": "x", "arg": "x", "semantic": "data", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
+      { "name": "y", "arg": "y", "semantic": "reduced", "buffer": { "type": "storage" }, "elementType": "$T" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "d0", "type": "u32", "value": "dim(shapes.data, 0)" },
+            { "name": "d1", "type": "u32", "value": "dim(shapes.data, 1)" },
+            { "name": "d2", "type": "u32", "value": "dim(shapes.data, 2)" },
+            { "name": "outCount", "type": "u32", "value": "numel(shapes.reduced)" }
+          ]
+        }
+      }
+    ],
+    "suffixVec4": [
+      {
+        "name": "x",
+        "arg": "x",
+        "semantic": "data",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      { "name": "y", "arg": "y", "semantic": "reduced", "buffer": { "type": "storage" }, "elementType": "$T" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "numel(shapes.reduced)" },
+            {
+              "name": "chunkCount",
+              "type": "u32",
+              "value": "numel(shapes.data) / numel(shapes.reduced) / tunables.VECTOR_WIDTH"
+            }
+          ]
+        }
+      }
+    ],
+    "suffixScalar": [
+      { "name": "x", "arg": "x", "semantic": "data", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
+      { "name": "y", "arg": "y", "semantic": "reduced", "buffer": { "type": "storage" }, "elementType": "$T" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "numel(shapes.reduced)" },
+            { "name": "cols", "type": "u32", "value": "numel(shapes.data) / numel(shapes.reduced)" }
+          ]
+        }
+      }
+    ],
+    "elementwise": [
+      { "name": "x", "arg": "x", "semantic": "data", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
+      { "name": "y", "arg": "y", "semantic": "reduced", "buffer": { "type": "storage" }, "elementType": "$T" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.reduced)" }] }
+      }
+    ],
+    "lastAxisVec4": [
+      {
+        "name": "x",
+        "arg": "x",
+        "semantic": "data",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      { "name": "y", "arg": "y", "semantic": "reduced", "buffer": { "type": "storage" }, "elementType": "$T" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "rows(shapes.data, ranks.data - 1)" },
+            { "name": "chunkCount", "type": "u32", "value": "dim(shapes.data, ranks.data - 1) / tunables.VECTOR_WIDTH" }
+          ]
+        }
+      }
+    ],
+    "lastAxisScalar": [
+      { "name": "x", "arg": "x", "semantic": "data", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
+      { "name": "y", "arg": "y", "semantic": "reduced", "buffer": { "type": "storage" }, "elementType": "$T" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "rows(shapes.data, ranks.data - 1)" },
+            { "name": "cols", "type": "u32", "value": "dim(shapes.data, ranks.data - 1)" }
+          ]
+        }
+      }
+    ],
+    "lastAxisScalarSubgroup": [
+      { "name": "x", "arg": "x", "semantic": "data", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
+      { "name": "y", "arg": "y", "semantic": "reduced", "buffer": { "type": "storage" }, "elementType": "$T" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "rows(shapes.data, ranks.data - 1)" },
+            { "name": "chunkCount", "type": "u32", "value": "dim(shapes.data, ranks.data - 1)" }
+          ]
+        }
+      }
+    ],
+    "scalar": [
+      { "name": "x", "arg": "x", "semantic": "data", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
+      { "name": "y", "arg": "y", "semantic": "reduced", "buffer": { "type": "storage" }, "elementType": "$T" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "1" },
+            { "name": "cols", "type": "u32", "value": "1" },
+            { "name": "outCount", "type": "u32", "value": "1" }
+          ]
+        }
+      }
+    ],
+    "rank1Axis0": [
+      { "name": "x", "arg": "x", "semantic": "data", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
+      { "name": "y", "arg": "y", "semantic": "reduced", "buffer": { "type": "storage" }, "elementType": "$T" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "dim(shapes.data, 0)" },
+            { "name": "cols", "type": "u32", "value": "1" },
+            { "name": "outCount", "type": "u32", "value": "numel(shapes.reduced)" }
+          ]
+        }
+      }
+    ],
+    "rank2Serial": [
+      { "name": "x", "arg": "x", "semantic": "data", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
+      { "name": "y", "arg": "y", "semantic": "reduced", "buffer": { "type": "storage" }, "elementType": "$T" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "dim(shapes.data, 0)" },
+            { "name": "cols", "type": "u32", "value": "dim(shapes.data, 1)" },
+            { "name": "outCount", "type": "u32", "value": "numel(shapes.reduced)" }
+          ]
+        }
+      }
+    ],
+    "rank2SerialAxis1": [
+      { "name": "x", "arg": "x", "semantic": "data", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
+      { "name": "y", "arg": "y", "semantic": "reduced", "buffer": { "type": "storage" }, "elementType": "$T" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "cols", "type": "u32", "value": "dim(shapes.data, 1)" },
+            { "name": "outCount", "type": "u32", "value": "numel(shapes.reduced)" }
+          ]
+        }
+      }
+    ],
+    "axis0Parallel": [
+      { "name": "x", "arg": "x", "semantic": "data", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
+      { "name": "y", "arg": "y", "semantic": "reduced", "buffer": { "type": "storage" }, "elementType": "$T" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "dim(shapes.data, 0)" },
+            { "name": "cols", "type": "u32", "value": "dim(shapes.data, 1)" }
+          ]
+        }
+      }
+    ],
+    "fullReduceSerial": [
+      { "name": "x", "arg": "x", "semantic": "data", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
+      { "name": "y", "arg": "y", "semantic": "reduced", "buffer": { "type": "storage" }, "elementType": "$T" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "numel(shapes.data)" },
+            { "name": "cols", "type": "u32", "value": "1" },
+            { "name": "outCount", "type": "u32", "value": "numel(shapes.reduced)" }
+          ]
+        }
+      }
+    ],
+    "axisSplitReduce": [
+      { "name": "x", "arg": "x", "semantic": "data", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
+      { "name": "partials", "semantic": "partials", "buffer": { "type": "storage" }, "elementType": "$partialElement" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "axisDim", "type": "u32", "value": "axisSplitDim" },
+            { "name": "inner", "type": "u32", "value": "axisSplitInner" },
+            { "name": "outputs", "type": "u32", "value": "axisSplitOutputs" }
+          ]
+        }
+      }
+    ],
+    "axisSplitCombine": [
+      {
+        "name": "partials",
+        "semantic": "partials",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$partialElement"
+      },
+      { "name": "y", "arg": "y", "semantic": "reduced", "buffer": { "type": "storage" }, "elementType": "$T" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "cols", "type": "u32", "value": "axisSplitOutputs" }] }
+      }
+    ],
+    "axis0SplitReduce": [
+      { "name": "x", "arg": "x", "semantic": "data", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
+      { "name": "partials", "semantic": "partials", "buffer": { "type": "storage" }, "elementType": "$partialElement" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "dim(shapes.data, 0)" },
+            { "name": "cols", "type": "u32", "value": "dim(shapes.data, 1)" }
+          ]
+        }
+      }
+    ],
+    "axis0SplitCombine": [
+      {
+        "name": "partials",
+        "semantic": "partials",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$partialElement"
+      },
+      { "name": "y", "arg": "y", "semantic": "reduced", "buffer": { "type": "storage" }, "elementType": "$T" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "cols", "type": "u32", "value": "dim(shapes.data, 1)" }] }
+      }
+    ],
+    "rankNAxis": [
+      { "name": "x", "arg": "x", "semantic": "data", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
+      { "name": "y", "arg": "y", "semantic": "reduced", "buffer": { "type": "storage" }, "elementType": "$T" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "axisDim", "type": "u32", "value": "axisSplitDim" },
+            { "name": "outCount", "type": "u32", "value": "numel(shapes.reduced)" }
+          ]
+        }
+      }
+    ],
+    "flatPartialF32": [
+      {
+        "name": "x",
+        "arg": "x",
+        "semantic": "data",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$flatScalar"
+      },
+      { "name": "partials", "semantic": "partials", "buffer": { "type": "storage" }, "elementType": "f32" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "flatItems" }] }
+      }
+    ],
+    "flatCombineF32": [
+      { "name": "partials", "semantic": "partials", "buffer": { "type": "read-only-storage" }, "elementType": "f32" },
+      { "name": "y", "arg": "y", "semantic": "reduced", "buffer": { "type": "storage" }, "elementType": "$T" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "cols", "type": "u32", "value": "1" }] }
+      }
+    ],
+    "multiAxis": [
+      { "name": "x", "arg": "x", "semantic": "data", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
+      { "name": "y", "arg": "y", "semantic": "reduced", "buffer": { "type": "storage" }, "elementType": "$T" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [{ "name": "outCount", "type": "u32", "value": "numel(shapes.reduced)" }]
+        }
+      }
+    ]
+  },
+  "variants": [
+    {
+      "id": "contiguous_suffix_subgroup_vec4",
+      "priority": 30,
+      "requires": { "features": ["subgroups"] },
+      "when": ["device.wgslLanguageFeatures.has(\"subgroup_id\")", "not flatParallelCovered", "contiguousSuffixParallelCovered", "(numel(shapes.data) / numel(shapes.reduced)) % tunables.VECTOR_WIDTH == 0"],
+      "constants": {
+        "scalar": "dtypes.T",
+        "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
+        "workgroupSize": "min(reduceWorkgroupSize, max(subgroupWorkgroupFloor, pow2ceil(ceilDiv(numel(shapes.data) / numel(shapes.reduced), tunables.VECTOR_WIDTH))))"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ReduceLogSumExp.ContiguousSuffixSubgroupVec4",
+          "source": {
+            "shader": "reduce-row-subgroup.wgsl.jinja",
+            "inputs": {
+              "op": "\"logsumexp\"",
+              "vec4": true,
+              "castF32": "dtypes.T == \"f16\"",
+              "usesF16": "dtypes.T == \"f16\""
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "suffixVec4",
+          "dispatch": { "workgroups": "numel(shapes.reduced)" }
+        }
+      ]
+    },
+    {
+      "id": "contiguous_suffix_tree_vec4",
+      "priority": 22,
+      "when": ["not flatParallelCovered", "contiguousSuffixParallelCovered", "(numel(shapes.data) / numel(shapes.reduced)) % tunables.VECTOR_WIDTH == 0", "treeWorkgroupOk"],
+      "constants": {
+        "scalar": "dtypes.T",
+        "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
+        "workgroupSize": "min(reduceWorkgroupSize, pow2ceil(ceilDiv(numel(shapes.data) / numel(shapes.reduced), tunables.VECTOR_WIDTH)))"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ReduceLogSumExp.ContiguousSuffixTreeVec4",
+          "source": {
+            "shader": "reduce-row-tree.wgsl.jinja",
+            "inputs": {
+              "op": "\"logsumexp\"",
+              "vec4": true,
+              "castF32": "dtypes.T == \"f16\"",
+              "usesF16": "dtypes.T == \"f16\""
+            }
+          },
+          "bindings": "suffixVec4",
+          "dispatch": { "workgroups": "numel(shapes.reduced)" }
+        }
+      ]
+    },
+    {
+      "id": "contiguous_suffix_tree",
+      "priority": 21,
+      "when": ["not flatParallelCovered", "contiguousSuffixParallelCovered", "treeWorkgroupOk"],
+      "constants": {
+        "workgroupSize": "min(reduceWorkgroupSize, pow2ceil(numel(shapes.data) / numel(shapes.reduced)))",
+        "scalar": "dtypes.T"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ReduceLogSumExp.ContiguousSuffixTree",
+          "source": {
+            "shader": "reduce-row-tree.wgsl.jinja",
+            "inputs": { "op": "\"logsumexp\"", "castF32": "dtypes.T == \"f16\"", "usesF16": "dtypes.T == \"f16\"" }
+          },
+          "bindings": "suffixScalar",
+          "dispatch": { "workgroups": "numel(shapes.reduced)" }
+        }
+      ]
+    },
+    {
+      "id": "multi_axis_rank3",
+      "priority": 8,
+      "when": ["not flatParallelCovered", "not contiguousSuffixParallelCovered", "f16Ok(dtypes.T)", "ranks.data == 3", "(attrs.keepdims == 1 and ranks.reduced == 3 and (dim(shapes.reduced, 0) == 1 if hasAxis(attrs.axes, 0, 3) else dim(shapes.reduced, 0) == dim(shapes.data, 0)) and (dim(shapes.reduced, 1) == 1 if hasAxis(attrs.axes, 1, 3) else dim(shapes.reduced, 1) == dim(shapes.data, 1)) and (dim(shapes.reduced, 2) == 1 if hasAxis(attrs.axes, 2, 3) else dim(shapes.reduced, 2) == dim(shapes.data, 2))) or (attrs.keepdims == 0 and ranks.reduced == 1)"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "ReduceLogSumExp.MultiAxisRank3",
+          "source": {
+            "shader": "reduce-serial-axis.wgsl.jinja",
+            "inputs": {
+              "op": "\"logsumexp\"",
+              "indexing": "\"multiaxis\"",
+              "rank": 3,
+              "reduce": ["hasAxis(attrs.axes, 0, 3)", "hasAxis(attrs.axes, 1, 3)", "hasAxis(attrs.axes, 2, 3)"],
+              "dataShape": "shapes.data",
+              "outputShape": "shapes.reduced",
+              "outputRank": "ranks.reduced",
+              "keepDims": "attrs.keepdims != 0",
+              "intMode": "dtypes.T == \"i32\"",
+              "castF32": "dtypes.T == \"f16\"",
+              "usesF16": "dtypes.T == \"f16\""
+            }
+          },
+          "bindings": "multiAxis",
+          "dispatch": { "threads": "numel(shapes.reduced)", "workgroupSize": "reduceWorkgroupSize" }
+        }
+      ],
+      "constants": { "scalar": "dtypes.T" }
+    },
+    {
+      "id": "multi_axis_rank4",
+      "priority": 8,
+      "when": ["not flatParallelCovered", "not contiguousSuffixParallelCovered", "f16Ok(dtypes.T)", "ranks.data == 4", "attrs.noop_with_empty_axes == 0", "numel(shapes.reduced) == (1 if hasAxis(attrs.axes, 0, 4) else dim(shapes.data, 0)) * (1 if hasAxis(attrs.axes, 1, 4) else dim(shapes.data, 1)) * (1 if hasAxis(attrs.axes, 2, 4) else dim(shapes.data, 2)) * (1 if hasAxis(attrs.axes, 3, 4) else dim(shapes.data, 3))", "((attrs.keepdims == 1 and ranks.reduced == 4) or (attrs.keepdims == 0 and ranks.reduced < 4))"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "ReduceLogSumExp.MultiAxisRank4",
+          "source": {
+            "shader": "reduce-serial-axis.wgsl.jinja",
+            "inputs": {
+              "op": "\"logsumexp\"",
+              "indexing": "\"multiaxis\"",
+              "rank": 4,
+              "reduce": ["hasAxis(attrs.axes, 0, 4)", "hasAxis(attrs.axes, 1, 4)", "hasAxis(attrs.axes, 2, 4)", "hasAxis(attrs.axes, 3, 4)"],
+              "dataShape": "shapes.data",
+              "outputShape": "shapes.reduced",
+              "outputRank": "ranks.reduced",
+              "keepDims": "attrs.keepdims != 0",
+              "intMode": "dtypes.T == \"i32\"",
+              "castF32": "dtypes.T == \"f16\"",
+              "usesF16": "dtypes.T == \"f16\""
+            }
+          },
+          "bindings": "multiAxis",
+          "dispatch": { "threads": "numel(shapes.reduced)", "workgroupSize": "reduceWorkgroupSize" }
+        }
+      ],
+      "constants": { "scalar": "dtypes.T" }
+    },
+    {
+      "id": "int32_rank3_axes02_keepdims",
+      "priority": 30,
+      "when": ["dtypes.T == \"i32\"", "ranks.data == 3", "attrs.keepdims == 1", "hasAxis(attrs.axes, 0, 3)", "hasAxis(attrs.axes, 2, 3)", "hasAxis(attrs.axes, 1, 3) == false", "dim(shapes.data, 0) > 0", "dim(shapes.data, 2) > 0", "ranks.reduced == 3", "dim(shapes.reduced, 0) == 1", "dim(shapes.reduced, 1) == dim(shapes.data, 1)", "dim(shapes.reduced, 2) == 1"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "ReduceLogSumExp.Int32Rank3Axes02Keepdims",
+          "bindings": "axes02",
+          "dispatch": { "threads": "numel(shapes.reduced)", "workgroupSize": "axes02WorkgroupSize" },
+          "source": {
+            "shader": "reduce-i32-axes02.wgsl.jinja",
+            "inputs": { "op": "\"logsumexp\"", "workgroupSize": "axes02WorkgroupSize" }
+          }
+        }
+      ],
+      "derive": {
+        "axes02WorkgroupSize": "min(tunables.AXES02_WORKGROUP_SIZE, device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)"
+      }
+    },
+    {
+      "id": "noop_empty_axes",
+      "priority": 40,
+      "when": ["dtypes.T == \"f32\"", "attrs.noop_with_empty_axes == 1", "(attrs.axes | length) == 0", "sameShape(shapes.data, shapes.reduced)"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "ReduceLogSumExp.NoopEmptyAxes",
+          "source": { "shader": "reduce-noop-empty-axes.wgsl.jinja", "inputs": { "op": "\"identity\"" } },
+          "bindings": "elementwise",
+          "dispatch": { "threads": "numel(shapes.reduced)", "workgroupSize": "reduceWorkgroupSize" }
+        }
+      ]
+    },
+    {
+      "id": "tree_last_axis_vec4",
+      "priority": 23,
+      "demoteWhen": ["rowSerialPreferred"],
+      "when": ["not flatParallelCovered", "(dtypes.T == \"f32\" or dtypes.T == \"f16\")", "f16Ok(dtypes.T)", "ranks.data >= 1", "reduceAxis == ranks.data - 1", "numel(shapes.reduced) == rows(shapes.data, ranks.data - 1)", "attrs.noop_with_empty_axes == 0", "lastAxisCols >= tunables.ROW_PARALLEL_MIN_COLS", "lastAxisCols % tunables.VECTOR_WIDTH == 0", "treeWorkgroupOk"],
+      "constants": {
+        "scalar": "dtypes.T",
+        "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
+        "workgroupSize": "min(reduceWorkgroupSize, pow2ceil(ceilDiv(lastAxisCols, tunables.VECTOR_WIDTH)))"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ReduceLogSumExp.TreeRowVec4",
+          "source": {
+            "shader": "reduce-row-tree.wgsl.jinja",
+            "inputs": {
+              "op": "\"logsumexp\"",
+              "vec4": true,
+              "castF32": "dtypes.T == \"f16\"",
+              "usesF16": "dtypes.T == \"f16\""
+            }
+          },
+          "bindings": "lastAxisVec4",
+          "dispatch": { "workgroups": "lastAxisRows" }
+        }
+      ]
+    },
+    {
+      "id": "rank0_scalar",
+      "priority": 40,
+      "constants": { "axis": 0, "scalar": "dtypes.T" },
+      "when": ["f16Ok(dtypes.T)", "ranks.data == 0", "ranks.reduced == 0"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "ReduceLogSumExp.Rank0Scalar",
+          "source": {
+            "shader": "reduce-serial-axis.wgsl.jinja",
+            "inputs": {
+              "op": "\"logsumexp\"",
+              "indexing": "\"axis2d\"",
+              "intMode": "dtypes.T == \"i32\"",
+              "castF32": "dtypes.T == \"f16\"",
+              "usesF16": "dtypes.T == \"f16\"",
+              "logicalBool": "tensorDtypes.data == \"bool\""
+            }
+          },
+          "bindings": "scalar",
+          "dispatch": { "x": 1 }
+        }
+      ]
+    },
+    {
+      "id": "rank1_axis0",
+      "constants": { "axis": 0, "scalar": "dtypes.T" },
+      "when": ["not flatParallelCovered", "f16Ok(dtypes.T)", "ranks.data == 1", "reduceAxis == 0", "((attrs.keepdims == 0 and ranks.reduced == 0) or (attrs.keepdims == 1 and ranks.reduced == 1 and dim(shapes.reduced, 0) == 1))"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "ReduceLogSumExp.Rank1Axis0",
+          "source": {
+            "shader": "reduce-serial-axis.wgsl.jinja",
+            "inputs": {
+              "op": "\"logsumexp\"",
+              "indexing": "\"axis2d\"",
+              "intMode": "dtypes.T == \"i32\"",
+              "castF32": "dtypes.T == \"f16\"",
+              "usesF16": "dtypes.T == \"f16\"",
+              "logicalBool": "tensorDtypes.data == \"bool\""
+            }
+          },
+          "bindings": "rank1Axis0",
+          "dispatch": { "threads": "numel(shapes.reduced)", "workgroupSize": "reduceWorkgroupSize" }
+        }
+      ]
+    },
+    {
+      "id": "axis1_parallel",
+      "priority": 20,
+      "demoteWhen": ["rowSerialPreferred"],
+      "when": ["not flatParallelCovered", "(dtypes.T == \"f32\" or dtypes.T == \"f16\")", "f16Ok(dtypes.T)", "ranks.data >= 2", "reduceAxis == ranks.data - 1", "numel(shapes.reduced) == rows(shapes.data, ranks.data - 1)", "lastAxisCols >= tunables.ROW_PARALLEL_MIN_COLS", "treeWorkgroupOk"],
+      "constants": { "workgroupSize": "min(reduceWorkgroupSize, pow2ceil(dim(shapes.data, ranks.data - 1)))" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ReduceLogSumExp.Axis1Parallel",
+          "source": {
+            "shader": "reduce-row-tree.wgsl.jinja",
+            "inputs": { "op": "\"logsumexp\"", "castF32": "dtypes.T == \"f16\"", "usesF16": "dtypes.T == \"f16\"" }
+          },
+          "bindings": "lastAxisScalar",
+          "dispatch": { "workgroups": "rows(shapes.data, ranks.data - 1)" }
+        }
+      ]
+    },
+    {
+      "id": "axis_split",
+      "priority": 24,
+      "when": ["not flatParallelCovered", "(dtypes.T == \"f32\" or dtypes.T == \"f16\")", "f16Ok(dtypes.T)", "attrs.noop_with_empty_axes == 0", "ranks.data >= 2", "reduceAxis < ranks.data - 1", "not (ranks.data == 2 and reduceAxis == 0)", "axisSplitDim >= tunables.AXIS0_SPLIT_MIN_ROWS", "axisSplitOutputs >= 1", "axisSplitOutputs <= 4096", "axisSplitOutputs == rows(shapes.data, reduceAxis)", "axisSplitPathFits"],
+      "derive": { "splitCount": "axisSplitCount" },
+      "constants": { "partialElement": "\"f32\"", "workgroupSize": "reduceWorkgroupSize", "split": "splitCount" },
+      "intermediates": [{ "id": "partials", "dtype": "float32", "shape": "[3 * splitCount * axisSplitOutputs]" }],
+      "passes": [
+        {
+          "id": "split_reduce",
+          "name": "ReduceLogSumExp.AxisSplitReduce",
+          "source": {
+            "shader": "reduce-axis-split-reduce.wgsl.jinja",
+            "inputs": {
+              "op": "\"logsumexp\"",
+              "workgroupSize": "reduceWorkgroupSize",
+              "split": "splitCount",
+              "castF32": "dtypes.T == \"f16\"",
+              "usesF16": "dtypes.T == \"f16\""
+            }
+          },
+          "bindings": "axisSplitReduce",
+          "dispatch": { "threads": "axisSplitOutputs", "workgroupSize": "reduceWorkgroupSize", "y": "splitCount" }
+        },
+        {
+          "id": "combine",
+          "name": "ReduceLogSumExp.AxisSplitCombine",
+          "source": {
+            "shader": "reduce-axis0-splitk-combine.wgsl.jinja",
+            "inputs": {
+              "op": "\"logsumexp\"",
+              "workgroupSize": "reduceWorkgroupSize",
+              "split": "splitCount",
+              "outputF16": "dtypes.T == \"f16\""
+            }
+          },
+          "bindings": "axisSplitCombine",
+          "dispatch": { "threads": "axisSplitOutputs", "workgroupSize": "reduceWorkgroupSize" }
+        }
+      ]
+    },
+    {
+      "id": "axis_split_tiled_narrow",
+      "priority": 25,
+      "when": ["not flatParallelCovered", "(dtypes.T == \"f32\" or dtypes.T == \"f16\")", "f16Ok(dtypes.T)", "attrs.noop_with_empty_axes == 0", "ranks.data >= 2", "reduceAxis < ranks.data - 1", "axisSplitDim >= tunables.AXIS0_SPLIT_MIN_ROWS", "axisSplitOutputs >= 1", "axisSplitOutputs <= 2 * tunables.AXIS_SPLIT_TILE_COLS", "reduceWorkgroupSize % tunables.AXIS_SPLIT_TILE_COLS == 0", "axisSplitOutputs == rows(shapes.data, reduceAxis)", "axisSplitPathFits"],
+      "derive": { "splitCount": "axisSplitCount" },
+      "constants": {
+        "partialElement": "\"f32\"",
+        "scalar": "dtypes.T",
+        "workgroupSize": "reduceWorkgroupSize",
+        "split": "splitCount",
+        "tileCols": "tunables.AXIS_SPLIT_TILE_COLS"
+      },
+      "intermediates": [{ "id": "partials", "dtype": "float32", "shape": "[3 * splitCount * axisSplitOutputs]" }],
+      "passes": [
+        {
+          "id": "split_reduce",
+          "name": "ReduceLogSumExp.AxisSplitTiledReduce",
+          "source": {
+            "shader": "reduce-axis0-tilecols.wgsl.jinja",
+            "inputs": {
+              "op": "\"logsumexp\"",
+              "workgroupSize": "reduceWorkgroupSize",
+              "split": "splitCount",
+              "tileCols": "tunables.AXIS_SPLIT_TILE_COLS",
+              "castF32": "dtypes.T == \"f16\"",
+              "usesF16": "dtypes.T == \"f16\""
+            }
+          },
+          "bindings": "axisSplitReduce",
+          "dispatch": { "workgroups": "ceilDiv((axisSplitOutputs), (constants.tileCols))", "y": "splitCount" }
+        },
+        {
+          "id": "combine",
+          "name": "ReduceLogSumExp.AxisSplitCombine",
+          "source": {
+            "shader": "reduce-axis0-splitk-combine.wgsl.jinja",
+            "inputs": {
+              "op": "\"logsumexp\"",
+              "workgroupSize": "reduceWorkgroupSize",
+              "split": "splitCount",
+              "outputF16": "dtypes.T == \"f16\""
+            }
+          },
+          "bindings": "axisSplitCombine",
+          "dispatch": { "threads": "axisSplitOutputs", "workgroupSize": "reduceWorkgroupSize" }
+        }
+      ]
+    },
+    {
+      "id": "axis0_splitk",
+      "priority": 22,
+      "when": ["not flatParallelCovered", "(dtypes.T == \"f32\" or dtypes.T == \"f16\")", "f16Ok(dtypes.T)", "ranks.data == 2", "reduceAxis == 0", "axis0Rows >= tunables.AXIS0_SPLIT_MIN_ROWS", "dim(shapes.data, 1) > 0", "((attrs.keepdims == 0 and ranks.reduced == 1 and dim(shapes.reduced, 0) == dim(shapes.data, 1)) or (attrs.keepdims == 1 and ranks.reduced == 2 and dim(shapes.reduced, 0) == 1 and dim(shapes.reduced, 1) == dim(shapes.data, 1)))", "axis0SplitPathFits"],
+      "derive": { "splitCount": "axis0SplitCount" },
+      "constants": { "partialElement": "\"f32\"", "workgroupSize": "reduceWorkgroupSize", "split": "splitCount" },
+      "intermediates": [{ "id": "partials", "dtype": "float32", "shape": "[3 * splitCount * dim(shapes.data, 1)]" }],
+      "passes": [
+        {
+          "id": "split_reduce",
+          "name": "ReduceLogSumExp.Axis0SplitKReduce",
+          "source": {
+            "shader": "reduce-axis0-splitk-reduce.wgsl.jinja",
+            "inputs": {
+              "op": "\"logsumexp\"",
+              "workgroupSize": "reduceWorkgroupSize",
+              "split": "splitCount",
+              "castF32": "dtypes.T == \"f16\"",
+              "usesF16": "dtypes.T == \"f16\""
+            }
+          },
+          "bindings": "axis0SplitReduce",
+          "dispatch": { "threads": "dim(shapes.data, 1)", "workgroupSize": "reduceWorkgroupSize", "y": "splitCount" }
+        },
+        {
+          "id": "combine",
+          "name": "ReduceLogSumExp.Axis0SplitKCombine",
+          "source": {
+            "shader": "reduce-axis0-splitk-combine.wgsl.jinja",
+            "inputs": {
+              "op": "\"logsumexp\"",
+              "workgroupSize": "reduceWorkgroupSize",
+              "split": "splitCount",
+              "outputF16": "dtypes.T == \"f16\""
+            }
+          },
+          "bindings": "axis0SplitCombine",
+          "dispatch": { "threads": "dim(shapes.data, 1)", "workgroupSize": "reduceWorkgroupSize" }
+        }
+      ]
+    },
+    {
+      "id": "axis0_tilecols",
+      "priority": 20,
+      "when": ["not flatParallelCovered", "(dtypes.T == \"f32\" or dtypes.T == \"f16\")", "f16Ok(dtypes.T)", "ranks.data == 2", "reduceAxis == 0", "axis0Rows >= tunables.AXIS0_TILE_MIN_ROWS", "axis0Cols >= tunables.AXIS0_TILE_MIN_COLS", "((attrs.keepdims == 0 and ranks.reduced == 1 and dim(shapes.reduced, 0) == dim(shapes.data, 1)) or (attrs.keepdims == 1 and ranks.reduced == 2 and dim(shapes.reduced, 0) == 1 and dim(shapes.reduced, 1) == dim(shapes.data, 1)))", "axis0TilePathFits"],
+      "constants": { "workgroupSize": "reduceWorkgroupSize", "tileCols": "tunables.AXIS0_TILE_COLS" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ReduceLogSumExp.Axis0TileCols",
+          "source": {
+            "shader": "reduce-axis0-tilecols.wgsl.jinja",
+            "inputs": { "op": "\"logsumexp\"", "castF32": "dtypes.T == \"f16\"", "usesF16": "dtypes.T == \"f16\"" }
+          },
+          "bindings": "axis0Parallel",
+          "dispatch": { "workgroups": "ceilDiv((dim(shapes.data, 1)), (constants.tileCols))" }
+        }
+      ]
+    },
+    {
+      "id": "all_axes_flat",
+      "priority": 31,
+      "constants": {
+        "scalar": "dtypes.T",
+        "workgroupSize": "reduceWorkgroupSize",
+        "flatScalar": "\"vec4<\" ~ dtypes.T ~ \">\" if numel(shapes.data) % tunables.VECTOR_WIDTH == 0 else dtypes.T",
+        "split": "flatSplitCount"
+      },
+      "when": ["flatParallelCovered"],
+      "intermediates": [{ "id": "partials", "dtype": "float32", "shape": "[3 * flatSplitCount]" }],
+      "passes": [
+        {
+          "id": "flat_partial",
+          "name": "ReduceLogSumExp.AllAxesFlatPartial",
+          "source": {
+            "shader": "reduce-flat-partial-logsumexp.wgsl.jinja",
+            "inputs": {
+              "vec4": "numel(shapes.data) % tunables.VECTOR_WIDTH == 0",
+              "castF32": "dtypes.T == \"f16\"",
+              "usesF16": "dtypes.T == \"f16\""
+            }
+          },
+          "bindings": "flatPartialF32",
+          "dispatch": { "x": "flatSplitCount" }
+        },
+        {
+          "id": "combine",
+          "name": "ReduceLogSumExp.AllAxesFlatCombine",
+          "source": {
+            "shader": "reduce-flat-combine-logsumexp.wgsl.jinja",
+            "inputs": { "outputF16": "dtypes.T == \"f16\"" }
+          },
+          "bindings": "flatCombineF32",
+          "dispatch": { "x": 1 }
+        }
+      ]
+    },
+    {
+      "id": "rankn_single_axis_generic",
+      "priority": 12,
+      "supersededBy": ["axis_split_tiled_narrow", "axis_split", "subgroup_last_axis_vec4", "subgroup_last_axis", "tree_last_axis_vec4"],
+      "when": ["not flatParallelCovered", "f16Ok(dtypes.T)", "ranks.data >= 3", "attrs.noop_with_empty_axes == 0", "reduceAxis < ranks.data", "numel(shapes.reduced) == rows(shapes.data, reduceAxis)", "((attrs.keepdims == 0 and ranks.reduced == ranks.data - 1) or (attrs.keepdims == 1 and ranks.reduced == ranks.data and dim(shapes.reduced, reduceAxis) == 1))"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "ReduceLogSumExp.RankNSingleAxisGeneric",
+          "source": {
+            "shader": "reduce-serial-axis.wgsl.jinja",
+            "inputs": {
+              "op": "\"logsumexp\"",
+              "indexing": "\"rankn\"",
+              "rank": "ranks.data",
+              "axis": "reduceAxis",
+              "dataShape": "shapes.data",
+              "outputShape": "shapes.reduced",
+              "outputRank": "ranks.reduced",
+              "keepDims": "attrs.keepdims != 0",
+              "intMode": "dtypes.T == \"i32\"",
+              "castF32": "dtypes.T == \"f16\"",
+              "usesF16": "dtypes.T == \"f16\""
+            }
+          },
+          "bindings": "rankNAxis",
+          "dispatch": { "threads": "numel(shapes.reduced)", "workgroupSize": "reduceWorkgroupSize" }
+        }
+      ],
+      "constants": { "scalar": "dtypes.T" }
+    },
+    {
+      "id": "subgroup_last_axis_vec4",
+      "priority": 25,
+      "requires": { "features": ["subgroups"] },
+      "when": ["device.wgslLanguageFeatures.has(\"subgroup_id\")", "not flatParallelCovered", "(dtypes.T == \"f32\" or dtypes.T == \"f16\")", "f16Ok(dtypes.T)", "ranks.data >= 1", "reduceAxis == ranks.data - 1", "numel(shapes.reduced) == rows(shapes.data, ranks.data - 1)", "dim(shapes.data, ranks.data - 1) >= 4", "dim(shapes.data, ranks.data - 1) % tunables.VECTOR_WIDTH == 0", "(lastAxisCols >= tunables.SUBGROUP_MIN_COLS or lastAxisRows < tunables.SUBGROUP_SMALL_ROW_LIMIT)", "not rowSerialPreferred"],
+      "constants": {
+        "scalar": "dtypes.T",
+        "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
+        "workgroupSize": "min(reduceWorkgroupSize, max(subgroupWorkgroupFloor, pow2ceil(ceilDiv(lastAxisCols, tunables.VECTOR_WIDTH))))"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ReduceLogSumExp.SubgroupRowVec4",
+          "source": {
+            "shader": "reduce-row-subgroup.wgsl.jinja",
+            "inputs": {
+              "op": "\"logsumexp\"",
+              "vec4": true,
+              "castF32": "dtypes.T == \"f16\"",
+              "usesF16": "dtypes.T == \"f16\""
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "lastAxisVec4",
+          "dispatch": { "workgroups": "rows(shapes.data, ranks.data - 1)" }
+        }
+      ]
+    },
+    {
+      "id": "subgroup_last_axis",
+      "priority": 24,
+      "requires": { "features": ["subgroups"] },
+      "when": ["device.wgslLanguageFeatures.has(\"subgroup_id\")", "not flatParallelCovered", "(dtypes.T == \"f32\" or dtypes.T == \"f16\")", "f16Ok(dtypes.T)", "ranks.data >= 1", "reduceAxis == ranks.data - 1", "numel(shapes.reduced) == rows(shapes.data, ranks.data - 1)", "dim(shapes.data, ranks.data - 1) > 0", "dim(shapes.data, ranks.data - 1) % tunables.VECTOR_WIDTH != 0", "(lastAxisCols >= tunables.SUBGROUP_MIN_COLS or lastAxisRows < tunables.SUBGROUP_SMALL_ROW_LIMIT)", "not rowSerialPreferred"],
+      "constants": {
+        "scalar": "dtypes.T",
+        "workgroupSize": "min(reduceWorkgroupSize, max(subgroupWorkgroupFloor, pow2ceil(lastAxisCols)))"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ReduceLogSumExp.SubgroupRow",
+          "source": {
+            "shader": "reduce-row-subgroup.wgsl.jinja",
+            "inputs": {
+              "op": "\"logsumexp\"",
+              "vec4": false,
+              "castF32": "dtypes.T == \"f16\"",
+              "usesF16": "dtypes.T == \"f16\""
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "lastAxisScalarSubgroup",
+          "dispatch": { "workgroups": "rows(shapes.data, ranks.data - 1)" }
+        }
+      ]
+    },
+    {
+      "id": "axis0",
+      "priority": 0,
+      "supersededBy": ["axis_split_tiled_narrow", "axis0_splitk", "axis0_tilecols"],
+      "constants": { "axis": 0, "scalar": "dtypes.T" },
+      "when": ["not flatParallelCovered", "f16Ok(dtypes.T)", "ranks.data == 2", "reduceAxis == 0", "((attrs.keepdims == 0 and ranks.reduced == 1 and dim(shapes.reduced, 0) == dim(shapes.data, 1)) or (attrs.keepdims == 1 and ranks.reduced == 2 and dim(shapes.reduced, 0) == 1 and dim(shapes.reduced, 1) == dim(shapes.data, 1)))"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "axis0",
+          "source": {
+            "shader": "reduce-serial-axis.wgsl.jinja",
+            "inputs": {
+              "op": "\"logsumexp\"",
+              "indexing": "\"axis2d\"",
+              "intMode": "dtypes.T == \"i32\"",
+              "castF32": "dtypes.T == \"f16\"",
+              "usesF16": "dtypes.T == \"f16\""
+            }
+          },
+          "bindings": "rank2Serial",
+          "constants": { "axis": 0 },
+          "dispatch": { "threads": "numel(shapes.reduced)", "workgroupSize": "reduceWorkgroupSize" }
+        }
+      ]
+    },
+    {
+      "id": "axis1",
+      "priority": 0,
+      "constants": { "axis": 1, "scalar": "dtypes.T" },
+      "when": ["not flatParallelCovered", "f16Ok(dtypes.T)", "ranks.data == 2", "reduceAxis == 1", "((attrs.keepdims == 0 and ranks.reduced == 1 and dim(shapes.reduced, 0) == dim(shapes.data, 0)) or (attrs.keepdims == 1 and ranks.reduced == 2 and dim(shapes.reduced, 0) == dim(shapes.data, 0) and dim(shapes.reduced, 1) == 1))"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "axis1",
+          "source": {
+            "shader": "reduce-serial-axis.wgsl.jinja",
+            "inputs": {
+              "op": "\"logsumexp\"",
+              "indexing": "\"axis2d\"",
+              "intMode": "dtypes.T == \"i32\"",
+              "castF32": "dtypes.T == \"f16\"",
+              "usesF16": "dtypes.T == \"f16\""
+            }
+          },
+          "bindings": "rank2SerialAxis1",
+          "constants": { "axis": 1 },
+          "dispatch": { "threads": "numel(shapes.reduced)", "workgroupSize": "reduceWorkgroupSize" }
+        }
+      ]
+    },
+    {
+      "id": "all_axes_keepdims",
+      "priority": 30,
+      "constants": { "axis": 0, "scalar": "dtypes.T" },
+      "when": ["not flatParallelCovered", "f16Ok(dtypes.T)", "ranks.data >= 3", "attrs.keepdims == 1", "ranks.reduced == ranks.data", "numel(shapes.reduced) == 1"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "ReduceLogSumExp.Rank3AllAxesKeepdims",
+          "source": {
+            "shader": "reduce-serial-axis.wgsl.jinja",
+            "inputs": {
+              "op": "\"logsumexp\"",
+              "indexing": "\"axis2d\"",
+              "intMode": "dtypes.T == \"i32\"",
+              "castF32": "dtypes.T == \"f16\"",
+              "usesF16": "dtypes.T == \"f16\""
+            }
+          },
+          "bindings": "fullReduceSerial",
+          "dispatch": { "threads": "numel(shapes.reduced)", "workgroupSize": "reduceWorkgroupSize" }
+        }
+      ]
+    },
+    {
+      "id": "all_axes_no_keepdims",
+      "priority": 30,
+      "constants": { "axis": 0, "scalar": "dtypes.T" },
+      "when": ["not flatParallelCovered", "f16Ok(dtypes.T)", "ranks.data >= 3", "attrs.keepdims == 0", "attrs.noop_with_empty_axes == 0", "ranks.reduced == 0"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "ReduceLogSumExp.Rank3AllAxesNoKeepdims",
+          "source": {
+            "shader": "reduce-serial-axis.wgsl.jinja",
+            "inputs": {
+              "op": "\"logsumexp\"",
+              "indexing": "\"axis2d\"",
+              "intMode": "dtypes.T == \"i32\"",
+              "castF32": "dtypes.T == \"f16\"",
+              "usesF16": "dtypes.T == \"f16\""
+            }
+          },
+          "bindings": "fullReduceSerial",
+          "dispatch": { "threads": "numel(shapes.reduced)", "workgroupSize": "reduceWorkgroupSize" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,28 @@

+{
+  "name": "ai.onnx.ReduceLogSumExp",
+  "id": "_ai_onnx_reducelogsumexp_webgpu_85769e2",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "osiIimng7rf8c4+lWBGf/4BUtBJ1pMrNA2I08w6ygf0=",
+      "manifest.json": "rbJKRnqa38U/J/QPQmwV9MWpWxBCCQQoVVrRzy1fRI0=",
+      "reduce-axis-split-reduce.wgsl.jinja": "CKODGEG8mXIkTNFzjDJZwfExy1KzYxG50C3LaUiYd3U=",
+      "reduce-axis0-splitk-combine.wgsl.jinja": "P3U2gCak59UZSyakvyuWCDhW27IPcilPg+yE1BPY5nM=",
+      "reduce-axis0-splitk-reduce.wgsl.jinja": "cowJ5TK8xFHBwnCPYUYEBScEIKXFiCN4mlu0B2lWaw4=",
+      "reduce-axis0-tilecols.wgsl.jinja": "CNwWK3NVE27T0IJo4+HoDfLFDGMuJ0PVDJRk5Jmwq7o=",
+      "reduce-flat-combine-logsumexp.wgsl.jinja": "qrI8ma9jUPm+p+oWJejet/DvaNNTgjw4AoUBKIgoaPQ=",
+      "reduce-flat-partial-logsumexp.wgsl.jinja": "dl9CTTHHIAjnI0QvdFpNkv4IFWD62Cf0Q3jpgsAfpPo=",
+      "reduce-i32-axes02.wgsl.jinja": "qSNKLHjGG9CFO6DRIlBsQ7XETdKrQbbnn3O7NJPxvY0=",
+      "reduce-noop-empty-axes.wgsl.jinja": "IO2MEzyBGJr3Z8AVP7ELJXqfhwcG/vtBfa0wCdkbuwo=",
+      "reduce-row-subgroup.wgsl.jinja": "GqRcdIa6KA7BPnGGO9ApvQPizeinFAG1XDOSLsutZ6Y=",
+      "reduce-row-tree.wgsl.jinja": "yX9bu0RBuTGybmUUcK4bZ9Gf0CyqDsMThC9jIXoXxNU=",
+      "reduce-serial-axis.wgsl.jinja": "BGP2SlwMidLS+xcUX/FBC9lAwCi5gLM6DWOZYiqxYQo=",
+      "test.json": "nvmxvijOcgGOSHCGvqrJ6/iAdaTeHTSgnc4Ghf7IGCE="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.ReduceLogSumExp" }
+}

build/webgpu/reduce-axis-split-reduce.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,71 @@

+// Pass 1 of a split-K reduction for an (outer, axis, inner) flattening. Each
+// (output, segment) invocation reduces one axis slice at stride inner and writes
+// partials[segment * outputs + output]. Adjacent output threads read adjacent
+// inner-axis elements. The combine pass folds the segments and finalizes the
+// selected reduction.
+//
+// logsumexp writes three partial planes per output and segment: the segment
+// maximum, the sum of exp(x - maximum), and a packed NaN marker.
+{% set castF32 = source.castF32 is defined and source.castF32 %}
+{% set xa = "f32(" if castF32 else "" %}
+{% set ax = ")" if castF32 else "" %}
+{% if source.usesF16 is defined and source.usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+/* Exact max/min reduction identities. WGSL rejects infinity during constant
+ * evaluation, so f32 identities are constructed at runtime from their
+ * IEEE-754 bit patterns. */
+const WG: u32 = {{ workgroupSize }}u;
+const SPLIT: u32 = {{ split }}u;
+const F32_MIN: f32 = -3.4028234663852886e38;
+fn is_nan_f32(value: f32) -> bool {
+  let bits = bitcast<u32>(value);
+  return (bits & 0x7f800000u) == 0x7f800000u && (bits & 0x007fffffu) != 0u;
+}
+@compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>,
+        @builtin(workgroup_id) wg: vec3<u32>,
+        @builtin(num_workgroups) nwg: vec3<u32>) {
+  // 2D-folded output index: wg.z carries the high bits past the
+  // per-dimension dispatch limit on the x dimension.
+  let output_index = (wg.x + wg.z * nwg.x) * WG + (gid.x % WG);
+  let seg = wg.y;
+  if (output_index >= params.outputs) { return; }
+  let outer_index = output_index / params.inner;
+  let inner_index = output_index % params.inner;
+  let input_base = outer_index * params.axisDim * params.inner + inner_index;
+  // Even split of [0, axisDim) across SPLIT segments (last may be shorter).
+  let chunk = (params.axisDim + SPLIT - 1u) / SPLIT;
+  let a0 = seg * chunk;
+  var a1 = a0 + chunk;
+  if (a1 > params.axisDim) { a1 = params.axisDim; }
+  var local_max = F32_MIN;
+  var local_nan_count = 0.0;
+  var local_nan_value = 0.0;
+  for (var axis_index = a0; axis_index < a1; axis_index = axis_index + 1u) {
+    let value = {{ xa }}x[input_base + axis_index * params.inner]{{ ax }};
+    if (is_nan_f32(value)) {
+      local_nan_count = local_nan_count + 1.0;
+      local_nan_value = value;
+    } else {
+      local_max = max(local_max, value);
+    }
+  }
+  var acc = 0.0;
+  if (local_nan_count == 0.0) {
+    for (var axis_index = a0; axis_index < a1; axis_index = axis_index + 1u) {
+      acc = acc + exp({{ xa }}x[input_base + axis_index * params.inner]{{ ax }} - local_max);
+    }
+  }
+  partials[seg * params.outputs + output_index] = local_max;
+  partials[(SPLIT + seg) * params.outputs + output_index] = acc;
+  partials[(2u * SPLIT + seg) * params.outputs + output_index] = select(0.0, local_nan_value, local_nan_count > 0.0);
+}

build/webgpu/reduce-axis0-splitk-combine.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,54 @@

+// Pass 2 of the split-K column-wise reduction. One invocation per output column
+// folds the segment partials and applies the selected reduction's final step.
+// Segments are folded in ascending order for deterministic results. This order
+// differs from the single-pass reduction but remains within the f32 tolerance.
+{% set yv = "f16(" if source.outputF16 else "" %}
+{% set vy = ")" if source.outputF16 else "" %}
+{% if source.outputF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+/* Exact max/min reduction identities. WGSL rejects infinity during constant
+ * evaluation, so f32 identities are constructed at runtime from their
+ * IEEE-754 bit patterns. */
+const WG: u32 = {{ workgroupSize }}u;
+const SPLIT: u32 = {{ split }}u;
+const F32_MIN: f32 = -3.4028234663852886e38;
+const F32_MAX: f32 = 3.4028234663852886e38;
+fn is_nan_f32(value: f32) -> bool {
+  let bits = bitcast<u32>(value);
+  return (bits & 0x7f800000u) == 0x7f800000u && (bits & 0x007fffffu) != 0u;
+}
+@compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>,
+        @builtin(num_workgroups) nwg: vec3<u32>) {
+  let stride = nwg.x * WG;
+  let start = (gid.y * nwg.x * WG) + gid.x;
+  for (var col = start; col < params.cols; col = col + stride) {
+    // Merge SPLIT (segMax, segSumExp) pairs stably; carry NaN / +Inf markers.
+    var nan_value = 0.0;
+    var has_nan = false;
+    var global_max = F32_MIN;
+    for (var seg = 0u; seg < SPLIT; seg = seg + 1u) {
+      let nv = partials[(2u * SPLIT + seg) * params.cols + col];
+      if (nv != 0.0 || is_nan_f32(nv)) {
+        has_nan = true;
+        nan_value = nv;
+      }
+      global_max = max(global_max, partials[seg * params.cols + col]);
+    }
+    var sum = 0.0;
+    for (var seg = 0u; seg < SPLIT; seg = seg + 1u) {
+      let seg_max = partials[seg * params.cols + col];
+      let seg_sum = partials[(SPLIT + seg) * params.cols + col];
+      sum = sum + seg_sum * exp(seg_max - global_max);
+    }
+    let has_positive_inf = global_max > F32_MAX;
+    let finite_or_inf = select(global_max + log(sum), global_max, has_positive_inf);
+    y[col] = {{ yv }}select(finite_or_inf, nan_value, has_nan){{ vy }};
+  }
+}

build/webgpu/reduce-axis0-splitk-reduce.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,68 @@

+// Pass 1 of a split-K column-wise reduction. Splitting rows across workgroup
+// segments increases residency for tall matrices. Each (column, segment)
+// invocation reduces one row slice and writes partials[segment * columns +
+// column]. Adjacent column threads keep row reads coalesced.
+//
+// logsumexp writes the segment maximum and sum of exp(x - maximum) as separate
+// partial planes; the combine pass merges them stably and handles NaN and +Inf.
+{% set castF32 = source.castF32 is defined and source.castF32 %}
+{% set xa = "f32(" if castF32 else "" %}
+{% set ax = ")" if castF32 else "" %}
+{% if source.usesF16 is defined and source.usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+/* Exact max/min reduction identities. WGSL rejects infinity during constant
+ * evaluation, so f32 identities are constructed at runtime from their
+ * IEEE-754 bit patterns. */
+const WG: u32 = {{ workgroupSize }}u;
+const SPLIT: u32 = {{ split }}u;
+const F32_MIN: f32 = -3.4028234663852886e38;
+fn is_nan_f32(value: f32) -> bool {
+  let bits = bitcast<u32>(value);
+  return (bits & 0x7f800000u) == 0x7f800000u && (bits & 0x007fffffu) != 0u;
+}
+@compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>,
+        @builtin(workgroup_id) wg: vec3<u32>,
+        @builtin(num_workgroups) nwg: vec3<u32>) {
+  // 2D-folded column index: wg.z carries the high bits past the maxComputeWorkgroupsPerDimension
+  // workgroup-per-dimension dispatch limit on the x dimension.
+  let col = (wg.x + wg.z * nwg.x) * WG + (gid.x % WG);
+  let seg = wg.y;
+  if (col >= params.cols) { return; }
+  // Even split of [0, rows) across SPLIT segments (last segment may be shorter).
+  let chunk = (params.rows + SPLIT - 1u) / SPLIT;
+  let r0 = seg * chunk;
+  var r1 = r0 + chunk;
+  if (r1 > params.rows) { r1 = params.rows; }
+  var local_max = F32_MIN;
+  var local_nan_count = 0.0;
+  var local_nan_value = 0.0;
+  for (var row = r0; row < r1; row = row + 1u) {
+    let value = {{ xa }}x[row * params.cols + col]{{ ax }};
+    if (is_nan_f32(value)) {
+      local_nan_count = local_nan_count + 1.0;
+      local_nan_value = value;
+    } else {
+      local_max = max(local_max, value);
+    }
+  }
+  var acc = 0.0;
+  if (local_nan_count == 0.0) {
+    for (var row = r0; row < r1; row = row + 1u) {
+      acc = acc + exp({{ xa }}x[row * params.cols + col]{{ ax }} - local_max);
+    }
+  }
+  // Three partial planes: seg max, seg sumexp(x - max), and a packed nan
+  // marker (count in the low slot, the nan bit-pattern smuggled as f32).
+  partials[seg * params.cols + col] = local_max;
+  partials[(SPLIT + seg) * params.cols + col] = acc;
+  partials[(2u * SPLIT + seg) * params.cols + col] = select(0.0, local_nan_value, local_nan_count > 0.0);
+}

build/webgpu/reduce-axis0-tilecols.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,174 @@

+// Tiled column-wise reduction. Each workgroup owns TILE_COLS columns;
+// ROW_LANES threads stride the rows of one column, then lane 0 folds their
+// partials. In split mode, workgroup y selects an axis segment and finalization
+// is deferred to the combine pass. Cooperative row lanes preserve occupancy
+// when the flattened output has only a few elements.
+{% set splitMode = source.split is defined %}
+{% if splitMode %}
+{% set rowBegin = "row_begin + row_lane" %}
+{% set rowEnd = "row_end" %}
+{% set elem = "x[input_base + row * params.inner]" %}
+{% else %}
+{% set rowBegin = "row_lane" %}
+{% set rowEnd = "params.rows" %}
+{% set elem = "x[inputBase + row * params.cols + col]" %}
+{% endif %}
+{% set castF32 = source.castF32 is defined and source.castF32 %}
+{% set intMode = source.intMode is defined and source.intMode %}
+{% set scalar = "f32" if castF32 else scalar %}
+{% if castF32 %}
+{% set elem = "f32(" ~ elem ~ ")" %}
+{% endif %}
+{% set yv = "f16(" if castF32 else "" %}
+{% set vy = ")" if castF32 else "" %}
+{% if source.usesF16 is defined and source.usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+/* Exact max/min reduction identities. WGSL rejects infinity during constant
+ * evaluation, so f32 identities are constructed at runtime from their
+ * IEEE-754 bit patterns. */
+{% if not splitMode and not intMode and (source.op == "logsum" or source.op == "logsumexp") %}
+fn negative_infinity() -> f32 {
+  var bits = 0xff800000u;
+  return bitcast<f32>(bits);
+}
+{% endif %}
+const WG: u32 = {{ workgroupSize }}u;
+const TILE_COLS: u32 = {{ tileCols }}u;
+const ROW_LANES: u32 = WG / TILE_COLS;
+{% if splitMode %}
+const SPLIT: u32 = {{ source.split }}u;
+{% endif %}
+const F32_MIN: f32 = -3.4028234663852886e38;
+const F32_MAX: f32 = 3.4028234663852886e38;
+var<workgroup> partial: array<{{ scalar if (source.op == "max" or source.op == "min" or intMode) else "f32" }}, WG>;
+fn is_nan_f32(value: f32) -> bool {
+  let bits = bitcast<u32>(value);
+  return (bits & 0x7f800000u) == 0x7f800000u && (bits & 0x007fffffu) != 0u;
+}
+@compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid: vec3<u32>{% if not splitMode %},
+        @builtin(num_workgroups) nwg: vec3<u32>{% endif %}) {
+  let tid = lid.x;
+  let col_lane = tid % TILE_COLS;
+  let row_lane = tid / TILE_COLS;
+{% if splitMode %}
+  // Narrow outputs: wg.x covers every column tile, wg.y is the axis segment.
+  let col = wg.x * TILE_COLS + col_lane;
+  let outputIndex = col;
+  let in_bounds = col < params.outputs;
+  let seg = wg.y;
+  let outer_index = col / params.inner;
+  let inner_index = col % params.inner;
+  let input_base = outer_index * params.axisDim * params.inner + inner_index;
+  // Even split of [0, axisDim) across SPLIT segments (last may be shorter).
+  let chunk = (params.axisDim + SPLIT - 1u) / SPLIT;
+  let row_begin = seg * chunk;
+  let row_end = min(row_begin + chunk, params.axisDim);
+{% else %}
+  // 2D-folded tile index: wg.y carries the high bits past the dispatch limit.
+  // The batched form reuses this same coalesced axis-0 reduction for a middle
+  // axis by assigning consecutive tiles to each outer slice.
+  let tile = wg.x + wg.y * nwg.x;
+  let col = tile * TILE_COLS + col_lane;
+  let inputBase = 0u;
+  let outputIndex = col;
+  let in_bounds = col < params.cols;
+{% endif %}
+{% if not splitMode %}
+  if (params.rows == 0u) {
+    if (row_lane == 0u && in_bounds) {
+      y[outputIndex] = {{ yv }}negative_infinity(){{ vy }};
+    }
+    return;
+  }
+{% endif %}
+  var local_max = F32_MIN;
+  var local_nan_count = 0.0;
+  var local_nan_value = 0.0;
+  if (in_bounds) {
+    for (var row = {{ rowBegin }}; row < {{ rowEnd }}; row = row + ROW_LANES) {
+      let value = {{ elem }};
+      if (is_nan_f32(value)) {
+        local_nan_count = local_nan_count + 1.0;
+        local_nan_value = value;
+      } else {
+        local_max = max(local_max, value);
+      }
+    }
+  }
+  partial[tid] = local_max;
+  workgroupBarrier();
+  if (row_lane == 0u && in_bounds) {
+    var max_lanes = partial[col_lane];
+    for (var lane = 1u; lane < ROW_LANES; lane = lane + 1u) {
+      max_lanes = max(max_lanes, partial[lane * TILE_COLS + col_lane]);
+    }
+    partial[col_lane] = max_lanes;
+  }
+  workgroupBarrier();
+  let max_value = partial[col_lane];
+  workgroupBarrier();
+  partial[tid] = local_nan_count;
+  workgroupBarrier();
+  if (row_lane == 0u && in_bounds) {
+    var nan_lanes = partial[col_lane];
+    for (var lane = 1u; lane < ROW_LANES; lane = lane + 1u) {
+      nan_lanes = nan_lanes + partial[lane * TILE_COLS + col_lane];
+    }
+    partial[col_lane] = nan_lanes;
+  }
+  workgroupBarrier();
+  let nan_count = partial[col_lane];
+  workgroupBarrier();
+  partial[tid] = local_nan_value;
+  workgroupBarrier();
+  if (row_lane == 0u && in_bounds) {
+    var nan_value_lanes = partial[col_lane];
+    for (var lane = 1u; lane < ROW_LANES; lane = lane + 1u) {
+      nan_value_lanes = nan_value_lanes + partial[lane * TILE_COLS + col_lane];
+    }
+    partial[col_lane] = nan_value_lanes;
+  }
+  workgroupBarrier();
+  let nan_value = partial[col_lane];
+  let has_nan = nan_count > 0.0;
+  let has_positive_inf = max_value > F32_MAX;
+  workgroupBarrier();
+  var acc = 0.0;
+  if (in_bounds && !has_nan && !has_positive_inf) {
+    for (var row = {{ rowBegin }}; row < {{ rowEnd }}; row = row + ROW_LANES) {
+      acc = acc + exp({{ elem }} - max_value);
+    }
+  }
+  partial[tid] = acc;
+  workgroupBarrier();
+  if (row_lane == 0u && in_bounds) {
+    var sum = partial[col_lane];
+    for (var lane = 1u; lane < ROW_LANES; lane = lane + 1u) {
+      sum = sum + partial[lane * TILE_COLS + col_lane];
+    }
+{% if splitMode %}
+    partials[seg * params.outputs + outputIndex] = max_value;
+    partials[(SPLIT + seg) * params.outputs + outputIndex] = sum;
+    partials[(2u * SPLIT + seg) * params.outputs + outputIndex] = select(0.0, nan_value, has_nan);
+{% else %}
+    let finite_or_inf = select(max_value + log(sum), max_value, has_positive_inf);
+    y[outputIndex] = {{ yv }}select(finite_or_inf, nan_value, has_nan){{ vy }};
+{% endif %}
+  }
+}

build/webgpu/reduce-flat-combine-logsumexp.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,88 @@

+// Parallel scalar combine for the flat LogSumExp full reduction. The partial
+// pass writes SPLIT stable (segment maximum, shifted exponential sum, NaN
+// marker) triples. Unlike the generic axis-0 combine, which assigns one lane
+// to each output column, this scalar-only path has a full workgroup fold the
+// partial planes. That avoids leaving one lane to execute 2*SPLIT
+// transcendental merges when the output has exactly one element.
+{% set yv = "f16(" if source.outputF16 else "" %}
+{% set vy = ")" if source.outputF16 else "" %}
+{% if source.outputF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const WG: u32 = {{ workgroupSize }}u;
+const SPLIT: u32 = {{ split }}u;
+const F32_MIN: f32 = -3.4028234663852886e38;
+const F32_MAX: f32 = 3.4028234663852886e38;
+var<workgroup> wmax: array<f32, WG>;
+var<workgroup> wsegmax: array<f32, WG>;
+var<workgroup> wsum: array<f32, WG>;
+var<workgroup> wnan: array<f32, WG>;
+fn is_nan_f32(value: f32) -> bool {
+  let bits = bitcast<u32>(value);
+  return (bits & 0x7f800000u) == 0x7f800000u && (bits & 0x007fffffu) != 0u;
+}
+@compute @workgroup_size(WG)
+fn main(@builtin(local_invocation_id) lid: vec3<u32>) {
+  let tid = lid.x;
+  var seg_max = F32_MIN;
+  var seg_sum = 0.0;
+  var nan_value = 0.0;
+  if (tid < SPLIT) {
+    // params.cols is one for this scalar-only path. Keep it in the addressing
+    // expression so the binding contract stays identical to the shared axis-0
+    // combine used by the other flat reductions.
+    let col = 0u;
+    seg_max = partials[tid * params.cols + col];
+    seg_sum = partials[(SPLIT + tid) * params.cols + col];
+    nan_value = partials[(2u * SPLIT + tid) * params.cols + col];
+  }
+  wmax[tid] = seg_max;
+  wsegmax[tid] = seg_max;
+  wsum[tid] = seg_sum;
+  wnan[tid] = nan_value;
+  workgroupBarrier();
+  var stride = WG / 2u;
+  loop {
+    if (stride == 0u) { break; }
+    if (tid < stride) {
+      wmax[tid] = max(wmax[tid], wmax[tid + stride]);
+      let other_nan = wnan[tid + stride];
+      if (other_nan != 0.0 || is_nan_f32(other_nan)) {
+        wnan[tid] = other_nan;
+      }
+    }
+    stride = stride / 2u;
+    workgroupBarrier();
+  }
+  let global_max = wmax[0];
+  if (tid < SPLIT) {
+    wsum[tid] = wsum[tid] * exp(wsegmax[tid] - global_max);
+  } else {
+    wsum[tid] = 0.0;
+  }
+  workgroupBarrier();
+  stride = WG / 2u;
+  loop {
+    if (stride == 0u) { break; }
+    if (tid < stride) {
+      wsum[tid] = wsum[tid] + wsum[tid + stride];
+    }
+    stride = stride / 2u;
+    workgroupBarrier();
+  }
+  if (tid == 0u) {
+    let has_positive_inf = global_max > F32_MAX;
+    let finite_or_inf = select(global_max + log(wsum[0]), global_max, has_positive_inf);
+    let has_nan = wnan[0] != 0.0 || is_nan_f32(wnan[0]);
+    y[0] = {{ yv }}select(finite_or_inf, wnan[0], has_nan){{ vy }};
+  }
+}

build/webgpu/reduce-flat-partial-logsumexp.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,101 @@

+// Pass 1 of a flat LogSumExp reduction to one scalar. SPLIT workgroups traverse
+// disjoint partial segments. Each first finds its segment maximum while
+// tracking NaN, then rereads the segment to sum exp(x - segmentMaximum).
+// Workgroups write three partial planes: segment maximum, shifted exponential
+// sum, and a NaN marker. The combine pass merges them stably and emits
+// globalMaximum + log(sum), preserving NaN and positive infinity.
+{% set vec4 = source.vec4 | default(true) %}
+{% set castF32 = source.castF32 is defined and source.castF32 %}
+{% set xa = "f32(" if castF32 else "" %}
+{% set ax = ")" if castF32 else "" %}
+{% if source.usesF16 is defined and source.usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const WG: u32 = {{ workgroupSize }}u;
+const SPLIT: u32 = {{ split }}u;
+const F32_MIN: f32 = -3.4028234663852886e38;
+var<workgroup> wmax: array<f32, WG>;
+var<workgroup> wnan: array<f32, WG>;
+var<workgroup> wsum: array<f32, WG>;
+fn is_nan_f32(value: f32) -> bool {
+  let bits = bitcast<u32>(value);
+  return (bits & 0x7f800000u) == 0x7f800000u && (bits & 0x007fffffu) != 0u;
+}
+@compute @workgroup_size(WG)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>,
+        @builtin(local_invocation_id) lid: vec3<u32>,
+        @builtin(workgroup_id) wg: vec3<u32>,
+        @builtin(num_workgroups) nwg: vec3<u32>) {
+  let tid = lid.x;
+  let gstride = nwg.x * WG;
+  // Pass A: per-lane max over its grid-stride slice of the flat view,
+  // tracking NaN separately (NaN never participates in max).
+  var local_max = F32_MIN;
+  var nan_value = 0.0;
+  var has_nan = false;
+  for (var i = gid.x; i < params.count; i = i + gstride) {
+{% if vec4 %}
+    let v = x[i];
+    for (var c = 0u; c < 4u; c = c + 1u) {
+      let e = {{ xa }}v[c]{{ ax }};
+      if (is_nan_f32(e)) { has_nan = true; nan_value = e; }
+      else { local_max = max(local_max, e); }
+    }
+{% else %}
+    let e = {{ xa }}x[i]{{ ax }};
+    if (is_nan_f32(e)) { has_nan = true; nan_value = e; }
+    else { local_max = max(local_max, e); }
+{% endif %}
+  }
+  wmax[tid] = local_max;
+  wnan[tid] = select(0.0, nan_value, has_nan);
+  workgroupBarrier();
+  var stride: u32 = WG / 2u;
+  loop {
+    if (stride == 0u) { break; }
+    if (tid < stride) {
+      wmax[tid] = max(wmax[tid], wmax[tid + stride]);
+      let other = wnan[tid + stride];
+      if (other != 0.0 || is_nan_f32(other)) { wnan[tid] = other; }
+    }
+    stride = stride / 2u;
+    workgroupBarrier();
+  }
+  let seg_max = wmax[0];
+  let seg_nan = wnan[0];
+  // Pass B: sum exp(x - segMax) over the same slice. (When the segment holds a
+  // NaN the combine discards this sum, so its value is irrelevant then.)
+  var local_sum = 0.0;
+  for (var i = gid.x; i < params.count; i = i + gstride) {
+{% if vec4 %}
+    let v = x[i];
+    local_sum = local_sum
+      + exp({{ xa }}v.x{{ ax }} - seg_max) + exp({{ xa }}v.y{{ ax }} - seg_max)
+      + exp({{ xa }}v.z{{ ax }} - seg_max) + exp({{ xa }}v.w{{ ax }} - seg_max);
+{% else %}
+    local_sum = local_sum + exp({{ xa }}x[i]{{ ax }} - seg_max);
+{% endif %}
+  }
+  wsum[tid] = local_sum;
+  workgroupBarrier();
+  stride = WG / 2u;
+  loop {
+    if (stride == 0u) { break; }
+    if (tid < stride) { wsum[tid] = wsum[tid] + wsum[tid + stride]; }
+    stride = stride / 2u;
+    workgroupBarrier();
+  }
+  if (tid == 0u) {
+    partials[wg.x] = seg_max;
+    partials[SPLIT + wg.x] = wsum[0];
+    partials[2u * SPLIT + wg.x] = seg_nan;
+  }
+}

build/webgpu/reduce-i32-axes02.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,33 @@

+// Rank-3 i32 reduction over axes {0, 2}, leaving the middle dimension.
+{{ env.wgsl.resourceDeclarations }}
+const WG: u32 = {{ source.workgroupSize }}u;
+@compute @workgroup_size(WG)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>,
+        @builtin(num_workgroups) nwg: vec3<u32>) {
+  let outputIndex = gid.x + gid.y * nwg.x * WG;
+  let d0 = params.d0;
+  let d1 = params.d1;
+  let d2 = params.d2;
+  let outputCount = params.outCount;
+  if (outputIndex >= outputCount) {
+    return;
+  }
+  var maxValue = f32(x[outputIndex * d2]);
+  for (var i = 0u; i < d0; i = i + 1u) {
+    for (var k = 0u; k < d2; k = k + 1u) {
+      maxValue = max(maxValue, f32(x[(i * d1 + outputIndex) * d2 + k]));
+    }
+  }
+  var sum = 0.0;
+  for (var i = 0u; i < d0; i = i + 1u) {
+    for (var k = 0u; k < d2; k = k + 1u) {
+      let value = f32(x[(i * d1 + outputIndex) * d2 + k]);
+      sum = sum + exp(value - maxValue);
+    }
+  }
+  y[outputIndex] = i32(maxValue + log(sum));
+}

build/webgpu/reduce-noop-empty-axes.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,13 @@

+{{ env.wgsl.resourceDeclarations }}
+@compute @workgroup_size({{ reduceWorkgroupSize }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
+  let i = gid.x + gid.y * nwg.x * {{ reduceWorkgroupSize }}u;
+  if (i >= params.count) {
+    return;
+  }
+  let v = x[i];
+  y[i] = v;
+}

build/webgpu/reduce-row-subgroup.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,110 @@

+// Subgroup row reduction for a contiguous last axis. One workgroup owns each
+// output row. Threads reduce strided chunks, subgroup leaders deposit their
+// partials in workgroup memory, and thread 0 folds those slots and finalizes.
+//
+// Accumulators are f32 except for integer max/min. Mean, l2, and logsum apply
+// division, square root, or logarithm only at finalization. logsumexp uses
+// max-subtraction, bit-exact NaN detection, NaN propagation, and a +Inf
+// short-circuit. Runtime bitcasts provide max/min infinity identities because
+// WGSL constant evaluation rejects infinite constants. f16 storage is widened
+// before accumulation and narrowed only for the final store.
+{% set castF32 = source.castF32 is defined and source.castF32 %}
+{% set scalar = "f32" if castF32 else scalar %}
+{% set xv = ("vec4<f32>(" if source.vec4 else "f32(") if castF32 else "" %}
+{% set vx = ")" if castF32 else "" %}
+{% set yv = "f16(" if castF32 else "" %}
+{% set vy = ")" if castF32 else "" %}
+enable subgroups;
+{% if source.usesF16 is defined and source.usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+/* Exact max/min reduction identities. WGSL rejects infinity during constant
+ * evaluation, so f32 identities are constructed at runtime from their
+ * IEEE-754 bit patterns. */
+const WG: u32 = {{ workgroupSize }}u;
+const F32_MIN: f32 = -3.4028234663852886e38;
+const F32_MAX: f32 = 3.4028234663852886e38;
+fn is_nan_f32(value: f32) -> bool {
+  let bits = bitcast<u32>(value);
+  return (bits & 0x7f800000u) == 0x7f800000u && (bits & 0x007fffffu) != 0u;
+}
+var<workgroup> wgPartial: array<{{ scalar }}, WG>;
+{% macro emit_reduce(name, collective, combine) %}
+fn {{ name }}(value: {{ scalar }}, sgLid: u32, sgId: u32, numSg: u32) -> {{ scalar }} {
+  let sgVal = {{ collective }}(value);
+  if (numSg == 1u) {
+    return sgVal;
+  }
+  if (sgLid == 0u) {
+    wgPartial[sgId] = sgVal;
+  }
+  workgroupBarrier();
+  var total = wgPartial[0];
+  for (var i = 1u; i < numSg; i = i + 1u) {
+    {{ combine }}
+  }
+  workgroupBarrier();
+  return total;
+}
+{%- endmacro %}{{ emit_reduce("reduce_row_add", "subgroupAdd", "total = total + wgPartial[i];") }}
+{{ emit_reduce("reduce_row_max", "subgroupMax", "total = max(total, wgPartial[i]);") }}
+@compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(workgroup_id) wg: vec3<u32>,
+        @builtin(num_workgroups) nwg: vec3<u32>,
+        @builtin(local_invocation_id) lid: vec3<u32>,
+        @builtin(subgroup_invocation_id) sgLid: u32,
+        @builtin(subgroup_id) sgId: u32,
+        @builtin(num_subgroups) numSg: u32) {
+  let row = wg.x + wg.y * nwg.x;
+  if (row >= params.rows) {
+    return;
+  }
+  let tid = lid.x;
+  let base = row * params.chunkCount;  var localMax = F32_MIN;
+  var localNan = 0.0;
+  var localNanValue = 0.0;
+  for (var c = tid; c < params.chunkCount; c = c + WG) {
+    let v = {{ xv }}x[base + c]{{ vx }};
+{%- if source.vec4 %}
+{% for comp in ["x", "y", "z", "w"] %}
+    if (is_nan_f32(v.{{ comp }})) {
+      localNan = 1.0;
+      localNanValue = v.{{ comp }};
+    } else {
+      localMax = max(localMax, v.{{ comp }});
+    }
+{%- endfor %}
+{% else %}
+    if (is_nan_f32(v)) {
+      localNan = 1.0;
+      localNanValue = v;
+    } else {
+      localMax = max(localMax, v);
+    }
+{%- endif %}
+  }
+  let rowMax = reduce_row_max(localMax, sgLid, sgId, numSg);
+  let nanCount = reduce_row_add(localNan, sgLid, sgId, numSg);
+  let nanValue = reduce_row_add(localNanValue, sgLid, sgId, numSg);
+  let hasPositiveInf = rowMax > F32_MAX;
+  let hasNan = nanCount > 0.0;
+  var acc = 0.0;
+  for (var c = tid; c < params.chunkCount; c = c + WG) {
+    let v = {{ xv }}x[base + c]{{ vx }};
+{%- if source.vec4 %}
+    let e = select(exp(v - vec4<f32>(rowMax)), vec4<f32>(0.0), hasPositiveInf || hasNan);
+    acc = acc + ((e.x + e.y) + (e.z + e.w));
+{%- else %}
+    acc = acc + select(exp(v - rowMax), 0.0, hasPositiveInf || hasNan);
+{%- endif %}
+  }
+  let sum = reduce_row_add(acc, sgLid, sgId, numSg);
+  if (tid == 0u) {
+    let finiteOrInf = select(rowMax + log(sum), rowMax, hasPositiveInf);
+    y[row] = {{ yv }}select(finiteOrInf, nanValue, hasNan){{ vy }};
+  }}

build/webgpu/reduce-row-tree.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,161 @@

+// Portable one-workgroup-per-row reduction for the Reduce value family.
+// Threads stride a contiguous row, accumulate locally, and fold their values
+// through a shared-memory tree without relying on subgroups.
+//
+// Max, min, and product use the native f32/i32/u32 value type. Additive
+// transforms accumulate float inputs in f32 but retain integer accumulation
+// for sum, L1, and sum-of-squares. Mean, L2, log-sum, and log-sum-exp are
+// f32-only because their finalizers divide, take a square root, or take a log.
+// f16 storage widens through f32 for both accumulation and the shared tree,
+// then narrows only at the final store.
+{% set isVec4 = source.vec4 is defined and source.vec4 %}
+{% set rowIsEmpty = "params.chunkCount == 0u" if isVec4 else "params.cols == 0u" %}
+{% set castF32 = source.castF32 is defined and source.castF32 %}
+{% set scalar = "f32" if castF32 else scalar %}
+{% set xv = ("vec4<f32>(" if source.vec4 else "f32(") if castF32 else "" %}
+{% set vx = ")" if castF32 else "" %}
+{% set yv = "f16(" if castF32 else "" %}
+{% set vy = ")" if castF32 else "" %}
+{% if source.usesF16 is defined and source.usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+/* Exact max/min reduction identities. WGSL rejects infinity during constant
+ * evaluation, so f32 identities are constructed at runtime from their
+ * IEEE-754 bit patterns. */
+{% if scalar != "i32" and scalar != "u32" and (source.op == "logsum" or source.op == "logsumexp") %}
+fn negative_infinity() -> f32 {
+  var bits = 0xff800000u;
+  return bitcast<f32>(bits);
+}
+{% endif %}
+const WG: u32 = {{ workgroupSize }}u;
+const F32_MIN: f32 = -3.4028234663852886e38;
+const F32_MAX: f32 = 3.4028234663852886e38;
+var<workgroup> partial: array<f32, WG>;
+{% macro wgsl_tree_reduce_f32(name, mode, buffer="partial", wg="WG", trailingBarrier=true) %}
+fn {{ name }}(value: f32, tid: u32) -> f32 {
+  {{ buffer }}[tid] = value;
+  workgroupBarrier();
+  // Ceil-halving keeps every lane when the workgroup size is not a power of
+  // two. For even n this matches the power-of-two tree order; for odd n, lanes
+  // [0, n-half) fold the upper tail while the middle lane carries forward.
+  var n: u32 = {{ wg }};
+  loop {
+    let half = (n + 1u) / 2u;
+    if (tid < n - half) {
+{% if mode == "max" %}
+      {{ buffer }}[tid] = max({{ buffer }}[tid], {{ buffer }}[tid + half]);
+{% else %}
+      {{ buffer }}[tid] = {{ buffer }}[tid] + {{ buffer }}[tid + half];
+{% endif %}
+    }
+    workgroupBarrier();
+    n = half;
+    if (n == 1u) {
+      break;
+    }
+  }
+  // The default trailing barrier makes this helper safe for back-to-back calls: every lane reads
+  // slot 0 here, so the next call's first store must not run until all lanes have read it.
+  // `trailingBarrier=false` is safe only when the buffer is never written again before kernel exit.
+  let reduced = {{ buffer }}[0];
+{% if trailingBarrier %}
+  workgroupBarrier();
+{% endif %}
+  return reduced;
+}
+{% endmacro %}
+{{ wgsl_tree_reduce_f32("reduce_sum", "add", "partial", "WG") }}
+{{ wgsl_tree_reduce_f32("reduce_max", "max", "partial", "WG") }}
+fn is_nan_f32(value: f32) -> bool {
+  let bits = bitcast<u32>(value);
+  return (bits & 0x7f800000u) == 0x7f800000u
+      && (bits & 0x007fffffu) != 0u;
+}
+@compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(workgroup_id) wg: vec3<u32>,
+        @builtin(num_workgroups) nwg: vec3<u32>,
+        @builtin(local_invocation_id) lid: vec3<u32>) {
+  let row = wg.x + wg.y * nwg.x;
+  if (row >= params.rows) {
+    return;
+  }
+  let tid = lid.x;
+{% if source.vec4 %}
+  let base = row * params.chunkCount;
+{% else %}
+  let base = row * params.cols;
+{% endif %}
+  if ({{ rowIsEmpty }}) {
+    if (tid == 0u) {
+      y[row] = {{ yv }}negative_infinity(){{ vy }};
+    }
+    return;
+  }
+  var localMax = F32_MIN;
+  var localNan = 0.0;
+  var localNanValue = 0.0;
+{% if source.vec4 %}
+  for (var col = tid; col < params.chunkCount; col = col + WG) {
+    let value = {{ xv }}x[base + col]{{ vx }};
+{% for component in ["x", "y", "z", "w"] %}
+    if (is_nan_f32(value.{{ component }})) {
+      localNan = 1.0;
+      localNanValue = value.{{ component }};
+    } else {
+      localMax = max(localMax, value.{{ component }});
+    }
+{% endfor %}
+  }
+{% else %}
+  for (var col = tid; col < params.cols; col = col + WG) {
+    let value = {{ xv }}x[base + col]{{ vx }};
+    if (is_nan_f32(value)) {
+      localNan = 1.0;
+      localNanValue = value;
+    } else {
+      localMax = max(localMax, value);
+    }
+  }
+{% endif %}
+  let rowMax = reduce_max(localMax, tid);
+  workgroupBarrier();
+  let nanCount = reduce_sum(localNan, tid);
+  workgroupBarrier();
+  let nanValue = reduce_sum(localNanValue, tid);
+  workgroupBarrier();
+  let hasPositiveInf = rowMax > F32_MAX;
+  let hasNan = nanCount > 0.0;
+  var acc = 0.0;
+{% if source.vec4 %}
+  for (var col = tid; col < params.chunkCount; col = col + WG) {
+    let value = {{ xv }}x[base + col]{{ vx }};
+    let exponentials = select(exp(value - vec4<f32>(rowMax)), vec4<f32>(0.0),
+                              hasPositiveInf || hasNan);
+    acc = acc + (exponentials.x + exponentials.y)
+              + (exponentials.z + exponentials.w);
+  }
+{% else %}
+  for (var col = tid; col < params.cols; col = col + WG) {
+    acc = acc + select(exp({{ xv }}x[base + col]{{ vx }} - rowMax), 0.0,
+                       hasPositiveInf || hasNan);
+  }
+{% endif %}
+  let sum = reduce_sum(acc, tid);
+  if (tid == 0u) {
+    let finiteOrInf = select(rowMax + log(sum), rowMax, hasPositiveInf);
+    y[row] = {{ yv }}select(finiteOrInf, nanValue, hasNan){{ vy }};
+  }
+}

build/webgpu/reduce-serial-axis.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,196 @@

+// Serial one-thread-per-output reduction for the no-feature tier. f16 storage
+// is widened before every accumulation and narrowed only for the final store.
+{% set castF32 = source.castF32 is defined and source.castF32 %}
+{% set intMode = source.intMode is defined and source.intMode %}
+{% set yv = "f16(" if castF32 else "" %}
+{% set vy = ")" if castF32 else "" %}
+{% if source.usesF16 is defined and source.usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+/* Exact max/min reduction identities. WGSL rejects infinity during constant
+ * evaluation, so f32 identities are constructed at runtime from their
+ * IEEE-754 bit patterns. */
+{% if not intMode and (source.op == "logsum" or source.op == "logsumexp") %}
+fn negative_infinity() -> f32 {
+  var bits = 0xff800000u;
+  return bitcast<f32>(bits);
+}
+{% endif %}
+const F32_MIN: f32 = -3.4028234663852886e38;
+{% if not intMode %}
+const F32_MAX: f32 = 3.4028234663852886e38;
+fn is_nan_f32(value: f32) -> bool {
+  let bits = bitcast<u32>(value);
+  return (bits & 0x7f800000u) == 0x7f800000u && (bits & 0x007fffffu) != 0u;
+}
+{% endif %}
+{% if source.indexing == "rankn" %}
+fn input_offset(out_index: u32, reduce_index: u32) -> u32 {
+  var rem = out_index;
+{% for axis in range(source.outputRank) %}
+{% set out_stride = namespace(value=1) %}
+{% for j in range(axis + 1, source.outputRank) %}
+{% set out_stride.value = out_stride.value * source.outputShape[j] %}
+{% endfor %}
+{% set safe_out_stride = 1 if out_stride.value == 0 else out_stride.value %}
+{% if not source.keepDims or axis != source.axis %}
+  let out_coord{{ axis }} = rem / {{ safe_out_stride }}u;
+{% endif %}
+  rem = rem % {{ safe_out_stride }}u;
+{% endfor %}
+{% for axis in range(source.rank) %}
+{% if axis == source.axis %}
+  let coord{{ axis }} = reduce_index;
+{% elif source.keepDims %}
+  let coord{{ axis }} = out_coord{{ axis }};
+{% elif axis < source.axis %}
+  let coord{{ axis }} = out_coord{{ axis }};
+{% else %}
+  let coord{{ axis }} = out_coord{{ axis - 1 }};
+{% endif %}
+{% endfor %}
+{% set src = namespace(value="coord0") %}
+{% for axis in range(1, source.rank) %}
+{% set src.value = "(" ~ src.value ~ " * " ~ source.dataShape[axis] ~ "u + coord" ~ axis ~ ")" %}
+{% endfor %}
+  return {{ src.value }};
+}
+{% endif %}
+{% if source.indexing == "multiaxis" %}
+{% set hasReducedAxis = namespace(value=false) %}
+{% for a in range(source.rank) %}{% if source.reduce[a] %}{% set hasReducedAxis.value = true %}{% endif %}{% endfor %}
+// One thread per output element walks the Cartesian product of the reduced axes,
+// linearized as reduce_linear. Specialized shapes make every input offset a sum
+// of coordinate-times-constant terms.
+fn input_offset(out_index: u32{% if hasReducedAxis.value %}, reduce_linear: u32{% endif %}) -> u32 {
+  var rem = out_index;
+{% for oaxis in range(source.outputRank) %}
+{% set ostride = namespace(value=1) %}
+{% for j in range(oaxis + 1, source.outputRank) %}
+{% set ostride.value = ostride.value * source.outputShape[j] %}
+{% endfor %}
+{% set osafe = 1 if ostride.value == 0 else ostride.value %}
+{% if not source.keepDims or not source.reduce[oaxis] %}
+  let out_coord{{ oaxis }} = rem / {{ osafe }}u;
+{% endif %}
+  rem = rem % {{ osafe }}u;
+{% endfor %}
+{% if hasReducedAxis.value %}
+  var rrem = reduce_linear;
+{% endif %}
+{% for a in range(source.rank) if source.reduce[a] %}
+{% set rstride = namespace(value=1) %}
+{% for b in range(a + 1, source.rank) if source.reduce[b] %}
+{% set rstride.value = rstride.value * source.dataShape[b] %}
+{% endfor %}
+{% set rsafe = 1 if rstride.value == 0 else rstride.value %}
+  let red_coord{{ a }} = rrem / {{ rsafe }}u;
+  rrem = rrem % {{ rsafe }}u;
+{% endfor %}
+{% set oc = namespace(i=0) %}
+{% for a in range(source.rank) %}
+{% if source.reduce[a] %}
+  let coord{{ a }} = red_coord{{ a }};
+{% elif source.keepDims %}
+  let coord{{ a }} = out_coord{{ a }};
+{% else %}
+  let coord{{ a }} = out_coord{{ oc.i }};
+{% set oc.i = oc.i + 1 %}
+{% endif %}
+{% endfor %}
+{% set src = namespace(value="coord0") %}
+{% for a in range(1, source.rank) %}
+{% set src.value = "(" ~ src.value ~ " * " ~ source.dataShape[a] ~ "u + coord" ~ a ~ ")" %}
+{% endfor %}
+  return {{ src.value }};
+}
+{% endif %}
+{% if source.indexing == "multiaxis" %}
+{% set mcount = namespace(value=1) %}
+{% for a in range(source.rank) if source.reduce[a] %}
+{% set mcount.value = mcount.value * source.dataShape[a] %}
+{% endfor %}
+{% set count = mcount.value ~ "u" %}
+{% if hasReducedAxis.value %}
+{% set at = "x[input_offset(i, r)]" %}
+{% else %}
+{% set at = "x[input_offset(i)]" %}
+{% endif %}
+{% elif source.indexing == "rankn" %}
+{% set count = "params.axisDim" %}
+{% set at = "x[input_offset(i, r)]" %}
+{% elif axis == 0 %}
+{% set count = "params.rows" %}
+{% set at = "x[r * params.cols + i]" %}
+{% else %}
+{% set count = "params.cols" %}
+{% set at = "x[i * params.cols + r]" %}
+{% endif %}
+{% if castF32 %}
+{% set at = "f32(" ~ at ~ ")" %}
+{% endif %}
+@compute @workgroup_size({{ reduceWorkgroupSize }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
+  let i = gid.x + gid.y * nwg.x * {{ reduceWorkgroupSize }}u;
+  if (i >= params.outCount) {
+    return;
+  }
+{% if intMode %}
+  // Integer logsumexp widens each element for exp/log, then truncates the result
+  // back to the integer output type.
+  // Integers are never NaN, so the f32 NaN-propagation path is unnecessary here.
+  if ({{ count }} == 0u) {
+    y[i] = {{ scalar }}(F32_MIN);
+    return;
+  }
+  var m = F32_MIN;
+  for (var r = 0u; r < {{ count }}; r = r + 1u) {
+    m = max(m, f32({{ at }}));
+  }
+  var acc = 0.0;
+  for (var r = 0u; r < {{ count }}; r = r + 1u) {
+    acc = acc + exp(f32({{ at }}) - m);
+  }
+  y[i] = {{ scalar }}(m + log(acc));
+{% else %}
+  if ({{ count }} == 0u) {
+    y[i] = {{ yv }}negative_infinity(){{ vy }};
+    return;
+  }
+  var m = F32_MIN;
+  var has_nan = false;
+  var nan_value = 0.0;
+  for (var r = 0u; r < {{ count }}; r = r + 1u) {
+    let value = {{ at }};
+    if (is_nan_f32(value)) {
+      has_nan = true;
+      nan_value = value;
+    } else {
+      m = max(m, value);
+    }
+  }
+  if (has_nan) {
+    y[i] = {{ yv }}nan_value{{ vy }};
+    return;
+  }
+  if (m > F32_MAX) {
+    y[i] = {{ yv }}m{{ vy }};
+    return;
+  }
+  var acc = 0.0;
+  for (var r = 0u; r < {{ count }}; r = r + 1u) {
+    acc = acc + exp({{ at }} - m);
+  }
+  y[i] = {{ yv }}m + log(acc){{ vy }};
+{% endif %}
+}

build/webgpu/test.json ADDED Viewed

The diff for this file is too large to render. See raw diff