sync 2e7068faf55e

Browse files

Files changed (8) hide show

README.md +59 -0
build/webgpu/bench.json +67 -0
build/webgpu/binary-broadcast-vec4.wgsl.jinja +172 -0
build/webgpu/binary-broadcast.wgsl.jinja +120 -0
build/webgpu/binary-vec4.wgsl.jinja +32 -0
build/webgpu/manifest.json +234 -0
build/webgpu/metadata.json +20 -0
build/webgpu/test.json +719 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.Add
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 14
+## Description
+Performs elementwise binary addition of two tensors with multidirectional (NumPy-style) broadcasting support. The output has the same element type as the inputs.
+See the [ONNX `Add` spec](https://onnx.ai/onnx/operators/onnx__Add.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `A` | `a` | `T` | — | — | First operand. | required |
+| `B` | `b` | `T` | — | — | Second operand. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `C` | `c` | `T` | derived | broadcast result of `A` and `B` | Elementwise sum of `A` and `B`; has the same element type as the inputs. | required |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16`, `int32`, `uint32`, `int8`, `uint8` |
+## 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
+- [`binary-broadcast-vec4.wgsl.jinja`](build/webgpu/binary-broadcast-vec4.wgsl.jinja)
+- [`binary-broadcast.wgsl.jinja`](build/webgpu/binary-broadcast.wgsl.jinja)
+- [`binary-vec4.wgsl.jinja`](build/webgpu/binary-vec4.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.Add", { version: 1 });
+const { c } = await kernel({ a: { data: aData, shape: [] }, b: { data: bData, shape: [] } });
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,67 @@

+{
+  "op": "ai.onnx.Add",
+  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
+  "cases": [
+    {
+      "name": "add-f32-4m",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "count": 4194304 },
+      "inputs": {
+        "a": { "shape": [4194304], "dtype": "float32", "dist": "normal", "seed": 102, "scale": 2 },
+        "b": { "shape": [4194304], "dtype": "float32", "dist": "normal", "seed": 103, "scale": 2 }
+      },
+      "outputs": { "c": { "shape": [4194304], "dtype": "float32" } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 3" }]
+      }
+    },
+    {
+      "name": "add-f32-biasadd-2048x2048",
+      "preset": "model",
+      "vars": { "dtype": "float32" },
+      "inputs": {
+        "a": { "shape": [2048, 2048], "dtype": "float32", "dist": "normal", "seed": 202, "scale": 2 },
+        "b": { "shape": [2048], "dtype": "float32", "dist": "normal", "seed": 203, "scale": 2 }
+      },
+      "outputs": { "c": { "shape": [2048, 2048], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2048 * 2048 * 2 * dtypeBytes(args.dtype)" }] }
+    },
+    {
+      "name": "add-f32-channelscale-1x1024x64x64",
+      "preset": "model",
+      "vars": { "dtype": "float32" },
+      "inputs": {
+        "a": { "shape": [1, 1024, 64, 64], "dtype": "float32", "dist": "normal", "seed": 302, "scale": 2 },
+        "b": { "shape": [1, 1024, 1, 1], "dtype": "float32", "dist": "normal", "seed": 303, "scale": 2 }
+      },
+      "outputs": { "c": { "shape": [1, 1024, 64, 64], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "1024 * 64 * 64 * 2 * dtypeBytes(args.dtype)" }] }
+    },
+    {
+      "name": "add-f32-broadcast-scalar-fallback-4m",
+      "preset": "stress",
+      "vars": { "dtype": "float32", "count": 4190208 },
+      "inputs": {
+        "a": { "shape": [4096, 1023], "dtype": "float32", "dist": "normal", "seed": 142, "scale": 2 },
+        "b": { "shape": [1, 1023], "dtype": "float32", "dist": "normal", "seed": 143, "scale": 2 }
+      },
+      "outputs": { "c": { "shape": [4096, 1023], "dtype": "float32", "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "add-f32-odd-numel-same-shape-4m",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "count": 4194303 },
+      "inputs": {
+        "a": { "shape": [4194303], "dtype": "float32", "dist": "normal", "seed": 102, "scale": 2 },
+        "b": { "shape": [4194303], "dtype": "float32", "dist": "normal", "seed": 103, "scale": 2 }
+      },
+      "outputs": { "c": { "shape": [4194303], "dtype": "float32" } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 3" }]
+      }
+    }
+  ]
+}

build/webgpu/binary-broadcast-vec4.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,172 @@

+{% macro offset_fn(fn_name, opShape, opRank, op_same, op_numel, outShape, outRank, out_numel) %}
+fn {{ fn_name }}({% if out_numel != 0 and op_numel != 1 %}out_index: u32{% endif %}) -> u32 {
+{% if out_numel == 0 %}
+  return 0u;
+{% elif op_numel == 1 %}
+  return 0u;
+{% elif op_same %}
+  return out_index;
+{% else %}
+  var offset = 0u;
+{% for axis in range(outRank) %}
+{% set op_axis = axis - (outRank - opRank) %}
+{% if op_axis >= 0 and opShape[op_axis] != 1 %}
+{% set c_stride = namespace(value=1) %}
+{% for j in range(axis + 1, outRank) %}
+{% set c_stride.value = c_stride.value * outShape[j] %}
+{% endfor %}
+{% set op_stride = namespace(value=1) %}
+{% for j in range(op_axis + 1, opRank) %}
+{% set op_stride.value = op_stride.value * opShape[j] %}
+{% endfor %}
+{% if c_stride.value == 1 %}
+  let coord{{ axis }} = out_index % {{ outShape[axis] }}u;
+{% else %}
+  let coord{{ axis }} = (out_index / {{ c_stride.value }}u) % {{ outShape[axis] }}u;
+{% endif %}
+{% if op_stride.value == 1 %}
+  offset = offset + coord{{ axis }};
+{% else %}
+  offset = offset + coord{{ axis }} * {{ op_stride.value }}u;
+{% endif %}
+{% endif %}
+{% endfor %}
+  return offset;
+{% endif %}
+}
+{%- endmacro %}{% macro broadcast_offset_call(fn_name, opShape, outShape, out_index) %}
+{% set op_numel = namespace(value=1) %}
+{% for d in opShape %}{% set op_numel.value = op_numel.value * d %}{% endfor %}
+{% set out_numel = namespace(value=1) %}
+{% for d in outShape %}{% set out_numel.value = out_numel.value * d %}{% endfor %}
+{{ fn_name }}({% if out_numel.value != 0 and op_numel.value != 1 %}{{ out_index }}{% endif %})
+{%- endmacro %}
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+// Vec4 broadcast binary op. The same compiled-stride folding as
+// binary-broadcast (offsets are compile-time-strength-reducible), but each
+// thread writes a vec4 of the output. If the output's innermost axis is a
+// multiple of four, each operand uses one base offset. For an
+// odd innermost extent, a vec4 may cross a row boundary: same-shape operands
+// remain contiguous, while broadcast operands compute the four lane offsets
+// independently so their outer broadcast coordinates wrap correctly. Under
+// broadcasting each operand axis is either matched (== output) or 1, so the
+// innermost axis of each operand is exactly one of:
+//   - "scalar"     : the whole operand is one element  -> splat a[0]
+//   - "splat"      : innermost axis is 1 (broadcast)   -> splat a[offset(base)]
+//   - "contiguous" : innermost axis matches the output -> vec4 a[offset(base)..+3]
+// Outer-axis broadcasting is captured by offset(base). A fully same-shape
+// operand is vector-bound and loaded directly; broadcast operands remain
+// scalar-bound so their independent lane offsets stay addressable.
+{% set a_numel = namespace(value=1) %}
+{% for d in source.aShape %}
+{% set a_numel.value = a_numel.value * d %}
+{% endfor %}
+{% set b_numel = namespace(value=1) %}
+{% for d in source.bShape %}
+{% set b_numel.value = b_numel.value * d %}
+{% endfor %}
+{% set c_numel = namespace(value=1) %}
+{% for d in source.cShape %}
+{% set c_numel.value = c_numel.value * d %}
+{% endfor %}
+{% set a_same = namespace(value=(source.aRank == source.cRank)) %}
+{% if a_same.value %}
+{% for axis in range(source.cRank) %}
+{% if source.aShape[axis] != source.cShape[axis] %}
+{% set a_same.value = false %}
+{% endif %}
+{% endfor %}
+{% endif %}
+{% set b_same = namespace(value=(source.bRank == source.cRank)) %}
+{% if b_same.value %}
+{% for axis in range(source.cRank) %}
+{% if source.bShape[axis] != source.cShape[axis] %}
+{% set b_same.value = false %}
+{% endif %}
+{% endfor %}
+{% endif %}
+{% set a_inner = source.aShape[source.aRank - 1] if source.aRank >= 1 else 1 %}
+{% set b_inner = source.bShape[source.bRank - 1] if source.bRank >= 1 else 1 %}
+{% set c_inner = source.cShape[source.cRank - 1] %}
+{% set crosses_inner_rows = c_inner % 4 != 0 %}
+{% if a_numel.value == 1 %}{% set a_mode = "scalar" %}
+{% elif a_inner == 1 %}{% set a_mode = "splat" %}
+{% else %}{% set a_mode = "contig" %}{% endif %}
+{% if b_numel.value == 1 %}{% set b_mode = "scalar" %}
+{% elif b_inner == 1 %}{% set b_mode = "splat" %}
+{% else %}{% set b_mode = "contig" %}{% endif %}
+{% if not a_same.value and a_mode != "scalar" %}
+{{ offset_fn("a_offset", source.aShape, source.aRank, false, a_numel.value, source.cShape, source.cRank, c_numel.value) }}
+{% endif %}
+{% if not b_same.value and b_mode != "scalar" %}
+{{ offset_fn("b_offset", source.bShape, source.bRank, false, b_numel.value, source.cShape, source.cRank, c_numel.value) }}
+{% endif %}
+{% set is_int = scalar == "i32" or scalar == "u32" %}
+{% set acc = scalar if is_int else "f32" %}
+// Narrow integer operations wrap modulo the logical dtype width; int8/uint8
+// use i32/u32 storage.
+{% if source.cDtype == "int8" %}
+fn wrap_dtype(v: vec4<i32>) -> vec4<i32> { return (v << vec4<u32>(24u)) >> vec4<u32>(24u); }
+{% set wrap = "wrap_dtype" %}
+{% elif source.cDtype == "uint8" %}
+fn wrap_dtype(v: vec4<u32>) -> vec4<u32> { return v & vec4<u32>(0xFFu); }
+{% set wrap = "wrap_dtype" %}
+{% else %}
+{% set wrap = "" %}
+{% endif %}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+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 when the element count exceeds the
+  // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills the rest into y).
+  let i4 = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if (i4 >= params.count) {
+    return;
+  }
+{% set needsBase = c_numel.value != 0 and ((not a_same.value and a_mode != "scalar") or (not b_same.value and b_mode != "scalar")) %}
+{% if needsBase %}
+  let base = i4 * 4u;
+{% endif %}
+{% if a_same.value %}
+  let av = vec4<{{ acc }}>(a[i4]);
+{% elif a_mode == "scalar" %}
+  let av = vec4<{{ acc }}>({{ acc }}(a[0]));
+{% elif crosses_inner_rows %}
+  let av = vec4<{{ acc }}>(
+    {{ acc }}(a[{{ broadcast_offset_call("a_offset", source.aShape, source.cShape, "base") }}]),
+    {{ acc }}(a[{{ broadcast_offset_call("a_offset", source.aShape, source.cShape, "base + 1u") }}]),
+    {{ acc }}(a[{{ broadcast_offset_call("a_offset", source.aShape, source.cShape, "base + 2u") }}]),
+    {{ acc }}(a[{{ broadcast_offset_call("a_offset", source.aShape, source.cShape, "base + 3u") }}])
+  );
+{% elif a_mode == "splat" %}
+  let av = vec4<{{ acc }}>({{ acc }}(a[{{ broadcast_offset_call("a_offset", source.aShape, source.cShape, "base") }}]));
+{% else %}
+  let ao = {{ broadcast_offset_call("a_offset", source.aShape, source.cShape, "base") }};
+  let av = vec4<{{ acc }}>({{ acc }}(a[ao]), {{ acc }}(a[ao + 1u]), {{ acc }}(a[ao + 2u]), {{ acc }}(a[ao + 3u]));
+{% endif %}
+{% if b_same.value %}
+  let bv = vec4<{{ acc }}>(b[i4]);
+{% elif b_mode == "scalar" %}
+  let bv = vec4<{{ acc }}>({{ acc }}(b[0]));
+{% elif crosses_inner_rows %}
+  let bv = vec4<{{ acc }}>(
+    {{ acc }}(b[{{ broadcast_offset_call("b_offset", source.bShape, source.cShape, "base") }}]),
+    {{ acc }}(b[{{ broadcast_offset_call("b_offset", source.bShape, source.cShape, "base + 1u") }}]),
+    {{ acc }}(b[{{ broadcast_offset_call("b_offset", source.bShape, source.cShape, "base + 2u") }}]),
+    {{ acc }}(b[{{ broadcast_offset_call("b_offset", source.bShape, source.cShape, "base + 3u") }}])
+  );
+{% elif b_mode == "splat" %}
+  let bv = vec4<{{ acc }}>({{ acc }}(b[{{ broadcast_offset_call("b_offset", source.bShape, source.cShape, "base") }}]));
+{% else %}
+  let bo = {{ broadcast_offset_call("b_offset", source.bShape, source.cShape, "base") }};
+  let bv = vec4<{{ acc }}>({{ acc }}(b[bo]), {{ acc }}(b[bo + 1u]), {{ acc }}(b[bo + 2u]), {{ acc }}(b[bo + 3u]));
+{% endif %}
+  c[i4] = {{ wrap }}(vec4<{{ scalar }}>(av + bv));
+}

build/webgpu/binary-broadcast.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,120 @@

+{% macro flat_tail_open() %}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+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 when the element count exceeds the
+  // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills the rest into y).
+  let invocation = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+{% if source.itemsPerInvocation is defined %}
+  // Tail-safe scalar x4 keeps vector-like dispatch density without requiring
+  // the logical tensor length (or its storage binding) to be vec4 aligned.
+  let begin = invocation * {{ source.itemsPerInvocation }}u;
+  let end = min(begin + {{ source.itemsPerInvocation }}u, params.count);
+  for (var i = begin; i < end; i = i + 1u) {
+{%- else %}
+  let i = invocation;
+  if (i >= params.count) {
+    return;
+  }
+{%- endif %}
+{% endmacro %}
+{% macro flat_tail_close() %}
+{% if source.itemsPerInvocation is defined %}
+  }
+{% endif %}
+{% endmacro %}
+{% macro offset_fn(fn_name, opShape, opRank, op_same, op_numel, outShape, outRank, out_numel) %}
+fn {{ fn_name }}({% if out_numel != 0 and op_numel != 1 %}out_index: u32{% endif %}) -> u32 {
+{% if out_numel == 0 %}
+  return 0u;
+{% elif op_numel == 1 %}
+  return 0u;
+{% elif op_same %}
+  return out_index;
+{% else %}
+  var offset = 0u;
+{% for axis in range(outRank) %}
+{% set op_axis = axis - (outRank - opRank) %}
+{% if op_axis >= 0 and opShape[op_axis] != 1 %}
+{% set c_stride = namespace(value=1) %}
+{% for j in range(axis + 1, outRank) %}
+{% set c_stride.value = c_stride.value * outShape[j] %}
+{% endfor %}
+{% set op_stride = namespace(value=1) %}
+{% for j in range(op_axis + 1, opRank) %}
+{% set op_stride.value = op_stride.value * opShape[j] %}
+{% endfor %}
+{% if c_stride.value == 1 %}
+  let coord{{ axis }} = out_index % {{ outShape[axis] }}u;
+{% else %}
+  let coord{{ axis }} = (out_index / {{ c_stride.value }}u) % {{ outShape[axis] }}u;
+{% endif %}
+{% if op_stride.value == 1 %}
+  offset = offset + coord{{ axis }};
+{% else %}
+  offset = offset + coord{{ axis }} * {{ op_stride.value }}u;
+{% endif %}
+{% endif %}
+{% endfor %}
+  return offset;
+{% endif %}
+}
+{%- endmacro %}{% macro broadcast_offset_call(fn_name, opShape, outShape, out_index) %}
+{% set op_numel = namespace(value=1) %}
+{% for d in opShape %}{% set op_numel.value = op_numel.value * d %}{% endfor %}
+{% set out_numel = namespace(value=1) %}
+{% for d in outShape %}{% set out_numel.value = out_numel.value * d %}{% endfor %}
+{{ fn_name }}({% if out_numel.value != 0 and op_numel.value != 1 %}{{ out_index }}{% endif %})
+{%- endmacro %}{% macro broadcast_offset_fn(fn_name, opShape, opRank, outShape, outRank) %}
+{% set op_numel = namespace(value=1) %}
+{% for d in opShape %}
+{% set op_numel.value = op_numel.value * d %}
+{% endfor %}
+{% set out_numel = namespace(value=1) %}
+{% for d in outShape %}
+{% set out_numel.value = out_numel.value * d %}
+{% endfor %}
+{% set op_same = namespace(value=(opRank == outRank)) %}
+{% if op_same.value %}
+{% for axis in range(outRank) %}
+{% if opShape[axis] != outShape[axis] %}
+{% set op_same.value = false %}
+{% endif %}
+{% endfor %}
+{% endif %}
+{{ offset_fn(fn_name, opShape, opRank, op_same.value, op_numel.value, outShape, outRank, out_numel.value) }}
+{%- endmacro %}{% macro binary_broadcast_offsets() %}
+{{ broadcast_offset_fn("a_offset", source.aShape, source.aRank, source.cShape, source.cRank) }}
+{{ broadcast_offset_fn("b_offset", source.bShape, source.bRank, source.cShape, source.cRank) }}
+{%- endmacro %}
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+{{ binary_broadcast_offsets() }}
+{{ flat_tail_open() }}
+{% if scalar == "i32" or scalar == "u32" %}
+  let av = a[{{ broadcast_offset_call("a_offset", source.aShape, source.cShape, "i") }}];
+  let bv = b[{{ broadcast_offset_call("b_offset", source.bShape, source.cShape, "i") }}];
+  let r = av + bv;
+// Narrow integer operations wrap modulo the logical dtype width; int8/uint8
+// use i32/u32 storage.
+{% if source.cDtype == "int8" %}
+  c[i] = (r << 24u) >> 24u;
+{% elif source.cDtype == "uint8" %}
+  c[i] = r & 0xFFu;
+{% else %}
+  c[i] = r;
+{% endif %}
+{% else %}
+  let av = f32(a[{{ broadcast_offset_call("a_offset", source.aShape, source.cShape, "i") }}]);
+  let bv = f32(b[{{ broadcast_offset_call("b_offset", source.bShape, source.cShape, "i") }}]);
+  c[i] = {{ scalar }}(av + bv);
+{% endif %}
+{{ flat_tail_close() -}}
+}

build/webgpu/binary-vec4.wgsl.jinja ADDED Viewed

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

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+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 when the element count exceeds the
+  // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills the rest into y).
+  let i = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if (i >= params.count) {
+    return;
+  }
+  let av = a[i];
+  let bv = b[i];
+{% if scalar == "i32" or scalar == "u32" %}
+  let r = av + bv;
+// Narrow integer operations wrap modulo the logical dtype width; int8/uint8
+// use i32/u32 storage.
+{% if source.cDtype == "int8" %}
+  c[i] = (r << vec4<u32>(24u)) >> vec4<u32>(24u);
+{% elif source.cDtype == "uint8" %}
+  c[i] = r & vec4<u32>(0xFFu);
+{% else %}
+  c[i] = r;
+{% endif %}
+{% elif scalar == "f16" %}
+  c[i] = vec4<f16>(vec4<f32>(av) + vec4<f32>(bv));
+{% else %}
+  c[i] = av + bv;
+{% endif %}
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,234 @@

+{
+  "domain": "ai.onnx",
+  "name": "Add",
+  "sinceVersion": 14,
+  "description": "Performs elementwise binary addition of two tensors with multidirectional (NumPy-style) broadcasting support. The output has the same element type as the inputs.",
+  "inputs": [
+    { "role": "A", "dtype": "T", "description": "First operand." },
+    { "role": "B", "dtype": "T", "description": "Second operand." }
+  ],
+  "outputs": [
+    {
+      "role": "C",
+      "dtype": "T",
+      "rank": "max(ranks.A, ranks.B)",
+      "description": "Elementwise sum of `A` and `B`; has the same element type as the inputs.",
+      "shape": "broadcastShape(shapes.A, shapes.B)"
+    }
+  ],
+  "typeConstraints": { "T": ["float32", "float16", "int32", "uint32", "int8", "uint8"] },
+  "args": {
+    "a": { "kind": "tensor", "semantic": "A", "role": "input" },
+    "b": { "kind": "tensor", "semantic": "B", "role": "input" },
+    "c": { "kind": "tensor", "semantic": "C", "role": "output" }
+  },
+  "tunables": { "WORKGROUP_SIZE": 256 },
+  "variants": [
+    {
+      "id": "same_shape_vec4",
+      "priority": 20,
+      "when": ["sameShape(shapes.A, shapes.C)", "sameShape(shapes.B, shapes.C)", "numel(shapes.C) > 0", "numel(shapes.C) % 4 == 0", "f16Ok(dtypes.T)"],
+      "constants": {
+        "scalar": "dtypes.T",
+        "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
+        "usesF16": "dtypes.T == \"f16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Add.vec4",
+          "source": { "shader": "binary-vec4.wgsl.jinja", "inputs": { "op": "\"add\"", "cDtype": "tensorDtypes.C" } },
+          "bindings": [
+            {
+              "name": "a",
+              "arg": "a",
+              "semantic": "A",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$vectorScalar"
+            },
+            {
+              "name": "b",
+              "arg": "b",
+              "semantic": "B",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$vectorScalar"
+            },
+            {
+              "name": "c",
+              "arg": "c",
+              "semantic": "C",
+              "buffer": { "type": "storage" },
+              "elementType": "$vectorScalar"
+            },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": {
+                "name": "Params",
+                "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.C) / 4" }]
+              }
+            }
+          ],
+          "dispatch": { "threads": "numel(shapes.C) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "broadcast_vec4",
+      "when": ["ranks.A <= ranks.C", "ranks.B <= ranks.C", "ranks.C >= 1", "numel(shapes.C) % 4 == 0", "numel(shapes.C) >= 4", "f16Ok(dtypes.T)"],
+      "constants": {
+        "scalar": "dtypes.T",
+        "aElement": "\"vec4<\" ~ dtypes.T ~ \">\" if sameShape(shapes.A, shapes.C) else dtypes.T",
+        "bElement": "\"vec4<\" ~ dtypes.T ~ \">\" if sameShape(shapes.B, shapes.C) else dtypes.T",
+        "vec4Scalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
+        "usesF16": "dtypes.T == \"f16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Add",
+          "source": {
+            "shader": "binary-broadcast-vec4.wgsl.jinja",
+            "inputs": {
+              "aShape": "shapes.A",
+              "bShape": "shapes.B",
+              "cShape": "shapes.C",
+              "aRank": "ranks.A",
+              "bRank": "ranks.B",
+              "cRank": "ranks.C",
+              "op": "\"add\"",
+              "cDtype": "tensorDtypes.C"
+            }
+          },
+          "bindings": [
+            {
+              "name": "a",
+              "arg": "a",
+              "semantic": "A",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$aElement"
+            },
+            {
+              "name": "b",
+              "arg": "b",
+              "semantic": "B",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$bElement"
+            },
+            { "name": "c", "arg": "c", "semantic": "C", "buffer": { "type": "storage" }, "elementType": "$vec4Scalar" },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": {
+                "name": "Params",
+                "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.C) / 4" }]
+              }
+            }
+          ],
+          "dispatch": { "threads": "numel(shapes.C) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ],
+      "priority": 10
+    },
+    {
+      "id": "same_shape_scalar_x4",
+      "priority": 15,
+      "when": ["sameShape(shapes.A, shapes.C)", "sameShape(shapes.B, shapes.C)", "numel(shapes.C) > 0", "numel(shapes.C) % 4 != 0", "f16Ok(dtypes.T)"],
+      "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Add",
+          "source": {
+            "shader": "binary-broadcast.wgsl.jinja",
+            "inputs": {
+              "aShape": "shapes.A",
+              "bShape": "shapes.B",
+              "cShape": "shapes.C",
+              "aRank": "ranks.A",
+              "bRank": "ranks.B",
+              "cRank": "ranks.C",
+              "op": "\"add\"",
+              "cDtype": "tensorDtypes.C",
+              "itemsPerInvocation": 4
+            }
+          },
+          "bindings": [
+            {
+              "name": "a",
+              "arg": "a",
+              "semantic": "A",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$scalar"
+            },
+            {
+              "name": "b",
+              "arg": "b",
+              "semantic": "B",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$scalar"
+            },
+            { "name": "c", "arg": "c", "semantic": "C", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.C)" }] }
+            }
+          ],
+          "dispatch": { "threads": "ceilDiv(numel(shapes.C), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "broadcast",
+      "when": ["ranks.A <= ranks.C", "ranks.B <= ranks.C", "f16Ok(dtypes.T)"],
+      "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Add",
+          "source": {
+            "shader": "binary-broadcast.wgsl.jinja",
+            "inputs": {
+              "aShape": "shapes.A",
+              "bShape": "shapes.B",
+              "cShape": "shapes.C",
+              "aRank": "ranks.A",
+              "bRank": "ranks.B",
+              "cRank": "ranks.C",
+              "op": "\"add\"",
+              "cDtype": "tensorDtypes.C"
+            }
+          },
+          "bindings": [
+            {
+              "name": "a",
+              "arg": "a",
+              "semantic": "A",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$scalar"
+            },
+            {
+              "name": "b",
+              "arg": "b",
+              "semantic": "B",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$scalar"
+            },
+            { "name": "c", "arg": "c", "semantic": "C", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.C)" }] }
+            }
+          ],
+          "dispatch": { "threads": "numel(shapes.C)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,20 @@

+{
+  "name": "ai.onnx.Add",
+  "id": "_ai_onnx_add_webgpu_a199e91",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "U0RcZ6cEwHaPzKmWjxkZIosg9WgaFSUCYGWbCefJhGY=",
+      "binary-broadcast-vec4.wgsl.jinja": "i2NgmeStm8zKVGrhtYRJoRwj1sT9hK7l23JOsnkSPaM=",
+      "binary-broadcast.wgsl.jinja": "JRtdDC2eKRvRIbOXPZ6/cnt1J6mD2xx/w+aBs5stGH8=",
+      "binary-vec4.wgsl.jinja": "mnGHNzK6EqW5t1KCp51LQgMM5tG3ufJTHiAGLGNtjqg=",
+      "manifest.json": "g9eLFcjl3csJ5qlpqOuRrAcpVBzLhRlrbZa2LHEH/tE=",
+      "test.json": "X1RwOC4YQxALtpaNkTZWgM4UQQqNU7llruve1sWX3uo="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Add" }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,719 @@

+{
+  "op": "ai.onnx.Add",
+  "fixtureArrays": {
+    "onnx_backend_add_input_a": [1.764052391052246, 0.40015721321105957, 0.978738009929657, 2.2408931255340576, 1.8675580024719238, -0.9772778749465942, 0.9500884413719177, -0.15135720372200012, -0.10321885347366333, 0.4105985164642334, 0.14404356479644775, 1.4542734622955322, 0.7610377073287964, 0.12167501449584961, 0.44386324286460876, 0.3336743414402008, 1.4940791130065918, -0.2051582634449005, 0.3130677044391632, -0.8540957570075989, -2.5529897212982178, 0.653618574142456, 0.8644362092018127, -0.7421650290489197, 2.269754648208618, -1.4543657302856445, 0.04575851559638977, -0.18718385696411133, 1.5327792167663574, 1.4693588018417358, 0.154947429895401, 0.37816253304481506, -0.8877857327461243, -1.980796456336975, -0.34791216254234314, 0.15634897351264954, 1.2302906513214111, 1.202379822731018, -0.38732680678367615, -0.302302747964859, -1.0485529899597168, -1.420017957687378, -1.7062702178955078, 1.950775384902954, -0.5096521973609924, -0.4380742907524109, -1.2527953386306763, 0.7774903774261475, -1.6138978004455566, -0.21274028718471527, -0.8954665660858154, 0.38690251111984253, -0.5108051300048828, -1.18063223361969, -0.02818222902715206, 0.4283318817615509, 0.06651721894741058, 0.30247190594673157, -0.6343221068382263, -0.3627411723136902]
+  },
+  "cases": [
+    {
+      "name": "vector_19",
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [19], "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29 } },
+        "b": { "dtype": "float32", "shape": [19], "data": { "kind": "fillFloat32", "sinStep": 0.31, "cosStep": 0.07 } }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [19], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "f32_subnormal_identity_and_cancellation_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 float32 subnormals to zero; the preserved-subnormal lanes cannot be reproduced on GPU."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Add_float",
+        "notes": "Finite subnormal addends are valid float32 values; zero-add identity and exact opposite cancellation should not flush unrelated lanes."
+      },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1e-40, -1e-40, 0.0, 1e-39] } },
+        "b": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.0, 0.0, 1e-40, -1e-39] } }
+      },
+      "outputs": {
+        "c": {
+          "dtype": "float32",
+          "shape": [4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [1e-40, -1e-40, 1e-40, 0.0] }
+        }
+      }
+    },
+    {
+      "name": "f32_subnormal_identity_and_cancellation_scalar_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 float32 subnormals to zero; the preserved-subnormal lanes cannot be reproduced on GPU."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Add_float",
+        "notes": "Scalar-path companion: zero-add identity should preserve finite subnormal lanes while exact opposites cancel."
+      },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1e-40, -1e-40, 1e-39] } },
+        "b": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.0, 0.0, -1e-39] } }
+      },
+      "outputs": {
+        "c": {
+          "dtype": "float32",
+          "shape": [3],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [1e-40, -1e-40, 0.0] }
+        }
+      }
+    },
+    {
+      "name": "float16_vec4_same_shape",
+      "inputs": {
+        "a": {
+          "dtype": "float16",
+          "shape": [8],
+          "data": { "kind": "values", "values": [1.0, -2.0, 3.5, -4.0, 0.25, 10.0, -100.0, 0.001] }
+        },
+        "b": {
+          "dtype": "float16",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0.5, 2.0, -1.5, 4.0, 0.75, -5.0, 100.0, -0.001] }
+        }
+      },
+      "outputs": { "c": { "dtype": "float16", "shape": [8], "tolerance": 0.001 } }
+    },
+    {
+      "name": "scalar_rhs_broadcast",
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [17], "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29 } },
+        "b": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.25] } }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [17], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "row_vector_rank4_broadcast",
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.19 }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [1, 1, 1, 5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.11, "scale": 0.2 }
+        }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [2, 3, 4, 5], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank7_alternating_broadcast",
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [2, 1, 3, 1, 2, 1, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.05, "cosStep": 0.17 }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [1, 4, 1, 5, 1, 3, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.11, "scale": 0.4 }
+        }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [2, 4, 3, 5, 2, 3, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "true_scalar_lhs_broadcast",
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.25] } },
+        "b": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [0.0, 1.0, -2.0, 3.5, 4.0, -5.0] }
+        }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [2, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank0_scalar_scalar_output",
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-1.25] } },
+        "b": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [3.5] } }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_broadcast_scalar_lhs_to_length1",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Add_Broadcast_0x1"
+      },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [10.0] } },
+        "b": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [2.0] } }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [1], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_broadcast_scalar_rhs_to_length1",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Add_Broadcast_1x0"
+      },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [10.0] } },
+        "b": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [1], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_broadcast_length1_length1",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Add_Broadcast_1x1"
+      },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [10.0] } },
+        "b": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [2.0] } }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [1], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_broadcast_3x2_3x1",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Add_Broadcast_3x2_3x1"
+      },
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [3, 2],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        },
+        "b": { "dtype": "float32", "shape": [3, 1], "data": { "kind": "values", "values": [1.0, 2.0, 3.0] } }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [3, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_broadcast_2x2x2_1x2x2",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Add_Broadcast_2x2x2_1x2x2"
+      },
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [2, 2, 2],
+          "data": { "kind": "values", "values": [101.0, 102.0, 103.0, 104.0, 201.0, 202.0, 203.0, 204.0] }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [1, 2, 2],
+          "data": { "kind": "values", "values": [10.0, 20.0, 30.0, 40.0] }
+        }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [2, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_broadcast_2x1x1_3x4",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Add_Broadcast_2x1x1_3x4"
+      },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [2, 1, 1], "data": { "kind": "values", "values": [100.0, 200.0] } },
+        "b": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": {
+            "kind": "values",
+            "values": [11.0, 12.0, 13.0, 14.0, 21.0, 22.0, 23.0, 24.0, 31.0, 32.0, 33.0, 34.0]
+          }
+        }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [2, 3, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "int32_exact_above_float24",
+      "inputs": {
+        "a": {
+          "dtype": "int32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [16777217, -16777217, 123456789, -123456789] }
+        },
+        "b": { "dtype": "int32", "shape": [4], "data": { "kind": "values", "values": [2, -2, 10, -10] } }
+      },
+      "outputs": { "c": { "dtype": "int32", "shape": [4] } }
+    },
+    {
+      "name": "ort_float_3x3",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Add_float"
+      },
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [3, 3],
+          "data": { "kind": "values", "values": [1.0, 2.0, -1.0, 0.0, 1.5, -100.0, -5.4, 9.3, -10000.0] }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [3, 3],
+          "data": { "kind": "values", "values": [-1.0, 4.4, 432.3, 0.0, 3.5, 64.0, -5.4, 9.3, 10000.0] }
+        }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [3, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_int8_vector",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Add_int8"
+      },
+      "inputs": {
+        "a": { "dtype": "int8", "shape": [3], "data": { "kind": "values", "values": [1, -2, 3] } },
+        "b": { "dtype": "int8", "shape": [3], "data": { "kind": "values", "values": [4, 5, 6] } }
+      },
+      "outputs": { "c": { "dtype": "int8", "shape": [3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int8_overflow_wrap",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Add_int8",
+        "notes": "Extends ORT's int8 Add coverage with signed overflow values that must wrap to logical int8 width."
+      },
+      "inputs": {
+        "a": { "dtype": "int8", "shape": [3], "data": { "kind": "values", "values": [127, -128, 100] } },
+        "b": { "dtype": "int8", "shape": [3], "data": { "kind": "values", "values": [1, -1, 30] } }
+      },
+      "outputs": { "c": { "dtype": "int8", "shape": [3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int32_vector",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Add_int32"
+      },
+      "inputs": {
+        "a": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [1, 2, 3] } },
+        "b": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [4, 5, 6] } }
+      },
+      "outputs": { "c": { "dtype": "int32", "shape": [3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_uint8_vector",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Add_uint8"
+      },
+      "inputs": {
+        "a": { "dtype": "uint8", "shape": [3], "data": { "kind": "values", "values": [1, 2, 3] } },
+        "b": { "dtype": "uint8", "shape": [3], "data": { "kind": "values", "values": [4, 5, 6] } }
+      },
+      "outputs": { "c": { "dtype": "uint8", "shape": [3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_uint8_overflow_wrap",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Add_uint8",
+        "notes": "Extends ORT's uint8 Add coverage with overflow values that must wrap to logical uint8 width."
+      },
+      "inputs": {
+        "a": { "dtype": "uint8", "shape": [3], "data": { "kind": "values", "values": [255, 0, 200] } },
+        "b": { "dtype": "uint8", "shape": [3], "data": { "kind": "values", "values": [1, 255, 100] } }
+      },
+      "outputs": { "c": { "dtype": "uint8", "shape": [3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_uint32_vector",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Add_uint32"
+      },
+      "inputs": {
+        "a": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [1, 2, 3] } },
+        "b": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [4, 5, 6] } }
+      },
+      "outputs": { "c": { "dtype": "uint32", "shape": [3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_broadcast_axis_3x3_3x1",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Add_Broadcast_Axis"
+      },
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [3, 3],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0] }
+        },
+        "b": { "dtype": "float32", "shape": [3, 1], "data": { "kind": "values", "values": [3.0, 2.0, 1.0] } }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [3, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_broadcast_multidirectional_ab",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Add_Broadcast_MultidirectionalAB"
+      },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [3, 1], "data": { "kind": "values", "values": [3.0, 2.0, 1.0] } },
+        "b": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 2.0, 3.0] } }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [3, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_broadcast_multidirectional_ba",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Add_Broadcast_MultidirectionalBA"
+      },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 2.0, 3.0] } },
+        "b": { "dtype": "float32", "shape": [3, 1], "data": { "kind": "values", "values": [3.0, 2.0, 1.0] } }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [3, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_broadcast_scalar_scalar",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Add_Broadcast_0x0"
+      },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [10.0] } },
+        "b": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_broadcast_2x1x4_1x3x1",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Add_Broadcast_2x1x4_1x3x1"
+      },
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [2, 1, 4],
+          "data": { "kind": "values", "values": [101.0, 102.0, 103.0, 104.0, 201.0, 202.0, 203.0, 204.0] }
+        },
+        "b": { "dtype": "float32", "shape": [1, 3, 1], "data": { "kind": "values", "values": [10.0, 20.0, 30.0] } }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [2, 3, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_dim_zero_equal_rank",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.DimWithZeroHandling",
+        "notes": "Float32 projection of ORT's equal-rank zero-dimension Add broadcast case."
+      },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [3, 1], "data": { "kind": "values", "values": [1.0, 2.0, 3.0] } },
+        "b": { "dtype": "float32", "shape": [3, 0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [3, 0], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_dim_zero_smaller_rank",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.DimWithZeroHandling",
+        "notes": "Float32 projection of ORT's smaller-rank zero-dimension Add broadcast case."
+      },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [2, 1, 2], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] } },
+        "b": { "dtype": "float32", "shape": [0, 2], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [2, 0, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_dim_zero_larger_rank",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.DimWithZeroHandling",
+        "notes": "Float32 projection of ORT's larger-rank zero-dimension Add broadcast case."
+      },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [0, 2, 2], "data": { "kind": "values", "values": [] } },
+        "b": { "dtype": "float32", "shape": [1, 2], "data": { "kind": "values", "values": [1.0, 2.0] } }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [0, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_dim_zero_scalar_broadcast",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.DimWithZeroHandling",
+        "notes": "Float32 projection of ORT's scalar-to-zero-dimension Add broadcast case."
+      },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] } },
+        "b": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [0], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_add_bcast",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_add_bcast" },
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_add_input_a" } }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": {
+            "kind": "values",
+            "values": [-0.6724604368209839, -0.35955315828323364, -0.8131462931632996, -1.7262825965881348, 0.17742614448070526]
+          }
+        }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [3, 4, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_add",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_add", "test": "test_add" },
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_add_input_a" } }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [-0.6724604368209839, -0.35955315828323364, -0.8131462931632996, -1.7262825965881348, 0.17742614448070526, -0.4017809331417084, -1.630198359489441, 0.46278226375579834, -0.9072983860969543, 0.05194539576768875, 0.7290905714035034, 0.12898291647434235, 1.1394007205963135, -1.234825849533081, 0.4023416340351105, -0.6848101019859314, -0.8707971572875977, -0.5788496732711792, -0.3115525245666504, 0.056165341287851334, -1.1651498079299927, 0.9008265137672424, 0.4656624495983124, -1.5362436771392822, 1.4882521629333496, 1.895889163017273, 1.1787796020507812, -0.1799248307943344, -1.0707526206970215, 1.0544517040252686, -0.4031769335269928, 1.222445011138916, 0.2082749754190445, 0.9766390323638916, 0.3563663959503174, 0.7065731883049011, 0.01050002034753561, 1.7858705520629883, 0.12691208720207214, 0.4019893705844879, 1.8831506967544556, -1.3477590084075928, -1.2704850435256958, 0.969396710395813, -1.1731233596801758, 1.9436211585998535, -0.4136189818382263, -0.747454822063446, 1.922942042350769, 1.4805147647857666, 1.8675589561462402, 0.9060446619987488, -0.8612256646156311, 1.910064935684204, -0.26800337433815, 0.8024563789367676, 0.9472519755363464, -0.15501008927822113, 0.6140793561935425, 0.922206699848175]
+          }
+        }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [3, 4, 5], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "onnx_backend_add_int8",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_add_int8",
+        "test": "test_add_int8"
+      },
+      "inputs": {
+        "a": {
+          "dtype": "int8",
+          "shape": [3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [0, 8, 11, 8, 22, 3, 11, 20, 17, 10, 4, 22, 22, 22, 0, 11, 15, 20, 5, 21, 20, 23, 2, 6, 22, 7, 15, 2, 8, 9, 2, 3, 9, 22, 0, 2, 19, 15, 21, 4, 22, 10, 23, 10, 18, 21, 3, 15, 7, 19, 17, 23, 14, 4, 12, 20, 20, 12, 14, 21]
+          }
+        },
+        "b": {
+          "dtype": "int8",
+          "shape": [3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [22, 16, 17, 3, 21, 3, 3, 10, 8, 7, 3, 1, 8, 9, 13, 20, 17, 9, 8, 10, 9, 17, 10, 18, 23, 15, 3, 20, 14, 15, 10, 11, 2, 18, 23, 19, 12, 12, 11, 3, 7, 17, 7, 9, 1, 4, 10, 13, 21, 14, 17, 13, 17, 2, 7, 0, 17, 3, 7, 5]
+          }
+        }
+      },
+      "outputs": { "c": { "dtype": "int8", "shape": [3, 4, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_add_uint8",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_add_uint8",
+        "test": "test_add_uint8"
+      },
+      "inputs": {
+        "a": {
+          "dtype": "uint8",
+          "shape": [3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [23, 21, 13, 9, 1, 9, 14, 7, 11, 12, 5, 17, 0, 22, 22, 9, 6, 23, 16, 0, 15, 6, 23, 22, 20, 9, 16, 4, 11, 8, 10, 9, 3, 4, 0, 0, 3, 2, 9, 5, 14, 20, 5, 2, 20, 15, 1, 3, 4, 12, 17, 19, 5, 8, 7, 8, 16, 10, 2, 3]
+          }
+        },
+        "b": {
+          "dtype": "uint8",
+          "shape": [3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [19, 22, 12, 19, 23, 20, 14, 10, 11, 5, 23, 20, 4, 1, 0, 0, 3, 0, 19, 8, 6, 8, 18, 22, 15, 7, 6, 6, 23, 4, 10, 12, 13, 12, 12, 12, 5, 2, 17, 22, 5, 15, 20, 12, 14, 20, 7, 10, 19, 23, 11, 4, 20, 0, 15, 11, 10, 0, 2, 11]
+          }
+        }
+      },
+      "outputs": { "c": { "dtype": "uint8", "shape": [3, 4, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_add_uint32",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_add_uint32",
+        "test": "test_add_uint32"
+      },
+      "inputs": {
+        "a": {
+          "dtype": "uint32",
+          "shape": [3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [17, 7, 20, 2, 16, 18, 13, 23, 4, 4, 12, 17, 23, 22, 20, 1, 5, 23, 1, 19, 21, 23, 19, 22, 22, 7, 9, 17, 9, 22, 0, 19, 4, 17, 20, 21, 16, 3, 10, 1, 20, 4, 4, 16, 16, 15, 10, 13, 8, 11, 20, 6, 9, 15, 3, 19, 12, 18, 17, 10]
+          }
+        },
+        "b": {
+          "dtype": "uint32",
+          "shape": [3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [2, 1, 11, 3, 12, 4, 1, 17, 0, 7, 4, 3, 5, 6, 3, 2, 8, 1, 4, 0, 13, 19, 9, 21, 5, 1, 7, 8, 22, 20, 7, 3, 21, 12, 3, 6, 20, 7, 19, 0, 5, 11, 9, 19, 7, 21, 5, 8, 11, 14, 0, 8, 19, 6, 10, 19, 2, 23, 0, 19]
+          }
+        }
+      },
+      "outputs": { "c": { "dtype": "uint32", "shape": [3, 4, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int16_basic_gpu_gap",
+      "skipGpu": {
+        "category": "todo",
+        "reason": "The widened-i32 Add path does not yet restore signed 16-bit wraparound after overflow; enable int16 only with explicit narrowing and overflow fixtures."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Add_int16"
+      },
+      "inputs": {
+        "a": { "dtype": "int16", "shape": [3], "data": { "kind": "values", "values": [1, -2, 3] } },
+        "b": { "dtype": "int16", "shape": [3], "data": { "kind": "values", "values": [4, 5, 6] } }
+      },
+      "outputs": { "c": { "dtype": "int16", "shape": [3], "tolerance": 0 } }
+    },
+    {
+      "name": "broadcast_vec4_biasadd_row",
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [3, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29 }
+        },
+        "b": { "dtype": "float32", "shape": [8], "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.19 } }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [3, 8], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "broadcast_vec4_channelscale_inner",
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [1, 3, 2, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23 }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [1, 3, 1, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.05, "cosStep": 0.17 }
+        }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [1, 3, 2, 8], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "broadcast_vec4_scalar_operand",
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [2, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "b": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [2.5] } }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [2, 8], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "broadcast_vec4_f16_biasadd",
+      "inputs": {
+        "a": {
+          "dtype": "float16",
+          "shape": [3, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29 }
+        },
+        "b": { "dtype": "float16", "shape": [8], "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.19 } }
+      },
+      "outputs": { "c": { "dtype": "float16", "shape": [3, 8], "tolerance": 0.01 } }
+    },
+    {
+      "name": "broadcast_vec4_2d_fold_boundary",
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [65537, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.01, "cosStep": 0.02, "scale": 0.5 }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [65537, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.03, "cosStep": 0.05, "scale": 0.5 }
+        }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [65537, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "broadcast_vec4_rank_mismatch_outer_broadcast",
+      "inputs": {
+        "a": {
+          "dtype": "int32",
+          "shape": [3, 4],
+          "data": { "kind": "values", "values": [100, 200, 300, 400, 500, 600, 700, 800, 900, 1000, 1100, 1200] }
+        },
+        "b": {
+          "dtype": "int32",
+          "shape": [4, 3, 4],
+          "data": {
+            "kind": "values",
+            "values": [1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31, 32, 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 43, 44, 45, 46, 47, 48]
+          }
+        }
+      },
+      "outputs": { "c": { "dtype": "int32", "shape": [4, 3, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "f16_result_overflows_to_inf",
+      "inputs": {
+        "a": {
+          "dtype": "float16",
+          "shape": [4],
+          "data": { "kind": "values", "values": [60000.0, -60000.0, 32768.0, 0.0] }
+        },
+        "b": {
+          "dtype": "float16",
+          "shape": [4],
+          "data": { "kind": "values", "values": [10000.0, -10000.0, 32768.0, 65504.0] }
+        }
+      },
+      "outputs": { "c": { "dtype": "float16", "shape": [4], "tolerance": 0, "allowNaN": false } }
+    },
+    {
+      "name": "int8_broadcast_vec4_wrap_splat",
+      "inputs": {
+        "a": {
+          "dtype": "int8",
+          "shape": [2, 8],
+          "data": {
+            "kind": "values",
+            "values": [120, 121, 122, 123, 124, 125, 126, 127, -128, -127, -126, -125, -124, -123, -122, -121]
+          }
+        },
+        "b": { "dtype": "int8", "shape": [2, 1], "data": { "kind": "values", "values": [10, -10] } }
+      },
+      "outputs": {
+        "c": {
+          "dtype": "int8",
+          "shape": [2, 8],
+          "tolerance": 0,
+          "data": {
+            "kind": "values",
+            "values": [-126, -125, -124, -123, -122, -121, -120, -119, 118, 119, 120, 121, 122, 123, 124, 125]
+          }
+        }
+      }
+    }
+  ]
+}