sync 2e7068faf55e

Browse files

Files changed (6) hide show

README.md +60 -0
build/webgpu/bench.json +73 -0
build/webgpu/bias-split-gelu.wgsl.jinja +116 -0
build/webgpu/manifest.json +194 -0
build/webgpu/metadata.json +18 -0
build/webgpu/test.json +266 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# com.microsoft.BiasSplitGelu
+`com.microsoft`  ·  ONNX Runtime contrib operator  ·  contrib since_version 1
+## Description
+Adds `bias` to `X`, splits the last dimension in half, then multiplies the left half elementwise by the GELU activation of the right half, producing an output with half the hidden dimension. Used as a fused activation in diffusion models.
+See the [ONNX Runtime `BiasSplitGelu` contrib-operator spec](https://github.com/microsoft/onnxruntime/blob/main/docs/ContribOperators.md#com.microsoft.BiasSplitGelu) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `X` | `X` | `T` | `3` | — | Input tensor of shape `(N, S, D)`, where `N` is the batch size, `S` is the spatial size, and `D` is the hidden dimension. | required |
+| `bias` | `bias` | `T` | `1` | — | 1-D bias tensor of length D, matching the hidden dimension of X. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `Y` | `Y` | `T` | `3` | derived; see description | Output tensor of shape `(N, S, D/2)` after the split-GELU activation. | 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
+- [`bias-split-gelu.wgsl.jinja`](build/webgpu/bias-split-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.BiasSplitGelu", { version: 1 });
+const { Y } = await kernel({
+  X: { data: XData, shape: [1, 2, 4] },
+  bias: { data: biasData, shape: [4] },
+});
+```

build/webgpu/bench.json ADDED Viewed

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

+{
+  "op": "com.microsoft.BiasSplitGelu",
+  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
+  "cases": [
+    {
+      "name": "biassplitgelu-f32-4096x5120",
+      "preset": "smoke",
+      "vars": { "dtype": "float32" },
+      "inputs": {
+        "X": { "shape": [1, 4096, 5120], "dtype": "float32", "dist": "normal", "seed": 610, "scale": 2 },
+        "bias": { "shape": [5120], "dtype": "float32", "dist": "normal", "seed": 611, "scale": 1 }
+      },
+      "outputs": { "Y": { "shape": [1, 4096, 2560], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "4096 * 5120 * dtypeBytes(args.dtype) + 4096 * 2560 * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "biassplitgelu-f16-4096x5120",
+      "preset": "model",
+      "vars": { "dtype": "float16" },
+      "inputs": {
+        "X": { "shape": [1, 4096, 5120], "dtype": "float16", "dist": "normal", "seed": 612, "scale": 2 },
+        "bias": { "shape": [5120], "dtype": "float16", "dist": "normal", "seed": 613, "scale": 1 }
+      },
+      "outputs": { "Y": { "shape": [1, 4096, 2560], "dtype": "float16" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "4096 * 5120 * dtypeBytes(args.dtype) + 4096 * 2560 * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "biassplitgelu-f16-compact-vec4-validation-gap",
+      "preset": "edge",
+      "vars": { "dtype": "float16" },
+      "inputs": {
+        "X": { "shape": [1, 1, 8], "dtype": "float16", "dist": "normal", "seed": 614, "scale": 2 },
+        "bias": { "shape": [8], "dtype": "float16", "dist": "normal", "seed": 615, "scale": 1 }
+      },
+      "outputs": { "Y": { "shape": [1, 1, 4], "dtype": "float16" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "8 * dtypeBytes(args.dtype) + 4 * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "biassplitgelu-f32-scalar-cliff-4096x5116",
+      "preset": "stress",
+      "vars": { "dtype": "float32" },
+      "inputs": {
+        "X": { "shape": [1, 4096, 5116], "dtype": "float32", "dist": "normal", "seed": 620, "scale": 2 },
+        "bias": { "shape": [5116], "dtype": "float32", "dist": "normal", "seed": 621, "scale": 1 }
+      },
+      "outputs": { "Y": { "shape": [1, 4096, 2558], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "4096 * 5116 * dtypeBytes(args.dtype) + 4096 * 2558 * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    }
+  ]
+}

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

	@@ -0,0 +1,116 @@

+{% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
+{% if note == "dispatch-limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
+{% elif note == "limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit.
+{% elif note == "device-axis" %}
+  // The flat dispatch is folded across x/y at the device's per-axis workgroup
+  // limit; gid.y carries the high portion of the output index.
+{% elif note == "vec4-limit" %}
+  // 2D-folded flat vec4 index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills into y).
+{% elif note == "element-limit" %}
+  // 2D-folded flat element index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit.
+{% elif note == "dispatch" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit.
+{% endif %}
+{% if bound == "" %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+{%- elif guardInline %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) { return; }
+{%- else %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) {
+    return;
+  }
+{%- endif %}
+{% endmacro %}
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+// com.microsoft.BiasSplitGelu GEGLU gate:
+//   The innermost (hidden) axis of (X + bias) has size 2*D. Split it into
+//   value = [.., 0:D] and gate = [.., D:2D], then
+//   Y[.., d] = value[.., d] * gelu_erf(gate[.., d])
+//   gelu_erf(g) = 0.5 * g * (1 + erf(g / sqrt2)).
+// bias is a rank-1 vector [2*D] broadcast over the hidden axis. The output
+// hidden axis is D (= HIDDEN below); the input hidden axis is 2*D (TWICE).
+// erf_approx uses an overflow-safe rational approximation.
+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;
+}
+fn gelu_value(v: f32) -> f32 {
+  return 0.5 * v * (1.0 + erf_approx(v * 0.7071067811865476));
+}
+const HIDDEN: u32 = {{ hidden }}u;   // D: output hidden axis
+const TWICE: u32 = {{ twice }}u;     // 2*D: input hidden axis
+{% if vec4Tail %}
+const GROUPS_PER_ROW: u32 = (HIDDEN + 3u) / 4u;
+{% endif %}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+{{ flat_index_2d() }}
+{% if vec4Tail %}
+  let row = i / GROUPS_PER_ROW;
+  let col = (i % GROUPS_PER_ROW) * 4u;
+{% for lane in range(4) %}
+  if (col + {{ lane }}u < HIDDEN) {
+    let vidx{{ lane }} = row * TWICE + col + {{ lane }}u;
+    let gidx{{ lane }} = vidx{{ lane }} + HIDDEN;
+    let val{{ lane }} = f32(x[vidx{{ lane }}]) + f32(bias[col + {{ lane }}u]);
+    let gate{{ lane }} = f32(x[gidx{{ lane }}]) + f32(bias[col + HIDDEN + {{ lane }}u]);
+    y[row * HIDDEN + col + {{ lane }}u] = {{ scalar }}(val{{ lane }} * gelu_value(gate{{ lane }}));
+  }
+{% endfor %}
+{% elif vec4 %}
+  // Each output vec4 covers 4 consecutive elements within one row of the D
+  // axis (D % 4 == 0 guarantees the group never crosses the hidden boundary).
+  let base = i * 4u;
+  let row = base / HIDDEN;
+  let col = base % HIDDEN;
+  let vbase = row * TWICE + col;          // value half: [row, col .. col+3]
+  let gbase = vbase + HIDDEN;             // gate half:  [row, D+col .. D+col+3]
+  // f32(...) each lane: X may be f16 storage, and WGSL has no implicit f16->f32
+  // conversion inside a vec4<f32> constructor (it is a no-op when X is f32).
+  let valv = vec4<f32>(f32(x[vbase]), f32(x[vbase + 1u]), f32(x[vbase + 2u]), f32(x[vbase + 3u]))
+           + vec4<f32>(f32(bias[col]), f32(bias[col + 1u]), f32(bias[col + 2u]), f32(bias[col + 3u]));
+  let gatev = vec4<f32>(f32(x[gbase]), f32(x[gbase + 1u]), f32(x[gbase + 2u]), f32(x[gbase + 3u]))
+            + vec4<f32>(f32(bias[col + HIDDEN]), f32(bias[col + HIDDEN + 1u]), f32(bias[col + HIDDEN + 2u]), f32(bias[col + HIDDEN + 3u]));
+  y[i] = vec4<{{ scalar }}>(vec4<f32>(
+    valv.x * gelu_value(gatev.x),
+    valv.y * gelu_value(gatev.y),
+    valv.z * gelu_value(gatev.z),
+    valv.w * gelu_value(gatev.w)));
+{% else %}
+  let row = i / HIDDEN;
+  let col = i % HIDDEN;
+  let vidx = row * TWICE + col;           // value element
+  let gidx = vidx + HIDDEN;               // gate element
+  let val = f32(x[vidx]) + f32(bias[col]);
+  let gate = f32(x[gidx]) + f32(bias[col + HIDDEN]);
+  y[i] = {{ scalar }}(val * gelu_value(gate));
+{% endif %}
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,194 @@

+{
+  "domain": "com.microsoft",
+  "name": "BiasSplitGelu",
+  "sinceVersion": 1,
+  "description": "Adds `bias` to `X`, splits the last dimension in half, then multiplies the left half elementwise by the GELU activation of the right half, producing an output with half the hidden dimension. Used as a fused activation in diffusion models.",
+  "inputs": [
+    {
+      "role": "X",
+      "dtype": "T",
+      "rank": 3,
+      "description": "Input tensor of shape `(N, S, D)`, where `N` is the batch size, `S` is the spatial size, and `D` is the hidden dimension."
+    },
+    {
+      "role": "bias",
+      "dtype": "T",
+      "rank": 1,
+      "description": "1-D bias tensor of length D, matching the hidden dimension of X."
+    }
+  ],
+  "outputs": [
+    {
+      "role": "Y",
+      "dtype": "T",
+      "rank": 3,
+      "shape": "[dim(shapes.X, 0), dim(shapes.X, 1), dim(shapes.X, 2) / 2]",
+      "description": "Output tensor of shape `(N, S, D/2)` after the split-GELU activation."
+    }
+  ],
+  "typeConstraints": { "T": ["float32", "float16"] },
+  "args": {
+    "X": { "kind": "tensor", "semantic": "X", "role": "input" },
+    "bias": { "kind": "tensor", "semantic": "bias", "role": "input" },
+    "Y": { "kind": "tensor", "semantic": "Y", "role": "output" }
+  },
+  "tunables": { "WORKGROUP_SIZE": 64 },
+  "variants": [
+    {
+      "id": "vec4",
+      "priority": 30,
+      "when": ["ranks.X == 3", "ranks.Y == 3", "dim(shapes.Y, 0) == dim(shapes.X, 0)", "dim(shapes.Y, 1) == dim(shapes.X, 1)", "dim(shapes.X, 2) % 2 == 0", "dim(shapes.Y, 2) == dim(shapes.X, 2) / 2", "f16Ok(dtypes.T)", "ranks.bias == 1", "dim(shapes.bias, 0) == dim(shapes.X, 2)", "numel(shapes.Y) > 0", "numel(shapes.Y) % 4 == 0", "dim(shapes.Y, 2) % 4 == 0"],
+      "constants": {
+        "scalar": "dtypes.T",
+        "usesF16": "dtypes.T == \"f16\"",
+        "vec4": true,
+        "vec4Tail": false,
+        "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
+        "hidden": "dim(shapes.X, ranks.X - 1) / 2 if dim(shapes.X, ranks.X - 1) > 0 else 1",
+        "twice": "dim(shapes.X, ranks.X - 1) if dim(shapes.X, ranks.X - 1) > 0 else 2"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "BiasSplitGelu.vec4",
+          "shader": "bias-split-gelu.wgsl.jinja",
+          "bindings": [
+            {
+              "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",
+              "length": "$twice"
+            },
+            {
+              "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.Y) / 4" }]
+              }
+            }
+          ],
+          "dispatch": { "threads": "numel(shapes.Y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "vec4_tail",
+      "priority": 20,
+      "when": ["ranks.X == 3", "ranks.Y == 3", "dim(shapes.Y, 0) == dim(shapes.X, 0)", "dim(shapes.Y, 1) == dim(shapes.X, 1)", "dim(shapes.X, 2) % 2 == 0", "dim(shapes.Y, 2) == dim(shapes.X, 2) / 2", "f16Ok(dtypes.T)", "ranks.bias == 1", "dim(shapes.bias, 0) == dim(shapes.X, 2)", "numel(shapes.Y) > 0"],
+      "constants": {
+        "scalar": "dtypes.T",
+        "usesF16": "dtypes.T == \"f16\"",
+        "vec4": false,
+        "vec4Tail": true,
+        "hidden": "dim(shapes.X, ranks.X - 1) / 2 if dim(shapes.X, ranks.X - 1) > 0 else 1",
+        "twice": "dim(shapes.X, ranks.X - 1) if dim(shapes.X, ranks.X - 1) > 0 else 2"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "BiasSplitGelu.vec4Tail",
+          "shader": "bias-split-gelu.wgsl.jinja",
+          "bindings": [
+            {
+              "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",
+              "length": "$twice"
+            },
+            { "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.Y) / (dim(shapes.X, ranks.X - 1) / 2)) * ceilDiv(dim(shapes.X, ranks.X - 1) / 2, 4)"
+                  }
+                ]
+              }
+            }
+          ],
+          "dispatch": {
+            "threads": "(numel(shapes.Y) / (dim(shapes.X, ranks.X - 1) / 2)) * ceilDiv(dim(shapes.X, ranks.X - 1) / 2, 4)",
+            "workgroupSize": "tunables.WORKGROUP_SIZE"
+          }
+        }
+      ]
+    },
+    {
+      "id": "scalar",
+      "priority": 0,
+      "when": ["ranks.X == 3", "ranks.Y == 3", "dim(shapes.Y, 0) == dim(shapes.X, 0)", "dim(shapes.Y, 1) == dim(shapes.X, 1)", "dim(shapes.X, 2) % 2 == 0", "dim(shapes.Y, 2) == dim(shapes.X, 2) / 2", "f16Ok(dtypes.T)", "ranks.bias == 1", "dim(shapes.bias, 0) == dim(shapes.X, 2)", "true"],
+      "constants": {
+        "scalar": "dtypes.T",
+        "usesF16": "dtypes.T == \"f16\"",
+        "vec4": false,
+        "vec4Tail": false,
+        "hidden": "dim(shapes.X, ranks.X - 1) / 2 if dim(shapes.X, ranks.X - 1) > 0 else 1",
+        "twice": "dim(shapes.X, ranks.X - 1) if dim(shapes.X, ranks.X - 1) > 0 else 2"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "BiasSplitGelu.scalar",
+          "shader": "bias-split-gelu.wgsl.jinja",
+          "bindings": [
+            {
+              "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.Y)" }] }
+            }
+          ],
+          "dispatch": { "threads": "numel(shapes.Y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

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

+{
+  "name": "com.microsoft.BiasSplitGelu",
+  "id": "_com_microsoft_biassplitgelu_webgpu_e56edec",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "oYxkYz1Ukonqiav8OJSz76L0uidTkGg049tpmfnC0Ms=",
+      "bias-split-gelu.wgsl.jinja": "YGnskVbzCTa0y/SN+O6Wq0Yp8SWpUav6Q930/EHEvsw=",
+      "manifest.json": "i40t5ZCHmqMoQrSWSWZlC3X/ERSmyoa5z4PkrzsWP/w=",
+      "test.json": "JVElcOkuNl49j2nfKWBgotONhDTNVT7ZbWTDW+5Of/Q="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/com.microsoft.BiasSplitGelu" }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,266 @@

+{
+  "op": "com.microsoft.BiasSplitGelu",
+  "cases": [
+    {
+      "name": "dispatch_cliff_vec4_over_16M",
+      "requires": { "limits": { "maxBufferSize": 536862752, "maxStorageBufferBindingSize": 536862752 } },
+      "provenance": {
+        "notes": "Its largest tensor is 536862752 bytes, so the case needs an adapter whose maxBufferSize and maxStorageBufferBindingSize both reach it — declared, because the WebGPU guaranteed minimums (256 MiB / 128 MiB) do not, and a device at them must report the case inapplicable rather than fail allocating it."
+      },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [1, 16776961, 8],
+          "data": { "kind": "linspace", "start": -1.0, "end": 1.0 }
+        },
+        "bias": { "dtype": "float32", "shape": [8], "data": { "kind": "linspace", "start": -0.5, "end": 0.5 } }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 16776961, 4], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_split_last_dim_then_gelu_gate",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/bias_split_gelu_op_test.cc",
+        "test": "BiasSplitGeluTest.BiasSplitGeluTest_HiddenSize_2560",
+        "notes": "Small deterministic projection of ORT's fused bias+split+Gelu gate pattern."
+      },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [1, 2, 4],
+          "data": { "kind": "values", "values": [1.0, 2.0, 0.0, 1.0, -1.0, 0.5, 2.0, -2.0] }
+        },
+        "bias": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.5, -0.5, 1.0, -1.0] } }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float32",
+          "shape": [1, 2, 2],
+          "tolerance": 0.000001,
+          "relTolerance": 0.000001,
+          "data": { "kind": "values", "values": [1.262017119103, 0.0, -1.497975152953, 0.0] }
+        }
+      }
+    },
+    {
+      "name": "ort_projection_sequence_len1",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/bias_split_gelu_op_test.cc",
+        "test": "BiasSplitGeluTest.BiasSplitGeluTest_HiddenSize_5120",
+        "notes": "Compact deterministic projection of ORT's batch=2, sequence_length=1 shape coverage."
+      },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [2, 1, 6],
+          "data": { "kind": "values", "values": [1.0, -1.0, 0.5, 2.0, -2.0, 0.0, -0.25, 0.75, -1.5, 1.5, 0.25, -0.5] }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [0.25, -0.5, 0.75, -1.0, 0.5, -0.25] }
+        }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float32",
+          "shape": [2, 1, 3],
+          "tolerance": 0.000001,
+          "relTolerance": 0.000001,
+          "data": {
+            "kind": "values",
+            "values": [1.05168092, 0.1503162697, -0.1254042904, 0.0, 0.1450073842, 0.1274778475]
+          }
+        }
+      }
+    },
+    {
+      "name": "ort_projection_batch1_sequence2",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/bias_split_gelu_op_test.cc",
+        "test": "BiasSplitGeluTest.BiasSplitGeluTest_HiddenSize_10240",
+        "notes": "Compact deterministic projection of ORT's batch=1, sequence_length=2 shape coverage."
+      },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [1, 2, 6],
+          "data": { "kind": "values", "values": [1.0, 2.0, -1.0, 0.0, 0.5, -0.5, -2.0, 1.5, 0.25, 2.5, -1.0, 0.0] }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [0.25, -0.5, 0.75, -1.0, 0.5, -0.25] }
+        }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float32",
+          "shape": [1, 2, 3],
+          "tolerance": 0.000001,
+          "relTolerance": 0.000001,
+          "data": {
+            "kind": "values",
+            "values": [-0.1983190798, 1.262017104, 0.04249261584, -2.449631019, -0.1542687686, -0.1003234323]
+          }
+        }
+      }
+    },
+    {
+      "name": "ort_projection_batch2_sequence3_vec4",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/bias_split_gelu_op_test.cc",
+        "test": "BiasSplitGeluTest.BiasSplitGeluTest_HiddenSize_6144",
+        "notes": "Compact deterministic projection of ORT's batch=2, sequence_length=3 shape coverage with output hidden divisible by 4."
+      },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [2, 3, 8],
+          "data": {
+            "kind": "values",
+            "values": [1.0, -1.0, 0.5, -0.5, 2.0, -2.0, 0.25, -0.25, -1.5, 1.5, 0.0, -0.75, 0.75, -0.5, 1.25, -1.25, 3.0, -3.0, 0.5, -0.5, -2.0, 2.0, 1.0, -1.0, 0.0, 1.0, -1.0, 0.5, -0.5, 2.5, -2.5, 1.5, 0.125, -0.125, 4.0, -4.0, 0.75, -0.75, 1.5, -1.5, 2.25, -2.25, 0.333, -0.333, 5.0, -5.0, 1.25, -1.25]
+          }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0.25, -0.5, 0.75, -1.0, 0.5, -0.25, 1.5, -1.5] }
+        }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float32",
+          "shape": [2, 3, 4],
+          "tolerance": 0.000001,
+          "relTolerance": 0.000001,
+          "data": {
+            "kind": "values",
+            "values": [3.1055946350097656, 0.04125746339559555, 2.0998706817626953, 0.10515517741441727, -1.3974220752716064, -0.16997046768665314, 2.05635404586792, 0.014340377412736416, -0.3256852626800537, -5.879637718200684, 3.1055946350097656, 0.023286299780011177, 0.0, 1.1112475395202637, 0.03966381773352623, 0.0, 0.41922664642333984, 0.09915953874588013, 14.23076343536377, 0.020249508321285248, 13.75, 0.0000011001571920132847, 2.9693753719329834, 0.010923270136117935]
+          }
+        }
+      }
+    },
+    {
+      "name": "split_gelu_d4_vec4_path",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/bias_split_gelu_op_test.cc",
+        "notes": "D=4 (input hidden 8) so the output hidden axis is a multiple of 4 and the vec4 manifest variant is exercised. Expected values from the same erf-form GEGLU gate computed at f64 then f32-rounded."
+      },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [1, 1, 8],
+          "data": { "kind": "values", "values": [0.5, -1.0, 2.0, -0.5, 1.5, 0.0, -2.0, 1.0] }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0.25, 0.5, -0.5, 1.0, -1.0, 0.5, 2.0, -0.25] }
+        }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float32",
+          "shape": [1, 1, 4],
+          "tolerance": 0.000001,
+          "relTolerance": 0.000001,
+          "data": { "kind": "values", "values": [0.2592984235214977, -0.17286561568099845, 0.0, 0.2900147683230033] }
+        }
+      }
+    },
+    {
+      "name": "empty_zero_dim",
+      "provenance": {
+        "notes": "Zero-length batch axis of X (empty input). ONNX treats empty tensors as first-class; output is the same broadcast shape with the 0 batch ⇒ EMPTY output. bias kept at its normal rank-1 non-zero shape."
+      },
+      "inputs": {
+        "X": { "dtype": "float32", "shape": [0, 2, 4], "data": { "kind": "values", "values": [] } },
+        "bias": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.5, -0.5, 1.0, -1.0] } }
+      },
+      "outputs": {
+        "Y": { "dtype": "float32", "shape": [0, 2, 2], "tolerance": 0, "data": { "kind": "values", "values": [] } }
+      }
+    },
+    {
+      "name": "empty_zero_dim_f16",
+      "provenance": { "notes": "f16 variant of the empty zero-batch-axis case." },
+      "inputs": {
+        "X": { "dtype": "float16", "shape": [0, 2, 4], "data": { "kind": "values", "values": [] } },
+        "bias": { "dtype": "float16", "shape": [4], "data": { "kind": "values", "values": [0.5, -0.5, 1.0, -1.0] } }
+      },
+      "outputs": {
+        "Y": { "dtype": "float16", "shape": [0, 2, 2], "tolerance": 0, "data": { "kind": "values", "values": [] } }
+      }
+    },
+    {
+      "name": "empty_zero_hidden",
+      "provenance": {
+        "notes": "Zero-length hidden axis: bias stays 2x the (zero) half-width, so both the bias vector and the split halves are empty and no kernel may size a binding from the hidden extent."
+      },
+      "inputs": {
+        "X": { "dtype": "float32", "shape": [1, 2, 0], "data": { "kind": "values", "values": [] } },
+        "bias": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": {
+        "Y": { "dtype": "float32", "shape": [1, 2, 0], "tolerance": 0, "data": { "kind": "values", "values": [] } }
+      }
+    },
+    {
+      "name": "f16_vec4_shader_validation",
+      "provenance": {
+        "notes": "Non-empty f16 D=4 case forces the vec4 variant; the empty f16 case cannot catch shader construction errors because it returns before compiling meaningful lane loads."
+      },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [1, 1, 8],
+          "data": { "kind": "values", "values": [0.5, -1.0, 2.0, -0.5, 1.5, 0.0, -2.0, 1.0] }
+        },
+        "bias": {
+          "dtype": "float16",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0.25, 0.5, -0.5, 1.0, -1.0, 0.5, 2.0, -0.25] }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 1, 4], "tolerance": 0.01 } }
+    },
+    {
+      "name": "scalar_path_dispatch_fold_over_16M",
+      "provenance": {
+        "notes": "D=6 (twice=12, 6%4!=0) forces the scalar variant; numel(Y)=16776966 > 65535*256 forces the scalar-branch 2D dispatch fold (gid.y high bits), an untested path (dispatch_cliff_vec4_over_16M covers only the vec4 branch). Validates the i=gid.x+gid.y*nwg.x*WG reconstruction and i>=count tail guard on the last rows."
+      },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [1, 2796161, 12],
+          "data": { "kind": "linspace", "start": -1.0, "end": 1.0 }
+        },
+        "bias": { "dtype": "float32", "shape": [12], "data": { "kind": "linspace", "start": -0.5, "end": 0.5 } }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 2796161, 6], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "scalar_path_odd_hidden_d6",
+      "provenance": {
+        "notes": "twice=12, D=6 (6%4!=0) routes to the scalar variant; small case verifying scalar split-index math (value half [0:6], gate half [6:12]) at a hidden size not covered by existing D=3 scalar cases. Reference supplies truth."
+      },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [1, 2, 12],
+          "data": {
+            "kind": "values",
+            "values": [1.0, -1.0, 0.5, -0.5, 2.0, -2.0, 0.25, -0.25, 1.5, -1.5, 0.75, -0.75, -2.0, 2.0, 1.0, -1.0, 0.5, -0.5, 3.0, -3.0, 0.125, -0.125, 0.0, 1.0]
+          }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [12],
+          "data": { "kind": "values", "values": [0.25, -0.5, 0.75, -1.0, 0.5, -0.25, 1.5, -1.5, 0.1, -0.1, 0.6, -0.6] }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 2, 6], "tolerance": 0.000001, "relTolerance": 0.000001 } }
+    }
+  ]
+}