sync 91d990483a17

Browse files

Files changed (7) hide show

README.md +14 -10
build/webgpu/bench.json +3 -4
build/webgpu/manifest.json +81 -217
build/webgpu/metadata.json +16 -8
build/webgpu/test.json +4 -7
build/webgpu/where-broadcast.wgsl.jinja +22 -33
build/webgpu/where-vec4.wgsl.jinja +4 -7

README.md CHANGED Viewed

@@ -18,17 +18,17 @@ See the [ONNX `Where` spec](https://onnx.ai/onnx/operators/onnx__Where.html) for
 ## Inputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
-| `condition` | `condition` | `C` | — | — | Boolean mask; nonzero entries select from X, zero entries select from Y. | required |
-| `X` | `x` | `T` | — | — | Values selected at indices where condition is true. | required |
-| `Y` | `y` | `T` | — | — | Values selected at indices where condition is false. | required |
 ## Outputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
-| --- | --- | --- | --- | --- | --- | --- |
-| `output` | `output` | `T` | derived | broadcast result of `condition`, `X`, and `Y` | Tensor of shape equal to the broadcasted shape of condition, X, and Y. | required |
 ## Type constraints
@@ -39,7 +39,7 @@ See the [ONNX `Where` spec](https://onnx.ai/onnx/operators/onnx__Where.html) for
 ## 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
@@ -48,10 +48,14 @@ See the [ONNX `Where` spec](https://onnx.ai/onnx/operators/onnx__Where.html) for
 ## 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 |
 | --- | --- | --- | --- | --- | --- | --- |
+| `condition` | — | `C` | — | — | Boolean mask; nonzero entries select from X, zero entries select from Y. | required |
+| `x` | `X` | `T` | — | — | Values selected at indices where condition is true. | required |
+| `y` | `Y` | `T` | — | — | Values selected at indices where condition is false. | required |
 ## Outputs
+| Name | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- |
+| `output` | `T` | derived | broadcast result of `condition`, `x`, and `y` | Tensor of shape equal to the broadcasted shape of condition, X, and Y. | 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.Where",
   "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
   "cases": [
     {
@@ -7,7 +6,7 @@
       "preset": "smoke",
       "vars": { "dtype": "float32", "count": 1048576 },
       "inputs": {
-        "condition": { "shape": [1048576], "dtype": "bool", "dist": "uniform", "seed": 870, "scale": 1 },
         "x": { "shape": [1048576], "dtype": "float32", "dist": "normal", "seed": 871, "scale": 1 },
         "y": { "shape": [1048576], "dtype": "float32", "dist": "normal", "seed": 872, "scale": 1 }
       },
@@ -59,7 +58,7 @@
       "preset": "smoke",
       "vars": { "dtype": "float16", "count": 1048576 },
       "inputs": {
-        "condition": { "shape": [1048576], "dtype": "bool", "dist": "uniform", "seed": 970, "scale": 1 },
         "x": { "shape": [1048576], "dtype": "float16", "dist": "normal", "seed": 971, "scale": 1 },
         "y": { "shape": [1048576], "dtype": "float16", "dist": "normal", "seed": 972, "scale": 1 }
       },
@@ -73,7 +72,7 @@
       "preset": "smoke",
       "vars": { "dtype": "float32", "count": 1048575 },
       "inputs": {
-        "condition": { "shape": [1048575], "dtype": "bool", "dist": "uniform", "seed": 973, "scale": 1 },
         "x": { "shape": [1048575], "dtype": "float32", "dist": "normal", "seed": 974, "scale": 1 },
         "y": { "shape": [1048575], "dtype": "float32", "dist": "normal", "seed": 975, "scale": 1 }
       },

 {
   "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
   "cases": [
     {
       "preset": "smoke",
       "vars": { "dtype": "float32", "count": 1048576 },
       "inputs": {
+        "condition": { "shape": [1048576], "dtype": "bool", "dist": "uniform", "seed": 870 },
         "x": { "shape": [1048576], "dtype": "float32", "dist": "normal", "seed": 871, "scale": 1 },
         "y": { "shape": [1048576], "dtype": "float32", "dist": "normal", "seed": 872, "scale": 1 }
       },
       "preset": "smoke",
       "vars": { "dtype": "float16", "count": 1048576 },
       "inputs": {
+        "condition": { "shape": [1048576], "dtype": "bool", "dist": "uniform", "seed": 970 },
         "x": { "shape": [1048576], "dtype": "float16", "dist": "normal", "seed": 971, "scale": 1 },
         "y": { "shape": [1048576], "dtype": "float16", "dist": "normal", "seed": 972, "scale": 1 }
       },
       "preset": "smoke",
       "vars": { "dtype": "float32", "count": 1048575 },
       "inputs": {
+        "condition": { "shape": [1048575], "dtype": "bool", "dist": "uniform", "seed": 973 },
         "x": { "shape": [1048575], "dtype": "float32", "dist": "normal", "seed": 974, "scale": 1 },
         "y": { "shape": [1048575], "dtype": "float32", "dist": "normal", "seed": 975, "scale": 1 }
       },

build/webgpu/manifest.json CHANGED Viewed

@@ -2,268 +2,132 @@
   "domain": "ai.onnx",
   "name": "Where",
   "sinceVersion": 16,
-  "description": "Selects elements from `X` or `Y` according to a boolean `condition`, following NumPy-style multidirectional broadcasting. Where `condition` is `true`, the output takes the value from `X`; otherwise it takes the value from `Y`.",
-  "inputs": [
-    {
-      "role": "condition",
-      "dtype": "C",
-      "description": "Boolean mask; nonzero entries select from X, zero entries select from Y."
-    },
-    { "role": "X", "dtype": "T", "description": "Values selected at indices where condition is true." },
-    { "role": "Y", "dtype": "T", "description": "Values selected at indices where condition is false." }
-  ],
-  "outputs": [
-    {
-      "role": "output",
       "dtype": "T",
-      "rank": "max(max(ranks.condition, ranks.X), ranks.Y)",
-      "description": "Tensor of shape equal to the broadcasted shape of condition, X, and Y.",
-      "shape": "broadcastShape(broadcastShape(shapes.condition, shapes.X), shapes.Y)"
     }
-  ],
   "typeConstraints": { "C": ["bool"], "T": ["float32", "float16", "uint32", "int32", "int16", "uint8", "int8", "bool"] },
-  "args": {
-    "condition": { "kind": "tensor", "semantic": "condition", "role": "input" },
-    "x": { "kind": "tensor", "semantic": "X", "role": "input" },
-    "y": { "kind": "tensor", "semantic": "Y", "role": "input" },
-    "output": { "kind": "tensor", "semantic": "output", "role": "output" }
   },
-  "tunables": { "WORKGROUP_SIZE": 256 },
-  "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
   "variants": [
     {
       "id": "broadcast",
-      "when": ["ranks.condition <= ranks.output", "ranks.X <= ranks.output", "ranks.Y <= ranks.output", "f16Ok(dtypes.T)"],
       "passes": [
         {
           "id": "main",
           "name": "Where",
-          "source": {
-            "shader": "where-broadcast.wgsl.jinja",
-            "inputs": {
-              "cShape": "shapes.condition",
-              "xShape": "shapes.X",
-              "yShape": "shapes.Y",
-              "oShape": "shapes.output",
-              "cRank": "ranks.condition",
-              "xRank": "ranks.X",
-              "yRank": "ranks.Y",
-              "oRank": "ranks.output",
-              "itemsPerInvocation": 4
-            }
           },
           "bindings": [
-            {
-              "name": "condition",
-              "arg": "condition",
-              "semantic": "condition",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "u32"
-            },
-            {
-              "name": "x",
-              "arg": "x",
-              "semantic": "X",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "$scalar"
-            },
-            {
-              "name": "y",
-              "arg": "y",
-              "semantic": "Y",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "$scalar"
-            },
-            {
-              "name": "output",
-              "arg": "output",
-              "semantic": "output",
-              "buffer": { "type": "storage" },
-              "elementType": "$scalar"
-            },
-            {
-              "name": "params",
-              "semantic": "kernel.params",
-              "buffer": { "type": "uniform" },
-              "struct": {
-                "name": "Params",
-                "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.output)" }]
-              }
-            }
           ],
-          "dispatch": { "threads": "ceilDiv(numel(shapes.output), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "broadcast_inner_vec4",
       "priority": 1,
-      "when": ["ranks.output >= 1", "numel(shapes.output) > 0", "dim(shapes.output, ranks.output - 1) % 4 == 0", "(ranks.condition == 0 or dim(shapes.condition, ranks.condition - 1) == 1 or dim(shapes.condition, ranks.condition - 1) == dim(shapes.output, ranks.output - 1))", "(ranks.X == 0 or dim(shapes.X, ranks.X - 1) == 1 or dim(shapes.X, ranks.X - 1) == dim(shapes.output, ranks.output - 1))", "(ranks.Y == 0 or dim(shapes.Y, ranks.Y - 1) == 1 or dim(shapes.Y, ranks.Y - 1) == dim(shapes.output, ranks.output - 1))", "ranks.condition <= ranks.output", "ranks.X <= ranks.output", "ranks.Y <= ranks.output", "f16Ok(dtypes.T)"],
-      "constants": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "Where",
-          "source": {
-            "shader": "where-broadcast.wgsl.jinja",
-            "inputs": {
-              "cShape": "shapes.condition",
-              "xShape": "shapes.X",
-              "yShape": "shapes.Y",
-              "oShape": "shapes.output",
-              "cRank": "ranks.condition",
-              "xRank": "ranks.X",
-              "yRank": "ranks.Y",
-              "oRank": "ranks.output",
-              "vec4Out": true
-            }
           },
-          "bindings": [
-            {
-              "name": "condition",
-              "arg": "condition",
-              "semantic": "condition",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "u32"
-            },
-            {
-              "name": "x",
-              "arg": "x",
-              "semantic": "X",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "$scalar"
-            },
-            {
-              "name": "y",
-              "arg": "y",
-              "semantic": "Y",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "$scalar"
-            },
-            {
-              "name": "output",
-              "arg": "output",
-              "semantic": "output",
-              "buffer": { "type": "storage" },
-              "elementType": "$vectorScalar"
-            },
-            {
-              "name": "params",
-              "semantic": "kernel.params",
-              "buffer": { "type": "uniform" },
-              "struct": {
-                "name": "Params",
-                "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.output) / 4" }]
-              }
-            }
-          ],
-          "dispatch": { "threads": "numel(shapes.output) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "all_same_shape_vec4",
       "priority": 30,
-      "when": ["sameShape(shapes.condition, shapes.output)", "sameShape(shapes.X, shapes.output)", "sameShape(shapes.Y, shapes.output)", "numel(shapes.output) > 0", "numel(shapes.output) % 4 == 0", "f16Ok(dtypes.T)"],
-      "constants": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "Where.vec4",
-          "source": { "shader": "where-vec4.wgsl.jinja", "inputs": { "condMode": "\"vector\"" } },
-          "bindings": [
-            {
-              "name": "condition",
-              "arg": "condition",
-              "semantic": "condition",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "vec4<u32>"
-            },
-            {
-              "name": "x",
-              "arg": "x",
-              "semantic": "X",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "$vectorScalar"
-            },
-            {
-              "name": "y",
-              "arg": "y",
-              "semantic": "Y",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "$vectorScalar"
-            },
-            {
-              "name": "output",
-              "arg": "output",
-              "semantic": "output",
-              "buffer": { "type": "storage" },
-              "elementType": "$vectorScalar"
-            },
-            {
-              "name": "params",
-              "semantic": "kernel.params",
-              "buffer": { "type": "uniform" },
-              "struct": {
-                "name": "Params",
-                "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.output) / 4" }]
-              }
-            }
-          ],
-          "dispatch": { "threads": "numel(shapes.output) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "scalar_condition_vec4",
       "priority": 25,
-      "when": ["numel(shapes.condition) == 1", "sameShape(shapes.X, shapes.output)", "sameShape(shapes.Y, shapes.output)", "numel(shapes.output) > 0", "numel(shapes.output) % 4 == 0", "f16Ok(dtypes.T)"],
-      "constants": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "Where.scalarCondVec4",
-          "source": { "shader": "where-vec4.wgsl.jinja", "inputs": { "condMode": "\"scalar\"" } },
-          "bindings": [
-            {
-              "name": "condition",
-              "arg": "condition",
-              "semantic": "condition",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "u32"
-            },
-            {
-              "name": "x",
-              "arg": "x",
-              "semantic": "X",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "$vectorScalar"
-            },
-            {
-              "name": "y",
-              "arg": "y",
-              "semantic": "Y",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "$vectorScalar"
-            },
-            {
-              "name": "output",
-              "arg": "output",
-              "semantic": "output",
-              "buffer": { "type": "storage" },
-              "elementType": "$vectorScalar"
-            },
-            {
-              "name": "params",
-              "semantic": "kernel.params",
-              "buffer": { "type": "uniform" },
-              "struct": {
-                "name": "Params",
-                "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.output) / 4" }]
-              }
-            }
-          ],
-          "dispatch": { "threads": "numel(shapes.output) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     }

   "domain": "ai.onnx",
   "name": "Where",
   "sinceVersion": 16,
+  "inputs": { "condition": { "dtype": "C" }, "x": { "onnx": "X", "dtype": "T" }, "y": { "onnx": "Y", "dtype": "T" } },
+  "outputs": {
+    "output": {
       "dtype": "T",
+      "rank": "max(max(ranks.condition, ranks.x), ranks.y)",
+      "shape": "broadcastShape(broadcastShape(shapes.condition, shapes.x), shapes.y)"
     }
+  },
   "typeConstraints": { "C": ["bool"], "T": ["float32", "float16", "uint32", "int32", "int16", "uint8", "int8", "bool"] },
+  "tunables": { "WORKGROUP_SIZE": { "default": 256 } },
+  "derive": { "scalar": "dtypes.T" },
+  "when": ["f16Ok(dtypes.T)"],
+  "bindings": {
+    "condition": { "buffer": "read-only-storage", "elementType": "u32" },
+    "output_2": { "name": "output", "buffer": "storage", "elementType": "$vectorScalar" },
+    "params_2": {
+      "name": "params",
+      "buffer": "uniform",
+      "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.output) / 4" }]
+    },
+    "condition_2": { "name": "condition", "buffer": "read-only-storage", "elementType": "vec4<u32>" },
+    "x_2": { "name": "x", "buffer": "read-only-storage", "elementType": "$vectorScalar" },
+    "y_2": { "name": "y", "buffer": "read-only-storage", "elementType": "$vectorScalar" }
   },
   "variants": [
     {
       "id": "broadcast",
+      "when": ["ranks.condition <= ranks.output", "ranks.x <= ranks.output", "ranks.y <= ranks.output"],
       "passes": [
         {
           "id": "main",
           "name": "Where",
+          "shader": "where-broadcast.wgsl.jinja",
+          "derive": {
+            "cShape": "shapes.condition",
+            "xShape": "shapes.x",
+            "yShape": "shapes.y",
+            "oShape": "shapes.output",
+            "cRank": "ranks.condition",
+            "xRank": "ranks.x",
+            "yRank": "ranks.y",
+            "oRank": "ranks.output",
+            "itemsPerInvocation": 4
           },
           "bindings": [
+            "condition",
+            "x",
+            "y",
+            "output",
+            { "name": "params", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.output)" }] }
           ],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.output), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.output), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "broadcast_inner_vec4",
       "priority": 1,
+      "when": ["ranks.output >= 1", "numel(shapes.output) > 0", "dim(shapes.output, ranks.output - 1) % 4 == 0", "(ranks.condition == 0 or dim(shapes.condition, ranks.condition - 1) == 1 or dim(shapes.condition, ranks.condition - 1) == dim(shapes.output, ranks.output - 1))", "(ranks.x == 0 or dim(shapes.x, ranks.x - 1) == 1 or dim(shapes.x, ranks.x - 1) == dim(shapes.output, ranks.output - 1))", "(ranks.y == 0 or dim(shapes.y, ranks.y - 1) == 1 or dim(shapes.y, ranks.y - 1) == dim(shapes.output, ranks.output - 1))", "ranks.condition <= ranks.output", "ranks.x <= ranks.output", "ranks.y <= ranks.output"],
+      "derive": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "Where",
+          "shader": "where-broadcast.wgsl.jinja",
+          "derive": {
+            "cShape": "shapes.condition",
+            "xShape": "shapes.x",
+            "yShape": "shapes.y",
+            "oShape": "shapes.output",
+            "cRank": "ranks.condition",
+            "xRank": "ranks.x",
+            "yRank": "ranks.y",
+            "oRank": "ranks.output",
+            "vec4Out": true
           },
+          "bindings": ["condition", "x", "y", "output_2", "params_2"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.output) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.output) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "all_same_shape_vec4",
       "priority": 30,
+      "when": ["sameShape(shapes.condition, shapes.output)", "sameShape(shapes.x, shapes.output)", "sameShape(shapes.y, shapes.output)", "numel(shapes.output) > 0", "numel(shapes.output) % 4 == 0"],
+      "derive": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "Where.vec4",
+          "shader": "where-vec4.wgsl.jinja",
+          "derive": { "condMode": "\"vector\"" },
+          "bindings": ["condition_2", "x_2", "y_2", "output_2", "params_2"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.output) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.output) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "scalar_condition_vec4",
       "priority": 25,
+      "when": ["numel(shapes.condition) == 1", "sameShape(shapes.x, shapes.output)", "sameShape(shapes.y, shapes.output)", "numel(shapes.output) > 0", "numel(shapes.output) % 4 == 0"],
+      "derive": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "Where.scalarCondVec4",
+          "shader": "where-vec4.wgsl.jinja",
+          "derive": { "condMode": "\"scalar\"" },
+          "bindings": ["condition", "x_2", "y_2", "output_2", "params_2"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.output) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.output) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     }

build/webgpu/metadata.json CHANGED Viewed

@@ -1,19 +1,27 @@
 {
   "name": "ai.onnx.Where",
-  "id": "_ai_onnx_where_webgpu_ffdaf20",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "VyEb3ADRUxJJwyaSc9I4fNcWHPy0yqIKtT8Bm3H1WX8=",
-      "manifest.json": "q6/trJ9PjIgZtZLhHHnZeBm6zSHhY6xTpkTGJtQmsZY=",
-      "test.json": "BeLATZd7ICyT+vvjrY04kpYaDrb3BPPBZUo676qxX24=",
-      "where-broadcast.wgsl.jinja": "xPK1v4t7GBnKJl4YYPF0J6lf3u2pUilCgvefOgKlauE=",
-      "where-vec4.wgsl.jinja": "EXTMQw4HgDmNJp+QU4QyFhep584Ch0kd4x1OdB7KC8w="
     }
   },
-  "provenance": { "kernel": { "sha": "c928d21e6cc1310861cba3bafb75f5f679ecf5f3", "dirty": false } },
-  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Where" }
 }

 {
   "name": "ai.onnx.Where",
+  "id": "_ai_onnx_where_webgpu_5d1b103",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "skxPQI/vUW/oEssq1hdPgFGDyPZYLC0kDHmeyzLR2yY=",
+      "manifest.json": "64IxOIxpn6DCKhM59kpc8GKd98yH/HeAA/d/Gp+UZLo=",
+      "test.json": "OCgBwQVsPNU0gH93TL5AWMMsMy2aG9EGLiia2+Xi0GA=",
+      "where-broadcast.wgsl.jinja": "bYDdm4/2eF9cUf/EGEMys3FeP970NKGKJRafMO+ZxHQ=",
+      "where-vec4.wgsl.jinja": "u2PIkiP+3lcTS/Df7xkzylyTNEONTHgEHZSV3c14ZyE="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": {
+      "broadcast": ["where-broadcast.wgsl.jinja"],
+      "broadcast_inner_vec4": ["where-broadcast.wgsl.jinja"],
+      "all_same_shape_vec4": ["where-vec4.wgsl.jinja"],
+      "scalar_condition_vec4": ["where-vec4.wgsl.jinja"]
+    }
+  }
 }

build/webgpu/test.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "ai.onnx.Where",
   "cases": [
     {
       "name": "int16_same_shape_boundaries",
@@ -144,7 +143,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/where_op_test.cc",
         "test": "WhereOpTest.BasicNumeric",
-        "notes": "ORT covers float and double; this framework also checks the same branch pattern for float16."
       },
       "inputs": {
         "condition": { "dtype": "bool", "shape": [2, 2], "data": { "kind": "values", "values": [0, 1, 1, 0] } },
@@ -184,7 +183,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/where_op_test.cc",
         "test": "WhereOpTest.BroadcastWithScalar",
-        "notes": "ORT uses int64 values; this framework covers the same scalar broadcast shape behavior with float32 data."
       },
       "inputs": {
         "condition": { "dtype": "bool", "shape": [3], "data": { "kind": "values", "values": [1, 0, 1] } },
@@ -198,7 +197,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/where_op_test.cc",
         "test": "WhereOpTest.BroadcastWithScalar",
-        "notes": "ORT uses int64 values; this framework covers the same scalar broadcast shape behavior with int32 data."
       },
       "inputs": {
         "condition": { "dtype": "bool", "shape": [3], "data": { "kind": "values", "values": [1, 0, 1] } },
@@ -308,9 +307,7 @@
     },
     {
       "name": "scalar_condition_true_vec4_4096",
-      "provenance": {
-        "notes": "Compact sibling for the scalar-condition Where benchmark; preserves a shape-[1] condition broadcasting over a vec4-aligned f32 payload."
-      },
       "inputs": {
         "condition": { "dtype": "bool", "shape": [1], "data": { "kind": "values", "values": [1] } },
         "x": {

 {
   "cases": [
     {
       "name": "int16_same_shape_boundaries",
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/where_op_test.cc",
         "test": "WhereOpTest.BasicNumeric",
+        "notes": "A float16 payload exercises the branch pattern with reduced-precision storage."
       },
       "inputs": {
         "condition": { "dtype": "bool", "shape": [2, 2], "data": { "kind": "values", "values": [0, 1, 1, 0] } },
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/where_op_test.cc",
         "test": "WhereOpTest.BroadcastWithScalar",
+        "notes": "A float32 payload exercises scalar broadcast shape behavior."
       },
       "inputs": {
         "condition": { "dtype": "bool", "shape": [3], "data": { "kind": "values", "values": [1, 0, 1] } },
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/where_op_test.cc",
         "test": "WhereOpTest.BroadcastWithScalar",
+        "notes": "An int32 payload exercises scalar broadcast shape behavior."
       },
       "inputs": {
         "condition": { "dtype": "bool", "shape": [3], "data": { "kind": "values", "values": [1, 0, 1] } },
     },
     {
       "name": "scalar_condition_true_vec4_4096",
+      "provenance": { "notes": "A shape-[1] condition broadcasts over a vec4-aligned float32 payload." },
       "inputs": {
         "condition": { "dtype": "bool", "shape": [1], "data": { "kind": "values", "values": [1] } },
         "x": {

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

@@ -1,31 +1,20 @@
 {% macro flat_tail_open() %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
-fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
   // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
-  // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills the rest into y).
-  let invocation = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
-{% if source.itemsPerInvocation is defined %}
   // Tail-safe scalar x4 keeps vector-like dispatch density without requiring
   // the logical tensor length (or its storage binding) to be vec4 aligned.
-  let begin = invocation * {{ source.itemsPerInvocation }}u;
-  let end = min(begin + {{ source.itemsPerInvocation }}u, params.count);
   for (var i = begin; i < end; i = i + 1u) {
-{%- else %}
-  let i = invocation;
-  if (i >= params.count) {
-    return;
-  }
-{%- endif %}
-{% endmacro %}
 {% macro flat_tail_close() %}
-{% if source.itemsPerInvocation is defined %}
   }
-{% endif %}
 {% endmacro %}
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 {% macro offset_fn(fn_name, opShape, opRank, op_same, op_numel, outShape, outRank, out_numel) %}
@@ -91,33 +80,33 @@ fn {{ fn_name }}({% if out_numel != 0 and op_numel != 1 %}out_index: u32{% endif
 {%- endmacro %}
-{{ broadcast_offset_fn("c_offset", source.cShape, source.cRank, source.oShape, source.oRank) }}
-{{ broadcast_offset_fn("x_offset", source.xShape, source.xRank, source.oShape, source.oRank) }}
-{{ broadcast_offset_fn("y_offset", source.yShape, source.yRank, source.oShape, source.oRank) }}
-{% set vec4Out = source.vec4Out | default(false) %}
 {% if vec4Out %}
 // One invocation handles one output vec4. Its four elements share every outer
 // coordinate. Each input offset either advances by one when its last axis
 // matches the output, or remains fixed when that axis is broadcast.
-{% set o_last = source.oShape[source.oRank - 1] %}
-{% set c_inner = source.cRank > 0 and source.cShape[source.cRank - 1] == o_last %}
-{% set x_inner = source.xRank > 0 and source.xShape[source.xRank - 1] == o_last %}
-{% set y_inner = source.yRank > 0 and source.yShape[source.yRank - 1] == o_last %}
 @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 vi = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
   if (vi >= params.count) {
     return;
   }
   let base = vi * 4u;
-  let cb = {{ broadcast_offset_call("c_offset", source.cShape, source.oShape, "base") }};
-  let xb = {{ broadcast_offset_call("x_offset", source.xShape, source.oShape, "base") }};
-  let yb = {{ broadcast_offset_call("y_offset", source.yShape, source.oShape, "base") }};
 {% if c_inner %}
   let cond = vec4<u32>(condition[cb], condition[cb + 1u], condition[cb + 2u], condition[cb + 3u]);
 {% else %}
@@ -137,7 +126,7 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
 }
 {% else %}
 {{ flat_tail_open() }}
-  output[i] = select(y[{{ broadcast_offset_call("y_offset", source.yShape, source.oShape, "i") }}], x[{{ broadcast_offset_call("x_offset", source.xShape, source.oShape, "i") }}], condition[{{ broadcast_offset_call("c_offset", source.cShape, source.oShape, "i") }}] != 0u);
 {{ flat_tail_close() -}}
 }
 {% endif %}

 {% macro flat_tail_open() %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
+  // dispatch's per-axis workgroup fold width (the dispatch caps x and spills the rest into y).
+  let invocation = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   // Tail-safe scalar x4 keeps vector-like dispatch density without requiring
   // the logical tensor length (or its storage binding) to be vec4 aligned.
+{% set itemsPerInvocation = itemsPerInvocation if itemsPerInvocation is defined else 4 %}
+  let begin = invocation * {{ itemsPerInvocation }}u;
+  let end = min(begin + {{ itemsPerInvocation }}u, params.count);
   for (var i = begin; i < end; i = i + 1u) {
+{%- endmacro %}
 {% macro flat_tail_close() %}
   }
 {% endmacro %}
 {{ env.wgsl.resourceDeclarations }}
 {% macro offset_fn(fn_name, opShape, opRank, op_same, op_numel, outShape, outRank, out_numel) %}
 {%- endmacro %}
+{{ broadcast_offset_fn("c_offset", cShape, cRank, oShape, oRank) }}
+{{ broadcast_offset_fn("x_offset", xShape, xRank, oShape, oRank) }}
+{{ broadcast_offset_fn("y_offset", yShape, yRank, oShape, oRank) }}
+{% set vec4Out = vec4Out | default(false) %}
 {% if vec4Out %}
 // One invocation handles one output vec4. Its four elements share every outer
 // coordinate. Each input offset either advances by one when its last axis
 // matches the output, or remains fixed when that axis is broadcast.
+{% set o_last = oShape[oRank - 1] %}
+{% set c_inner = cRank > 0 and cShape[cRank - 1] == o_last %}
+{% set x_inner = xRank > 0 and xShape[xRank - 1] == o_last %}
+{% set y_inner = yRank > 0 and yShape[yRank - 1] == o_last %}
 @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 vi = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if (vi >= params.count) {
     return;
   }
   let base = vi * 4u;
+  let cb = {{ broadcast_offset_call("c_offset", cShape, oShape, "base") }};
+  let xb = {{ broadcast_offset_call("x_offset", xShape, oShape, "base") }};
+  let yb = {{ broadcast_offset_call("y_offset", yShape, oShape, "base") }};
 {% if c_inner %}
   let cond = vec4<u32>(condition[cb], condition[cb + 1u], condition[cb + 2u], condition[cb + 3u]);
 {% else %}
 }
 {% else %}
 {{ flat_tail_open() }}
+  output[i] = select(y[{{ broadcast_offset_call("y_offset", yShape, oShape, "i") }}], x[{{ broadcast_offset_call("x_offset", xShape, oShape, "i") }}], condition[{{ broadcast_offset_call("c_offset", cShape, oShape, "i") }}] != 0u);
 {{ flat_tail_close() -}}
 }
 {% endif %}

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

@@ -1,17 +1,14 @@
-{% 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;
   }
-{% if source.condMode == "scalar" %}
   output[i] = select(y[i], x[i], condition[0] != 0u);
 {% else %}
   output[i] = select(y[i], x[i], condition[i] != vec4<u32>(0u));

 {{ 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;
   }
+{% if condMode == "scalar" %}
   output[i] = select(y[i], x[i], condition[0] != 0u);
 {% else %}
   output[i] = select(y[i], x[i], condition[i] != vec4<u32>(0u));