sync 2e7068faf55e

Browse files

Files changed (7) hide show

README.md +74 -0
build/webgpu/bench.json +246 -0
build/webgpu/manifest.json +1108 -0
build/webgpu/metadata.json +19 -0
build/webgpu/norm-skip-row-vec4.wgsl.jinja +145 -0
build/webgpu/norm-skip-row.wgsl.jinja +226 -0
build/webgpu/test.json +939 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# com.microsoft.SkipLayerNormalization
+`com.microsoft`  ·  ONNX Runtime contrib operator  ·  contrib since_version 1
+## Description
+Fuses skip addition with layer normalization. The rank-3 standard surface currently supports float32, required `beta`, no `bias` or residual output, hidden sizes divisible by four, and exact or documented broadcast skip shapes. The provider's rank-2 extension supports float32 output-only with optional `beta`, or `beta` with optional `bias` when emitting the residual; its float16 path requires `beta`, `bias`, a residual output, and four-wide hidden size. Other combinations, bfloat16, and training statistics are not implemented.
+See the [ONNX Runtime `SkipLayerNormalization` contrib-operator spec](https://github.com/microsoft/onnxruntime/blob/main/docs/ContribOperators.md#com.microsoft.SkipLayerNormalization) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `input` | `inputT` | `T` | — | — | Primary input normalized over the final hidden-size axis. Rank 3 is the public shape; rank 2 is an ONNX Runtime provider extension supported by this package. | required |
+| `skip` | `skipT` | `T` | — | — | Residual tensor. For rank-3 input it is exact shape, `(1, sequence_length, hidden_size)`, or `(sequence_length, hidden_size)`; rank-2 input requires exact shape. | required |
+| `gamma` | `gammaT` | `T` | `1` | — | Layer-norm scale weights of shape `(hidden_size)`. | required |
+| `beta` | `betaT` | `T` | `1` | — | Layer-norm bias weights of shape `(hidden_size)`. | optional |
+| `bias` | `biasT` | `T` | `1` | — | Optional additive bias of shape `(hidden_size)` added to `input + skip` before normalization. | optional |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `output` | `outputT` | `T` | same as `input` | same as `input` | Normalized output tensor with the same shape as `input`. | required |
+| `input_skip_bias_sum` | `residualT` | `T` | same as `input` | same as `input` | Sum of `input`, `skip`, and `bias` (when present) before normalization, with the same shape as `input`. | optional |
+## Attributes
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `epsilon` | `9.999999960041972e-13` | Non-negative epsilon added to the variance before taking the square root. |
+## 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
+- [`norm-skip-row-vec4.wgsl.jinja`](build/webgpu/norm-skip-row-vec4.wgsl.jinja)
+- [`norm-skip-row.wgsl.jinja`](build/webgpu/norm-skip-row.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.SkipLayerNormalization", { version: 1 });
+const { outputT } = await kernel({
+  inputT: { data: inputTData, shape: [2, 4] },
+  skipT: { data: skipTData, shape: [2, 4] },
+  gammaT: { data: gammaTData, shape: [4] },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,246 @@

+{
+  "op": "com.microsoft.SkipLayerNormalization",
+  "cases": [
+    {
+      "name": "skip-layernorm-f32-256x128",
+      "preset": "smoke",
+      "vars": { "rows": 256, "hidden": 128 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "shape": [256, 128], "dtype": "float32", "dist": "normal", "seed": 213, "scale": 0.2 },
+        "skipT": { "shape": [256, 128], "dtype": "float32", "dist": "normal", "seed": 214, "scale": 0.2 },
+        "gammaT": { "shape": [128], "dtype": "float32", "dist": "uniform", "seed": 215, "scale": 0.1, "offset": 1 },
+        "betaT": { "shape": [128], "dtype": "float32", "dist": "normal", "seed": 216, "scale": 0.1 }
+      },
+      "outputs": {
+        "outputT": { "shape": [256, 128], "dtype": "float32" },
+        "residualT": { "shape": [256, 128], "dtype": "float32" }
+      },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "4 * (4 * args.rows * args.hidden + 2 * args.hidden)" }]
+      }
+    },
+    {
+      "name": "skip-layernorm-f32-no-beta-output-256x128",
+      "preset": "smoke",
+      "vars": { "rows": 256, "hidden": 128 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "shape": [256, 128], "dtype": "float32", "dist": "normal", "seed": 219, "scale": 0.2 },
+        "skipT": { "shape": [256, 128], "dtype": "float32", "dist": "normal", "seed": 220, "scale": 0.2 },
+        "gammaT": { "shape": [128], "dtype": "float32", "dist": "uniform", "seed": 221, "scale": 0.1, "offset": 1 }
+      },
+      "outputs": { "outputT": { "shape": [256, 128], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (3 * args.rows * args.hidden + args.hidden)" }] }
+    },
+    {
+      "name": "skip-layernorm-f32-bias-256x128",
+      "preset": "smoke",
+      "vars": { "rows": 256, "hidden": 128 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "shape": [256, 128], "dtype": "float32", "dist": "normal", "seed": 214, "scale": 0.2 },
+        "skipT": { "shape": [256, 128], "dtype": "float32", "dist": "normal", "seed": 215, "scale": 0.2 },
+        "gammaT": { "shape": [128], "dtype": "float32", "dist": "uniform", "seed": 216, "scale": 0.1, "offset": 1 },
+        "betaT": { "shape": [128], "dtype": "float32", "dist": "normal", "seed": 217, "scale": 0.1 },
+        "biasT": { "shape": [128], "dtype": "float32", "dist": "normal", "seed": 218, "scale": 0.05 }
+      },
+      "outputs": {
+        "outputT": { "shape": [256, 128], "dtype": "float32" },
+        "residualT": { "shape": [256, 128], "dtype": "float32" }
+      },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "4 * (4 * args.rows * args.hidden + 3 * args.hidden)" }]
+      }
+    },
+    {
+      "name": "skip-layernorm-f32-4096x768",
+      "preset": "model",
+      "vars": { "rows": 4096, "hidden": 768 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "shape": [4096, 768], "dtype": "float32", "dist": "normal", "seed": 313, "scale": 0.2 },
+        "skipT": { "shape": [4096, 768], "dtype": "float32", "dist": "normal", "seed": 314, "scale": 0.2 },
+        "gammaT": { "shape": [768], "dtype": "float32", "dist": "uniform", "seed": 315, "scale": 0.1, "offset": 1 },
+        "betaT": { "shape": [768], "dtype": "float32", "dist": "normal", "seed": 316, "scale": 0.1 }
+      },
+      "outputs": {
+        "outputT": { "shape": [4096, 768], "dtype": "float32" },
+        "residualT": { "shape": [4096, 768], "dtype": "float32" }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (4 * args.rows * args.hidden + 2 * args.hidden)" }] }
+    },
+    {
+      "name": "skip-layernorm-f32-no-beta-output-4096x768",
+      "preset": "model",
+      "vars": { "rows": 4096, "hidden": 768 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "shape": [4096, 768], "dtype": "float32", "dist": "normal", "seed": 319, "scale": 0.2 },
+        "skipT": { "shape": [4096, 768], "dtype": "float32", "dist": "normal", "seed": 320, "scale": 0.2 },
+        "gammaT": { "shape": [768], "dtype": "float32", "dist": "uniform", "seed": 321, "scale": 0.1, "offset": 1 }
+      },
+      "outputs": { "outputT": { "shape": [4096, 768], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (3 * args.rows * args.hidden + args.hidden)" }] }
+    },
+    {
+      "name": "skip-layernorm-f32-bias-4096x768",
+      "preset": "model",
+      "vars": { "rows": 4096, "hidden": 768 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "shape": [4096, 768], "dtype": "float32", "dist": "normal", "seed": 314, "scale": 0.2 },
+        "skipT": { "shape": [4096, 768], "dtype": "float32", "dist": "normal", "seed": 315, "scale": 0.2 },
+        "gammaT": { "shape": [768], "dtype": "float32", "dist": "uniform", "seed": 316, "scale": 0.1, "offset": 1 },
+        "betaT": { "shape": [768], "dtype": "float32", "dist": "normal", "seed": 317, "scale": 0.1 },
+        "biasT": { "shape": [768], "dtype": "float32", "dist": "normal", "seed": 318, "scale": 0.05 }
+      },
+      "outputs": {
+        "outputT": { "shape": [4096, 768], "dtype": "float32" },
+        "residualT": { "shape": [4096, 768], "dtype": "float32" }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (4 * args.rows * args.hidden + 3 * args.hidden)" }] }
+    },
+    {
+      "name": "skip-layernorm-f32-bias-output-only-4096x768",
+      "preset": "model",
+      "provenance": {
+        "notes": "Ordinary inference shape with bias present and all three optional outputs omitted. It uses six storage buffers and writes only the primary output, isolating the output-only path from the full-output sibling at the same geometry."
+      },
+      "vars": { "rows": 4096, "hidden": 768 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "shape": [4096, 768], "dtype": "float32", "dist": "normal", "seed": 314, "scale": 0.2 },
+        "skipT": { "shape": [4096, 768], "dtype": "float32", "dist": "normal", "seed": 315, "scale": 0.2 },
+        "gammaT": { "shape": [768], "dtype": "float32", "dist": "uniform", "seed": 316, "scale": 0.1, "offset": 1 },
+        "betaT": { "shape": [768], "dtype": "float32", "dist": "normal", "seed": 317, "scale": 0.1 },
+        "biasT": { "shape": [768], "dtype": "float32", "dist": "normal", "seed": 318, "scale": 0.05 }
+      },
+      "outputs": { "outputT": { "shape": [4096, 768], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (3 * args.rows * args.hidden + 3 * args.hidden)" }] }
+    },
+    {
+      "name": "skip-layernorm-f32-4096x770-hidden770-unaligned",
+      "preset": "model",
+      "vars": { "rows": 4096, "hidden": 770 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "shape": [4096, 770], "dtype": "float32", "dist": "normal", "seed": 413, "scale": 0.2 },
+        "skipT": { "shape": [4096, 770], "dtype": "float32", "dist": "normal", "seed": 414, "scale": 0.2 },
+        "gammaT": { "shape": [770], "dtype": "float32", "dist": "uniform", "seed": 415, "scale": 0.1, "offset": 1 },
+        "betaT": { "shape": [770], "dtype": "float32", "dist": "normal", "seed": 416, "scale": 0.1 }
+      },
+      "outputs": {
+        "outputT": { "shape": [4096, 770], "dtype": "float32" },
+        "residualT": { "shape": [4096, 770], "dtype": "float32" }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (4 * args.rows * args.hidden + 2 * args.hidden)" }] }
+    },
+    {
+      "name": "skip-layernorm-f32-4096x768-hidden768-aligned-sibling",
+      "preset": "model",
+      "vars": { "rows": 4096, "hidden": 768 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "shape": [4096, 768], "dtype": "float32", "dist": "normal", "seed": 413, "scale": 0.2 },
+        "skipT": { "shape": [4096, 768], "dtype": "float32", "dist": "normal", "seed": 414, "scale": 0.2 },
+        "gammaT": { "shape": [768], "dtype": "float32", "dist": "uniform", "seed": 415, "scale": 0.1, "offset": 1 },
+        "betaT": { "shape": [768], "dtype": "float32", "dist": "normal", "seed": 416, "scale": 0.1 }
+      },
+      "outputs": {
+        "outputT": { "shape": [4096, 768], "dtype": "float32" },
+        "residualT": { "shape": [4096, 768], "dtype": "float32" }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (4 * args.rows * args.hidden + 2 * args.hidden)" }] }
+    },
+    {
+      "name": "skip-layernorm-f32-rows1-hidden4096-decode",
+      "preset": "model",
+      "vars": { "rows": 1, "hidden": 4096 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "shape": [1, 4096], "dtype": "float32", "dist": "normal", "seed": 513, "scale": 0.2 },
+        "skipT": { "shape": [1, 4096], "dtype": "float32", "dist": "normal", "seed": 514, "scale": 0.2 },
+        "gammaT": { "shape": [4096], "dtype": "float32", "dist": "uniform", "seed": 515, "scale": 0.1, "offset": 1 },
+        "betaT": { "shape": [4096], "dtype": "float32", "dist": "normal", "seed": 516, "scale": 0.1 }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 4096], "dtype": "float32" },
+        "residualT": { "shape": [1, 4096], "dtype": "float32" }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (4 * args.rows * args.hidden + 2 * args.hidden)" }] }
+    },
+    {
+      "name": "skip-layernorm-f32-rows4096-hidden4096-prefill-sibling",
+      "preset": "model",
+      "vars": { "rows": 4096, "hidden": 4096 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "shape": [4096, 4096], "dtype": "float32", "dist": "normal", "seed": 513, "scale": 0.2 },
+        "skipT": { "shape": [4096, 4096], "dtype": "float32", "dist": "normal", "seed": 514, "scale": 0.2 },
+        "gammaT": { "shape": [4096], "dtype": "float32", "dist": "uniform", "seed": 515, "scale": 0.1, "offset": 1 },
+        "betaT": { "shape": [4096], "dtype": "float32", "dist": "normal", "seed": 516, "scale": 0.1 }
+      },
+      "outputs": {
+        "outputT": { "shape": [4096, 4096], "dtype": "float32" },
+        "residualT": { "shape": [4096, 4096], "dtype": "float32" }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (4 * args.rows * args.hidden + 2 * args.hidden)" }] }
+    },
+    {
+      "name": "skip-layernorm-f32-4096x4095-hidden4095-unaligned-largewidth",
+      "preset": "model",
+      "vars": { "rows": 4096, "hidden": 4095 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "shape": [4096, 4095], "dtype": "float32", "dist": "normal", "seed": 613, "scale": 0.2 },
+        "skipT": { "shape": [4096, 4095], "dtype": "float32", "dist": "normal", "seed": 614, "scale": 0.2 },
+        "gammaT": { "shape": [4095], "dtype": "float32", "dist": "uniform", "seed": 615, "scale": 0.1, "offset": 1 },
+        "betaT": { "shape": [4095], "dtype": "float32", "dist": "normal", "seed": 616, "scale": 0.1 }
+      },
+      "outputs": {
+        "outputT": { "shape": [4096, 4095], "dtype": "float32" },
+        "residualT": { "shape": [4096, 4095], "dtype": "float32" }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (4 * args.rows * args.hidden + 2 * args.hidden)" }] }
+    },
+    {
+      "name": "skip-layernorm-f32-bias-8192x1025-unaligned-twopass-fallback",
+      "preset": "stress",
+      "vars": { "rows": 8192, "hidden": 1025 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "shape": [8192, 1025], "dtype": "float32", "dist": "normal", "seed": 713, "scale": 0.2 },
+        "skipT": { "shape": [8192, 1025], "dtype": "float32", "dist": "normal", "seed": 714, "scale": 0.2 },
+        "gammaT": { "shape": [1025], "dtype": "float32", "dist": "uniform", "seed": 715, "scale": 0.1, "offset": 1 },
+        "betaT": { "shape": [1025], "dtype": "float32", "dist": "normal", "seed": 716, "scale": 0.1 },
+        "biasT": { "shape": [1025], "dtype": "float32", "dist": "normal", "seed": 717, "scale": 0.05 }
+      },
+      "outputs": {
+        "outputT": { "shape": [8192, 1025], "dtype": "float32" },
+        "residualT": { "shape": [8192, 1025], "dtype": "float32" }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (4 * args.rows * args.hidden + 3 * args.hidden)" }] }
+    },
+    {
+      "name": "skip-layernorm-f32-bias-65535x1-wg1-lowoccupancy-fallback",
+      "preset": "edge",
+      "vars": { "rows": 65535, "hidden": 1 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "shape": [65535, 1], "dtype": "float32", "dist": "normal", "seed": 813, "scale": 0.2 },
+        "skipT": { "shape": [65535, 1], "dtype": "float32", "dist": "normal", "seed": 814, "scale": 0.2 },
+        "gammaT": { "shape": [1], "dtype": "float32", "dist": "uniform", "seed": 815, "scale": 0.1, "offset": 1 },
+        "betaT": { "shape": [1], "dtype": "float32", "dist": "normal", "seed": 816, "scale": 0.1 },
+        "biasT": { "shape": [1], "dtype": "float32", "dist": "normal", "seed": 817, "scale": 0.05 }
+      },
+      "outputs": {
+        "outputT": { "shape": [65535, 1], "dtype": "float32" },
+        "residualT": { "shape": [65535, 1], "dtype": "float32" }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (4 * args.rows * args.hidden + 3 * args.hidden)" }] }
+    }
+  ]
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,1108 @@

+{
+  "domain": "com.microsoft",
+  "name": "SkipLayerNormalization",
+  "sinceVersion": 1,
+  "description": "Fuses skip addition with layer normalization. The rank-3 standard surface currently supports float32, required `beta`, no `bias` or residual output, hidden sizes divisible by four, and exact or documented broadcast skip shapes. The provider's rank-2 extension supports float32 output-only with optional `beta`, or `beta` with optional `bias` when emitting the residual; its float16 path requires `beta`, `bias`, a residual output, and four-wide hidden size. Other combinations, bfloat16, and training statistics are not implemented.",
+  "inputs": [
+    {
+      "role": "input",
+      "dtype": "T",
+      "description": "Primary input normalized over the final hidden-size axis. Rank 3 is the public shape; rank 2 is an ONNX Runtime provider extension supported by this package."
+    },
+    {
+      "role": "skip",
+      "dtype": "T",
+      "description": "Residual tensor. For rank-3 input it is exact shape, `(1, sequence_length, hidden_size)`, or `(sequence_length, hidden_size)`; rank-2 input requires exact shape."
+    },
+    { "role": "gamma", "dtype": "T", "rank": 1, "description": "Layer-norm scale weights of shape `(hidden_size)`." },
+    {
+      "role": "beta",
+      "dtype": "T",
+      "rank": 1,
+      "optional": true,
+      "description": "Layer-norm bias weights of shape `(hidden_size)`."
+    },
+    {
+      "role": "bias",
+      "dtype": "T",
+      "rank": 1,
+      "optional": true,
+      "description": "Optional additive bias of shape `(hidden_size)` added to `input + skip` before normalization."
+    }
+  ],
+  "outputs": [
+    {
+      "role": "output",
+      "dtype": "T",
+      "rank": "ranks.inputT",
+      "shape": "shapes.inputT",
+      "description": "Normalized output tensor with the same shape as `input`."
+    },
+    {
+      "role": "input_skip_bias_sum",
+      "dtype": "T",
+      "rank": "ranks.inputT",
+      "optional": true,
+      "shape": "shapes.inputT",
+      "description": "Sum of `input`, `skip`, and `bias` (when present) before normalization, with the same shape as `input`."
+    }
+  ],
+  "attributes": { "epsilon": 9.999999960041972e-13 },
+  "attributeDescriptions": { "epsilon": "Non-negative epsilon added to the variance before taking the square root." },
+  "args": {
+    "inputT": { "kind": "tensor", "semantic": "input", "role": "input" },
+    "skipT": { "kind": "tensor", "semantic": "skip", "role": "input" },
+    "gammaT": { "kind": "tensor", "semantic": "gamma", "role": "input" },
+    "betaT": { "kind": "tensor", "semantic": "beta", "role": "input", "required": false },
+    "biasT": { "kind": "tensor", "semantic": "bias", "role": "input", "required": false },
+    "outputT": { "kind": "tensor", "semantic": "output", "role": "output" },
+    "residualT": { "kind": "tensor", "semantic": "input_skip_bias_sum", "role": "output", "required": false }
+  },
+  "typeConstraints": { "T": ["float32", "float16"] },
+  "derive": {
+    "rowCount": "dim(shapes.inputT, 0)",
+    "hiddenSize": "dim(shapes.inputT, 1)",
+    "skipWg": "max(1, min(tunables.MAX_WORKGROUP_SIZE, device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX, pow2ceil(hiddenSize)))",
+    "skipWgVec4": "max(1, min(tunables.MAX_WORKGROUP_SIZE, device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX, pow2ceil(ceilDiv(hiddenSize, 4))))",
+    "portableWideExecution": "not has(device.adapterInfo, \"subgroupMinSize\") or device.adapterInfo.subgroupMinSize >= 32",
+    "broadcastRows": "dim(shapes.inputT, 0) * dim(shapes.inputT, 1)",
+    "broadcastHiddenSize": "dim(shapes.inputT, 2)",
+    "broadcastSkipWgVec4": "max(1, min(tunables.MAX_WORKGROUP_SIZE, device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX, pow2ceil(ceilDiv(broadcastHiddenSize, 4))))",
+    "rowDispatchFits": "rowCount <= device.limits.maxComputeWorkgroupsPerDimension * device.limits.maxComputeWorkgroupsPerDimension",
+    "broadcastDispatchFits": "broadcastRows <= device.limits.maxComputeWorkgroupsPerDimension * device.limits.maxComputeWorkgroupsPerDimension",
+    "normResourcesFit": "skipWg * 8 <= device.limits.maxComputeWorkgroupStorageSize and skipWgVec4 * 8 <= device.limits.maxComputeWorkgroupStorageSize",
+    "broadcastResourcesFit": "broadcastSkipWgVec4 * 8 <= device.limits.maxComputeWorkgroupStorageSize",
+    "epsilonOk": "attrs.epsilon >= 0",
+    "rank2CoreContract": "epsilonOk and ranks.inputT == 2 and ranks.skipT == 2 and ranks.gammaT == 1 and ranks.outputT == 2 and dim(shapes.inputT, 0) == dim(shapes.skipT, 0) and dim(shapes.inputT, 1) == dim(shapes.skipT, 1) and dim(shapes.inputT, 1) > 0 and dim(shapes.gammaT, 0) == dim(shapes.inputT, 1) and dim(shapes.outputT, 0) == dim(shapes.inputT, 0) and dim(shapes.outputT, 1) == dim(shapes.inputT, 1)",
+    "residualOutputContract": "present.residualT and ranks.residualT == 2 and dim(shapes.residualT, 0) == dim(shapes.inputT, 0) and dim(shapes.residualT, 1) == dim(shapes.inputT, 1)",
+    "outputOnlyContract": "not present.residualT",
+    "betaContract": "false if not present.betaT else (ranks.betaT == 1 and dim(shapes.betaT, 0) == dim(shapes.inputT, 1))",
+    "noBetaContract": "not present.betaT",
+    "f32MainDtypes": "tensorDtypes.inputT == \"float32\" and tensorDtypes.skipT == \"float32\" and tensorDtypes.gammaT == \"float32\" and tensorDtypes.outputT == \"float32\"",
+    "f16MainDtypes": "tensorDtypes.inputT == \"float16\" and tensorDtypes.skipT == \"float16\" and tensorDtypes.gammaT == \"float16\" and tensorDtypes.outputT == \"float16\"",
+    "f32ResidualDtypes": "f32MainDtypes and tensorDtypes.residualT == \"float32\" if present.residualT else false",
+    "f16ResidualDtypes": "f16MainDtypes and tensorDtypes.residualT == \"float16\" if present.residualT else false",
+    "vec4Aligned": "dim(shapes.inputT, 1) % 4 == 0",
+    "broadcastSkipShapeOk": "(ranks.skipT == 2 and dim(shapes.skipT, 0) == dim(shapes.inputT, 1) and dim(shapes.skipT, 1) == dim(shapes.inputT, 2)) or (ranks.skipT == 3 and ((dim(shapes.skipT, 0) == 1 and dim(shapes.skipT, 1) == dim(shapes.inputT, 1) and dim(shapes.skipT, 2) == dim(shapes.inputT, 2)) or sameShape(shapes.skipT, shapes.inputT)))",
+    "broadcastOutputOnlyContract": "false if ranks.inputT != 3 or not present.betaT else (epsilonOk and not present.biasT and not present.residualT and dim(shapes.inputT, 2) % 4 == 0 and broadcastSkipShapeOk and ranks.gammaT == 1 and ranks.betaT == 1 and ranks.outputT == 3 and tensorDtypes.inputT == \"float32\" and tensorDtypes.skipT == \"float32\" and tensorDtypes.gammaT == \"float32\" and tensorDtypes.betaT == \"float32\" and tensorDtypes.outputT == \"float32\" and dim(shapes.inputT, 2) > 0 and dim(shapes.gammaT, 0) == dim(shapes.inputT, 2) and dim(shapes.betaT, 0) == dim(shapes.inputT, 2) and sameShape(shapes.outputT, shapes.inputT))",
+    "hasSubgroups": "device.features.has(\"subgroups\") and device.wgslLanguageFeatures.has(\"subgroup_id\")",
+    "hasF16": "device.features.has(\"shader-f16\")",
+    "f32_beta_no_bias_residual_contract": "rank2CoreContract and residualOutputContract and betaContract and f32ResidualDtypes and not present.biasT and tensorDtypes.betaT == \"float32\"",
+    "f32_beta_bias_residual_contract": "false if not present.biasT else (rank2CoreContract and residualOutputContract and betaContract and f32ResidualDtypes and ranks.biasT == 1 and tensorDtypes.betaT == \"float32\" and tensorDtypes.biasT == \"float32\" and dim(shapes.biasT, 0) == hiddenSize)",
+    "f16_beta_bias_residual_contract": "false if not present.biasT else (hasF16 and rank2CoreContract and residualOutputContract and betaContract and f16ResidualDtypes and ranks.biasT == 1 and tensorDtypes.betaT == \"float16\" and tensorDtypes.biasT == \"float16\" and dim(shapes.biasT, 0) == hiddenSize)",
+    "f32_no_beta_output_contract": "rank2CoreContract and outputOnlyContract and noBetaContract and f32MainDtypes and not present.biasT",
+    "f32_beta_no_bias_output_only_contract": "rank2CoreContract and outputOnlyContract and betaContract and f32MainDtypes and not present.biasT and tensorDtypes.betaT == \"float32\"",
+    "f32_beta_bias_output_only_contract": "false if not present.biasT else (rank2CoreContract and outputOnlyContract and betaContract and f32MainDtypes and ranks.biasT == 1 and tensorDtypes.betaT == \"float32\" and tensorDtypes.biasT == \"float32\" and dim(shapes.biasT, 0) == hiddenSize)"
+  },
+  "bindingSets": {
+    "scalar_no_bias_residual": [
+      {
+        "name": "input",
+        "arg": "inputT",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar"
+      },
+      {
+        "name": "skip",
+        "arg": "skipT",
+        "semantic": "skip",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar"
+      },
+      {
+        "name": "gamma",
+        "arg": "gammaT",
+        "semantic": "gamma",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar",
+        "length": "$HIDDEN_LEN"
+      },
+      {
+        "name": "beta",
+        "arg": "betaT",
+        "semantic": "beta",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar",
+        "length": "$HIDDEN_LEN"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$scalar"
+      },
+      {
+        "name": "input_skip_bias_sum",
+        "arg": "residualT",
+        "semantic": "input_skip_bias_sum",
+        "buffer": { "type": "storage" },
+        "elementType": "$scalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "rowCount" },
+            { "name": "epsilon", "type": "f32", "value": "attrs.epsilon" }
+          ]
+        }
+      }
+    ],
+    "scalar_bias_output_residual": [
+      {
+        "name": "input",
+        "arg": "inputT",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar"
+      },
+      {
+        "name": "skip",
+        "arg": "skipT",
+        "semantic": "skip",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar"
+      },
+      {
+        "name": "bias",
+        "arg": "biasT",
+        "semantic": "bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar",
+        "length": "$HIDDEN_LEN"
+      },
+      {
+        "name": "gamma",
+        "arg": "gammaT",
+        "semantic": "gamma",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar",
+        "length": "$HIDDEN_LEN"
+      },
+      {
+        "name": "beta",
+        "arg": "betaT",
+        "semantic": "beta",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar",
+        "length": "$HIDDEN_LEN"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$scalar"
+      },
+      {
+        "name": "input_skip_bias_sum",
+        "arg": "residualT",
+        "semantic": "input_skip_bias_sum",
+        "buffer": { "type": "storage" },
+        "elementType": "$scalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "rowCount" },
+            { "name": "epsilon", "type": "f32", "value": "attrs.epsilon" }
+          ]
+        }
+      }
+    ],
+    "vec4_bias_output_residual": [
+      {
+        "name": "input",
+        "arg": "inputT",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "skip",
+        "arg": "skipT",
+        "semantic": "skip",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "bias",
+        "arg": "biasT",
+        "semantic": "bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar",
+        "length": "$HIDDEN_LEN"
+      },
+      {
+        "name": "gamma",
+        "arg": "gammaT",
+        "semantic": "gamma",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar",
+        "length": "$HIDDEN_LEN"
+      },
+      {
+        "name": "beta",
+        "arg": "betaT",
+        "semantic": "beta",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar",
+        "length": "$HIDDEN_LEN"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "input_skip_bias_sum",
+        "arg": "residualT",
+        "semantic": "input_skip_bias_sum",
+        "buffer": { "type": "storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "rowCount" },
+            {
+              "name": "rowStride",
+              "type": "u32",
+              "value": "max(1, min(rowCount, device.limits.maxComputeWorkgroupsPerDimension))"
+            },
+            { "name": "epsilon", "type": "f32", "value": "attrs.epsilon" }
+          ]
+        }
+      }
+    ],
+    "scalar_no_bias_output_only": [
+      {
+        "name": "input",
+        "arg": "inputT",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar"
+      },
+      {
+        "name": "skip",
+        "arg": "skipT",
+        "semantic": "skip",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar"
+      },
+      {
+        "name": "gamma",
+        "arg": "gammaT",
+        "semantic": "gamma",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar",
+        "length": "$HIDDEN_LEN"
+      },
+      {
+        "name": "beta",
+        "arg": "betaT",
+        "semantic": "beta",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar",
+        "length": "$HIDDEN_LEN"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$scalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "rowCount" },
+            { "name": "epsilon", "type": "f32", "value": "attrs.epsilon" }
+          ]
+        }
+      }
+    ],
+    "scalar_bias_output_only": [
+      {
+        "name": "input",
+        "arg": "inputT",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar"
+      },
+      {
+        "name": "skip",
+        "arg": "skipT",
+        "semantic": "skip",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar"
+      },
+      {
+        "name": "bias",
+        "arg": "biasT",
+        "semantic": "bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar",
+        "length": "$HIDDEN_LEN"
+      },
+      {
+        "name": "gamma",
+        "arg": "gammaT",
+        "semantic": "gamma",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar",
+        "length": "$HIDDEN_LEN"
+      },
+      {
+        "name": "beta",
+        "arg": "betaT",
+        "semantic": "beta",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar",
+        "length": "$HIDDEN_LEN"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$scalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "rowCount" },
+            { "name": "epsilon", "type": "f32", "value": "attrs.epsilon" }
+          ]
+        }
+      }
+    ],
+    "scalar_no_beta_output_only": [
+      {
+        "name": "input",
+        "arg": "inputT",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar"
+      },
+      {
+        "name": "skip",
+        "arg": "skipT",
+        "semantic": "skip",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar"
+      },
+      {
+        "name": "gamma",
+        "arg": "gammaT",
+        "semantic": "gamma",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar",
+        "length": "$HIDDEN_LEN"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$scalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "rowCount" },
+            { "name": "epsilon", "type": "f32", "value": "attrs.epsilon" }
+          ]
+        }
+      }
+    ],
+    "vec4_no_bias_residual": [
+      {
+        "name": "input",
+        "arg": "inputT",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "skip",
+        "arg": "skipT",
+        "semantic": "skip",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "gamma",
+        "arg": "gammaT",
+        "semantic": "gamma",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar",
+        "length": "$HIDDEN_LEN"
+      },
+      {
+        "name": "beta",
+        "arg": "betaT",
+        "semantic": "beta",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar",
+        "length": "$HIDDEN_LEN"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "input_skip_bias_sum",
+        "arg": "residualT",
+        "semantic": "input_skip_bias_sum",
+        "buffer": { "type": "storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "rowCount" },
+            {
+              "name": "rowStride",
+              "type": "u32",
+              "value": "max(1, min(rowCount, device.limits.maxComputeWorkgroupsPerDimension))"
+            },
+            { "name": "epsilon", "type": "f32", "value": "attrs.epsilon" }
+          ]
+        }
+      }
+    ],
+    "vec4_no_bias_output_only": [
+      {
+        "name": "input",
+        "arg": "inputT",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "skip",
+        "arg": "skipT",
+        "semantic": "skip",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "gamma",
+        "arg": "gammaT",
+        "semantic": "gamma",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar",
+        "length": "$HIDDEN_LEN"
+      },
+      {
+        "name": "beta",
+        "arg": "betaT",
+        "semantic": "beta",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar",
+        "length": "$HIDDEN_LEN"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "rowCount" },
+            {
+              "name": "rowStride",
+              "type": "u32",
+              "value": "max(1, min(rowCount, device.limits.maxComputeWorkgroupsPerDimension))"
+            },
+            { "name": "epsilon", "type": "f32", "value": "attrs.epsilon" }
+          ]
+        }
+      }
+    ],
+    "vec4_bias_output_only": [
+      {
+        "name": "input",
+        "arg": "inputT",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "skip",
+        "arg": "skipT",
+        "semantic": "skip",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "gamma",
+        "arg": "gammaT",
+        "semantic": "gamma",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar",
+        "length": "$HIDDEN_LEN"
+      },
+      {
+        "name": "beta",
+        "arg": "betaT",
+        "semantic": "beta",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar",
+        "length": "$HIDDEN_LEN"
+      },
+      {
+        "name": "bias",
+        "arg": "biasT",
+        "semantic": "bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar",
+        "length": "$HIDDEN_LEN"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "rowCount" },
+            {
+              "name": "rowStride",
+              "type": "u32",
+              "value": "max(1, min(rowCount, device.limits.maxComputeWorkgroupsPerDimension))"
+            },
+            { "name": "epsilon", "type": "f32", "value": "attrs.epsilon" }
+          ]
+        }
+      }
+    ],
+    "vec4_no_beta_output_only": [
+      {
+        "name": "input",
+        "arg": "inputT",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "skip",
+        "arg": "skipT",
+        "semantic": "skip",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "gamma",
+        "arg": "gammaT",
+        "semantic": "gamma",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar",
+        "length": "$HIDDEN_LEN"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "rowCount" },
+            {
+              "name": "rowStride",
+              "type": "u32",
+              "value": "max(1, min(rowCount, device.limits.maxComputeWorkgroupsPerDimension))"
+            },
+            { "name": "epsilon", "type": "f32", "value": "attrs.epsilon" }
+          ]
+        }
+      }
+    ],
+    "vec4_beta_broadcast_output_only": [
+      {
+        "name": "input",
+        "arg": "inputT",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "skip",
+        "arg": "skipT",
+        "semantic": "skip",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "gamma",
+        "arg": "gammaT",
+        "semantic": "gamma",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar",
+        "length": "$HIDDEN_LEN"
+      },
+      {
+        "name": "beta",
+        "arg": "betaT",
+        "semantic": "beta",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar",
+        "length": "$HIDDEN_LEN"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "dim(shapes.inputT, 0) * dim(shapes.inputT, 1)" },
+            {
+              "name": "rowStride",
+              "type": "u32",
+              "value": "max(1, min(broadcastRows, device.limits.maxComputeWorkgroupsPerDimension))"
+            },
+            { "name": "epsilon", "type": "f32", "value": "attrs.epsilon" },
+            { "name": "skipRows", "type": "u32", "value": "numel(shapes.skipT) / broadcastHiddenSize" }
+          ]
+        }
+      }
+    ]
+  },
+  "variants": [
+    {
+      "id": "beta_output_only_vec4_broadcast",
+      "priority": 19,
+      "when": ["broadcastOutputOnlyContract", "broadcastResourcesFit", "broadcastDispatchFits"],
+      "constants": { "scalar": "\"f32\"", "vectorScalar": "\"vec4<f32>\"", "HIDDEN_LEN": "broadcastHiddenSize / 4" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipLayerNormalization.BroadcastSkip",
+          "source": {
+            "shader": "norm-skip-row-vec4.wgsl.jinja",
+            "inputs": {
+              "simplified": false,
+              "hasBias": false,
+              "hasBeta": true,
+              "writeResidualSum": false,
+              "usesF16": false,
+              "broadcastSkip": true,
+              "hidden": "broadcastHiddenSize",
+              "hiddenVec": "broadcastHiddenSize / 4",
+              "wg": "broadcastSkipWgVec4",
+              "vecType": "\"vec4<f32>\"",
+              "useSubgroups": "hasSubgroups"
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "vec4_beta_broadcast_output_only",
+          "dispatch": { "workgroups": "broadcastRows" }
+        }
+      ]
+    },
+    {
+      "id": "beta_bias_vec4",
+      "priority": 15,
+      "when": ["f32_beta_bias_residual_contract", "vec4Aligned", "hasSubgroups or \"bias\" == \"bias\" or portableWideExecution", "normResourcesFit", "rowDispatchFits"],
+      "constants": {
+        "scalar": "\"f32\"",
+        "vectorScalar": "\"vec4<f32>\"",
+        "hasBias": "\"bias\" == \"bias\"",
+        "hiddenSize": "hiddenSize",
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize / 4"
+      },
+      "passes": [
+        {
+          "id": "normalize",
+          "name": "SkipLayerNormalization.Vec4.Normalize",
+          "source": {
+            "shader": "norm-skip-row-vec4.wgsl.jinja",
+            "inputs": {
+              "simplified": false,
+              "hasBias": "\"bias\" == \"bias\"",
+              "hasBeta": true,
+              "writeResidualSum": true,
+              "usesF16": false,
+              "hidden": "hiddenSize",
+              "hiddenVec": "hiddenSize / 4",
+              "wg": "skipWgVec4",
+              "vecType": "\"vec4<f32>\"",
+              "useSubgroups": "hasSubgroups"
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "vec4_bias_output_residual",
+          "dispatch": { "workgroups": "rowCount" }
+        }
+      ]
+    },
+    {
+      "id": "beta_bias_row",
+      "priority": 5,
+      "when": ["f32_beta_bias_residual_contract", "normResourcesFit", "rowDispatchFits"],
+      "constants": {
+        "simplified": false,
+        "useSubgroups": "hasSubgroups",
+        "hasBeta": true,
+        "writeResidualSum": true,
+        "hasBias": "\"bias\" == \"bias\"",
+        "scalar": "\"f32\"",
+        "hiddenSize": "hiddenSize",
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize"
+      },
+      "passes": [
+        {
+          "id": "normalize",
+          "name": "SkipLayerNormalization.Row.Normalize",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "scalar_bias_output_residual",
+          "dispatch": { "workgroups": "rowCount" },
+          "constants": { "writeResidualSum": true }
+        }
+      ]
+    },
+    {
+      "id": "beta_bias_vec4_f16",
+      "priority": 21,
+      "when": ["f16_beta_bias_residual_contract", "vec4Aligned", "normResourcesFit", "rowDispatchFits"],
+      "constants": { "scalar": "\"f16\"", "vectorScalar": "\"vec4<f16>\"", "HIDDEN_LEN": "hiddenSize / 4" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipLayerNormalization.Vec4",
+          "source": {
+            "shader": "norm-skip-row-vec4.wgsl.jinja",
+            "inputs": {
+              "simplified": false,
+              "hasBias": true,
+              "hasBeta": true,
+              "writeResidualSum": true,
+              "usesF16": true,
+              "hidden": "hiddenSize",
+              "hiddenVec": "hiddenSize / 4",
+              "wg": "skipWgVec4",
+              "vecType": "\"vec4<f16>\"",
+              "useSubgroups": "hasSubgroups"
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "vec4_bias_output_residual",
+          "dispatch": { "workgroups": "rowCount" }
+        }
+      ]
+    },
+    {
+      "id": "no_beta_output_only_vec4",
+      "priority": 20,
+      "when": ["f32_no_beta_output_contract", "vec4Aligned", "normResourcesFit", "rowDispatchFits"],
+      "constants": { "scalar": "\"f32\"", "vectorScalar": "\"vec4<f32>\"", "HIDDEN_LEN": "hiddenSize / 4" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipLayerNormalization.NoBetaOutputOnly.Vec4",
+          "source": {
+            "shader": "norm-skip-row-vec4.wgsl.jinja",
+            "inputs": {
+              "simplified": false,
+              "hasBias": false,
+              "hasBeta": false,
+              "writeResidualSum": false,
+              "usesF16": false,
+              "hidden": "hiddenSize",
+              "hiddenVec": "hiddenSize / 4",
+              "wg": "skipWgVec4",
+              "vecType": "\"vec4<f32>\"",
+              "useSubgroups": "hasSubgroups"
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "vec4_no_beta_output_only",
+          "dispatch": { "workgroups": "rowCount" }
+        }
+      ]
+    },
+    {
+      "id": "no_beta_output_only_row",
+      "priority": 10,
+      "when": ["f32_no_beta_output_contract", "normResourcesFit", "rowDispatchFits"],
+      "constants": {
+        "simplified": false,
+        "hasBias": false,
+        "hasBeta": false,
+        "writeResidualSum": false,
+        "useSubgroups": "hasSubgroups",
+        "scalar": "\"f32\"",
+        "hiddenSize": "hiddenSize",
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipLayerNormalization.NoBetaOutputOnly.Row",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "scalar_no_beta_output_only",
+          "dispatch": { "workgroups": "rowCount" }
+        }
+      ]
+    },
+    {
+      "id": "beta_no_bias_vec4",
+      "priority": 20,
+      "when": ["f32_beta_no_bias_residual_contract", "vec4Aligned", "hasSubgroups or portableWideExecution", "normResourcesFit", "rowDispatchFits"],
+      "constants": {
+        "scalar": "\"f32\"",
+        "vectorScalar": "\"vec4<f32>\"",
+        "hasBias": "\"no_bias\" == \"bias\"",
+        "hiddenSize": "hiddenSize",
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize / 4"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipLayerNormalization.Vec4",
+          "source": {
+            "shader": "norm-skip-row-vec4.wgsl.jinja",
+            "inputs": {
+              "simplified": false,
+              "hasBias": "\"no_bias\" == \"bias\"",
+              "hasBeta": true,
+              "writeResidualSum": true,
+              "usesF16": false,
+              "hidden": "hiddenSize",
+              "hiddenVec": "hiddenSize / 4",
+              "wg": "skipWgVec4",
+              "vecType": "\"vec4<f32>\"",
+              "useSubgroups": "hasSubgroups"
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "vec4_no_bias_residual",
+          "dispatch": { "workgroups": "rowCount" }
+        }
+      ]
+    },
+    {
+      "id": "beta_no_bias_row",
+      "priority": 10,
+      "when": ["f32_beta_no_bias_residual_contract", "normResourcesFit", "rowDispatchFits"],
+      "constants": {
+        "simplified": false,
+        "useSubgroups": "hasSubgroups",
+        "hasBeta": true,
+        "writeResidualSum": true,
+        "hasBias": "\"no_bias\" == \"bias\"",
+        "scalar": "\"f32\"",
+        "hiddenSize": "hiddenSize",
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipLayerNormalization.Row",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "scalar_no_bias_residual",
+          "dispatch": { "workgroups": "rowCount" }
+        }
+      ]
+    },
+    {
+      "id": "beta_no_bias_output_only_vec4",
+      "priority": 20,
+      "when": ["f32_beta_no_bias_output_only_contract", "vec4Aligned", "hasSubgroups or \"no_bias\" == \"bias\" or portableWideExecution", "normResourcesFit", "rowDispatchFits"],
+      "constants": {
+        "scalar": "\"f32\"",
+        "vectorScalar": "\"vec4<f32>\"",
+        "hasBias": "\"no_bias\" == \"bias\"",
+        "hiddenSize": "hiddenSize",
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize / 4"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipLayerNormalization.Vec4",
+          "source": {
+            "shader": "norm-skip-row-vec4.wgsl.jinja",
+            "inputs": {
+              "simplified": false,
+              "hasBias": "\"no_bias\" == \"bias\"",
+              "hasBeta": true,
+              "writeResidualSum": false,
+              "usesF16": false,
+              "hidden": "hiddenSize",
+              "hiddenVec": "hiddenSize / 4",
+              "wg": "skipWgVec4",
+              "vecType": "\"vec4<f32>\"",
+              "useSubgroups": "hasSubgroups"
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "vec4_no_bias_output_only",
+          "dispatch": { "workgroups": "rowCount" }
+        }
+      ]
+    },
+    {
+      "id": "beta_no_bias_output_only_row",
+      "priority": 10,
+      "when": ["f32_beta_no_bias_output_only_contract", "normResourcesFit", "rowDispatchFits"],
+      "constants": {
+        "simplified": false,
+        "useSubgroups": "hasSubgroups",
+        "hasBeta": true,
+        "writeResidualSum": false,
+        "hasBias": "\"no_bias\" == \"bias\"",
+        "scalar": "\"f32\"",
+        "hiddenSize": "hiddenSize",
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipLayerNormalization.Row",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "scalar_no_bias_output_only",
+          "dispatch": { "workgroups": "rowCount" }
+        }
+      ]
+    },
+    {
+      "id": "beta_bias_output_only_vec4",
+      "priority": 20,
+      "when": ["f32_beta_bias_output_only_contract", "vec4Aligned", "hasSubgroups or \"bias\" == \"bias\" or portableWideExecution", "normResourcesFit", "rowDispatchFits"],
+      "constants": {
+        "scalar": "\"f32\"",
+        "vectorScalar": "\"vec4<f32>\"",
+        "hasBias": "\"bias\" == \"bias\"",
+        "hiddenSize": "hiddenSize",
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize / 4"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipLayerNormalization.Vec4",
+          "source": {
+            "shader": "norm-skip-row-vec4.wgsl.jinja",
+            "inputs": {
+              "simplified": false,
+              "hasBias": "\"bias\" == \"bias\"",
+              "hasBeta": true,
+              "writeResidualSum": false,
+              "usesF16": false,
+              "hidden": "hiddenSize",
+              "hiddenVec": "hiddenSize / 4",
+              "wg": "skipWgVec4",
+              "vecType": "\"vec4<f32>\"",
+              "useSubgroups": "hasSubgroups"
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "vec4_bias_output_only",
+          "dispatch": { "workgroups": "rowCount" }
+        }
+      ]
+    },
+    {
+      "id": "beta_bias_output_only_row",
+      "priority": 10,
+      "when": ["f32_beta_bias_output_only_contract", "normResourcesFit", "rowDispatchFits"],
+      "constants": {
+        "simplified": false,
+        "useSubgroups": "hasSubgroups",
+        "hasBeta": true,
+        "writeResidualSum": false,
+        "hasBias": "\"bias\" == \"bias\"",
+        "scalar": "\"f32\"",
+        "hiddenSize": "hiddenSize",
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipLayerNormalization.Row",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "scalar_bias_output_only",
+          "dispatch": { "workgroups": "rowCount" }
+        }
+      ]
+    }
+  ],
+  "tunables": { "MAX_WORKGROUP_SIZE": 256 }
+}

build/webgpu/metadata.json ADDED Viewed

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

+{
+  "name": "com.microsoft.SkipLayerNormalization",
+  "id": "_com_microsoft_skiplayernormalization_webgpu_1767bd7",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "STz1miuskWFMJijUblWNS3pG/DHJsRwQo0M+up/ng3s=",
+      "manifest.json": "MM5RYZ3xnC2rbelhnBH9kdS+qwkWaQmvEVXhBmag8jg=",
+      "norm-skip-row-vec4.wgsl.jinja": "qHoa8QjoRD1rgxTGGuJpfVb16Jou8PEg2CF1GjDt78s=",
+      "norm-skip-row.wgsl.jinja": "ck3s/G7WRph2YZRJs4LTft+R0IDjMPeE5Oy4UXDsqQs=",
+      "test.json": "Dkphsy4eO4Z8f2ZY1wfBJVmLcKNVWcOupJHOighE3RU="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/com.microsoft.SkipLayerNormalization" }
+}

build/webgpu/norm-skip-row-vec4.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,145 @@

+{% macro wgsl_tree_fold_stmt(a, op, idx, svar) %}
+{% if op == "max" %}
+{{ a }}[{{ idx }}] = max({{ a }}[{{ idx }}], {{ a }}[{{ idx }} + {{ svar }}]);
+{%- else %}
+{{ a }}[{{ idx }}] = {{ a }}[{{ idx }}] + {{ a }}[{{ idx }} + {{ svar }}];
+{%- endif %}
+{% endmacro %}
+{% macro wgsl_tree_fold(arrays, op="add", idx="lid", wg="WORKGROUP_SIZE", svar="stride", typed=false, form="tail", breakInline=false, bodyInline=false, barrierFirst=false) %}
+  var {{ svar }}{{ ": u32 " if typed else " " }}= {{ wg }} / 2u;
+  loop {
+{% if form == "head" %}
+{% if breakInline %}
+    if ({{ svar }} == 0u) { break; }
+{% else %}
+    if ({{ svar }} == 0u) {
+      break;
+    }
+{% endif %}
+{% endif %}
+{% if bodyInline %}
+    if ({{ idx }} < {{ svar }}) { {{ wgsl_tree_fold_stmt(arrays[0], op, idx, svar) }} }
+{% else %}
+    if ({{ idx }} < {{ svar }}) {
+{% for a in arrays %}
+      {{ wgsl_tree_fold_stmt(a, op, idx, svar) }}
+{% endfor %}
+    }
+{% endif %}
+{% if form == "head" %}
+{% if barrierFirst %}
+    workgroupBarrier();
+    {{ svar }} = {{ svar }} / 2u;
+{% else %}
+    {{ svar }} = {{ svar }} / 2u;
+    workgroupBarrier();
+{% endif %}
+{% else %}
+    workgroupBarrier();
+    if ({{ svar }} == 1u) {
+      break;
+    }
+    {{ svar }} = {{ svar }} / 2u;
+{% endif %}
+  }
+{%- endmacro %}{% set useSubgroups = source.useSubgroups %}
+{% if source.usesF16 %}
+enable f16;
+{% endif %}
+{% if useSubgroups %}
+enable subgroups;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const HIDDEN: u32 = {{ source.hidden }}u;
+const HIDDEN_V: u32 = {{ source.hiddenVec }}u;
+const WG: u32 = {{ source.wg }}u;
+var<workgroup> sg_partials: array<vec2<f32>, WG>;
+fn reduce_pair(value: vec2<f32>{% if useSubgroups %}, sg_lane: u32, sg_id: u32, num_sg: u32{% else %}, tid: u32{% endif %}) -> vec2<f32> {
+{% if useSubgroups %}
+  let s = vec2<f32>(subgroupAdd(value.x), subgroupAdd(value.y));
+  if (num_sg == 1u) {
+    return s;
+  }
+  if (sg_lane == 0u) {
+    sg_partials[sg_id] = s;
+  }
+  workgroupBarrier();
+  var total = vec2<f32>(0.0, 0.0);
+  for (var i = 0u; i < num_sg; i = i + 1u) {
+    total = total + sg_partials[i];
+  }
+  return total;
+{% else %}
+  // No-subgroup tier: workgroup barrier tree-reduction (WG is a power of two).
+  sg_partials[tid] = value;
+  workgroupBarrier();
+{{ wgsl_tree_fold(["sg_partials"], idx="tid", wg="WG", form="head", breakInline=true) }}
+  return sg_partials[0];
+{% endif %}
+}
+// 4 contiguous residual elements (input[idx] + skip[skip_idx] [+ bias]) at vec4
+// index `vi`. skip_idx == idx for the normal (non-broadcast) path; for a skip
+// that broadcasts across the leading/batch dim uses a folded index.
+fn residual_value(idx: u32, skip_idx: u32{% if source.hasBias %}, vi: u32{% endif %}) -> vec4<f32> {
+  var value = vec4<f32>(input[idx]) + vec4<f32>(skip[skip_idx]);
+{% if source.hasBias %}
+  value = value + vec4<f32>(bias[vi]);
+{% endif %}
+  return value;
+}
+@compute @workgroup_size(WG, 1, 1)
+fn main(
+  @builtin(workgroup_id) wg_id: vec3<u32>,
+  @builtin(local_invocation_id) lid: vec3<u32>{% if useSubgroups %},
+  @builtin(subgroup_invocation_id) sg_lane: u32,
+  @builtin(subgroup_id) sg_id: u32,
+  @builtin(num_subgroups) num_sg: u32{% endif %}
+) {
+  let row = wg_id.x + wg_id.y * params.rowStride;
+  if (row >= params.rows) {
+    return;
+  }
+  let tid = lid.x;
+  let base = row * HIDDEN_V;
+{% if source.broadcastSkip %}
+  // skip broadcasts across the batch dim: fold row into [0, skipRows) so every
+  // batch reuses the same skip row (skipRows == params.rows ⇒ identity).
+  let skip_base = (row % params.skipRows) * HIDDEN_V;
+{% else %}
+  let skip_base = base;
+{% endif %}
+  let shift = residual_value(base, skip_base{% if source.hasBias %}, 0u{% endif %}).x;
+  var acc = vec2<f32>(0.0, 0.0);
+  for (var i = tid; i < HIDDEN_V; i = i + WG) {
+    let v = residual_value(base + i, skip_base + i{% if source.hasBias %}, i{% endif %});
+    let d = v - vec4<f32>(shift);
+    acc.x = acc.x + d.x + d.y + d.z + d.w;
+    acc.y = acc.y + dot(d, d);
+  }
+  let totals = reduce_pair(acc{% if useSubgroups %}, sg_lane, sg_id, num_sg{% else %}, tid{% endif %});
+  let mean_d = totals.x / f32(HIDDEN);
+  let variance = max(totals.y / f32(HIDDEN) - mean_d * mean_d, 0.0);
+  let row_inv = inverseSqrt(variance + params.epsilon);
+  let row_mean = shift + mean_d;
+  for (var i = tid; i < HIDDEN_V; i = i + WG) {
+    let idx = base + i;
+    let residual = residual_value(idx, skip_base + i{% if source.hasBias %}, i{% endif %});
+{% if source.writeResidualSum %}
+    input_skip_bias_sum[idx] = {{ source.vecType }}(residual);
+{% endif %}
+    var value = (residual - vec4<f32>(row_mean)) * row_inv * vec4<f32>(gamma[i]);
+{% if source.hasBeta %}
+    value = value + vec4<f32>(beta[i]);
+{% endif %}
+    output[idx] = {{ source.vecType }}(value);
+  }
+}

build/webgpu/norm-skip-row.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,226 @@

+{% macro wgsl_tree_fold_stmt(a, op, idx, svar) %}
+{% if op == "max" %}
+{{ a }}[{{ idx }}] = max({{ a }}[{{ idx }}], {{ a }}[{{ idx }} + {{ svar }}]);
+{%- else %}
+{{ a }}[{{ idx }}] = {{ a }}[{{ idx }}] + {{ a }}[{{ idx }} + {{ svar }}];
+{%- endif %}
+{% endmacro %}
+{% macro wgsl_tree_fold(arrays, op="add", idx="lid", wg="WORKGROUP_SIZE", svar="stride", typed=false, form="tail", breakInline=false, bodyInline=false, barrierFirst=false) %}
+  var {{ svar }}{{ ": u32 " if typed else " " }}= {{ wg }} / 2u;
+  loop {
+{% if form == "head" %}
+{% if breakInline %}
+    if ({{ svar }} == 0u) { break; }
+{% else %}
+    if ({{ svar }} == 0u) {
+      break;
+    }
+{% endif %}
+{% endif %}
+{% if bodyInline %}
+    if ({{ idx }} < {{ svar }}) { {{ wgsl_tree_fold_stmt(arrays[0], op, idx, svar) }} }
+{% else %}
+    if ({{ idx }} < {{ svar }}) {
+{% for a in arrays %}
+      {{ wgsl_tree_fold_stmt(a, op, idx, svar) }}
+{% endfor %}
+    }
+{% endif %}
+{% if form == "head" %}
+{% if barrierFirst %}
+    workgroupBarrier();
+    {{ svar }} = {{ svar }} / 2u;
+{% else %}
+    {{ svar }} = {{ svar }} / 2u;
+    workgroupBarrier();
+{% endif %}
+{% else %}
+    workgroupBarrier();
+    if ({{ svar }} == 1u) {
+      break;
+    }
+    {{ svar }} = {{ svar }} / 2u;
+{% endif %}
+  }
+{%- endmacro %}
+/* One workgroup normalizes each row of residual = input + skip, with an
+ * optional bias. */
+{% set degenerateRow = (not simplified) and hiddenSize == 1 %}
+{% if useSubgroups and not degenerateRow %}
+enable subgroups;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+{% if not degenerateRow or writeResidualSum %}
+const HIDDEN: u32 = {{ hiddenSize }}u;
+{% endif %}
+const WG: u32 = {{ workgroupSize }}u;
+{% if simplified %}
+var<workgroup> partial: array<f32, WG>;
+{% macro wgsl_tree_reduce_f32(name, mode, buffer="partial", wg="WG", trailingBarrier=true) %}
+fn {{ name }}(value: f32, tid: u32) -> f32 {
+  {{ buffer }}[tid] = value;
+  workgroupBarrier();
+  // Ceil-halving keeps every lane when the workgroup size is not a power of
+  // two. For even n this matches the power-of-two tree order; for odd n, lanes
+  // [0, n-half) fold the upper tail while the middle lane carries forward.
+  var n: u32 = {{ wg }};
+  loop {
+    let half = (n + 1u) / 2u;
+    if (tid < n - half) {
+{% if mode == "max" %}
+      {{ buffer }}[tid] = max({{ buffer }}[tid], {{ buffer }}[tid + half]);
+{% else %}
+      {{ buffer }}[tid] = {{ buffer }}[tid] + {{ buffer }}[tid + half];
+{% endif %}
+    }
+    workgroupBarrier();
+    n = half;
+    if (n == 1u) {
+      break;
+    }
+  }
+  // The default trailing barrier makes this helper safe for back-to-back calls: every lane reads
+  // slot 0 here, so the next call's first store must not run until all lanes have read it.
+  // `trailingBarrier=false` is safe only when the buffer is never written again before kernel exit.
+  let reduced = {{ buffer }}[0];
+{% if trailingBarrier %}
+  workgroupBarrier();
+{% endif %}
+  return reduced;
+}
+{% endmacro %}
+{{ wgsl_tree_reduce_f32("reduce_sum", "add", "partial", "WG") }}
+var<workgroup> row_inv: f32;
+{% else %}
+{% if not degenerateRow %}
+var<workgroup> pair_partial: array<vec2<f32>, WG>;
+{% if useSubgroups %}
+fn reduce_pair(value: vec2<f32>, sg_lane: u32, sg_id: u32, num_sg: u32) -> vec2<f32> {
+  let s = vec2<f32>(subgroupAdd(value.x), subgroupAdd(value.y));
+  if (num_sg == 1u) {
+    return s;
+  }
+  if (sg_lane == 0u) {
+    pair_partial[sg_id] = s;
+  }
+  workgroupBarrier();
+  var total = vec2<f32>(0.0, 0.0);
+  for (var i = 0u; i < num_sg; i = i + 1u) {
+    total = total + pair_partial[i];
+  }
+  return total;
+}
+{% else %}
+fn reduce_pair(value: vec2<f32>, tid: u32) -> vec2<f32> {
+  pair_partial[tid] = value;
+  workgroupBarrier();
+{{ wgsl_tree_fold(["pair_partial"], idx="tid", wg="WG", form="head") }}
+  return pair_partial[0];
+}
+{% endif %}
+{% endif %}
+{% endif %}
+{% if not degenerateRow or writeResidualSum %}
+fn residual_value(row: u32, d: u32) -> f32 {
+  let index = row * HIDDEN + d;
+  var value = f32(input[index]) + f32(skip[index]);
+{% if hasBias %}
+  value = value + f32(bias[d]);
+{% endif %}
+  return value;
+}
+{% endif %}
+@compute @workgroup_size(WG, 1, 1)
+fn main(
+  @builtin(workgroup_id) wg: vec3<u32>,
+  @builtin(num_workgroups) nwg: vec3<u32>{% if not degenerateRow %},
+  @builtin(local_invocation_id) lid: vec3<u32>{% endif %}{% if useSubgroups and not degenerateRow %},
+  @builtin(subgroup_invocation_id) sg_lane: u32,
+  @builtin(subgroup_id) sg_id: u32,
+  @builtin(num_subgroups) num_sg: u32{% endif %}
+) {
+  // 2D-folded row index: wg.y carries the high bits past the maxComputeWorkgroupsPerDimension
+  // workgroup-per-dimension dispatch limit. Reduces to wg.x when nwg.y == 1;
+  // the row >= params.rows guard drops the over-dispatched tail.
+  let row = wg.x + wg.y * nwg.x;
+  if (row >= params.rows) {
+    return;
+  }
+{% if not degenerateRow %}
+  let tid = lid.x;
+{% endif %}
+{% if simplified %}
+  // RMS normalization uses one sum-of-squares sweep, without a mean or beta.
+  var local_sq = 0.0;
+  for (var d: u32 = tid; d < HIDDEN; d = d + WG) {
+    let value = residual_value(row, d);
+    local_sq = local_sq + value * value;
+  }
+  let sq = reduce_sum(local_sq, tid);
+  if (tid == 0u) {
+    row_inv = inverseSqrt(sq / f32(HIDDEN) + params.epsilon);
+  }
+  workgroupBarrier();
+  for (var d: u32 = tid; d < HIDDEN; d = d + WG) {
+    let index = row * HIDDEN + d;
+    let residual = residual_value(row, d);
+{% if writeResidualSum %}
+    input_skip_bias_sum[index] = {{ scalar }}(residual);
+{% endif %}
+    output[index] = {{ scalar }}(residual * row_inv * f32(gamma[d]));
+  }
+{% elif degenerateRow %}
+  // HIDDEN == 1: the row's mean is its only element, so the centered value and
+  // the variance are exactly zero and the output reduces to beta. The closed
+  // form avoids computing that zero by subtracting two equal rounded values.
+  let row_inv = inverseSqrt(params.epsilon);
+{% if writeResidualSum %}
+  let residual = residual_value(row, 0u);
+  input_skip_bias_sum[row] = {{ scalar }}(residual);
+{% endif %}
+  // 0.0 * row_inv keeps the IEEE result when epsilon == 0 makes row_inv +Inf.
+  output[row] = {{ scalar }}(0.0 * row_inv * f32(gamma[0]){% if hasBeta %} + f32(beta[0]){% endif %});
+{% else %}
+  // Shifted moments: accumulating (x - x[0], (x - x[0])^2) keeps the sums
+  // small for rows with a large common offset; every thread reconstructs the
+  // row mean and variance from the merged pair.
+  let shift = residual_value(row, 0u);
+  var acc = vec2<f32>(0.0, 0.0);
+  for (var d = tid; d < HIDDEN; d = d + WG) {
+    let centered = residual_value(row, d) - shift;
+    acc.x = acc.x + centered;
+    acc.y = acc.y + centered * centered;
+  }
+{% if useSubgroups %}
+  let totals = reduce_pair(acc, sg_lane, sg_id, num_sg);
+{% else %}
+  let totals = reduce_pair(acc, tid);
+{% endif %}
+  let mean_delta = totals.x / f32(HIDDEN);
+  let row_mean = shift + mean_delta;
+  let variance = max(totals.y / f32(HIDDEN) - mean_delta * mean_delta, 0.0);
+  let row_inv = inverseSqrt(variance + params.epsilon);
+  for (var d = tid; d < HIDDEN; d = d + WG) {
+    let index = row * HIDDEN + d;
+    let residual = residual_value(row, d);
+{% if writeResidualSum %}
+    input_skip_bias_sum[index] = {{ scalar }}(residual);
+{% endif %}
+    output[index] = {{ scalar }}((residual - row_mean) * row_inv * f32(gamma[d]){% if hasBeta %} + f32(beta[d]){% endif %});
+  }
+{% endif %}
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,939 @@

+{
+  "op": "com.microsoft.SkipLayerNormalization",
+  "fixtureArrays": {
+    "ort_batch2_bias_flattened_tokens_input_skipT": [0.1, -0.2, 0.3, 1, 0.5, 0.1, 0.4, 1.6, 1.8, -0.3, 0, 1, -0.5, 0.4, 0.8, -0.6],
+    "ort_batch2_flattened_tokens_input_inputT": [0.8, -0.5, 0, 1, 0.5, 0.2, 0.3, -0.6, 0.8, -0.5, 0, 1, 0.5, 0.2, 0.3, -0.6],
+    "ort_batch2_skip_broadcast_no_batch_size_output_outputT": [0.2843385934829712, -0.17090578377246857, -0.9289716482162476, 4.692415237426758, 0.4611165225505829, -0.21333980560302734, -0.29631003737449646, 3.5148544311523438, 0.2843385934829712, -0.17090578377246857, -0.9289716482162476, 4.692415237426758, 0.4611165225505829, -0.21333980560302734, -0.29631003737449646, 3.5148544311523438]
+  },
+  "cases": [
+    {
+      "name": "rank3_exact_skip_shape",
+      "provenance": {
+        "source": "onnxruntime/contrib_ops/cpu/bert/skip_layer_norm.cc",
+        "notes": "Pins the public rank-3 same-shape skip mode independently of the two documented broadcast forms."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [2, 2, 4], "data": { "kind": "linspace", "start": -1.5, "end": 1.5 } },
+        "skipT": { "dtype": "float32", "shape": [2, 2, 4], "data": { "kind": "linspace", "start": 0.75, "end": -0.25 } },
+        "gammaT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 0.5, 1.5, -0.25] } },
+        "betaT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.1, -0.2, 0.3, -0.4] } }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [2, 2, 4], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "no_bias",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [3, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [3, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41, "scale": 0.2 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.13, "scale": 0.1 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [3, 8], "tolerance": 0.00002 },
+        "residualT": { "dtype": "float32", "shape": [3, 8], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "f32_epsilon_zero_explicit_tiny_variance",
+      "attrs": { "epsilon": 0 },
+      "provenance": {
+        "source": "onnxruntime/contrib_ops/webgpu/bert/skip_layer_norm.h",
+        "test": "GetAttrOrDefault epsilon semantics",
+        "notes": "An explicit epsilon=0.0 must be honored, not replaced by the 1e-12 schema default via a truthiness fallback. The 1e-7-scale rows make that difference numerically observable."
+      },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [3, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 1e-7 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [3, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 1e-7 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41, "scale": 0.2 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.13, "scale": 0.1 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [3, 8], "tolerance": 0.0001 },
+        "residualT": { "dtype": "float32", "shape": [3, 8], "tolerance": 1e-9 }
+      }
+    },
+    {
+      "name": "bias",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [3, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [3, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41, "scale": 0.2 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.13, "scale": 0.1 }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.17, "scale": 0.08 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [3, 8], "tolerance": 0.00002 },
+        "residualT": { "dtype": "float32", "shape": [3, 8], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "zero_variance_returns_beta",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [5.0, 5.0, 5.0, 5.0, -3.0, -3.0, -3.0, -3.0] }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [1.0, 1.0, 1.0, 1.0, 7.0, 7.0, 7.0, 7.0] }
+        },
+        "gammaT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [10.0, -2.0, 3.0, 4.0] } },
+        "betaT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.5, -1.0, 2.0, -3.0] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 4], "tolerance": 0.000001 },
+        "residualT": { "dtype": "float32", "shape": [2, 4], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "hidden_size_one_bias_path",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [3, 1], "data": { "kind": "values", "values": [2.0, -4.0, 0.5] } },
+        "skipT": { "dtype": "float32", "shape": [3, 1], "data": { "kind": "values", "values": [3.0, 1.0, -0.5] } },
+        "gammaT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [7.0] } },
+        "betaT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [-2.0] } },
+        "biasT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.25] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [3, 1], "tolerance": 0.000001 },
+        "residualT": { "dtype": "float32", "shape": [3, 1], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "hidden_size_one_bias_output_only",
+      "provenance": {
+        "notes": "Pins the scalar hidden-size-one closed form when the optional residual sum is not requested."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [1, 1], "data": { "kind": "values", "values": [2.0] } },
+        "skipT": { "dtype": "float32", "shape": [1, 1], "data": { "kind": "values", "values": [3.0] } },
+        "gammaT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [7.0] } },
+        "betaT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [-2.0] } },
+        "biasT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.25] } }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [1, 1], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "hidden_size_one_no_beta_output_only",
+      "provenance": {
+        "notes": "Pins the supported no-beta hidden-size-one closed form: every centered value is zero, so the output is zero."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [1, 1], "data": { "kind": "values", "values": [2.0] } },
+        "skipT": { "dtype": "float32", "shape": [1, 1], "data": { "kind": "values", "values": [3.0] } },
+        "gammaT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [7.0] } }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [1, 1], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "hidden_size_one_bias_rows65535_dispatch_edge",
+      "provenance": {
+        "notes": "Correctness sibling for the low-occupancy hidden=1 benchmark at rows=65535, the maximum single-dimension workgroup count edge."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [65535, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.017, "cosStep": 0.031 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [65535, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.011, "cosStep": 0.023 }
+        },
+        "gammaT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.125] } },
+        "betaT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [-0.25] } },
+        "biasT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.125] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [65535, 1], "tolerance": 0.000001 },
+        "residualT": { "dtype": "float32", "shape": [65535, 1], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "large_mean_small_variance_centered",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 4],
+          "data": { "kind": "values", "values": [40000.0, 40001.0, 40002.0, 40003.0] }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [1, 4],
+          "data": { "kind": "values", "values": [-39999.0, -40000.0, -40001.0, -40002.0] }
+        },
+        "gammaT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 2.0, -1.0, 0.5] } },
+        "betaT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.0, 0.5, -0.25, 1.0] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 4], "tolerance": 0.000001 },
+        "residualT": { "dtype": "float32", "shape": [1, 4], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "ort_zero_tokens_null_input",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/skiplayernorm_op_test.cc",
+        "test": "SkipLayerNormTest.SkipLayerNormNullInput",
+        "notes": "ORT shape [1, 0, 4] is represented as lowered token rows [0, 4]."
+      },
+      "attrs": { "epsilon": 1e-12 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [0, 4], "data": { "kind": "values", "values": [] } },
+        "skipT": { "dtype": "float32", "shape": [0, 4], "data": { "kind": "values", "values": [] } },
+        "gammaT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.3, 0.2, 4.0, 2.2] } },
+        "betaT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.2, 0.1, 0.4, 1.6] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [0, 4], "tolerance": 0 },
+        "residualT": { "dtype": "float32", "shape": [0, 4], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "ort_batch1_flattened_tokens",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/skiplayernorm_op_test.cc",
+        "test": "SkipLayerNormTest.SkipLayerNormBatch1",
+        "notes": "ORT shape [1, 2, 4] is represented as [2, 4] tokens by this lowered kernel. Epsilon is omitted to exercise the schema default of 1e-12."
+      },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0.8, -0.5, 0.0, 1.0, 0.5, 0.2, 0.3, -0.6] }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0.1, -0.2, 0.3, 1.0, 0.5, 0.1, 0.4, 1.6] }
+        },
+        "gammaT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.3, 0.2, 4.0, 2.2] } },
+        "betaT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.2, 0.1, 0.4, 1.6] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 4], "tolerance": 0.00002 },
+        "residualT": { "dtype": "float32", "shape": [2, 4], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "ort_batch2_bias_flattened_tokens",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/skiplayernorm_op_test.cc",
+        "test": "SkipLayerNormTest.SkipLayerNormBatch2_Bias",
+        "notes": "ORT shape [2, 2, 4] is represented as [4, 4] tokens by this lowered kernel."
+      },
+      "attrs": { "epsilon": 1e-12 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [4, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.7, -0.4, -0.2, 1.2, 0.4, 0.3, 0.1, -0.4, 0.7, -0.4, -0.2, 1.2, 0.4, 0.3, 0.1, -0.4]
+          }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [4, 4],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_batch2_bias_flattened_tokens_input_skipT" }
+          }
+        },
+        "gammaT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.3, 0.2, 4.0, 2.2] } },
+        "betaT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.2, 0.1, 0.4, 1.6] } },
+        "biasT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.1, -0.1, 0.2, -0.2] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [4, 4], "tolerance": 0.00002 },
+        "residualT": { "dtype": "float32", "shape": [4, 4], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "ort_batch2_flattened_tokens",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/skiplayernorm_op_test.cc",
+        "test": "SkipLayerNormTest.SkipLayerNormBatch2",
+        "notes": "ORT shape [2, 2, 4] is represented as [4, 4] tokens by this lowered kernel."
+      },
+      "attrs": { "epsilon": 1e-12 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [4, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_batch2_flattened_tokens_input_inputT" } }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [4, 4],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_batch2_bias_flattened_tokens_input_skipT" }
+          }
+        },
+        "gammaT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.3, 0.2, 4.0, 2.2] } },
+        "betaT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.2, 0.1, 0.4, 1.6] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [4, 4], "tolerance": 0.00002 },
+        "residualT": { "dtype": "float32", "shape": [4, 4], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "large_hidden_320_no_bias",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 320],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.017, "cosStep": 0.031 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [2, 320],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.011, "cosStep": 0.023 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [320],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "offset": 1.0, "sinStep": 0.007, "cosStep": 0.041 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [320],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.019, "cosStep": 0.013 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 320], "tolerance": 0.0002 },
+        "residualT": { "dtype": "float32", "shape": [2, 320], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "large_hidden_320_bias",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 320],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.029 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [2, 320],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.019, "cosStep": 0.037 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [320],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "offset": 1.0, "sinStep": 0.029, "cosStep": 0.017 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [320],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.023, "cosStep": 0.011 }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [320],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.007 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 320], "tolerance": 0.0002 },
+        "residualT": { "dtype": "float32", "shape": [2, 320], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "ort_batch2_skip_broadcast_no_batch_size",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/skiplayernorm_op_test.cc",
+        "test": "SkipLayerNormTest.SkipLayerNormBatch2_Skip_Broadcast_No_Batch_Size"
+      },
+      "attrs": { "epsilon": 1e-12 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 2, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_batch2_flattened_tokens_input_inputT" } }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0.1, -0.2, 0.3, 1.0, 0.5, 0.1, 0.4, 1.6] }
+        },
+        "gammaT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.3, 0.2, 4.0, 2.2] } },
+        "betaT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.2, 0.1, 0.4, 1.6] } }
+      },
+      "outputs": {
+        "outputT": {
+          "dtype": "float32",
+          "shape": [2, 2, 4],
+          "tolerance": 0.00002,
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_batch2_skip_broadcast_no_batch_size_output_outputT" }
+          }
+        }
+      }
+    },
+    {
+      "name": "ort_batch1_no_beta_flattened_tokens",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/skiplayernorm_op_test.cc",
+        "test": "SkipLayerNormTest.SkipLayerNormBatch1_NoBeta",
+        "notes": "ORT shape [1, 2, 4] is represented as [2, 4] tokens by this lowered fixture, with beta omitted."
+      },
+      "attrs": { "epsilon": 1e-12 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0.8, -0.5, 0.0, 1.0, 0.5, 0.2, 0.3, -0.6] }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0.1, -0.2, 0.3, 1.0, 0.5, 0.1, 0.4, 1.6] }
+        },
+        "gammaT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.3, 0.2, 4.0, 2.2] } }
+      },
+      "outputs": {
+        "outputT": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "tolerance": 0.00002,
+          "data": {
+            "kind": "values",
+            "values": [0.0843385934829712, -0.27090578377246854, -1.3289716482162477, 3.092415237426758, 0.2611165225505829, -0.3133398056030273, -0.6963100373744965, 1.9148544311523439]
+          }
+        }
+      }
+    },
+    {
+      "name": "no_beta_output_only_hidden6_unaligned_row",
+      "provenance": {
+        "notes": "Only coverage for the scalar no_beta_output_only_row variant: beta omitted with output as the sole output picks the no-beta output-only pair, and hidden=6 is not a multiple of 4, so the vec4 sibling's vec4Aligned guard is false and the row kernel is selected on every tier."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [2, 6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41, "scale": 0.2 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [2, 6], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "beta_no_bias_output_only",
+      "provenance": {
+        "notes": "Exercises beta with all three optional auxiliary outputs absent, using a vec4-aligned hidden size of 8."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [3, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [3, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41, "scale": 0.2 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.1 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [3, 8], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "beta_no_bias_output_only_hidden6_unaligned_row",
+      "provenance": {
+        "notes": "Row-kernel arm of beta_no_bias_output_only: hidden=6 fails vec4Aligned, so the scalar sibling is selected on every tier."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.37 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [2, 6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.43 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41, "scale": 0.2 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.1 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [2, 6], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "beta_bias_output_only",
+      "provenance": {
+        "notes": "bias + beta with no optional outputs uses six storage buffers instead of the nine required when all optional outputs are present. It therefore remains valid at WebGPU's guaranteed minimum of eight storage buffers."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [3, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [3, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41, "scale": 0.2 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.1 }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.05, "cosStep": 0.47, "scale": 0.3 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [3, 8], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "beta_bias_output_only_hidden6_unaligned_row",
+      "provenance": {
+        "notes": "Row-kernel arm of beta_bias_output_only: hidden=6 fails vec4Aligned. 6 storage buffers."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.37 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [2, 6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.43 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41, "scale": 0.2 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.1 }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.05, "cosStep": 0.47, "scale": 0.3 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [2, 6], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "ort_batch2_skip_broadcast_batch_size_one",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/skiplayernorm_op_test.cc",
+        "test": "SkipLayerNormTest.SkipLayerNormBatch2_Skip_Broadcast_Batch_Size_1"
+      },
+      "attrs": { "epsilon": 1e-12 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 2, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_batch2_flattened_tokens_input_inputT" } }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [1, 2, 4],
+          "data": { "kind": "values", "values": [0.1, -0.2, 0.3, 1.0, 0.5, 0.1, 0.4, 1.6] }
+        },
+        "gammaT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.3, 0.2, 4.0, 2.2] } },
+        "betaT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.2, 0.1, 0.4, 1.6] } }
+      },
+      "outputs": {
+        "outputT": {
+          "dtype": "float32",
+          "shape": [2, 2, 4],
+          "tolerance": 0.00002,
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_batch2_skip_broadcast_no_batch_size_output_outputT" }
+          }
+        }
+      }
+    },
+    {
+      "name": "f16_hidden768_bias_residual",
+      "requires": { "features": ["shader-f16"] },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float16",
+          "shape": [4, 768],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.017, "cosStep": 0.031 }
+        },
+        "skipT": {
+          "dtype": "float16",
+          "shape": [4, 768],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.011, "cosStep": 0.023 }
+        },
+        "gammaT": {
+          "dtype": "float16",
+          "shape": [768],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "offset": 1.0, "sinStep": 0.007, "cosStep": 0.041 }
+        },
+        "betaT": {
+          "dtype": "float16",
+          "shape": [768],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.019, "cosStep": 0.013 }
+        },
+        "biasT": {
+          "dtype": "float16",
+          "shape": [768],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.007 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [4, 768], "tolerance": 0.01 },
+        "residualT": { "dtype": "float16", "shape": [4, 768], "tolerance": 0.005 }
+      }
+    },
+    {
+      "name": "f32_hidden768_no_bias_residual",
+      "provenance": {
+        "notes": "Compact companion for bench cases skip-layernorm-f32-4096x768 and skip-layernorm-f32-4096x768-hidden768-aligned-sibling, where reduced tiers expose a subgroup fallback gap."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [4, 768],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.017, "cosStep": 0.031 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [4, 768],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.011, "cosStep": 0.023 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [768],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "offset": 1.0, "sinStep": 0.007, "cosStep": 0.041 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [768],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.019, "cosStep": 0.013 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [4, 768], "tolerance": 0.0002 },
+        "residualT": { "dtype": "float32", "shape": [4, 768], "tolerance": 0.000002 }
+      }
+    },
+    {
+      "name": "f32_hidden2048_bias_residual",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 2048],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.029 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [2, 2048],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.019, "cosStep": 0.037 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [2048],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "offset": 1.0, "sinStep": 0.029, "cosStep": 0.017 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [2048],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.023, "cosStep": 0.011 }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [2048],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.007 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 2048], "tolerance": 0.0005 },
+        "residualT": { "dtype": "float32", "shape": [2, 2048], "tolerance": 0.000002 }
+      }
+    },
+    {
+      "name": "f32_hidden1025_bias_residual",
+      "provenance": {
+        "notes": "Compact companion for bench case skip-layernorm-f32-bias-8192x1025-unaligned-twopass-fallback, preserving the odd hidden-size bias path."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [3, 1025],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.029 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [3, 1025],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.019, "cosStep": 0.037 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [1025],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "offset": 1.0, "sinStep": 0.029, "cosStep": 0.017 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [1025],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.023, "cosStep": 0.011 }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [1025],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.007 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [3, 1025], "tolerance": 0.0005 },
+        "residualT": { "dtype": "float32", "shape": [3, 1025], "tolerance": 0.000002 }
+      }
+    },
+    {
+      "name": "f32_hidden770_unaligned_beta_bias_scalar_subgroup",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [3, 770],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.017, "cosStep": 0.031 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [3, 770],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.011, "cosStep": 0.023 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [770],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "offset": 1.0, "sinStep": 0.007, "cosStep": 0.041 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [770],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.019, "cosStep": 0.013 }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [770],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.007 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [3, 770], "tolerance": 0.0002 },
+        "residualT": { "dtype": "float32", "shape": [3, 770], "tolerance": 0.000002 }
+      }
+    },
+    {
+      "name": "f32_rows1_hidden4096_decode",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 4096],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.017, "cosStep": 0.031 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [1, 4096],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.011, "cosStep": 0.023 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [4096],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "offset": 1.0, "sinStep": 0.007, "cosStep": 0.041 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [4096],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.019, "cosStep": 0.013 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 4096], "tolerance": 0.0005 },
+        "residualT": { "dtype": "float32", "shape": [1, 4096], "tolerance": 0.000002 }
+      }
+    },
+    {
+      "name": "empty_tokens_bias_residual_twopass",
+      "attrs": { "epsilon": 1e-12 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [0, 4], "data": { "kind": "values", "values": [] } },
+        "skipT": { "dtype": "float32", "shape": [0, 4], "data": { "kind": "values", "values": [] } },
+        "gammaT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.3, 0.2, 4.0, 2.2] } },
+        "betaT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.2, 0.1, 0.4, 1.6] } },
+        "biasT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.1, -0.1, 0.2, -0.2] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [0, 4], "tolerance": 0 },
+        "residualT": { "dtype": "float32", "shape": [0, 4], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "rows65537_hidden3_fold_lastrow_guard",
+      "provenance": {
+        "notes": "Pins the two-dimensional dispatch fold and final-row guard with 65,537 hidden-size-3 rows; near-constant rows exercise float32 one-pass variance while the residual sum remains exact."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [65537, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.017, "cosStep": 0.031 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [65537, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.011, "cosStep": 0.023 }
+        },
+        "gammaT": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 2.0, -1.0] } },
+        "betaT": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.5, -0.25, 1.0] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [65537, 3], "tolerance": 0.0021 },
+        "residualT": { "dtype": "float32", "shape": [65537, 3], "tolerance": 0.000002 }
+      }
+    },
+    {
+      "name": "hidden_size_one_bias_many_rows",
+      "provenance": {
+        "notes": "Compact correctness sibling for the hidden-size-one low-occupancy benchmark; many rows with hidden=1 keep the variance-zero bias path honest without using benchmark-scale dimensions."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [257, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.017, "cosStep": 0.031 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [257, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.011, "cosStep": 0.023 }
+        },
+        "gammaT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.25] } },
+        "betaT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [-0.5] } },
+        "biasT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.125] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [257, 1], "tolerance": 0.000001 },
+        "residualT": { "dtype": "float32", "shape": [257, 1], "tolerance": 0.000001 }
+      }
+    }
+  ]
+}