sync 2e7068faf55e

Browse files

Files changed (7) hide show

README.md +57 -0
build/webgpu/bench.json +45 -0
build/webgpu/manifest.json +101 -0
build/webgpu/metadata.json +19 -0
build/webgpu/test.json +317 -0
build/webgpu/unary-scalar.wgsl.jinja +41 -0
build/webgpu/unary-vec4.wgsl.jinja +31 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.Sign
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 13
+## Description
+Computes the elementwise sign of the input tensor: outputs `1` for positive values, `-1` for negative values, and `0` where the input is zero. The output has the same shape and type as the input.
+See the [ONNX `Sign` spec](https://onnx.ai/onnx/operators/onnx__Sign.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `input` | `input` | `T` | — | — | Input tensor of any numeric type. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `output` | `output` | `T` | same as `input` | same as `input` | Elementwise sign of the input; same shape and type as the input. | required |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16`, `int32`, `uint32`, `int16`, `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.Sign", { version: 1 });
+const { output } = await kernel({ input: { data: inputData, shape: [] } });
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,45 @@

+{
+  "op": "ai.onnx.Sign",
+  "cases": [
+    {
+      "name": "1m_f32",
+      "inputs": { "input": { "dtype": "float32", "shape": [1048576] } },
+      "outputs": { "output": { "dtype": "float32", "shape": [1048576] } }
+    },
+    {
+      "name": "sign-f32-1m-vec4",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "count": 1048576 },
+      "inputs": { "input": { "shape": [1048576], "dtype": "float32", "dist": "normal", "seed": 7201, "scale": 1 } },
+      "outputs": { "output": { "shape": [1048576], "dtype": "float32" } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "numel(shapes.input) * dtypeBytes(args.dtype) * 2" }]
+      }
+    },
+    {
+      "name": "sign-f32-1m-minus1-scalar-fallback",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "count": 1048575 },
+      "inputs": { "input": { "shape": [1048575], "dtype": "float32", "dist": "normal", "seed": 7202, "scale": 1 } },
+      "outputs": { "output": { "shape": [1048575], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "numel(shapes.input) * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "sign-f16-8m-vec4",
+      "preset": "smoke",
+      "vars": { "dtype": "float16", "count": 8388608 },
+      "inputs": { "input": { "shape": [8388608], "dtype": "float16", "dist": "normal", "seed": 7203, "scale": 1 } },
+      "outputs": { "output": { "shape": [8388608], "dtype": "float16" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "numel(shapes.input) * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "sign-f16-8m-minus1-scalar-fallback",
+      "preset": "smoke",
+      "vars": { "dtype": "float16", "count": 8388607 },
+      "inputs": { "input": { "shape": [8388607], "dtype": "float16", "dist": "normal", "seed": 7204, "scale": 1 } },
+      "outputs": { "output": { "shape": [8388607], "dtype": "float16" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "numel(shapes.input) * dtypeBytes(args.dtype) * 2" }] }
+    }
+  ]
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,101 @@

+{
+  "domain": "ai.onnx",
+  "name": "Sign",
+  "sinceVersion": 13,
+  "description": "Computes the elementwise sign of the input tensor: outputs `1` for positive values, `-1` for negative values, and `0` where the input is zero. The output has the same shape and type as the input.",
+  "inputs": [{ "role": "input", "dtype": "T", "description": "Input tensor of any numeric type." }],
+  "outputs": [
+    {
+      "role": "output",
+      "dtype": "T",
+      "rank": "ranks.input",
+      "description": "Elementwise sign of the input; same shape and type as the input.",
+      "shape": "shapes.input"
+    }
+  ],
+  "typeConstraints": { "T": ["float32", "float16", "int32", "uint32", "int16", "int8", "uint8"] },
+  "args": {
+    "input": { "kind": "tensor", "semantic": "input", "role": "input" },
+    "output": { "kind": "tensor", "semantic": "output", "role": "output" }
+  },
+  "tunables": { "WORKGROUP_SIZE": 256 },
+  "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
+  "variants": [
+    {
+      "id": "same_layout_vec4",
+      "when": ["numel(shapes.input) > 0", "numel(shapes.input) % 4 == 0", "numel(shapes.input) == numel(shapes.output)", "f16Ok(dtypes.T)"],
+      "constants": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Sign.vec4",
+          "source": { "shader": "unary-vec4.wgsl.jinja", "inputs": { "op": "\"sign\"" } },
+          "bindings": [
+            {
+              "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" }]
+              }
+            }
+          ],
+          "dispatch": { "threads": "numel(shapes.output) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ],
+      "priority": 20
+    },
+    {
+      "id": "elementwise",
+      "when": "numel(shapes.input) == numel(shapes.output) and (f16Ok(dtypes.T))",
+      "passes": [
+        {
+          "id": "main",
+          "name": "Sign",
+          "source": { "shader": "unary-scalar.wgsl.jinja", "inputs": { "op": "\"sign\"", "itemsPerInvocation": 4 } },
+          "bindings": [
+            {
+              "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)" }]
+              }
+            }
+          ],
+          "dispatch": { "threads": "ceilDiv(numel(shapes.output), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

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

+{
+  "name": "ai.onnx.Sign",
+  "id": "_ai_onnx_sign_webgpu_6d0a2b1",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "4q8mZQrcY+ViFixjxj3zYSAeXdqwAnfqK6JL0eXfAeU=",
+      "manifest.json": "4+AgHdYHHwfbRxAG9bwppknCrSxXYIjvelehzFR1osw=",
+      "test.json": "xB83KjUJDVMn8MAFK2+ZMNngy6SCZSwPM0ALPFALosA=",
+      "unary-scalar.wgsl.jinja": "dCTx9hbmVW/fySnWmzBzZT9z4Ij9/iG6RFliizgwwxs=",
+      "unary-vec4.wgsl.jinja": "h0Y4ksIDPy57eIshgKVikk0FJSpUp/mXgHTtmxmVRNw="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Sign" }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,317 @@

+{
+  "op": "ai.onnx.Sign",
+  "cases": [
+    {
+      "name": "int16_scalar_boundaries",
+      "inputs": {
+        "input": { "dtype": "int16", "shape": [5], "data": { "kind": "values", "values": [-32768, -1, 0, 1, 32767] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "int16",
+          "shape": [5],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [-1, -1, 0, 1, 1] }
+        }
+      }
+    },
+    {
+      "name": "f32_signed_subnormal_sign_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/sign_test.cc",
+        "test": "MathOpTest.Sign_float",
+        "notes": "Signed subnormal values have nonzero signs; flushing them to zero incorrectly returns 0."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": { "kind": "values", "values": [-1e-39, -1e-40, 0.0, 1e-40, 1e-39] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [5], "tolerance": 0 } }
+    },
+    {
+      "name": "f32_signed_subnormal_sign_vec4_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/sign_test.cc",
+        "test": "MathOpTest.Sign_float",
+        "notes": "Vec4 companion for signed subnormal Sign behavior; optimized unary paths must preserve nonzero signs."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [-1e-40, 0.0, 1e-40, 1e-39] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4], "tolerance": 0 } }
+    },
+    {
+      "name": "f32_values",
+      "provenance": { "source": "onnxruntime/test/providers/cpu/math/sign_test.cc", "test": "MathOpTest.Sign_float" },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": { "kind": "values", "values": [-2.0, -0.1, 0.0, 0.1, 2.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [5] } }
+    },
+    {
+      "name": "f32_zero_and_infinities",
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": ["-Infinity", 0.0, 0.0, 1e-7, -1e-7, "Infinity"] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [6], "tolerance": 0 } }
+    },
+    {
+      "name": "f16_values",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/sign_test.cc",
+        "test": "MathOpTest.Sign_MLFloat16"
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float16",
+          "shape": [6],
+          "data": { "kind": "values", "values": [-4.0, -0.5, 0.0, 0.0, 0.5, 4.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [6], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int32_sequence",
+      "provenance": { "source": "onnxruntime/test/providers/cpu/math/sign_test.cc", "test": "MathOpTest.Sign_int64" },
+      "inputs": {
+        "input": { "dtype": "int32", "shape": [7], "data": { "kind": "values", "values": [-5, -4, -3, -2, -1, 0, 1] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [7], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_uint32_sequence",
+      "provenance": { "source": "onnxruntime/test/providers/cpu/math/sign_test.cc", "test": "MathOpTest.Sign_uint64" },
+      "inputs": {
+        "input": { "dtype": "uint32", "shape": [7], "data": { "kind": "values", "values": [0, 1, 2, 3, 4, 5, 6] } }
+      },
+      "outputs": { "output": { "dtype": "uint32", "shape": [7], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int8_sequence",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/sign_test.cc",
+        "test": "MathOpTest.Sign_int64",
+        "notes": "Uses the same signed integer sequence semantics as MathOpTest.Sign_int64 with logical int8 storage."
+      },
+      "inputs": {
+        "input": { "dtype": "int8", "shape": [7], "data": { "kind": "values", "values": [-5, -4, -3, -2, -1, 0, 1] } }
+      },
+      "outputs": { "output": { "dtype": "int8", "shape": [7], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_uint8_sequence",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/sign_test.cc",
+        "test": "MathOpTest.Sign_uint64",
+        "notes": "Uses the same unsigned integer sequence semantics as MathOpTest.Sign_uint64 with logical uint8 storage."
+      },
+      "inputs": {
+        "input": { "dtype": "uint8", "shape": [7], "data": { "kind": "values", "values": [0, 1, 2, 3, 4, 5, 6] } }
+      },
+      "outputs": { "output": { "dtype": "uint8", "shape": [7], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_float32_range",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_sign" },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [11],
+          "data": { "kind": "values", "values": [-5.0, -4.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, 4.0, 5.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [11], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_sign",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_sign" },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [11],
+          "data": { "kind": "values", "values": [-5.0, -4.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, 4.0, 5.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [11], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "vec4_f32_lanes",
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [-5.0, -0.1, 0.0, 0.1, 5.0, "-Infinity", "Infinity", 2.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [8], "tolerance": 0 } }
+    },
+    {
+      "name": "vec4_f16_lanes",
+      "inputs": {
+        "input": {
+          "dtype": "float16",
+          "shape": [8],
+          "data": { "kind": "values", "values": [-4.0, -0.5, 0.0, 0.0, 0.5, 4.0, -2.0, 2.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [8], "tolerance": 0 } }
+    },
+    {
+      "name": "vec4_i32_lanes",
+      "inputs": {
+        "input": { "dtype": "int32", "shape": [8], "data": { "kind": "values", "values": [-5, -1, 0, 1, 5, -2, 2, 0] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [8] } }
+    },
+    {
+      "name": "vec4_u32_lanes",
+      "inputs": {
+        "input": { "dtype": "uint32", "shape": [8], "data": { "kind": "values", "values": [0, 1, 2, 3, 4, 5, 0, 7] } }
+      },
+      "outputs": { "output": { "dtype": "uint32", "shape": [8] } }
+    },
+    {
+      "name": "vec4_i8_lanes",
+      "inputs": {
+        "input": {
+          "dtype": "int8",
+          "shape": [8],
+          "data": { "kind": "values", "values": [-128, -5, -1, 0, 1, 5, 100, 127] }
+        }
+      },
+      "outputs": { "output": { "dtype": "int8", "shape": [8] } }
+    },
+    {
+      "name": "vec4_u8_lanes",
+      "inputs": {
+        "input": {
+          "dtype": "uint8",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0, 1, 2, 3, 128, 200, 254, 255] }
+        }
+      },
+      "outputs": { "output": { "dtype": "uint8", "shape": [8] } }
+    },
+    {
+      "name": "ort_float_nan_preserved",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/sign_test.cc",
+        "test": "MathOpTest.Sign_float",
+        "notes": "Extends the ORT float Sign coverage with ORT CPU-verified NaN preservation."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": ["NaN", "-Infinity", 0.0, "Infinity"] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4], "tolerance": 0, "allowNaN": true } }
+    },
+    {
+      "name": "ort_float_nan_preserved_scalar",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/sign_test.cc",
+        "test": "MathOpTest.Sign_float",
+        "notes": "Scalar-path companion for ORT CPU-verified NaN preservation."
+      },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": ["NaN", -2.0, 2.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3], "tolerance": 0, "allowNaN": true } }
+    },
+    {
+      "name": "empty_input_zero_dim",
+      "inputs": { "input": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } } },
+      "outputs": { "output": { "dtype": "float32", "shape": [0], "tolerance": 0 } }
+    },
+    {
+      "name": "f32_rank0_scalar",
+      "inputs": { "input": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-3.5] } } },
+      "outputs": { "output": { "dtype": "float32", "shape": [], "tolerance": 0 } }
+    },
+    {
+      "name": "f32_rank2_3x5",
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [3, 5],
+          "data": {
+            "kind": "values",
+            "values": [-7.0, -0.5, 0.0, 0.5, 7.0, -1e-7, 1e-7, -100.0, 100.0, 0.0, -2.0, -1.0, 1.0, 2.0, 3.0]
+          }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "f16_rank2_realistic",
+      "inputs": {
+        "input": {
+          "dtype": "float16",
+          "shape": [4, 16],
+          "data": {
+            "kind": "values",
+            "values": [-8.0, -4.0, -2.0, -1.0, -0.5, -0.25, -0.125, -0.0625, 0.0, 0.0625, 0.125, 0.25, 0.5, 1.0, 2.0, 4.0, -7.5, -3.5, -1.5, -0.75, -0.375, -0.1875, -0.09375, -0.046875, 0.046875, 0.09375, 0.1875, 0.375, 0.75, 1.5, 3.5, 7.5, 0.0, 0.0, -6.0, 6.0, -5.0, 5.0, -3.0, 3.0, -2.5, 2.5, -1.25, 1.25, -0.625, 0.625, -0.3125, 0.3125, 8.0, -8.0, 4.0, -4.0, 2.0, -2.0, 1.0, -1.0, 0.5, -0.5, 0.25, -0.25, 0.125, -0.125, 0.0625, -0.0625]
+          }
+        }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [4, 16], "tolerance": 0 } }
+    },
+    {
+      "name": "f16_nan_preserved_scalar_path",
+      "inputs": {
+        "input": {
+          "dtype": "float16",
+          "shape": [5],
+          "data": { "kind": "values", "values": ["NaN", -2.0, 0.0, 2.0, "NaN"] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [5], "tolerance": 0, "allowNaN": true } }
+    },
+    {
+      "name": "int8_min_value_sign_scalar_path",
+      "inputs": {
+        "input": { "dtype": "int8", "shape": [5], "data": { "kind": "values", "values": [-128, -1, 0, 1, 127] } }
+      },
+      "outputs": { "output": { "dtype": "int8", "shape": [5], "tolerance": 0 } }
+    },
+    {
+      "name": "f32_scalar_tail_4095",
+      "provenance": {
+        "notes": "Compact sibling for the large f32 Sign scalar-fallback benchmark; odd numel forces the elementwise path while preserving exact sign semantics."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [4095],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.027, "scale": 3.0 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4095], "tolerance": 0 } }
+    }
+  ]
+}

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

	@@ -0,0 +1,41 @@

+{% macro flat_tail_open() %}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+  // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
+  // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills the rest into y).
+  let invocation = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  // Tail-safe scalar x4 keeps vector-like dispatch density without requiring
+  // the logical tensor length (or its storage binding) to be vec4 aligned.
+  let begin = invocation * {{ source.itemsPerInvocation }}u;
+  let end = min(begin + {{ source.itemsPerInvocation }}u, params.count);
+  for (var i = begin; i < end; i = i + 1u) {
+{%- endmacro %}
+{% macro flat_tail_close() %}
+  }
+{% endmacro %}
+// 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 }}
+{{ flat_tail_open() }}
+{% if scalar == "i32" %}
+  let v = input[i];
+  output[i] = select(select(0i, 1i, v > 0i), -1i, v < 0i);
+{% elif scalar == "u32" %}
+  let v = input[i];
+  output[i] = select(0u, 1u, v > 0u);
+{% else %}
+  let v = f32(input[i]);
+  // ONNX Sign preserves NaN. A bitcast test cannot be folded away by no-NaN
+  // assumptions; it isolates NaN, whose ordered comparisons would collapse to 0.
+  let v_bits = bitcast<u32>(v);
+  let v_is_nan = (v_bits & 0x7f800000u) == 0x7f800000u && (v_bits & 0x007fffffu) != 0u;
+  var s = select(select(0.0, 1.0, v > 0.0), -1.0, v < 0.0);
+  if (v_is_nan) { s = v; }
+  output[i] = {{ scalar }}(s);
+{% endif %}
+{{ flat_tail_close() -}}
+}

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

	@@ -0,0 +1,31 @@

+// 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 }}
+@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];
+{% if scalar == "i32" %}
+  y[i] = select(select(vec4<i32>(0i), vec4<i32>(1i), xv > vec4<i32>(0i)), vec4<i32>(-1i), xv < vec4<i32>(0i));
+{% elif scalar == "u32" %}
+  y[i] = select(vec4<u32>(0u), vec4<u32>(1u), xv > vec4<u32>(0u));
+{% else %}
+  let fv = vec4<f32>(xv);
+  // ONNX Sign preserves NaN. A bitcast test cannot be folded away by no-NaN
+  // assumptions; it isolates lanes whose ordered comparisons would collapse to 0.
+  let fv_bits = bitcast<vec4<u32>>(fv);
+  let fv_is_nan = ((fv_bits & vec4<u32>(0x7f800000u)) == vec4<u32>(0x7f800000u)) & ((fv_bits & vec4<u32>(0x007fffffu)) != vec4<u32>(0u));
+  let signed = select(select(vec4<f32>(0.0), vec4<f32>(1.0), fv > vec4<f32>(0.0)), vec4<f32>(-1.0), fv < vec4<f32>(0.0));
+  y[i] = {{ vectorScalar }}(select(signed, fv, fv_is_nan));
+{% endif %}
+}