sync 91d990483a17

Browse files

Files changed (8) hide show

README.md +14 -10
build/webgpu/bench.json +0 -1
build/webgpu/datamove-elementwise-copy.wgsl.jinja +0 -3
build/webgpu/manifest.json +189 -291
build/webgpu/metadata.json +22 -9
build/webgpu/summean-broadcast.wgsl.jinja +13 -16
build/webgpu/summean-vec4.wgsl.jinja +4 -7
build/webgpu/test.json +1 -2

README.md CHANGED Viewed

@@ -18,18 +18,18 @@ See the [ONNX `Sum` spec](https://onnx.ai/onnx/operators/onnx__Sum.html) for the
 ## 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
@@ -39,7 +39,7 @@ See the [ONNX `Sum` spec](https://onnx.ai/onnx/operators/onnx__Sum.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
@@ -49,10 +49,14 @@ See the [ONNX `Sum` spec](https://onnx.ai/onnx/operators/onnx__Sum.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 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 | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
+| `y` | `sum` | `T` | derived | derived | Elementwise sum of all provided input tensors. | 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.Sum",
   "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
   "cases": [
     {

 {
   "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
   "cases": [
     {

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

@@ -1,6 +1,3 @@
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;





1	{{ env.wgsl.resourceDeclarations }}
2
3	const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;

build/webgpu/manifest.json CHANGED Viewed

@@ -2,200 +2,71 @@
   "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 ~ \">\"",
@@ -205,17 +76,27 @@
         {
           "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 ~ \">\"",
@@ -225,17 +106,22 @@
         {
           "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 ~ \">\"",
@@ -246,37 +132,41 @@
         {
           "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" }
         }
       ]
     },
@@ -284,96 +174,102 @@
       "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" }
         }
       ]
     },
@@ -381,42 +277,44 @@
       "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" }
         }
       ]
     }

   "domain": "ai.onnx",
   "name": "Sum",
   "sinceVersion": 13,
+  "inputs": {
+    "a": { "onnx": "A", "dtype": "T" },
+    "b": { "onnx": "B", "dtype": "T", "optional": true },
+    "c": { "onnx": "C", "dtype": "T", "optional": true },
+    "d": { "onnx": "D", "dtype": "T", "optional": true }
+  },
+  "outputs": {
+    "y": {
+      "onnx": "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"
     }
   },
+  "typeConstraints": { "T": ["float32", "float16"] },
+  "tunables": { "WORKGROUP_SIZE": { "default": 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)))",
+    "scalar": "dtypes.T",
+    "usesF16": "dtypes.T == \"f16\""
   },
+  "bindings": {
+    "params": { "buffer": "uniform", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }] },
+    "y_2": { "name": "y", "buffer": "storage", "elementType": "$vectorScalar" },
+    "params_2": {
+      "name": "params",
+      "buffer": "uniform",
+      "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.y) / 4" }]
+    },
+    "a_2": { "name": "a", "buffer": "read-only-storage", "elementType": "$aElement" },
+    "b_2": { "name": "b", "buffer": "read-only-storage", "elementType": "$bElement" },
+    "c": { "buffer": "read-only-storage", "elementType": "$cElement" },
+    "d": { "buffer": "read-only-storage", "elementType": "$dElement" }
   },
   "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": [{ "arg": "a", "name": "x", "elementType": "$scalar" }, "y", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
+            "y": 1,
+            "z": 1
+          }
         }
       ]
     },
     {
       "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"],
+      "derive": {
         "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
         "aElement": "\"vec4<\" ~ dtypes.T ~ \">\"",
         "bElement": "\"vec4<\" ~ dtypes.T ~ \">\"",
         {
           "id": "main",
           "name": "Sum.vec4",
+          "shader": "summean-vec4.wgsl.jinja",
+          "derive": { "op": "\"sum\"", "hasC": "false" },
+          "bindings": [
+            { "arg": "a", "elementType": "$vectorScalar" },
+            { "arg": "b", "elementType": "$vectorScalar" },
+            "y_2",
+            "params_2"
+          ],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "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"],
+      "derive": {
         "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
         "aElement": "\"vec4<\" ~ dtypes.T ~ \">\"",
         "bElement": "\"vec4<\" ~ dtypes.T ~ \">\"",
         {
           "id": "main",
           "name": "Sum.vec4_3",
+          "shader": "summean-vec4.wgsl.jinja",
+          "derive": { "op": "\"sum\"", "hasC": "true" },
+          "bindings": ["a_2", "b_2", "c", "y_2", "params_2"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "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"],
+      "derive": {
         "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
         "aElement": "\"vec4<\" ~ dtypes.T ~ \">\"",
         "bElement": "\"vec4<\" ~ dtypes.T ~ \">\"",
         {
           "id": "main",
           "name": "Sum.vec4_3",
+          "shader": "summean-vec4.wgsl.jinja",
+          "derive": { "op": "\"sum\"", "hasC": "true", "hasD": "true" },
+          "bindings": ["a_2", "b_2", "c", "d", "y_2", "params_2"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "broadcast_two_input",
+      "when": ["present.b", "not present.c", "ranks.a <= ranks.y", "ranks.b <= ranks.y", "broadcastOutputOk"],
       "passes": [
         {
           "id": "main",
           "name": "Sum",
+          "shader": "summean-broadcast.wgsl.jinja",
+          "derive": {
+            "aShape": "shapes.a",
+            "bShape": "shapes.b",
+            "yShape": "shapes.y",
+            "aRank": "ranks.a",
+            "bRank": "ranks.b",
+            "yRank": "ranks.y",
+            "hasC": "false",
+            "op": "\"sum\""
           },
+          "bindings": ["a", "b", "y", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
       "id": "broadcast_three_input_vec4",
       "priority": 24,
       "when": ["present.b", "present.c", "not present.d", "broadcastVec4OutputOk", "aBroadcastVec4Ok", "bBroadcastVec4Ok", "cBroadcastVec4Ok"],
+      "derive": {
         "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",
+          "shader": "summean-broadcast.wgsl.jinja",
+          "derive": {
+            "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,
+            "vectorizedSpec": true,
+            "op": "\"sum\""
           },
+          "bindings": ["a_2", "b_2", "c", "y_2", "params_2"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "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",
+          "shader": "summean-broadcast.wgsl.jinja",
+          "derive": {
+            "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": ["a", "b", { "arg": "c" }, "y", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "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",
+          "shader": "summean-broadcast.wgsl.jinja",
+          "derive": {
+            "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": ["a", "b", { "arg": "c" }, { "arg": "d" }, "y", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
       "id": "broadcast_four_input_vec4",
       "priority": 26,
       "when": ["present.b", "present.c", "present.d", "broadcastVec4OutputOk", "aBroadcastVec4Ok", "bBroadcastVec4Ok", "cBroadcastVec4Ok", "dBroadcastVec4Ok"],
+      "derive": {
         "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",
+          "shader": "summean-broadcast.wgsl.jinja",
+          "derive": {
+            "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,
+            "vectorizedSpec": true,
+            "op": "\"sum\""
           },
+          "bindings": ["a_2", "b_2", "c", "d", "y_2", "params_2"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     }

build/webgpu/metadata.json CHANGED Viewed

@@ -1,20 +1,33 @@
 {
   "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" }
 }

 {
   "name": "ai.onnx.Sum",
+  "id": "_ai_onnx_sum_webgpu_897002a",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "6IfhKTJQtyCXiNDSgCS1fq/K6MnfDkTN4OOlHyk+tAA=",
+      "datamove-elementwise-copy.wgsl.jinja": "Q1WuZCqcDf9b6cbT4rLilpeWj6327t/bSFmKr6IVBoQ=",
+      "manifest.json": "U1BSd6o0oGGE29P1Qs81OguhhV+/J3ahevzT4CTe1+Q=",
+      "summean-broadcast.wgsl.jinja": "PoxdcI7l73NJlE9atrddvWDgeaeBi6F92AFEcqSA+sA=",
+      "summean-vec4.wgsl.jinja": "3v0ep58YgrM8lkb5b3jqOiDR/FxnKdH8ORIFrWEcLqM=",
+      "test.json": "rjLE2xy25wmywX1N+nSDlWvfBH5AfD1J2/R5WWSx83M="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": {
+      "single_input_identity": ["datamove-elementwise-copy.wgsl.jinja"],
+      "same_shape_vec4_two_input": ["summean-vec4.wgsl.jinja"],
+      "same_shape_vec4_three_input": ["summean-vec4.wgsl.jinja"],
+      "same_shape_vec4_four_input": ["summean-vec4.wgsl.jinja"],
+      "broadcast_two_input": ["summean-broadcast.wgsl.jinja"],
+      "broadcast_three_input_vec4": ["summean-broadcast.wgsl.jinja"],
+      "broadcast_three_input": ["summean-broadcast.wgsl.jinja"],
+      "broadcast_four_input": ["summean-broadcast.wgsl.jinja"],
+      "broadcast_four_input_vec4": ["summean-broadcast.wgsl.jinja"]
+    }
+  }
 }

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

@@ -1,8 +1,5 @@
-{% 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 {
@@ -68,40 +65,40 @@ fn {{ fn_name }}({% if out_numel != 0 and op_numel != 1 %}out_index: u32{% endif
-{{ 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 %}
 }

+{% set allInputs = ["a", "b"] + (extraInputs if extraInputs is defined else (["c"] if hasC else []) + (["d"] if hasD else [])) %}
+{% set extraInputs = extraInputs if extraInputs is defined else (["c"] if hasC else []) + (["d"] if hasD else []) %}
 {{ 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 {
+{{ broadcast_offset_fn("a_offset", aShape, aRank, yShape, yRank) }}
+{{ broadcast_offset_fn("b_offset", bShape, bRank, yShape, yRank) }}
 {% for n in extraInputs %}
+{{ broadcast_offset_fn(n ~ "_offset", derive[n ~ "Shape"], derive[n ~ "Rank"], yShape, yRank) }}
 {% endfor %}
 @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 past the
+  // per-axis dispatch fold width.
+  let i = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if (i >= params.count) {
     return;
   }
+{% if vectorizedSpec %}
   // 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 derive[n ~ "Vector"] %}
+  let {{ n }}v = vec4<f32>({{ n }}[{{ broadcast_offset_call(n ~ "_offset", derive[n ~ "Shape"], yShape, "base") }} / 4u]);
 {% else %}
+  let {{ n }}v = vec4<f32>(f32({{ n }}[{{ broadcast_offset_call(n ~ "_offset", derive[n ~ "Shape"], 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", derive[n ~ "Shape"], yShape, "i") }}]){% if not loop.last %} + {% endif %}{% endfor %};
   y[i] = {{ scalar }}(total);
 {% endif %}
 }

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

@@ -1,14 +1,11 @@
-{% 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;
   }

+{% set extraInputs = extraInputs if extraInputs is defined else (["c"] if hasC else []) + (["d"] if hasD else []) %}
 {{ env.wgsl.resourceDeclarations }}
 @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 past the
+  // per-axis dispatch fold width.
+  let i = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if (i >= params.count) {
     return;
   }

build/webgpu/test.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "ai.onnx.Sum",
   "cases": [
     {
       "name": "max_arity_float16_positions",
@@ -71,7 +70,7 @@
       "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] } },

 {
   "cases": [
     {
       "name": "max_arity_float16_positions",
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
         "test": "MathOpTest.Sum_8_Test1",
+        "notes": "Vectorized variadic Sum must preserve finite subnormal totals."
       },
       "inputs": {
         "a": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1e-40, -1e-40, 0.0, 1e-39] } },