sync 91d990483a17

Browse files

Files changed (8) hide show

README.md +15 -11
build/webgpu/bench.json +0 -1
build/webgpu/datamove-elementwise-copy.wgsl.jinja +0 -3
build/webgpu/manifest.json +181 -360
build/webgpu/metadata.json +22 -9
build/webgpu/minmax-broadcast.wgsl.jinja +16 -18
build/webgpu/minmax-vec4.wgsl.jinja +4 -7
build/webgpu/test.json +15 -13

README.md CHANGED Viewed

@@ -18,19 +18,19 @@ See the [ONNX `Min` spec](https://onnx.ai/onnx/operators/onnx__Min.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; broadcast-compatible with the other inputs. | optional |
-| `C` | `c` | `T` | — | — | Third input tensor; broadcast-compatible with the other inputs. | optional |
-| `D` | `d` | `T` | — | — | Fourth input tensor; broadcast-compatible with the other inputs. | optional |
-| `E` | `e` | `T` | — | — | Fifth input tensor, broadcast-compatible with all other inputs. | optional |
 ## Outputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
-| `min` | `y` | `T` | derived | derived; see description | Elementwise minimum of all input tensors, broadcast to the output shape. | required |
 ## Type constraints
@@ -40,7 +40,7 @@ See the [ONNX `Min` spec](https://onnx.ai/onnx/operators/onnx__Min.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
@@ -50,10 +50,14 @@ See the [ONNX `Min` spec](https://onnx.ai/onnx/operators/onnx__Min.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; broadcast-compatible with the other inputs. | optional |
+| `c` | `C` | `T` | — | — | Third input tensor; broadcast-compatible with the other inputs. | optional |
+| `d` | `D` | `T` | — | — | Fourth input tensor; broadcast-compatible with the other inputs. | optional |
+| `e` | `E` | `T` | — | — | Fifth input tensor, broadcast-compatible with all other inputs. | optional |
 ## Outputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
+| `y` | `min` | `T` | derived | derived | Elementwise minimum of all input tensors, broadcast to the output shape. | 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.Min",
   "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,277 +2,76 @@
   "domain": "ai.onnx",
   "name": "Min",
   "sinceVersion": 13,
-  "description": "Computes the elementwise minimum across one or more input tensors with NumPy-style multidirectional broadcasting. All inputs and the output share the same data type.",
-  "inputs": [
-    { "role": "A", "dtype": "T", "description": "First input tensor." },
-    {
-      "role": "B",
-      "dtype": "T",
-      "optional": true,
-      "description": "Second input tensor; broadcast-compatible with the other inputs."
-    },
-    {
-      "role": "C",
-      "dtype": "T",
-      "optional": true,
-      "description": "Third input tensor; broadcast-compatible with the other inputs."
-    },
-    {
-      "role": "D",
-      "dtype": "T",
-      "optional": true,
-      "description": "Fourth input tensor; broadcast-compatible with the other inputs."
-    },
-    {
-      "role": "E",
-      "dtype": "T",
-      "optional": true,
-      "description": "Fifth input tensor, broadcast-compatible with all other inputs."
-    }
-  ],
-  "outputs": [
-    {
-      "role": "min",
       "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, ranks.E if present.b and present.c and present.d and present.e else 0)",
-      "shape": "variadicShape",
-      "description": "Elementwise minimum of all input tensors, broadcast to the output shape."
     }
-  ],
-  "typeConstraints": { "T": ["float32", "float16", "int32", "uint32", "int16", "int8", "uint8"] },
-  "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": "min", "role": "output" },
-    "d": { "kind": "tensor", "semantic": "D", "role": "input3", "required": false },
-    "e": { "kind": "tensor", "semantic": "E", "role": "input4", "required": false }
   },
-  "tunables": { "WORKGROUP_SIZE": 256 },
   "derive": {
     "variadicInputCount": "1 + (1 if present.b else 0) + (1 if (present.b and present.c) else 0) + (1 if (present.b and present.c and present.d) else 0) + (1 if (present.b and present.c and present.d and present.e) else 0)",
-    "variadicShape": "broadcastShape(broadcastShape(broadcastShape(broadcastShape(shapes.A, shapes.B), shapes.C), shapes.D), shapes.E) if present.b and present.c and present.d and present.e else (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)",
     "narrowNoSubgroupExecution": "not device.features.has(\"subgroups\") and has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize < 32 and device.adapterInfo.subgroupMaxSize <= 32"
   },
-  "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": "$vectorScalar"
-      },
-      {
-        "name": "b",
-        "arg": "b",
-        "semantic": "B",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$vectorScalar"
-      },
-      {
-        "name": "c",
-        "arg": "c",
-        "semantic": "C",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$vectorScalar"
-      },
-      { "name": "y", "arg": "y", "semantic": "min", "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": "$vectorScalar"
-      },
-      {
-        "name": "b",
-        "arg": "b",
-        "semantic": "B",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$vectorScalar"
-      },
-      {
-        "name": "c",
-        "arg": "c",
-        "semantic": "C",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$vectorScalar"
-      },
-      {
-        "name": "d",
-        "arg": "d",
-        "semantic": "D",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$vectorScalar"
-      },
-      { "name": "y", "arg": "y", "semantic": "min", "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)" }] }
-      }
-    ],
-    "vec4Five": [
-      {
-        "name": "a",
-        "arg": "a",
-        "semantic": "A",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$vectorScalar"
-      },
-      {
-        "name": "b",
-        "arg": "b",
-        "semantic": "B",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$vectorScalar"
-      },
-      {
-        "name": "c",
-        "arg": "c",
-        "semantic": "C",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$vectorScalar"
-      },
-      {
-        "name": "d",
-        "arg": "d",
-        "semantic": "D",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$vectorScalar"
-      },
-      {
-        "name": "e",
-        "arg": "e",
-        "semantic": "E",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$vectorScalar"
-      },
-      { "name": "y", "arg": "y", "semantic": "min", "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" }] }
-      }
-    ],
-    "scalarFive": [
-      { "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": "e", "arg": "e", "semantic": "E", "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": ["variadicInputCount == 1", "ranks.A == ranks.y", "numel(shapes.A) == numel(shapes.y)", "f16Ok(dtypes.T)"],
-      "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
       "passes": [
         {
           "id": "main",
           "name": "Min.Identity",
-          "source": { "shader": "datamove-elementwise-copy.wgsl.jinja" },
-          "bindings": "identity",
-          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "same_shape_vec4_two_input",
       "priority": 20,
-      "when": ["variadicInputCount == 2", "sameShape(shapes.A, shapes.y)", "sameShape(shapes.B, shapes.y)", "flatVec4OutputOk", "not narrowNoSubgroupExecution"],
-      "constants": {
         "scalar": "dtypes.T",
         "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
         "usesF16": "dtypes.T == \"f16\""
@@ -281,17 +80,22 @@
         {
           "id": "main",
           "name": "Min.vec4",
-          "source": { "shader": "minmax-vec4.wgsl.jinja", "inputs": { "op": "\"min\"", "hasC": "false" } },
-          "bindings": "vec4Two",
-          "dispatch": { "threads": "numel(shapes.y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "same_shape_vec4_three_input",
       "priority": 25,
-      "when": ["variadicInputCount == 3", "sameShape(shapes.A, shapes.y)", "sameShape(shapes.B, shapes.y)", "sameShape(shapes.C, shapes.y)", "flatVec4OutputOk"],
-      "constants": {
         "scalar": "dtypes.T",
         "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
         "usesF16": "dtypes.T == \"f16\""
@@ -300,17 +104,22 @@
         {
           "id": "main",
           "name": "Min.vec4_3",
-          "source": { "shader": "minmax-vec4.wgsl.jinja", "inputs": { "op": "\"min\"", "hasC": "true" } },
-          "bindings": "vec4Three",
-          "dispatch": { "threads": "numel(shapes.y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "same_shape_vec4_four_input",
       "priority": 27,
-      "when": ["variadicInputCount == 4", "sameShape(shapes.A, shapes.y)", "sameShape(shapes.B, shapes.y)", "sameShape(shapes.C, shapes.y)", "sameShape(shapes.D, shapes.y)", "flatVec4OutputOk"],
-      "constants": {
         "scalar": "dtypes.T",
         "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
         "usesF16": "dtypes.T == \"f16\""
@@ -319,20 +128,22 @@
         {
           "id": "main",
           "name": "Min.vec4_3",
-          "source": {
-            "shader": "minmax-vec4.wgsl.jinja",
-            "inputs": { "op": "\"min\"", "hasC": "true", "hasD": "true" }
-          },
-          "bindings": "vec4Four",
-          "dispatch": { "threads": "numel(shapes.y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "same_shape_vec4_five_input",
       "priority": 28,
-      "when": ["variadicInputCount == 5", "sameShape(shapes.A, shapes.y)", "sameShape(shapes.B, shapes.y)", "sameShape(shapes.C, shapes.y)", "sameShape(shapes.D, shapes.y)", "sameShape(shapes.E, shapes.y)", "flatVec4OutputOk"],
-      "constants": {
         "scalar": "dtypes.T",
         "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
         "usesF16": "dtypes.T == \"f16\""
@@ -341,145 +152,155 @@
         {
           "id": "main",
           "name": "Min.vec4_4",
-          "source": {
-            "shader": "minmax-vec4.wgsl.jinja",
-            "inputs": {
-              "op": "\"min\"",
-              "hasC": "true",
-              "hasD": "true",
-              "hasE": "true",
-              "extraInputs": "[\"c\", \"d\", \"e\"]"
-            }
           },
-          "bindings": "vec4Five",
-          "dispatch": { "threads": "numel(shapes.y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "broadcast_two_input",
-      "when": ["variadicInputCount == 2", "ranks.A <= ranks.y", "ranks.B <= ranks.y", "broadcastOutputOk"],
-      "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
       "passes": [
         {
           "id": "main",
           "name": "Min",
-          "source": {
-            "shader": "minmax-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": "\"min\"",
-              "itemsPerInvocation": 4
-            }
           },
-          "bindings": "scalarTwo",
-          "dispatch": { "threads": "ceilDiv(numel(shapes.y), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "broadcast_three_input",
       "priority": 20,
-      "when": ["variadicInputCount == 3", "ranks.A <= ranks.y", "ranks.B <= ranks.y", "ranks.C <= ranks.y", "broadcastOutputOk"],
-      "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
       "passes": [
         {
           "id": "main",
           "name": "Min",
-          "source": {
-            "shader": "minmax-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": "\"min\"",
-              "itemsPerInvocation": 4
-            }
           },
-          "bindings": "scalarThree",
-          "dispatch": { "threads": "ceilDiv(numel(shapes.y), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "broadcast_four_input",
       "priority": 22,
-      "when": ["variadicInputCount == 4", "ranks.A <= ranks.y", "ranks.B <= ranks.y", "ranks.C <= ranks.y", "ranks.D <= ranks.y", "broadcastOutputOk"],
-      "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
       "passes": [
         {
           "id": "main",
           "name": "Min",
-          "source": {
-            "shader": "minmax-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": "\"min\"",
-              "dShape": "shapes.D",
-              "dRank": "ranks.D",
-              "hasD": "true",
-              "itemsPerInvocation": 4
-            }
           },
-          "bindings": "scalarFour",
-          "dispatch": { "threads": "ceilDiv(numel(shapes.y), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "broadcast_five_input",
       "priority": 23,
-      "when": ["variadicInputCount == 5", "ranks.A <= ranks.y", "ranks.B <= ranks.y", "ranks.C <= ranks.y", "ranks.D <= ranks.y", "ranks.E <= ranks.y", "broadcastOutputOk"],
-      "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
       "passes": [
         {
           "id": "main",
           "name": "Min",
-          "source": {
-            "shader": "minmax-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": "\"min\"",
-              "dShape": "shapes.D",
-              "dRank": "ranks.D",
-              "hasD": "true",
-              "hasE": "true",
-              "extraInputs": "[\"c\", \"d\", \"e\"]",
-              "eRank": "ranks.E",
-              "eShape": "shapes.E",
-              "itemsPerInvocation": 4
-            }
           },
-          "bindings": "scalarFive",
-          "dispatch": { "threads": "ceilDiv(numel(shapes.y), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     }

   "domain": "ai.onnx",
   "name": "Min",
   "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 },
+    "e": { "onnx": "E", "dtype": "T", "optional": true }
+  },
+  "outputs": {
+    "y": {
+      "onnx": "min",
       "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, ranks.e if present.b and present.c and present.d and present.e else 0)",
+      "shape": "variadicShape"
     }
   },
+  "typeConstraints": { "T": ["float32", "float16", "int32", "uint32", "int16", "int8", "uint8"] },
+  "tunables": { "WORKGROUP_SIZE": { "default": 256 } },
   "derive": {
     "variadicInputCount": "1 + (1 if present.b else 0) + (1 if (present.b and present.c) else 0) + (1 if (present.b and present.c and present.d) else 0) + (1 if (present.b and present.c and present.d and present.e) else 0)",
+    "variadicShape": "broadcastShape(broadcastShape(broadcastShape(broadcastShape(shapes.a, shapes.b), shapes.c), shapes.d), shapes.e) if present.b and present.c and present.d and present.e else (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)",
     "narrowNoSubgroupExecution": "not device.features.has(\"subgroups\") and has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize < 32 and device.adapterInfo.subgroupMaxSize <= 32"
   },
+  "bindings": {
+    "x": { "arg": "a", "buffer": "read-only-storage", "elementType": "$scalar" },
+    "y": { "buffer": "storage", "elementType": "$scalar" },
+    "params": { "buffer": "uniform", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }] },
+    "a": { "buffer": "read-only-storage", "elementType": "$vectorScalar" },
+    "b": { "buffer": "read-only-storage", "elementType": "$vectorScalar" },
+    "y_2": { "name": "y", "buffer": "storage", "elementType": "$vectorScalar" },
+    "params_2": {
+      "name": "params",
+      "buffer": "uniform",
+      "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.y) / 4" }]
+    },
+    "c": { "buffer": "read-only-storage", "elementType": "$vectorScalar" },
+    "d": { "buffer": "read-only-storage", "elementType": "$vectorScalar" },
+    "e": { "buffer": "read-only-storage", "elementType": "$vectorScalar" },
+    "a_2": { "name": "a", "buffer": "read-only-storage", "elementType": "$scalar" },
+    "b_2": { "name": "b", "buffer": "read-only-storage", "elementType": "$scalar" },
+    "c_2": { "name": "c", "buffer": "read-only-storage", "elementType": "$scalar" },
+    "d_2": { "name": "d", "buffer": "read-only-storage", "elementType": "$scalar" },
+    "e_2": { "name": "e", "buffer": "read-only-storage", "elementType": "$scalar" }
   },
   "variants": [
     {
       "id": "single_input_identity",
       "priority": 30,
+      "when": ["variadicInputCount == 1", "ranks.a == ranks.y", "numel(shapes.a) == numel(shapes.y)", "f16Ok(dtypes.T)"],
+      "derive": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
       "passes": [
         {
           "id": "main",
           "name": "Min.Identity",
+          "shader": "datamove-elementwise-copy.wgsl.jinja",
+          "bindings": ["x", "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": "same_shape_vec4_two_input",
       "priority": 20,
+      "when": ["variadicInputCount == 2", "sameShape(shapes.a, shapes.y)", "sameShape(shapes.b, shapes.y)", "flatVec4OutputOk", "not narrowNoSubgroupExecution"],
+      "derive": {
         "scalar": "dtypes.T",
         "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
         "usesF16": "dtypes.T == \"f16\""
         {
           "id": "main",
           "name": "Min.vec4",
+          "shader": "minmax-vec4.wgsl.jinja",
+          "derive": { "op": "\"min\"", "hasC": "false" },
+          "bindings": ["a", "b", "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": ["variadicInputCount == 3", "sameShape(shapes.a, shapes.y)", "sameShape(shapes.b, shapes.y)", "sameShape(shapes.c, shapes.y)", "flatVec4OutputOk"],
+      "derive": {
         "scalar": "dtypes.T",
         "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
         "usesF16": "dtypes.T == \"f16\""
         {
           "id": "main",
           "name": "Min.vec4_3",
+          "shader": "minmax-vec4.wgsl.jinja",
+          "derive": { "op": "\"min\"", "hasC": "true" },
+          "bindings": ["a", "b", "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": ["variadicInputCount == 4", "sameShape(shapes.a, shapes.y)", "sameShape(shapes.b, shapes.y)", "sameShape(shapes.c, shapes.y)", "sameShape(shapes.d, shapes.y)", "flatVec4OutputOk"],
+      "derive": {
         "scalar": "dtypes.T",
         "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
         "usesF16": "dtypes.T == \"f16\""
         {
           "id": "main",
           "name": "Min.vec4_3",
+          "shader": "minmax-vec4.wgsl.jinja",
+          "derive": { "op": "\"min\"", "hasC": "true", "hasD": "true" },
+          "bindings": ["a", "b", "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": "same_shape_vec4_five_input",
       "priority": 28,
+      "when": ["variadicInputCount == 5", "sameShape(shapes.a, shapes.y)", "sameShape(shapes.b, shapes.y)", "sameShape(shapes.c, shapes.y)", "sameShape(shapes.d, shapes.y)", "sameShape(shapes.e, shapes.y)", "flatVec4OutputOk"],
+      "derive": {
         "scalar": "dtypes.T",
         "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
         "usesF16": "dtypes.T == \"f16\""
         {
           "id": "main",
           "name": "Min.vec4_4",
+          "shader": "minmax-vec4.wgsl.jinja",
+          "derive": {
+            "op": "\"min\"",
+            "hasC": "true",
+            "hasD": "true",
+            "hasE": "true",
+            "extraInputs": "[\"c\", \"d\", \"e\"]"
           },
+          "bindings": ["a", "b", "c", "d", "e", "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": ["variadicInputCount == 2", "ranks.a <= ranks.y", "ranks.b <= ranks.y", "broadcastOutputOk"],
+      "derive": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
       "passes": [
         {
           "id": "main",
           "name": "Min",
+          "shader": "minmax-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": "\"min\"",
+            "itemsPerInvocation": 4
           },
+          "bindings": ["a_2", "b_2", "y", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.y), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.y), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "broadcast_three_input",
       "priority": 20,
+      "when": ["variadicInputCount == 3", "ranks.a <= ranks.y", "ranks.b <= ranks.y", "ranks.c <= ranks.y", "broadcastOutputOk"],
+      "derive": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
       "passes": [
         {
           "id": "main",
           "name": "Min",
+          "shader": "minmax-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": "\"min\"",
+            "itemsPerInvocation": 4
           },
+          "bindings": ["a_2", "b_2", "c_2", "y", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.y), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.y), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "broadcast_four_input",
       "priority": 22,
+      "when": ["variadicInputCount == 4", "ranks.a <= ranks.y", "ranks.b <= ranks.y", "ranks.c <= ranks.y", "ranks.d <= ranks.y", "broadcastOutputOk"],
+      "derive": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
       "passes": [
         {
           "id": "main",
           "name": "Min",
+          "shader": "minmax-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": "\"min\"",
+            "dShape": "shapes.d",
+            "dRank": "ranks.d",
+            "hasD": "true",
+            "itemsPerInvocation": 4
           },
+          "bindings": ["a_2", "b_2", "c_2", "d_2", "y", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.y), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.y), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "broadcast_five_input",
       "priority": 23,
+      "when": ["variadicInputCount == 5", "ranks.a <= ranks.y", "ranks.b <= ranks.y", "ranks.c <= ranks.y", "ranks.d <= ranks.y", "ranks.e <= ranks.y", "broadcastOutputOk"],
+      "derive": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
       "passes": [
         {
           "id": "main",
           "name": "Min",
+          "shader": "minmax-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": "\"min\"",
+            "dShape": "shapes.d",
+            "dRank": "ranks.d",
+            "hasD": "true",
+            "hasE": "true",
+            "extraInputs": "[\"c\", \"d\", \"e\"]",
+            "eRank": "ranks.e",
+            "eShape": "shapes.e",
+            "itemsPerInvocation": 4
           },
+          "bindings": ["a_2", "b_2", "c_2", "d_2", "e_2", "y", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.y), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(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.Min",
-  "id": "_ai_onnx_min_webgpu_30c0530",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "4AfCZQh/G8SMpQs4IgS9e2U8PUeiDMUzt8TWOgNvQsg=",
-      "datamove-elementwise-copy.wgsl.jinja": "J5yC2bAddPiP+odXLgVGS3TJ9jeNsfRTvedKrj/fhZg=",
-      "manifest.json": "cnSm6aRoRIHenuDqHLRwEXoPyRnT/7qmAEYYEBDzkAY=",
-      "minmax-broadcast.wgsl.jinja": "dDLFWoUuV4Q+I/R6Ee+zGCbeStKZ2h/ObmLzlK1XdP0=",
-      "minmax-vec4.wgsl.jinja": "oS72DBVc9P7XoADR4T16+G9FALJyQzFR4hxnagD8VB4=",
-      "test.json": "pZblIGVry/Tin/efCEfKrchJ7gk9kdu/yLZO93hefBg="
     }
   },
-  "provenance": { "kernel": { "sha": "c928d21e6cc1310861cba3bafb75f5f679ecf5f3", "dirty": false } },
-  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Min" }
 }

 {
   "name": "ai.onnx.Min",
+  "id": "_ai_onnx_min_webgpu_25bdece",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "11Avk6oIAzN8DKPeDG+sjNh7o3JWi7VQRkjfOrFX1Ns=",
+      "datamove-elementwise-copy.wgsl.jinja": "Q1WuZCqcDf9b6cbT4rLilpeWj6327t/bSFmKr6IVBoQ=",
+      "manifest.json": "tpGUvXih1rgUtp9pH1559keeuj97lh9IHo+AmGSf+uE=",
+      "minmax-broadcast.wgsl.jinja": "VfDVtRI1nLx4PtZ05JccODRFVvZXCjfe9T/TYUKJzsw=",
+      "minmax-vec4.wgsl.jinja": "ykNWo2M5sTSFRXgwTWuULv4PPZOHiRtkqR9OuJEgu5s=",
+      "test.json": "5BzaEWeLO5YgM4RltLBpE2RH76bJHtEjlxB4dx+NLgA="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": {
+      "single_input_identity": ["datamove-elementwise-copy.wgsl.jinja"],
+      "same_shape_vec4_two_input": ["minmax-vec4.wgsl.jinja"],
+      "same_shape_vec4_three_input": ["minmax-vec4.wgsl.jinja"],
+      "same_shape_vec4_four_input": ["minmax-vec4.wgsl.jinja"],
+      "same_shape_vec4_five_input": ["minmax-vec4.wgsl.jinja"],
+      "broadcast_two_input": ["minmax-broadcast.wgsl.jinja"],
+      "broadcast_three_input": ["minmax-broadcast.wgsl.jinja"],
+      "broadcast_four_input": ["minmax-broadcast.wgsl.jinja"],
+      "broadcast_five_input": ["minmax-broadcast.wgsl.jinja"]
+    }
+  }
 }

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

@@ -1,23 +1,21 @@
 {% macro flat_tail_open() %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
-fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
   // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
-  // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills the rest into y).
-  let invocation = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
   // Tail-safe scalar x4 keeps vector-like dispatch density without requiring
   // the logical tensor length (or its storage binding) to be vec4 aligned.
-  let begin = invocation * {{ source.itemsPerInvocation }}u;
-  let end = min(begin + {{ source.itemsPerInvocation }}u, params.count);
   for (var i = begin; i < end; i = i + 1u) {
 {%- endmacro %}
 {% macro flat_tail_close() %}
   }
 {% endmacro %}
-{% 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 }}
 {% if scalar != "i32" and scalar != "u32" %}
@@ -91,34 +89,34 @@ 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 %}
 {{ flat_tail_open() }}
 {% if scalar == "i32" or scalar == "u32" %}
-  let av = a[{{ broadcast_offset_call("a_offset", source.aShape, source.yShape, "i") }}];
-  let bv = b[{{ broadcast_offset_call("b_offset", source.bShape, source.yShape, "i") }}];
   var out = select(bv, av, av < bv);
 {% for n in extraInputs %}
-  let {{ n }}v = {{ n }}[{{ broadcast_offset_call(n ~ "_offset", source[n ~ "Shape"], source.yShape, "i") }}];
   out = select({{ n }}v, out, out < {{ n }}v);
 {% endfor %}
   y[i] = out;
 {% else %}
-  let av = f32(a[{{ broadcast_offset_call("a_offset", source.aShape, source.yShape, "i") }}]);
-  let bv = f32(b[{{ broadcast_offset_call("b_offset", source.bShape, source.yShape, "i") }}]);
   var out = min(av, bv);
   out = select(out, av, is_nan_f32(av));
   out = select(out, bv, is_nan_f32(bv));
 {% for n in extraInputs %}
-  let {{ n }}v = f32({{ n }}[{{ broadcast_offset_call(n ~ "_offset", source[n ~ "Shape"], source.yShape, "i") }}]);
   let m{{ n }} = min(out, {{ n }}v);
 {% if loop.first %}
   // Built-in min/max may drop a NaN operand. Re-inject {{ n }}v's NaN and

 {% 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 %}
+{% set extraInputs = extraInputs if extraInputs is defined else (["c"] if hasC else []) + (["d"] if hasD else []) %}
 {{ env.wgsl.resourceDeclarations }}
 {% if scalar != "i32" and scalar != "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 %}
 {{ flat_tail_open() }}
 {% if scalar == "i32" or scalar == "u32" %}
+  let av = a[{{ broadcast_offset_call("a_offset", aShape, yShape, "i") }}];
+  let bv = b[{{ broadcast_offset_call("b_offset", bShape, yShape, "i") }}];
   var out = select(bv, av, av < bv);
 {% for n in extraInputs %}
+  let {{ n }}v = {{ n }}[{{ broadcast_offset_call(n ~ "_offset", derive[n ~ "Shape"], yShape, "i") }}];
   out = select({{ n }}v, out, out < {{ n }}v);
 {% endfor %}
   y[i] = out;
 {% else %}
+  let av = f32(a[{{ broadcast_offset_call("a_offset", aShape, yShape, "i") }}]);
+  let bv = f32(b[{{ broadcast_offset_call("b_offset", bShape, yShape, "i") }}]);
   var out = min(av, bv);
   out = select(out, av, is_nan_f32(av));
   out = select(out, bv, is_nan_f32(bv));
 {% for n in extraInputs %}
+  let {{ n }}v = f32({{ n }}[{{ broadcast_offset_call(n ~ "_offset", derive[n ~ "Shape"], yShape, "i") }}]);
   let m{{ n }} = min(out, {{ n }}v);
 {% if loop.first %}
   // Built-in min/max may drop a NaN operand. Re-inject {{ n }}v's NaN and

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

@@ -1,7 +1,4 @@
-{% 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 }}
 {% if scalar != "i32" and scalar != "u32" %}
@@ -22,10 +19,10 @@ fn is_nan_vec4(value: vec4<f32>) -> vec4<bool> {
 {% endif %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
-fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
   // 2D-folded flat index: gid.y carries the high bits 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 }}
 {% if scalar != "i32" and scalar != "u32" %}
 {% endif %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   // 2D-folded flat index: gid.y carries the high bits 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.Min",
   "cases": [
     {
       "name": "int16_max_arity_boundaries",
@@ -147,7 +146,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
         "test": "MathOpTest.Min_8",
-        "notes": "Scalar broadcast companion: negative subnormal inputs remain less than zero."
       },
       "inputs": {
         "a": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-1e-40, 0.0, 1e-40] } },
@@ -284,7 +283,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
         "test": "MathOpTest.Min_8",
-        "notes": "Extends ORT's same-shape Min coverage to a valid four-input ONNX variadic node."
       },
       "inputs": {
         "a": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "constant", "value": -1.0 } },
@@ -299,7 +298,7 @@
     {
       "name": "four_input_same_shape_vec4_reference_generated",
       "provenance": {
-        "notes": "Reference-generated four-operand coverage; see the Max sibling. Equal-amplitude interleaved sinusoids mean D supplies the minimum on a substantial share of lanes, so a reference that stops at C is visible here."
       },
       "inputs": {
         "a": {
@@ -330,7 +329,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
         "test": "MathOpTest.Min_8",
-        "notes": "Extends ORT's variadic Min coverage with a fourth input that carries NaNs; any NaN input should propagate at that element."
       },
       "inputs": {
         "a": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, "NaN", 5.0, -4.0] } },
@@ -457,10 +456,11 @@
       "name": "ort_float_scalar_nan_broadcast",
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
-        "test": "MathOpTest.Min_12_Float_with_scalar_Nan"
       },
       "inputs": {
-        "a": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "values", "values": [0.25, -0.25, -0.5, 0.5] } },
         "b": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": ["NaN"] } }
       },
       "outputs": { "y": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.000001, "allowNaN": true } }
@@ -525,10 +525,11 @@
       "name": "ort_f16_scalar_nan_broadcast",
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
-        "test": "MathOpTest.Min_13_Float16_with_scalar_Nan"
       },
       "inputs": {
-        "a": { "dtype": "float16", "shape": [3, 1], "data": { "kind": "values", "values": [-0.5, 1.0, 1.5] } },
         "b": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": ["NaN"] } }
       },
       "outputs": { "y": { "dtype": "float16", "shape": [3, 1], "tolerance": 0.002, "allowNaN": true } }
@@ -577,12 +578,13 @@
       "name": "ort_f16_scalar1_three_input",
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
-        "test": "MathOpTest.Min_12_MLFloat16_Scalar1"
       },
       "inputs": {
-        "a": { "dtype": "float16", "shape": [1, 3], "data": { "kind": "values", "values": [2.0, 3.0, 4.0] } },
         "b": { "dtype": "float16", "shape": [], "data": { "kind": "values", "values": [-10.0] } },
-        "c": { "dtype": "float16", "shape": [1, 3], "data": { "kind": "values", "values": [3.0, 2.0, -3.0] } }
       },
       "outputs": {
         "y": {
@@ -972,7 +974,7 @@
     {
       "name": "f16_four_input_same_shape_vec4",
       "provenance": {
-        "notes": "float16 at arity 4 on the flat vec4 same-shape path (same_shape_vec4_four_input, numel 8). Two extra-operand folds in the f16 lane of minmax-vec4, so the second fold consumes the first fold's result and the loop.first-guarded comment renders on the first pass only. Valid ONNX Min-13 for T = float16 with four same-shape operands. Every value is exactly representable in float16 and each of a, b, c, d supplies the minimum on exactly two lanes, so no operand can be dropped."
       },
       "inputs": {
         "a": {

 {
   "cases": [
     {
       "name": "int16_max_arity_boundaries",
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
         "test": "MathOpTest.Min_8",
+        "notes": "On the scalar broadcast path, negative subnormal inputs remain less than zero."
       },
       "inputs": {
         "a": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-1e-40, 0.0, 1e-40] } },
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
         "test": "MathOpTest.Min_8",
+        "notes": "Four same-shaped inputs exercise the valid maximum arity of this variadic Min package."
       },
       "inputs": {
         "a": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "constant", "value": -1.0 } },
     {
       "name": "four_input_same_shape_vec4_reference_generated",
       "provenance": {
+        "notes": "Four equal-amplitude interleaved sinusoids make the fourth operand supply the minimum on many lanes, so omitting any operand changes the output well beyond tolerance."
       },
       "inputs": {
         "a": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
         "test": "MathOpTest.Min_8",
+        "notes": "A fourth input carrying NaNs verifies that any NaN operand propagates at that element."
       },
       "inputs": {
         "a": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, "NaN", 5.0, -4.0] } },
       "name": "ort_float_scalar_nan_broadcast",
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Min_12_Float_with_scalar_Nan",
+        "notes": "Diverges from the upstream test's inputs (inputs.a values [0.25, -0.25, -0.5, 0.5] -> constant 0.25); the expected output is recomputed by the CPU reference for the new inputs. A NaN scalar operand propagates to every element whatever the other operand holds, so the tensor operand is uniform: any finite result fails."
       },
       "inputs": {
+        "a": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "constant", "value": 0.25 } },
         "b": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": ["NaN"] } }
       },
       "outputs": { "y": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.000001, "allowNaN": true } }
       "name": "ort_f16_scalar_nan_broadcast",
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Min_13_Float16_with_scalar_Nan",
+        "notes": "Diverges from the upstream test's inputs (inputs.a values [-0.5, 1.0, 1.5] -> constant 1.0); the expected output is recomputed by the CPU reference for the new inputs. A NaN scalar operand propagates to every element whatever the other operand holds, so the tensor operand is uniform: any finite result fails."
       },
       "inputs": {
+        "a": { "dtype": "float16", "shape": [3, 1], "data": { "kind": "constant", "value": 1.0 } },
         "b": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": ["NaN"] } }
       },
       "outputs": { "y": { "dtype": "float16", "shape": [3, 1], "tolerance": 0.002, "allowNaN": true } }
       "name": "ort_f16_scalar1_three_input",
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Min_12_MLFloat16_Scalar1",
+        "notes": "Diverges from the upstream test's inputs (inputs.a values [2.0, 3.0, 4.0] -> constant 2.0; inputs.c values [3.0, 2.0, -3.0] -> constant 3.0); the expected output is recomputed by the CPU reference for the new inputs. The rank-0 operand dominates all three inputs, so the tensor operands are uniform: a variadic fold that drops the scalar returns the tensor value instead."
       },
       "inputs": {
+        "a": { "dtype": "float16", "shape": [1, 3], "data": { "kind": "constant", "value": 2.0 } },
         "b": { "dtype": "float16", "shape": [], "data": { "kind": "values", "values": [-10.0] } },
+        "c": { "dtype": "float16", "shape": [1, 3], "data": { "kind": "constant", "value": 3.0 } }
       },
       "outputs": {
         "y": {
     {
       "name": "f16_four_input_same_shape_vec4",
       "provenance": {
+        "notes": "Four same-shaped float16 inputs require two successive minimum folds on the vectorized path. Every value is exactly representable, and each operand supplies the minimum on exactly two lanes."
       },
       "inputs": {
         "a": {