sync 91d990483a17

Browse files

Files changed (7) hide show

README.md +15 -11
build/webgpu/bench.json +4 -5
build/webgpu/manifest.json +132 -327
build/webgpu/metadata.json +20 -8
build/webgpu/test.json +11 -20
build/webgpu/unary-scalar.wgsl.jinja +21 -16
build/webgpu/unary-vec4.wgsl.jinja +36 -17

README.md CHANGED Viewed

@@ -18,17 +18,17 @@ See the [ONNX `Clip` spec](https://onnx.ai/onnx/operators/onnx__Clip.html) for t
 ## Inputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
-| --- | --- | --- | --- | --- | --- | --- |
-| `input` | `input` | `T` | — | — | Input tensor whose elements are to be clipped. | required |
-| `min` | `min` | `T` | `0` | `[]` | Scalar lower bound; elements below this value are replaced by it. | optional |
-| `max` | `max` | `T` | `0` | `[]` | Scalar upper bound; elements above this value are replaced by it. | optional |
 ## Outputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
-| --- | --- | --- | --- | --- | --- | --- |
-| `output` | `output` | `T` | same as `input` | same as `input` | Output tensor with each element clipped to the specified interval. | required |
 ## Type constraints
@@ -38,7 +38,7 @@ See the [ONNX `Clip` spec](https://onnx.ai/onnx/operators/onnx__Clip.html) for t
 ## Files
-- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
 - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
 - [`test.json`](build/webgpu/test.json) — correctness cases
 - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
@@ -47,10 +47,14 @@ See the [ONNX `Clip` spec](https://onnx.ai/onnx/operators/onnx__Clip.html) for t
 ## 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 | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- |
+| `input` | `T` | — | — | Input tensor whose elements are to be clipped. | required |
+| `min` | `T` | `0` | `[]` | Scalar lower bound; elements below this value are replaced by it. | optional |
+| `max` | `T` | `0` | `[]` | Scalar upper bound; elements above this value are replaced by it. | optional |
 ## Outputs
+| Name | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- |
+| `output` | `T` | same as `input` | same as `input` | Output tensor with each element clipped to the specified interval. | 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.Clip",
   "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
   "cases": [
     {
@@ -38,8 +37,8 @@
     {
       "name": "clip-f32-input-bounds-large-scalar-path",
       "provenance": {
-        "source": "authored for performance coverage",
-        "notes": "Named for the scalar route it was authored to exercise. Its element count is divisible by four, so it now takes the vectorized input-bounds route; the scalar route stays live for counts that are not, and its fixtures still cover it."
       },
       "preset": "stress",
       "vars": { "dtype": "float32", "count": 8388608 },
@@ -54,8 +53,8 @@
     {
       "name": "clip-f16-8m-default-bounds-vec4",
       "provenance": {
-        "source": "authored for performance coverage",
-        "notes": "Clip with no min or max inputs, the only form that reaches the vec4 unary path. Nothing benchmarked that path, so the scalar input-bounds route had no vectorized sibling to be measured against."
       },
       "preset": "smoke",
       "vars": { "dtype": "float16", "count": 8388608 },

 {
   "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
   "cases": [
     {
     {
       "name": "clip-f32-input-bounds-large-scalar-path",
       "provenance": {
+        "source": "local fixture",
+        "notes": "The 8,388,608-element float32 input is divisible by four and selects the vectorized input-bounds route."
       },
       "preset": "stress",
       "vars": { "dtype": "float32", "count": 8388608 },
     {
       "name": "clip-f16-8m-default-bounds-vec4",
       "provenance": {
+        "source": "local fixture",
+        "notes": "Clip without min or max inputs measures the vec4 unary path independently of scalar input-bound handling."
       },
       "preset": "smoke",
       "vars": { "dtype": "float16", "count": 8388608 },

build/webgpu/manifest.json CHANGED Viewed

@@ -2,43 +2,14 @@
   "domain": "ai.onnx",
   "name": "Clip",
   "sinceVersion": 13,
-  "description": "Limits each element of the input tensor to the interval `[min, max]`, equivalent to `Min(max, Max(input, min))`. Each bound is an optional scalar input tensor; an omitted bound defaults to the input dtype's lowest or highest numeric value. When `min` exceeds `max`, all elements are set to `max`.",
-  "inputs": [
-    { "role": "input", "dtype": "T", "description": "Input tensor whose elements are to be clipped." },
-    {
-      "role": "min",
-      "dtype": "T",
-      "rank": 0,
-      "optional": true,
-      "shape": [],
-      "description": "Scalar lower bound; elements below this value are replaced by it."
-    },
-    {
-      "role": "max",
-      "dtype": "T",
-      "rank": 0,
-      "optional": true,
-      "shape": [],
-      "description": "Scalar upper bound; elements above this value are replaced by it."
-    }
-  ],
-  "outputs": [
-    {
-      "role": "output",
-      "dtype": "T",
-      "rank": "ranks.input",
-      "shape": "shapes.input",
-      "description": "Output tensor with each element clipped to the specified interval."
-    }
-  ],
-  "typeConstraints": { "T": ["float32", "float16", "int32", "uint32", "int8", "uint8"] },
-  "args": {
-    "input": { "kind": "tensor", "semantic": "input", "role": "input" },
-    "output": { "kind": "tensor", "semantic": "output", "role": "output" },
-    "min": { "kind": "tensor", "semantic": "min", "role": "input1", "required": false },
-    "max": { "kind": "tensor", "semantic": "max", "role": "input2", "required": false }
   },
-  "tunables": { "WORKGROUP_SIZE": 256 },
   "derive": {
     "shapeContract": "numel(shapes.input) == numel(shapes.output) and f16Ok(dtypes.T)",
     "defaultBoundsContract": "shapeContract and not present.min and not present.max",
@@ -47,271 +18,49 @@
     "inputBoundsContract": "shapeContract and present.min and present.max and ranks.min == 0 and numel(shapes.min) == 1 and ranks.max == 0 and numel(shapes.max) == 1",
     "defaultMin": "-65504 if tensorDtypes.input == \"float16\" else (-340282346638528859811704183484516925440 if tensorDtypes.input == \"float32\" else (-128 if tensorDtypes.input == \"int8\" else (-2147483648 if tensorDtypes.input == \"int32\" else 0)))",
     "defaultMax": "65504 if tensorDtypes.input == \"float16\" else (340282346638528859811704183484516925440 if tensorDtypes.input == \"float32\" else (127 if tensorDtypes.input == \"int8\" else (2147483647 if tensorDtypes.input == \"int32\" else (255 if tensorDtypes.input == \"uint8\" else 4294967295))))",
-    "wideVec4StorageOk": "dtypes.T != \"f32\" or device.features.has(\"subgroups\") or not (has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize == 8 and device.adapterInfo.subgroupMaxSize <= 32)"
   },
-  "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
-  "bindingSets": {
-    "defaultVector": [
-      {
-        "name": "x",
-        "arg": "input",
-        "semantic": "input",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$vectorScalar"
-      },
-      {
-        "name": "y",
-        "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" }]
-        }
-      }
-    ],
-    "defaultScalar": [
-      {
-        "name": "input",
-        "arg": "input",
-        "semantic": "input",
-        "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)" }] }
-      }
-    ],
-    "inputBounds": [
-      {
-        "name": "input",
-        "arg": "input",
-        "semantic": "input",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scalar"
-      },
-      {
-        "name": "output",
-        "arg": "output",
-        "semantic": "output",
-        "buffer": { "type": "storage" },
-        "elementType": "$scalar"
-      },
-      {
-        "name": "clipMin",
-        "arg": "min",
-        "semantic": "min",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scalar"
-      },
-      {
-        "name": "clipMax",
-        "arg": "max",
-        "semantic": "max",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scalar"
-      },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.output)" }] }
-      }
-    ],
-    "minInput": [
-      {
-        "name": "input",
-        "arg": "input",
-        "semantic": "input",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scalar"
-      },
-      {
-        "name": "output",
-        "arg": "output",
-        "semantic": "output",
-        "buffer": { "type": "storage" },
-        "elementType": "$scalar"
-      },
-      {
-        "name": "clipMin",
-        "arg": "min",
-        "semantic": "min",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scalar"
-      },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.output)" }] }
-      }
-    ],
-    "maxInput": [
-      {
-        "name": "input",
-        "arg": "input",
-        "semantic": "input",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scalar"
-      },
-      {
-        "name": "output",
-        "arg": "output",
-        "semantic": "output",
-        "buffer": { "type": "storage" },
-        "elementType": "$scalar"
-      },
-      {
-        "name": "clipMax",
-        "arg": "max",
-        "semantic": "max",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scalar"
-      },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.output)" }] }
-      }
-    ],
-    "inputBoundsVector": [
-      {
-        "name": "x",
-        "arg": "input",
-        "semantic": "input",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$vectorScalar"
-      },
-      {
-        "name": "y",
-        "arg": "output",
-        "semantic": "output",
-        "buffer": { "type": "storage" },
-        "elementType": "$vectorScalar"
-      },
-      {
-        "name": "clipMin",
-        "arg": "min",
-        "semantic": "min",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scalar"
-      },
-      {
-        "name": "clipMax",
-        "arg": "max",
-        "semantic": "max",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scalar"
-      },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": {
-          "name": "Params",
-          "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.output) / 4" }]
-        }
-      }
-    ],
-    "minInputVector": [
-      {
-        "name": "x",
-        "arg": "input",
-        "semantic": "input",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$vectorScalar"
-      },
-      {
-        "name": "y",
-        "arg": "output",
-        "semantic": "output",
-        "buffer": { "type": "storage" },
-        "elementType": "$vectorScalar"
-      },
-      {
-        "name": "clipMin",
-        "arg": "min",
-        "semantic": "min",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scalar"
-      },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": {
-          "name": "Params",
-          "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.output) / 4" }]
-        }
-      }
-    ],
-    "maxInputVector": [
-      {
-        "name": "x",
-        "arg": "input",
-        "semantic": "input",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$vectorScalar"
-      },
-      {
-        "name": "y",
-        "arg": "output",
-        "semantic": "output",
-        "buffer": { "type": "storage" },
-        "elementType": "$vectorScalar"
-      },
-      {
-        "name": "clipMax",
-        "arg": "max",
-        "semantic": "max",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scalar"
-      },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": {
-          "name": "Params",
-          "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.output) / 4" }]
-        }
-      }
-    ]
   },
   "variants": [
     {
       "id": "default_bounds_vec4",
       "priority": 20,
       "when": ["defaultBoundsContract", "numel(shapes.input) > 0", "numel(shapes.input) % 4 == 0", "wideVec4StorageOk"],
-      "constants": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "Clip.defaultBoundsVec4",
-          "source": {
-            "shader": "unary-vec4.wgsl.jinja",
-            "inputs": { "op": "\"clip\"", "boundsFromInput": false, "minValue": "defaultMin", "maxValue": "defaultMax" }
           },
-          "bindings": "defaultVector",
-          "dispatch": { "threads": "numel(shapes.output) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
@@ -322,12 +71,20 @@
         {
           "id": "main",
           "name": "Clip.defaultBounds",
-          "source": {
-            "shader": "unary-scalar.wgsl.jinja",
-            "inputs": { "op": "\"clip\"", "boundsFromInput": false, "minValue": "defaultMin", "maxValue": "defaultMax" }
           },
-          "bindings": "defaultScalar",
-          "dispatch": { "threads": "numel(shapes.output)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
@@ -335,17 +92,25 @@
       "id": "input_bounds_vec4",
       "priority": 40,
       "when": ["inputBoundsContract", "numel(shapes.input) > 0", "numel(shapes.input) % 4 == 0", "wideVec4StorageOk"],
-      "constants": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "Clip.inputBoundsVec4",
-          "source": {
-            "shader": "unary-vec4.wgsl.jinja",
-            "inputs": { "op": "\"clip\"", "boundsFromInput": true, "clipMinFromInput": true, "clipMaxFromInput": true }
           },
-          "bindings": "inputBoundsVector",
-          "dispatch": { "threads": "numel(shapes.output) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
@@ -353,17 +118,25 @@
       "id": "min_input_vec4",
       "priority": 40,
       "when": ["minInputContract", "numel(shapes.input) > 0", "numel(shapes.input) % 4 == 0", "wideVec4StorageOk"],
-      "constants": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "Clip.minInputVec4",
-          "source": {
-            "shader": "unary-vec4.wgsl.jinja",
-            "inputs": { "op": "\"clip\"", "boundsFromInput": true, "clipMinFromInput": true, "clipMaxFromInput": false }
           },
-          "bindings": "minInputVector",
-          "dispatch": { "threads": "numel(shapes.output) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
@@ -371,17 +144,25 @@
       "id": "max_input_vec4",
       "priority": 40,
       "when": ["maxInputContract", "numel(shapes.input) > 0", "numel(shapes.input) % 4 == 0", "wideVec4StorageOk"],
-      "constants": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "Clip.maxInputVec4",
-          "source": {
-            "shader": "unary-vec4.wgsl.jinja",
-            "inputs": { "op": "\"clip\"", "boundsFromInput": true, "clipMinFromInput": false, "clipMaxFromInput": true }
           },
-          "bindings": "maxInputVector",
-          "dispatch": { "threads": "numel(shapes.output) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
@@ -393,12 +174,20 @@
         {
           "id": "main",
           "name": "Clip.inputBounds",
-          "source": {
-            "shader": "unary-scalar.wgsl.jinja",
-            "inputs": { "op": "\"clip\"", "boundsFromInput": true, "clipMinFromInput": true, "clipMaxFromInput": true }
           },
-          "bindings": "inputBounds",
-          "dispatch": { "threads": "numel(shapes.output)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
@@ -410,12 +199,20 @@
         {
           "id": "main",
           "name": "Clip.minInput",
-          "source": {
-            "shader": "unary-scalar.wgsl.jinja",
-            "inputs": { "op": "\"clip\"", "boundsFromInput": true, "clipMinFromInput": true, "clipMaxFromInput": false }
           },
-          "bindings": "minInput",
-          "dispatch": { "threads": "numel(shapes.output)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
@@ -427,12 +224,20 @@
         {
           "id": "main",
           "name": "Clip.maxInput",
-          "source": {
-            "shader": "unary-scalar.wgsl.jinja",
-            "inputs": { "op": "\"clip\"", "boundsFromInput": true, "clipMinFromInput": false, "clipMaxFromInput": true }
           },
-          "bindings": "maxInput",
-          "dispatch": { "threads": "numel(shapes.output)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     }

   "domain": "ai.onnx",
   "name": "Clip",
   "sinceVersion": 13,
+  "inputs": {
+    "input": { "dtype": "T" },
+    "min": { "dtype": "T", "rank": 0, "optional": true, "shape": [] },
+    "max": { "dtype": "T", "rank": 0, "optional": true, "shape": [] }
   },
+  "outputs": { "output": { "dtype": "T", "rank": "ranks.input", "shape": "shapes.input" } },
+  "typeConstraints": { "T": ["float32", "float16", "int32", "uint32", "int8", "uint8"] },
+  "tunables": { "WORKGROUP_SIZE": { "default": 256 } },
   "derive": {
     "shapeContract": "numel(shapes.input) == numel(shapes.output) and f16Ok(dtypes.T)",
     "defaultBoundsContract": "shapeContract and not present.min and not present.max",
     "inputBoundsContract": "shapeContract and present.min and present.max and ranks.min == 0 and numel(shapes.min) == 1 and ranks.max == 0 and numel(shapes.max) == 1",
     "defaultMin": "-65504 if tensorDtypes.input == \"float16\" else (-340282346638528859811704183484516925440 if tensorDtypes.input == \"float32\" else (-128 if tensorDtypes.input == \"int8\" else (-2147483648 if tensorDtypes.input == \"int32\" else 0)))",
     "defaultMax": "65504 if tensorDtypes.input == \"float16\" else (340282346638528859811704183484516925440 if tensorDtypes.input == \"float32\" else (127 if tensorDtypes.input == \"int8\" else (2147483647 if tensorDtypes.input == \"int32\" else (255 if tensorDtypes.input == \"uint8\" else 4294967295))))",
+    "wideVec4StorageOk": "dtypes.T != \"f32\" or device.features.has(\"subgroups\") or not (has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize == 8 and device.adapterInfo.subgroupMaxSize <= 32)",
+    "scalar": "dtypes.T",
+    "usesF16": "dtypes.T == \"f16\""
   },
+  "bindings": {
+    "x": { "arg": "input", "buffer": "read-only-storage", "elementType": "$vectorScalar" },
+    "y": { "arg": "output", "buffer": "storage", "elementType": "$vectorScalar" },
+    "params": {
+      "buffer": "uniform",
+      "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.output) / 4" }]
+    },
+    "params_2": {
+      "name": "params",
+      "buffer": "uniform",
+      "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.output)" }]
+    },
+    "clipMin": { "arg": "min", "buffer": "read-only-storage", "elementType": "$scalar" },
+    "clipMax": { "arg": "max", "buffer": "read-only-storage", "elementType": "$scalar" }
   },
   "variants": [
     {
       "id": "default_bounds_vec4",
       "priority": 20,
       "when": ["defaultBoundsContract", "numel(shapes.input) > 0", "numel(shapes.input) % 4 == 0", "wideVec4StorageOk"],
+      "derive": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "Clip.defaultBoundsVec4",
+          "shader": "unary-vec4.wgsl.jinja",
+          "derive": {
+            "op": "\"clip\"",
+            "boundsFromInput": false,
+            "minValue": "defaultMin",
+            "maxValue": "defaultMax",
+            "vec4PerThread": "4 if numel(shapes.output) * dtypeBytes(tensorDtypes.output) <= 16777216 else 1"
           },
+          "bindings": ["x", "y", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.output) / 4, 4 if numel(shapes.output) * dtypeBytes(tensorDtypes.output) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.output) / 4, 4 if numel(shapes.output) * dtypeBytes(tensorDtypes.output) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
         {
           "id": "main",
           "name": "Clip.defaultBounds",
+          "shader": "unary-scalar.wgsl.jinja",
+          "derive": {
+            "op": "\"clip\"",
+            "boundsFromInput": false,
+            "minValue": "defaultMin",
+            "maxValue": "defaultMax",
+            "itemsPerInvocation": 4
           },
+          "bindings": ["input", "output", "params_2"],
+          "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": "input_bounds_vec4",
       "priority": 40,
       "when": ["inputBoundsContract", "numel(shapes.input) > 0", "numel(shapes.input) % 4 == 0", "wideVec4StorageOk"],
+      "derive": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "Clip.inputBoundsVec4",
+          "shader": "unary-vec4.wgsl.jinja",
+          "derive": {
+            "op": "\"clip\"",
+            "boundsFromInput": true,
+            "clipMinFromInput": true,
+            "clipMaxFromInput": true,
+            "vec4PerThread": "4 if numel(shapes.output) * dtypeBytes(tensorDtypes.output) <= 16777216 else 1"
           },
+          "bindings": ["x", "y", "clipMin", "clipMax", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.output) / 4, 4 if numel(shapes.output) * dtypeBytes(tensorDtypes.output) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.output) / 4, 4 if numel(shapes.output) * dtypeBytes(tensorDtypes.output) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
       "id": "min_input_vec4",
       "priority": 40,
       "when": ["minInputContract", "numel(shapes.input) > 0", "numel(shapes.input) % 4 == 0", "wideVec4StorageOk"],
+      "derive": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "Clip.minInputVec4",
+          "shader": "unary-vec4.wgsl.jinja",
+          "derive": {
+            "op": "\"clip\"",
+            "boundsFromInput": true,
+            "clipMinFromInput": true,
+            "clipMaxFromInput": false,
+            "vec4PerThread": "4 if numel(shapes.output) * dtypeBytes(tensorDtypes.output) <= 16777216 else 1"
           },
+          "bindings": ["x", "y", "clipMin", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.output) / 4, 4 if numel(shapes.output) * dtypeBytes(tensorDtypes.output) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.output) / 4, 4 if numel(shapes.output) * dtypeBytes(tensorDtypes.output) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
       "id": "max_input_vec4",
       "priority": 40,
       "when": ["maxInputContract", "numel(shapes.input) > 0", "numel(shapes.input) % 4 == 0", "wideVec4StorageOk"],
+      "derive": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "Clip.maxInputVec4",
+          "shader": "unary-vec4.wgsl.jinja",
+          "derive": {
+            "op": "\"clip\"",
+            "boundsFromInput": true,
+            "clipMinFromInput": false,
+            "clipMaxFromInput": true,
+            "vec4PerThread": "4 if numel(shapes.output) * dtypeBytes(tensorDtypes.output) <= 16777216 else 1"
           },
+          "bindings": ["x", "y", "clipMax", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.output) / 4, 4 if numel(shapes.output) * dtypeBytes(tensorDtypes.output) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.output) / 4, 4 if numel(shapes.output) * dtypeBytes(tensorDtypes.output) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
         {
           "id": "main",
           "name": "Clip.inputBounds",
+          "shader": "unary-scalar.wgsl.jinja",
+          "derive": {
+            "op": "\"clip\"",
+            "boundsFromInput": true,
+            "clipMinFromInput": true,
+            "clipMaxFromInput": true,
+            "itemsPerInvocation": 4
           },
+          "bindings": ["input", "output", "clipMin", "clipMax", "params_2"],
+          "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": "main",
           "name": "Clip.minInput",
+          "shader": "unary-scalar.wgsl.jinja",
+          "derive": {
+            "op": "\"clip\"",
+            "boundsFromInput": true,
+            "clipMinFromInput": true,
+            "clipMaxFromInput": false,
+            "itemsPerInvocation": 4
           },
+          "bindings": ["input", "output", "clipMin", "params_2"],
+          "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": "main",
           "name": "Clip.maxInput",
+          "shader": "unary-scalar.wgsl.jinja",
+          "derive": {
+            "op": "\"clip\"",
+            "boundsFromInput": true,
+            "clipMinFromInput": false,
+            "clipMaxFromInput": true,
+            "itemsPerInvocation": 4
           },
+          "bindings": ["input", "output", "clipMax", "params_2"],
+          "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
+          }
         }
       ]
     }

build/webgpu/metadata.json CHANGED Viewed

@@ -1,19 +1,31 @@
 {
   "name": "ai.onnx.Clip",
-  "id": "_ai_onnx_clip_webgpu_ea7567b",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "SiGZqqpALf6ubJItLzIkYagu90oP+HzScN4MdEyZgc8=",
-      "manifest.json": "us9b3p/dfzQqD40BHqWMZiVp6xaoru4kvNbEXWhfevM=",
-      "test.json": "ymGZ3BiNHy1Nye9xpVQCWrfPh5sbCvMpsaE5E8lWq/c=",
-      "unary-scalar.wgsl.jinja": "NeWScwhOG4N4iqKUuLMB0vL0DbMQqgUCFSp+g+YoKCY=",
-      "unary-vec4.wgsl.jinja": "QKAUBbXPhcNyTR7St0KyqUqJjuS++OgbomsrxdMCrTw="
     }
   },
-  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
-  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Clip" }
 }

 {
   "name": "ai.onnx.Clip",
+  "id": "_ai_onnx_clip_webgpu_de72431",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "F12x3+s84ktuo167e5515+F2hnUMtcowAYHmEgoDQOg=",
+      "manifest.json": "kwVE+OfpodbDYkXNlYpHzflNPmAefXobA+ugNpwKfME=",
+      "test.json": "a604ZK+4B6A13Q4yccOjW3dVe84WzyJ57YjvMttgjuQ=",
+      "unary-scalar.wgsl.jinja": "rsQQXT2OWXvamkOuvViLDBj3K2kZY/CpmMOvsDNXU7Y=",
+      "unary-vec4.wgsl.jinja": "/UXGg3mvbuFOUdyjrpM1npD9tJ5fovKFZ6dt0LFgRRc="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": {
+      "default_bounds_vec4": ["unary-vec4.wgsl.jinja"],
+      "default_bounds": ["unary-scalar.wgsl.jinja"],
+      "input_bounds_vec4": ["unary-vec4.wgsl.jinja"],
+      "min_input_vec4": ["unary-vec4.wgsl.jinja"],
+      "max_input_vec4": ["unary-vec4.wgsl.jinja"],
+      "input_bounds": ["unary-scalar.wgsl.jinja"],
+      "min_input": ["unary-scalar.wgsl.jinja"],
+      "max_input": ["unary-scalar.wgsl.jinja"]
+    }
+  }
 }

build/webgpu/test.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "ai.onnx.Clip",
   "fixtureArrays": {
     "onnx_backend_clip_input": [1.764052391052246, 0.40015721321105957, 0.978738009929657, 2.2408931255340576, 1.8675580024719238, -0.9772778749465942, 0.9500884413719177, -0.15135720372200012, -0.10321885347366333, 0.4105985164642334, 0.14404356479644775, 1.4542734622955322, 0.7610377073287964, 0.12167501449584961, 0.44386324286460876, 0.3336743414402008, 1.4940791130065918, -0.2051582634449005, 0.3130677044391632, -0.8540957570075989, -2.5529897212982178, 0.653618574142456, 0.8644362092018127, -0.7421650290489197, 2.269754648208618, -1.4543657302856445, 0.04575851559638977, -0.18718385696411133, 1.5327792167663574, 1.4693588018417358, 0.154947429895401, 0.37816253304481506, -0.8877857327461243, -1.980796456336975, -0.34791216254234314, 0.15634897351264954, 1.2302906513214111, 1.202379822731018, -0.38732680678367615, -0.302302747964859, -1.0485529899597168, -1.420017957687378, -1.7062702178955078, 1.950775384902954, -0.5096521973609924, -0.4380742907524109, -1.2527953386306763, 0.7774903774261475, -1.6138978004455566, -0.21274028718471527, -0.8954665660858154, 0.38690251111984253, -0.5108051300048828, -1.18063223361969, -0.02818222902715206, 0.4283318817615509, 0.06651721894741058, 0.30247190594673157, -0.6343221068382263, -0.3627411723136902]
   },
@@ -35,7 +34,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/math/clip_test.cc",
         "test": "MathOpTest.Clip_Relu",
-        "notes": "Scalar-path companion: Clip with min=0 should pass through positive subnormal inputs instead of flushing them to zero."
       },
       "inputs": {
         "input": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-1e-40, 0.0, 1e-40] } },
@@ -370,16 +369,6 @@
       },
       "outputs": { "output": { "dtype": "float32", "shape": [4], "tolerance": 0.000001, "allowNaN": true } }
     },
-    {
-      "name": "backend_min_greater_than_max",
-      "inputs": {
-        "input": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-2.0, 0.0, 6.0] } },
-        "min": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } },
-        "max": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] } }
-      },
-      "outputs": { "output": { "dtype": "float32", "shape": [3], "tolerance": 0 } },
-      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_clip_min_greater_than_max" }
-    },
     {
       "name": "default_max_only",
       "inputs": {
@@ -844,14 +833,14 @@
     {
       "name": "onnx_backend_clip_min_greater_than_max",
       "inputs": {
-        "input": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-2.0, 0.0, 6.0] } },
         "min": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } },
         "max": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] } }
       },
       "outputs": { "output": { "dtype": "float32", "shape": [3] } },
       "provenance": {
         "source": "cmake/external/onnx/onnx/backend/test/data/node/test_clip_min_greater_than_max",
-        "notes": "ONNX Clip min/max bounds are supplied as scalar tensor inputs."
       }
     },
     {
@@ -931,11 +920,7 @@
     {
       "name": "input_bounds_i32_min_greater_than_max",
       "inputs": {
-        "input": {
-          "dtype": "int32",
-          "shape": [6],
-          "data": { "kind": "values", "values": [-2147483648, -10, 0, 5, 10, 2147483647] }
-        },
         "min": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [5] } },
         "max": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [-5] } }
       },
@@ -946,6 +931,9 @@
           "data": { "kind": "values", "values": [-5, -5, -5, -5, -5, -5] },
           "tolerance": 0
         }
       }
     },
     {
@@ -968,7 +956,7 @@
     {
       "name": "uint8_input_bounds_min_greater_than_max",
       "inputs": {
-        "input": { "dtype": "uint8", "shape": [5], "data": { "kind": "values", "values": [0, 50, 100, 150, 255] } },
         "min": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [200] } },
         "max": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [100] } }
       },
@@ -979,6 +967,9 @@
           "data": { "kind": "values", "values": [100, 100, 100, 100, 100] },
           "tolerance": 0
         }
       }
     }
   ]

 {
   "fixtureArrays": {
     "onnx_backend_clip_input": [1.764052391052246, 0.40015721321105957, 0.978738009929657, 2.2408931255340576, 1.8675580024719238, -0.9772778749465942, 0.9500884413719177, -0.15135720372200012, -0.10321885347366333, 0.4105985164642334, 0.14404356479644775, 1.4542734622955322, 0.7610377073287964, 0.12167501449584961, 0.44386324286460876, 0.3336743414402008, 1.4940791130065918, -0.2051582634449005, 0.3130677044391632, -0.8540957570075989, -2.5529897212982178, 0.653618574142456, 0.8644362092018127, -0.7421650290489197, 2.269754648208618, -1.4543657302856445, 0.04575851559638977, -0.18718385696411133, 1.5327792167663574, 1.4693588018417358, 0.154947429895401, 0.37816253304481506, -0.8877857327461243, -1.980796456336975, -0.34791216254234314, 0.15634897351264954, 1.2302906513214111, 1.202379822731018, -0.38732680678367615, -0.302302747964859, -1.0485529899597168, -1.420017957687378, -1.7062702178955078, 1.950775384902954, -0.5096521973609924, -0.4380742907524109, -1.2527953386306763, 0.7774903774261475, -1.6138978004455566, -0.21274028718471527, -0.8954665660858154, 0.38690251111984253, -0.5108051300048828, -1.18063223361969, -0.02818222902715206, 0.4283318817615509, 0.06651721894741058, 0.30247190594673157, -0.6343221068382263, -0.3627411723136902]
   },
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/math/clip_test.cc",
         "test": "MathOpTest.Clip_Relu",
+        "notes": "On an unaligned scalar path, Clip with min=0 must preserve positive subnormal inputs."
       },
       "inputs": {
         "input": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-1e-40, 0.0, 1e-40] } },
       },
       "outputs": { "output": { "dtype": "float32", "shape": [4], "tolerance": 0.000001, "allowNaN": true } }
     },
     {
       "name": "default_max_only",
       "inputs": {
     {
       "name": "onnx_backend_clip_min_greater_than_max",
       "inputs": {
+        "input": { "dtype": "float32", "shape": [3], "data": { "kind": "constant", "value": -2.0 } },
         "min": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } },
         "max": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] } }
       },
       "outputs": { "output": { "dtype": "float32", "shape": [3] } },
       "provenance": {
         "source": "cmake/external/onnx/onnx/backend/test/data/node/test_clip_min_greater_than_max",
+        "notes": "Diverges from the upstream test's inputs (inputs.input values [-2.0, 0.0, 6.0] -> constant -2.0); the expected output is recomputed by the CPU reference for the new inputs. ONNX Clip min/max bounds are supplied as scalar tensor inputs. With min above max every element becomes max whatever the input holds, so the input is uniform at a value below both bounds: min(x,max) would return -2 and max(x,min) would return 2."
       }
     },
     {
     {
       "name": "input_bounds_i32_min_greater_than_max",
       "inputs": {
+        "input": { "dtype": "int32", "shape": [6], "data": { "kind": "constant", "value": -2147483648 } },
         "min": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [5] } },
         "max": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [-5] } }
       },
           "data": { "kind": "values", "values": [-5, -5, -5, -5, -5, -5] },
           "tolerance": 0
         }
+      },
+      "provenance": {
+        "notes": "With min above max every element becomes max whatever the input holds, so the input is uniform at int32 min: clamping in the wrong order returns -2147483648 or 5 instead of -5."
       }
     },
     {
     {
       "name": "uint8_input_bounds_min_greater_than_max",
       "inputs": {
+        "input": { "dtype": "uint8", "shape": [5], "data": { "kind": "constant", "value": 0 } },
         "min": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [200] } },
         "max": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [100] } }
       },
           "data": { "kind": "values", "values": [100, 100, 100, 100, 100] },
           "tolerance": 0
         }
+      },
+      "provenance": {
+        "notes": "With min above max every element becomes max whatever the input holds, so the input is uniform at 0: clamping in the wrong order returns 0 or 200 instead of 100."
       }
     }
   ]

build/webgpu/unary-scalar.wgsl.jinja CHANGED Viewed

@@ -1,18 +1,23 @@
 {% 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;
-  let i = invocation;
-  if (i >= params.count) {
-    return;
-  }{% endmacro %}
 {% macro flat_tail_close() %}
 {% endmacro %}
-// Scalar unary fallback. Each branch retains the operation's numeric hardening,
-// including Payne-Hanek trigonometric range reduction and NaN/overflow guards.
 {% if usesF16 %}
 enable f16;
 {% endif %}
@@ -22,18 +27,18 @@ fn is_nan_f32(value: f32) -> bool {
   let bits = bitcast<u32>(value);
   return (bits & 0x7f800000u) == 0x7f800000u && (bits & 0x007fffffu) != 0u;
 }{% endmacro %}
-{% if source.op == "clip" and scalar != "i32" and scalar != "u32" %}
 {{ emit_is_nan_f32() }}
 {% endif %}
 {{ flat_tail_open() }}
-{% set clipMinFromInput = source.clipMinFromInput | default(false) %}
-{% set clipMaxFromInput = source.clipMaxFromInput | default(false) %}
-{% set clipMinValue = source.minValue | default(0) %}
-{% set clipMaxValue = source.maxValue | default(0) %}
 {% set clipScalar = scalar | default("f32") %}
 {% if scalar == "i32" or scalar == "u32" %}
   let value = input[i];
-{% if source.boundsFromInput %}
 {% if clipMinFromInput and clipMaxFromInput %}
   output[i] = min(max(value, clipMin[0]), clipMax[0]);
 {% elif clipMinFromInput %}
@@ -46,7 +51,7 @@ fn is_nan_f32(value: f32) -> bool {
 {% endif %}
 {% else %}
   let value = f32(input[i]);
-{% if source.boundsFromInput %}
 {% if clipMinFromInput and clipMaxFromInput %}
   var out = min(max(value, f32(clipMin[0])), f32(clipMax[0]));
 {% elif clipMinFromInput %}

 {% 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 boundsFromInput = boundsFromInput is defined and boundsFromInput %}
+// Scalar unary elementwise implementation. Specialization emits only the
+// selected operation and any numerical helper it requires.
 {% if usesF16 %}
 enable f16;
 {% endif %}
   let bits = bitcast<u32>(value);
   return (bits & 0x7f800000u) == 0x7f800000u && (bits & 0x007fffffu) != 0u;
 }{% endmacro %}
+{% if op == "clip" and scalar != "i32" and scalar != "u32" %}
 {{ emit_is_nan_f32() }}
 {% endif %}
 {{ flat_tail_open() }}
+{% set clipMinFromInput = clipMinFromInput | default(false) %}
+{% set clipMaxFromInput = clipMaxFromInput | default(false) %}
+{% set clipMinValue = minValue | default(0) %}
+{% set clipMaxValue = maxValue | default(0) %}
 {% set clipScalar = scalar | default("f32") %}
 {% if scalar == "i32" or scalar == "u32" %}
   let value = input[i];
+{% if boundsFromInput %}
 {% if clipMinFromInput and clipMaxFromInput %}
   output[i] = min(max(value, clipMin[0]), clipMax[0]);
 {% elif clipMinFromInput %}
 {% endif %}
 {% else %}
   let value = f32(input[i]);
+{% if boundsFromInput %}
 {% if clipMinFromInput and clipMaxFromInput %}
   var out = min(max(value, f32(clipMin[0])), f32(clipMax[0]));
 {% elif clipMinFromInput %}

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

@@ -1,8 +1,5 @@
-// Loads and stores vec4<T> (128 bits) while retaining scalar per-component
-// arithmetic, including per-component helper calls for guard-heavy operations.
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 {% macro emit_is_nan_f32() %}
@@ -10,12 +7,12 @@ fn is_nan_f32(value: f32) -> bool {
   let bits = bitcast<u32>(value);
   return (bits & 0x7f800000u) == 0x7f800000u && (bits & 0x007fffffu) != 0u;
 }{% endmacro %}
-{% if (source.op == "clip" and scalar != "i32" and scalar != "u32") or (source.op == "relu" and scalar != "i32") or source.op == "sin" or source.op == "tan" %}
 {{ emit_is_nan_f32() }}
 {% endif %}
-{% if not (source.boundsFromInput | default(false)) and scalar != "i32" and scalar != "u32" %}
-{% set clipMinValue = source.minValue | default(0) %}
-{% set clipMaxValue = source.maxValue | default(0) %}
 fn clip_value(v: f32) -> f32 {
   var out = min(max(v, f32({{ clipMinValue }})), f32({{ clipMaxValue }}));
   if (is_nan_f32(v)) {
@@ -25,20 +22,39 @@ fn clip_value(v: f32) -> f32 {
 }
 {% endif %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
-fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
   // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
-  // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills the rest into y).
-  let i = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
   if (i >= params.count) {
     return;
   }
   let xv = x[i];
-{% set clipMinValue = source.minValue | default(0) %}
-{% set clipMaxValue = source.maxValue | default(0) %}
-{% if source.boundsFromInput | default(false) %}
-{% set clipMinFromInput = source.clipMinFromInput | default(false) %}
-{% set clipMaxFromInput = source.clipMaxFromInput | default(false) %}
 {% if scalar == "i32" or scalar == "u32" %}
 {% if clipMinFromInput and clipMaxFromInput %}
   y[i] = min(max(xv, vec4<{{ scalar }}>(clipMin[0])), vec4<{{ scalar }}>(clipMax[0]));
@@ -70,4 +86,7 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
   let fv = vec4<f32>(xv);
   y[i] = {{ vectorScalar }}(vec4<f32>(clip_value(fv.x), clip_value(fv.y), clip_value(fv.z), clip_value(fv.w)));
 {% endif %}
 }

+// Loads and stores vec4<T> while evaluating the selected unary operation per
+// component.
 {{ env.wgsl.resourceDeclarations }}
 {% macro emit_is_nan_f32() %}
   let bits = bitcast<u32>(value);
   return (bits & 0x7f800000u) == 0x7f800000u && (bits & 0x007fffffu) != 0u;
 }{% endmacro %}
+{% if (op == "clip" and scalar != "i32" and scalar != "u32") or (op == "relu" and scalar != "i32") or op == "sin" or op == "tan" %}
 {{ emit_is_nan_f32() }}
 {% endif %}
+{% if not (boundsFromInput | default(false)) and scalar != "i32" and scalar != "u32" %}
+{% set clipMinValue = minValue | default(0) %}
+{% set clipMaxValue = maxValue | default(0) %}
 fn clip_value(v: f32) -> f32 {
   var out = min(max(v, f32({{ clipMinValue }})), f32({{ clipMaxValue }}));
   if (is_nan_f32(v)) {
 }
 {% endif %}
+{% set vec4PerThread = vec4PerThread %}
+{% if vec4PerThread > 1 %}
+const ITEMS: u32 = {{ vec4PerThread }}u;
+{% endif %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
+  // per-axis dispatch fold width (the dispatch caps x and spills the rest into y).
+{% if vec4PerThread > 1 %}
+  // Each invocation walks ITEMS vec4 groups a span apart. Consecutive lanes
+  // access consecutive words on every step, while each lane can keep several
+  // independent loads in flight.
+  let tid = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
+  let span = (params.count + ITEMS - 1u) / ITEMS;
+  for (var j = 0u; j < ITEMS; j = j + 1u) {
+    let i = tid + j * span;
+    if (i >= params.count) {
+      break;
+    }
+{% else %}
+  let i = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if (i >= params.count) {
     return;
   }
+{% endif %}
   let xv = x[i];
+{% set clipMinValue = minValue | default(0) %}
+{% set clipMaxValue = maxValue | default(0) %}
+{% if boundsFromInput | default(false) %}
+{% set clipMinFromInput = clipMinFromInput | default(false) %}
+{% set clipMaxFromInput = clipMaxFromInput | default(false) %}
 {% if scalar == "i32" or scalar == "u32" %}
 {% if clipMinFromInput and clipMaxFromInput %}
   y[i] = min(max(xv, vec4<{{ scalar }}>(clipMin[0])), vec4<{{ scalar }}>(clipMax[0]));
   let fv = vec4<f32>(xv);
   y[i] = {{ vectorScalar }}(vec4<f32>(clip_value(fv.x), clip_value(fv.y), clip_value(fv.z), clip_value(fv.w)));
 {% endif %}
+{% if vec4PerThread > 1 %}
+  }
+{% endif %}
 }