sync 2e7068faf55e

Browse files

Files changed (14) hide show

README.md +79 -0
build/webgpu/bench.json +292 -0
build/webgpu/datamove-scatter-row.wgsl.jinja +50 -0
build/webgpu/manifest.json +591 -0
build/webgpu/metadata.json +26 -0
build/webgpu/scatter-f16-f32-convert.wgsl.jinja +17 -0
build/webgpu/scatter-flat-copy.wgsl.jinja +25 -0
build/webgpu/scatter-narrow-wrap.wgsl.jinja +32 -0
build/webgpu/scatter-nd-reduction-atomic.wgsl.jinja +327 -0
build/webgpu/scatter-nd-reduction-f16-packed-atomic.wgsl.jinja +124 -0
build/webgpu/scatter-nd-reduction-slice.wgsl.jinja +149 -0
build/webgpu/scatter-nd.wgsl.jinja +94 -0
build/webgpu/scatter-scalar-add.wgsl.jinja +63 -0
build/webgpu/test.json +1326 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.ScatterND
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 18
+## Description
+Copies `data` and then scatters values from `updates` into the copy at positions specified by `indices`, producing an output with the same shape as `data`. Each entry in `indices` is a k-tuple that addresses either a single element (when k equals the rank of `data`) or a slice; an optional `reduction` attribute replaces the default overwrite with an elementwise `add`, `mul`, `max`, or `min` accumulation.
+See the [ONNX `ScatterND` spec](https://onnx.ai/onnx/operators/onnx__ScatterND.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | WebGPU storage | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- | --- |
+| `data` | `data` | `T` | runtime-selected; narrow integers and bool use 32-bit slots | — | — | Source tensor of rank r >= 1 that is copied to form the base of the output. | required |
+| `indices` | `indices` | `I` | `int32` | — | — | Logical int64 index tensor of rank q >= 1 whose last dimension gives signed k-tuples that address into `data`; WebGPU stores it as int32. | required |
+| `updates` | `updates` | `T` | runtime-selected; narrow integers and bool use 32-bit slots | — | — | Values to scatter into the output, of rank q + r - k - 1, matching the leading (q-1) index dimensions and the trailing (r-k) data slice dimensions. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `output` | `output` | `T` | same as `data` | same as `data` | Result tensor of rank r with the same shape as `data`, with selected positions overwritten or reduced by `updates`. | required |
+## Attributes
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `reduction` | `"none"` | Reduction to apply when writing each update value: `none` (overwrite), `add`, `mul`, `max`, or `min`. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16`, `int32`, `uint32`, `int8`, `uint8`, `int16`, `bool` |
+| `I` | `int64` |
+## 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-scatter-row.wgsl.jinja`](build/webgpu/datamove-scatter-row.wgsl.jinja)
+- [`scatter-f16-f32-convert.wgsl.jinja`](build/webgpu/scatter-f16-f32-convert.wgsl.jinja)
+- [`scatter-flat-copy.wgsl.jinja`](build/webgpu/scatter-flat-copy.wgsl.jinja)
+- [`scatter-narrow-wrap.wgsl.jinja`](build/webgpu/scatter-narrow-wrap.wgsl.jinja)
+- [`scatter-nd-reduction-atomic.wgsl.jinja`](build/webgpu/scatter-nd-reduction-atomic.wgsl.jinja)
+- [`scatter-nd-reduction-f16-packed-atomic.wgsl.jinja`](build/webgpu/scatter-nd-reduction-f16-packed-atomic.wgsl.jinja)
+- [`scatter-nd-reduction-slice.wgsl.jinja`](build/webgpu/scatter-nd-reduction-slice.wgsl.jinja)
+- [`scatter-nd.wgsl.jinja`](build/webgpu/scatter-nd.wgsl.jinja)
+- [`scatter-scalar-add.wgsl.jinja`](build/webgpu/scatter-scalar-add.wgsl.jinja)
+## Use with `@huggingface/kernels`
+The loader derives every required output's shape and logical dtype from the manifest contract and this call.
+It then allocates the result tensors automatically.
+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.ScatterND", { version: 1 });
+const { output } = await kernel({
+  data: { data: dataData, shape: [2, 2] },
+  indices: { data: indicesData, shape: [2, 1] },
+  updates: { data: updatesData, shape: [2, 2] },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,292 @@

+{
+  "op": "ai.onnx.ScatterND",
+  "cases": [
+    {
+      "name": "row_updates",
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [4096, 768] },
+        "indices": { "dtype": "int32", "shape": [128, 1] },
+        "updates": { "dtype": "float32", "shape": [128, 768] }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4096, 768] } }
+    },
+    {
+      "name": "row_updates_reduction_add",
+      "preset": "smoke",
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [4096, 768], "dist": "normal", "seed": 1 },
+        "indices": { "dtype": "int32", "shape": [512, 1], "dist": "linearMod", "mod": 4096, "step": 9 },
+        "updates": { "dtype": "float32", "shape": [512, 768], "dist": "normal", "seed": 2 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4096, 768], "dist": "empty" } },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "4 * (2 * numel(shapes.data) + 3 * numel(shapes.updates)) + 4 * numel(shapes.indices)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "row_updates_reduction_max",
+      "preset": "smoke",
+      "attrs": { "reduction": "max" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [4096, 768], "dist": "normal", "seed": 1 },
+        "indices": { "dtype": "int32", "shape": [512, 1], "dist": "linearMod", "mod": 4096, "step": 9 },
+        "updates": { "dtype": "float32", "shape": [512, 768], "dist": "normal", "seed": 2 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4096, 768], "dist": "empty" } }
+    },
+    {
+      "name": "scalar_updates_reduction_add_contended_1m",
+      "preset": "stress",
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [1024], "dist": "normal", "seed": 1 },
+        "indices": { "dtype": "int32", "shape": [1048576, 1], "dist": "randint", "min": 0, "max": 1024, "seed": 3 },
+        "updates": { "dtype": "float32", "shape": [1048576], "dist": "normal", "seed": 2 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1024], "dist": "empty" } }
+    },
+    {
+      "name": "scalar_updates_reduction_add_single_slot_1m",
+      "preset": "stress",
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [1024], "dist": "normal", "seed": 1 },
+        "indices": { "dtype": "int32", "shape": [1048576, 1], "dist": "constant", "value": 0 },
+        "updates": { "dtype": "float32", "shape": [1048576], "dist": "normal", "seed": 2 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1024], "dist": "empty" } }
+    },
+    {
+      "name": "row_updates_dense_8192x1024",
+      "preset": "model",
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [8192, 1024], "dist": "normal", "seed": 701 },
+        "indices": { "dtype": "int32", "shape": [8192, 1], "dist": "linearMod", "mod": 8192, "step": 7 },
+        "updates": { "dtype": "float32", "shape": [8192, 1024], "dist": "normal", "seed": 703 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [8192, 1024], "dist": "empty" } },
+      "bench": {
+        "primary": false,
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "4 * (2 * numel(shapes.data) + numel(shapes.updates)) + 4 * numel(shapes.indices)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "f16_reduction_add_single_lane_serial_cliff",
+      "preset": "smoke",
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": { "dtype": "float16", "shape": [4096, 768], "dist": "normal", "seed": 1 },
+        "indices": { "dtype": "int32", "shape": [512, 1], "dist": "linearMod", "mod": 4096, "step": 9 },
+        "updates": { "dtype": "float16", "shape": [512, 768], "dist": "normal", "seed": 2 }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [4096, 768], "dist": "empty" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "2 * (2 * numel(shapes.data) + 3 * numel(shapes.updates)) + 4 * numel(shapes.indices)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "f32_reduction_add_atomic_parallel_healthy",
+      "preset": "smoke",
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [4096, 768], "dist": "normal", "seed": 1 },
+        "indices": { "dtype": "int32", "shape": [512, 1], "dist": "linearMod", "mod": 4096, "step": 9 },
+        "updates": { "dtype": "float32", "shape": [512, 768], "dist": "normal", "seed": 2 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4096, 768], "dist": "empty" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "4 * (2 * numel(shapes.data) + 3 * numel(shapes.updates)) + 4 * numel(shapes.indices)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "int8_reduction_add_atomic_path_healthy",
+      "preset": "smoke",
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": { "dtype": "int8", "shape": [1024], "dist": "constant", "value": 0 },
+        "indices": { "dtype": "int32", "shape": [262144, 1], "dist": "randint", "min": 0, "max": 1024, "seed": 3 },
+        "updates": { "dtype": "int8", "shape": [262144], "dist": "constant", "value": 1 }
+      },
+      "outputs": { "output": { "dtype": "int8", "shape": [1024], "dist": "empty" } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "numel(shapes.data) + 3 * numel(shapes.updates) + 4 * numel(shapes.indices)" }
+        ]
+      }
+    },
+    {
+      "name": "f16_row_update_none_unaligned770_batch0_cliff",
+      "preset": "smoke",
+      "inputs": {
+        "data": { "dtype": "float16", "shape": [4096, 770], "dist": "normal", "seed": 11 },
+        "indices": { "dtype": "int32", "shape": [128, 1], "dist": "linearMod", "mod": 4096, "step": 9 },
+        "updates": { "dtype": "float16", "shape": [128, 770], "dist": "normal", "seed": 12 }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [4096, 770], "dist": "empty" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "2 * (2 * numel(shapes.data) + numel(shapes.updates)) + 4 * numel(shapes.indices)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "f16_row_update_none_aligned768_rowfast_healthy",
+      "preset": "smoke",
+      "inputs": {
+        "data": { "dtype": "float16", "shape": [4096, 768], "dist": "normal", "seed": 11 },
+        "indices": { "dtype": "int32", "shape": [128, 1], "dist": "linearMod", "mod": 4096, "step": 9 },
+        "updates": { "dtype": "float16", "shape": [128, 768], "dist": "normal", "seed": 12 }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [4096, 768], "dist": "empty" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "2 * (2 * numel(shapes.data) + numel(shapes.updates)) + 4 * numel(shapes.indices)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "f16_reduction_add_scalar_updates_rank1_single_thread",
+      "preset": "stress",
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": { "dtype": "float16", "shape": [16384], "dist": "normal", "seed": 20 },
+        "indices": { "dtype": "int32", "shape": [16384, 1], "dist": "randint", "min": 0, "max": 16384, "seed": 21 },
+        "updates": { "dtype": "float16", "shape": [16384], "dist": "normal", "seed": 22 }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [16384], "dist": "empty" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "2 * (2 * numel(shapes.data) + 3 * numel(shapes.updates)) + 4 * numel(shapes.indices)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "f32_sparse_copy_dominates_tiny_scatter",
+      "preset": "stress",
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [307200, 13], "dist": "normal", "seed": 30 },
+        "indices": { "dtype": "int32", "shape": [8, 1], "dist": "randint", "min": 0, "max": 307200, "seed": 31 },
+        "updates": { "dtype": "float32", "shape": [8, 13], "dist": "normal", "seed": 32 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [307200, 13], "dist": "empty" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "4 * (2 * numel(shapes.data) + numel(shapes.updates)) + 4 * numel(shapes.indices)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "scalar-loss-accumulation-add-4m-single-output-pathology",
+      "preset": "model",
+      "attrs": { "reduction": "add" },
+      "vars": { "count": 4194304, "dtype": "float32" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [1], "dist": "constant", "value": 0 },
+        "indices": { "dtype": "int32", "shape": [4194304, 1], "dist": "constant", "value": 0 },
+        "updates": { "dtype": "float32", "shape": [4194304], "dist": "normal", "seed": 33, "scale": 0.001 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1], "dist": "empty" } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.dtype) + 4) + 2 * dtypeBytes(args.dtype)" }
+        ]
+      }
+    },
+    {
+      "name": "padding-token-hotspot-add-2m-vocab131072",
+      "preset": "stress",
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [131072], "dist": "constant", "value": 0 },
+        "indices": { "dtype": "int32", "shape": [2097152, 1], "dist": "constant", "value": 0 },
+        "updates": { "dtype": "float32", "shape": [2097152], "dist": "constant", "value": 1 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [131072], "dist": "empty" } },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "4 * (2 * numel(shapes.data) + 3 * numel(shapes.updates)) + 4 * numel(shapes.indices)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "f16-sparse-router-gradient-65536-into16384",
+      "preset": "stress",
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": { "dtype": "float16", "shape": [16384], "dist": "constant", "value": 0 },
+        "indices": { "dtype": "int32", "shape": [65536, 1], "dist": "randint", "min": 0, "max": 16384, "seed": 41 },
+        "updates": { "dtype": "float16", "shape": [65536], "dist": "normal", "seed": 42, "scale": 0.01 }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [16384], "dist": "empty" } },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "2 * (2 * numel(shapes.data) + 3 * numel(shapes.updates)) + 4 * numel(shapes.indices)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "f16-router-max-262144-into16384-single-column",
+      "preset": "stress",
+      "provenance": {
+        "notes": "Large f16 max reduction that exercises the serialized reduction_slice route; the atomic-add control distinguishes the non-add reduction gap."
+      },
+      "attrs": { "reduction": "max" },
+      "inputs": {
+        "data": { "dtype": "float16", "shape": [16384], "dist": "constant", "value": 0 },
+        "indices": { "dtype": "int32", "shape": [262144, 1], "dist": "randint", "min": 0, "max": 16384, "seed": 43 },
+        "updates": { "dtype": "float16", "shape": [262144], "dist": "normal", "seed": 44 }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [16384], "dist": "empty" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "2 * (2 * numel(shapes.data) + 3 * numel(shapes.updates)) + 4 * numel(shapes.indices)"
+          }
+        ]
+      }
+    }
+  ]
+}

build/webgpu/datamove-scatter-row.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,50 @@

+// One workgroup per update row computes its normalized output offset and
+// vec4-copies the contiguous inner slice from updates[row] into
+// output[scatter_index]. This path handles indexDepth == 1 (scatter along
+// data axis 0) with an inner size divisible by four. It performs no reduction;
+// duplicate indices are last-write-wins, matching ONNX's nondeterministic
+// duplicate-index semantics.
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
+const NUM_INDICES: u32 = {{ source.numIndices }}u;
+const AXIS_DIM: u32 = {{ source.axisDim }}u;
+const ROW_VECS: u32 = {{ source.rowVecs }}u;
+const INVALID_INDEX: u32 = 0xffffffffu;
+fn normalize_index(raw: {{ indexScalar }}, dim: u32) -> u32 {
+  let dim_i = i32(dim);
+  var index = raw;
+  if (index < 0) {
+    index = index + dim_i;
+  }
+  if (index < 0 || index >= dim_i) {
+    return INVALID_INDEX;
+  }
+  return u32(index);
+}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(
+  @builtin(workgroup_id) wid: vec3<u32>,
+  @builtin(num_workgroups) nwg: vec3<u32>,
+  @builtin(local_invocation_id) lid: vec3<u32>
+) {
+  for (var row = wid.x; row < NUM_INDICES; row += nwg.x) {
+    let scatter_index = normalize_index(indices[row], AXIS_DIM);
+    // Out-of-range index: skip the row (ONNX leaves it to the copied `data`).
+    // The continue is uniform across the workgroup (all threads share `row`).
+    if (scatter_index == INVALID_INDEX) {
+      continue;
+    }
+    let src_base = row * ROW_VECS;
+    let dst_base = scatter_index * ROW_VECS;
+    for (var j = lid.x; j < ROW_VECS; j += WG) {
+      output[dst_base + j] = updates[src_base + j];
+    }
+  }
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,591 @@

+{
+  "domain": "ai.onnx",
+  "name": "ScatterND",
+  "sinceVersion": 18,
+  "description": "Copies `data` and then scatters values from `updates` into the copy at positions specified by `indices`, producing an output with the same shape as `data`. Each entry in `indices` is a k-tuple that addresses either a single element (when k equals the rank of `data`) or a slice; an optional `reduction` attribute replaces the default overwrite with an elementwise `add`, `mul`, `max`, or `min` accumulation.",
+  "inputs": [
+    {
+      "role": "data",
+      "dtype": "T",
+      "description": "Source tensor of rank r >= 1 that is copied to form the base of the output."
+    },
+    {
+      "role": "indices",
+      "dtype": "I",
+      "description": "Logical int64 index tensor of rank q >= 1 whose last dimension gives signed k-tuples that address into `data`; WebGPU stores it as int32."
+    },
+    {
+      "role": "updates",
+      "dtype": "T",
+      "description": "Values to scatter into the output, of rank q + r - k - 1, matching the leading (q-1) index dimensions and the trailing (r-k) data slice dimensions."
+    }
+  ],
+  "outputs": [
+    {
+      "role": "output",
+      "dtype": "T",
+      "rank": "ranks.data",
+      "description": "Result tensor of rank r with the same shape as `data`, with selected positions overwritten or reduced by `updates`.",
+      "shape": "shapes.data"
+    }
+  ],
+  "attributes": { "reduction": "none" },
+  "attributeDescriptions": {
+    "reduction": "Reduction to apply when writing each update value: `none` (overwrite), `add`, `mul`, `max`, or `min`."
+  },
+  "attributeConstraints": { "reduction": { "values": ["none", "add", "mul", "max", "min"] } },
+  "typeConstraints": {
+    "T": ["float32", "float16", "int32", "uint32", "int8", "uint8", "int16", "bool"],
+    "I": ["int64"]
+  },
+  "args": {
+    "data": { "kind": "tensor", "semantic": "data", "role": "input" },
+    "indices": { "kind": "tensor", "semantic": "indices", "role": "indices", "dtype": "int32", "narrowing": "checked" },
+    "updates": { "kind": "tensor", "semantic": "updates", "role": "updates" },
+    "output": { "kind": "tensor", "semantic": "output", "role": "output" }
+  },
+  "tunables": {
+    "WORKGROUP_SIZE": 256,
+    "F16_ATOMIC_MIN_UPDATES": 4096,
+    "SCALAR_ADD_MIN_UPDATES": 256,
+    "PREAGG_ITEMS_PER_THREAD": 8,
+    "PREAGG_MIN_DUPLICATION": 8,
+    "PREAGG_SUBGROUP_ROUNDS": 2,
+    "SCALAR_ADD_ITEMS_PER_THREAD": 16
+  },
+  "derive": {
+    "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
+    "foldedDispatchCapacity": "device.limits.maxComputeWorkgroupsPerDimension * device.limits.maxComputeWorkgroupsPerDimension",
+    "indexDepth": "dim(shapes.indices, ranks.indices - 1)",
+    "narrowIntData": "tensorDtypes.data == \"int8\" or tensorDtypes.data == \"uint8\" or tensorDtypes.data == \"int16\"",
+    "validLayout": "(dtypes.I == \"u32\" or dtypes.I == \"i32\") and ranks.data == ranks.output and ranks.data >= 1 and ranks.indices >= 2 and indexDepth <= ranks.data and ranks.updates == ranks.indices - 1 + ranks.data - indexDepth and numel(shapes.data) == numel(shapes.output)",
+    "workgroupLimit": "min(tunables.WORKGROUP_SIZE, deviceWorkgroupCap)",
+    "workgroupSize": "max(1, pow2ceil(workgroupLimit + 1) / 2)",
+    "foldedThreadCapacity": "foldedDispatchCapacity * workgroupSize",
+    "updateDispatchFits": "numel(shapes.updates) <= foldedThreadCapacity",
+    "outputDispatchFits": "numel(shapes.output) <= foldedThreadCapacity",
+    "f16AccumBytes": "numel(shapes.output) * dtypeBytes(\"uint32\")",
+    "f16AccumFits": "f16AccumBytes <= device.limits.maxStorageBufferBindingSize and f16AccumBytes <= device.limits.maxBufferSize",
+    "preaggStorageFits": "workgroupSize * dtypeBytes(\"float32\") + 16 <= device.limits.maxComputeWorkgroupStorageSize",
+    "scalarAddStorageFits": "workgroupSize * dtypeBytes(\"float32\") <= device.limits.maxComputeWorkgroupStorageSize"
+  },
+  "bindingSets": {
+    "copy": [
+      {
+        "name": "src",
+        "arg": "data",
+        "semantic": "data",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$copyElementType"
+      },
+      {
+        "name": "dst",
+        "arg": "output",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$copyElementType"
+      }
+    ],
+    "rowVec4": [
+      {
+        "name": "indices",
+        "arg": "indices",
+        "semantic": "indices",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$indexScalar"
+      },
+      {
+        "name": "updates",
+        "arg": "updates",
+        "semantic": "updates",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "output",
+        "arg": "output",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$vectorScalar"
+      }
+    ],
+    "scatter": [
+      {
+        "name": "indices",
+        "arg": "indices",
+        "semantic": "indices",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$indexScalar"
+      },
+      {
+        "name": "updates",
+        "arg": "updates",
+        "semantic": "updates",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$updateElement"
+      },
+      {
+        "name": "output",
+        "arg": "output",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputElement"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.updates)" }] }
+      }
+    ],
+    "f16WideInit": [
+      {
+        "name": "narrow",
+        "arg": "data",
+        "semantic": "data",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f16"
+      },
+      { "name": "wide", "semantic": "accum", "buffer": { "type": "storage" }, "elementType": "u32" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.output)" }] }
+      }
+    ],
+    "f16AtomicAccumulate": [
+      {
+        "name": "indices",
+        "arg": "indices",
+        "semantic": "indices",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$indexScalar"
+      },
+      {
+        "name": "updates",
+        "arg": "updates",
+        "semantic": "updates",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f16"
+      },
+      { "name": "output", "semantic": "accum", "buffer": { "type": "storage" }, "elementType": "atomic<u32>" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.updates)" }] }
+      }
+    ],
+    "f16NarrowFinalize": [
+      { "name": "wide", "semantic": "accum", "buffer": { "type": "read-only-storage" }, "elementType": "u32" },
+      { "name": "narrow", "arg": "output", "semantic": "output", "buffer": { "type": "storage" }, "elementType": "f16" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.output)" }] }
+      }
+    ],
+    "narrowWrap": [
+      { "name": "values", "arg": "output", "semantic": "output", "buffer": { "type": "storage" }, "elementType": "u32" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.output)" }] }
+      }
+    ]
+  },
+  "variants": [
+    {
+      "id": "reduction_f16_atomic_f32",
+      "priority": 35,
+      "when": ["(attrs.reduction == \"add\" or attrs.reduction == \"min\" or attrs.reduction == \"max\")", "dtypes.T == \"f16\"", "validLayout", "indexDepth == ranks.data", "numel(shapes.updates) >= tunables.F16_ATOMIC_MIN_UPDATES", "f16Ok(dtypes.T)", "f16AccumFits", "outputDispatchFits", "updateDispatchFits"],
+      "constants": {
+        "scalar": "\"f32\"",
+        "indexScalar": "dtypes.I",
+        "usesF16": true,
+        "reductionOp": "attrs.reduction",
+        "useSubgroups": "device.features.has(\"subgroups\") and attrs.reduction == \"add\"",
+        "useWorkgroupPreagg": false,
+        "preaggRounds": "tunables.PREAGG_SUBGROUP_ROUNDS",
+        "preaggItemsPerThread": "tunables.PREAGG_ITEMS_PER_THREAD"
+      },
+      "intermediates": [{ "id": "accum", "dtype": "uint32", "shape": "[numel(shapes.output)]" }],
+      "passes": [
+        {
+          "id": "init",
+          "name": "ScatterND.F16InitF32",
+          "source": { "shader": "scatter-f16-f32-convert.wgsl.jinja", "inputs": { "toWide": true } },
+          "bindings": "f16WideInit",
+          "dispatch": { "threads": "numel(shapes.output)", "workgroupSize": "workgroupSize" }
+        },
+        {
+          "id": "scatter",
+          "name": "ScatterND.F16AtomicF32",
+          "subgroupCollectivesWidth": "portable",
+          "source": {
+            "shader": "scatter-nd-reduction-atomic.wgsl.jinja",
+            "inputs": {
+              "dataShape": "shapes.data",
+              "indicesShape": "shapes.indices",
+              "updatesShape": "shapes.updates",
+              "dataRank": "ranks.data",
+              "indicesRank": "ranks.indices",
+              "updatesRank": "ranks.updates",
+              "indexDepth": "indexDepth"
+            }
+          },
+          "bindings": "f16AtomicAccumulate",
+          "dispatch": { "threads": "numel(shapes.updates)", "workgroupSize": "workgroupSize" }
+        },
+        {
+          "id": "finalize",
+          "name": "ScatterND.F16FinalizeF32",
+          "source": { "shader": "scatter-f16-f32-convert.wgsl.jinja", "inputs": { "toWide": false } },
+          "bindings": "f16NarrowFinalize",
+          "dispatch": { "threads": "numel(shapes.output)", "workgroupSize": "workgroupSize" }
+        }
+      ]
+    },
+    {
+      "id": "scalar_output_add_parallel_reduce",
+      "priority": 30,
+      "when": ["attrs.reduction == \"add\"", "validLayout", "dtypes.T == \"f32\"", "ranks.data == 1", "ranks.output == 1", "numel(shapes.data) == 1", "numel(shapes.output) == 1", "ranks.indices == 2", "dim(shapes.indices, 1) == 1", "ranks.updates == 1", "dim(shapes.updates, 0) == dim(shapes.indices, 0)", "numel(shapes.updates) >= tunables.SCALAR_ADD_MIN_UPDATES", "scalarAddStorageFits"],
+      "constants": {
+        "scalar": "dtypes.T",
+        "indexScalar": "dtypes.I",
+        "copyElementType": "dtypes.T",
+        "updateElement": "dtypes.T",
+        "outputElement": "\"atomic<u32>\""
+      },
+      "passes": [
+        {
+          "id": "copy",
+          "name": "ScatterND.Copy",
+          "source": { "shader": "scatter-flat-copy.wgsl.jinja", "inputs": { "count": "numel(shapes.output)" } },
+          "bindings": "copy",
+          "dispatch": { "gridStride": "numel(shapes.output)", "workgroupSize": "workgroupSize" }
+        },
+        {
+          "id": "reduce",
+          "name": "ScatterND.ScalarAddParallelReduce",
+          "shader": "scatter-scalar-add.wgsl.jinja",
+          "bindings": "scatter",
+          "dispatch": {
+            "x": "min(ceilDiv(numel(shapes.updates), workgroupSize * tunables.SCALAR_ADD_ITEMS_PER_THREAD), device.limits.maxComputeWorkgroupsPerDimension)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "reduction_narrow_wrap_atomic",
+      "priority": 26,
+      "when": ["(attrs.reduction == \"add\" or attrs.reduction == \"mul\")", "validLayout", "narrowIntData", "updateDispatchFits", "outputDispatchFits"],
+      "constants": {
+        "scalar": "\"u32\"",
+        "indexScalar": "dtypes.I",
+        "usesF16": false,
+        "reductionOp": "attrs.reduction",
+        "useSubgroups": "device.features.has(\"subgroups\")",
+        "useWorkgroupPreagg": false,
+        "preaggRounds": "tunables.PREAGG_SUBGROUP_ROUNDS",
+        "preaggItemsPerThread": "tunables.PREAGG_ITEMS_PER_THREAD",
+        "workgroupSize": "workgroupSize",
+        "narrowBits": "16 if tensorDtypes.data == \"int16\" else 8",
+        "narrowSigned": "tensorDtypes.data == \"int8\" or tensorDtypes.data == \"int16\"",
+        "copyElementType": "(\"vec4<\" ~ dtypes.T ~ \">\") if numel(shapes.output) % 4 == 0 else dtypes.T",
+        "updateElement": "\"u32\"",
+        "outputElement": "\"atomic<u32>\""
+      },
+      "passes": [
+        {
+          "id": "copy",
+          "name": "ScatterND.Copy",
+          "source": {
+            "shader": "scatter-flat-copy.wgsl.jinja",
+            "inputs": { "count": "numel(shapes.output) / 4 if numel(shapes.output) % 4 == 0 else numel(shapes.output)" }
+          },
+          "bindings": "copy",
+          "dispatch": {
+            "gridStride": "numel(shapes.output) / 4 if numel(shapes.output) % 4 == 0 else numel(shapes.output)",
+            "workgroupSize": "workgroupSize"
+          }
+        },
+        {
+          "id": "scatter",
+          "name": "ScatterND.NarrowWrapAtomic",
+          "subgroupCollectivesWidth": "portable",
+          "source": {
+            "shader": "scatter-nd-reduction-atomic.wgsl.jinja",
+            "inputs": {
+              "dataShape": "shapes.data",
+              "indicesShape": "shapes.indices",
+              "updatesShape": "shapes.updates",
+              "dataRank": "ranks.data",
+              "indicesRank": "ranks.indices",
+              "updatesRank": "ranks.updates",
+              "indexDepth": "indexDepth"
+            }
+          },
+          "bindings": "scatter",
+          "dispatch": { "threads": "numel(shapes.updates)", "workgroupSize": "workgroupSize" }
+        },
+        {
+          "id": "narrow",
+          "name": "ScatterND.NarrowWrapFinalize",
+          "shader": "scatter-narrow-wrap.wgsl.jinja",
+          "bindings": "narrowWrap",
+          "dispatch": { "threads": "numel(shapes.output)", "workgroupSize": "workgroupSize" }
+        }
+      ]
+    },
+    {
+      "id": "reduction_f16_packed_atomic_add",
+      "priority": 36,
+      "when": ["attrs.reduction == \"add\"", "dtypes.T == \"f16\"", "validLayout", "indexDepth < ranks.data", "numel(shapes.output) % 2 == 0", "numel(shapes.updates) >= tunables.F16_ATOMIC_MIN_UPDATES", "f16Ok(dtypes.T)", "outputDispatchFits", "updateDispatchFits"],
+      "constants": {
+        "indexScalar": "dtypes.I",
+        "usesF16": true,
+        "copyElementType": "(\"vec4<\" ~ dtypes.T ~ \">\") if numel(shapes.output) % 4 == 0 else dtypes.T",
+        "workgroupSize": "workgroupSize",
+        "updateElement": "\"f16\"",
+        "outputElement": "\"atomic<u32>\""
+      },
+      "passes": [
+        {
+          "id": "copy",
+          "name": "ScatterND.Copy",
+          "source": {
+            "shader": "scatter-flat-copy.wgsl.jinja",
+            "inputs": { "count": "numel(shapes.output) / 4 if numel(shapes.output) % 4 == 0 else numel(shapes.output)" }
+          },
+          "bindings": "copy",
+          "dispatch": {
+            "gridStride": "numel(shapes.output) / 4 if numel(shapes.output) % 4 == 0 else numel(shapes.output)",
+            "workgroupSize": "workgroupSize"
+          }
+        },
+        {
+          "id": "scatter",
+          "name": "ScatterND.F16PackedAtomicAdd",
+          "source": {
+            "shader": "scatter-nd-reduction-f16-packed-atomic.wgsl.jinja",
+            "inputs": {
+              "dataShape": "shapes.data",
+              "indicesShape": "shapes.indices",
+              "updatesShape": "shapes.updates",
+              "dataRank": "ranks.data",
+              "indicesRank": "ranks.indices",
+              "updatesRank": "ranks.updates",
+              "indexDepth": "indexDepth"
+            }
+          },
+          "bindings": "scatter",
+          "dispatch": { "threads": "numel(shapes.updates)", "workgroupSize": "workgroupSize" }
+        }
+      ]
+    },
+    {
+      "id": "reduction_atomic",
+      "priority": 25,
+      "when": ["(attrs.reduction == \"add\" or attrs.reduction == \"mul\" or attrs.reduction == \"min\" or attrs.reduction == \"max\")", "validLayout", "(tensorDtypes.data == \"float32\" or tensorDtypes.data == \"int32\" or tensorDtypes.data == \"uint32\" or (narrowIntData and (attrs.reduction == \"min\" or attrs.reduction == \"max\")))", "updateDispatchFits"],
+      "constants": {
+        "scalar": "dtypes.T",
+        "indexScalar": "dtypes.I",
+        "usesF16": "dtypes.T == \"f16\"",
+        "reductionOp": "attrs.reduction",
+        "useSubgroups": "device.features.has(\"subgroups\") and (attrs.reduction == \"add\" or attrs.reduction == \"mul\")",
+        "useWorkgroupPreagg": "attrs.reduction == \"add\" and dtypes.T == \"f32\" and indexDepth == ranks.data and numel(shapes.updates) >= numel(shapes.output) * tunables.PREAGG_MIN_DUPLICATION and preaggStorageFits",
+        "preaggRounds": "tunables.PREAGG_SUBGROUP_ROUNDS",
+        "workgroupSize": "workgroupSize",
+        "preaggItemsPerThread": "tunables.PREAGG_ITEMS_PER_THREAD",
+        "copyElementType": "(\"vec4<\" ~ dtypes.T ~ \">\") if numel(shapes.output) % 4 == 0 else dtypes.T",
+        "updateElement": "dtypes.T",
+        "outputElement": "\"atomic<i32>\" if dtypes.T == \"i32\" else \"atomic<u32>\""
+      },
+      "passes": [
+        {
+          "id": "copy",
+          "name": "ScatterND.Copy",
+          "source": {
+            "shader": "scatter-flat-copy.wgsl.jinja",
+            "inputs": { "count": "numel(shapes.output) / 4 if numel(shapes.output) % 4 == 0 else numel(shapes.output)" }
+          },
+          "bindings": "copy",
+          "dispatch": {
+            "gridStride": "numel(shapes.output) / 4 if numel(shapes.output) % 4 == 0 else numel(shapes.output)",
+            "workgroupSize": "workgroupSize"
+          }
+        },
+        {
+          "id": "scatter",
+          "name": "ScatterND.ReductionAtomic",
+          "subgroupCollectivesWidth": "portable",
+          "source": {
+            "shader": "scatter-nd-reduction-atomic.wgsl.jinja",
+            "inputs": {
+              "dataShape": "shapes.data",
+              "indicesShape": "shapes.indices",
+              "updatesShape": "shapes.updates",
+              "dataRank": "ranks.data",
+              "indicesRank": "ranks.indices",
+              "updatesRank": "ranks.updates",
+              "indexDepth": "indexDepth"
+            }
+          },
+          "bindings": "scatter",
+          "dispatch": {
+            "threads": "ceilDiv(numel(shapes.updates), tunables.PREAGG_ITEMS_PER_THREAD) if attrs.reduction == \"add\" and dtypes.T == \"f32\" and indexDepth == ranks.data and numel(shapes.updates) >= numel(shapes.output) * tunables.PREAGG_MIN_DUPLICATION and preaggStorageFits else numel(shapes.updates)",
+            "workgroupSize": "workgroupSize"
+          }
+        }
+      ]
+    },
+    {
+      "id": "reduction_slice",
+      "priority": 21,
+      "supersededBy": ["scalar_output_add_parallel_reduce"],
+      "when": ["(attrs.reduction == \"add\" or attrs.reduction == \"mul\" or attrs.reduction == \"min\" or attrs.reduction == \"max\")", "validLayout", "f16Ok(dtypes.T)"],
+      "constants": {
+        "scalar": "dtypes.T",
+        "indexScalar": "dtypes.I",
+        "usesF16": "dtypes.T == \"f16\"",
+        "reductionOp": "attrs.reduction",
+        "narrowBits": "8 if tensorDtypes.data == \"int8\" or tensorDtypes.data == \"uint8\" else (16 if tensorDtypes.data == \"int16\" else 0)",
+        "narrowSigned": "tensorDtypes.data == \"int8\" or tensorDtypes.data == \"int16\"",
+        "copyElementType": "(\"vec4<\" ~ dtypes.T ~ \">\") if numel(shapes.output) % 4 == 0 else dtypes.T",
+        "workgroupSize": "workgroupSize",
+        "updateElement": "dtypes.T",
+        "outputElement": "dtypes.T"
+      },
+      "passes": [
+        {
+          "id": "copy",
+          "name": "ScatterND.Copy",
+          "source": {
+            "shader": "scatter-flat-copy.wgsl.jinja",
+            "inputs": { "count": "numel(shapes.output) / 4 if numel(shapes.output) % 4 == 0 else numel(shapes.output)" }
+          },
+          "bindings": "copy",
+          "dispatch": {
+            "gridStride": "numel(shapes.output) / 4 if numel(shapes.output) % 4 == 0 else numel(shapes.output)",
+            "workgroupSize": "workgroupSize"
+          }
+        },
+        {
+          "id": "scatter",
+          "name": "ScatterND.ReductionSlice",
+          "source": {
+            "shader": "scatter-nd-reduction-slice.wgsl.jinja",
+            "inputs": {
+              "dataShape": "shapes.data",
+              "indicesShape": "shapes.indices",
+              "updatesShape": "shapes.updates",
+              "dataRank": "ranks.data",
+              "indicesRank": "ranks.indices",
+              "updatesRank": "ranks.updates",
+              "indexDepth": "indexDepth"
+            }
+          },
+          "bindings": "scatter",
+          "dispatch": {
+            "gridStride": "numel(shapes.updates) * indexDepth / numel(shapes.indices)",
+            "workgroupSize": "workgroupSize"
+          }
+        }
+      ]
+    },
+    {
+      "id": "row_fast_vec4",
+      "priority": 10,
+      "when": ["attrs.reduction == \"none\"", "validLayout", "f16Ok(dtypes.T)", "indexDepth == 1", "inner(shapes.data, 0) % 4 == 0", "numel(shapes.indices) > 0"],
+      "constants": {
+        "scalar": "dtypes.T",
+        "indexScalar": "dtypes.I",
+        "usesF16": "dtypes.T == \"f16\"",
+        "copyElementType": "(\"vec4<\" ~ dtypes.T ~ \">\") if numel(shapes.output) % 4 == 0 else dtypes.T",
+        "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\""
+      },
+      "passes": [
+        {
+          "id": "copy",
+          "name": "ScatterND.Copy",
+          "source": {
+            "shader": "scatter-flat-copy.wgsl.jinja",
+            "inputs": { "count": "numel(shapes.output) / 4 if numel(shapes.output) % 4 == 0 else numel(shapes.output)" }
+          },
+          "bindings": "copy",
+          "dispatch": {
+            "gridStride": "numel(shapes.output) / 4 if numel(shapes.output) % 4 == 0 else numel(shapes.output)",
+            "workgroupSize": "workgroupSize"
+          },
+          "viewAlias": [{ "input": "src", "output": "dst" }]
+        },
+        {
+          "id": "scatter",
+          "name": "ScatterND.Scatter",
+          "source": {
+            "shader": "datamove-scatter-row.wgsl.jinja",
+            "inputs": {
+              "numIndices": "numel(shapes.indices)",
+              "axisDim": "dim(shapes.data, 0)",
+              "rowVecs": "inner(shapes.data, 0) / 4"
+            }
+          },
+          "bindings": "rowVec4",
+          "dispatch": { "gridStride": "numel(shapes.indices)" }
+        }
+      ]
+    },
+    {
+      "id": "batch0",
+      "priority": 0,
+      "when": ["attrs.reduction == \"none\"", "validLayout", "f16Ok(dtypes.T)", "updateDispatchFits"],
+      "constants": {
+        "scalar": "dtypes.T",
+        "indexScalar": "dtypes.I",
+        "usesF16": "dtypes.T == \"f16\"",
+        "copyElementType": "(\"vec4<\" ~ dtypes.T ~ \">\") if numel(shapes.output) % 4 == 0 else dtypes.T",
+        "workgroupSize": "workgroupSize",
+        "updateElement": "dtypes.T",
+        "outputElement": "dtypes.T"
+      },
+      "passes": [
+        {
+          "id": "copy",
+          "name": "ScatterND.Copy",
+          "source": {
+            "shader": "scatter-flat-copy.wgsl.jinja",
+            "inputs": { "count": "numel(shapes.output) / 4 if numel(shapes.output) % 4 == 0 else numel(shapes.output)" }
+          },
+          "bindings": "copy",
+          "dispatch": {
+            "gridStride": "numel(shapes.output) / 4 if numel(shapes.output) % 4 == 0 else numel(shapes.output)",
+            "workgroupSize": "workgroupSize"
+          },
+          "viewAlias": [{ "input": "src", "output": "dst" }]
+        },
+        {
+          "id": "scatter",
+          "name": "ScatterND.Scatter",
+          "source": {
+            "shader": "scatter-nd.wgsl.jinja",
+            "inputs": {
+              "dataShape": "shapes.data",
+              "indicesShape": "shapes.indices",
+              "updatesShape": "shapes.updates",
+              "dataRank": "ranks.data",
+              "indicesRank": "ranks.indices",
+              "updatesRank": "ranks.updates",
+              "indexDepth": "indexDepth"
+            }
+          },
+          "bindings": "scatter",
+          "dispatch": { "threads": "numel(shapes.updates)", "workgroupSize": "workgroupSize" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,26 @@

+{
+  "name": "ai.onnx.ScatterND",
+  "id": "_ai_onnx_scatternd_webgpu_3220727",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "GrY1TM5HIEywYLyjniLuwycr0czxfwYG/ac8PZ+IE9E=",
+      "datamove-scatter-row.wgsl.jinja": "nfTmOf2O7JCLw3aMpwqmEnUiLlraLVBklQLxkqhVUNg=",
+      "manifest.json": "ISRWqj8bAyNt1z7/35pzxXtQvaPRq4fRwm92neVyjgo=",
+      "scatter-f16-f32-convert.wgsl.jinja": "6ZiaoKGXqY9I7aiy2QY76t8N9f8tQTVed2GoDwBD9/o=",
+      "scatter-flat-copy.wgsl.jinja": "j95cywCwuvbfDJJxf9+SLt2Li3CvLllAesHpPykppFk=",
+      "scatter-narrow-wrap.wgsl.jinja": "skKLW560sgq8K0xFclFx8eEyk5tSDoPy11OgZhEQUfI=",
+      "scatter-nd-reduction-atomic.wgsl.jinja": "EEOc7xPQZcqWGQGN9sjbiHFaDFFm+FZqj0y3CsEtDog=",
+      "scatter-nd-reduction-f16-packed-atomic.wgsl.jinja": "B8AVBcP3R7qpgdMIZvwc1wFbZhanlnEib0kXY+4Ovj0=",
+      "scatter-nd-reduction-slice.wgsl.jinja": "N9MUXlG46j0ZO144klWLlT2tq8U8yeSkI8UIih6y2Q8=",
+      "scatter-nd.wgsl.jinja": "PJRmnpsn+b33y/t3VrpSqg1VuPA3Y1VZU1O9tl0eXWc=",
+      "scatter-scalar-add.wgsl.jinja": "GUpLZ/XGwu51w1vl/GEjAKdL3+iiWzSrcWqnb6qBwyQ=",
+      "test.json": "8D1gbC8s0AyoDrU6AG/OMKPiSSOe8uJnnZXOxy8uxwE="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.ScatterND" }
+}

build/webgpu/scatter-f16-f32-convert.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,17 @@

+enable f16;
+{{ env.wgsl.resourceDeclarations }}
+// Narrow/wide conversion around f16 scatter reductions. `wide` stores f32
+// payload bits in u32 words so the middle pass can bind the same allocation as
+// atomic<u32> and perform portable float CAS operations.
+@compute @workgroup_size({{ workgroupSize }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>,
+        @builtin(num_workgroups) nwg: vec3<u32>) {
+  let i = gid.x + gid.y * nwg.x * {{ workgroupSize }}u;
+  if (i >= params.count) { return; }
+{% if source.toWide %}
+  wide[i] = bitcast<u32>(f32(narrow[i]));
+{% else %}
+  narrow[i] = f16(bitcast<f32>(wide[i]));
+{% endif %}
+}

build/webgpu/scatter-flat-copy.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,25 @@

+// Flat src -> dst copy for the scatter-family "materialize output from data"
+// pass. The element width is owned by the binding element type: when
+// numel(output) % 4 == 0 the bindings rebind as vec4<T> (128-bit loads).
+// Tail-safe modes keep scalar bindings but copy a compile-time group of adjacent
+// elements per invocation, guarding only the final partial group. A grid-stride
+// loop lets the dispatch stay within maxComputeWorkgroupsPerDimension.
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+{% set wg = workgroupSize if workgroupSize is defined else tunables.WORKGROUP_SIZE %}
+const COUNT: u32 = {{ source.count }}u;
+const WG: u32 = {{ wg }}u;
+@compute @workgroup_size({{ wg }})
+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 < COUNT; i += stride) {
+    dst[i] = src[i];
+  }
+}

build/webgpu/scatter-narrow-wrap.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,32 @@

+{{ env.wgsl.resourceDeclarations }}
+{% set mask = 255 if narrowBits == 8 else 65535 %}
+{% set sign_bit = 128 if narrowBits == 8 else 32768 %}
+// Truncates an atomically accumulated scatter result back to its narrow integer
+// width, in place, after an add or mul reduction over int8/uint8/int16.
+//
+// Those types have no native atomic, so the reduction runs on the raw
+// two's-complement bits held in u32. u32 add and mul are modulo 2^32 and
+// truncation to the narrow width commutes with both, so narrowing once here is
+// exactly the sequence of per-update narrowings the serial kernel performs --
+// whatever order the atomics landed in.
+//
+// The values buffer stays bound as u32 for this pass, so the sign extension is
+// written as a bit pattern rather than an i32 conversion.
+const MASK: u32 = {{ mask }}u;
+{% if narrowSigned %}
+const SIGN_BIT: u32 = {{ sign_bit }}u;
+{% endif %}
+@compute @workgroup_size({{ workgroupSize }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>,
+        @builtin(num_workgroups) nwg: vec3<u32>) {
+  let i = gid.x + gid.y * nwg.x * {{ workgroupSize }}u;
+  if (i >= params.count) { return; }
+  let low = values[i] & MASK;
+{% if narrowSigned %}
+  values[i] = select(low, low | ~MASK, low >= SIGN_BIT);
+{% else %}
+  values[i] = low;
+{% endif %}
+}

build/webgpu/scatter-nd-reduction-atomic.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,327 @@

+// Parallel ScatterND reduction: one thread per update and an atomic
+// read-modify-write at the destination. Native integer atomics cover add/min/max;
+// integer mul uses compare-exchange, f32 uses a CAS bitcast loop, and eligible
+// f16 reductions use that f32 path through an upcast scratch before narrowing.
+//
+// Subgroup lanes targeting one destination may pre-aggregate CAS-based add/mul
+// updates before an elected lane issues the atomic. Uniform match-loop control
+// ensures every lane reaches each collective. Floating-point grouping and atomic
+// update order can both change rounding. Without subgroups, an all-equal-key f32
+// add workgroup reduces locally and emits one CAS; mixed keys use per-update CAS.
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{% if useSubgroups %}
+enable subgroups;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const WG: u32 = {{ workgroupSize }}u;
+const INVALID_INDEX: u32 = 0xffffffffu;
+{% if useWorkgroupPreagg %}
+var<workgroup> wg_first_key: u32;
+var<workgroup> wg_mismatch: atomic<u32>;
+var<workgroup> wg_all_same: u32;
+var<workgroup> wg_vals: array<f32, {{ workgroupSize }}>;
+{% endif %}
+fn normalize_index(raw: {{ indexScalar }}, dim: u32) -> u32 {
+  let dim_i = i32(dim);
+  var index = raw;
+  if (index < 0) {
+    index = index + dim_i;
+  }
+  if (index < 0 || index >= dim_i) {
+    return INVALID_INDEX;
+  }
+  return u32(index);
+}
+fn indices_base(update_index: u32) -> u32 {
+  var rem = update_index;
+  var offset = 0u;
+{% for axis in range(source.updatesRank) %}
+{% set update_stride = namespace(value=1) %}
+{% for j in range(axis + 1, source.updatesRank) %}
+{% set update_stride.value = update_stride.value * source.updatesShape[j] %}
+{% endfor %}
+{% set update_stride_divisor = update_stride.value if update_stride.value != 0 else 1 %}
+{% if axis < source.indicesRank - 1 %}
+  let coord{{ axis }} = rem / {{ update_stride_divisor }}u;
+{% endif %}
+  rem = rem % {{ update_stride_divisor }}u;
+{% if axis < source.indicesRank - 1 %}
+{% set index_stride = namespace(value=source.indexDepth) %}
+{% for j in range(axis + 1, source.indicesRank - 1) %}
+{% set index_stride.value = index_stride.value * source.indicesShape[j] %}
+{% endfor %}
+  offset = offset + coord{{ axis }} * {{ index_stride.value }}u;
+{% endif %}
+{% endfor %}
+  return offset;
+}
+fn output_offset(update_index: u32) -> u32 {
+  let index_base = indices_base(update_index);
+  var rem = update_index;
+{% for axis in range(source.updatesRank) %}
+{% set update_stride = namespace(value=1) %}
+{% for j in range(axis + 1, source.updatesRank) %}
+{% set update_stride.value = update_stride.value * source.updatesShape[j] %}
+{% endfor %}
+{% set update_stride_divisor = update_stride.value if update_stride.value != 0 else 1 %}
+{% set suffixBegin = source.indicesRank - 1 %}
+{% set suffixEnd = suffixBegin + source.dataRank - source.indexDepth %}
+{% if axis >= suffixBegin and axis < suffixEnd %}
+  let ucoord{{ axis }} = rem / {{ update_stride_divisor }}u;
+{% endif %}
+  rem = rem % {{ update_stride_divisor }}u;
+{% endfor %}
+  var offset = 0u;
+{% for axis in range(source.dataRank) %}
+{% set data_stride = namespace(value=1) %}
+{% for j in range(axis + 1, source.dataRank) %}
+{% set data_stride.value = data_stride.value * source.dataShape[j] %}
+{% endfor %}
+{% if axis < source.indexDepth %}
+  let gathered{{ axis }} = normalize_index(indices[index_base + {{ axis }}u], {{ source.dataShape[axis] }}u);
+  if (gathered{{ axis }} == INVALID_INDEX) {
+    return INVALID_INDEX;
+  }
+  offset = offset + gathered{{ axis }} * {{ data_stride.value }}u;
+{% else %}
+{% set suffix_axis = (source.indicesRank - 1) + (axis - source.indexDepth) %}
+  offset = offset + ucoord{{ suffix_axis }} * {{ data_stride.value }}u;
+{% endif %}
+{% endfor %}
+  return offset;
+}
+{% if scalar == "f32" %}
+// CAS-bitcast read-modify-write: output is bound as atomic<u32>; the f32
+// payload is bitcast in-shader because WGSL has no float atomics.
+fn apply_update(dst: u32, update: f32) {
+  var old_bits = atomicLoad(&output[dst]);
+  loop {
+    let current = bitcast<f32>(old_bits);
+{% if reductionOp == "add" %}
+    let combined = current + update;
+{% elif reductionOp == "mul" %}
+    let combined = current * update;
+{% elif reductionOp == "min" %}
+    let combined = min(current, update);
+{% else %}
+    let combined = max(current, update);
+{% endif %}
+    let result = atomicCompareExchangeWeak(&output[dst], old_bits, bitcast<u32>(combined));
+    if (result.exchanged) {
+      break;
+    }
+    old_bits = result.old_value;
+  }
+}
+{% elif reductionOp == "mul" %}
+// No atomicMul in WGSL: CAS loop on the native integer payload (wrapping
+// integer multiply, same as the serial shader's current * update).
+fn apply_update(dst: u32, update: {{ scalar }}) {
+  var old = atomicLoad(&output[dst]);
+  loop {
+    let result = atomicCompareExchangeWeak(&output[dst], old, old * update);
+    if (result.exchanged) {
+      break;
+    }
+    old = result.old_value;
+  }
+}
+{% else %}
+fn apply_update(dst: u32, update: {{ scalar }}) {
+{% if reductionOp == "add" %}
+  atomicAdd(&output[dst], update);
+{% elif reductionOp == "min" %}
+  atomicMin(&output[dst], update);
+{% else %}
+  atomicMax(&output[dst], update);
+{% endif %}
+}
+{% endif %}
+{% if useSubgroups %}
+// Index of the lowest set bit across a 128-bit subgroup ballot. The ballot is
+// subgroup-uniform, so every lane computes the same leader lane id.
+fn ballot_first_lane(b: vec4<u32>) -> u32 {
+  if (b.x != 0u) { return firstTrailingBit(b.x); }
+  if (b.y != 0u) { return 32u + firstTrailingBit(b.y); }
+  if (b.z != 0u) { return 64u + firstTrailingBit(b.z); }
+  return 96u + firstTrailingBit(b.w);
+}
+// How many distinct output slots per subgroup are pre-aggregated before falling
+// back to per-lane atomics for the remainder. A compile-time constant so the
+// round loop has a uniform trip count (no data-dependent break): the subgroup
+// builtins inside it are reached by every lane on every round, i.e. in
+// subgroup-uniform control flow. ROUNDS covers the common contention shapes
+// (few hot slots per subgroup, incl. the all-same-slot worst case in round 0);
+// rare highly-scattered subgroups simply spill to the per-lane path.
+const PREAGG_ROUNDS: u32 = {{ preaggRounds }}u;
+// Warp-aggregated atomic: in each round, all still-pending lanes elect the
+// smallest target slot (subgroupMin), the lanes hitting that slot combine their
+// updates locally with subgroupAdd/subgroupMul, and one leader lane issues a
+// single atomic for it. After PREAGG_ROUNDS rounds any lane still pending does a
+// direct per-lane atomic. `valid` lanes whose dst is INVALID_INDEX never
+// participate and never write.
+//
+// The algorithm is width-independent and makes no assumption about the subgroup
+// width (subgroupMin/Add/Mul/Ballot reduce over whatever the compiled width is,
+// and subgroupBallot is always a 128-bit vec4 regardless). `sg_size` only gates
+// the cheap degenerate case: a 1-wide subgroup has no peers to combine, so it
+// goes straight to the per-lane atomic and skips the reduction overhead.
+fn aggregate_update(dst_in: u32, update_in: {{ scalar }}, valid_in: bool, sg_lane: u32, sg_size: u32) {
+  var pending = valid_in && (dst_in != INVALID_INDEX);
+  if (sg_size <= 1u) {
+    if (pending) {
+      apply_update(dst_in, update_in);
+    }
+    return;
+  }
+  for (var r = 0u; r < PREAGG_ROUNDS; r = r + 1u) {
+    // Elect a candidate dst: the smallest dst among still-pending lanes.
+    // Non-pending lanes contribute the sentinel so they cannot win.
+    let key = select(INVALID_INDEX, dst_in, pending);
+    let cand = subgroupMin(key);
+    let matched = pending && (dst_in == cand);
+{% if reductionOp == "add" %}
+    let contrib = select({{ scalar }}(0), update_in, matched);
+    let total = subgroupAdd(contrib);
+{% else %}
+    let contrib = select({{ scalar }}(1), update_in, matched);
+    let total = subgroupMul(contrib);
+{% endif %}
+    // Elect the lowest matched lane as leader to
+    // issue the single atomic for `cand`. Ballot + firstTrailingBit is uniform.
+    // cand == INVALID_INDEX means no lane is pending this round -> skip.
+    let leader = ballot_first_lane(subgroupBallot(matched));
+    if (matched && sg_lane == leader && cand != INVALID_INDEX) {
+      apply_update(cand, total);
+    }
+    if (matched) {
+      pending = false;
+    }
+  }
+  // Spill: lanes that still have work after the fixed rounds fall back to a
+  // direct per-lane atomic (correct, just unaggregated for that remainder).
+  if (pending) {
+    apply_update(dst_in, update_in);
+  }
+}
+{% endif %}
+@compute @workgroup_size(WG)
+fn main(
+  @builtin(global_invocation_id) gid: vec3<u32>,
+  @builtin(num_workgroups) nwg: vec3<u32>
+{%- if useSubgroups %},
+  @builtin(subgroup_invocation_id) sg_lane: u32,
+  @builtin(subgroup_size) sg_size: u32
+{%- endif %}
+{% if useWorkgroupPreagg %},
+  @builtin(local_invocation_id) lid: vec3<u32>
+{%- endif %}
+) {
+{% if useWorkgroupPreagg %}
+  // Coalesce several updates per lane, then test whether every live lane maps
+  // to the same valid destination. Equal-key workgroups reduce locally; mixed
+  // keys replay through the subgroup or direct-atomic path after the barriers.
+  const ITEMS: u32 = {{ preaggItemsPerThread }}u;
+  let thread = gid.y * (nwg.x * WG) + gid.x;
+  let totalThreads = nwg.x * nwg.y * WG;
+  if (lid.x == 0u) {
+    let firstIndex = thread - lid.x;
+    var firstKey = INVALID_INDEX;
+    if (firstIndex < params.count) { firstKey = output_offset(firstIndex); }
+    wg_first_key = firstKey;
+    atomicStore(&wg_mismatch, 0u);
+  }
+  let firstKey = workgroupUniformLoad(&wg_first_key);
+  var localSum = 0.0;
+  for (var item = 0u; item < ITEMS; item += 1u) {
+    let i = thread + item * totalThreads;
+    if (i < params.count) {
+      let dst = output_offset(i);
+      if (dst != INVALID_INDEX) {
+        if (dst == firstKey) {
+          localSum += updates[i];
+        } else {
+          atomicStore(&wg_mismatch, 1u);
+        }
+      }
+    }
+  }
+  wg_vals[lid.x] = localSum;
+  workgroupBarrier();
+  if (lid.x == 0u) {
+    wg_all_same = select(0u, 1u, firstKey != INVALID_INDEX && atomicLoad(&wg_mismatch) == 0u);
+  }
+  let allSame = workgroupUniformLoad(&wg_all_same) != 0u;
+  if (allSame) {
+    var offset = WG / 2u;
+    loop {
+      if (lid.x < offset) {
+        wg_vals[lid.x] += wg_vals[lid.x + offset];
+      }
+      workgroupBarrier();
+      if (offset <= 1u) { break; }
+      offset /= 2u;
+    }
+    if (lid.x == 0u) {
+      apply_update(firstKey, wg_vals[0]);
+    }
+  } else {
+    for (var item = 0u; item < ITEMS; item += 1u) {
+      let i = thread + item * totalThreads;
+      let inRange = i < params.count;
+      var dst = INVALID_INDEX;
+      var upd = 0.0;
+      if (inRange) {
+        dst = output_offset(i);
+        upd = {{ scalar }}(updates[i]);
+      }
+{% if useSubgroups %}
+      aggregate_update(dst, upd, inRange, sg_lane, sg_size);
+{% else %}
+      if (inRange && dst != INVALID_INDEX) {
+        apply_update(dst, upd);
+      }
+{% endif %}
+    }
+  }
+{% elif useSubgroups %}
+  // The folded 2D dispatch assigns one index to every lane. Every lane calls
+  // `aggregate_update` once, keeping its subgroup builtins uniform; lanes past
+  // the count join as non-pending and neither read nor write.
+  let i = gid.y * (nwg.x * WG) + gid.x;
+  let in_range = i < params.count;
+  var dst = INVALID_INDEX;
+  var upd = {{ scalar }}(0);
+  if (in_range) {
+    dst = output_offset(i);
+    upd = {{ scalar }}(updates[i]);
+  }
+  aggregate_update(dst, upd, in_range, sg_lane, sg_size);
+{% else %}
+  // Without subgroups, grid-stride over the folded 2D dispatch and issue one
+  // atomic for each valid update.
+  let row = nwg.x * WG;
+  let stride = row * nwg.y;
+  for (var i = gid.y * row + gid.x; i < params.count; i += stride) {
+    let dst = output_offset(i);
+    if (dst == INVALID_INDEX) {
+      continue;
+    }
+    apply_update(dst, {{ scalar }}(updates[i]));
+  }
+{% endif %}
+}

build/webgpu/scatter-nd-reduction-f16-packed-atomic.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,124 @@

+enable f16;
+{{ env.wgsl.resourceDeclarations }}
+// Direct f16 reduction via CAS on the packed pair containing each output
+// element. Updating either half preserves the other half from the successful
+// observed word, so adjacent columns and duplicate indices remain race-free
+// without widening the output into scratch storage.
+// Add only: the one case that selects this shader gates on
+// attrs.reduction == "add", so min/max never reach it. The general
+// atomic and slice routes carry the other reductions.
+//
+// The same case requires numel(updates) >= F16_ATOMIC_MIN_UPDATES, so no
+// update extent is zero and every stride below divides.
+const WG: u32 = {{ workgroupSize }}u;
+const INVALID_INDEX: u32 = 0xffffffffu;
+fn normalize_index(raw: {{ indexScalar }}, dim: u32) -> u32 {
+  let dim_i = i32(dim);
+  var index = raw;
+  if (index < 0) {
+    index = index + dim_i;
+  }
+  if (index < 0 || index >= dim_i) {
+    return INVALID_INDEX;
+  }
+  return u32(index);
+}
+fn indices_base(update_index: u32) -> u32 {
+  var rem = update_index;
+  var offset = 0u;
+{% for axis in range(source.updatesRank) %}
+{% set update_stride = namespace(value=1) %}
+{% for j in range(axis + 1, source.updatesRank) %}
+{% set update_stride.value = update_stride.value * source.updatesShape[j] %}
+{% endfor %}
+{% set update_stride_divisor = update_stride.value %}
+{% if axis < source.indicesRank - 1 %}
+  let coord{{ axis }} = rem / {{ update_stride_divisor }}u;
+{% endif %}
+  rem %= {{ update_stride_divisor }}u;
+{% if axis < source.indicesRank - 1 %}
+{% set index_stride = namespace(value=source.indexDepth) %}
+{% for j in range(axis + 1, source.indicesRank - 1) %}
+{% set index_stride.value = index_stride.value * source.indicesShape[j] %}
+{% endfor %}
+  offset += coord{{ axis }} * {{ index_stride.value }}u;
+{% endif %}
+{% endfor %}
+  return offset;
+}
+fn output_offset(update_index: u32) -> u32 {
+  let index_base = indices_base(update_index);
+  var rem = update_index;
+{% for axis in range(source.updatesRank) %}
+{% set update_stride = namespace(value=1) %}
+{% for j in range(axis + 1, source.updatesRank) %}
+{% set update_stride.value = update_stride.value * source.updatesShape[j] %}
+{% endfor %}
+{% set update_stride_divisor = update_stride.value %}
+{% set suffixBegin = source.indicesRank - 1 %}
+{% set suffixEnd = suffixBegin + source.dataRank - source.indexDepth %}
+{% if axis >= suffixBegin and axis < suffixEnd %}
+  let ucoord{{ axis }} = rem / {{ update_stride_divisor }}u;
+{% endif %}
+  rem %= {{ update_stride_divisor }}u;
+{% endfor %}
+  var offset = 0u;
+{% for axis in range(source.dataRank) %}
+{% set data_stride = namespace(value=1) %}
+{% for j in range(axis + 1, source.dataRank) %}
+{% set data_stride.value = data_stride.value * source.dataShape[j] %}
+{% endfor %}
+{% if axis < source.indexDepth %}
+  let gathered{{ axis }} = normalize_index(indices[index_base + {{ axis }}u], {{ source.dataShape[axis] }}u);
+  if (gathered{{ axis }} == INVALID_INDEX) {
+    return INVALID_INDEX;
+  }
+  offset += gathered{{ axis }} * {{ data_stride.value }}u;
+{% else %}
+{% set suffix_axis = (source.indicesRank - 1) + (axis - source.indexDepth) %}
+  offset += ucoord{{ suffix_axis }} * {{ data_stride.value }}u;
+{% endif %}
+{% endfor %}
+  return offset;
+}
+fn apply_update(dst: u32, update: f32) {
+  let word_index = dst >> 1u;
+  let high_half = (dst & 1u) != 0u;
+  var observed = atomicLoad(&output[word_index]);
+  loop {
+    var pair = unpack2x16float(observed);
+    let current = select(pair.x, pair.y, high_half);
+    let reduced = current + update;
+    if (high_half) {
+      pair.y = reduced;
+    } else {
+      pair.x = reduced;
+    }
+    let exchanged = atomicCompareExchangeWeak(&output[word_index], observed, pack2x16float(pair));
+    if (exchanged.exchanged) {
+      return;
+    }
+    observed = exchanged.old_value;
+  }
+}
+@compute @workgroup_size(WG, 1, 1)
+fn main(
+  @builtin(global_invocation_id) gid: vec3<u32>,
+  @builtin(num_workgroups) nwg: vec3<u32>,
+) {
+  let row = nwg.x * WG;
+  let stride = row * nwg.y;
+  for (var i = gid.y * row + gid.x; i < params.count; i += stride) {
+    let dst = output_offset(i);
+    if (dst != INVALID_INDEX) {
+      apply_update(dst, f32(updates[i]));
+    }
+  }
+}

build/webgpu/scatter-nd-reduction-slice.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,149 @@

+// Column-parallel ScatterND reduction without atomics for f16/int8. Duplicate
+// indices can collide only within one inner-width column: each update retains
+// its trailing coordinate `w`, while indices remap only its leading slice.
+// Different `w` values write disjoint cells, so columns run in parallel and the
+// slice sweep within each column stays serial to order read-modify-write safely.
+// A grid-stride loop covers any width.
+// INNER_WIDTH = product(updatesShape[indicesRank-1:]) = data slice width.
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const WG: u32 = {{ workgroupSize }}u;
+{% set inner_width = namespace(value=1) %}
+{% for j in range(source.indicesRank - 1, source.updatesRank) %}
+{% set inner_width.value = inner_width.value * source.updatesShape[j] %}
+{% endfor %}
+const INNER_WIDTH: u32 = {{ inner_width.value if inner_width.value != 0 else 1 }}u;
+const INVALID_INDEX: u32 = 0xffffffffu;
+fn normalize_index(raw: {{ indexScalar }}, dim: u32) -> u32 {
+  let dim_i = i32(dim);
+  var index = raw;
+  if (index < 0) {
+    index = index + dim_i;
+  }
+  if (index < 0 || index >= dim_i) {
+    return INVALID_INDEX;
+  }
+  return u32(index);
+}
+fn indices_base(update_index: u32) -> u32 {
+  var rem = update_index;
+  var offset = 0u;
+{% for axis in range(source.updatesRank) %}
+{% set update_stride = namespace(value=1) %}
+{% for j in range(axis + 1, source.updatesRank) %}
+{% set update_stride.value = update_stride.value * source.updatesShape[j] %}
+{% endfor %}
+{% set update_stride_divisor = update_stride.value if update_stride.value != 0 else 1 %}
+{% if axis < source.indicesRank - 1 %}
+  let coord{{ axis }} = rem / {{ update_stride_divisor }}u;
+{% endif %}
+  rem = rem % {{ update_stride_divisor }}u;
+{% if axis < source.indicesRank - 1 %}
+{% set index_stride = namespace(value=source.indexDepth) %}
+{% for j in range(axis + 1, source.indicesRank - 1) %}
+{% set index_stride.value = index_stride.value * source.indicesShape[j] %}
+{% endfor %}
+  offset = offset + coord{{ axis }} * {{ index_stride.value }}u;
+{% endif %}
+{% endfor %}
+  return offset;
+}
+fn output_offset(update_index: u32) -> u32 {
+  let index_base = indices_base(update_index);
+  var rem = update_index;
+{% for axis in range(source.updatesRank) %}
+{% set update_stride = namespace(value=1) %}
+{% for j in range(axis + 1, source.updatesRank) %}
+{% set update_stride.value = update_stride.value * source.updatesShape[j] %}
+{% endfor %}
+{% set update_stride_divisor = update_stride.value if update_stride.value != 0 else 1 %}
+{% set suffixBegin = source.indicesRank - 1 %}
+{% set suffixEnd = suffixBegin + source.dataRank - source.indexDepth %}
+{% if axis >= suffixBegin and axis < suffixEnd %}
+  let ucoord{{ axis }} = rem / {{ update_stride_divisor }}u;
+{% endif %}
+  rem = rem % {{ update_stride_divisor }}u;
+{% endfor %}
+  var offset = 0u;
+{% for axis in range(source.dataRank) %}
+{% set data_stride = namespace(value=1) %}
+{% for j in range(axis + 1, source.dataRank) %}
+{% set data_stride.value = data_stride.value * source.dataShape[j] %}
+{% endfor %}
+{% if axis < source.indexDepth %}
+  let gathered{{ axis }} = normalize_index(indices[index_base + {{ axis }}u], {{ source.dataShape[axis] }}u);
+  if (gathered{{ axis }} == INVALID_INDEX) {
+    return INVALID_INDEX;
+  }
+  offset = offset + gathered{{ axis }} * {{ data_stride.value }}u;
+{% else %}
+{% set suffix_axis = (source.indicesRank - 1) + (axis - source.indexDepth) %}
+  offset = offset + ucoord{{ suffix_axis }} * {{ data_stride.value }}u;
+{% endif %}
+{% endfor %}
+  return offset;
+}
+@compute @workgroup_size(WG)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>,
+        @builtin(num_workgroups) nwg: vec3<u32>) {
+  let num_slices = params.count / INNER_WIDTH;
+  let stride = nwg.x * WG;
+  // One thread per inner-width column; sweep slices for that column so
+  // duplicate-index collisions stay on a single lane (disjoint columns never race).
+  for (var w = gid.x; w < INNER_WIDTH; w = w + stride) {
+    for (var s = 0u; s < num_slices; s = s + 1u) {
+      let i = s * INNER_WIDTH + w;
+      let dst = output_offset(i);
+      if (dst == INVALID_INDEX) {
+        continue;
+      }
+{% set scatterIndent = "      " %}
+{% if scalar == "i32" or scalar == "u32" %}
+{{ scatterIndent }}let current = output[dst];
+{{ scatterIndent }}let update = updates[i];
+{% if reductionOp == "add" %}
+{{ scatterIndent }}let combined = current + update;
+{% elif reductionOp == "mul" %}
+{{ scatterIndent }}let combined = current * update;
+{% elif reductionOp == "min" %}
+{{ scatterIndent }}output[dst] = min(current, update);
+{% else %}
+{{ scatterIndent }}output[dst] = max(current, update);
+{% endif %}
+{% if reductionOp == "add" or reductionOp == "mul" %}
+{% if narrowBits == 8 and narrowSigned %}
+{{ scatterIndent }}let low = combined & 255;
+{{ scatterIndent }}output[dst] = select(low, low - 256, low > 127);
+{% elif narrowBits == 16 and narrowSigned %}
+{{ scatterIndent }}let low = combined & 65535;
+{{ scatterIndent }}output[dst] = select(low, low - 65536, low > 32767);
+{% elif narrowBits == 8 %}
+{{ scatterIndent }}output[dst] = combined & 255u;
+{% else %}
+{{ scatterIndent }}output[dst] = combined;
+{% endif %}
+{% endif %}
+{% else %}
+{{ scatterIndent }}let current = f32(output[dst]);
+{{ scatterIndent }}let update = f32(updates[i]);
+{% if reductionOp == "add" %}
+{{ scatterIndent }}output[dst] = {{ scalar }}(current + update);
+{% elif reductionOp == "mul" %}
+{{ scatterIndent }}output[dst] = {{ scalar }}(current * update);
+{% elif reductionOp == "min" %}
+{{ scatterIndent }}output[dst] = {{ scalar }}(min(current, update));
+{% else %}
+{{ scatterIndent }}output[dst] = {{ scalar }}(max(current, update));
+{% endif %}
+{% endif %}
+    }
+  }
+}

build/webgpu/scatter-nd.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,94 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const INVALID_INDEX: u32 = 0xffffffffu;
+const WG: u32 = {{ workgroupSize }}u;
+fn normalize_index(raw: {{ indexScalar }}, dim: u32) -> u32 {
+  let dim_i = i32(dim);
+  var index = raw;
+  if (index < 0) {
+    index = index + dim_i;
+  }
+  if (index < 0 || index >= dim_i) {
+    return INVALID_INDEX;
+  }
+  return u32(index);
+}
+fn indices_base(update_index: u32) -> u32 {
+  var rem = update_index;
+  var offset = 0u;
+{% for axis in range(source.updatesRank) %}
+{% set update_stride = namespace(value=1) %}
+{% for j in range(axis + 1, source.updatesRank) %}
+{% set update_stride.value = update_stride.value * source.updatesShape[j] %}
+{% endfor %}
+{% set update_stride_divisor = update_stride.value if update_stride.value != 0 else 1 %}
+{% if axis < source.indicesRank - 1 %}
+  let coord{{ axis }} = rem / {{ update_stride_divisor }}u;
+{% endif %}
+  rem = rem % {{ update_stride_divisor }}u;
+{% if axis < source.indicesRank - 1 %}
+{% set index_stride = namespace(value=source.indexDepth) %}
+{% for j in range(axis + 1, source.indicesRank - 1) %}
+{% set index_stride.value = index_stride.value * source.indicesShape[j] %}
+{% endfor %}
+  offset = offset + coord{{ axis }} * {{ index_stride.value }}u;
+{% endif %}
+{% endfor %}
+  return offset;
+}
+fn output_offset(update_index: u32) -> u32 {
+  let index_base = indices_base(update_index);
+  var rem = update_index;
+{% for axis in range(source.updatesRank) %}
+{% set update_stride = namespace(value=1) %}
+{% for j in range(axis + 1, source.updatesRank) %}
+{% set update_stride.value = update_stride.value * source.updatesShape[j] %}
+{% endfor %}
+{% set update_stride_divisor = update_stride.value if update_stride.value != 0 else 1 %}
+{% set suffixBegin = source.indicesRank - 1 %}
+{% set suffixEnd = suffixBegin + source.dataRank - source.indexDepth %}
+{% if axis >= suffixBegin and axis < suffixEnd %}
+  let ucoord{{ axis }} = rem / {{ update_stride_divisor }}u;
+{% endif %}
+  rem = rem % {{ update_stride_divisor }}u;
+{% endfor %}
+  var offset = 0u;
+{% for axis in range(source.dataRank) %}
+{% set data_stride = namespace(value=1) %}
+{% for j in range(axis + 1, source.dataRank) %}
+{% set data_stride.value = data_stride.value * source.dataShape[j] %}
+{% endfor %}
+{% if axis < source.indexDepth %}
+  let gathered{{ axis }} = normalize_index(indices[index_base + {{ axis }}u], {{ source.dataShape[axis] }}u);
+  if (gathered{{ axis }} == INVALID_INDEX) {
+    return INVALID_INDEX;
+  }
+  offset = offset + gathered{{ axis }} * {{ data_stride.value }}u;
+{% else %}
+{% set suffix_axis = (source.indicesRank - 1) + (axis - source.indexDepth) %}
+  offset = offset + ucoord{{ suffix_axis }} * {{ data_stride.value }}u;
+{% endif %}
+{% endfor %}
+  return offset;
+}
+@compute @workgroup_size(WG)
+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 device's
+  // per-dimension dispatch limit.
+  let i = gid.x + gid.y * nwg.x * WG;
+  if (i >= params.count) {
+    return;
+  }
+  let offset = output_offset(i);
+  if (offset != INVALID_INDEX) {
+    output[offset] = updates[i];
+  }
+}

build/webgpu/scatter-scalar-add.wgsl.jinja ADDED Viewed

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

+// Feature-independent high-contention ScatterND add for a one-element output.
+// Each lane sums a grid-stride stream, each workgroup reduces locally, and one
+// lane emits one f32 CAS. This replaces O(updates) globally serialized CAS
+// operations with O(workgroups) CAS operations.
+{{ env.wgsl.resourceDeclarations }}
+const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
+const INVALID_INDEX: u32 = 0xffffffffu;
+var<workgroup> partials: array<f32, {{ tunables.WORKGROUP_SIZE }}>;
+fn normalize_index(raw: {{ indexScalar }}, dim: u32) -> u32 {
+  let dim_i = i32(dim);
+  var index = raw;
+  if (index < 0) {
+    index = index + dim_i;
+  }
+  if (index < 0 || index >= dim_i) {
+    return INVALID_INDEX;
+  }
+  return u32(index);
+}
+fn atomic_add_f32(update: f32) {
+  var old_bits = atomicLoad(&output[0]);
+  loop {
+    let result = atomicCompareExchangeWeak(
+      &output[0], old_bits, bitcast<u32>(bitcast<f32>(old_bits) + update)
+    );
+    if (result.exchanged) { break; }
+    old_bits = result.old_value;
+  }
+}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(
+  @builtin(global_invocation_id) gid: vec3<u32>,
+  @builtin(local_invocation_id) lid: vec3<u32>,
+  @builtin(num_workgroups) nwg: vec3<u32>
+) {
+  let total_threads = nwg.x * WG;
+  var sum = 0.0;
+  for (var i = gid.x; i < params.count; i += total_threads) {
+    if (normalize_index(indices[i], 1u) != INVALID_INDEX) {
+      sum += updates[i];
+    }
+  }
+  partials[lid.x] = sum;
+  workgroupBarrier();
+  var offset = WG / 2u;
+  loop {
+    if (lid.x < offset) {
+      partials[lid.x] += partials[lid.x + offset];
+    }
+    workgroupBarrier();
+    if (offset <= 1u) { break; }
+    offset /= 2u;
+  }
+  if (lid.x == 0u) {
+    atomic_add_f32(partials[0]);
+  }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,1326 @@

+{
+  "op": "ai.onnx.ScatterND",
+  "fixtureArrays": {
+    "scatter_nd_data_4x4x4": [1, 2, 3, 4, 5, 6, 7, 8, 8, 7, 6, 5, 4, 3, 2, 1, 1, 2, 3, 4, 5, 6, 7, 8, 8, 7, 6, 5, 4, 3, 2, 1, 8, 7, 6, 5, 4, 3, 2, 1, 1, 2, 3, 4, 5, 6, 7, 8, 8, 7, 6, 5, 4, 3, 2, 1, 1, 2, 3, 4, 5, 6, 7, 8],
+    "backend_row_slice_updates_input_updates": [5, 5, 5, 5, 6, 6, 6, 6, 7, 7, 7, 7, 8, 8, 8, 8, 1, 1, 1, 1, 2, 2, 2, 2, 3, 3, 3, 3, 4, 4, 4, 4],
+    "ort_reduction_add_slice_duplicates_input_updates": [2, 4, 8, 16, 32, 64, 128, 256, 512, 1024, 2048, 4096, 8192, 16384, 32768, 65536, 131072, 262144],
+    "f16_reduction_add_sparse_router_gradient_input_indices": [0, 17, 34, 51, 4, 21, 38, 55, 8, 25, 42, 59, 12, 29, 46, 63, 16, 33, 50, 3, 20, 37, 54, 7, 24, 41, 58, 11, 28, 45, 62, 15, 32, 49, 2, 19, 36, 53, 6, 23, 40, 57, 10, 27, 44, 61, 14, 31, 48, 1, 18, 35, 52, 5, 22, 39, 56, 9, 26, 43, 60, 13, 30, 47]
+  },
+  "cases": [
+    {
+      "name": "dispatch_cliff_batch0_row_update",
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [1, 16777219], "data": { "kind": "constant", "value": 1.0 } },
+        "indices": { "dtype": "int32", "shape": [1, 1], "data": { "kind": "values", "values": [0] } },
+        "updates": { "dtype": "float32", "shape": [1, 16777219], "data": { "kind": "constant", "value": 7.0 } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 16777219], "tolerance": 0 } }
+    },
+    {
+      "name": "scalar_updates",
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [3, 3],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 10.0, 11.0, 12.0, 20.0, 21.0, 22.0] }
+        },
+        "indices": { "dtype": "int32", "shape": [2, 2], "data": { "kind": "values", "values": [0, 2, 2, 1] } },
+        "updates": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [99.0, 77.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 3] } }
+    },
+    {
+      "name": "row_updates_f16",
+      "inputs": {
+        "data": { "dtype": "float16", "shape": [3, 4] },
+        "indices": { "dtype": "int32", "shape": [2, 1], "data": { "kind": "values", "values": [2, 0] } },
+        "updates": {
+          "dtype": "float16",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [9.0, 8.0, 7.0, 6.0, 5.0, 4.0, 3.0, 2.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [3, 4] } },
+      "tolerance": 0.001
+    },
+    {
+      "name": "rank5_indices_scalar_updates",
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 10.0, 11.0, 12.0] }
+        },
+        "indices": { "dtype": "int32", "shape": [1, 1, 1, 1, 2], "data": { "kind": "values", "values": [1, 2] } },
+        "updates": { "dtype": "float32", "shape": [1, 1, 1, 1], "data": { "kind": "values", "values": [99.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank6_indices_scalar_updates",
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [3, 3],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 10.0, 11.0, 12.0, 20.0, 21.0, 22.0] }
+        },
+        "indices": {
+          "dtype": "int32",
+          "shape": [1, 1, 1, 1, 2, 2],
+          "data": { "kind": "values", "values": [0, 2, 2, 0] }
+        },
+        "updates": {
+          "dtype": "float32",
+          "shape": [1, 1, 1, 1, 2],
+          "data": { "kind": "values", "values": [99.0, 77.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank6_full_depth_scalar_updates",
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [1, 2, 1, 2, 1, 3],
+          "data": {
+            "kind": "values",
+            "values": [0.0, 1.0, 2.0, 10.0, 11.0, 12.0, 100.0, 101.0, 102.0, 110.0, 111.0, 112.0]
+          }
+        },
+        "indices": {
+          "dtype": "int32",
+          "shape": [2, 6],
+          "data": { "kind": "values", "values": [0, 0, 0, 0, 0, 2, 0, 1, 0, 1, 0, 0] }
+        },
+        "updates": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [900.0, 901.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 1, 2, 1, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "reduction_add_duplicate_indices",
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [10.0, 20.0, 30.0] } },
+        "indices": { "dtype": "int32", "shape": [3, 1], "data": { "kind": "values", "values": [1, 1, 2] } },
+        "updates": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [2.0, 3.0, 4.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "reduction_add_subnormal_duplicate_indices_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: Metal flushes denormals to zero in the ALU, so this reduction's subnormal sum/product result is flushed on GPU; permanent FTZ limitation."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/scatter_nd_op_test.cc",
+        "test": "ScatterNDOpTest.ScatterND_18_add",
+        "notes": "Duplicate-index add reduction with finite subnormal scalar updates; valid tiny updates must accumulate to nonzero subnormal outputs."
+      },
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.0, 0.0, 0.0, 0.0] } },
+        "indices": { "dtype": "int32", "shape": [4, 1], "data": { "kind": "values", "values": [1, 1, 3, 3] } },
+        "updates": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [1e-40, 1e-40, -1e-40, -1e-40] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4], "tolerance": 1e-44 } }
+    },
+    {
+      "name": "reduction_mul_normal_inputs_subnormal_product_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: Metal flushes denormals to zero in the ALU, so this reduction's subnormal sum/product result is flushed on GPU; permanent FTZ limitation."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/scatter_nd_op_test.cc",
+        "test": "ScatterNDOpTest.ScatterND_18_mul",
+        "notes": "Duplicate-index mul reduction where tiny normal updates produce a finite subnormal product."
+      },
+      "attrs": { "reduction": "mul" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.0, 1.0, 0.0, 0.0] } },
+        "indices": { "dtype": "int32", "shape": [2, 1], "data": { "kind": "values", "values": [1, 1] } },
+        "updates": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1e-20, 1e-20] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4], "tolerance": 1e-44 } }
+    },
+    {
+      "name": "rank6_indices_reduction_add_duplicate_indices",
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [10.0, 20.0, 30.0] } },
+        "indices": { "dtype": "int32", "shape": [1, 1, 1, 1, 3, 1], "data": { "kind": "values", "values": [1, 1, 2] } },
+        "updates": {
+          "dtype": "float32",
+          "shape": [1, 1, 1, 1, 3],
+          "data": { "kind": "values", "values": [2.0, 3.0, 4.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank6_reduction_add_duplicate_indices",
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [1, 1, 1, 1, 1, 4],
+          "data": { "kind": "values", "values": [0.0, 10.0, 20.0, 30.0] }
+        },
+        "indices": {
+          "dtype": "int32",
+          "shape": [3, 6],
+          "data": { "kind": "values", "values": [0, 0, 0, 0, 0, 1, 0, 0, 0, 0, 0, 1, 0, 0, 0, 0, 0, 3] }
+        },
+        "updates": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [2.0, 3.0, 4.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 1, 1, 1, 1, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "backend_row_slice_updates",
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [4, 4, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/scatter_nd_data_4x4x4" } }
+        },
+        "indices": { "dtype": "int32", "shape": [2, 1], "data": { "kind": "values", "values": [0, 2] } },
+        "updates": {
+          "dtype": "float32",
+          "shape": [2, 4, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/backend_row_slice_updates_input_updates" } }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 4, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "backend_row_slice_reduction_add_duplicate",
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [4, 4, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/scatter_nd_data_4x4x4" } }
+        },
+        "indices": { "dtype": "int32", "shape": [2, 1], "data": { "kind": "values", "values": [0, 0] } },
+        "updates": {
+          "dtype": "float32",
+          "shape": [2, 4, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/backend_row_slice_updates_input_updates" } }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 4, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_slice_float_rank2",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/scatter_nd_op_test.cc",
+        "test": "ScatterNDOpTest.ScatterND_slice_float_int64_t"
+      },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "values", "values": [0.0, 0.1, 0.1, 0.1] } },
+        "indices": { "dtype": "int32", "shape": [2, 1], "data": { "kind": "values", "values": [1, 0] } },
+        "updates": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "values", "values": [0.2, 0.3, 0.0, 0.1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_3tensor_int16_slice_update",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/scatter_nd_op_test.cc",
+        "test": "ScatterNDOpTest.ScatterND_3tensor_int64"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "int16",
+          "shape": [2, 2, 2],
+          "data": { "kind": "values", "values": [0, 1, 2, 3, 0, 1, 2, 3] }
+        },
+        "indices": { "dtype": "int32", "shape": [1, 1], "data": { "kind": "values", "values": [1] } },
+        "updates": { "dtype": "int16", "shape": [1, 2, 2], "data": { "kind": "values", "values": [4, 5, 6, 7] } }
+      },
+      "outputs": { "output": { "dtype": "int16", "shape": [2, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_batched_index_float_adapted",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/scatter_nd_op_test.cc",
+        "test": "ScatterNDOpTest.ScatterND_batched_index_int64",
+        "notes": "ORT uses int64 data; this framework covers the same batched-index pattern with float32 data."
+      },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "values", "values": [2.0, 3.0, 2.0, 3.0] } },
+        "indices": { "dtype": "int32", "shape": [2, 1, 2], "data": { "kind": "values", "values": [0, 0, 0, 1] } },
+        "updates": { "dtype": "float32", "shape": [2, 1], "data": { "kind": "values", "values": [0.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_batched_index_bool",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/scatter_nd_op_test.cc",
+        "test": "ScatterNDOpTest.ScatterND_batched_index_bool_int64"
+      },
+      "inputs": {
+        "data": { "dtype": "bool", "shape": [2, 2], "data": { "kind": "values", "values": [0, 1, 0, 1] } },
+        "indices": { "dtype": "int32", "shape": [2, 1, 2], "data": { "kind": "values", "values": [0, 0, 0, 1] } },
+        "updates": { "dtype": "bool", "shape": [2, 1], "data": { "kind": "values", "values": [1, 0] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "bool",
+          "shape": [2, 2],
+          "data": { "kind": "values", "values": [1, 0, 0, 1] },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "ort_sliced_index_float_adapted",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/scatter_nd_op_test.cc",
+        "test": "ScatterNDOpTest.ScatterND_sliced_index_int64",
+        "notes": "ORT uses int64 data; this framework covers the same sliced-index update pattern with float32 data."
+      },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "constant", "value": 0.0 } },
+        "indices": { "dtype": "int32", "shape": [2, 1, 1], "data": { "kind": "values", "values": [1, 0] } },
+        "updates": {
+          "dtype": "float32",
+          "shape": [2, 1, 2],
+          "data": { "kind": "values", "values": [2.0, 3.0, 0.0, 1.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_uint32_batched_index_depth2",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/scatter_nd_op_test.cc",
+        "test": "ScatterNDOpTest.ScatterND_batched_3tensor_int64",
+        "notes": "Compact uint32 batched ScatterND fixture using the same depth-2 index pattern as ORT."
+      },
+      "inputs": {
+        "data": {
+          "dtype": "uint32",
+          "shape": [2, 2, 2],
+          "data": { "kind": "values", "values": [0, 0, 0, 0, 0, 0, 0, 0] }
+        },
+        "indices": {
+          "dtype": "int32",
+          "shape": [2, 2, 2],
+          "data": { "kind": "values", "values": [0, 1, 1, 0, 0, 0, 1, 1] }
+        },
+        "updates": {
+          "dtype": "uint32",
+          "shape": [2, 2, 2],
+          "data": { "kind": "values", "values": [2, 3, 4, 5, 0, 1, 6, 7] }
+        }
+      },
+      "outputs": { "output": { "dtype": "uint32", "shape": [2, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int8_negative_indices_scalar_updates",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/scatter_nd_op_test.cc",
+        "test": "ScatterNDOpTest.ScatterND_3tensor_int64",
+        "notes": "Compact signed-byte scalar-update fixture with negative indices and edge payload values."
+      },
+      "inputs": {
+        "data": {
+          "dtype": "int8",
+          "shape": [2, 2, 2],
+          "data": { "kind": "values", "values": [0, 0, 2, 3, 4, 0, 6, 7] }
+        },
+        "indices": { "dtype": "int32", "shape": [2, 3], "data": { "kind": "values", "values": [0, 0, 1, -1, 0, -1] } },
+        "updates": { "dtype": "int8", "shape": [2], "data": { "kind": "values", "values": [-128, 127] } }
+      },
+      "outputs": { "output": { "dtype": "int8", "shape": [2, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "empty_indices_zero_width_update_identity",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/scatter_nd_op_test.cc",
+        "test": "ScatterNDOpTest.ScatterND_empty_indices",
+        "notes": "Valid empty-index ScatterND identity where both the data suffix and update suffix have zero width."
+      },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [2, 0], "data": { "kind": "values", "values": [] } },
+        "indices": { "dtype": "int32", "shape": [0, 1], "data": { "kind": "values", "values": [] } },
+        "updates": { "dtype": "float32", "shape": [0, 0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": {
+        "output": { "dtype": "float32", "shape": [2, 0], "tolerance": 0, "data": { "kind": "values", "values": [] } }
+      }
+    },
+    {
+      "name": "empty_indices_zero_width_reduction_add_identity",
+      "provenance": {
+        "notes": "Reduction twin of the zero-width identity above. The slice route floors its inner width at one so a zero-width data suffix cannot divide by zero, and only a reduction case reaches that route."
+      },
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [2, 0], "data": { "kind": "values", "values": [] } },
+        "indices": { "dtype": "int32", "shape": [0, 1], "data": { "kind": "values", "values": [] } },
+        "updates": { "dtype": "float32", "shape": [0, 0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": {
+        "output": { "dtype": "float32", "shape": [2, 0], "tolerance": 0, "data": { "kind": "values", "values": [] } }
+      }
+    },
+    {
+      "name": "ort_empty_indices_identity",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/scatter_nd_op_test.cc",
+        "test": "ScatterNDOpTest.ScatterND_empty_indices"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        },
+        "indices": { "dtype": "int32", "shape": [0, 1], "data": { "kind": "values", "values": [] } },
+        "updates": { "dtype": "float32", "shape": [0, 3], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_reduction_add_slice_duplicates",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/scatter_nd_op_test.cc",
+        "test": "ScatterNDOpTest.ScatterND_18_add"
+      },
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 2, 3],
+          "data": { "kind": "values", "values": [0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1] }
+        },
+        "indices": { "dtype": "int32", "shape": [3, 1], "data": { "kind": "values", "values": [0, 1, 0] } },
+        "updates": {
+          "dtype": "float32",
+          "shape": [3, 2, 3],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_reduction_add_slice_duplicates_input_updates" }
+          }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2, 3], "tolerance": 0.01 } }
+    },
+    {
+      "name": "ort_reduction_mul_slice_duplicates",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/scatter_nd_op_test.cc",
+        "test": "ScatterNDOpTest.ScatterND_18_mul"
+      },
+      "attrs": { "reduction": "mul" },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 2, 3],
+          "data": { "kind": "values", "values": [0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1] }
+        },
+        "indices": { "dtype": "int32", "shape": [3, 1], "data": { "kind": "values", "values": [0, 1, 0] } },
+        "updates": {
+          "dtype": "float32",
+          "shape": [3, 2, 3],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_reduction_add_slice_duplicates_input_updates" }
+          }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2, 3], "tolerance": 0.01 } }
+    },
+    {
+      "name": "ort_reduction_min_slice_duplicates",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/scatter_nd_op_test.cc",
+        "test": "ScatterNDOpTest.ScatterND_18_min"
+      },
+      "attrs": { "reduction": "min" },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 2, 3],
+          "data": { "kind": "values", "values": [0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1] }
+        },
+        "indices": { "dtype": "int32", "shape": [3, 1], "data": { "kind": "values", "values": [0, 1, 0] } },
+        "updates": {
+          "dtype": "float32",
+          "shape": [3, 2, 3],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_reduction_add_slice_duplicates_input_updates" }
+          }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_reduction_max_slice_duplicates",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/scatter_nd_op_test.cc",
+        "test": "ScatterNDOpTest.ScatterND_18_max"
+      },
+      "attrs": { "reduction": "max" },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 2, 3],
+          "data": { "kind": "values", "values": [0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1] }
+        },
+        "indices": { "dtype": "int32", "shape": [3, 1], "data": { "kind": "values", "values": [0, 1, 0] } },
+        "updates": {
+          "dtype": "float32",
+          "shape": [3, 2, 3],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_reduction_add_slice_duplicates_input_updates" }
+          }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_int32_negative_indices_matrix",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/scatter_nd_op_test.cc",
+        "test": "ScatterNDOpTest.ScatterND_matrice_int64_int64_neg_indices",
+        "notes": "ORT uses int64 data and indices; this framework covers the same negative-index pattern with float32 data and int32 index storage."
+      },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "values", "values": [1.0, 1.0, 2.0, 2.0] } },
+        "indices": { "dtype": "int32", "shape": [2, 2], "data": { "kind": "values", "values": [0, 0, -1, -1] } },
+        "updates": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 3.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_int32_negative_indices_slice_update_rank3",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/scatter_nd_op_test.cc",
+        "test": "ScatterNDOpTest.ScatterND_3tensor_int64",
+        "notes": "ORT uses integer data; this framework covers the same negative-index slice update with float32 data and int32 index storage."
+      },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 2, 2],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0] }
+        },
+        "indices": { "dtype": "int32", "shape": [2, 2], "data": { "kind": "values", "values": [0, -1, -1, 0] } },
+        "updates": {
+          "dtype": "float32",
+          "shape": [2, 2],
+          "data": { "kind": "values", "values": [20.0, 21.0, 30.0, 31.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_reduction_mul_rank6_slice_duplicates",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/scatter_nd_op_test.cc",
+        "test": "ScatterNDOpTest.ScatterND_18_mul_long_shape",
+        "notes": "Same duplicate-index reduction pattern as ORT's long-shape case, constrained to this manifest's rank-6 limit."
+      },
+      "attrs": { "reduction": "mul" },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 2, 3, 1, 1, 1],
+          "data": { "kind": "values", "values": [0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1] }
+        },
+        "indices": { "dtype": "int32", "shape": [3, 1], "data": { "kind": "values", "values": [0, 1, 0] } },
+        "updates": {
+          "dtype": "float32",
+          "shape": [3, 2, 3, 1, 1, 1],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_reduction_add_slice_duplicates_input_updates" }
+          }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2, 3, 1, 1, 1], "tolerance": 0.01 } }
+    },
+    {
+      "name": "ort_reduction_min_updates_win_duplicate_indices",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/scatter_nd_op_test.cc",
+        "test": "ScatterNDOpTest.ScatterND_18_min",
+        "notes": "Uses ORT's reduction=min semantics with values chosen so both data and duplicate updates affect the answer."
+      },
+      "attrs": { "reduction": "min" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [5.0, -1.0, 3.0] } },
+        "indices": { "dtype": "int32", "shape": [3, 1], "data": { "kind": "values", "values": [0, 0, 2] } },
+        "updates": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [2.0, -4.0, 10.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_reduction_max_updates_win_duplicate_indices",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/scatter_nd_op_test.cc",
+        "test": "ScatterNDOpTest.ScatterND_18_max",
+        "notes": "Uses ORT's reduction=max semantics with values chosen so both data and duplicate updates affect the answer."
+      },
+      "attrs": { "reduction": "max" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-5.0, 1.0, 3.0] } },
+        "indices": { "dtype": "int32", "shape": [3, 1], "data": { "kind": "values", "values": [0, 0, 2] } },
+        "updates": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-2.0, 4.0, 2.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "onnx_backend_scatternd",
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [4, 4, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/scatter_nd_data_4x4x4" } }
+        },
+        "indices": { "dtype": "int32", "shape": [2, 1], "data": { "kind": "values", "values": [0, 2] } },
+        "updates": {
+          "dtype": "float32",
+          "shape": [2, 4, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/backend_row_slice_updates_input_updates" } }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 4, 4] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_scatternd",
+        "notes": "ONNX int64 tensors use framework int32/uint32 slots where representable."
+      }
+    },
+    {
+      "name": "onnx_backend_scatternd_add",
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [4, 4, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/scatter_nd_data_4x4x4" } }
+        },
+        "indices": { "dtype": "int32", "shape": [2, 1], "data": { "kind": "values", "values": [0, 0] } },
+        "updates": {
+          "dtype": "float32",
+          "shape": [2, 4, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/backend_row_slice_updates_input_updates" } }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 4, 4] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_scatternd_add",
+        "notes": "ONNX int64 tensors use framework int32/uint32 slots where representable."
+      }
+    },
+    {
+      "name": "onnx_backend_scatternd_max",
+      "attrs": { "reduction": "max" },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [4, 4, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/scatter_nd_data_4x4x4" } }
+        },
+        "indices": { "dtype": "int32", "shape": [2, 1], "data": { "kind": "values", "values": [0, 0] } },
+        "updates": {
+          "dtype": "float32",
+          "shape": [2, 4, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/backend_row_slice_updates_input_updates" } }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 4, 4] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_scatternd_max",
+        "notes": "ONNX int64 tensors use framework int32/uint32 slots where representable."
+      }
+    },
+    {
+      "name": "onnx_backend_scatternd_min",
+      "attrs": { "reduction": "min" },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [4, 4, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/scatter_nd_data_4x4x4" } }
+        },
+        "indices": { "dtype": "int32", "shape": [2, 1], "data": { "kind": "values", "values": [0, 0] } },
+        "updates": {
+          "dtype": "float32",
+          "shape": [2, 4, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/backend_row_slice_updates_input_updates" } }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 4, 4] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_scatternd_min",
+        "notes": "ONNX int64 tensors use framework int32/uint32 slots where representable."
+      }
+    },
+    {
+      "name": "onnx_backend_scatternd_multiply",
+      "attrs": { "reduction": "mul" },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [4, 4, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/scatter_nd_data_4x4x4" } }
+        },
+        "indices": { "dtype": "int32", "shape": [2, 1], "data": { "kind": "values", "values": [0, 0] } },
+        "updates": {
+          "dtype": "float32",
+          "shape": [2, 4, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/backend_row_slice_updates_input_updates" } }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 4, 4] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_scatternd_multiply",
+        "notes": "ONNX int64 tensors use framework int32/uint32 slots where representable."
+      }
+    },
+    {
+      "name": "ort_int32_slice_updates_negative_indices_exact_above_float24",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/scatter_nd_op_test.cc",
+        "test": "ScatterNDOpTest.ScatterND_3tensor_int64",
+        "notes": "Native int32 projection of ORT's integer slice-update case with negative indices and values above f32's exact integer range."
+      },
+      "inputs": {
+        "data": {
+          "dtype": "int32",
+          "shape": [2, 2, 2],
+          "data": { "kind": "values", "values": [16777217, 1, 2, 3, 4, 5, -16777217, 7] }
+        },
+        "indices": { "dtype": "int32", "shape": [2, 2], "data": { "kind": "values", "values": [0, 1, -1, 0] } },
+        "updates": {
+          "dtype": "int32",
+          "shape": [2, 2],
+          "data": { "kind": "values", "values": [20000001, -20000001, 30000003, -30000003] }
+        }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [2, 2, 2] } }
+    },
+    {
+      "name": "int32_reduction_add_exact_above_float24",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/scatter_nd_op_test.cc",
+        "test": "ScatterNDOpTest.ScatterND_18_add",
+        "notes": "Integer projection of ORT's add-reduction coverage using values above f32's exact integer range."
+      },
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [0, 0] } },
+        "indices": { "dtype": "int32", "shape": [2, 1], "data": { "kind": "values", "values": [0, 0] } },
+        "updates": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [16777217, 1] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [2] } }
+    },
+    {
+      "name": "int32_reduction_min_duplicate_indices_negative",
+      "attrs": { "reduction": "min" },
+      "inputs": {
+        "data": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [5, -7, 3] } },
+        "indices": { "dtype": "int32", "shape": [3, 1], "data": { "kind": "values", "values": [0, 0, 2] } },
+        "updates": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [-2, 4, 10] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [-2, -7, 3] } } }
+    },
+    {
+      "name": "int32_reduction_max_duplicate_indices",
+      "attrs": { "reduction": "max" },
+      "inputs": {
+        "data": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [-5, -7, 3] } },
+        "indices": { "dtype": "int32", "shape": [3, 1], "data": { "kind": "values", "values": [0, 0, 2] } },
+        "updates": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [-2, 4, 1] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [4, -7, 3] } } }
+    },
+    {
+      "name": "int32_reduction_mul_duplicate_indices",
+      "attrs": { "reduction": "mul" },
+      "inputs": {
+        "data": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [2, 3, -1] } },
+        "indices": { "dtype": "int32", "shape": [3, 1], "data": { "kind": "values", "values": [0, 0, 2] } },
+        "updates": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [3, 5, -4] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [30, 3, 4] } } }
+    },
+    {
+      "name": "uint32_reduction_add_duplicate_indices",
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [1, 2] } },
+        "indices": { "dtype": "int32", "shape": [3, 1], "data": { "kind": "values", "values": [0, 0, 0] } },
+        "updates": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [2, 3, 4] } }
+      },
+      "outputs": { "output": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [10, 2] } } }
+    },
+    {
+      "name": "uint32_reduction_min_duplicate_indices",
+      "attrs": { "reduction": "min" },
+      "inputs": {
+        "data": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [8, 1] } },
+        "indices": { "dtype": "int32", "shape": [3, 1], "data": { "kind": "values", "values": [0, 0, 1] } },
+        "updates": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [5, 3, 9] } }
+      },
+      "outputs": { "output": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 1] } } }
+    },
+    {
+      "name": "uint32_reduction_max_duplicate_indices",
+      "attrs": { "reduction": "max" },
+      "inputs": {
+        "data": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [1, 7] } },
+        "indices": { "dtype": "int32", "shape": [3, 1], "data": { "kind": "values", "values": [0, 0, 1] } },
+        "updates": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [5, 9, 2] } }
+      },
+      "outputs": { "output": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [9, 7] } } }
+    },
+    {
+      "name": "uint32_reduction_mul_duplicate_indices",
+      "attrs": { "reduction": "mul" },
+      "inputs": {
+        "data": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 5] } },
+        "indices": { "dtype": "int32", "shape": [3, 1], "data": { "kind": "values", "values": [0, 0, 0] } },
+        "updates": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [2, 3, 4] } }
+      },
+      "outputs": { "output": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [72, 5] } } }
+    },
+    {
+      "name": "reduction_add_f32_high_contention_4096",
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.5, 0.0, 0.0, 0.0] } },
+        "indices": { "dtype": "int32", "shape": [4096, 1], "data": { "kind": "constant", "value": 0 } },
+        "updates": { "dtype": "float32", "shape": [4096], "data": { "kind": "constant", "value": 0.25 } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [1024.5, 0.0, 0.0, 0.0] },
+          "tolerance": 0.000001
+        }
+      }
+    },
+    {
+      "name": "reduction_min_f32_high_contention_2048",
+      "attrs": { "reduction": "min" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1000.0, -3.0] } },
+        "indices": { "dtype": "int32", "shape": [2048, 1], "data": { "kind": "constant", "value": 0 } },
+        "updates": {
+          "dtype": "float32",
+          "shape": [2048],
+          "data": { "kind": "cycle", "values": [9.0, 4.0, 7.0, 2.5, 8.0] }
+        }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [2],
+          "data": { "kind": "values", "values": [2.5, -3.0] },
+          "tolerance": 0.000001
+        }
+      }
+    },
+    {
+      "name": "ort_reduction_add_slice_negative_indices",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/scatter_nd_op_test.cc",
+        "test": "ScatterNDOpTest.ScatterND_18_add",
+        "notes": "Same duplicate slice-add reduction as ORT, with signed negative indices that normalize to rows 0, 1, 0."
+      },
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 2, 3],
+          "data": { "kind": "values", "values": [0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1] }
+        },
+        "indices": { "dtype": "int32", "shape": [3, 1], "data": { "kind": "values", "values": [-2, -1, -2] } },
+        "updates": {
+          "dtype": "float32",
+          "shape": [3, 2, 3],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_reduction_add_slice_duplicates_input_updates" }
+          }
+        }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [2, 2, 3],
+          "tolerance": 0.01,
+          "data": {
+            "kind": "values",
+            "values": [8194.1, 16388.1, 32776.1, 65552.1, 131104.1, 262208.1, 128.1, 256.1, 512.1, 1024.1, 2048.1, 4096.1]
+          }
+        }
+      }
+    },
+    {
+      "name": "reduction_mul_f32_high_contention_1024",
+      "attrs": { "reduction": "mul" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [7.0, 2.0] } },
+        "indices": { "dtype": "int32", "shape": [1024, 1], "data": { "kind": "constant", "value": 0 } },
+        "updates": { "dtype": "float32", "shape": [1024], "data": { "kind": "constant", "value": 1.0 } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [2],
+          "data": { "kind": "values", "values": [7.0, 2.0] },
+          "tolerance": 0.000001
+        }
+      }
+    },
+    {
+      "name": "f16_reduction_add_row_slice_duplicate_indices",
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": {
+          "dtype": "float16",
+          "shape": [4, 4],
+          "data": {
+            "kind": "values",
+            "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0]
+          }
+        },
+        "indices": { "dtype": "int32", "shape": [3, 1], "data": { "kind": "values", "values": [0, 0, 2] } },
+        "updates": {
+          "dtype": "float16",
+          "shape": [3, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.5, 1.5, 2.5, 3.5, 10.0, 20.0, 30.0, 40.0, 100.0, 200.0, 300.0, 400.0]
+          }
+        }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [4, 4], "tolerance": 0.01 } }
+    },
+    {
+      "name": "f16_reduction_mul_scalar_updates_duplicate_indices",
+      "attrs": { "reduction": "mul" },
+      "inputs": {
+        "data": { "dtype": "float16", "shape": [4], "data": { "kind": "values", "values": [2.0, 1.0, 3.0, 1.0] } },
+        "indices": { "dtype": "int32", "shape": [3, 1], "data": { "kind": "values", "values": [0, 0, 2] } },
+        "updates": { "dtype": "float16", "shape": [3], "data": { "kind": "values", "values": [1.5, 2.0, 4.0] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [4], "tolerance": 0.01 } }
+    },
+    {
+      "name": "f16_reduction_max_row_slice_duplicate_indices",
+      "attrs": { "reduction": "max" },
+      "inputs": {
+        "data": {
+          "dtype": "float16",
+          "shape": [3, 4],
+          "data": { "kind": "values", "values": [1.0, -1.0, 0.5, 2.0, -3.0, -3.0, -3.0, -3.0, 0.0, 0.0, 0.0, 0.0] }
+        },
+        "indices": { "dtype": "int32", "shape": [3, 1], "data": { "kind": "values", "values": [0, 0, 1] } },
+        "updates": {
+          "dtype": "float16",
+          "shape": [3, 4],
+          "data": { "kind": "values", "values": [0.0, -2.0, 1.5, 1.0, -0.5, 4.0, -1.0, 3.0, -1.0, -2.0, -1.0, -2.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [3, 4], "tolerance": 0.001 } }
+    },
+    {
+      "name": "int8_reduction_add_scalar_updates_duplicate_indices_in_range",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/scatter_nd_op_test.cc",
+        "test": "ScatterNDOpTest.ScatterND_18_add",
+        "notes": "int8 add reduction with duplicate indices kept within int8 range (no overflow/wrap), exercising the untested serial f32-cast int8 reduction branch."
+      },
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": { "dtype": "int8", "shape": [4], "data": { "kind": "values", "values": [10, 0, -20, 0] } },
+        "indices": { "dtype": "int32", "shape": [3, 1], "data": { "kind": "values", "values": [0, 0, 2] } },
+        "updates": { "dtype": "int8", "shape": [3], "data": { "kind": "values", "values": [5, 7, -30] } }
+      },
+      "outputs": { "output": { "dtype": "int8", "shape": [4], "tolerance": 0 } }
+    },
+    {
+      "name": "batch0_non_vec4_aligned_inner_row",
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [6, 7], "data": { "kind": "linspace", "start": 0.0, "end": 41.0 } },
+        "indices": { "dtype": "int32", "shape": [3, 1], "data": { "kind": "values", "values": [0, 4, 5] } },
+        "updates": {
+          "dtype": "float32",
+          "shape": [3, 7],
+          "data": {
+            "kind": "values",
+            "values": [100.0, 101.0, 102.0, 103.0, 104.0, 105.0, 106.0, 200.0, 201.0, 202.0, 203.0, 204.0, 205.0, 206.0, 300.0, 301.0, 302.0, 303.0, 304.0, 305.0, 306.0]
+          }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [6, 7] } }
+    },
+    {
+      "name": "reduction_slice_f16_multi_outer_batch_duplicate_depth2",
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": {
+          "dtype": "float16",
+          "shape": [3, 4],
+          "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] }
+        },
+        "indices": {
+          "dtype": "int32",
+          "shape": [2, 2, 2],
+          "data": { "kind": "values", "values": [0, 1, 0, 2, 1, 0, 1, 3] }
+        },
+        "updates": {
+          "dtype": "float16",
+          "shape": [2, 2],
+          "data": { "kind": "values", "values": [10.0, 20.0, 30.0, 40.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [3, 4], "tolerance": 0.1 } }
+    },
+    {
+      "name": "row_fast_vec4_last_row_edge_large_axis_dim",
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [65537, 8], "data": { "kind": "constant", "value": 0.0 } },
+        "indices": { "dtype": "int32", "shape": [3, 1], "data": { "kind": "values", "values": [0, 65536, 32768] } },
+        "updates": {
+          "dtype": "float32",
+          "shape": [3, 8],
+          "data": { "kind": "cycle", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [65537, 8] } }
+    },
+    {
+      "name": "reduction_add_single_hot_slot_many_updates",
+      "provenance": {
+        "notes": "Compact semantic sibling for the hot-slot atomic benchmark: many duplicate scalar updates target one output slot and must accumulate deterministically against the CPU reference."
+      },
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [10.0, 20.0, 30.0, 40.0, 50.0, 60.0, 70.0, 80.0] }
+        },
+        "indices": {
+          "dtype": "int32",
+          "shape": [16, 1],
+          "data": { "kind": "values", "values": [0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0] }
+        },
+        "updates": {
+          "dtype": "float32",
+          "shape": [16],
+          "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, 13.0, -14.0, 15.0, -16.0]
+          }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [8], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "scalar_output_add_many_updates_parallel_reduce",
+      "provenance": {
+        "notes": "Correctness lock for scalar accumulation with int32 indices, where every valid index targets the sole output element."
+      },
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [3.5] } },
+        "indices": { "dtype": "int32", "shape": [1024, 1], "data": { "kind": "constant", "value": 0 } },
+        "updates": {
+          "dtype": "float32",
+          "shape": [1024],
+          "data": { "kind": "cycle", "values": [0.25, -0.5, 1.0, -0.125, 0.0625, 0.75, -0.25, 0.5] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1], "tolerance": 0.001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "scalar_output_add_many_updates_parallel_reduce_i32_negative",
+      "provenance": {
+        "notes": "Negative indices on the same route. ONNX admits them and normalizeGatherIndex wraps them, so -1 addresses the sole output element exactly as 0 does; alternating the two makes a route that drops one of the spellings differ from the reference rather than merely halving the sum."
+      },
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [3.5] } },
+        "indices": { "dtype": "int32", "shape": [1024, 1], "data": { "kind": "cycle", "values": [0, -1] } },
+        "updates": {
+          "dtype": "float32",
+          "shape": [1024],
+          "data": { "kind": "cycle", "values": [0.25, -0.5, 1.0, -0.125, 0.0625, 0.75, -0.25, 0.5] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1], "tolerance": 0.001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "reduction_add_padding_hotspot_many_updates",
+      "provenance": {
+        "notes": "Correctness sibling for sparse-routing and padding-token workloads where many updates target one hot key in a much larger output table."
+      },
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [128], "data": { "kind": "constant", "value": 0.0 } },
+        "indices": { "dtype": "int32", "shape": [2048, 1], "data": { "kind": "constant", "value": 0 } },
+        "updates": { "dtype": "float32", "shape": [2048], "data": { "kind": "constant", "value": 1.0 } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [128], "tolerance": 0.001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "f16_reduction_add_sparse_router_gradient",
+      "provenance": {
+        "notes": "Correctness lock for f16 scalar sparse-gradient accumulation through the f32 atomic scratch path."
+      },
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": { "dtype": "float16", "shape": [64], "data": { "kind": "constant", "value": 0.0 } },
+        "indices": {
+          "dtype": "int32",
+          "shape": [4096, 1],
+          "data": {
+            "kind": "cycle",
+            "values": { "$ref": "#/fixtureArrays/f16_reduction_add_sparse_router_gradient_input_indices" }
+          }
+        },
+        "updates": {
+          "dtype": "float16",
+          "shape": [4096],
+          "data": { "kind": "fillFloat32", "sinStep": 0.031, "cosStep": 0.047, "scale": 0.01 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [64], "tolerance": 0.1, "relTolerance": 0.02 } }
+    },
+    {
+      "name": "f16_reduction_add_sparse_router_gradient_offset_scale_lock",
+      "provenance": {
+        "notes": "Scale lock for f16 reduction-add through atomic f32 scratch. Each of 64 slots receives 64 updates offset around +0.02, preventing zero-mean cancellation and making a lost, duplicated, or uniformly rescaled contribution observable. Tolerance covers the expected difference between per-add f16 rounding in the reference and one final f16 rounding after f32 scratch accumulation."
+      },
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": { "dtype": "float16", "shape": [64], "data": { "kind": "constant", "value": 0.0 } },
+        "indices": {
+          "dtype": "int32",
+          "shape": [4096, 1],
+          "data": {
+            "kind": "cycle",
+            "values": { "$ref": "#/fixtureArrays/f16_reduction_add_sparse_router_gradient_input_indices" }
+          }
+        },
+        "updates": {
+          "dtype": "float16",
+          "shape": [4096],
+          "data": { "kind": "fillFloat32", "sinStep": 0.031, "cosStep": 0.047, "scale": 0.01, "offset": 0.02 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [64], "tolerance": 0.012, "relTolerance": 0.004 } }
+    },
+    {
+      "name": "f16_reduction_add_row_slice_packed_atomic_route",
+      "provenance": {
+        "notes": "Correctness lock for direct packed-f16 CAS row-slice accumulation. Duplicate rows and adjacent columns contend on both halves of each output word."
+      },
+      "requires": { "features": ["shader-f16"] },
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": { "dtype": "float16", "shape": [16, 4], "data": { "kind": "constant", "value": 0.0 } },
+        "indices": {
+          "dtype": "int32",
+          "shape": [4096, 1],
+          "data": { "kind": "cycle", "values": [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15] }
+        },
+        "updates": { "dtype": "float16", "shape": [4096, 4], "data": { "kind": "constant", "value": 1.0 } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [16, 4], "tolerance": 0.01 } }
+    },
+    {
+      "name": "f16_reduction_add_rank3_batched_indices_packed_atomic",
+      "provenance": {
+        "notes": "Rank-3 batched indices through the direct packed-f16 CAS route, where indexDepth is less than rank(data). indices[2,8,1] pairs ascending rows in batch 0 with descending rows in batch 1, so dropping the interior index stride changes the update multiplicities. Data and updates are quarter-integers whose periods are coprime with both row and batch strides; every partial sum stays exactly representable in f16 regardless of atomic retirement order, allowing zero tolerance."
+      },
+      "requires": { "features": ["shader-f16"] },
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": {
+          "dtype": "float16",
+          "shape": [8, 256],
+          "data": { "kind": "cycle", "values": [-1.0, 0.5, -0.25, 2.0, 0.0, -1.5, 1.25] }
+        },
+        "indices": {
+          "dtype": "int32",
+          "shape": [2, 8, 1],
+          "data": { "kind": "values", "values": [0, 1, 2, 3, 4, 5, 6, 7, 7, 6, 5, 4, 3, 2, 1, 0] }
+        },
+        "updates": {
+          "dtype": "float16",
+          "shape": [2, 8, 256],
+          "data": { "kind": "cycle", "values": [0.25, 1.5, 0.75, 2.0, 1.25] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [8, 256], "tolerance": 0 } }
+    },
+    {
+      "name": "f16_reduction_max_full_depth_atomic_f32_route",
+      "provenance": {
+        "notes": "Correctness lock for the generalized full-index-depth f16 scratch-atomic route. Four thousand duplicate scalar updates contend on one destination, and max remains exact regardless of atomic order."
+      },
+      "attrs": { "reduction": "max" },
+      "inputs": {
+        "data": { "dtype": "float16", "shape": [16384], "data": { "kind": "constant", "value": 0.0 } },
+        "indices": { "dtype": "int32", "shape": [4096, 1], "data": { "kind": "constant", "value": 7 } },
+        "updates": { "dtype": "float16", "shape": [4096], "data": { "kind": "cycle", "values": [-3.0, 2.0, 1.0, 4.0] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [16384], "tolerance": 0 } }
+    },
+    {
+      "name": "rank7_scalar_updates",
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [1, 1, 1, 1, 1, 2, 3], "data": { "kind": "constant", "value": 0.0 } },
+        "indices": {
+          "dtype": "int32",
+          "shape": [2, 7],
+          "data": { "kind": "values", "values": [0, 0, 0, 0, 0, 0, 2, 0, 0, 0, 0, 0, 1, 0] }
+        },
+        "updates": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [9.0, 7.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 1, 1, 1, 1, 2, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "rank8_scalar_updates",
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [1, 1, 1, 1, 1, 1, 2, 3],
+          "data": { "kind": "linspace", "start": 1.0, "end": 6.0 }
+        },
+        "indices": {
+          "dtype": "int32",
+          "shape": [2, 8],
+          "data": { "kind": "values", "values": [0, 0, 0, 0, 0, 0, 1, 2, 0, 0, 0, 0, 0, 0, 0, 1] }
+        },
+        "updates": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [9.0, 7.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 1, 1, 1, 1, 1, 2, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int8_reduction_mul_wrap8",
+      "provenance": {
+        "source": "ONNX Runtime CPUExecutionProvider",
+        "notes": "The existing INT8 route previously accumulated the physical i32 product without narrowing it."
+      },
+      "attrs": { "reduction": "mul" },
+      "inputs": {
+        "data": { "dtype": "int8", "shape": [2], "data": { "kind": "values", "values": [100, -7] } },
+        "indices": { "dtype": "int32", "shape": [2, 1], "data": { "kind": "values", "values": [0, 0] } },
+        "updates": { "dtype": "int8", "shape": [2], "data": { "kind": "values", "values": [20, 20] } }
+      },
+      "outputs": {
+        "output": { "dtype": "int8", "shape": [2], "tolerance": 0, "data": { "kind": "values", "values": [64, -7] } }
+      }
+    },
+    {
+      "name": "ort_uint8_reduction_add_wrap8",
+      "provenance": {
+        "source": "ONNX Runtime CPUExecutionProvider",
+        "notes": "Covers standard UINT8 payload support and modulo-256 add reduction."
+      },
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": { "dtype": "uint8", "shape": [2], "data": { "kind": "values", "values": [250, 7] } },
+        "indices": { "dtype": "int32", "shape": [2, 1], "data": { "kind": "values", "values": [0, 0] } },
+        "updates": { "dtype": "uint8", "shape": [2], "data": { "kind": "values", "values": [10, 10] } }
+      },
+      "outputs": {
+        "output": { "dtype": "uint8", "shape": [2], "tolerance": 0, "data": { "kind": "values", "values": [14, 7] } }
+      }
+    },
+    {
+      "name": "ort_int16_reduction_mul_wrap16",
+      "provenance": {
+        "source": "ONNX Runtime CPUExecutionProvider",
+        "notes": "Duplicate signed updates lock 16-bit two's-complement narrowing after multiplication."
+      },
+      "attrs": { "reduction": "mul" },
+      "inputs": {
+        "data": { "dtype": "int16", "shape": [2], "data": { "kind": "values", "values": [20000, -7] } },
+        "indices": { "dtype": "int32", "shape": [2, 1], "data": { "kind": "values", "values": [0, 0] } },
+        "updates": { "dtype": "int16", "shape": [2], "data": { "kind": "values", "values": [2, 2] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "int16",
+          "shape": [2],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [14464, -7] }
+        }
+      }
+    },
+    {
+      "name": "int8_reduction_add_high_contention_wrap",
+      "provenance": {
+        "notes": "3000 updates over 4 slots, every slot wrapping several times. The parallel narrow-int route reduces on the raw two's-complement bits and narrows once at the end; this pins that against the reference's narrow-every-step arithmetic at a contention level a two-update fixture cannot reach. Index and update cycles share a period so each slot's total is distinct, which also catches a mispaired index/update lane."
+      },
+      "attrs": { "reduction": "add" },
+      "inputs": {
+        "data": { "dtype": "int8", "shape": [4], "data": { "kind": "values", "values": [7, -100, 0, 3] } },
+        "indices": { "dtype": "int32", "shape": [3000, 1], "data": { "kind": "cycle", "values": [0, 1, 2, 3] } },
+        "updates": { "dtype": "int8", "shape": [3000], "data": { "kind": "cycle", "values": [5, -3, 1, 2] } }
+      },
+      "outputs": { "output": { "dtype": "int8", "shape": [4], "tolerance": 0 } }
+    },
+    {
+      "name": "int8_reduction_min_duplicate_indices",
+      "provenance": {
+        "notes": "Signed narrow min reduces through a native atomicMin on the i32-widened output. The update cycle length shares a factor with the index cycle so each slot sees a disjoint half of the values and the two slots reduce to different results, which a mispaired index/update lane would not reproduce; one slot is left untouched by any update."
+      },
+      "attrs": { "reduction": "min" },
+      "inputs": {
+        "data": { "dtype": "int8", "shape": [3], "data": { "kind": "values", "values": [100, -5, 42] } },
+        "indices": { "dtype": "int32", "shape": [512, 1], "data": { "kind": "cycle", "values": [0, 1] } },
+        "updates": { "dtype": "int8", "shape": [512], "data": { "kind": "cycle", "values": [9, -128, 7, 3] } }
+      },
+      "outputs": { "output": { "dtype": "int8", "shape": [3], "tolerance": 0 } }
+    },
+    {
+      "name": "uint8_reduction_max_duplicate_indices",
+      "provenance": {
+        "notes": "Unsigned narrow max reduces through a native atomicMax on the u32-widened output. Each slot sees a disjoint half of the update cycle, so the two slots reduce to different results and one of them is decided by the data rather than by an update."
+      },
+      "attrs": { "reduction": "max" },
+      "inputs": {
+        "data": { "dtype": "uint8", "shape": [3], "data": { "kind": "values", "values": [0, 200, 17] } },
+        "indices": { "dtype": "int32", "shape": [512, 1], "data": { "kind": "cycle", "values": [0, 1] } },
+        "updates": { "dtype": "uint8", "shape": [512], "data": { "kind": "cycle", "values": [3, 255, 9, 1] } }
+      },
+      "outputs": { "output": { "dtype": "uint8", "shape": [3], "tolerance": 0 } }
+    },
+    {
+      "name": "int16_reduction_max_duplicate_indices",
+      "provenance": {
+        "notes": "16-bit signed max over the same i32-widened atomic path as int8, at values no 8-bit fixture can represent."
+      },
+      "attrs": { "reduction": "max" },
+      "inputs": {
+        "data": { "dtype": "int16", "shape": [3], "data": { "kind": "values", "values": [-30000, 1, -7] } },
+        "indices": { "dtype": "int32", "shape": [512, 1], "data": { "kind": "cycle", "values": [0, 1] } },
+        "updates": { "dtype": "int16", "shape": [512], "data": { "kind": "cycle", "values": [-20000, 32767, -1, 5] } }
+      },
+      "outputs": { "output": { "dtype": "int16", "shape": [3], "tolerance": 0 } }
+    }
+  ]
+}