sync 2e7068faf55e

Browse files

Files changed (6) hide show

README.md +57 -0
build/webgpu/bench.json +39 -0
build/webgpu/elementwise-bias-gelu.wgsl.jinja +90 -0
build/webgpu/manifest.json +221 -0
build/webgpu/metadata.json +18 -0
build/webgpu/test.json +448 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# com.microsoft.FastGelu
+`com.microsoft`  ·  ONNX Runtime contrib operator  ·  contrib since_version 1
+## Description
+Applies the GELU (Gaussian Error Linear Unit) activation using a tanh approximation: `Y = 0.5 * X * (1 + tanh(0.797885 * X + 0.035677 * X^3))`. An optional `bias` is added to `X` before the activation is computed. This WebGPU package implements float16 and float32; the schema-allowed double and bfloat16 types are not supported.
+See the [ONNX Runtime `FastGelu` contrib-operator spec](https://github.com/microsoft/onnxruntime/blob/main/docs/ContribOperators.md#com.microsoft.FastGelu) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `X` | `X` | `T` | — | — | Values transformed by FastGelu after adding the optional `bias`. | required |
+| `bias` | `bias` | `T` | `1` | — | Optional 1-D bias added to `X` along the last dimension before the GELU activation. | optional |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `Y` | `Y` | `T` | same as `X` | same as `X` | Output tensor after applying the GELU activation; same shape as `X`. | required |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16` |
+## 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
+- [`elementwise-bias-gelu.wgsl.jinja`](build/webgpu/elementwise-bias-gelu.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/com.microsoft.FastGelu", { version: 1 });
+const { Y } = await kernel({ X: { data: XData, shape: [5] } });
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,39 @@

+{
+  "op": "com.microsoft.FastGelu",
+  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
+  "cases": [
+    {
+      "name": "fastgelu-f32-bias-4096x3072",
+      "preset": "smoke",
+      "vars": { "dtype": "float32" },
+      "inputs": {
+        "X": { "shape": [4096, 3072], "dtype": "float32", "dist": "normal", "seed": 410, "scale": 2 },
+        "bias": { "shape": [3072], "dtype": "float32", "dist": "normal", "seed": 411, "scale": 1 }
+      },
+      "outputs": { "Y": { "shape": [4096, 3072], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4096 * 3072 * 2 * dtypeBytes(args.dtype)" }] }
+    },
+    {
+      "name": "fastgelu-f16-bias-4096x3072",
+      "preset": "model",
+      "vars": { "dtype": "float16" },
+      "inputs": {
+        "X": { "shape": [4096, 3072], "dtype": "float16", "dist": "normal", "seed": 412, "scale": 2 },
+        "bias": { "shape": [3072], "dtype": "float16", "dist": "normal", "seed": 413, "scale": 1 }
+      },
+      "outputs": { "Y": { "shape": [4096, 3072], "dtype": "float16" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4096 * 3072 * 2 * dtypeBytes(args.dtype)" }] }
+    },
+    {
+      "name": "fastgelu-f32-scalar-bias-odd-hidden",
+      "preset": "stress",
+      "vars": { "dtype": "float32" },
+      "inputs": {
+        "X": { "shape": [2097152, 3], "dtype": "float32", "dist": "normal", "seed": 440, "scale": 2 },
+        "bias": { "shape": [3], "dtype": "float32", "dist": "normal", "seed": 441, "scale": 1 }
+      },
+      "outputs": { "Y": { "shape": [2097152, 3], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2097152 * 3 * 2 * dtypeBytes(args.dtype)" }] }
+    }
+  ]
+}

build/webgpu/elementwise-bias-gelu.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,90 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+{% set wg = workgroupSize if workgroupSize is defined else tunables.WORKGROUP_SIZE %}
+// Bias plus GELU, with a specialization-selected tanh or erf approximation.
+// The optional bias is a rank-1 vector broadcast over the innermost (hidden)
+// axis: bias index = element_index % HIDDEN. The `vec4` path requires
+// HIDDEN % 4 == 0 and numel % 4 == 0 so a vec4 group never crosses the hidden
+// axis (the bias slice is then contiguous). `vec4Tail` keeps scalar bindings but
+// evaluates four guarded lanes per invocation, so odd hidden sizes retain the
+// same parallel efficiency without crossing row/bias boundaries. Gelu math and overflow guards match
+// the vectorized unary implementation: tanh saturates to +/-1 by |x|~9, and the
+// erf path uses the same rational approximation as Gelu.
+{% if approximate == "erf" %}
+fn erf_approx(x: f32) -> f32 {
+  let ax = abs(x);
+  // The polynomial has a small nonzero floor near zero. Use erf(x) ~=
+  // 2/sqrt(pi)*x below 2^-20 to preserve erf(0) == 0, odd symmetry, and the
+  // correctly rounded f32 result. The exactly representable threshold keeps
+  // scalar and vector branching identical. NaN falls through to the polynomial
+  // and propagates.
+  if (ax < 9.5367431640625e-7) {
+    return 1.1283791670955126 * x;
+  }
+  let sign = select(-1.0, 1.0, x >= 0.0);
+  let t = 1.0 / (1.0 + 0.3275911 * ax);
+  let y = 1.0 - (((((1.061405429 * t - 1.453152027) * t) + 1.421413741) * t - 0.284496736) * t + 0.254829592) * t * exp(-(ax * ax));
+  return sign * y;
+}
+{% else %}
+fn tanh_safe(x: f32) -> f32 {
+  if (x > 10.0) { return 1.0; }
+  if (x < -10.0) { return -1.0; }
+  return tanh(x);
+}
+{% endif %}
+fn gelu_value(v: f32) -> f32 {
+{% if approximate == "erf" %}
+  return 0.5 * v * (1.0 + erf_approx(v * 0.7071067811865476));
+{% else %}
+  return 0.5 * v * (1.0 + tanh_safe(0.7978845608028654 * (v + 0.044715 * v * v * v)));
+{% endif %}
+}
+{% if hasBias %}
+const HIDDEN: u32 = {{ hidden }}u;
+{% endif %}
+@compute @workgroup_size({{ wg }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
+  let i = gid.x + gid.y * nwg.x * {{ wg }}u;
+  if (i >= params.count) {
+    return;
+  }
+{% if vec4Tail %}
+  let base = i * 4u;
+{% for lane in range(4) %}
+  if (base + {{ lane }}u < params.count) {
+    let xv{{ lane }} = f32(x[base + {{ lane }}u]);
+{% if hasBias %}
+    let v{{ lane }} = xv{{ lane }} + f32(bias[(base + {{ lane }}u) % HIDDEN]);
+{% else %}
+    let v{{ lane }} = xv{{ lane }};
+{% endif %}
+    y[base + {{ lane }}u] = {{ scalar }}(gelu_value(v{{ lane }}));
+  }
+{% endfor %}
+{% elif vec4 %}
+  let xv = vec4<f32>(x[i]);
+{% if hasBias %}
+  let bcol = (i * 4u) % HIDDEN;
+  let v = xv + vec4<f32>(f32(bias[bcol]), f32(bias[bcol + 1u]), f32(bias[bcol + 2u]), f32(bias[bcol + 3u]));
+{% else %}
+  let v = xv;
+{% endif %}
+  y[i] = vec4<{{ scalar }}>(vec4<f32>(gelu_value(v.x), gelu_value(v.y), gelu_value(v.z), gelu_value(v.w)));
+{% else %}
+  let xv = f32(x[i]);
+{% if hasBias %}
+  let v = xv + f32(bias[i % HIDDEN]);
+{% else %}
+  let v = xv;
+{% endif %}
+  y[i] = {{ scalar }}(gelu_value(v));
+{% endif %}
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,221 @@

+{
+  "domain": "com.microsoft",
+  "name": "FastGelu",
+  "sinceVersion": 1,
+  "description": "Applies the GELU (Gaussian Error Linear Unit) activation using a tanh approximation: `Y = 0.5 * X * (1 + tanh(0.797885 * X + 0.035677 * X^3))`. An optional `bias` is added to `X` before the activation is computed. This WebGPU package implements float16 and float32; the schema-allowed double and bfloat16 types are not supported.",
+  "inputs": [
+    { "role": "X", "dtype": "T", "description": "Values transformed by FastGelu after adding the optional `bias`." },
+    {
+      "role": "bias",
+      "dtype": "T",
+      "rank": 1,
+      "optional": true,
+      "description": "Optional 1-D bias added to `X` along the last dimension before the GELU activation."
+    }
+  ],
+  "outputs": [
+    {
+      "role": "Y",
+      "dtype": "T",
+      "rank": "ranks.X",
+      "shape": "shapes.X",
+      "description": "Output tensor after applying the GELU activation; same shape as `X`."
+    }
+  ],
+  "typeConstraints": { "T": ["float32", "float16"] },
+  "args": {
+    "X": { "kind": "tensor", "semantic": "X", "role": "input" },
+    "bias": { "kind": "tensor", "semantic": "bias", "role": "input", "required": false },
+    "Y": { "kind": "tensor", "semantic": "Y", "role": "output" }
+  },
+  "tunables": { "WORKGROUP_SIZE": 256 },
+  "derive": {
+    "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
+    "storageBufferLimit": "min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
+    "workgroupSize": "min(tunables.WORKGROUP_SIZE, deviceWorkgroupCap)",
+    "buffersFit": "numel(shapes.X) * dtypeBytes(dtypes.T) <= storageBufferLimit and numel(shapes.Y) * dtypeBytes(dtypes.T) <= storageBufferLimit and (not present.bias or numel(shapes.bias) * dtypeBytes(dtypes.T) <= storageBufferLimit)",
+    "dispatchFits": "numel(shapes.X) <= device.limits.maxComputeWorkgroupsPerDimension * device.limits.maxComputeWorkgroupsPerDimension * workgroupSize * 4",
+    "baseOk": "workgroupSize > 0 and ranks.X >= 1 and numel(shapes.X) == numel(shapes.Y) and f16Ok(dtypes.T) and buffersFit and dispatchFits",
+    "biasOk": "present.bias and ranks.bias == 1 and dim(shapes.bias, 0) == dim(shapes.X, ranks.X - 1)",
+    "noBiasOk": "not present.bias",
+    "vec4Ok": "numel(shapes.X) > 0 and numel(shapes.X) % 4 == 0 and dim(shapes.X, ranks.X - 1) % 4 == 0"
+  },
+  "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"", "approximate": "\"tanh\"" },
+  "bindingSets": {
+    "vec4Bias": [
+      {
+        "name": "x",
+        "arg": "X",
+        "semantic": "X",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "bias",
+        "arg": "bias",
+        "semantic": "bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar",
+        "length": "$hidden"
+      },
+      { "name": "y", "arg": "Y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$vectorScalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.X) / 4" }] }
+      }
+    ],
+    "vec4": [
+      {
+        "name": "x",
+        "arg": "X",
+        "semantic": "X",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      { "name": "y", "arg": "Y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$vectorScalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.X) / 4" }] }
+      }
+    ],
+    "scalarBias": [
+      { "name": "x", "arg": "X", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
+      {
+        "name": "bias",
+        "arg": "bias",
+        "semantic": "bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar"
+      },
+      { "name": "y", "arg": "Y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.X)" }] }
+      }
+    ],
+    "scalar": [
+      { "name": "x", "arg": "X", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
+      { "name": "y", "arg": "Y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.X)" }] }
+      }
+    ]
+  },
+  "variants": [
+    {
+      "id": "vec4_bias",
+      "priority": 30,
+      "when": ["baseOk", "biasOk", "vec4Ok"],
+      "constants": {
+        "vec4": true,
+        "vec4Tail": false,
+        "hasBias": true,
+        "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
+        "hidden": "dim(shapes.X, ranks.X - 1) if dim(shapes.X, ranks.X - 1) > 0 else 1"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "FastGelu.vec4Bias",
+          "shader": "elementwise-bias-gelu.wgsl.jinja",
+          "bindings": "vec4Bias",
+          "dispatch": { "threads": "numel(shapes.X) / 4", "workgroupSize": "workgroupSize" }
+        }
+      ]
+    },
+    {
+      "id": "vec4_no_bias",
+      "priority": 25,
+      "when": ["baseOk", "noBiasOk", "vec4Ok"],
+      "constants": { "vec4": true, "vec4Tail": false, "hasBias": false, "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "FastGelu.vec4",
+          "shader": "elementwise-bias-gelu.wgsl.jinja",
+          "bindings": "vec4",
+          "dispatch": { "threads": "numel(shapes.X) / 4", "workgroupSize": "workgroupSize" }
+        }
+      ]
+    },
+    {
+      "id": "vec4_tail_bias",
+      "priority": 20,
+      "when": ["baseOk", "biasOk", "numel(shapes.X) > 0"],
+      "constants": {
+        "vec4": false,
+        "vec4Tail": true,
+        "hasBias": true,
+        "hidden": "dim(shapes.X, ranks.X - 1) if dim(shapes.X, ranks.X - 1) > 0 else 1"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "FastGelu.vec4TailBias",
+          "shader": "elementwise-bias-gelu.wgsl.jinja",
+          "bindings": "scalarBias",
+          "dispatch": { "threads": "ceilDiv(numel(shapes.X), 4)", "workgroupSize": "workgroupSize" }
+        }
+      ]
+    },
+    {
+      "id": "vec4_tail_no_bias",
+      "priority": 15,
+      "when": ["baseOk", "noBiasOk", "numel(shapes.X) > 0"],
+      "constants": { "vec4": false, "vec4Tail": true, "hasBias": false },
+      "passes": [
+        {
+          "id": "main",
+          "name": "FastGelu.vec4Tail",
+          "shader": "elementwise-bias-gelu.wgsl.jinja",
+          "bindings": "scalar",
+          "dispatch": { "threads": "ceilDiv(numel(shapes.X), 4)", "workgroupSize": "workgroupSize" }
+        }
+      ]
+    },
+    {
+      "id": "scalar_bias",
+      "priority": 10,
+      "when": ["baseOk", "biasOk", "true"],
+      "constants": {
+        "vec4": false,
+        "vec4Tail": false,
+        "hasBias": true,
+        "hidden": "dim(shapes.X, ranks.X - 1) if dim(shapes.X, ranks.X - 1) > 0 else 1"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "FastGelu.scalarBias",
+          "shader": "elementwise-bias-gelu.wgsl.jinja",
+          "bindings": "scalarBias",
+          "dispatch": { "threads": "numel(shapes.X)", "workgroupSize": "workgroupSize" }
+        }
+      ]
+    },
+    {
+      "id": "scalar_no_bias",
+      "priority": 0,
+      "when": ["baseOk", "noBiasOk", "true"],
+      "constants": { "vec4": false, "vec4Tail": false, "hasBias": false },
+      "passes": [
+        {
+          "id": "main",
+          "name": "FastGelu.scalar",
+          "shader": "elementwise-bias-gelu.wgsl.jinja",
+          "bindings": "scalar",
+          "dispatch": { "threads": "numel(shapes.X)", "workgroupSize": "workgroupSize" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,18 @@

+{
+  "name": "com.microsoft.FastGelu",
+  "id": "_com_microsoft_fastgelu_webgpu_efaff66",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "m93bKUUMLGZxG/F5G6AFYMyDIKSE1Pk4hOyYNKrO+y0=",
+      "elementwise-bias-gelu.wgsl.jinja": "Eg8N2jJCMce+IsYNcCzuxvs8CSv7yTQXdorvv+c0L58=",
+      "manifest.json": "WckeD7rVTEJ1Lb6FDdpYC9n//UUq4+QWwd7vyAGc5m8=",
+      "test.json": "shIb34pVGz4K2kmzCn/yOmL6OuRbXD+UaRJ/PADEFtM="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/com.microsoft.FastGelu" }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,448 @@

+{
+  "op": "com.microsoft.FastGelu",
+  "fixtureArrays": {
+    "ort_float16_hidden8_with_bias_input_X": [0.8, -0.5, 0, 1, 1.3, 2.1, -0.2, 1.1, 0.5, 0.2, 0.3, -0.6, 3.1, 2.2, -1.1, 0]
+  },
+  "cases": [
+    {
+      "name": "dispatch_cliff_scalar_no_bias",
+      "inputs": {
+        "X": { "dtype": "float32", "shape": [16777, 1001], "data": { "kind": "linspace", "start": -3.0, "end": 3.0 } }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [16777, 1001], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_float32_with_bias",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/fastgelu_op_test.cc",
+        "test": "FastGeluTest.FastGeluWithBiasFloat32"
+      },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [1, 2, 4],
+          "data": { "kind": "values", "values": [0.8, -0.5, 0.0, 1.0, 0.5, 0.2, 0.3, -0.6] }
+        },
+        "bias": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [-0.5, 0.6, 1.2, 2.1] } }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float32",
+          "shape": [1, 2, 4],
+          "tolerance": 0.000001,
+          "data": {
+            "kind": "values",
+            "values": [0.18537092, 0.05398276, 1.0617028, 3.0973732, 0.0, 0.6304317, 1.3995715, 1.3995714]
+          }
+        }
+      }
+    },
+    {
+      "name": "ort_float32_without_bias",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/fastgelu_op_test.cc",
+        "test": "FastGeluTest.FastGeluWithoutBiasFloat32"
+      },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [1, 2, 4],
+          "data": { "kind": "values", "values": [0.8, -0.5, 0.0, 1.0, 0.5, 0.2, 0.3, -0.6] }
+        }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float32",
+          "shape": [1, 2, 4],
+          "tolerance": 0.000001,
+          "data": {
+            "kind": "values",
+            "values": [0.6304317, -0.154286, 0.0, 0.841192, 0.345714, 0.11585142, 0.18537092, -0.16458479]
+          }
+        }
+      }
+    },
+    {
+      "name": "ort_float32_zero_sequence_with_bias",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/fastgelu_op_test.cc",
+        "test": "FastGeluTest.FastGeluWithNullInput"
+      },
+      "inputs": {
+        "X": { "dtype": "float32", "shape": [1, 0, 4], "data": { "kind": "values", "values": [] } },
+        "bias": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [-0.5, 0.6, 1.2, 2.1] } }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 0, 4], "data": { "kind": "values", "values": [] } } }
+    },
+    {
+      "name": "f32_scalar_no_bias_zero_sequence",
+      "provenance": {
+        "notes": "Bias-free twin of ort_float32_zero_sequence_with_bias. An empty X fails the numel > 0 guard both vec4 routes carry, so the plain scalar no-bias kernel is the only one left."
+      },
+      "inputs": { "X": { "dtype": "float32", "shape": [1, 0, 4], "data": { "kind": "values", "values": [] } } },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 0, 4], "data": { "kind": "values", "values": [] } } }
+    },
+    {
+      "name": "empty_zero_hidden",
+      "provenance": {
+        "notes": "Zero-length last (bias-broadcast) axis: the bias vector itself is empty, so no kernel may size a binding from the hidden extent."
+      },
+      "inputs": {
+        "X": { "dtype": "float32", "shape": [2, 0], "data": { "kind": "values", "values": [] } },
+        "bias": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [2, 0], "data": { "kind": "values", "values": [] } } }
+    },
+    {
+      "name": "ort_float16_hidden2_with_bias",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/fastgelu_op_test.cc",
+        "test": "FastGeluTest.FastGeluWithBiasFloat16_2"
+      },
+      "inputs": {
+        "X": { "dtype": "float16", "shape": [1, 2, 2], "data": { "kind": "values", "values": [0.8, -0.5, 0.5, 0.2] } },
+        "bias": { "dtype": "float16", "shape": [2], "data": { "kind": "values", "values": [-0.5, 0.6] } }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float16",
+          "shape": [1, 2, 2],
+          "tolerance": 0.001,
+          "data": { "kind": "values", "values": [0.1851806640625, 0.054046630859375, 0.0, 0.63037109375] }
+        }
+      }
+    },
+    {
+      "name": "ort_float16_hidden2_without_bias",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/fastgelu_op_test.cc",
+        "test": "FastGeluTest.FastGeluWithoutBiasFloat16_2"
+      },
+      "inputs": {
+        "X": { "dtype": "float16", "shape": [1, 2, 2], "data": { "kind": "values", "values": [0.8, -0.5, 0.5, 0.2] } }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float16",
+          "shape": [1, 2, 2],
+          "tolerance": 0.001,
+          "data": { "kind": "values", "values": [0.63037109375, -0.154296875, 0.345703125, 0.11578369140625] }
+        }
+      }
+    },
+    {
+      "name": "ort_float16_hidden4_with_bias",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/fastgelu_op_test.cc",
+        "test": "FastGeluTest.FastGeluWithBiasFloat16_4"
+      },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [1, 2, 4],
+          "data": { "kind": "values", "values": [0.8, -0.5, 0.0, 1.0, 0.5, 0.2, 0.3, -0.6] }
+        },
+        "bias": { "dtype": "float16", "shape": [4], "data": { "kind": "values", "values": [-0.5, 0.6, 1.2, 2.1] } }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float16",
+          "shape": [1, 2, 4],
+          "tolerance": 0.001,
+          "data": {
+            "kind": "values",
+            "values": [0.1851806640625, 0.054046630859375, 1.0615234375, 3.09765625, 0.0, 0.63037109375, 1.3994140625, 1.3994140625]
+          }
+        }
+      }
+    },
+    {
+      "name": "ort_float16_hidden4_without_bias",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/fastgelu_op_test.cc",
+        "test": "FastGeluTest.FastGeluWithoutBiasFloat16_4"
+      },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [1, 2, 4],
+          "data": { "kind": "values", "values": [0.8, -0.5, 0.0, 1.0, 0.5, 0.2, 0.3, -0.6] }
+        }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float16",
+          "shape": [1, 2, 4],
+          "tolerance": 0.001,
+          "data": {
+            "kind": "values",
+            "values": [0.63037109375, -0.154296875, 0.0, 0.84130859375, 0.345703125, 0.1158447265625, 0.1854248046875, -0.16455078125]
+          }
+        }
+      }
+    },
+    {
+      "name": "ort_float16_hidden8_with_bias",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/fastgelu_op_test.cc",
+        "test": "FastGeluTest.FastGeluWithBiasFloat16_8"
+      },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [1, 2, 8],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_float16_hidden8_with_bias_input_X" } }
+        },
+        "bias": {
+          "dtype": "float16",
+          "shape": [8],
+          "data": { "kind": "values", "values": [-0.5, 0.6, 1.2, 2.1, 1.3, -1.0, 0.0, 3.1] }
+        }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float16",
+          "shape": [1, 2, 8],
+          "tolerance": 0.001,
+          "data": {
+            "kind": "values",
+            "values": [0.1851806640625, 0.054046630859375, 1.0615234375, 3.09765625, 2.587890625, 0.9501953125, -0.0841064453125, 4.19921875, 0.0, 0.63037109375, 1.3994140625, 1.3994140625, 4.3984375, 1.060546875, -0.1494140625, 3.09765625]
+          }
+        }
+      }
+    },
+    {
+      "name": "ort_float16_hidden8_without_bias",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/fastgelu_op_test.cc",
+        "test": "FastGeluTest.FastGeluWithoutBiasFloat16_8"
+      },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [1, 2, 8],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_float16_hidden8_with_bias_input_X" } }
+        }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float16",
+          "shape": [1, 2, 8],
+          "tolerance": 0.001,
+          "data": {
+            "kind": "values",
+            "values": [0.63037109375, -0.154296875, 0.0, 0.84130859375, 1.173828125, 2.0625, -0.0841064453125, 0.9501953125, 0.345703125, 0.1158447265625, 0.1854248046875, -0.16455078125, 3.09765625, 2.16796875, -0.1494140625, 0.0]
+          }
+        }
+      }
+    },
+    {
+      "name": "float32_extreme_saturation_edges",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/fastgelu_op_test.cc",
+        "notes": "Extra edge case for the tanh approximation: large negative values saturate to signed-zero-ish outputs while large positives pass through."
+      },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [1, 9],
+          "data": { "kind": "values", "values": [-20.0, -10.0, -5.0, -2.0, 0.0, 2.0, 5.0, 10.0, 20.0] }
+        }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float32",
+          "shape": [1, 9],
+          "tolerance": 0.000001,
+          "data": {
+            "kind": "values",
+            "values": [0.0, 0.0, -2.9802322e-7, -0.045402348, 0.0, 1.9545977, 4.9999995, 10.0, 20.0]
+          }
+        }
+      }
+    },
+    {
+      "name": "f32_subnormal_linear_region_no_bias_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/contrib_ops/fastgelu_op_test.cc",
+        "test": "FastGeluTest.FastGeluWithoutBiasFloat32",
+        "notes": "Near zero, tanh-approx FastGelu is approximately x/2; finite subnormal tails should survive the vec4 no-bias path."
+      },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [1e-40, -1e-40, 1e-38, -1e-38] }
+        }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float32",
+          "shape": [4],
+          "tolerance": 2e-45,
+          "data": {
+            "kind": "values",
+            "values": [4.99997305055738e-41, -4.99997305055738e-41, 4.999999675228202e-39, -4.999999675228202e-39]
+          }
+        }
+      }
+    },
+    {
+      "name": "f32_subnormal_linear_region_no_bias_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/contrib_ops/fastgelu_op_test.cc",
+        "test": "FastGeluTest.FastGeluWithoutBiasFloat32",
+        "notes": "Scalar-path companion for FastGelu's near-zero x/2 subnormal behavior."
+      },
+      "inputs": {
+        "X": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1e-40, -1e-40, 1e-38] } }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float32",
+          "shape": [3],
+          "tolerance": 2e-45,
+          "data": { "kind": "values", "values": [4.99997305055738e-41, -4.99997305055738e-41, 4.999999675228202e-39] }
+        }
+      }
+    },
+    {
+      "name": "f32_subnormal_linear_region_zero_bias_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/contrib_ops/fastgelu_op_test.cc",
+        "test": "FastGeluTest.FastGeluWithBiasFloat32",
+        "notes": "Bias-path companion: zero bias reduces FastGelu to the same near-zero x/2 behavior, but exercises the vec4 bias-broadcast kernel variant."
+      },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [1, 4],
+          "data": { "kind": "values", "values": [1e-40, -1e-40, 1e-38, -1e-38] }
+        },
+        "bias": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.0, 0.0, 0.0, 0.0] } }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float32",
+          "shape": [1, 4],
+          "tolerance": 2e-45,
+          "data": {
+            "kind": "values",
+            "values": [4.99997305055738e-41, -4.99997305055738e-41, 4.999999675228202e-39, -4.999999675228202e-39]
+          }
+        }
+      }
+    },
+    {
+      "name": "f32_subnormal_linear_region_zero_bias_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/contrib_ops/fastgelu_op_test.cc",
+        "test": "FastGeluTest.FastGeluWithBiasFloat32",
+        "notes": "Scalar bias-broadcast companion for FastGelu subnormal linear-region behavior with zero bias."
+      },
+      "inputs": {
+        "X": { "dtype": "float32", "shape": [1, 3], "data": { "kind": "values", "values": [1e-40, -1e-40, 1e-38] } },
+        "bias": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.0, 0.0, 0.0] } }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float32",
+          "shape": [1, 3],
+          "tolerance": 2e-45,
+          "data": { "kind": "values", "values": [4.99997305055738e-41, -4.99997305055738e-41, 4.999999675228202e-39] }
+        }
+      }
+    },
+    {
+      "name": "rank4_last_dim_bias_broadcast",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/fastgelu_op_test.cc",
+        "test": "FastGeluTest.FastGeluWithBiasFloat32",
+        "notes": "Compact rank-4 projection of ORT's last-dimension bias behavior."
+      },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [1, 1, 2, 3],
+          "data": { "kind": "values", "values": [-2.0, -1.0, 0.0, 1.0, 2.0, 3.0] }
+        },
+        "bias": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.5, -0.25, 1.0] } }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float32",
+          "shape": [1, 1, 2, 3],
+          "tolerance": 0.000001,
+          "data": { "kind": "values", "values": [-0.10042842, -0.1322858, 0.841192, 1.3995715, 1.6797954, 3.9999297] }
+        }
+      }
+    },
+    {
+      "name": "f32_scalar_bias_odd_hidden",
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [0.8, -0.5, 0.0, 0.5, 0.2, 0.3] }
+        },
+        "bias": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-0.5, 0.6, 1.2] } }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "tolerance": 0.000001,
+          "data": { "kind": "values", "values": [0.18537092, 0.05398275, 1.06170277, 0.0, 0.63043169, 1.39957158] }
+        }
+      }
+    },
+    {
+      "name": "f32_scalar_no_bias_numel_not_div4",
+      "inputs": {
+        "X": { "dtype": "float32", "shape": [5], "data": { "kind": "values", "values": [0.8, -0.5, 0.0, 1.0, -2.0] } }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float32",
+          "shape": [5],
+          "tolerance": 0.000001,
+          "data": { "kind": "values", "values": [0.63043169, -0.15428599, 0.0, 0.84119199, -0.04540231] }
+        }
+      }
+    },
+    {
+      "name": "f32_tanh_vs_erf_distinguisher",
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [1, 5],
+          "data": { "kind": "values", "values": [-2.0, -1.0, 0.0, 1.0, 2.0] }
+        }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float32",
+          "shape": [1, 5],
+          "tolerance": 0.0001,
+          "data": { "kind": "values", "values": [-0.04540231, -0.15880801, 0.0, 0.84119199, 1.95459769] }
+        }
+      }
+    }
+  ]
+}