sync 91d990483a17

Browse files

Files changed (8) hide show

README.md +12 -8
build/webgpu/bench.json +14 -1
build/webgpu/binary-broadcast-vec4.wgsl.jinja +33 -37
build/webgpu/binary-broadcast.wgsl.jinja +20 -16
build/webgpu/binary-vec4.wgsl.jinja +37 -8
build/webgpu/manifest.json +148 -186
build/webgpu/metadata.json +19 -9
build/webgpu/test.json +41 -4

README.md CHANGED Viewed

@@ -18,16 +18,16 @@ See the [ONNX `Add` spec](https://onnx.ai/onnx/operators/onnx__Add.html) for the
 ## 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
@@ -37,7 +37,7 @@ See the [ONNX `Add` spec](https://onnx.ai/onnx/operators/onnx__Add.html) for the
 ## 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
@@ -47,10 +47,14 @@ See the [ONNX `Add` spec](https://onnx.ai/onnx/operators/onnx__Add.html) for the
 ## 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.

 ## Inputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
+| `a` | `A` | `T` | — | — | First operand. | required |
+| `b` | `B` | `T` | — | — | Second operand. | required |
 ## Outputs
+| Name | Upstream name | 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
 ## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, 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
 ## Use with `@huggingface/kernels`
+```sh
+npm install --save-exact @huggingface/kernels@0.0.1-preview.2
+```
+Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
 The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+It follows the `v1` branch as fixes land. To pin exact artifact bytes, pass a 40-character commit `revision` instead of `version`.
 Replace each `*Data` placeholder with a typed array containing the corresponding input data.

build/webgpu/bench.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "ai.onnx.Add",
   "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
   "cases": [
     {
@@ -16,6 +15,20 @@
         "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 3" }]
       }
     },
     {
       "name": "add-f32-biasadd-2048x2048",
       "preset": "model",

 {
   "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
   "cases": [
     {
         "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 3" }]
       }
     },
+    {
+      "name": "add-f32-8m-dram-bound",
+      "preset": "stress",
+      "vars": { "dtype": "float32", "count": 8388608 },
+      "inputs": {
+        "a": { "shape": [8388608], "dtype": "float32", "dist": "normal", "seed": 102, "scale": 2 },
+        "b": { "shape": [8388608], "dtype": "float32", "dist": "normal", "seed": 103, "scale": 2 }
+      },
+      "outputs": { "c": { "shape": [8388608], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 3" }] },
+      "provenance": {
+        "notes": "A 32 MB output lies beyond the 16 MB threshold where the vec4 kernel changes from four groups per lane to one, exercising the large-output DRAM-bound configuration."
+      }
+    },
     {
       "name": "add-f32-biasadd-2048x2048",
       "preset": "model",

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

@@ -42,14 +42,10 @@ fn {{ fn_name }}({% if out_numel != 0 and op_numel != 1 %}out_index: u32{% endif
 {{ 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
@@ -63,36 +59,36 @@ enable f16;
 // 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" %}
@@ -102,31 +98,31 @@ enable f16;
 {% 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;
   }
@@ -140,15 +136,15 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
   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 %}
@@ -157,15 +153,15 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
   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));

 {{ fn_name }}({% if out_numel.value != 0 and op_numel.value != 1 %}{{ out_index }}{% endif %})
 {%- endmacro %}
 {{ env.wgsl.resourceDeclarations }}
+// Vec4 broadcast binary op. Offsets use compile-time strides, and each
+// invocation writes four output values. 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
 // 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 aShape %}
 {% set a_numel.value = a_numel.value * d %}
 {% endfor %}
 {% set b_numel = namespace(value=1) %}
+{% for d in bShape %}
 {% set b_numel.value = b_numel.value * d %}
 {% endfor %}
 {% set c_numel = namespace(value=1) %}
+{% for d in cShape %}
 {% set c_numel.value = c_numel.value * d %}
 {% endfor %}
+{% set a_same = namespace(value=(aRank == cRank)) %}
 {% if a_same.value %}
+{% for axis in range(cRank) %}
+{% if aShape[axis] != cShape[axis] %}
 {% set a_same.value = false %}
 {% endif %}
 {% endfor %}
 {% endif %}
+{% set b_same = namespace(value=(bRank == cRank)) %}
 {% if b_same.value %}
+{% for axis in range(cRank) %}
+{% if bShape[axis] != cShape[axis] %}
 {% set b_same.value = false %}
 {% endif %}
 {% endfor %}
 {% endif %}
+{% set a_inner = aShape[aRank - 1] if aRank >= 1 else 1 %}
+{% set b_inner = bShape[bRank - 1] if bRank >= 1 else 1 %}
+{% set c_inner = cShape[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 b_mode = "contig" %}{% endif %}
 {% if not a_same.value and a_mode != "scalar" %}
+{{ offset_fn("a_offset", aShape, aRank, false, a_numel.value, cShape, cRank, c_numel.value) }}
 {% endif %}
 {% if not b_same.value and b_mode != "scalar" %}
+{{ offset_fn("b_offset", bShape, bRank, false, b_numel.value, cShape, 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 cDtype == "int8" %}
 fn wrap_dtype(v: vec4<i32>) -> vec4<i32> { return (v << vec4<u32>(24u)) >> vec4<u32>(24u); }
 {% set wrap = "wrap_dtype" %}
+{% elif 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>) {
   // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
+  // per-axis dispatch fold width (the dispatch caps x and spills the rest into y).
+  let i4 = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if (i4 >= params.count) {
     return;
   }
   let av = vec4<{{ acc }}>({{ acc }}(a[0]));
 {% elif crosses_inner_rows %}
   let av = vec4<{{ acc }}>(
+    {{ acc }}(a[{{ broadcast_offset_call("a_offset", aShape, cShape, "base") }}]),
+    {{ acc }}(a[{{ broadcast_offset_call("a_offset", aShape, cShape, "base + 1u") }}]),
+    {{ acc }}(a[{{ broadcast_offset_call("a_offset", aShape, cShape, "base + 2u") }}]),
+    {{ acc }}(a[{{ broadcast_offset_call("a_offset", aShape, cShape, "base + 3u") }}])
   );
 {% elif a_mode == "splat" %}
+  let av = vec4<{{ acc }}>({{ acc }}(a[{{ broadcast_offset_call("a_offset", aShape, cShape, "base") }}]));
 {% else %}
+  let ao = {{ broadcast_offset_call("a_offset", aShape, 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 }}>({{ acc }}(b[0]));
 {% elif crosses_inner_rows %}
   let bv = vec4<{{ acc }}>(
+    {{ acc }}(b[{{ broadcast_offset_call("b_offset", bShape, cShape, "base") }}]),
+    {{ acc }}(b[{{ broadcast_offset_call("b_offset", bShape, cShape, "base + 1u") }}]),
+    {{ acc }}(b[{{ broadcast_offset_call("b_offset", bShape, cShape, "base + 2u") }}]),
+    {{ acc }}(b[{{ broadcast_offset_call("b_offset", bShape, cShape, "base + 3u") }}])
   );
 {% elif b_mode == "splat" %}
+  let bv = vec4<{{ acc }}>({{ acc }}(b[{{ broadcast_offset_call("b_offset", bShape, cShape, "base") }}]));
 {% else %}
+  let bo = {{ broadcast_offset_call("b_offset", bShape, 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 CHANGED Viewed

@@ -1,13 +1,14 @@
 {% 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;
   // 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) {
 {%- endmacro %}
 {% macro flat_tail_close() %}
@@ -75,14 +76,17 @@ fn {{ fn_name }}({% if out_numel != 0 and op_numel != 1 %}out_index: u32{% endif
 {% 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 }}
@@ -90,21 +94,21 @@ enable f16;
 {{ 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() -}}

 {% macro flat_tail_open() %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
+  // dispatch's per-axis workgroup fold width (the dispatch caps x and spills the rest into y).
+  let invocation = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   // Tail-safe scalar x4 keeps vector-like dispatch density without requiring
   // the logical tensor length (or its storage binding) to be vec4 aligned.
+{% set itemsPerInvocation = itemsPerInvocation if itemsPerInvocation is defined else 4 %}
+  let begin = invocation * {{ itemsPerInvocation }}u;
+  let end = min(begin + {{ itemsPerInvocation }}u, params.count);
   for (var i = begin; i < end; i = i + 1u) {
 {%- endmacro %}
 {% macro flat_tail_close() %}
 {% endif %}
 {{ offset_fn(fn_name, opShape, opRank, op_same.value, op_numel.value, outShape, outRank, out_numel.value) }}
 {%- endmacro %}{% macro binary_broadcast_offsets() %}
+{% set aShape = aShape | default([]) %}
+{% set aRank = aRank | default(0) %}
+{% set bShape = bShape | default([]) %}
+{% set bRank = bRank | default(0) %}
+{% set cShape = cShape | default([]) %}
+{% set cRank = cRank | default(0) %}
+{{ broadcast_offset_fn("a_offset", aShape, aRank, cShape, cRank) }}
+{{ broadcast_offset_fn("b_offset", bShape, bRank, cShape, cRank) }}
 {%- endmacro %}
 {{ env.wgsl.resourceDeclarations }}
 {{ flat_tail_open() }}
 {% if scalar == "i32" or scalar == "u32" %}
+  let av = a[{{ broadcast_offset_call("a_offset", aShape, cShape, "i") }}];
+  let bv = b[{{ broadcast_offset_call("b_offset", bShape, cShape, "i") }}];
   let r = av + bv;
 // Narrow integer operations wrap modulo the logical dtype width; int8/uint8
 // use i32/u32 storage.
+{% if cDtype == "int8" %}
   c[i] = (r << 24u) >> 24u;
+{% elif cDtype == "uint8" %}
   c[i] = r & 0xFFu;
 {% else %}
   c[i] = r;
 {% endif %}
 {% else %}
+  let av = f32(a[{{ broadcast_offset_call("a_offset", aShape, cShape, "i") }}]);
+  let bv = f32(b[{{ broadcast_offset_call("b_offset", bShape, cShape, "i") }}]);
   c[i] = {{ scalar }}(av + bv);
 {% endif %}
 {{ flat_tail_close() -}}

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

@@ -1,25 +1,51 @@
-{% 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;
@@ -29,4 +55,7 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
 {% else %}
   c[i] = av + bv;
 {% endif %}
 }

 {{ env.wgsl.resourceDeclarations }}
+{% set vec4PerThread = vec4PerThread %}
+{% if vec4PerThread > 1 %}
+const ITEMS: u32 = {{ vec4PerThread }}u;
+{% endif %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
+  // per-axis dispatch fold width (the dispatch caps x and spills the rest into y).
+{% if vec4PerThread > 1 %}
+  // Each invocation walks ITEMS vec4 groups a span apart. Consecutive lanes
+  // access consecutive words on every step, while each lane can keep several
+  // independent loads in flight.
+  let tid = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
+  let span = (params.count + ITEMS - 1u) / ITEMS;
+  for (var j = 0u; j < ITEMS; j = j + 1u) {
+    let i = tid + j * span;
+    if (i >= params.count) {
+      break;
+    }
+{% else %}
+  let i = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if (i >= params.count) {
     return;
   }
+{% endif %}
+{% set scalarOperand = scalarOperand if scalarOperand is defined else "" %}
+{% if scalarOperand == "a" %}
+  // One-element operand: read once and splat across the vector.
+  let av = {{ vectorScalar }}(a[0]);
+{% else %}
   let av = a[i];
+{% endif %}
+{% if scalarOperand == "b" %}
+  let bv = {{ vectorScalar }}(b[0]);
+{% else %}
   let bv = b[i];
+{% endif %}
 {% 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 cDtype == "int8" %}
   c[i] = (r << vec4<u32>(24u)) >> vec4<u32>(24u);
+{% elif cDtype == "uint8" %}
   c[i] = r & vec4<u32>(0xFFu);
 {% else %}
   c[i] = r;
 {% else %}
   c[i] = av + bv;
 {% endif %}
+{% if vec4PerThread > 1 %}
+  }
+{% endif %}
 }

build/webgpu/manifest.json CHANGED Viewed

@@ -2,234 +2,196 @@
   "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",
-              "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" }
         }
       ]
     }
-  ]
 }

   "domain": "ai.onnx",
   "name": "Add",
   "sinceVersion": 14,
+  "inputs": { "a": { "onnx": "A", "dtype": "T" }, "b": { "onnx": "B", "dtype": "T" } },
+  "outputs": {
+    "c": { "onnx": "C", "dtype": "T", "rank": "max(ranks.a, ranks.b)", "shape": "broadcastShape(shapes.a, shapes.b)" }
   },
+  "typeConstraints": { "T": ["float32", "float16", "int32", "uint32", "int8", "uint8"] },
+  "tunables": { "WORKGROUP_SIZE": { "default": 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)"],
+      "derive": { "scalar": "dtypes.T", "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "Add.vec4",
+          "shader": "binary-vec4.wgsl.jinja",
+          "derive": {
+            "op": "\"add\"",
+            "cDtype": "tensorDtypes.c",
+            "vec4PerThread": "4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1"
+          },
+          "bindings": ["a", "b", "c_binary", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.c) / 4, 4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.c) / 4, 4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
+        }
+      ]
+    },
+    {
+      "id": "scalar_b_vec4",
+      "priority": 18,
+      "when": ["sameShape(shapes.a, shapes.c)", "numel(shapes.b) == 1", "numel(shapes.c) > 0", "numel(shapes.c) % 4 == 0", "f16Ok(dtypes.T)"],
+      "derive": { "scalar": "dtypes.T", "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Add.scalarBVec4",
+          "shader": "binary-vec4.wgsl.jinja",
+          "derive": {
+            "op": "\"add\"",
+            "cDtype": "tensorDtypes.c",
+            "scalarOperand": "\"b\"",
+            "vec4PerThread": "4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1"
+          },
+          "bindings": ["a", "b_2", "c_binary", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.c) / 4, 4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.c) / 4, 4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
+        }
+      ]
+    },
+    {
+      "id": "scalar_a_vec4",
+      "priority": 18,
+      "when": ["sameShape(shapes.b, shapes.c)", "numel(shapes.a) == 1", "numel(shapes.c) > 0", "numel(shapes.c) % 4 == 0", "f16Ok(dtypes.T)"],
+      "derive": { "scalar": "dtypes.T", "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Add.scalarAVec4",
+          "shader": "binary-vec4.wgsl.jinja",
+          "derive": {
+            "op": "\"add\"",
+            "cDtype": "tensorDtypes.c",
+            "scalarOperand": "\"a\"",
+            "vec4PerThread": "4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1"
+          },
+          "bindings": ["a_2", "b", "c_binary", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.c) / 4, 4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.c) / 4, 4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "broadcast_vec4",
+      "priority": 10,
+      "when": ["ranks.a <= ranks.c", "ranks.b <= ranks.c", "ranks.c >= 1", "numel(shapes.c) % 4 == 0", "numel(shapes.c) >= 4", "f16Ok(dtypes.T)"],
+      "derive": {
         "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 ~ \">\""
       },
       "passes": [
         {
           "id": "main",
           "name": "Add",
+          "shader": "binary-broadcast-vec4.wgsl.jinja",
+          "derive": {
+            "aShape": "shapes.a",
+            "bShape": "shapes.b",
+            "cShape": "shapes.c",
+            "aRank": "ranks.a",
+            "bRank": "ranks.b",
+            "cRank": "ranks.c",
+            "op": "\"add\"",
+            "cDtype": "tensorDtypes.c"
           },
+          "bindings": ["a_3", "b_3", "c_2_binary", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.c) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.c) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
+      ]
     },
     {
       "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)"],
+      "derive": { "scalar": "dtypes.T" },
       "passes": [
         {
           "id": "main",
           "name": "Add",
+          "shader": "binary-broadcast.wgsl.jinja",
+          "derive": {
+            "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": ["a_2", "b_2", "c_3", "params_2"],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.c), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.c), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "broadcast",
+      "when": ["ranks.a <= ranks.c", "ranks.b <= ranks.c", "f16Ok(dtypes.T)"],
+      "derive": { "scalar": "dtypes.T" },
       "passes": [
         {
           "id": "main",
           "name": "Add",
+          "shader": "binary-broadcast.wgsl.jinja",
+          "derive": {
+            "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": ["a_2", "b_2", "c_3", "params_2"],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.c), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.c), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     }
+  ],
+  "bindings": {
+    "a": { "buffer": "read-only-storage", "elementType": "$vectorScalar" },
+    "b": { "buffer": "read-only-storage", "elementType": "$vectorScalar" },
+    "c_binary": { "buffer": "storage", "elementType": "$vectorScalar", "name": "c" },
+    "params": { "buffer": "uniform", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.c) / 4" }] },
+    "b_2": { "buffer": "read-only-storage", "name": "b", "elementType": "$scalar" },
+    "a_2": { "buffer": "read-only-storage", "name": "a", "elementType": "$scalar" },
+    "a_3": { "buffer": "read-only-storage", "name": "a", "elementType": "$aElement" },
+    "b_3": { "buffer": "read-only-storage", "name": "b", "elementType": "$bElement" },
+    "c_2_binary": { "buffer": "storage", "name": "c", "elementType": "$vec4Scalar" },
+    "c_3": { "buffer": "storage", "name": "c", "elementType": "$scalar" },
+    "params_2": {
+      "buffer": "uniform",
+      "name": "params",
+      "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.c)" }]
+    }
+  }
 }

build/webgpu/metadata.json CHANGED Viewed

@@ -1,20 +1,30 @@
 {
   "name": "ai.onnx.Add",
-  "id": "_ai_onnx_add_webgpu_c702c37",
   "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": "OARLrDsJW/gFuLicGept3Zg81d2iusgKd+FNBx4NqNU=",
-      "binary-vec4.wgsl.jinja": "mnGHNzK6EqW5t1KCp51LQgMM5tG3ufJTHiAGLGNtjqg=",
-      "manifest.json": "H+A5C9nJwxRWenku9NZ2kRYej0VXZ3x3sFai7RjQ5Io=",
-      "test.json": "X1RwOC4YQxALtpaNkTZWgM4UQQqNU7llruve1sWX3uo="
     }
   },
-  "provenance": { "kernel": { "sha": "c928d21e6cc1310861cba3bafb75f5f679ecf5f3", "dirty": false } },
-  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Add" }
 }

 {
   "name": "ai.onnx.Add",
+  "id": "_ai_onnx_add_webgpu_cf0ca6f",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "U6k4fuUAGJgoYWWbs5IbVMOK8X7eztmuSigDquudz5o=",
+      "binary-broadcast-vec4.wgsl.jinja": "tea2njOdNlRqQUK6jbsdepyNl/Lz9dmTycQZJOiuvzs=",
+      "binary-broadcast.wgsl.jinja": "GvnQZK2gkzwRTXhxCcOrZanBUvRSqPyqiTVMRF5a7gc=",
+      "binary-vec4.wgsl.jinja": "HowqRn81nC/ise93aa+XYqgA5oak9ddqfKbZ2wbVziI=",
+      "manifest.json": "63x5nCiu+fKZ1tFeDhC0Yhd1ziGqxw6Bv7KKrjGMlbQ=",
+      "test.json": "mwBGLiuw4ehiD/U0fa1hcnfLU0kUik2AVY19syL9Re4="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": {
+      "same_shape_vec4": ["binary-vec4.wgsl.jinja"],
+      "scalar_b_vec4": ["binary-vec4.wgsl.jinja"],
+      "scalar_a_vec4": ["binary-vec4.wgsl.jinja"],
+      "broadcast_vec4": ["binary-broadcast-vec4.wgsl.jinja"],
+      "same_shape_scalar_x4": ["binary-broadcast.wgsl.jinja"],
+      "broadcast": ["binary-broadcast.wgsl.jinja"]
+    }
+  }
 }

build/webgpu/test.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "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]
   },
@@ -45,7 +44,7 @@
       "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] } },
@@ -276,7 +275,7 @@
       "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] } },
@@ -313,7 +312,7 @@
       "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] } },
@@ -714,6 +713,44 @@
           }
         }
       }
     }
   ]
 }

 {
   "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]
   },
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
         "test": "MathOpTest.Add_float",
+        "notes": "On an unaligned scalar path, addition by zero must preserve finite subnormal values while exact opposites cancel."
       },
       "inputs": {
         "a": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1e-40, -1e-40, 1e-39] } },
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
         "test": "MathOpTest.Add_int8",
+        "notes": "Signed overflow must wrap to the logical int8 width."
       },
       "inputs": {
         "a": { "dtype": "int8", "shape": [3], "data": { "kind": "values", "values": [127, -128, 100] } },
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
         "test": "MathOpTest.Add_uint8",
+        "notes": "Unsigned overflow must wrap to the logical uint8 width."
       },
       "inputs": {
         "a": { "dtype": "uint8", "shape": [3], "data": { "kind": "values", "values": [255, 0, 200] } },
           }
         }
       }
+    },
+    {
+      "name": "scalar_b_vec4_route",
+      "provenance": {
+        "notes": "Route lock for the one-element-operand vec4 kernel: the other operand matches the output and the vector count is a multiple of four, so the scalar reads once and splats."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [2, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.23, "scale": 1.5 }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.31, "cosStep": 0.11, "scale": 1.5 }
+        }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [2, 8], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "scalar_a_vec4_route",
+      "provenance": {
+        "notes": "Route lock for the one-element-operand vec4 kernel: the other operand matches the output and the vector count is a multiple of four, so the scalar reads once and splats."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.23, "scale": 1.5 }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [2, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.31, "cosStep": 0.11, "scale": 1.5 }
+        }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [2, 8], "tolerance": 0.000001 } }
     }
   ]
 }