sync 2e7068faf55e

Browse files

Files changed (7) hide show

README.md +59 -0
build/webgpu/bench.json +67 -0
build/webgpu/manifest.json +440 -0
build/webgpu/metadata.json +19 -0
build/webgpu/test.json +985 -0
build/webgpu/unary-scalar.wgsl.jinja +66 -0
build/webgpu/unary-vec4.wgsl.jinja +73 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.Clip
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 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`.
+See the [ONNX `Clip` spec](https://onnx.ai/onnx/operators/onnx__Clip.html) for the reference semantics.
+## 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
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16`, `int32`, `uint32`, `int8`, `uint8` |
+## 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
+- [`unary-scalar.wgsl.jinja`](build/webgpu/unary-scalar.wgsl.jinja)
+- [`unary-vec4.wgsl.jinja`](build/webgpu/unary-vec4.wgsl.jinja)
+## Use with `@huggingface/kernels`
+The loader derives every required output's shape and logical dtype from the manifest contract and this call.
+It then allocates the result tensors automatically.
+The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+Replace each `*Data` placeholder with a typed array containing the corresponding input data.
+```js
+import { getKernel } from "@huggingface/kernels";
+const kernel = await getKernel("webgpu-kernels/ai.onnx.Clip", { version: 1 });
+const { output } = await kernel({ input: { data: inputData, shape: [3] } });
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,67 @@

+{
+  "op": "ai.onnx.Clip",
+  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
+  "cases": [
+    {
+      "name": "1m_f32",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [1048576] },
+        "min": { "dtype": "float32", "shape": [], "dist": "constant", "value": -1 },
+        "max": { "dtype": "float32", "shape": [], "dist": "constant", "value": 1 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1048576] } }
+    },
+    {
+      "name": "clip-f32-8m",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "count": 8388608 },
+      "inputs": {
+        "input": { "shape": [8388608], "dtype": "float32", "seed": 7025, "dist": "normal", "scale": 2 },
+        "min": { "dtype": "float32", "shape": [], "dist": "constant", "value": -1 },
+        "max": { "dtype": "float32", "shape": [], "dist": "constant", "value": 1 }
+      },
+      "outputs": { "output": { "shape": [8388608], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "clip-f16-8m",
+      "preset": "smoke",
+      "vars": { "dtype": "float16", "count": 8388608 },
+      "inputs": {
+        "input": { "shape": [8388608], "dtype": "float16", "seed": 7026, "dist": "normal", "scale": 2 },
+        "min": { "dtype": "float16", "shape": [], "dist": "constant", "value": -1 },
+        "max": { "dtype": "float16", "shape": [], "dist": "constant", "value": 1 }
+      },
+      "outputs": { "output": { "shape": [8388608], "dtype": "float16" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "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 },
+      "inputs": {
+        "input": { "shape": [8388608], "dtype": "float32", "dist": "normal", "seed": 7101, "scale": 2 },
+        "min": { "shape": [], "dtype": "float32", "dist": "constant", "value": -1 },
+        "max": { "shape": [], "dtype": "float32", "dist": "constant", "value": 1 }
+      },
+      "outputs": { "output": { "shape": [8388608], "dtype": "float32", "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "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 },
+      "inputs": { "input": { "shape": [8388608], "dtype": "float16", "seed": 7026, "dist": "normal", "scale": 2 } },
+      "outputs": { "output": { "shape": [8388608], "dtype": "float16" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
+    }
+  ]
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,440 @@

+{
+  "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",
+    "minInputContract": "shapeContract and present.min and not present.max and ranks.min == 0 and numel(shapes.min) == 1",
+    "maxInputContract": "shapeContract and not present.min and present.max and ranks.max == 0 and numel(shapes.max) == 1",
+    "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" }
+        }
+      ]
+    },
+    {
+      "id": "default_bounds",
+      "when": ["defaultBoundsContract"],
+      "passes": [
+        {
+          "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" }
+        }
+      ]
+    },
+    {
+      "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" }
+        }
+      ]
+    },
+    {
+      "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" }
+        }
+      ]
+    },
+    {
+      "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" }
+        }
+      ]
+    },
+    {
+      "id": "input_bounds",
+      "priority": 30,
+      "when": ["inputBoundsContract"],
+      "passes": [
+        {
+          "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" }
+        }
+      ]
+    },
+    {
+      "id": "min_input",
+      "priority": 30,
+      "when": ["minInputContract"],
+      "passes": [
+        {
+          "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" }
+        }
+      ]
+    },
+    {
+      "id": "max_input",
+      "priority": 30,
+      "when": ["maxInputContract"],
+      "passes": [
+        {
+          "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" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,19 @@

+{
+  "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" }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,985 @@

+{
+  "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]
+  },
+  "cases": [
+    {
+      "name": "f32_relu_style_positive_subnormal_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: WebGPU/Metal flushes subnormals to zero (f32 and f16); the kernel cannot preserve denormal inputs/outputs bit-exactly."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/clip_test.cc",
+        "test": "MathOpTest.Clip_Relu",
+        "notes": "Clip with min=0 is ReLU-like: positive subnormal inputs should pass through instead of flushing to zero."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [-1e-40, 0.0, 1e-40, 1e-39] }
+        },
+        "min": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } },
+        "max": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4], "tolerance": 0 } }
+    },
+    {
+      "name": "f32_relu_style_positive_subnormal_scalar_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: WebGPU/Metal flushes subnormals to zero (f32 and f16); the kernel cannot preserve denormal inputs/outputs bit-exactly."
+      },
+      "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] } },
+        "min": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } },
+        "max": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3], "tolerance": 0 } }
+    },
+    {
+      "name": "input_bounds_f32_small",
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [-3.0, -1.0, 0.5, 1.5, 3.0, 9.0] }
+        },
+        "min": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-1.5] } },
+        "max": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [6] } }
+    },
+    {
+      "name": "input_bounds_f32_1024",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [1024], "data": { "kind": "linspace", "start": -4.0, "end": 4.0 } },
+        "min": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-1.5] } },
+        "max": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1024], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "input_bounds_f32_nonfinite",
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": ["-Infinity", -1.0, 1.0, "Infinity"] }
+        },
+        "min": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": ["-Infinity"] } },
+        "max": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": ["Infinity"] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": ["-Infinity", -1.0, 1.0, "Infinity"] }
+        }
+      }
+    },
+    {
+      "name": "f16_relu6_style",
+      "inputs": {
+        "input": { "dtype": "float16", "shape": [4], "data": { "kind": "values", "values": [-1.0, 0.5, 7.0, 3.0] } },
+        "min": { "dtype": "float16", "shape": [], "data": { "kind": "values", "values": [0.0] } },
+        "max": { "dtype": "float16", "shape": [], "data": { "kind": "values", "values": [6.0] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [4] } },
+      "tolerance": 0.001
+    },
+    {
+      "name": "input_bounds_float32",
+      "provenance": {
+        "source": "onnx/onnx/backend/test/case/node/clip.py",
+        "notes": "Clip-13 expresses min/max as optional scalar inputs; this fixture exercises both tensor-bound inputs."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": { "kind": "values", "values": [-3.0, -1.0, 0.5, 2.0, 9.0] }
+        },
+        "min": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-1.0] } },
+        "max": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [5], "tolerance": 0 } }
+    },
+    {
+      "name": "input_bounds_float16",
+      "inputs": {
+        "input": {
+          "dtype": "float16",
+          "shape": [5],
+          "data": { "kind": "values", "values": [-4.0, -1.0, 0.5, 2.0, 8.0] }
+        },
+        "min": { "dtype": "float16", "shape": [], "data": { "kind": "values", "values": [-1.0] } },
+        "max": { "dtype": "float16", "shape": [], "data": { "kind": "values", "values": [2.0] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float16",
+          "shape": [5],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [-1.0, -1.0, 0.5, 2.0, 2.0] }
+        }
+      }
+    },
+    {
+      "name": "input_bounds_int8",
+      "inputs": {
+        "input": { "dtype": "int8", "shape": [6], "data": { "kind": "values", "values": [-128, -12, -3, 4, 17, 127] } },
+        "min": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [-3] } },
+        "max": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [17] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "int8",
+          "shape": [6],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [-3, -3, -3, 4, 17, 17] }
+        }
+      }
+    },
+    {
+      "name": "input_bounds_uint8",
+      "inputs": {
+        "input": { "dtype": "uint8", "shape": [6], "data": { "kind": "values", "values": [0, 5, 17, 18, 199, 255] } },
+        "min": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [17] } },
+        "max": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [199] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "uint8",
+          "shape": [6],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [17, 17, 17, 18, 199, 199] }
+        }
+      }
+    },
+    {
+      "name": "input_min_only_float16",
+      "inputs": {
+        "input": { "dtype": "float16", "shape": [4], "data": { "kind": "values", "values": [-4.0, -1.0, 0.5, 8.0] } },
+        "min": { "dtype": "float16", "shape": [], "data": { "kind": "values", "values": [-1.0] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float16",
+          "shape": [4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [-1.0, -1.0, 0.5, 8.0] }
+        }
+      }
+    },
+    {
+      "name": "input_max_only_float16",
+      "inputs": {
+        "input": { "dtype": "float16", "shape": [4], "data": { "kind": "values", "values": [-4.0, -1.0, 0.5, 8.0] } },
+        "max": { "dtype": "float16", "shape": [], "data": { "kind": "values", "values": [2.0] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float16",
+          "shape": [4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [-4.0, -1.0, 0.5, 2.0] }
+        }
+      }
+    },
+    {
+      "name": "input_min_only_int8",
+      "inputs": {
+        "input": { "dtype": "int8", "shape": [4], "data": { "kind": "values", "values": [-128, -3, 4, 127] } },
+        "min": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [-3] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "int8",
+          "shape": [4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [-3, -3, 4, 127] }
+        }
+      }
+    },
+    {
+      "name": "input_max_only_int8",
+      "inputs": {
+        "input": { "dtype": "int8", "shape": [4], "data": { "kind": "values", "values": [-128, -3, 17, 127] } },
+        "max": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [17] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "int8",
+          "shape": [4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [-128, -3, 17, 17] }
+        }
+      }
+    },
+    {
+      "name": "input_min_only_uint8",
+      "inputs": {
+        "input": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [0, 17, 18, 255] } },
+        "min": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [17] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "uint8",
+          "shape": [4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [17, 17, 18, 255] }
+        }
+      }
+    },
+    {
+      "name": "input_max_only_uint8",
+      "inputs": {
+        "input": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [0, 17, 199, 255] } },
+        "max": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [199] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "uint8",
+          "shape": [4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0, 17, 199, 199] }
+        }
+      }
+    },
+    {
+      "name": "input_min_only_float32",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [-4.0, -1.0, 0.5, 8.0] } },
+        "min": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-1.0] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [-1.0, -1.0, 0.5, 8.0] }
+        }
+      }
+    },
+    {
+      "name": "input_max_only_float32",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [-4.0, -1.0, 0.5, 8.0] } },
+        "max": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [-4.0, -1.0, 0.5, 2.0] }
+        }
+      }
+    },
+    {
+      "name": "input_min_only_int32",
+      "inputs": {
+        "input": {
+          "dtype": "int32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [-2147483648, -16777217, 0, 2147483647] }
+        },
+        "min": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [-16777217] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "int32",
+          "shape": [4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [-16777217, -16777217, 0, 2147483647] }
+        }
+      }
+    },
+    {
+      "name": "input_max_only_int32",
+      "inputs": {
+        "input": {
+          "dtype": "int32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [-2147483648, 0, 16777217, 2147483647] }
+        },
+        "max": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [16777217] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "int32",
+          "shape": [4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [-2147483648, 0, 16777217, 16777217] }
+        }
+      }
+    },
+    {
+      "name": "input_min_only_uint32",
+      "inputs": {
+        "input": {
+          "dtype": "uint32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [0, 16777216, 16777217, 4000000001] }
+        },
+        "min": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [16777217] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "uint32",
+          "shape": [4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [16777217, 16777217, 16777217, 4000000001] }
+        }
+      }
+    },
+    {
+      "name": "input_max_only_uint32",
+      "inputs": {
+        "input": {
+          "dtype": "uint32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [0, 16777217, 16777218, 4000000001] }
+        },
+        "max": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [16777217] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "uint32",
+          "shape": [4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0, 16777217, 16777217, 16777217] }
+        }
+      }
+    },
+    {
+      "name": "nan_input_propagates",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": ["NaN", -2.0, 0.0, 3.0] } },
+        "min": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-1.0] } },
+        "max": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] } }
+      },
+      "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": {
+        "input": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [-3.0, -0.5, 0.0, 2.0] } },
+        "max": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4], "tolerance": 0 } }
+    },
+    {
+      "name": "default_min_only",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [-3.0, 0.0, 0.5, 2.0] } },
+        "min": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_example_input_bounds",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-2.0, 0.0, 2.0] } },
+        "min": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-1.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_example" }
+    },
+    {
+      "name": "default_no_bounds_uses_float_limits",
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": ["-Infinity", -3.0, 3.0, "Infinity"] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4], "tolerance": 0 } }
+    },
+    {
+      "name": "default_int32_vec4_full_range",
+      "provenance": {
+        "source": "ONNX Clip-13 specification",
+        "test": "omitted integer bounds use the dtype range on the packed path"
+      },
+      "inputs": {
+        "input": {
+          "dtype": "int32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [-2147483648, -1, 1, 2147483647] }
+        }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [4], "tolerance": 0 } }
+    },
+    {
+      "name": "zero_size_noop",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [0, 3], "data": { "kind": "values", "values": [] } },
+        "min": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-1.0] } },
+        "max": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [0, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int32_input_bounds",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/clip_test.cc",
+        "test": "MathOpTest.Clip_int32",
+        "notes": "ORT supplies min/max as scalar tensor inputs; the input_bounds route consumes both values directly."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "int32",
+          "shape": [3, 3],
+          "data": { "kind": "values", "values": [-1, 0, 1, -16, 12, -6, -5, 2, 16] }
+        },
+        "min": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [-10] } },
+        "max": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [10] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [3, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_uint32_input_bounds",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/clip_test.cc",
+        "test": "MathOpTest.Clip_uint32",
+        "notes": "ORT supplies min/max as scalar tensor inputs; the input_bounds route consumes both values directly."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "uint32",
+          "shape": [3, 3],
+          "data": { "kind": "values", "values": [0, 0, 1, 5, 12, 3, 2, 7, 16] }
+        },
+        "min": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [3] } },
+        "max": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [10] } }
+      },
+      "outputs": { "output": { "dtype": "uint32", "shape": [3, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "int32_exact_above_float24_no_float_cast",
+      "inputs": {
+        "input": {
+          "dtype": "int32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [-16777218, -16777217, -16777216, 16777216, 16777217, 16777218] }
+        },
+        "min": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [-16777217] } },
+        "max": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [16777217] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [6], "tolerance": 0 } }
+    },
+    {
+      "name": "uint32_exact_above_float24_no_float_cast",
+      "inputs": {
+        "input": {
+          "dtype": "uint32",
+          "shape": [5],
+          "data": { "kind": "values", "values": [16777215, 16777216, 16777217, 16777218, 20000000] }
+        },
+        "min": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [16777216] } },
+        "max": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [16777217] } }
+      },
+      "outputs": { "output": { "dtype": "uint32", "shape": [5], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_input_bounds_f32_m10_p10",
+      "provenance": { "source": "onnxruntime/test/providers/cpu/math/clip_test.cc", "test": "MathOpTest.Clip_6" },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [3, 3],
+          "data": { "kind": "values", "values": [11.0, 4.4, 432.3, -1.3, 3.5, 64.0, -5.4, 9.3, 82.4] }
+        },
+        "min": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-10.0] } },
+        "max": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [10.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 3], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "ort_default_f32_3x3",
+      "provenance": { "source": "onnxruntime/test/providers/cpu/math/clip_test.cc", "test": "MathOpTest.Clip_Default" },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [3, 3],
+          "data": { "kind": "values", "values": [11.0, 4.4, 432.3, -1.3, 3.5, 64.0, -5.4, 9.3, 82.4] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_default_int8_3x3",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/clip_test.cc",
+        "test": "MathOpTest.Clip_Default_int8"
+      },
+      "inputs": {
+        "input": {
+          "dtype": "int8",
+          "shape": [3, 3],
+          "data": { "kind": "values", "values": [11, 4, 127, -1, 3, 64, -5, 9, 82] }
+        }
+      },
+      "outputs": { "output": { "dtype": "int8", "shape": [3, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_default_uint8_3x3",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/clip_test.cc",
+        "test": "MathOpTest.Clip_Default_uint8"
+      },
+      "inputs": {
+        "input": {
+          "dtype": "uint8",
+          "shape": [3, 3],
+          "data": { "kind": "values", "values": [11, 4, 255, 1, 3, 64, 5, 9, 82] }
+        }
+      },
+      "outputs": { "output": { "dtype": "uint8", "shape": [3, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_f16_input_bounds",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/clip_test.cc",
+        "test": "MathOpTest.Clip_MLFloat16"
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float16",
+          "shape": [3, 3],
+          "data": { "kind": "values", "values": [-1.0, -2.0, -3.0, -4.0, 0.0, 2.0, 4.0, 6.0, 8.0] }
+        },
+        "min": { "dtype": "float16", "shape": [], "data": { "kind": "values", "values": [0.0] } },
+        "max": { "dtype": "float16", "shape": [], "data": { "kind": "values", "values": [6.0] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [3, 3], "tolerance": 0.001 } }
+    },
+    {
+      "name": "ort_input_bounds_f32_3x3",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/clip_test.cc",
+        "test": "MathOpTest.Clip",
+        "notes": "ORT supplies min/max as scalar inputs; this fixture supplies both bounds through Clip-13 scalar tensor inputs."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [3, 3],
+          "data": { "kind": "values", "values": [-1.0, 0.0, 1.0, -6.0, 0.0, 6.0, -5.4, 2.0, 6.0] }
+        },
+        "min": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-5.0] } },
+        "max": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [5.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_f16_default_bounds",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/clip_test.cc",
+        "test": "MathOpTest.Clip_MLFloat16_NoMin_NoMax"
+      },
+      "inputs": {
+        "input": { "dtype": "float16", "shape": [3], "data": { "kind": "values", "values": [-1.0, -2.0, 3.0] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [3], "tolerance": 0.001 } }
+    },
+    {
+      "name": "ort_f16_min_only",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/clip_test.cc",
+        "test": "MathOpTest.Clip_MLFloat16_NoMax"
+      },
+      "inputs": {
+        "input": { "dtype": "float16", "shape": [3], "data": { "kind": "values", "values": [-1.0, -2.0, 3.0] } },
+        "min": { "dtype": "float16", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [3], "tolerance": 0.001 } }
+    },
+    {
+      "name": "ort_f16_max_only",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/clip_test.cc",
+        "test": "MathOpTest.Clip_MLFloat16_NoMin"
+      },
+      "inputs": {
+        "input": { "dtype": "float16", "shape": [3], "data": { "kind": "values", "values": [-1.0, -2.0, 3.0] } },
+        "max": { "dtype": "float16", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [3], "tolerance": 0.001 } }
+    },
+    {
+      "name": "ort_relu6_f32_3x3",
+      "provenance": { "source": "onnxruntime/test/providers/cpu/math/clip_test.cc", "test": "MathOpTest.Clip_Relu6" },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [3, 3],
+          "data": { "kind": "values", "values": [-1.0, 0.0, 1.0, -6.0, 3.5, 6.0, -5.4, 2.0, 8.0] }
+        },
+        "min": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } },
+        "max": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [6.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_relu_min_only_f32_3x3",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/clip_test.cc",
+        "test": "MathOpTest.Clip_Relu",
+        "notes": "ORT supplies the lower bound as a scalar input; this fixture exercises the same Clip-13 input directly."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [3, 3],
+          "data": { "kind": "values", "values": [-1.0, 0.0, 1.0, -6.0, 3.5, 6.0, -5.4, 2.0, 8.0] }
+        },
+        "min": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [3, 3],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0.0, 0.0, 1.0, 0.0, 3.5, 6.0, 0.0, 2.0, 8.0] }
+        }
+      }
+    },
+    {
+      "name": "ort_relu1_bounds_f32_3x3",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/clip_test.cc",
+        "test": "MathOpTest.Clip_Relu1",
+        "notes": "ORT supplies both bounds as scalar inputs; this fixture exercises the same Clip-13 inputs directly."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [3, 3],
+          "data": { "kind": "values", "values": [-1.0, 0.0, 1.0, -6.0, 3.5, 6.0, -5.4, 2.0, 8.0] }
+        },
+        "min": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-1.0] } },
+        "max": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [3, 3],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [-1.0, 0.0, 1.0, -1.0, 1.0, 1.0, -1.0, 1.0, 1.0] }
+        }
+      }
+    },
+    {
+      "name": "ort_empty_dim_input_bounds",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/clip_test.cc",
+        "test": "MathOpTest.ClipDimWithZero"
+      },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [3, 0], "data": { "kind": "values", "values": [] } },
+        "min": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-5.0] } },
+        "max": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [5.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 0], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_clip",
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_clip_input" } }
+        },
+        "min": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-1.0] } },
+        "max": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 4, 5] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_clip",
+        "notes": "ONNX Clip min/max bounds are supplied as scalar tensor inputs."
+      }
+    },
+    {
+      "name": "onnx_backend_clip_default_inbounds",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-1.0, 0.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_clip_default_inbounds",
+        "notes": "ONNX Clip min/max bounds are supplied as scalar tensor inputs."
+      }
+    },
+    {
+      "name": "onnx_backend_clip_default_int8_inbounds",
+      "inputs": { "input": { "dtype": "int8", "shape": [3], "data": { "kind": "values", "values": [-1, 0, 1] } } },
+      "outputs": { "output": { "dtype": "int8", "shape": [3] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_clip_default_int8_inbounds",
+        "notes": "ONNX Clip min/max bounds are supplied as scalar tensor inputs."
+      }
+    },
+    {
+      "name": "onnx_backend_clip_default_int8_max",
+      "inputs": {
+        "input": {
+          "dtype": "int8",
+          "shape": [3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [0, 0, 0, -1, 0, 0, -1, 0, 0, 0, 0, 0, 1, -1, 0, 0, 0, 0, 0, 0, -1, 0, 0, -1, 1, 1, 1, 0, -1, 1, 0, 1, 0, 0, 0, 0, 0, 1, 0, 0, 1, -1, -1, 0, -1, 1, 0, 0, 1, 1, 1, 0, 0, 1, 0, 0, 0, 0, 0, 0]
+          }
+        },
+        "max": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "output": { "dtype": "int8", "shape": [3, 4, 5] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_clip_default_int8_max",
+        "notes": "ONNX Clip min/max bounds are supplied as scalar tensor inputs."
+      }
+    },
+    {
+      "name": "onnx_backend_clip_default_int8_min",
+      "inputs": {
+        "input": {
+          "dtype": "int8",
+          "shape": [3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [1, 0, 0, 2, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 0, -2, 0, 0, 0, 2, -1, 0, 0, 1, 1, 0, 0, 0, -1, 0, 0, 1, 1, 0, 0, -1, -1, -1, 1, 0, 0, -1, 0, -1, 0, 0, 0, 0, -1, 0, 0, 0, 0, 0, 0]
+          }
+        },
+        "min": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "output": { "dtype": "int8", "shape": [3, 4, 5] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_clip_default_int8_min",
+        "notes": "ONNX Clip min/max bounds are supplied as scalar tensor inputs."
+      }
+    },
+    {
+      "name": "onnx_backend_clip_default_max",
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [-0.6724604368209839, -0.35955315828323364, -0.8131462931632996, -1.7262825965881348, 0.17742614448070526, -0.4017809331417084, -1.630198359489441, 0.46278226375579834, -0.9072983860969543, 0.05194539576768875, 0.7290905714035034, 0.12898291647434235, 1.1394007205963135, -1.234825849533081, 0.4023416340351105, -0.6848101019859314, -0.8707971572875977, -0.5788496732711792, -0.3115525245666504, 0.056165341287851334, -1.1651498079299927, 0.9008265137672424, 0.4656624495983124, -1.5362436771392822, 1.4882521629333496, 1.895889163017273, 1.1787796020507812, -0.1799248307943344, -1.0707526206970215, 1.0544517040252686, -0.4031769335269928, 1.222445011138916, 0.2082749754190445, 0.9766390323638916, 0.3563663959503174, 0.7065731883049011, 0.01050002034753561, 1.7858705520629883, 0.12691208720207214, 0.4019893705844879, 1.8831506967544556, -1.3477590084075928, -1.2704850435256958, 0.969396710395813, -1.1731233596801758, 1.9436211585998535, -0.4136189818382263, -0.747454822063446, 1.922942042350769, 1.4805147647857666, 1.8675589561462402, 0.9060446619987488, -0.8612256646156311, 1.910064935684204, -0.26800337433815, 0.8024563789367676, 0.9472519755363464, -0.15501008927822113, 0.6140793561935425, 0.922206699848175]
+          }
+        },
+        "max": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 4, 5] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_clip_default_max",
+        "notes": "ONNX Clip min/max bounds are supplied as scalar tensor inputs."
+      }
+    },
+    {
+      "name": "onnx_backend_clip_default_min",
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_clip_input" } }
+        },
+        "min": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 4, 5] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_clip_default_min",
+        "notes": "ONNX Clip min/max bounds are supplied as scalar tensor inputs."
+      }
+    },
+    {
+      "name": "onnx_backend_clip_example",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-2.0, 0.0, 2.0] } },
+        "min": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-1.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_example",
+        "notes": "ONNX Clip min/max bounds are supplied as scalar tensor inputs."
+      }
+    },
+    {
+      "name": "onnx_backend_clip_inbounds",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-1.0, 0.0, 1.0] } },
+        "min": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-5.0] } },
+        "max": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [5.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_clip_inbounds",
+        "notes": "ONNX Clip min/max bounds are supplied as scalar tensor inputs."
+      }
+    },
+    {
+      "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."
+      }
+    },
+    {
+      "name": "onnx_backend_clip_outbounds",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-6.0, 0.0, 6.0] } },
+        "min": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-5.0] } },
+        "max": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [5.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_clip_outbounds",
+        "notes": "ONNX Clip min/max bounds are supplied as scalar tensor inputs."
+      }
+    },
+    {
+      "name": "onnx_backend_clip_splitbounds",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-1.0, 0.0, 6.0] } },
+        "min": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-5.0] } },
+        "max": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [5.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_clip_splitbounds",
+        "notes": "ONNX Clip min/max bounds are supplied as scalar tensor inputs."
+      }
+    },
+    {
+      "name": "vec4_u32_lanes",
+      "inputs": {
+        "input": {
+          "dtype": "uint32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0, 1, 5, 7, 100, 200, 201, 250] }
+        },
+        "min": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [5] } },
+        "max": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [200] } }
+      },
+      "outputs": { "output": { "dtype": "uint32", "shape": [8] } }
+    },
+    {
+      "name": "vec4_f32_nan_passthrough",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": ["NaN", 0.2, 5.0, -5.0] } },
+        "min": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } },
+        "max": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4], "tolerance": 0.000001, "allowNaN": true } }
+    },
+    {
+      "name": "vec4_i32_lanes",
+      "inputs": {
+        "input": {
+          "dtype": "int32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [-2147483648, -100, -50, -1, 0, 49, 50, 2147483647] }
+        },
+        "min": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [-50] } },
+        "max": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [50] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [8] } }
+    },
+    {
+      "name": "vec4_u8_lanes",
+      "inputs": {
+        "input": {
+          "dtype": "uint8",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0, 1, 5, 7, 100, 200, 201, 255] }
+        },
+        "min": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [5] } },
+        "max": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [200] } }
+      },
+      "outputs": { "output": { "dtype": "uint8", "shape": [8] } }
+    },
+    {
+      "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] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "int32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [-5, -5, -5, -5, -5, -5] },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "f32_nan_passthrough_with_input_bounds",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": ["NaN", -3.0, 0.5, 3.0] } },
+        "min": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-1.0] } },
+        "max": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": ["NaN", -1.0, 0.5, 1.0] },
+          "tolerance": 0,
+          "allowNaN": true
+        }
+      }
+    },
+    {
+      "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] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "uint8",
+          "shape": [5],
+          "data": { "kind": "values", "values": [100, 100, 100, 100, 100] },
+          "tolerance": 0
+        }
+      }
+    }
+  ]
+}

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

	@@ -0,0 +1,66 @@

+{% 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 %}
+{{ env.wgsl.resourceDeclarations }}
+{% macro emit_is_nan_f32() %}
+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 %}
+  output[i] = max(value, clipMin[0]);
+{% else %}
+  output[i] = min(value, clipMax[0]);
+{% endif %}
+{% else %}
+  output[i] = min(max(value, {{ clipScalar }}({{ clipMinValue }})), {{ clipScalar }}({{ clipMaxValue }}));
+{% 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 %}
+  var out = max(value, f32(clipMin[0]));
+{% else %}
+  var out = min(value, f32(clipMax[0]));
+{% endif %}
+{% else %}
+  var out = min(max(value, f32({{ clipMinValue }})), f32({{ clipMaxValue }}));
+{% endif %}
+  if (is_nan_f32(value)) {
+    out = value;
+  }
+  output[i] = {{ clipScalar }}(out);
+{% endif %}
+{{ flat_tail_close() -}}
+}

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

	@@ -0,0 +1,73 @@

+// 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() %}
+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)) {
+    out = v;
+  }
+  return out;
+}
+{% 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]));
+{% elif clipMinFromInput %}
+  y[i] = max(xv, vec4<{{ scalar }}>(clipMin[0]));
+{% else %}
+  y[i] = min(xv, vec4<{{ scalar }}>(clipMax[0]));
+{% endif %}
+{% else %}
+  let fv = vec4<f32>(xv);
+{% if clipMinFromInput and clipMaxFromInput %}
+  var outv = min(max(fv, vec4<f32>(f32(clipMin[0]))), vec4<f32>(f32(clipMax[0])));
+{% elif clipMinFromInput %}
+  var outv = max(fv, vec4<f32>(f32(clipMin[0])));
+{% else %}
+  var outv = min(fv, vec4<f32>(f32(clipMax[0])));
+{% endif %}
+  if (is_nan_f32(fv.x)) { outv.x = fv.x; }
+  if (is_nan_f32(fv.y)) { outv.y = fv.y; }
+  if (is_nan_f32(fv.z)) { outv.z = fv.z; }
+  if (is_nan_f32(fv.w)) { outv.w = fv.w; }
+  y[i] = {{ vectorScalar }}(outv);
+{% endif %}
+{% elif scalar == "i32" or scalar == "u32" %}
+  let lower = {{ scalar }}({{ clipMinValue }});
+  let upper = {{ scalar }}({{ clipMaxValue }});
+  y[i] = min(max(xv, vec4<{{ scalar }}>(lower)), vec4<{{ scalar }}>(upper));
+{% else %}
+  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 %}
+}