sync 2e7068faf55e

Browse files

Files changed (16) hide show

README.md +88 -0
build/webgpu/bench.json +108 -0
build/webgpu/datamove-flat-copy-x4-tail.wgsl.jinja +31 -0
build/webgpu/manifest.json +1074 -0
build/webgpu/metadata.json +28 -0
build/webgpu/reduce-axis-split-reduce.wgsl.jinja +63 -0
build/webgpu/reduce-axis0-splitk-combine.wgsl.jinja +48 -0
build/webgpu/reduce-axis0-splitk-reduce.wgsl.jinja +58 -0
build/webgpu/reduce-axis0-tilecols.wgsl.jinja +105 -0
build/webgpu/reduce-flat-partial.wgsl.jinja +75 -0
build/webgpu/reduce-narrow-empty-identity.wgsl.jinja +16 -0
build/webgpu/reduce-noop-empty-axes.wgsl.jinja +13 -0
build/webgpu/reduce-row-subgroup.wgsl.jinja +82 -0
build/webgpu/reduce-row-tree.wgsl.jinja +96 -0
build/webgpu/reduce-serial-axis.wgsl.jinja +161 -0
build/webgpu/test.json +1177 -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.ReduceMin
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 20
+## Description
+Computes the minimum of input tensor elements along the specified axes. The output rank matches the input when `keepdims` is 1; otherwise reduced dimensions are pruned. Reduction over an empty set yields positive infinity when the dtype supports it, or the dtype's maximum value otherwise. For Boolean inputs, `false` is less than `true`.
+See the [ONNX `ReduceMin` spec](https://onnx.ai/onnx/operators/onnx__ReduceMin.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 containing minimum values. | required |
+## Attributes
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `keepdims` | `1` | If 1, retains the reduced dimensions with size 1 in the output; if 0, those dimensions are pruned. |
+| `noop_with_empty_axes` | `0` | When axes is empty: if 0 (default) reduces over all axes; if 1 the op acts as an identity (no-op reduction). |
+| `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`, `uint32`, `int8`, `uint8`, `bool` |
+## 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
+- [`datamove-flat-copy-x4-tail.wgsl.jinja`](build/webgpu/datamove-flat-copy-x4-tail.wgsl.jinja)
+- [`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-partial.wgsl.jinja`](build/webgpu/reduce-flat-partial.wgsl.jinja)
+- [`reduce-narrow-empty-identity.wgsl.jinja`](build/webgpu/reduce-narrow-empty-identity.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.ReduceMin", { 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,108 @@

+{
+  "op": "ai.onnx.ReduceMin",
+  "cases": [
+    {
+      "name": "reduce-min-axis0-1024x512",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "rows": 1024, "cols": 512 },
+      "attrs": { "axes": [0], "keepdims": 0 },
+      "inputs": { "x": { "shape": [1024, 512], "dtype": "float32", "dist": "normal", "seed": 953, "scale": 1 } },
+      "outputs": { "y": { "shape": [512], "dtype": "float32" } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "args.rows * args.cols * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "rows_1024_cols_768",
+      "attrs": { "axes": [1], "keepdims": 0 },
+      "inputs": { "x": { "dtype": "float32", "shape": [1024, 768] } },
+      "outputs": { "y": { "dtype": "float32", "shape": [1024] } }
+    },
+    {
+      "name": "reducemin-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", "value": "args.rows * args.cols * 4" }] }
+    },
+    {
+      "name": "reducemin-r3-last-64x64x4096",
+      "preset": "smoke",
+      "vars": { "rows": 4096, "cols": 4096 },
+      "attrs": { "axes": [2], "keepdims": 0 },
+      "inputs": { "x": { "shape": [64, 64, 4096], "dtype": "float32", "dist": "normal", "seed": 72, "scale": 0.2 } },
+      "outputs": { "y": { "shape": [64, 64], "dtype": "float32" } },
+      "bench": { "primary": true, "metrics": [{ "type": "bandwidth", "value": "args.rows * args.cols * 4" }] }
+    },
+    {
+      "name": "reducemin-lastaxis-f32-4096x4095-scalar-subgroup",
+      "vars": { "dtype": "float32", "rows": 4096, "cols": 4095 },
+      "attrs": { "axes": [1], "keepdims": 0 },
+      "inputs": { "x": { "shape": [4096, 4095], "dtype": "float32", "dist": "normal", "seed": 414, "scale": 0.2 } },
+      "outputs": { "y": { "shape": [4096], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.rows * args.cols * dtypeBytes(args.dtype)" }] }
+    },
+    {
+      "name": "reducemin-axis0-splitk-f32-8192x512",
+      "vars": { "dtype": "float32", "rows": 8192, "cols": 512 },
+      "attrs": { "axes": [0], "keepdims": 0 },
+      "inputs": { "x": { "shape": [8192, 512], "dtype": "float32", "dist": "normal", "seed": 818, "scale": 0.2 } },
+      "outputs": { "y": { "shape": [512], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.rows * args.cols * dtypeBytes(args.dtype)" }] }
+    },
+    {
+      "name": "reducemin-axis1-f32-16776961x1-degenerate-row",
+      "vars": { "dtype": "float32", "rows": 16776961, "cols": 1 },
+      "attrs": { "axes": [1], "keepdims": 0 },
+      "inputs": { "x": { "shape": [16776961, 1], "dtype": "float32", "dist": "normal", "seed": 167, "scale": 0.2 } },
+      "outputs": { "y": { "shape": [16776961], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.rows * args.cols * dtypeBytes(args.dtype)" }] }
+    },
+    {
+      "name": "reducemin-all-axes-flat-miss-single-thread-2M",
+      "preset": "stress",
+      "vars": { "rows": 2146689, "cols": 1 },
+      "attrs": { "keepdims": 0 },
+      "inputs": { "x": { "shape": [129, 129, 129], "dtype": "float32", "dist": "normal", "seed": 722, "scale": 1 } },
+      "outputs": { "y": { "shape": [], "dtype": "float32", "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "129 * 129 * 129 * 4" }] }
+    },
+    {
+      "name": "reducemin-rank3-spatial-axes12-f32-128x256x256-pathology",
+      "preset": "stress",
+      "provenance": {
+        "source": "authored for branch coverage",
+        "notes": "Rank-3 multi-axis reduction with 128 output lanes, each serially scanning a 256x256 plane."
+      },
+      "attrs": { "axes": [1, 2], "keepdims": 1 },
+      "inputs": { "x": { "shape": [128, 256, 256], "dtype": "float32", "dist": "normal", "seed": 824, "scale": 0.2 } },
+      "outputs": { "y": { "shape": [128, 1, 1], "dtype": "float32", "dist": "empty" } },
+      "bench": { "primary": true, "metrics": [{ "type": "bandwidth", "value": "128 * 256 * 256 * 4" }] }
+    },
+    {
+      "name": "reducemin-rank4-spatial-axes23-f32-2x64x256x256-pathology",
+      "preset": "stress",
+      "provenance": {
+        "source": "authored for branch coverage",
+        "notes": "Rank-4 multi-axis reduction with 128 output lanes, each serially scanning a 256x256 plane."
+      },
+      "attrs": { "axes": [2, 3], "keepdims": 1 },
+      "inputs": { "x": { "shape": [2, 64, 256, 256], "dtype": "float32", "dist": "normal", "seed": 825, "scale": 0.2 } },
+      "outputs": { "y": { "shape": [2, 64, 1, 1], "dtype": "float32", "dist": "empty" } },
+      "bench": { "primary": true, "metrics": [{ "type": "bandwidth", "value": "2 * 64 * 256 * 256 * 4" }] }
+    },
+    {
+      "name": "reducemin-axis0-cols3-ultra-low-parallelism",
+      "preset": "stress",
+      "vars": { "rows": 65536, "cols": 3 },
+      "attrs": { "axes": [0], "keepdims": 0 },
+      "inputs": { "x": { "shape": [65536, 3], "dtype": "float32", "dist": "normal", "seed": 823, "scale": 1 } },
+      "outputs": { "y": { "shape": [3], "dtype": "float32", "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "65536 * 3 * 4" }] }
+    }
+  ],
+  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] }
+}

build/webgpu/datamove-flat-copy-x4-tail.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,31 @@

+// One-dispatch flat copy that coarsens four scalar elements per invocation and
+// handles a final one-to-three-element tail in the same kernel. Scalar bindings
+// keep an arbitrary element count legal; the four explicit adjacent accesses
+// let the backend combine full groups into wide transactions.
+{{ env.wgsl.resourceDeclarations }}
+const COUNT: u32 = {{ source.count }}u;
+const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>,
+        @builtin(num_workgroups) nwg: vec3<u32>) {
+  let q = gid.x + gid.y * nwg.x * WG;
+  let i = q * 4u;
+  if (i >= COUNT) {
+    return;
+  }
+  if (i + 3u < COUNT) {
+    let v = vec4<{{ scalar }}>(src[i], src[i + 1u], src[i + 2u], src[i + 3u]);
+    dst[i] = v.x;
+    dst[i + 1u] = v.y;
+    dst[i + 2u] = v.z;
+    dst[i + 3u] = v.w;
+    return;
+  }
+  for (var j = i; j < COUNT; j = j + 1u) {
+    dst[j] = src[j];
+  }
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,1074 @@

+{
+  "domain": "ai.onnx",
+  "name": "ReduceMin",
+  "sinceVersion": 20,
+  "description": "Computes the minimum of input tensor elements along the specified axes. The output rank matches the input when `keepdims` is 1; otherwise reduced dimensions are pruned. Reduction over an empty set yields positive infinity when the dtype supports it, or the dtype's maximum value otherwise. For Boolean inputs, `false` is less than `true`.",
+  "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 containing minimum values."
+    }
+  ],
+  "attributes": { "keepdims": 1, "noop_with_empty_axes": 0, "axes": [] },
+  "attributeDescriptions": {
+    "keepdims": "If 1, retains the reduced dimensions with size 1 in the output; if 0, those dimensions are pruned.",
+    "noop_with_empty_axes": "When axes is empty: if 0 (default) reduces over all axes; if 1 the op acts as an identity (no-op reduction).",
+    "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", "uint32", "int8", "uint8", "bool"] },
+  "args": {
+    "x": { "kind": "tensor", "semantic": "data", "role": "input" },
+    "y": { "kind": "tensor", "semantic": "reduced", "role": "output" }
+  },
+  "derive": {
+    "emptyNarrowInputOk": "numel(shapes.data) > 0 or (tensorDtypes.data != \"int8\" and tensorDtypes.data != \"uint8\")",
+    "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": "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": "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": "floor(numel(shapes.data) / tunables.VECTOR_WIDTH)",
+    "flatSplitCount": "max(1, min(tunables.FULL_REDUCE_MAX_SPLITS, ceilDiv(flatItems, reduceWorkgroupSize)))",
+    "flatScratchBytes": "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,
+    "ROW_SERIAL_MIN_ROWS": 8192,
+    "ROW_SERIAL_MAX_COLS": 1024
+  },
+  "bindingSets": {
+    "narrowEmpty": [
+      { "name": "y", "arg": "y", "semantic": "reduced", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "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)" }
+          ]
+        }
+      }
+    ],
+    "copyX4Tail": [
+      {
+        "name": "src",
+        "arg": "x",
+        "semantic": "data",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar"
+      },
+      { "name": "dst", "arg": "y", "semantic": "reduced", "buffer": { "type": "storage" }, "elementType": "$scalar" }
+    ],
+    "flatPartialF32": [
+      { "name": "x", "arg": "x", "semantic": "data", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
+      { "name": "partials", "semantic": "partials", "buffer": { "type": "storage" }, "elementType": "f32" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "count4", "type": "u32", "value": "floor(numel(shapes.data) / tunables.VECTOR_WIDTH)" },
+            { "name": "numel", "type": "u32", "value": "numel(shapes.data)" }
+          ]
+        }
+      }
+    ],
+    "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" }] }
+      }
+    ],
+    "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)" }
+          ]
+        }
+      }
+    ],
+    "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": "narrow_empty_identity",
+      "priority": 100,
+      "when": ["numel(shapes.data) == 0", "numel(shapes.reduced) > 0", "tensorDtypes.data == \"int8\" or tensorDtypes.data == \"uint8\" or tensorDtypes.data == \"bool\""],
+      "constants": { "scalar": "dtypes.T" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ReduceMin.NarrowEmptyIdentity",
+          "source": {
+            "shader": "reduce-narrow-empty-identity.wgsl.jinja",
+            "inputs": {
+              "identity": "1 if tensorDtypes.data == \"bool\" else 127 if tensorDtypes.data == \"int8\" else 255"
+            }
+          },
+          "bindings": "narrowEmpty",
+          "dispatch": { "gridStride": "numel(shapes.reduced)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "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": "ReduceMin.ContiguousSuffixSubgroupVec4",
+          "source": {
+            "shader": "reduce-row-subgroup.wgsl.jinja",
+            "inputs": {
+              "op": "\"min\"",
+              "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": "ReduceMin.ContiguousSuffixTreeVec4",
+          "source": {
+            "shader": "reduce-row-tree.wgsl.jinja",
+            "inputs": {
+              "op": "\"min\"",
+              "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": "ReduceMin.ContiguousSuffixTree",
+          "source": {
+            "shader": "reduce-row-tree.wgsl.jinja",
+            "inputs": { "op": "\"min\"", "castF32": "dtypes.T == \"f16\"", "usesF16": "dtypes.T == \"f16\"" }
+          },
+          "bindings": "suffixScalar",
+          "dispatch": { "workgroups": "numel(shapes.reduced)" }
+        }
+      ]
+    },
+    {
+      "id": "degenerate_copy_x4_tail",
+      "priority": 46,
+      "when": ["(dtypes.T == \"f32\" or dtypes.T == \"i32\" or dtypes.T == \"u32\")", "numel(shapes.data) >= 4", "numel(shapes.data) == numel(shapes.reduced)"],
+      "constants": { "scalar": "dtypes.T", "usesF16": false },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ReduceMin.DegenerateCopyX4Tail",
+          "source": {
+            "shader": "datamove-flat-copy-x4-tail.wgsl.jinja",
+            "inputs": { "count": "numel(shapes.reduced)" }
+          },
+          "bindings": "copyX4Tail",
+          "dispatch": { "threads": "ceilDiv(numel(shapes.reduced), 4)", "workgroupSize": "reduceWorkgroupSize" },
+          "viewAlias": [{ "input": "src", "output": "dst" }]
+        }
+      ]
+    },
+    {
+      "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": "ReduceMin.NoopEmptyAxes",
+          "source": { "shader": "reduce-noop-empty-axes.wgsl.jinja" },
+          "bindings": "elementwise",
+          "dispatch": { "threads": "numel(shapes.reduced)", "workgroupSize": "reduceWorkgroupSize" }
+        }
+      ]
+    },
+    {
+      "id": "multi_axis_rank3",
+      "priority": 8,
+      "constants": { "scalar": "dtypes.T" },
+      "when": ["not flatParallelCovered", "not contiguousSuffixParallelCovered", "emptyNarrowInputOk", "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": "ReduceMin.MultiAxisRank3",
+          "source": {
+            "shader": "reduce-serial-axis.wgsl.jinja",
+            "inputs": {
+              "op": "\"min\"",
+              "indexing": "\"multiaxis\"",
+              "castF32": "dtypes.T == \"f16\"",
+              "usesF16": "dtypes.T == \"f16\"",
+              "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",
+              "logicalBool": "tensorDtypes.data == \"bool\""
+            }
+          },
+          "bindings": "multiAxis",
+          "dispatch": { "threads": "numel(shapes.reduced)", "workgroupSize": "reduceWorkgroupSize" }
+        }
+      ]
+    },
+    {
+      "id": "multi_axis_rank4",
+      "priority": 8,
+      "when": ["not flatParallelCovered", "not contiguousSuffixParallelCovered", "emptyNarrowInputOk", "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": "ReduceMin.MultiAxisRank4",
+          "source": {
+            "shader": "reduce-serial-axis.wgsl.jinja",
+            "inputs": {
+              "op": "\"min\"",
+              "indexing": "\"multiaxis\"",
+              "castF32": "dtypes.T == \"f16\"",
+              "usesF16": "dtypes.T == \"f16\"",
+              "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",
+              "logicalBool": "tensorDtypes.data == \"bool\""
+            }
+          },
+          "bindings": "multiAxis",
+          "dispatch": { "threads": "numel(shapes.reduced)", "workgroupSize": "reduceWorkgroupSize" }
+        }
+      ],
+      "constants": { "scalar": "dtypes.T" }
+    },
+    {
+      "id": "tree_last_axis",
+      "priority": 20,
+      "demoteWhen": ["rowSerialPreferred"],
+      "when": ["not flatParallelCovered", "(dtypes.T == \"f32\" or dtypes.T == \"i32\" or dtypes.T == \"u32\")", "ranks.data >= 2", "reduceAxis == ranks.data - 1", "numel(shapes.reduced) == rows(shapes.data, ranks.data - 1)", "lastAxisCols >= tunables.ROW_PARALLEL_MIN_COLS", "treeWorkgroupOk"],
+      "constants": {
+        "scalar": "dtypes.T",
+        "workgroupSize": "min(reduceWorkgroupSize, pow2ceil(dim(shapes.data, ranks.data - 1)))"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ReduceMin.TreeLastAxis",
+          "source": { "shader": "reduce-row-tree.wgsl.jinja", "inputs": { "op": "\"min\"" } },
+          "bindings": "lastAxisScalar",
+          "dispatch": { "workgroups": "rows(shapes.data, ranks.data - 1)" }
+        }
+      ]
+    },
+    {
+      "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": "ReduceMin.TreeRowVec4",
+          "source": {
+            "shader": "reduce-row-tree.wgsl.jinja",
+            "inputs": {
+              "op": "\"min\"",
+              "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": ["emptyNarrowInputOk", "f16Ok(dtypes.T)", "ranks.data == 0", "ranks.reduced == 0"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "ReduceMin.Rank0Scalar",
+          "source": {
+            "shader": "reduce-serial-axis.wgsl.jinja",
+            "inputs": {
+              "op": "\"min\"",
+              "indexing": "\"axis2d\"",
+              "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", "emptyNarrowInputOk", "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": "ReduceMin.Rank1Axis0",
+          "source": {
+            "shader": "reduce-serial-axis.wgsl.jinja",
+            "inputs": {
+              "op": "\"min\"",
+              "indexing": "\"axis2d\"",
+              "castF32": "dtypes.T == \"f16\"",
+              "usesF16": "dtypes.T == \"f16\"",
+              "logicalBool": "tensorDtypes.data == \"bool\""
+            }
+          },
+          "bindings": "rank1Axis0",
+          "dispatch": { "threads": "numel(shapes.reduced)", "workgroupSize": "reduceWorkgroupSize" }
+        }
+      ]
+    },
+    {
+      "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": "[splitCount * axisSplitOutputs]" }],
+      "passes": [
+        {
+          "id": "split_reduce",
+          "name": "ReduceMin.AxisSplitReduce",
+          "source": {
+            "shader": "reduce-axis-split-reduce.wgsl.jinja",
+            "inputs": {
+              "op": "\"min\"",
+              "workgroupSize": "reduceWorkgroupSize",
+              "split": "splitCount",
+              "castF32": "dtypes.T == \"f16\"",
+              "usesF16": "dtypes.T == \"f16\""
+            }
+          },
+          "bindings": "axisSplitReduce",
+          "dispatch": { "threads": "axisSplitOutputs", "workgroupSize": "reduceWorkgroupSize", "y": "splitCount" }
+        },
+        {
+          "id": "combine",
+          "name": "ReduceMin.AxisSplitCombine",
+          "source": {
+            "shader": "reduce-axis0-splitk-combine.wgsl.jinja",
+            "inputs": {
+              "op": "\"min\"",
+              "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": "[splitCount * axisSplitOutputs]" }],
+      "passes": [
+        {
+          "id": "split_reduce",
+          "name": "ReduceMin.AxisSplitTiledReduce",
+          "source": {
+            "shader": "reduce-axis0-tilecols.wgsl.jinja",
+            "inputs": {
+              "op": "\"min\"",
+              "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": "ReduceMin.AxisSplitCombine",
+          "source": {
+            "shader": "reduce-axis0-splitk-combine.wgsl.jinja",
+            "inputs": {
+              "op": "\"min\"",
+              "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": "[splitCount * dim(shapes.data, 1)]" }],
+      "passes": [
+        {
+          "id": "split_reduce",
+          "name": "ReduceMin.Axis0SplitKReduce",
+          "source": {
+            "shader": "reduce-axis0-splitk-reduce.wgsl.jinja",
+            "inputs": {
+              "op": "\"min\"",
+              "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": "ReduceMin.Axis0SplitKCombine",
+          "source": {
+            "shader": "reduce-axis0-splitk-combine.wgsl.jinja",
+            "inputs": {
+              "op": "\"min\"",
+              "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", "emptyNarrowInputOk", "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": {
+        "scalar": "dtypes.T",
+        "workgroupSize": "reduceWorkgroupSize",
+        "tileCols": "tunables.AXIS0_TILE_COLS"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ReduceMin.Axis0TileCols",
+          "source": {
+            "shader": "reduce-axis0-tilecols.wgsl.jinja",
+            "inputs": { "op": "\"min\"", "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", "split": "flatSplitCount" },
+      "when": ["flatParallelCovered"],
+      "intermediates": [{ "id": "partials", "dtype": "float32", "shape": "[flatSplitCount]" }],
+      "passes": [
+        {
+          "id": "flat_partial",
+          "name": "ReduceMin.AllAxesFlatPartial",
+          "source": {
+            "shader": "reduce-flat-partial.wgsl.jinja",
+            "inputs": { "op": "\"min\"", "castF32": "dtypes.T == \"f16\"", "usesF16": "dtypes.T == \"f16\"" }
+          },
+          "bindings": "flatPartialF32",
+          "dispatch": { "x": "flatSplitCount" }
+        },
+        {
+          "id": "combine",
+          "name": "ReduceMin.AllAxesFlatCombine",
+          "source": {
+            "shader": "reduce-axis0-splitk-combine.wgsl.jinja",
+            "inputs": { "op": "\"min\"", "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", "tree_last_axis"],
+      "when": ["not flatParallelCovered", "emptyNarrowInputOk", "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": "ReduceMin.RankNSingleAxisGeneric",
+          "source": {
+            "shader": "reduce-serial-axis.wgsl.jinja",
+            "inputs": {
+              "op": "\"min\"",
+              "indexing": "\"rankn\"",
+              "castF32": "dtypes.T == \"f16\"",
+              "usesF16": "dtypes.T == \"f16\"",
+              "rank": "ranks.data",
+              "axis": "reduceAxis",
+              "dataShape": "shapes.data",
+              "outputShape": "shapes.reduced",
+              "outputRank": "ranks.reduced",
+              "keepDims": "attrs.keepdims != 0",
+              "logicalBool": "tensorDtypes.data == \"bool\""
+            }
+          },
+          "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", "emptyNarrowInputOk", "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": "ReduceMin.SubgroupRowVec4",
+          "source": {
+            "shader": "reduce-row-subgroup.wgsl.jinja",
+            "inputs": {
+              "op": "\"min\"",
+              "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", "emptyNarrowInputOk", "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": "ReduceMin.SubgroupRow",
+          "source": {
+            "shader": "reduce-row-subgroup.wgsl.jinja",
+            "inputs": {
+              "op": "\"min\"",
+              "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", "emptyNarrowInputOk", "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": "\"min\"",
+              "indexing": "\"axis2d\"",
+              "castF32": "dtypes.T == \"f16\"",
+              "usesF16": "dtypes.T == \"f16\"",
+              "logicalBool": "tensorDtypes.data == \"bool\""
+            }
+          },
+          "bindings": "rank2Serial",
+          "constants": { "axis": 0, "scalar": "dtypes.T" },
+          "dispatch": { "threads": "numel(shapes.reduced)", "workgroupSize": "reduceWorkgroupSize" }
+        }
+      ]
+    },
+    {
+      "id": "axis1",
+      "priority": 0,
+      "constants": { "axis": 1, "scalar": "dtypes.T" },
+      "when": ["not flatParallelCovered", "emptyNarrowInputOk", "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": "\"min\"",
+              "indexing": "\"axis2d\"",
+              "castF32": "dtypes.T == \"f16\"",
+              "usesF16": "dtypes.T == \"f16\"",
+              "logicalBool": "tensorDtypes.data == \"bool\""
+            }
+          },
+          "bindings": "rank2SerialAxis1",
+          "constants": { "axis": 1, "scalar": "dtypes.T" },
+          "dispatch": { "threads": "numel(shapes.reduced)", "workgroupSize": "reduceWorkgroupSize" }
+        }
+      ]
+    },
+    {
+      "id": "all_axes_keepdims",
+      "priority": 30,
+      "constants": { "axis": 0, "scalar": "dtypes.T" },
+      "when": ["not flatParallelCovered", "emptyNarrowInputOk", "f16Ok(dtypes.T)", "ranks.data >= 3", "attrs.keepdims == 1 and ranks.reduced == ranks.data and numel(shapes.reduced) == 1"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "ReduceMin.Rank3AllAxesKeepdims",
+          "source": {
+            "shader": "reduce-serial-axis.wgsl.jinja",
+            "inputs": {
+              "op": "\"min\"",
+              "indexing": "\"axis2d\"",
+              "castF32": "dtypes.T == \"f16\"",
+              "usesF16": "dtypes.T == \"f16\"",
+              "logicalBool": "tensorDtypes.data == \"bool\""
+            }
+          },
+          "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", "emptyNarrowInputOk", "f16Ok(dtypes.T)", "ranks.data >= 3", "attrs.keepdims == 0 and ranks.reduced == 0"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "ReduceMin.Rank3AllAxesNoKeepdims",
+          "source": {
+            "shader": "reduce-serial-axis.wgsl.jinja",
+            "inputs": {
+              "op": "\"min\"",
+              "indexing": "\"axis2d\"",
+              "castF32": "dtypes.T == \"f16\"",
+              "usesF16": "dtypes.T == \"f16\"",
+              "logicalBool": "tensorDtypes.data == \"bool\""
+            }
+          },
+          "bindings": "fullReduceSerial",
+          "dispatch": { "threads": "numel(shapes.reduced)", "workgroupSize": "reduceWorkgroupSize" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

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

+{
+  "name": "ai.onnx.ReduceMin",
+  "id": "_ai_onnx_reducemin_webgpu_a5e3e3d",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "6KR8yCFLrmY6PdH2GnqpEOjK72FRANRAv+s5kcNoDnk=",
+      "datamove-flat-copy-x4-tail.wgsl.jinja": "zqQLkq4mysszKvFcAIzayrcBx0qTuJf2zavso5HN/+A=",
+      "manifest.json": "VUjwoC99OWimXu2iEgqL+ufHyaEZ60INXqPW3dw3Q6Q=",
+      "reduce-axis-split-reduce.wgsl.jinja": "rjy37ckoB9/9a88KxgKKjpZYjmIjy8vmI34yLvPzxuo=",
+      "reduce-axis0-splitk-combine.wgsl.jinja": "YOriZzB3yMakRsErMyOJhirJBBtgSenl+Gd4rBsHZsM=",
+      "reduce-axis0-splitk-reduce.wgsl.jinja": "QbHcxtUfEvqMJQRstIC89pBdlD7gDzTFZSdYH1Wldn8=",
+      "reduce-axis0-tilecols.wgsl.jinja": "ldhjG9yaosrcAFdq6OCglYkvJEaQWi38KdCNqx4N4Yw=",
+      "reduce-flat-partial.wgsl.jinja": "Esf7shfEzRgtyIkjVRYHlqU9Rc1lYCj5WqNLZgBxsSo=",
+      "reduce-narrow-empty-identity.wgsl.jinja": "S6UnLW+6dyCb9ciRezkgXE9mlYmqdfBCjPRO9iPm3Jk=",
+      "reduce-noop-empty-axes.wgsl.jinja": "IO2MEzyBGJr3Z8AVP7ELJXqfhwcG/vtBfa0wCdkbuwo=",
+      "reduce-row-subgroup.wgsl.jinja": "szYjtA6TUPJ/I4mbVX7WZUoXiseQ3JfulzqgGb6fdt4=",
+      "reduce-row-tree.wgsl.jinja": "srFhdHteraJrXxSPG56B56yK9qjUjRAzCEakQKBsuLY=",
+      "reduce-serial-axis.wgsl.jinja": "cLrTc0ite8YZmx/wn7PBwaxTp9dLWWUIUdQQ3tH8+4o=",
+      "test.json": "G8Sv1RfEOyXjQTYVWqiWogZbCuoav1pCHqswBcjcM6E="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.ReduceMin" }
+}

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

	@@ -0,0 +1,63 @@

+// 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. */
+{% macro wgsl_minmax_identity(name, op, scalar="f32") %}
+{% set logicalBoolIdentity = logicalBool is defined and logicalBool %}
+fn {{ name }}() -> {{ scalar }} {
+{% if scalar == "i32" %}
+  return {{ "-2147483647i - 1i" if op == "max" else "2147483647i" }};
+{% elif scalar == "u32" %}
+  return {{ "1u" if logicalBoolIdentity and op == "min" else "0u" if op == "max" else "4294967295u" }};
+{% else %}
+  var bits = {{ "0xff800000u" if op == "max" else "0x7f800000u" }};
+  return bitcast<f32>(bits);
+{% endif %}
+}
+{%- endmacro %}
+const WG: u32 = {{ workgroupSize }}u;
+const SPLIT: u32 = {{ split }}u;
+{{ wgsl_minmax_identity("reduction_identity", source.op) }}
+@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 acc = reduction_identity();
+  for (var axis_index = a0; axis_index < a1; axis_index = axis_index + 1u) {
+    acc = {{ source.op }}(acc, {{ xa }}x[input_base + axis_index * params.inner]{{ ax }});
+  }
+  partials[seg * params.outputs + output_index] = acc;
+}

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

	@@ -0,0 +1,48 @@

+// 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.
+{% 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. */
+{% macro wgsl_minmax_identity(name, op, scalar="f32") %}
+{% set logicalBoolIdentity = logicalBool is defined and logicalBool %}
+fn {{ name }}() -> {{ scalar }} {
+{% if scalar == "i32" %}
+  return {{ "-2147483647i - 1i" if op == "max" else "2147483647i" }};
+{% elif scalar == "u32" %}
+  return {{ "1u" if logicalBoolIdentity and op == "min" else "0u" if op == "max" else "4294967295u" }};
+{% else %}
+  var bits = {{ "0xff800000u" if op == "max" else "0x7f800000u" }};
+  return bitcast<f32>(bits);
+{% endif %}
+}
+{%- endmacro %}
+const WG: u32 = {{ workgroupSize }}u;
+const SPLIT: u32 = {{ split }}u;
+{{ wgsl_minmax_identity("reduction_identity", source.op) }}
+@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) {
+    var total = reduction_identity();
+    for (var seg = 0u; seg < SPLIT; seg = seg + 1u) {
+      let p = partials[seg * params.cols + col];
+      total = {{ source.op }}(total, p);
+    }
+{% if source.outputF16 %}
+    y[col] = f16(total);
+{% else %}
+    y[col] = total;
+{% endif %}
+  }
+}

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

	@@ -0,0 +1,58 @@

+// 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. */
+{% macro wgsl_minmax_identity(name, op, scalar="f32") %}
+{% set logicalBoolIdentity = logicalBool is defined and logicalBool %}
+fn {{ name }}() -> {{ scalar }} {
+{% if scalar == "i32" %}
+  return {{ "-2147483647i - 1i" if op == "max" else "2147483647i" }};
+{% elif scalar == "u32" %}
+  return {{ "1u" if logicalBoolIdentity and op == "min" else "0u" if op == "max" else "4294967295u" }};
+{% else %}
+  var bits = {{ "0xff800000u" if op == "max" else "0x7f800000u" }};
+  return bitcast<f32>(bits);
+{% endif %}
+}
+{%- endmacro %}
+const WG: u32 = {{ workgroupSize }}u;
+const SPLIT: u32 = {{ split }}u;
+{{ wgsl_minmax_identity("reduction_identity", source.op) }}
+@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 acc = reduction_identity();
+  for (var row = r0; row < r1; row = row + 1u) {
+    acc = {{ source.op }}(acc, {{ xa }}x[row * params.cols + col]{{ ax }});
+  }
+  partials[seg * params.cols + col] = acc;
+}

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

	@@ -0,0 +1,105 @@

+// 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. */
+{% macro wgsl_minmax_identity(name, op, scalar="f32") %}
+{% set logicalBoolIdentity = logicalBool is defined and logicalBool %}
+fn {{ name }}() -> {{ scalar }} {
+{% if scalar == "i32" %}
+  return {{ "-2147483647i - 1i" if op == "max" else "2147483647i" }};
+{% elif scalar == "u32" %}
+  return {{ "1u" if logicalBoolIdentity and op == "min" else "0u" if op == "max" else "4294967295u" }};
+{% else %}
+  var bits = {{ "0xff800000u" if op == "max" else "0x7f800000u" }};
+  return bitcast<f32>(bits);
+{% endif %}
+}
+{%- endmacro %}
+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 %}
+{{ wgsl_minmax_identity("lane_identity", source.op, scalar) }}
+var<workgroup> partial: array<{{ scalar if (source.op == "max" or source.op == "min" or intMode) else "f32" }}, WG>;
+@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 %}
+  var acc = lane_identity();
+  if (in_bounds) {
+    for (var row = {{ rowBegin }}; row < {{ rowEnd }}; row = row + ROW_LANES) {
+      acc = {{ source.op }}(acc, {{ elem }});
+    }
+  }
+  partial[tid] = acc;
+  workgroupBarrier();
+  if (row_lane == 0u && in_bounds) {
+    var total = partial[col_lane];
+    for (var lane = 1u; lane < ROW_LANES; lane = lane + 1u) {
+      total = {{ source.op }}(total, partial[lane * TILE_COLS + col_lane]);
+    }
+{% if splitMode %}
+    partials[seg * params.outputs + outputIndex] = total;
+{% else %}
+    y[outputIndex] = {{ yv }}total{{ vy }};
+{% endif %}
+  }
+}

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

	@@ -0,0 +1,75 @@

+// Pass 1 of a flat reduction to one scalar. Instead of assigning the whole
+// tensor to one invocation, SPLIT workgroups traverse the flat input, accumulate
+// in registers, reduce within each workgroup, and write one partial each. A
+// following combine pass folds the partials and applies the operation finalizer.
+//
+// Scalar f32 bindings keep arbitrary element counts legal. The grid-stride loop
+// manually assembles full vec4 groups from contiguous scalars, and one global
+// thread folds the final zero-to-three scalar elements exactly once.
+{% 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. */
+{% macro wgsl_minmax_identity(name, op, scalar="f32") %}
+{% set logicalBoolIdentity = logicalBool is defined and logicalBool %}
+fn {{ name }}() -> {{ scalar }} {
+{% if scalar == "i32" %}
+  return {{ "-2147483647i - 1i" if op == "max" else "2147483647i" }};
+{% elif scalar == "u32" %}
+  return {{ "1u" if logicalBoolIdentity and op == "min" else "0u" if op == "max" else "4294967295u" }};
+{% else %}
+  var bits = {{ "0xff800000u" if op == "max" else "0x7f800000u" }};
+  return bitcast<f32>(bits);
+{% endif %}
+}
+{%- endmacro %}
+const WG: u32 = {{ workgroupSize }}u;
+var<workgroup> red: array<{{ "i32" if source.intMode else "f32" }}, WG>;
+{{ wgsl_minmax_identity("reduction_identity", source.op) }}
+@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;
+  var acc = reduction_identity();
+  // Grid-stride over the flat vec4 groups (params.count4 = numel / 4, floored).
+  for (var i = gid.x; i < params.count4; i = i + gstride) {
+    let b = 4u * i;
+    let v = vec4<{{ "i32" if source.intMode else "f32" }}>({{ xa }}x[b]{{ ax }}, {{ xa }}x[b + 1u]{{ ax }}, {{ xa }}x[b + 2u]{{ ax }}, {{ xa }}x[b + 3u]{{ ax }});
+    acc = min(acc, min(min(v.x, v.y), min(v.z, v.w)));
+  }
+  // Scalar tail (the 0..3 elements past the last full vec4). One global thread
+  // folds it so it is counted exactly once; the count is tiny so serializing it
+  // is negligible. numel % 4 == 0 makes this loop empty.
+  if (gid.x == 0u) {
+    for (var i = 4u * params.count4; i < params.numel; i = i + 1u) {
+      let s = {{ xa }}x[i]{{ ax }};
+      acc = min(acc, s);
+    }
+  }
+  red[tid] = acc;
+  workgroupBarrier();
+  var stride: u32 = WG / 2u;
+  loop {
+    if (stride == 0u) { break; }
+    if (tid < stride) {
+      red[tid] = min(red[tid], red[tid + stride]);
+    }
+    stride = stride / 2u;
+    workgroupBarrier();
+  }
+  if (tid == 0u) {
+    partials[wg.x] = red[0];
+  }
+}

build/webgpu/reduce-narrow-empty-identity.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,16 @@

+{{ env.wgsl.resourceDeclarations }}
+const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
+// Empty min/max reductions use the logical narrow dtype's bound, while their
+// WebGPU storage is widened to i32/u32.
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(
+  @builtin(global_invocation_id) gid: vec3<u32>,
+  @builtin(num_workgroups) nwg: vec3<u32>
+) {
+  let stride = nwg.x * WG;
+  for (var i = gid.x; i < params.count; i += stride) {
+    y[i] = {{ scalar }}({{ source.identity }});
+  }
+}

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,82 @@

+// 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. */
+{% macro wgsl_minmax_identity(name, op, scalar="f32") %}
+{% set logicalBoolIdentity = logicalBool is defined and logicalBool %}
+fn {{ name }}() -> {{ scalar }} {
+{% if scalar == "i32" %}
+  return {{ "-2147483647i - 1i" if op == "max" else "2147483647i" }};
+{% elif scalar == "u32" %}
+  return {{ "1u" if logicalBoolIdentity and op == "min" else "0u" if op == "max" else "4294967295u" }};
+{% else %}
+  var bits = {{ "0xff800000u" if op == "max" else "0x7f800000u" }};
+  return bitcast<f32>(bits);
+{% endif %}
+}
+{%- endmacro %}
+const WG: u32 = {{ workgroupSize }}u;{{ wgsl_minmax_identity("reduction_identity", source.op, scalar) }}
+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", "subgroupMin", "total = min(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;  let INIT: {{ scalar }} = reduction_identity();{% if source.vec4 %}
+  var acc4 = vec4<{{ scalar }}>(INIT);
+  for (var c = tid; c < params.chunkCount; c = c + WG) {
+    let v = {{ xv }}x[base + c]{{ vx }};    acc4 = min(acc4, v);  }  let acc = min(min(acc4.x, acc4.y), min(acc4.z, acc4.w));{% else %}
+  var acc = INIT;
+  for (var c = tid; c < params.chunkCount; c = c + WG) {
+    let v = {{ xv }}x[base + c]{{ vx }};    acc = min(acc, v);  }
+{%- endif %}
+  let total = reduce_row(acc, sgLid, sgId, numSg);
+  if (tid == 0u) {    y[row] = {{ yv }}total{{ vy }};  }}

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

	@@ -0,0 +1,96 @@

+// 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 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. */
+{% macro wgsl_minmax_identity(name, op, scalar="f32") %}
+{% set logicalBoolIdentity = logicalBool is defined and logicalBool %}
+fn {{ name }}() -> {{ scalar }} {
+{% if scalar == "i32" %}
+  return {{ "-2147483647i - 1i" if op == "max" else "2147483647i" }};
+{% elif scalar == "u32" %}
+  return {{ "1u" if logicalBoolIdentity and op == "min" else "0u" if op == "max" else "4294967295u" }};
+{% else %}
+  var bits = {{ "0xff800000u" if op == "max" else "0x7f800000u" }};
+  return bitcast<f32>(bits);
+{% endif %}
+}
+{%- endmacro %}
+const WG: u32 = {{ workgroupSize }}u;
+{% set is_int = scalar == "i32" or scalar == "u32" %}
+{% set intAdditive = is_int and (source.op == "sum" or source.op == "l1" or source.op == "sumsquare") %}
+{% set accType = scalar if source.op == "max" or source.op == "min" or source.op == "prod" or intAdditive else "f32" %}
+{{ wgsl_minmax_identity("identity", source.op, accType) }}
+fn combine(a: {{ accType }}, b: {{ accType }}) -> {{ accType }} {
+  return min(a, b);
+}
+var<workgroup> partial: array<{{ accType }}, WG>;
+@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 source.vec4 %}
+  var acc4 = vec4<{{ accType }}>(identity());
+  for (var col = tid; col < params.chunkCount; col = col + WG) {
+    let value = {{ xv }}x[base + col]{{ vx }};
+    acc4 = min(acc4, value);
+  }
+  let acc = min(min(acc4.x, acc4.y), min(acc4.z, acc4.w));
+{% else %}
+  var acc = identity();
+  for (var col = tid; col < params.cols; col = col + WG) {
+    let value = {{ xv }}x[base + col]{{ vx }};
+    acc = combine(acc, value);
+  }
+{% endif %}
+  partial[tid] = acc;
+  workgroupBarrier();
+  for (var step = WG >> 1u; step > 0u; step = step >> 1u) {
+    if (tid < step) {
+      partial[tid] = combine(partial[tid], partial[tid + step]);
+    }
+    workgroupBarrier();
+  }
+  if (tid == 0u) {
+    let total = partial[0];
+    y[row] = {{ yv }}total{{ vy }};
+  }
+}

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

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

+// 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 logicalBool = source.logicalBool is defined and source.logicalBool %}
+{% 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. */
+{% macro wgsl_minmax_identity(name, op, scalar="f32") %}
+{% set logicalBoolIdentity = logicalBool is defined and logicalBool %}
+fn {{ name }}() -> {{ scalar }} {
+{% if scalar == "i32" %}
+  return {{ "-2147483647i - 1i" if op == "max" else "2147483647i" }};
+{% elif scalar == "u32" %}
+  return {{ "1u" if logicalBoolIdentity and op == "min" else "0u" if op == "max" else "4294967295u" }};
+{% else %}
+  var bits = {{ "0xff800000u" if op == "max" else "0x7f800000u" }};
+  return bitcast<f32>(bits);
+{% endif %}
+}
+{%- endmacro %}
+{{ wgsl_minmax_identity("empty_value", source.op, "f32" if castF32 else scalar) }}
+{% 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)]" %}
+{% set at_first = "x[input_offset(i, 0u)]" %}
+{% else %}
+{% set at = "x[input_offset(i)]" %}
+{% set at_first = "x[input_offset(i)]" %}
+{% endif %}
+{% elif source.indexing == "rankn" %}
+{% set count = "params.axisDim" %}
+{% set at = "x[input_offset(i, r)]" %}
+{% set at_first = "x[input_offset(i, 0u)]" %}
+{% elif axis == 0 %}
+{% set count = "params.rows" %}
+{% set at = "x[r * params.cols + i]" %}
+{% set at_first = "x[i]" %}
+{% else %}
+{% set count = "params.cols" %}
+{% set at = "x[i * params.cols + r]" %}
+{% set at_first = "x[i * params.cols]" %}
+{% endif %}
+{% if castF32 %}
+{% set at = "f32(" ~ at ~ ")" %}
+{% set at_first = "f32(" ~ at_first ~ ")" %}
+{% 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 ({{ count }} == 0u) {
+    y[i] = {{ yv }}empty_value(){{ vy }};
+    return;
+  }
+  var acc = {{ at_first }};
+  for (var r = 1u; r < {{ count }}; r = r + 1u) {
+    acc = {{ source.op }}(acc, {{ at }});
+  }
+  y[i] = {{ yv }}acc{{ vy }};
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,1177 @@

+{
+  "op": "ai.onnx.ReduceMin",
+  "fixtureArrays": {
+    "rank3_axis2_last_keepdims_input_x": [4, -1, 2, -1, 8, 7, 6, 5, -3, -4, -4, 2, 0, 9, -8, 3, -2, -6, 1, 10, 5, -5, -7, 4]
+  },
+  "cases": [
+    {
+      "name": "all_axes_flat_positive_infinity_identity_8192",
+      "provenance": {
+        "notes": "The parallel full-reduction threshold must supersede the rank1 serial and row-reduction fallbacks, while preserving the true +Infinity identity."
+      },
+      "attrs": { "axes": [0], "keepdims": 0 },
+      "inputs": { "x": { "dtype": "float32", "shape": [8192], "data": { "kind": "cycle", "values": ["Infinity"] } } },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [],
+          "data": { "kind": "cycle", "values": ["Infinity"] },
+          "tolerance": 0,
+          "relTolerance": 0
+        }
+      }
+    },
+    {
+      "name": "all_axes_flat_fullreduce_32x32x32_keepdims",
+      "attrs": { "keepdims": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [32, 32, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.027, "scale": 0.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 1], "tolerance": 0 } }
+    },
+    {
+      "name": "f32_negative_subnormal_axis1_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: WebGPU/Metal flushes subnormals to zero in f32; bit-exact subnormal preservation is unattainable on GPU."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/reduction/reduction_ops_test.cc",
+        "test": "ReduceMin",
+        "notes": "A negative subnormal is the row minimum over zero; flushing it changes the selected value."
+      },
+      "attrs": { "axes": [1], "keepdims": 0 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 3], "data": { "kind": "values", "values": [1e-40, 0.0, -1e-40] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1], "tolerance": 0 } }
+    },
+    {
+      "name": "dispatch_cliff_axis1_16776961x1",
+      "attrs": { "axes": [1], "keepdims": 0 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [16776961, 1], "data": { "kind": "linspace", "start": -5.0, "end": 5.0 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [16776961], "tolerance": 0 } }
+    },
+    {
+      "name": "degenerate_axis1_vec4_tail_17x1",
+      "provenance": {
+        "notes": "Compact correctness coverage for the vectorized singleton-axis copy plus its one-element scalar tail."
+      },
+      "attrs": { "axes": [1], "keepdims": 0 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [17, 1], "data": { "kind": "linspace", "start": -8.0, "end": 8.0 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [17], "tolerance": 0 } }
+    },
+    {
+      "name": "degenerate_axis1_vec4_aligned_i32_16x1",
+      "provenance": { "notes": "Compact native-integer coverage for the aligned vectorized singleton-axis copy." },
+      "attrs": { "axes": [1], "keepdims": 0 },
+      "inputs": {
+        "x": { "dtype": "int32", "shape": [16, 1], "data": { "kind": "cycle", "values": [16777217, -16777219, 7, -3] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [16], "tolerance": 0 } }
+    },
+    {
+      "name": "axis0",
+      "attrs": { "axes": [0], "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": { "kind": "values", "values": [1.0, 9.0, 3.0, 4.0, -1.0, 2.0, 7.0, 8.0, 0.0, 5.0, -3.0, 6.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4] } }
+    },
+    {
+      "name": "axis0_tilecols_positive_infinity_identity_64x16",
+      "attrs": { "axes": [0], "keepdims": 0 },
+      "inputs": { "x": { "dtype": "float32", "shape": [64, 16], "data": { "kind": "cycle", "values": ["Infinity"] } } },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [16],
+          "data": { "kind": "cycle", "values": ["Infinity"] },
+          "tolerance": 0,
+          "relTolerance": 0
+        }
+      }
+    },
+    {
+      "name": "int32_axis0_tiled_64x32",
+      "attrs": { "axes": [0], "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "int32",
+          "shape": [64, 32],
+          "data": {
+            "kind": "cycle",
+            "values": [-16777216, -16777245, -16777274, -16777236, -16777265, -16777227, -16777256, -16777218, -16777247, -16777276, -16777238, -16777267, -16777229, -16777258, -16777220, -16777249, -16777278, -16777240, -16777269, -16777231, -16777260, -16777222, -16777251, -16777280, -16777242, -16777271, -16777233, -16777262, -16777224, -16777253, -16777282, -16777244, -16777273, -16777235, -16777264, -16777226, -16777255, -16777217, -16777246, -16777275, -16777237, -16777266, -16777228, -16777257, -16777219, -16777248, -16777277, -16777239, -16777268, -16777230, -16777259, -16777221, -16777250, -16777279, -16777241, -16777270, -16777232, -16777261, -16777223, -16777252, -16777281, -16777243, -16777272, -16777234, -16777263, -16777225, -16777254]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [32], "tolerance": 0 } }
+    },
+    {
+      "name": "axis1",
+      "attrs": { "axes": [1], "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": { "kind": "values", "values": [1.0, 9.0, 3.0, 4.0, -1.0, 2.0, 7.0, 8.0, 0.0, 5.0, -3.0, 6.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3] } }
+    },
+    {
+      "name": "ort_axis1_nan_seed_comparison_gpu_gap",
+      "skipGpu": {
+        "category": "todo",
+        "reason": "The parallel min/max reduction routes do not yet preserve the reference's first-element NaN incumbent semantics; explicit NaN tracking is implementable in WGSL."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/reduction/reduction_ops_test.cc",
+        "test": "ReductionOpTest.ReduceMin",
+        "notes": "NaN extension: ORT seeds ReduceMin from the first reduced element, so a leading NaN remains the result while later NaNs are ignored after a finite incumbent."
+      },
+      "attrs": { "axes": [1], "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": { "kind": "values", "values": ["NaN", 1.0, -2.0, 0.0, 1.0, "NaN", -2.0, 0.0, 1.0, -2.0, "NaN", 0.0] }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [3],
+          "data": { "kind": "values", "values": ["NaN", -2.0, -2.0] },
+          "tolerance": 0,
+          "allowNaN": true
+        }
+      }
+    },
+    {
+      "name": "axis0_empty_rows_positive_infinity_identity_f32",
+      "attrs": { "axes": [0], "keepdims": 0 },
+      "inputs": { "x": { "dtype": "float32", "shape": [0, 3], "data": { "kind": "values", "values": [] } } },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [3], "data": { "kind": "cycle", "values": ["Infinity"] }, "tolerance": 0 }
+      }
+    },
+    {
+      "name": "axis1_empty_cols_positive_infinity_identity_f32",
+      "attrs": { "axes": [1], "keepdims": 0 },
+      "inputs": { "x": { "dtype": "float32", "shape": [2, 0], "data": { "kind": "values", "values": [] } } },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2], "data": { "kind": "cycle", "values": ["Infinity"] }, "tolerance": 0 }
+      }
+    },
+    {
+      "name": "axis1_keepdims",
+      "attrs": { "axes": [1], "keepdims": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": { "kind": "values", "values": [4.0, -1.0, 2.0, -1.0, 8.0, 7.0, 6.0, 5.0, -3.0, -4.0, -4.0, 2.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 1], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "axis1_all_positive_infinity_returns_positive_infinity",
+      "attrs": { "axes": [1], "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": {
+            "kind": "values",
+            "values": ["Infinity", "Infinity", "Infinity", "Infinity", "Infinity", 7.0, "Infinity", 9.0, 3.0, "Infinity", 5.0, "Infinity"]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_mixed_infinities_axis1",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/reduction/reduction_ops_test.cc",
+        "test": "ReductionOpTest.ReduceInfMin"
+      },
+      "attrs": { "axes": [1], "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [6, 2],
+          "data": {
+            "kind": "values",
+            "values": [1.0, "Infinity", "Infinity", 4.0, "Infinity", "-Infinity", "-Infinity", "Infinity", 1.0, "-Infinity", "-Infinity", 4.0]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [6], "tolerance": 0 } }
+    },
+    {
+      "name": "rank3_axis2_last_keepdims",
+      "attrs": { "axes": [2], "keepdims": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/rank3_axis2_last_keepdims_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank4_axis1_channel_no_keepdims",
+      "attrs": { "axes": [1], "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 2, 2],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/rank3_axis2_last_keepdims_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 2, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank1_axis0_scalar_output",
+      "attrs": { "axes": [0], "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [7],
+          "data": { "kind": "values", "values": [4.0, -8.0, 2.0, -8.0, 0.0, 5.0, -1.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_int32_axis0_no_keepdims",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/reduction/reduction_ops_test.cc",
+        "test": "ReductionOpTest.ReduceMin_int32",
+        "notes": "Single-axis subset adapted from ORT's multi-axis integer reduction fixture because this framework models one reduction axis per case."
+      },
+      "attrs": { "axes": [0], "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "int32",
+          "shape": [3, 2, 2],
+          "data": { "kind": "values", "values": [1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "int32_axis1_exact_above_float24",
+      "attrs": { "axes": [1], "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "int32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [16777217, 16777216, 5, -16777216, -16777217, 3] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_axis1_rank3_no_keepdims",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/reduction/reduction_ops_test.cc",
+        "test": "ReductionOpTest.ReduceMin_do_not_keepdims"
+      },
+      "attrs": { "axes": [1], "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 2, 2],
+          "data": { "kind": "values", "values": [5.0, 1.0, 20.0, 2.0, 30.0, 1.0, 40.0, 2.0, 55.0, 1.0, 60.0, 2.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_axis1_rank3_keepdims",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/reduction/reduction_ops_test.cc",
+        "test": "ReductionOpTest.ReduceMin_keepdims"
+      },
+      "attrs": { "axes": [1], "keepdims": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 2, 2],
+          "data": { "kind": "values", "values": [5.0, 1.0, 20.0, 2.0, 30.0, 1.0, 40.0, 2.0, 55.0, 1.0, 60.0, 2.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 1, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_axis0_rank1_scalar",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/reduction/reduction_ops_test.cc",
+        "test": "ReductionOpTest.ReduceMin_do_not_keepdims_2"
+      },
+      "attrs": { "axes": [0], "keepdims": 0 },
+      "inputs": { "x": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [5.0, 1.0, 20.0] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_float_axes02_rank3_keepdims1",
+      "attrs": { "axes": [0, 2], "keepdims": 1 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/reduction/reduction_ops_test.cc",
+        "test": "ReductionOpTest.ReduceMin",
+        "notes": "Direct multi-axis reduction from ORT."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 2, 2],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0] }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 2, 1],
+          "data": { "kind": "values", "values": [1.0, 3.0] },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "ort_rank0_scalar",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/reduction/reduction_ops_test.cc",
+        "test": "ReductionOpTest.ReduceMin0DTensor"
+      },
+      "inputs": { "x": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_reduce_min_do_not_keepdims_example",
+      "attrs": { "keepdims": 0, "axes": [1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 2, 2],
+          "data": { "kind": "values", "values": [5.0, 1.0, 20.0, 2.0, 30.0, 1.0, 40.0, 2.0, 55.0, 1.0, 60.0, 2.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 2] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_reduce_min_do_not_keepdims_example",
+        "notes": "The ONNX int64 axes input is materialized as this compile-time axes list."
+      }
+    },
+    {
+      "name": "onnx_backend_reduce_min_do_not_keepdims_random",
+      "attrs": { "keepdims": 0, "axes": [1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 2, 2],
+          "data": {
+            "kind": "values",
+            "values": [0.9762700796127319, 4.3037872314453125, 2.055267572402954, 0.8976636528968811, -1.5269039869308472, 2.917882204055786, -1.248255729675293, 7.835460186004639, 9.273255348205566, -2.331169605255127, 5.834500789642334, 0.577898383140564]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 2] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_reduce_min_do_not_keepdims_random",
+        "notes": "The ONNX int64 axes input is materialized as this compile-time axes list."
+      }
+    },
+    {
+      "name": "onnx_backend_reduce_min_empty_set",
+      "attrs": { "keepdims": 1, "axes": [1] },
+      "inputs": { "x": { "dtype": "float32", "shape": [2, 0, 4], "data": { "kind": "values", "values": [] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 4] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_reduce_min_empty_set",
+        "notes": "The ONNX int64 axes input is materialized as this compile-time axes list."
+      }
+    },
+    {
+      "name": "onnx_backend_reduce_min_keepdims_example",
+      "attrs": { "keepdims": 1, "axes": [1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 2, 2],
+          "data": { "kind": "values", "values": [5.0, 1.0, 20.0, 2.0, 30.0, 1.0, 40.0, 2.0, 55.0, 1.0, 60.0, 2.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 1, 2] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_reduce_min_keepdims_example",
+        "notes": "The ONNX int64 axes input is materialized as this compile-time axes list."
+      }
+    },
+    {
+      "name": "onnx_backend_reduce_min_keepdims_random",
+      "attrs": { "keepdims": 1, "axes": [1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 2, 2],
+          "data": {
+            "kind": "values",
+            "values": [0.9762700796127319, 4.3037872314453125, 2.055267572402954, 0.8976636528968811, -1.5269039869308472, 2.917882204055786, -1.248255729675293, 7.835460186004639, 9.273255348205566, -2.331169605255127, 5.834500789642334, 0.577898383140564]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 1, 2] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_reduce_min_keepdims_random",
+        "notes": "The ONNX int64 axes input is materialized as this compile-time axes list."
+      }
+    },
+    {
+      "name": "onnx_backend_reduce_min_negative_axes_keepdims_example",
+      "attrs": { "keepdims": 1, "axes": [-2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 2, 2],
+          "data": { "kind": "values", "values": [5.0, 1.0, 20.0, 2.0, 30.0, 1.0, 40.0, 2.0, 55.0, 1.0, 60.0, 2.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 1, 2] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_reduce_min_negative_axes_keepdims_example",
+        "notes": "The ONNX int64 axes input is materialized as this compile-time axes list."
+      }
+    },
+    {
+      "name": "onnx_backend_reduce_min_negative_axes_keepdims_random",
+      "attrs": { "keepdims": 1, "axes": [-2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 2, 2],
+          "data": {
+            "kind": "values",
+            "values": [0.9762700796127319, 4.3037872314453125, 2.055267572402954, 0.8976636528968811, -1.5269039869308472, 2.917882204055786, -1.248255729675293, 7.835460186004639, 9.273255348205566, -2.331169605255127, 5.834500789642334, 0.577898383140564]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 1, 2] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_reduce_min_negative_axes_keepdims_random",
+        "notes": "The ONNX int64 axes input is materialized as this compile-time axes list."
+      }
+    },
+    {
+      "name": "ort_int32_empty_axis0_keepdims_max_identity",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/reduction/reduction_ops_test.cc",
+        "test": "ReductionOpTest.ReduceMin_int32_EmptySet",
+        "notes": "The ONNX int64 axes input is materialized as this compile-time axes list."
+      },
+      "attrs": { "axes": [0], "keepdims": 1 },
+      "inputs": { "x": { "dtype": "int32", "shape": [0, 3], "data": { "kind": "values", "values": [] } } },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_empty_set_default_axes_keepdims_f32",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/reduction/reduction_ops_test.cc",
+        "test": "ReductionOpTest.empty_set_ReduceMin",
+        "notes": "An omitted axes input reduces all dimensions when noop_with_empty_axes is 0."
+      },
+      "attrs": { "keepdims": 1 },
+      "inputs": { "x": { "dtype": "float32", "shape": [2, 0, 3], "data": { "kind": "values", "values": [] } } },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 1, 1],
+          "data": { "kind": "values", "values": ["Infinity"] },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "onnx_backend_reduce_min_default_axes_keepdims_example",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_reduce_min_default_axes_keepdims_example"
+      },
+      "attrs": { "keepdims": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 2, 2],
+          "data": { "kind": "values", "values": [5.0, 1.0, 20.0, 2.0, 30.0, 1.0, 40.0, 2.0, 55.0, 1.0, 60.0, 2.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 1], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_default_axes_rank3_no_keepdims_scalar",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/reduction/reduction_ops_test.cc",
+        "test": "ReductionOpTest.ReduceMin_default_axes_do_not_keep_dims",
+        "notes": "Default axes reduce all input dimensions to a rank-0 scalar when keepdims=0."
+      },
+      "attrs": { "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 2, 2],
+          "data": { "kind": "values", "values": [5.0, 1.0, 20.0, 2.0, 30.0, 1.0, 40.0, 2.0, 55.0, 1.0, 60.0, 2.0] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] }, "tolerance": 0 }
+      }
+    },
+    {
+      "name": "onnx_backend_reduce_min_default_axes_keepdims_random",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_reduce_min_default_axes_keepdims_random"
+      },
+      "attrs": { "keepdims": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 2, 2],
+          "data": {
+            "kind": "values",
+            "values": [0.9762700796127319, 4.3037872314453125, 2.055267572402954, 0.8976636528968811, -1.5269039869308472, 2.917882204055786, -1.248255729675293, 7.835460186004639, 9.273255348205566, -2.331169605255127, 5.834500789642334, 0.577898383140564]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 1], "tolerance": 0 } }
+    },
+    {
+      "name": "subgroup_vec4_last_axis_2x256",
+      "attrs": { "axes": [-1], "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 256],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2], "tolerance": 0 } }
+    },
+    {
+      "name": "subgroup_scalar_last_axis_2x65",
+      "attrs": { "axes": [1], "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 65],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.11 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2], "tolerance": 0 } }
+    },
+    {
+      "name": "subgroup_vec4_int32_above_float24",
+      "attrs": { "axes": [1], "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": {
+            "kind": "values",
+            "values": [16777217, 16777219, 16777218, 16777220, -16777217, -16777219, -16777218, -16777220]
+          }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [16777217, -16777220] } }
+      }
+    },
+    {
+      "name": "ort_bool_axis0_keepdims0",
+      "attrs": { "axes": [0], "keepdims": 0 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/reduction/reduction_ops_test.cc",
+        "test": "test_bool_ReduceMin_9"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "bool",
+          "shape": [2, 3, 2],
+          "data": { "kind": "values", "values": [0, 0, 1, 1, 0, 1, 0, 1, 0, 1, 0, 1] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "bool", "shape": [3, 2], "data": { "kind": "values", "values": [0, 0, 0, 1, 0, 1] } }
+      }
+    },
+    {
+      "name": "ort_bool_axis2_keepdims1",
+      "attrs": { "axes": [2], "keepdims": 1 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/reduction/reduction_ops_test.cc",
+        "test": "test_bool_ReduceMin_15"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "bool",
+          "shape": [2, 3, 2],
+          "data": { "kind": "values", "values": [0, 0, 1, 1, 0, 1, 0, 1, 0, 1, 0, 1] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "bool", "shape": [2, 3, 1], "data": { "kind": "values", "values": [0, 1, 0, 0, 0, 0] } }
+      }
+    },
+    {
+      "name": "ort_bool_all_axes_keepdims1",
+      "attrs": { "keepdims": 1 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/reduction/reduction_ops_test.cc",
+        "test": "test_bool_ReduceMin_19",
+        "notes": "The fixture now uses the standard Boolean route introduced by ONNX ReduceMin-20."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "bool",
+          "shape": [2, 3, 2],
+          "data": { "kind": "values", "values": [0, 0, 1, 1, 0, 1, 0, 1, 0, 1, 0, 1] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "bool", "shape": [1, 1, 1], "data": { "kind": "values", "values": [0] }, "tolerance": 0 }
+      }
+    },
+    {
+      "name": "ort_bool_all_axes_keepdims0_scalar",
+      "attrs": { "keepdims": 0 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/reduction/reduction_ops_test.cc",
+        "test": "test_bool_ReduceMin_17",
+        "notes": "The fixture now uses the standard Boolean route introduced by ONNX ReduceMin-20."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "bool",
+          "shape": [2, 3, 2],
+          "data": { "kind": "values", "values": [0, 0, 1, 1, 0, 1, 0, 1, 0, 1, 0, 1] }
+        }
+      },
+      "outputs": { "y": { "dtype": "bool", "shape": [], "data": { "kind": "values", "values": [0] }, "tolerance": 0 } }
+    },
+    {
+      "name": "ort_noop_empty_axes_identity",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/reduction/reduction_ops_test.cc",
+        "test": "ReductionOpTest.ReduceMin_noop_axes_input_initializer_opset_18",
+        "notes": "The omitted axes input exercises empty-axes behavior."
+      },
+      "attrs": { "keepdims": 0, "noop_with_empty_axes": 1 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 2, 2], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "rank3_lastaxis_cols1024_tree_f32",
+      "attrs": { "axes": [2], "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 2, 1024],
+          "data": {
+            "kind": "cycle",
+            "values": [1.0, -2.0, 0.5, 3.25, -1.5, 2.0, -0.75, 4.0, -3.5, 1.25, 0.0, -2.25, 5.0, -4.0, 2.75, -1.0]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "rank3_lastaxis_cols1024_tree_i32",
+      "attrs": { "axes": [2], "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "int32",
+          "shape": [2, 2, 1024],
+          "data": { "kind": "cycle", "values": [3, -7, 11, -2, 5, -13, 8, -1, 19, -4, 6, -9, 2, -15, 10, -6] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 2], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "axis0_splitk_positive_infinity_identity_8192x2",
+      "attrs": { "axes": [0], "keepdims": 0 },
+      "inputs": { "x": { "dtype": "float32", "shape": [8192, 2], "data": { "kind": "cycle", "values": ["Infinity"] } } },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [2],
+          "data": { "kind": "cycle", "values": ["Infinity"] },
+          "tolerance": 0,
+          "relTolerance": 0
+        }
+      }
+    },
+    {
+      "name": "axis0_splitk_8192x48_keepdims",
+      "attrs": { "axes": [0], "keepdims": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [8192, 48],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.07, "scale": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 48], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "int32_rank4_axis0_no_keepdims",
+      "attrs": { "axes": [0], "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "int32",
+          "shape": [2, 3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [-3, 18, -25, 17, 33, -27, 42, 7, -36, -27, 22, 39, -8, 40, -42, -11, 18, -2, -43, -6, -50, 25, 5, -44, -31, 10, -6, 13, 19, 6, -26, 5, 3, 11, 14, -16, 6, 23, 28, -12, -46, -41, 37, 49, 17, 22, 33, -2, -49, 14, -34, -19, 43, -6, 42, 21, -27, -40, -15, 14, 11, -43, -27, 42, 46, -29, -15, 47, 17, 41, 47, -31, -23, -46, -31, -17, -20, 7, 23, 43, 31, -19, -48, 33, 0, -32, -29, 17, 25, 8, -3, -34, -38, 43, -41, 26, 17, -36, -30, 3, -35, 9, -15, -2, -46, -7, 22, 40, 4, -7, 37, 27, 0, -11, -47, -21, -38, 26, 29, -21]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [3, 4, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "f32_rank4_axis2_mid_no_keepdims",
+      "attrs": { "axes": [2], "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [5.248, -0.858, -1.454, -7.96, -0.025, -0.959, -1.61, 0.946, 1.263, -3.197, -2.659, -1.427, 2.069, 1.684, -3.917, -3.358, 2.211, 4.724, -0.093, -2.05, 3.287, -0.929, 2.177, 4.647, 1.89, 0.22, 2.197, -1.928, -0.534, -1.722, -0.613, -1.459, -0.557, -1.142, 0.267, 0.191, 0.889, 4.208, -4.641, 3.887, -0.712, -3.697, -0.517, 0.276, 3.203, -3.185, 0.652, 0.353, -5.052, -3.557, 1.8, 2.087, 3.263, 1.601, 1.186, 0.369, 3.627, -2.529, -0.426, 1.156, -4.732, 3.933, -2.379, -0.215, 6.47, -2.488, -1.588, 4.685, -3.241, -1.296, 1.548, 1.366, 1.781, 1.112, 4.036, 3.048, 1.786, -2.048, -2.141, -5.725, 2.062, -5.469, 2.637, 5.539, -3.185, -2.053, -1.429, 2.491, -2.59, -0.392, -1.569, -0.754, 3.874, -2.893, 0.215, 0.815, 2.576, -3.792, 3.345, 1.304, 2.623, 0.215, -4.917, -1.942, 2.453, 0.11, -0.146, 5.375, 6.606, -0.111, 5.799, -5.981, -6.147, 2.601, -0.786, 1.737, 1.55, -0.268, 2.046, 0.452]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 5] } }
+    },
+    {
+      "name": "f32_rank4_default_all_axes_no_keepdims_scalar",
+      "attrs": { "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [-1.425, 1.508, -0.089, 0.904, 2.69, 1.065, 2.7, 1.722, 2.957, -2.091, -1.578, -2.523, 1.126, -0.487, 1.827, 0.635, 0.255, 4.301, 1.213, -0.054, -1.968, 2.381, 1.906, -2.174, -0.29, 0.476, -3.278, -0.556, 2.798, -3.23, 0.982, 3.785, -1.242, -0.908, 0.435, 1.029, 0.794, -3.026, -1.528, 0.203, -0.635, 2.277, -0.644, -0.74, 3.699, -0.163, -1.961, -0.005, -0.465, 1.422, 0.906, -0.665, -0.769, -1.061, 2.701, 4.03, -0.115, 1.129, 0.719, -1.024, -1.046, -0.963, -3.848, -1.515, 0.907, 1.845, -0.415, -0.752, 1.28, 4.107, -1.857, -1.549, -0.874, 1.096, -1.708, 0.931, -1.442, 1.125, -0.63, -0.86, 1.63, -2.371, -1.537, 1.458, 1.532, 0.416, -4.374, 1.444, 0.698, -0.677, -3.891, 1.662, 3.498, 1.002, 2.661, -1.315, 0.37, -1.302, -2.241, -0.693, -1.823, -1.492, -1.542, -0.828, -0.584, 0.254, -2.495, 0.634, -3.637, 0.111, 0.162, -4.64, -0.682, 2.18, 0.901, -0.989, -1.093, 0.584, 0.308, 1.635]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [] } }
+    },
+    {
+      "name": "rank4_multi_axis_12_keepdims",
+      "attrs": { "axes": [1, 2], "keepdims": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 2, 2],
+          "data": {
+            "kind": "values",
+            "values": [0.5, -1.0, 2.0, -0.25, 1.5, 0.75, -2.0, 1.0, 0.125, -0.5, 3.0, -1.5, 0.25, 2.5, -0.75, 1.25, -3.0, 0.5, 2.0, -1.0, 0.75, -0.25, 1.5, -2.5]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 1, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "uint32_axis0_tiled_64x32",
+      "attrs": { "axes": [0], "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "uint32",
+          "shape": [64, 32],
+          "data": {
+            "kind": "cycle",
+            "values": [16777216, 16777245, 16777274, 16777236, 16777265, 16777227, 16777256, 16777218, 16777247, 16777276, 16777238, 16777267, 16777229, 16777258, 16777220, 16777249, 16777278, 16777240, 16777269, 16777231, 16777260, 16777222, 16777251, 16777280, 16777242, 16777271, 16777233, 16777262, 16777224, 16777253, 16777282, 16777244, 16777273, 16777235, 16777264, 16777226, 16777255, 16777217, 16777246, 16777275, 16777237, 16777266, 16777228, 16777257, 16777219, 16777248, 16777277, 16777239, 16777268, 16777230, 16777259, 16777221, 16777250, 16777279, 16777241, 16777270, 16777232, 16777261, 16777223, 16777252, 16777281, 16777243, 16777272, 16777234, 16777263, 16777225, 16777254]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "uint32", "shape": [32], "tolerance": 0 } }
+    },
+    {
+      "name": "rank3_all_axes_keepdims0_small_mod4_miss_f32",
+      "attrs": { "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 5, 7],
+          "data": {
+            "kind": "cycle",
+            "values": [-2.5, 1.0, 0.0, -3.5, 4.0, -1.5, 2.5, -0.5, -4.0, 0.5, -1.0, 3.0, -5.5, 1.5, -2.0, 2.0, -3.0, 5.0, -0.25, 0.25]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [] } }
+    },
+    {
+      "name": "rank3_axis2_dedicated_variant_correctness",
+      "attrs": { "axes": [2], "keepdims": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4, 3, 5],
+          "data": {
+            "kind": "values",
+            "values": [1.0, -2.0, 3.0, -4.0, 5.0, -1.0, 2.0, -3.0, 4.0, -5.0, 0.5, -0.5, 1.5, -1.5, 2.5, -2.5, 3.5, -3.5, 4.5, -4.5, 0.25, -0.25, 0.75, -0.75, 1.25, -1.25, 1.75, -1.75, 2.25, -2.25, 0.1, -0.1, 0.3, -0.3, 0.6, -0.6, 0.9, -0.9, 1.2, -1.2, 2.0, -2.0, 3.0, -3.0, 4.0, -4.0, 5.0, -5.0, 6.0, -6.0, 7.0, -7.0, 8.0, -8.0, 9.0, -9.0, 10.0, -10.0, 11.0, -11.0]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4, 3, 1] } }
+    },
+    {
+      "name": "axis1_rows32768_cols30_subgroup_gate_cliff_f32",
+      "attrs": { "axes": [1], "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [32768, 30],
+          "data": { "kind": "fillFloat32", "sinStep": 0.031, "cosStep": 0.017, "scale": 1.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [32768] } }
+    },
+    {
+      "name": "axis0_cols15_i32_serial_uncoalesced",
+      "attrs": { "axes": [0], "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "int32",
+          "shape": [256, 15],
+          "data": {
+            "kind": "cycle",
+            "values": [2147483647, -2147483648, 0, 1, -1, 1000000, -1000000, 16777217, -16777217, 42, -42, 100, -100, 7, -7]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [15] } }
+    },
+    {
+      "name": "all_axes_flat_fullreduce_101x103_nonmul4_keepdims",
+      "provenance": {
+        "notes": "numel=10403 (%4=3, >=8192) full-reduce: previously fell to the single-lane reduce-serial-axis kernel; now stays on all_axes_flat (f32 binding, manual vec4 groups + 3-element scalar tail). Cross-checks the flat path against the serial all_axes_keepdims reference — a dropped tail would change the result."
+      },
+      "attrs": { "keepdims": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [101, 103],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.027, "scale": 0.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1], "tolerance": 0.01, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "contiguous_suffix_axes12_parallel",
+      "provenance": {
+        "notes": "Contiguous axes {1,2} exercise the shared cooperative suffix reduction instead of one serial lane per output."
+      },
+      "attrs": { "axes": [1, 2], "keepdims": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 16, 16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.07, "scale": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 1, 1], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "axis_split_rank3_axis1_2x8192x4",
+      "attrs": { "axes": [1], "keepdims": 0 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [2, 8192, 4], "data": { "kind": "cycle", "values": ["Infinity"] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "cycle", "values": ["Infinity"] },
+          "tolerance": 0,
+          "relTolerance": 0
+        }
+      }
+    },
+    {
+      "name": "f16_axis_split_tiled_narrow_2x8192x4",
+      "attrs": { "axes": [1], "keepdims": 0 },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [2, 8192, 4], "data": { "kind": "cycle", "values": ["Infinity"] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float16",
+          "shape": [2, 4],
+          "data": { "kind": "cycle", "values": ["Infinity"] },
+          "tolerance": 0,
+          "relTolerance": 0
+        }
+      }
+    },
+    {
+      "name": "axis_split_rank3_axis1_wide_2x8192x32",
+      "attrs": { "axes": [1], "keepdims": 0 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [2, 8192, 32], "data": { "kind": "cycle", "values": ["Infinity"] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [2, 32],
+          "data": { "kind": "cycle", "values": ["Infinity"] },
+          "tolerance": 0,
+          "relTolerance": 0
+        }
+      }
+    },
+    {
+      "name": "f16_axis_split_wide_2x8192x32",
+      "attrs": { "axes": [1], "keepdims": 0 },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [2, 8192, 32], "data": { "kind": "cycle", "values": ["Infinity"] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float16",
+          "shape": [2, 32],
+          "data": { "kind": "cycle", "values": ["Infinity"] },
+          "tolerance": 0,
+          "relTolerance": 0
+        }
+      }
+    },
+    {
+      "name": "f16_rank3_axis1_serial",
+      "attrs": { "axes": [1], "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [3, 2, 2],
+          "data": { "kind": "values", "values": [5.0, 1.0, 20.0, 2.0, 30.0, 1.0, 40.0, 2.0, 55.0, 1.0, 60.0, 2.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [3, 2], "tolerance": 0.01 } }
+    },
+    {
+      "name": "f16_last_axis_serial_fallback",
+      "attrs": { "axes": [1], "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 65],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.11 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2], "tolerance": 0.01 } }
+    },
+    {
+      "name": "f16_all_axes",
+      "attrs": { "keepdims": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [32, 32, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.027, "scale": 0.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 1, 1], "tolerance": 0.01 } }
+    },
+    {
+      "name": "f16_axis0_splitk_8192x8",
+      "attrs": { "axes": [0], "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [8192, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.11, "scale": 4.0 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [8], "tolerance": 0.05, "relTolerance": 0.002 } }
+    },
+    {
+      "name": "f16_last_axis_vec4_8x1024",
+      "attrs": { "axes": [1], "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [8, 1024],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.11, "scale": 4.0 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [8], "tolerance": 0.05, "relTolerance": 0.002 } }
+    },
+    {
+      "name": "f16_last_axis_scalar_8x1023",
+      "attrs": { "axes": [1], "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [8, 1023],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.11, "scale": 4.0 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [8], "tolerance": 0.05, "relTolerance": 0.002 } }
+    },
+    {
+      "name": "f16_all_axes_flat_65543",
+      "attrs": { "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [65543],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.11, "scale": 4.0 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [], "tolerance": 0.05, "relTolerance": 0.002 } }
+    },
+    {
+      "name": "f16_suffix_vec4_4x8x128",
+      "attrs": { "axes": [1, 2], "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [4, 8, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.11, "scale": 4.0 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [4], "tolerance": 0.05, "relTolerance": 0.002 } }
+    },
+    {
+      "name": "f16_suffix_scalar_4x7x37",
+      "attrs": { "axes": [1, 2], "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [4, 7, 37],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.11, "scale": 4.0 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [4], "tolerance": 0.05, "relTolerance": 0.002 } }
+    },
+    {
+      "name": "f16_axis0_tilecols_4096x64",
+      "attrs": { "axes": [0], "keepdims": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [4096, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.11, "scale": 4.0 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [64], "tolerance": 0.05, "relTolerance": 0.002 } }
+    },
+    {
+      "name": "noop_empty_axes_under_x4_copy_floor",
+      "provenance": {
+        "notes": "Identity-shaped noop_with_empty_axes reduction below the four-element floor of the vec4 degenerate copy, which otherwise outranks the noop route on every identity case (the ORT [1,2,2] fixture has exactly four elements)."
+      },
+      "attrs": { "keepdims": 0, "noop_with_empty_axes": 1 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 1, 3], "data": { "kind": "values", "values": [3.0, -1.0, 2.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int8_axis1_extrema",
+      "provenance": {
+        "source": "ONNX Runtime CPUExecutionProvider",
+        "notes": "Covers the signed narrow-integer reduction route."
+      },
+      "attrs": { "axes": [1], "keepdims": 0 },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [2, 3], "data": { "kind": "values", "values": [-128, -5, 127, 12, -3, 11] } }
+      },
+      "outputs": {
+        "y": { "dtype": "int8", "shape": [2], "tolerance": 0, "data": { "kind": "values", "values": [-128, -3] } }
+      }
+    },
+    {
+      "name": "ort_uint8_axis0_extrema",
+      "provenance": {
+        "source": "ONNX Runtime CPUExecutionProvider",
+        "notes": "Covers the unsigned narrow-integer reduction route."
+      },
+      "attrs": { "axes": [0], "keepdims": 0 },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [2, 3], "data": { "kind": "values", "values": [1, 200, 255, 250, 3, 4] } }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [3], "tolerance": 0, "data": { "kind": "values", "values": [1, 3, 4] } }
+      }
+    },
+    {
+      "name": "ort_int8_empty_axis1_highest_identity",
+      "provenance": {
+        "source": "ONNX Runtime CPUExecutionProvider",
+        "notes": "An empty INT8 minimum is the logical dtype maximum (127), not the physical i32 storage maximum."
+      },
+      "attrs": { "axes": [1], "keepdims": 1 },
+      "inputs": { "x": { "dtype": "int8", "shape": [2, 0, 3], "data": { "kind": "values", "values": [] } } },
+      "outputs": {
+        "y": {
+          "dtype": "int8",
+          "shape": [2, 1, 3],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [127, 127, 127, 127, 127, 127] }
+        }
+      }
+    },
+    {
+      "name": "ort_uint8_empty_axis1_highest_identity",
+      "provenance": {
+        "source": "ONNX Runtime CPUExecutionProvider",
+        "notes": "An empty UINT8 minimum is the logical dtype maximum (255), not the physical u32 storage maximum."
+      },
+      "attrs": { "axes": [1], "keepdims": 1 },
+      "inputs": { "x": { "dtype": "uint8", "shape": [2, 0, 3], "data": { "kind": "values", "values": [] } } },
+      "outputs": {
+        "y": {
+          "dtype": "uint8",
+          "shape": [2, 1, 3],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [255, 255, 255, 255, 255, 255] }
+        }
+      }
+    },
+    {
+      "name": "onnx_v20_bool_empty_axis1_true_identity",
+      "provenance": {
+        "source": "https://onnx.ai/onnx/operators/onnx__ReduceMin.html",
+        "notes": "An empty Boolean minimum is true, the logical dtype maximum."
+      },
+      "attrs": { "axes": [1], "keepdims": 1 },
+      "inputs": { "x": { "dtype": "bool", "shape": [2, 0, 3], "data": { "kind": "values", "values": [] } } },
+      "outputs": {
+        "y": {
+          "dtype": "bool",
+          "shape": [2, 1, 3],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [1, 1, 1, 1, 1, 1] }
+        }
+      }
+    }
+  ]
+}