sync 2e7068faf55e

Browse files

Files changed (8) hide show

README.md +61 -0
build/webgpu/bench.json +71 -0
build/webgpu/datamove-elementwise-copy.wgsl.jinja +20 -0
build/webgpu/manifest.json +424 -0
build/webgpu/metadata.json +20 -0
build/webgpu/summean-broadcast.wgsl.jinja +107 -0
build/webgpu/summean-vec4.wgsl.jinja +30 -0
build/webgpu/test.json +578 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.Sum
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 13
+## Description
+Computes the elementwise sum of one or more input tensors with multidirectional (NumPy-style) broadcasting. All inputs and the output must share the same data type.
+See the [ONNX `Sum` spec](https://onnx.ai/onnx/operators/onnx__Sum.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `A` | `a` | `T` | — | — | First input tensor. | required |
+| `B` | `b` | `T` | — | — | Second input tensor (optional). | optional |
+| `C` | `c` | `T` | — | — | Third input tensor (optional). | optional |
+| `D` | `d` | `T` | — | — | Fourth input tensor (optional). | optional |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `sum` | `y` | `T` | derived | derived; see description | Elementwise sum of all provided input tensors. | required |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16` |
+## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
+- [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
+- [`test.json`](build/webgpu/test.json) — correctness cases
+- [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
+- [`datamove-elementwise-copy.wgsl.jinja`](build/webgpu/datamove-elementwise-copy.wgsl.jinja)
+- [`summean-broadcast.wgsl.jinja`](build/webgpu/summean-broadcast.wgsl.jinja)
+- [`summean-vec4.wgsl.jinja`](build/webgpu/summean-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.Sum", { version: 1 });
+const { y } = await kernel({ a: { data: aData, shape: [3] } });
+```

build/webgpu/bench.json ADDED Viewed

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

+{
+  "op": "ai.onnx.Sum",
+  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
+  "cases": [
+    {
+      "name": "sum-f32-4m",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "count": 4194304 },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [4194304], "dist": "normal", "seed": 702, "scale": 2 },
+        "b": { "dtype": "float32", "shape": [4194304], "dist": "normal", "seed": 703, "scale": 2 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4194304] } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 3" }]
+      }
+    },
+    {
+      "name": "sum-f32-broadcast-scalar-4m-unaligned",
+      "preset": "stress",
+      "vars": { "dtype": "float32", "count": 4194302 },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [4194302], "dist": "normal", "seed": 711, "scale": 2 },
+        "b": { "dtype": "float32", "shape": [4194302], "dist": "normal", "seed": 712, "scale": 2 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4194302] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 3" }] }
+    },
+    {
+      "name": "sum-f32-four-input-rank4-mixed-broadcast-pathology",
+      "preset": "stress",
+      "vars": { "dtype": "float32", "count": 2097152 },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [8, 64, 64, 64], "dist": "normal", "seed": 713, "scale": 2 },
+        "b": { "dtype": "float32", "shape": [1, 64, 1, 1], "dist": "normal", "seed": 714, "scale": 0.25 },
+        "c": { "dtype": "float32", "shape": [8, 1, 64, 64], "dist": "normal", "seed": 715, "scale": 0.5 },
+        "d": { "dtype": "float32", "shape": [1], "data": { "kind": "constant", "value": 0.125 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [8, 64, 64, 64], "dist": "empty" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "dtypeBytes(args.dtype) * (numel(shapes.a) + numel(shapes.b) + numel(shapes.c) + numel(shapes.d) + numel(shapes.y))"
+          }
+        ]
+      }
+    },
+    {
+      "name": "sum-f32-three-input-rank4-mixed-broadcast-pathology",
+      "preset": "stress",
+      "vars": { "dtype": "float32", "count": 2097152 },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [8, 64, 64, 64], "dist": "normal", "seed": 716, "scale": 2 },
+        "b": { "dtype": "float32", "shape": [1, 64, 1, 1], "dist": "normal", "seed": 717, "scale": 0.25 },
+        "c": { "dtype": "float32", "shape": [8, 1, 64, 64], "dist": "normal", "seed": 718, "scale": 0.5 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [8, 64, 64, 64], "dist": "empty" } },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "dtypeBytes(args.dtype) * (numel(shapes.a) + numel(shapes.b) + numel(shapes.c) + numel(shapes.y))"
+          }
+        ]
+      }
+    }
+  ]
+}

build/webgpu/datamove-elementwise-copy.wgsl.jinja ADDED Viewed

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

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(
+  @builtin(global_invocation_id) gid: vec3<u32>,
+  @builtin(num_workgroups) nwg: vec3<u32>
+) {
+  // Grid-stride loop: the dispatch is clamped to the maxComputeWorkgroupsPerDimension
+  // workgroups-per-dimension limit, so a thread may copy more than one
+  // element for very large tensors.
+  let stride = nwg.x * WG;
+  for (var i = gid.x; i < params.count; i += stride) {
+    y[i] = x[i];
+  }
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,424 @@

+{
+  "domain": "ai.onnx",
+  "name": "Sum",
+  "sinceVersion": 13,
+  "description": "Computes the elementwise sum of one or more input tensors with multidirectional (NumPy-style) broadcasting. All inputs and the output must share the same data type.",
+  "inputs": [
+    { "role": "A", "dtype": "T", "description": "First input tensor." },
+    { "role": "B", "dtype": "T", "optional": true, "description": "Second input tensor (optional)." },
+    { "role": "C", "dtype": "T", "optional": true, "description": "Third input tensor (optional)." },
+    { "role": "D", "dtype": "T", "optional": true, "description": "Fourth input tensor (optional)." }
+  ],
+  "outputs": [
+    {
+      "role": "sum",
+      "dtype": "T",
+      "rank": "max(ranks.A, ranks.B if present.b else 0, ranks.C if present.b and present.c else 0, ranks.D if present.b and present.c and present.d else 0)",
+      "shape": "variadicShape",
+      "description": "Elementwise sum of all provided input tensors."
+    }
+  ],
+  "typeConstraints": { "T": ["float32", "float16"] },
+  "args": {
+    "a": { "kind": "tensor", "semantic": "A", "role": "input" },
+    "b": { "kind": "tensor", "semantic": "B", "role": "input", "required": false },
+    "c": { "kind": "tensor", "semantic": "C", "role": "input2", "required": false },
+    "y": { "kind": "tensor", "semantic": "sum", "role": "output" },
+    "d": { "kind": "tensor", "semantic": "D", "role": "input3", "required": false }
+  },
+  "tunables": { "WORKGROUP_SIZE": 256 },
+  "derive": {
+    "variadicShape": "broadcastShape(broadcastShape(broadcastShape(shapes.A, shapes.B), shapes.C), shapes.D) if present.b and present.c and present.d else (broadcastShape(broadcastShape(shapes.A, shapes.B), shapes.C) if present.b and present.c else (broadcastShape(shapes.A, shapes.B) if present.b else shapes.A))",
+    "flatVec4OutputOk": "numel(shapes.y) > 0 and numel(shapes.y) % 4 == 0 and f16Ok(dtypes.T)",
+    "broadcastOutputOk": "f16Ok(dtypes.T)",
+    "broadcastVec4OutputOk": "ranks.y >= 1 and dim(shapes.y, ranks.y - 1) > 0 and dim(shapes.y, ranks.y - 1) % 4 == 0 and f16Ok(dtypes.T)",
+    "aBroadcastVec4Ok": "ranks.A <= ranks.y and (ranks.A == 0 or dim(shapes.A, ranks.A - 1) == 1 or dim(shapes.A, ranks.A - 1) == dim(shapes.y, ranks.y - 1))",
+    "bBroadcastVec4Ok": "not present.b or (ranks.B <= ranks.y and (ranks.B == 0 or dim(shapes.B, ranks.B - 1) == 1 or dim(shapes.B, ranks.B - 1) == dim(shapes.y, ranks.y - 1)))",
+    "cBroadcastVec4Ok": "not present.c or (ranks.C <= ranks.y and (ranks.C == 0 or dim(shapes.C, ranks.C - 1) == 1 or dim(shapes.C, ranks.C - 1) == dim(shapes.y, ranks.y - 1)))",
+    "dBroadcastVec4Ok": "not present.d or (ranks.D <= ranks.y and (ranks.D == 0 or dim(shapes.D, ranks.D - 1) == 1 or dim(shapes.D, ranks.D - 1) == dim(shapes.y, ranks.y - 1)))"
+  },
+  "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
+  "bindingSets": {
+    "identity": [
+      { "name": "x", "arg": "a", "semantic": "A", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
+      { "name": "y", "arg": "y", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }] }
+      }
+    ],
+    "vec4Two": [
+      {
+        "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": "y", "arg": "y", "buffer": { "type": "storage" }, "elementType": "$vectorScalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.y) / 4" }] }
+      }
+    ],
+    "vec4Three": [
+      {
+        "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": "read-only-storage" },
+        "elementType": "$cElement"
+      },
+      { "name": "y", "arg": "y", "semantic": "sum", "buffer": { "type": "storage" }, "elementType": "$vectorScalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.y) / 4" }] }
+      }
+    ],
+    "vec4Four": [
+      {
+        "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": "read-only-storage" },
+        "elementType": "$cElement"
+      },
+      {
+        "name": "d",
+        "arg": "d",
+        "semantic": "D",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$dElement"
+      },
+      { "name": "y", "arg": "y", "semantic": "sum", "buffer": { "type": "storage" }, "elementType": "$vectorScalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.y) / 4" }] }
+      }
+    ],
+    "scalarTwo": [
+      { "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": "y", "arg": "y", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }] }
+      }
+    ],
+    "scalarThree": [
+      { "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": "read-only-storage" }, "elementType": "$scalar" },
+      { "name": "y", "arg": "y", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }] }
+      }
+    ],
+    "scalarFour": [
+      { "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": "read-only-storage" }, "elementType": "$scalar" },
+      { "name": "d", "arg": "d", "semantic": "D", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
+      { "name": "y", "arg": "y", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }] }
+      }
+    ]
+  },
+  "variants": [
+    {
+      "id": "single_input_identity",
+      "priority": 30,
+      "when": ["not present.b", "not present.c", "ranks.A == ranks.y", "numel(shapes.A) == numel(shapes.y)", "f16Ok(dtypes.T)"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "Sum.Identity",
+          "shader": "datamove-elementwise-copy.wgsl.jinja",
+          "bindings": "identity",
+          "dispatch": { "gridStride": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "same_shape_vec4_two_input",
+      "priority": 20,
+      "when": ["present.b", "not present.c", "sameShape(shapes.A, shapes.y)", "sameShape(shapes.B, shapes.y)", "flatVec4OutputOk"],
+      "constants": {
+        "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
+        "aElement": "\"vec4<\" ~ dtypes.T ~ \">\"",
+        "bElement": "\"vec4<\" ~ dtypes.T ~ \">\"",
+        "cElement": "\"vec4<\" ~ dtypes.T ~ \">\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Sum.vec4",
+          "source": { "shader": "summean-vec4.wgsl.jinja", "inputs": { "op": "\"sum\"", "hasC": "false" } },
+          "bindings": "vec4Two",
+          "dispatch": { "threads": "numel(shapes.y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "same_shape_vec4_three_input",
+      "priority": 25,
+      "when": ["present.b", "present.c", "sameShape(shapes.A, shapes.y)", "sameShape(shapes.B, shapes.y)", "sameShape(shapes.C, shapes.y)", "flatVec4OutputOk", "not present.d"],
+      "constants": {
+        "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
+        "aElement": "\"vec4<\" ~ dtypes.T ~ \">\"",
+        "bElement": "\"vec4<\" ~ dtypes.T ~ \">\"",
+        "cElement": "\"vec4<\" ~ dtypes.T ~ \">\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Sum.vec4_3",
+          "source": { "shader": "summean-vec4.wgsl.jinja", "inputs": { "op": "\"sum\"", "hasC": "true" } },
+          "bindings": "vec4Three",
+          "dispatch": { "threads": "numel(shapes.y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "same_shape_vec4_four_input",
+      "priority": 27,
+      "when": ["present.b", "present.c", "present.d", "sameShape(shapes.A, shapes.y)", "sameShape(shapes.B, shapes.y)", "sameShape(shapes.C, shapes.y)", "sameShape(shapes.D, shapes.y)", "flatVec4OutputOk"],
+      "constants": {
+        "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
+        "aElement": "\"vec4<\" ~ dtypes.T ~ \">\"",
+        "bElement": "\"vec4<\" ~ dtypes.T ~ \">\"",
+        "cElement": "\"vec4<\" ~ dtypes.T ~ \">\"",
+        "dElement": "\"vec4<\" ~ dtypes.T ~ \">\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Sum.vec4_3",
+          "source": {
+            "shader": "summean-vec4.wgsl.jinja",
+            "inputs": { "op": "\"sum\"", "hasC": "true", "hasD": "true" }
+          },
+          "bindings": "vec4Four",
+          "dispatch": { "threads": "numel(shapes.y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "broadcast_two_input",
+      "when": ["present.b", "not present.c", "ranks.A <= ranks.y", "ranks.B <= ranks.y", "broadcastOutputOk"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "Sum",
+          "source": {
+            "shader": "summean-broadcast.wgsl.jinja",
+            "inputs": {
+              "aShape": "shapes.A",
+              "bShape": "shapes.B",
+              "yShape": "shapes.y",
+              "aRank": "ranks.A",
+              "bRank": "ranks.B",
+              "yRank": "ranks.y",
+              "hasC": "false",
+              "op": "\"sum\""
+            }
+          },
+          "bindings": "scalarTwo",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "broadcast_three_input_vec4",
+      "priority": 24,
+      "when": ["present.b", "present.c", "not present.d", "broadcastVec4OutputOk", "aBroadcastVec4Ok", "bBroadcastVec4Ok", "cBroadcastVec4Ok"],
+      "constants": {
+        "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
+        "aElement": "(\"vec4<\" ~ dtypes.T ~ \">\") if ranks.A > 0 and dim(shapes.A, ranks.A - 1) != 1 else dtypes.T",
+        "bElement": "(\"vec4<\" ~ dtypes.T ~ \">\") if ranks.B > 0 and dim(shapes.B, ranks.B - 1) != 1 else dtypes.T",
+        "cElement": "(\"vec4<\" ~ dtypes.T ~ \">\") if ranks.C > 0 and dim(shapes.C, ranks.C - 1) != 1 else dtypes.T"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Sum.Broadcast3Vec4",
+          "source": {
+            "shader": "summean-broadcast.wgsl.jinja",
+            "inputs": {
+              "aShape": "shapes.A",
+              "bShape": "shapes.B",
+              "cShape": "shapes.C",
+              "yShape": "shapes.y",
+              "aRank": "ranks.A",
+              "bRank": "ranks.B",
+              "cRank": "ranks.C",
+              "yRank": "ranks.y",
+              "aVector": "ranks.A > 0 and dim(shapes.A, ranks.A - 1) != 1",
+              "bVector": "ranks.B > 0 and dim(shapes.B, ranks.B - 1) != 1",
+              "cVector": "ranks.C > 0 and dim(shapes.C, ranks.C - 1) != 1",
+              "hasC": true,
+              "vectorized": true,
+              "op": "\"sum\""
+            }
+          },
+          "bindings": "vec4Three",
+          "dispatch": { "threads": "numel(shapes.y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "broadcast_three_input",
+      "priority": 20,
+      "when": ["present.b", "present.c", "ranks.A <= ranks.y", "ranks.B <= ranks.y", "ranks.C <= ranks.y", "broadcastOutputOk", "not present.d"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "Sum",
+          "source": {
+            "shader": "summean-broadcast.wgsl.jinja",
+            "inputs": {
+              "aShape": "shapes.A",
+              "bShape": "shapes.B",
+              "cShape": "shapes.C",
+              "yShape": "shapes.y",
+              "aRank": "ranks.A",
+              "bRank": "ranks.B",
+              "cRank": "ranks.C",
+              "yRank": "ranks.y",
+              "hasC": "true",
+              "op": "\"sum\""
+            }
+          },
+          "bindings": "scalarThree",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "broadcast_four_input",
+      "priority": 22,
+      "when": ["present.b", "present.c", "present.d", "ranks.A <= ranks.y", "ranks.B <= ranks.y", "ranks.C <= ranks.y", "ranks.D <= ranks.y", "broadcastOutputOk"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "Sum",
+          "source": {
+            "shader": "summean-broadcast.wgsl.jinja",
+            "inputs": {
+              "aShape": "shapes.A",
+              "bShape": "shapes.B",
+              "cShape": "shapes.C",
+              "yShape": "shapes.y",
+              "aRank": "ranks.A",
+              "bRank": "ranks.B",
+              "cRank": "ranks.C",
+              "yRank": "ranks.y",
+              "hasC": "true",
+              "op": "\"sum\"",
+              "dShape": "shapes.D",
+              "dRank": "ranks.D",
+              "hasD": "true"
+            }
+          },
+          "bindings": "scalarFour",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "broadcast_four_input_vec4",
+      "priority": 26,
+      "when": ["present.b", "present.c", "present.d", "broadcastVec4OutputOk", "aBroadcastVec4Ok", "bBroadcastVec4Ok", "cBroadcastVec4Ok", "dBroadcastVec4Ok"],
+      "constants": {
+        "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
+        "aElement": "(\"vec4<\" ~ dtypes.T ~ \">\") if ranks.A > 0 and dim(shapes.A, ranks.A - 1) != 1 else dtypes.T",
+        "bElement": "(\"vec4<\" ~ dtypes.T ~ \">\") if ranks.B > 0 and dim(shapes.B, ranks.B - 1) != 1 else dtypes.T",
+        "cElement": "(\"vec4<\" ~ dtypes.T ~ \">\") if ranks.C > 0 and dim(shapes.C, ranks.C - 1) != 1 else dtypes.T",
+        "dElement": "(\"vec4<\" ~ dtypes.T ~ \">\") if ranks.D > 0 and dim(shapes.D, ranks.D - 1) != 1 else dtypes.T"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Sum.Broadcast4Vec4",
+          "source": {
+            "shader": "summean-broadcast.wgsl.jinja",
+            "inputs": {
+              "aShape": "shapes.A",
+              "bShape": "shapes.B",
+              "cShape": "shapes.C",
+              "dShape": "shapes.D",
+              "yShape": "shapes.y",
+              "aRank": "ranks.A",
+              "bRank": "ranks.B",
+              "cRank": "ranks.C",
+              "dRank": "ranks.D",
+              "yRank": "ranks.y",
+              "aVector": "ranks.A > 0 and dim(shapes.A, ranks.A - 1) != 1",
+              "bVector": "ranks.B > 0 and dim(shapes.B, ranks.B - 1) != 1",
+              "cVector": "ranks.C > 0 and dim(shapes.C, ranks.C - 1) != 1",
+              "dVector": "ranks.D > 0 and dim(shapes.D, ranks.D - 1) != 1",
+              "hasC": true,
+              "hasD": true,
+              "vectorized": true,
+              "op": "\"sum\""
+            }
+          },
+          "bindings": "vec4Four",
+          "dispatch": { "threads": "numel(shapes.y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

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

+{
+  "name": "ai.onnx.Sum",
+  "id": "_ai_onnx_sum_webgpu_2c2501b",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "BZH67s4WnBU60TILXvJnrpu86w673p/81oS0ITOTOjw=",
+      "datamove-elementwise-copy.wgsl.jinja": "J5yC2bAddPiP+odXLgVGS3TJ9jeNsfRTvedKrj/fhZg=",
+      "manifest.json": "n3EnmVEajNhmmHf1ueSFLaPzePrvHY41jyuk9AfSxWM=",
+      "summean-broadcast.wgsl.jinja": "6BaP0d/1jnidmzD5A+BolzRw9fhm4PaQa4nAjjGCVmI=",
+      "summean-vec4.wgsl.jinja": "kS8/lN2swjqu28ALxVcF4rkhpEfMfqUGRxyMf770veU=",
+      "test.json": "DUAhzo/YUbzJAmNXu5b92KKxvQCxcDDuuNeen2eTyfI="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Sum" }
+}

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

	@@ -0,0 +1,107 @@

+{% set allInputs = ["a", "b"] + (source.extraInputs if source.extraInputs is defined else (["c"] if source.hasC else []) + (["d"] if source.hasD else [])) %}
+{% set extraInputs = source.extraInputs if source.extraInputs is defined else (["c"] if source.hasC else []) + (["d"] if source.hasD else []) %}
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+{% 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 %}
+{{ broadcast_offset_fn("a_offset", source.aShape, source.aRank, source.yShape, source.yRank) }}
+{{ broadcast_offset_fn("b_offset", source.bShape, source.bRank, source.yShape, source.yRank) }}
+{% for n in extraInputs %}
+{{ broadcast_offset_fn(n ~ "_offset", source[n ~ "Shape"], source[n ~ "Rank"], source.yShape, source.yRank) }}
+{% endfor %}
+@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 past the
+  // maxComputeWorkgroupsPerDimension limit.
+  let i = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if (i >= params.count) {
+    return;
+  }
+{% if source.vectorized %}
+  // The output's innermost dimension is four-aligned. Each input either keeps
+  // that dimension (one aligned vec4 load) or broadcasts it (one scalar splat).
+  // Offset folding is therefore paid once per four outputs without changing
+  // multidirectional broadcast semantics on any outer dimension.
+  let base = i * 4u;
+{% for n in allInputs %}
+{% if source[n ~ "Vector"] %}
+  let {{ n }}v = vec4<f32>({{ n }}[{{ broadcast_offset_call(n ~ "_offset", source[n ~ "Shape"], source.yShape, "base") }} / 4u]);
+{% else %}
+  let {{ n }}v = vec4<f32>(f32({{ n }}[{{ broadcast_offset_call(n ~ "_offset", source[n ~ "Shape"], source.yShape, "base") }}]));
+{% endif %}
+{% endfor %}
+  let total = {% for n in allInputs %}{{ n }}v{% if not loop.last %} + {% endif %}{% endfor %};
+  y[i] = {{ vectorScalar }}(total);
+{% else %}
+  let total = {% for n in allInputs %}f32({{ n }}[{{ broadcast_offset_call(n ~ "_offset", source[n ~ "Shape"], source.yShape, "i") }}]){% if not loop.last %} + {% endif %}{% endfor %};
+  y[i] = {{ scalar }}(total);
+{% endif %}
+}

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

	@@ -0,0 +1,30 @@

+{% set extraInputs = source.extraInputs if source.extraInputs is defined else (["c"] if source.hasC else []) + (["d"] if source.hasD else []) %}
+{% 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 past the
+  // maxComputeWorkgroupsPerDimension limit.
+  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 == "f16" %}
+  var total = vec4<f32>(av) + vec4<f32>(bv);
+{% for n in extraInputs %}
+  total = total + vec4<f32>({{ n }}[i]);
+{% endfor %}
+  y[i] = vec4<f16>(total);
+{% else %}
+  var total = av + bv;
+{% for n in extraInputs %}
+  total = total + {{ n }}[i];
+{% endfor %}
+  y[i] = total;
+{% endif %}
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,578 @@

+{
+  "op": "ai.onnx.Sum",
+  "cases": [
+    {
+      "name": "max_arity_float16_positions",
+      "provenance": {
+        "notes": "Synthetic four-input float16 Sum contract fixture; element-varying operands make every bounded input position contribute to a non-constant exact sum."
+      },
+      "inputs": {
+        "a": { "dtype": "float16", "shape": [5], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0] } },
+        "b": { "dtype": "float16", "shape": [5], "data": { "kind": "values", "values": [2.0, 4.0, 6.0, 8.0, 10.0] } },
+        "c": { "dtype": "float16", "shape": [5], "data": { "kind": "values", "values": [4.0, 1.0, 3.0, 2.0, 5.0] } },
+        "d": { "dtype": "float16", "shape": [5], "data": { "kind": "values", "values": [8.0, 2.0, 1.0, 4.0, 3.0] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float16",
+          "shape": [5],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [15.0, 9.0, 13.0, 18.0, 23.0] }
+        }
+      }
+    },
+    {
+      "name": "same_shape",
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29 }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.31, "cosStep": 0.07 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "f32_subnormal_three_input_sum_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 subnormal addends/sum cannot be reproduced on GPU."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Sum_8_Test1",
+        "notes": "Variadic Sum must preserve finite subnormal totals; a flush-to-zero path erases residual-like lanes before downstream scaling can amplify them."
+      },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1e-40, -1e-40, 0.0] } },
+        "b": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1e-40, -1e-40, 1e-40] } },
+        "c": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1e-40, -1e-40, -1e-40] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [3],
+          "tolerance": 2e-45,
+          "data": { "kind": "values", "values": [3e-40, -3e-40, 0.0] }
+        }
+      }
+    },
+    {
+      "name": "f32_subnormal_three_input_sum_vec4_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 subnormal addends/sum cannot be reproduced on GPU."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Sum_8_Test1",
+        "notes": "Vec4 companion: variadic Sum must preserve finite subnormal totals in the vectorized path."
+      },
+      "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": [1e-40, -1e-40, 1e-40, -1e-39] }
+        },
+        "c": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1e-40, -1e-40, -1e-40, 0.0] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [4],
+          "tolerance": 2e-45,
+          "data": { "kind": "values", "values": [3e-40, -3e-40, 0.0, 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": {
+        "y": {
+          "dtype": "float16",
+          "shape": [8],
+          "tolerance": 0.001,
+          "data": { "kind": "values", "values": [1.5, 0.0, 2.0, 0.0, 1.0, 5.0, 0.0, 0.0] }
+        }
+      }
+    },
+    {
+      "name": "same_shape_vec4_three_input",
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [1.0, 5.0, -2.0, 4.0, 0.0, 6.0, 10.0, -10.0] }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [3.0, 2.0, -4.0, 8.0, 1.0, 1.0, 9.0, -9.0] }
+        },
+        "c": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0.0, 7.0, -3.0, 2.0, -1.0, 8.0, 11.0, -11.0] }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [8],
+          "tolerance": 0.000001,
+          "data": { "kind": "values", "values": [4.0, 14.0, -9.0, 14.0, 0.0, 15.0, 30.0, -30.0] }
+        }
+      }
+    },
+    {
+      "name": "float16_vec4_three_input",
+      "provenance": {
+        "notes": "The float16 three-input vec4 path widens each operand to f32 before the extra-operand fold. Every operand and partial sum is an exactly representable dyadic float16 value, so dropping or double-counting C changes an exact zero-tolerance result."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "float16",
+          "shape": [8],
+          "data": { "kind": "values", "values": [1.0, -2.0, 0.5, 16.0, -0.25, 6.0, -32.0, 0.125] }
+        },
+        "b": {
+          "dtype": "float16",
+          "shape": [8],
+          "data": { "kind": "values", "values": [2.0, -1.5, -1.0, 8.0, 0.125, 3.0, -8.0, 0.375] }
+        },
+        "c": {
+          "dtype": "float16",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0.5, -4.0, 3.0, -4.0, -0.5, 12.0, 24.0, 0.5] }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float16",
+          "shape": [8],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [3.5, -7.5, 2.5, 20.0, -0.625, 21.0, -16.0, 1.0] }
+        }
+      }
+    },
+    {
+      "name": "broadcast_rank4",
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29 }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [1, 3, 1, 5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.31, "cosStep": 0.07 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 4, 5], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank0_rhs_scalar_broadcast",
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        },
+        "b": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-1.5] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_two_inputs_no_broadcasting_3d",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.SumMultipleInputsNoBroadcasting(num_inputs=2)"
+      },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [3, 3, 3], "data": { "kind": "constant", "value": 1.0 } },
+        "b": { "dtype": "float32", "shape": [3, 3, 3], "data": { "kind": "constant", "value": 1.0 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 3, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_four_inputs_no_broadcasting_3d_variadic",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.SumMultipleInputsNoBroadcasting(num_inputs=4)",
+        "notes": "ONNX Sum is variadic; this catches implementations capped at three inputs."
+      },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [3, 3, 3], "data": { "kind": "constant", "value": 1.0 } },
+        "b": { "dtype": "float32", "shape": [3, 3, 3], "data": { "kind": "constant", "value": 1.0 } },
+        "c": { "dtype": "float32", "shape": [3, 3, 3], "data": { "kind": "constant", "value": 1.0 } },
+        "d": { "dtype": "float32", "shape": [3, 3, 3], "data": { "kind": "constant", "value": 1.0 } }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [3, 3, 3], "tolerance": 0, "data": { "kind": "constant", "value": 4.0 } }
+      }
+    },
+    {
+      "name": "ort_four_inputs_nan_infinity_variadic",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.SumMultipleInputsNoBroadcasting(num_inputs=4)",
+        "notes": "Extends ORT's four-input Sum coverage with NaN propagation and inf-minus-inf cancellation."
+      },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, "Infinity", 1.0, 8.0] } },
+        "b": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [3.0, "-Infinity", 5.0, 4.0] } },
+        "c": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [5.0, 2.0, "NaN", -4.0] } },
+        "d": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [7.0, 6.0, 9.0, "NaN"] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [4],
+          "tolerance": 0,
+          "allowNaN": true,
+          "data": { "kind": "values", "values": [16.0, "NaN", "NaN", "NaN"] }
+        }
+      }
+    },
+    {
+      "name": "ort_fifteen_inputs_no_broadcasting_3d_variadic_gpu_gap",
+      "skipGpu": {
+        "category": "todo",
+        "reason": "A single-pass implementation needs 15 input bindings plus one output binding, which exceeds WebGPU's guaranteed storage-buffer limit. A multi-pass in-place accumulator can implement this portable high-arity route but is not yet available."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.SumMultipleInputsNoBroadcasting(num_inputs=15)",
+        "notes": "ORT explicitly covers high-arity no-broadcast Sum paths; this catches implementations capped at a small fixed arity."
+      },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [3, 3, 3], "data": { "kind": "constant", "value": 1.0 } },
+        "b": { "dtype": "float32", "shape": [3, 3, 3], "data": { "kind": "constant", "value": 1.0 } },
+        "c": { "dtype": "float32", "shape": [3, 3, 3], "data": { "kind": "constant", "value": 1.0 } },
+        "d": { "dtype": "float32", "shape": [3, 3, 3], "data": { "kind": "constant", "value": 1.0 } },
+        "e": { "dtype": "float32", "shape": [3, 3, 3], "data": { "kind": "constant", "value": 1.0 } },
+        "f": { "dtype": "float32", "shape": [3, 3, 3], "data": { "kind": "constant", "value": 1.0 } },
+        "g": { "dtype": "float32", "shape": [3, 3, 3], "data": { "kind": "constant", "value": 1.0 } },
+        "h": { "dtype": "float32", "shape": [3, 3, 3], "data": { "kind": "constant", "value": 1.0 } },
+        "i": { "dtype": "float32", "shape": [3, 3, 3], "data": { "kind": "constant", "value": 1.0 } },
+        "j": { "dtype": "float32", "shape": [3, 3, 3], "data": { "kind": "constant", "value": 1.0 } },
+        "k": { "dtype": "float32", "shape": [3, 3, 3], "data": { "kind": "constant", "value": 1.0 } },
+        "l": { "dtype": "float32", "shape": [3, 3, 3], "data": { "kind": "constant", "value": 1.0 } },
+        "m": { "dtype": "float32", "shape": [3, 3, 3], "data": { "kind": "constant", "value": 1.0 } },
+        "n": { "dtype": "float32", "shape": [3, 3, 3], "data": { "kind": "constant", "value": 1.0 } },
+        "o": { "dtype": "float32", "shape": [3, 3, 3], "data": { "kind": "constant", "value": 1.0 } }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [3, 3, 3], "tolerance": 0, "data": { "kind": "constant", "value": 15.0 } }
+      }
+    },
+    {
+      "name": "ort_opset6_two_input_same_shape_projection",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Sum_6",
+        "notes": "Two-input projection of ORT's same-shape three-input case."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [3, 3],
+          "data": { "kind": "values", "values": [1.0, 0.0, 1.0, -1.0, 1.1, -100.0, -5.4, 0.01, -10000.0] }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [3, 3],
+          "data": { "kind": "values", "values": [1.0, 0.0, 2.0, -2.0, 2.2, 64.0, -1.0, 0.02, 0.25] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_opset6_three_input_same_shape_exact",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Sum_6"
+      },
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [3, 3],
+          "data": { "kind": "values", "values": [1.0, 0.0, 1.0, -1.0, 1.1, -100.0, -5.4, 0.01, -10000.0] }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [3, 3],
+          "data": { "kind": "values", "values": [1.0, 0.0, 2.0, -2.0, 2.2, 64.0, -1.0, 0.02, 0.25] }
+        },
+        "c": {
+          "dtype": "float32",
+          "shape": [3, 3],
+          "data": { "kind": "values", "values": [1.0, 0.0, 3.0, -3.0, 3.3, 64.0, 5.4, 0.03, 10000.0] }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [3, 3],
+          "tolerance": 0.000001,
+          "data": { "kind": "values", "values": [3.0, 0.0, 6.0, -6.0, 6.6, 28.0, -1.0, 0.06, 0.25] }
+        }
+      }
+    },
+    {
+      "name": "ort_opset8_vector_column_broadcast_projection",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Sum_8_Test1",
+        "notes": "Two-input projection exercising vector plus column multidirectional broadcasting."
+      },
+      "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": [10.0, 20.0, 30.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_opset8_three_input_vector_column_depth_broadcast_exact",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Sum_8_Test1"
+      },
+      "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": [10.0, 20.0, 30.0] } },
+        "c": { "dtype": "float32", "shape": [3, 1, 1], "data": { "kind": "values", "values": [100.0, 200.0, 300.0] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [3, 3, 3],
+          "tolerance": 0,
+          "data": {
+            "kind": "values",
+            "values": [111.0, 112.0, 113.0, 121.0, 122.0, 123.0, 131.0, 132.0, 133.0, 211.0, 212.0, 213.0, 221.0, 222.0, 223.0, 231.0, 232.0, 233.0, 311.0, 312.0, 313.0, 321.0, 322.0, 323.0, 331.0, 332.0, 333.0]
+          }
+        }
+      }
+    },
+    {
+      "name": "ort_opset8_four_input_vector_column_depth_broadcast",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Sum_8_Test1",
+        "notes": "Extends ORT's multidirectional broadcast case to a valid four-input ONNX variadic Sum node."
+      },
+      "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": [10.0, 20.0, 30.0] } },
+        "c": { "dtype": "float32", "shape": [3, 1, 1], "data": { "kind": "values", "values": [100.0, 200.0, 300.0] } },
+        "d": {
+          "dtype": "float32",
+          "shape": [1, 1, 3],
+          "data": { "kind": "values", "values": [1000.0, 2000.0, 3000.0] }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [3, 3, 3],
+          "tolerance": 0,
+          "data": {
+            "kind": "values",
+            "values": [1111.0, 2112.0, 3113.0, 1121.0, 2122.0, 3123.0, 1131.0, 2132.0, 3133.0, 1211.0, 2212.0, 3213.0, 1221.0, 2222.0, 3223.0, 1231.0, 2232.0, 3233.0, 1311.0, 2312.0, 3313.0, 1321.0, 2322.0, 3323.0, 1331.0, 2332.0, 3333.0]
+          }
+        }
+      }
+    },
+    {
+      "name": "ort_opset8_matrix_vector_broadcast_projection",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Sum_8_Test2",
+        "notes": "Two-input projection of ORT's matrix plus trailing-vector broadcast case."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [3, 3],
+          "data": { "kind": "values", "values": [1.0, 0.0, 1.0, -1.0, 1.1, -100.0, -5.4, 0.01, -74.0] }
+        },
+        "b": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 0.0, 2.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "onnx_backend_sum_two_inputs",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_sum_two_inputs" },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [3.0, 0.0, 2.0] } },
+        "b": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 3.0, 4.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_sum_one_input_identity",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_sum_one_input" },
+      "inputs": { "a": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [3.0, 0.0, 2.0] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [3], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_sum_example_three_inputs",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_sum_example" },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [3.0, 0.0, 2.0] } },
+        "b": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 3.0, 4.0] } },
+        "c": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [2.0, 6.0, 6.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_dim_zero_equal_rank",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.DimWithZeroHandling",
+        "notes": "Projected from ORT's binary elementwise zero-dimension Add coverage to generic ONNX multidirectional broadcasting."
+      },
+      "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": { "y": { "dtype": "float32", "shape": [3, 0], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_dim_zero_scalar_broadcast",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.DimWithZeroHandling",
+        "notes": "Projected from ORT's binary elementwise zero-dimension Add coverage to generic ONNX multidirectional broadcasting."
+      },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] } },
+        "b": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [0], "tolerance": 0 } }
+    },
+    {
+      "name": "single_input_empty_shape",
+      "inputs": { "a": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [0], "tolerance": 0 } }
+    },
+    {
+      "name": "f16_broadcast_differing_shapes",
+      "inputs": {
+        "a": { "dtype": "float16", "shape": [1, 4], "data": { "kind": "values", "values": [1.0, 2.0, 4.0, 8.0] } },
+        "b": {
+          "dtype": "float16",
+          "shape": [4, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.41 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [4, 4], "tolerance": 0.001 } }
+    },
+    {
+      "name": "rank7_broadcast_two_input",
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [1, 2, 1, 2, 1, 2, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23 }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [2, 1, 2, 1, 2, 1, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.19 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 2, 2, 2, 2, 2, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "three_input_empty_dim_broadcast",
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [2, 1], "data": { "kind": "values", "values": [1.0, 2.0] } },
+        "b": { "dtype": "float32", "shape": [2, 0], "data": { "kind": "values", "values": [] } },
+        "c": { "dtype": "float32", "shape": [1, 0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 0], "tolerance": 0 } }
+    },
+    {
+      "name": "four_input_mixed_broadcast_vec4",
+      "provenance": {
+        "notes": "Route lock for mixed four-input vec4 broadcasting: A and C retain the eight-wide innermost dimension, while B broadcasts channel scalars and D broadcasts one scalar across every output lane."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [2, 3, 2, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.25 }
+        },
+        "b": { "dtype": "float32", "shape": [1, 3, 1, 1], "data": { "kind": "values", "values": [1.0, 2.0, 4.0] } },
+        "c": {
+          "dtype": "float32",
+          "shape": [2, 1, 2, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.29, "scale": 0.125 }
+        },
+        "d": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.5] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 2, 8], "tolerance": 0.000001, "relTolerance": 0.000001 } }
+    },
+    {
+      "name": "three_input_mixed_broadcast_vec4_candidate",
+      "provenance": {
+        "notes": "Three-input mixed-broadcast lock: A is dense, B varies over channel and the penultimate axis but broadcasts its innermost scalar, and C retains vector lanes while broadcasting both middle axes."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [2, 3, 2, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.25 }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [1, 3, 2, 1],
+          "data": { "kind": "values", "values": [1.0, 2.0, 4.0, 8.0, 16.0, 32.0] }
+        },
+        "c": {
+          "dtype": "float32",
+          "shape": [2, 1, 1, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.29, "scale": 0.125 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 2, 8], "tolerance": 0.000001, "relTolerance": 0.000001 } }
+    },
+    {
+      "name": "rank8_broadcast_two_input",
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [1, 2, 1, 2, 1, 2, 2, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23 }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [2, 1, 2, 1, 2, 1, 2, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.19 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 2, 2, 2, 2, 2, 2, 3], "tolerance": 0.000001 } }
+    }
+  ]
+}