sync 2e7068faf55e

Browse files

Files changed (7) hide show

README.md +77 -0
build/webgpu/bench.json +259 -0
build/webgpu/manifest.json +1006 -0
build/webgpu/metadata.json +23 -0
build/webgpu/norm-skip-row-vec4.wgsl.jinja +129 -0
build/webgpu/norm-skip-row.wgsl.jinja +229 -0
build/webgpu/test.json +837 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# com.microsoft.SkipSimplifiedLayerNormalization
+`com.microsoft`  ·  ONNX Runtime contrib operator  ·  contrib since_version 1
+## Description
+Adds `input` and `skip` (plus optional `bias`), then applies RMS normalization scaled by `gamma`. The optional second output exposes the pre-normalization sum. The schema's training-only mean and inverse-standard-deviation outputs are not implemented.
+See the [ONNX Runtime `SkipSimplifiedLayerNormalization` contrib-operator spec](https://github.com/microsoft/onnxruntime/blob/main/docs/ContribOperators.md#com.microsoft.SkipSimplifiedLayerNormalization) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `input` | `inputT` | `T` | — | — | Input tensor of shape `(token_count, hidden_size)` or `(batch, sequence, hidden_size)`, normalized over the last axis. | required |
+| `skip` | `skipT` | `T` | — | — | Residual tensor of the same shape as `input`, added before normalization. | required |
+| `gamma` | `gammaT` | `T` | `1` | — | 1-D scale tensor with shape `(hidden_size)` applied after normalization. | required |
+| `bias` | `biasT` | `T` | `1` | — | Optional 1-D bias tensor with shape `(hidden_size)` added to the `input + skip` sum. | 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 optional `bias` 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 mean square before taking the square root. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16` |
+## Device requirements
+Some implementation variants require `shader-f16`. These are route-specific capabilities, not package-wide requirements; availability also depends on the request shape and dtype.
+## 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.SkipSimplifiedLayerNormalization", { 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,259 @@

+{
+  "op": "com.microsoft.SkipSimplifiedLayerNormalization",
+  "cases": [
+    {
+      "name": "skip-rmsnorm-f32-256x128",
+      "preset": "smoke",
+      "vars": { "rows": 256, "hidden": 128 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "shape": [256, 128], "dtype": "float32", "dist": "normal", "seed": 215, "scale": 0.2 },
+        "skipT": { "shape": [256, 128], "dtype": "float32", "dist": "normal", "seed": 216, "scale": 0.2 },
+        "gammaT": { "shape": [128], "dtype": "float32", "dist": "uniform", "seed": 217, "scale": 0.1, "offset": 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 + args.hidden)" }]
+      }
+    },
+    {
+      "name": "kimi-linear-decode-f32-1x2304",
+      "preset": "model",
+      "provenance": {
+        "notes": "Original Kimi Linear hidden_size=2304 decode residual-plus-RMSNorm shape; occurs twice per decoder layer."
+      },
+      "vars": { "rows": 1, "hidden": 2304 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "shape": [1, 2304], "dtype": "float32", "dist": "normal", "seed": 1201, "scale": 0.2 },
+        "skipT": { "shape": [1, 2304], "dtype": "float32", "dist": "normal", "seed": 1202, "scale": 0.2 },
+        "gammaT": { "shape": [2304], "dtype": "float32", "dist": "uniform", "seed": 1203, "scale": 0.1, "offset": 1 }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 2304], "dtype": "float32" },
+        "residualT": { "shape": [1, 2304], "dtype": "float32" }
+      },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "4 * (4 * args.rows * args.hidden + args.hidden)" }]
+      }
+    },
+    {
+      "name": "kimi-linear-prefill-f32-64x2304",
+      "preset": "model",
+      "provenance": {
+        "notes": "Original Kimi Linear hidden_size=2304 at the 64-token recurrent chunk boundary; this is a future vectorized-prefill target because the current graph replays T=1."
+      },
+      "vars": { "rows": 64, "hidden": 2304 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "shape": [64, 2304], "dtype": "float32", "dist": "normal", "seed": 1211, "scale": 0.2 },
+        "skipT": { "shape": [64, 2304], "dtype": "float32", "dist": "normal", "seed": 1212, "scale": 0.2 },
+        "gammaT": { "shape": [2304], "dtype": "float32", "dist": "uniform", "seed": 1213, "scale": 0.1, "offset": 1 }
+      },
+      "outputs": {
+        "outputT": { "shape": [64, 2304], "dtype": "float32" },
+        "residualT": { "shape": [64, 2304], "dtype": "float32" }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (4 * args.rows * args.hidden + args.hidden)" }] }
+    },
+    {
+      "name": "kimi-linear-prefill-f32-512x2304",
+      "preset": "model",
+      "provenance": {
+        "notes": "Original Kimi Linear hidden_size=2304 representative future vectorized-prefill shape; the current graph replays T=1."
+      },
+      "vars": { "rows": 512, "hidden": 2304 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "shape": [512, 2304], "dtype": "float32", "dist": "normal", "seed": 1221, "scale": 0.2 },
+        "skipT": { "shape": [512, 2304], "dtype": "float32", "dist": "normal", "seed": 1222, "scale": 0.2 },
+        "gammaT": { "shape": [2304], "dtype": "float32", "dist": "uniform", "seed": 1223, "scale": 0.1, "offset": 1 }
+      },
+      "outputs": {
+        "outputT": { "shape": [512, 2304], "dtype": "float32" },
+        "residualT": { "shape": [512, 2304], "dtype": "float32" }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (4 * args.rows * args.hidden + args.hidden)" }] }
+    },
+    {
+      "name": "kimi-linear-prefill-f32-1024x2304",
+      "preset": "model",
+      "provenance": {
+        "notes": "Original Kimi Linear hidden_size=2304 representative future vectorized-prefill shape; the current graph replays T=1."
+      },
+      "vars": { "rows": 1024, "hidden": 2304 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "shape": [1024, 2304], "dtype": "float32", "dist": "normal", "seed": 1231, "scale": 0.2 },
+        "skipT": { "shape": [1024, 2304], "dtype": "float32", "dist": "normal", "seed": 1232, "scale": 0.2 },
+        "gammaT": { "shape": [2304], "dtype": "float32", "dist": "uniform", "seed": 1233, "scale": 0.1, "offset": 1 }
+      },
+      "outputs": {
+        "outputT": { "shape": [1024, 2304], "dtype": "float32" },
+        "residualT": { "shape": [1024, 2304], "dtype": "float32" }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (4 * args.rows * args.hidden + args.hidden)" }] }
+    },
+    {
+      "name": "skip-rmsnorm-f32-bias-256x128",
+      "preset": "smoke",
+      "vars": { "rows": 256, "hidden": 128 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "shape": [256, 128], "dtype": "float32", "dist": "normal", "seed": 216, "scale": 0.2 },
+        "skipT": { "shape": [256, 128], "dtype": "float32", "dist": "normal", "seed": 217, "scale": 0.2 },
+        "gammaT": { "shape": [128], "dtype": "float32", "dist": "uniform", "seed": 218, "scale": 0.1, "offset": 1 },
+        "biasT": { "shape": [128], "dtype": "float32", "dist": "normal", "seed": 219, "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 + 2 * args.hidden)" }]
+      }
+    },
+    {
+      "name": "skip-rmsnorm-f32-output-256x128",
+      "preset": "smoke",
+      "vars": { "rows": 256, "hidden": 128 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "shape": [256, 128], "dtype": "float32", "dist": "normal", "seed": 224, "scale": 0.2 },
+        "skipT": { "shape": [256, 128], "dtype": "float32", "dist": "normal", "seed": 225, "scale": 0.2 },
+        "gammaT": { "shape": [128], "dtype": "float32", "dist": "uniform", "seed": 226, "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-rmsnorm-f32-bias-output-256x128",
+      "preset": "smoke",
+      "vars": { "rows": 256, "hidden": 128 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "shape": [256, 128], "dtype": "float32", "dist": "normal", "seed": 227, "scale": 0.2 },
+        "skipT": { "shape": [256, 128], "dtype": "float32", "dist": "normal", "seed": 228, "scale": 0.2 },
+        "gammaT": { "shape": [128], "dtype": "float32", "dist": "uniform", "seed": 229, "scale": 0.1, "offset": 1 },
+        "biasT": { "shape": [128], "dtype": "float32", "dist": "normal", "seed": 230, "scale": 0.05 }
+      },
+      "outputs": { "outputT": { "shape": [256, 128], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (3 * args.rows * args.hidden + 2 * args.hidden)" }] }
+    },
+    {
+      "name": "skip-rmsnorm-f32-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" },
+        "residualT": { "shape": [4096, 768], "dtype": "float32" }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (4 * args.rows * args.hidden + args.hidden)" }] }
+    },
+    {
+      "name": "skip-rmsnorm-f32-8192x4096",
+      "preset": "model",
+      "vars": { "rows": 8192, "hidden": 4096 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "shape": [8192, 4096], "dtype": "float32", "dist": "normal", "seed": 419, "scale": 0.2 },
+        "skipT": { "shape": [8192, 4096], "dtype": "float32", "dist": "normal", "seed": 420, "scale": 0.2 },
+        "gammaT": { "shape": [4096], "dtype": "float32", "dist": "uniform", "seed": 421, "scale": 0.1, "offset": 1 }
+      },
+      "outputs": {
+        "outputT": { "shape": [8192, 4096], "dtype": "float32" },
+        "residualT": { "shape": [8192, 4096], "dtype": "float32" }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (4 * args.rows * args.hidden + args.hidden)" }] }
+    },
+    {
+      "name": "skip-rmsnorm-f32-output-4096x768",
+      "preset": "model",
+      "vars": { "rows": 4096, "hidden": 768 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "shape": [4096, 768], "dtype": "float32", "dist": "normal", "seed": 324, "scale": 0.2 },
+        "skipT": { "shape": [4096, 768], "dtype": "float32", "dist": "normal", "seed": 325, "scale": 0.2 },
+        "gammaT": { "shape": [768], "dtype": "float32", "dist": "uniform", "seed": 326, "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-rmsnorm-f32-bias-4096x768",
+      "preset": "model",
+      "vars": { "rows": 4096, "hidden": 768 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "shape": [4096, 768], "dtype": "float32", "dist": "normal", "seed": 320, "scale": 0.2 },
+        "skipT": { "shape": [4096, 768], "dtype": "float32", "dist": "normal", "seed": 321, "scale": 0.2 },
+        "gammaT": { "shape": [768], "dtype": "float32", "dist": "uniform", "seed": 322, "scale": 0.1, "offset": 1 },
+        "biasT": { "shape": [768], "dtype": "float32", "dist": "normal", "seed": 323, "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 + 2 * args.hidden)" }] }
+    },
+    {
+      "name": "skip-rmsnorm-f32-bias-output-4096x768",
+      "preset": "model",
+      "vars": { "rows": 4096, "hidden": 768 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "shape": [4096, 768], "dtype": "float32", "dist": "normal", "seed": 327, "scale": 0.2 },
+        "skipT": { "shape": [4096, 768], "dtype": "float32", "dist": "normal", "seed": 328, "scale": 0.2 },
+        "gammaT": { "shape": [768], "dtype": "float32", "dist": "uniform", "seed": 329, "scale": 0.1, "offset": 1 },
+        "biasT": { "shape": [768], "dtype": "float32", "dist": "normal", "seed": 330, "scale": 0.05 }
+      },
+      "outputs": { "outputT": { "shape": [4096, 768], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (3 * args.rows * args.hidden + 2 * args.hidden)" }] }
+    },
+    {
+      "name": "skip-rmsnorm-f32-scalar-fallback-8192x4095",
+      "preset": "stress",
+      "vars": { "rows": 8192, "hidden": 4095 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "shape": [8192, 4095], "dtype": "float32", "dist": "normal", "seed": 621, "scale": 0.2 },
+        "skipT": { "shape": [8192, 4095], "dtype": "float32", "dist": "normal", "seed": 622, "scale": 0.2 },
+        "gammaT": { "shape": [4095], "dtype": "float32", "dist": "uniform", "seed": 623, "scale": 0.1, "offset": 1 }
+      },
+      "outputs": {
+        "outputT": { "shape": [8192, 4095], "dtype": "float32", "dist": "empty" },
+        "residualT": { "shape": [8192, 4095], "dtype": "float32", "dist": "empty" }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (4 * args.rows * args.hidden + args.hidden)" }] }
+    },
+    {
+      "name": "skip-rmsnorm-f32-bias-scalar-fallback-8192x2049",
+      "preset": "stress",
+      "vars": { "rows": 8192, "hidden": 2049 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "shape": [8192, 2049], "dtype": "float32", "dist": "normal", "seed": 631, "scale": 0.2 },
+        "skipT": { "shape": [8192, 2049], "dtype": "float32", "dist": "normal", "seed": 632, "scale": 0.2 },
+        "gammaT": { "shape": [2049], "dtype": "float32", "dist": "uniform", "seed": 633, "scale": 0.1, "offset": 1 },
+        "biasT": { "shape": [2049], "dtype": "float32", "dist": "normal", "seed": 634, "scale": 0.05 }
+      },
+      "outputs": {
+        "outputT": { "shape": [8192, 2049], "dtype": "float32", "dist": "empty" },
+        "residualT": { "shape": [8192, 2049], "dtype": "float32", "dist": "empty" }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (4 * args.rows * args.hidden + 2 * args.hidden)" }] }
+    }
+  ]
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,1006 @@

+{
+  "domain": "com.microsoft",
+  "name": "SkipSimplifiedLayerNormalization",
+  "sinceVersion": 1,
+  "description": "Adds `input` and `skip` (plus optional `bias`), then applies RMS normalization scaled by `gamma`. The optional second output exposes the pre-normalization sum. The schema's training-only mean and inverse-standard-deviation outputs are not implemented.",
+  "inputs": [
+    {
+      "role": "input",
+      "dtype": "T",
+      "description": "Input tensor of shape `(token_count, hidden_size)` or `(batch, sequence, hidden_size)`, normalized over the last axis."
+    },
+    {
+      "role": "skip",
+      "dtype": "T",
+      "description": "Residual tensor of the same shape as `input`, added before normalization."
+    },
+    {
+      "role": "gamma",
+      "dtype": "T",
+      "rank": 1,
+      "description": "1-D scale tensor with shape `(hidden_size)` applied after normalization."
+    },
+    {
+      "role": "bias",
+      "dtype": "T",
+      "rank": 1,
+      "optional": true,
+      "description": "Optional 1-D bias tensor with shape `(hidden_size)` added to the `input + skip` sum."
+    }
+  ],
+  "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",
+      "optional": true,
+      "rank": "ranks.inputT",
+      "shape": "shapes.inputT",
+      "description": "Sum of `input`, `skip`, and optional `bias` before normalization, with the same shape as `input`."
+    }
+  ],
+  "attributes": { "epsilon": 9.999999960041972e-13 },
+  "attributeDescriptions": { "epsilon": "Non-negative epsilon added to the mean square 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" },
+    "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": "numel(shapes.inputT) / max(1, dim(shapes.inputT, -1))",
+    "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))))",
+    "rowDispatchFits": "rowCount <= device.limits.maxComputeWorkgroupsPerDimension * device.limits.maxComputeWorkgroupsPerDimension",
+    "normResourcesFit": "skipWg * 8 <= device.limits.maxComputeWorkgroupStorageSize and skipWgVec4 * 8 <= device.limits.maxComputeWorkgroupStorageSize",
+    "epsilonOk": "attrs.epsilon >= 0",
+    "coreContract": "epsilonOk and (ranks.inputT == 2 or ranks.inputT == 3) and ranks.skipT == ranks.inputT and ranks.gammaT == 1 and ranks.outputT == ranks.inputT and sameShape(shapes.inputT, shapes.skipT) and sameShape(shapes.outputT, shapes.inputT) and dim(shapes.inputT, -1) > 0 and dim(shapes.gammaT, 0) == dim(shapes.inputT, -1)",
+    "residualOutputContract": "present.residualT and sameShape(shapes.residualT, shapes.inputT)",
+    "outputOnlyContract": "not present.residualT",
+    "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",
+    "hasSubgroups": "device.features.has(\"subgroups\") and device.wgslLanguageFeatures.has(\"subgroup_id\")",
+    "hasF16": "device.features.has(\"shader-f16\")",
+    "no_bias_contract": "not present.biasT",
+    "f32_bias_contract": "false if not present.biasT else (ranks.biasT == 1 and tensorDtypes.biasT == \"float32\" and dim(shapes.biasT, 0) == hiddenSize)",
+    "f16_bias_contract": "false if not present.biasT else (ranks.biasT == 1 and tensorDtypes.biasT == \"float16\" and dim(shapes.biasT, 0) == hiddenSize)",
+    "f32_no_bias_residual_contract": "coreContract and residualOutputContract and f32ResidualDtypes and no_bias_contract",
+    "f32_bias_residual_contract": "coreContract and residualOutputContract and f32ResidualDtypes and f32_bias_contract",
+    "f16_no_bias_residual_contract": "hasF16 and coreContract and residualOutputContract and f16ResidualDtypes and no_bias_contract",
+    "f16_bias_residual_contract": "hasF16 and coreContract and residualOutputContract and f16ResidualDtypes and f16_bias_contract",
+    "f32_no_bias_output_contract": "coreContract and outputOnlyContract and f32MainDtypes and no_bias_contract",
+    "f32_bias_output_contract": "coreContract and outputOnlyContract and f32MainDtypes and f32_bias_contract",
+    "f16_no_bias_output_contract": "hasF16 and coreContract and outputOnlyContract and f16MainDtypes and no_bias_contract",
+    "f16_bias_output_contract": "hasF16 and coreContract and outputOnlyContract and f16MainDtypes and f16_bias_contract"
+  },
+  "bindingSets": {
+    "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": "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_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": "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": "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": "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": "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" }
+          ]
+        }
+      }
+    ],
+    "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": "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_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": "bias",
+        "arg": "biasT",
+        "semantic": "bias",
+        "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_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": "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": "gamma",
+        "arg": "gammaT",
+        "semantic": "gamma",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar",
+        "length": "$HIDDEN_LEN"
+      },
+      {
+        "name": "bias",
+        "arg": "biasT",
+        "semantic": "bias",
+        "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" }
+          ]
+        }
+      }
+    ]
+  },
+  "tunables": { "MAX_WORKGROUP_SIZE": 256 },
+  "variants": [
+    {
+      "id": "no_bias_vec4_f16",
+      "priority": 21,
+      "when": ["f16_no_bias_residual_contract", "vec4Aligned", "normResourcesFit", "rowDispatchFits"],
+      "constants": {
+        "scalar": "\"f16\"",
+        "vectorScalar": "\"vec4<f16>\"",
+        "hasBias": "\"no_bias\" == \"bias\"",
+        "HIDDEN_LEN": "hiddenSize / 4"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipSimplifiedLayerNormalization.Vec4",
+          "source": {
+            "shader": "norm-skip-row-vec4.wgsl.jinja",
+            "inputs": {
+              "simplified": true,
+              "hasBias": "\"no_bias\" == \"bias\"",
+              "hasBeta": false,
+              "writeResidualSum": true,
+              "usesF16": true,
+              "hidden": "hiddenSize",
+              "hiddenVec": "hiddenSize / 4",
+              "wg": "skipWgVec4",
+              "vecType": "\"vec4<f16>\"",
+              "useSubgroups": "hasSubgroups"
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "vec4_no_bias_residual",
+          "dispatch": { "workgroups": "rowCount" }
+        }
+      ]
+    },
+    {
+      "id": "no_bias_vec4",
+      "priority": 20,
+      "when": ["f32_no_bias_residual_contract", "vec4Aligned", "normResourcesFit", "rowDispatchFits"],
+      "constants": {
+        "scalar": "\"f32\"",
+        "vectorScalar": "\"vec4<f32>\"",
+        "hasBias": "\"no_bias\" == \"bias\"",
+        "HIDDEN_LEN": "hiddenSize / 4"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipSimplifiedLayerNormalization.Vec4",
+          "source": {
+            "shader": "norm-skip-row-vec4.wgsl.jinja",
+            "inputs": {
+              "simplified": true,
+              "hasBias": "\"no_bias\" == \"bias\"",
+              "hasBeta": false,
+              "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": "no_bias_output_only_vec4",
+      "priority": 20,
+      "when": ["f32_no_bias_output_contract", "vec4Aligned", "normResourcesFit", "rowDispatchFits"],
+      "constants": {
+        "scalar": "\"f32\"",
+        "vectorScalar": "\"vec4<f32>\"",
+        "hasBias": "\"no_bias\" == \"bias\"",
+        "HIDDEN_LEN": "hiddenSize / 4"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipSimplifiedLayerNormalization.Vec4OutputOnly",
+          "source": {
+            "shader": "norm-skip-row-vec4.wgsl.jinja",
+            "inputs": {
+              "simplified": true,
+              "hasBias": "\"no_bias\" == \"bias\"",
+              "hasBeta": false,
+              "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": "no_bias_output_only_vec4_f16",
+      "priority": 21,
+      "when": ["f16_no_bias_output_contract", "vec4Aligned", "normResourcesFit", "rowDispatchFits"],
+      "constants": {
+        "scalar": "\"f16\"",
+        "vectorScalar": "\"vec4<f16>\"",
+        "hasBias": "\"no_bias\" == \"bias\"",
+        "HIDDEN_LEN": "hiddenSize / 4"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipSimplifiedLayerNormalization.Vec4OutputOnly",
+          "source": {
+            "shader": "norm-skip-row-vec4.wgsl.jinja",
+            "inputs": {
+              "simplified": true,
+              "hasBias": "\"no_bias\" == \"bias\"",
+              "hasBeta": false,
+              "writeResidualSum": false,
+              "usesF16": true,
+              "hidden": "hiddenSize",
+              "hiddenVec": "hiddenSize / 4",
+              "wg": "skipWgVec4",
+              "vecType": "\"vec4<f16>\"",
+              "useSubgroups": "hasSubgroups"
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "vec4_no_bias_output_only",
+          "dispatch": { "workgroups": "rowCount" }
+        }
+      ]
+    },
+    {
+      "id": "no_bias",
+      "priority": 0,
+      "when": ["f32_no_bias_residual_contract", "normResourcesFit", "rowDispatchFits"],
+      "constants": {
+        "simplified": true,
+        "useSubgroups": false,
+        "hasBeta": false,
+        "writeResidualSum": true,
+        "hasBias": "\"no_bias\" == \"bias\"",
+        "scalar": "\"f32\"",
+        "hiddenSize": "dim(shapes.inputT, -1)",
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipSimplifiedLayerNormalization",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": "scalar_no_bias_residual",
+          "dispatch": { "workgroups": "rowCount" }
+        }
+      ]
+    },
+    {
+      "id": "no_bias_f16",
+      "requires": { "features": ["shader-f16"] },
+      "priority": 0,
+      "when": ["f16_no_bias_residual_contract", "normResourcesFit", "rowDispatchFits"],
+      "constants": {
+        "simplified": true,
+        "useSubgroups": false,
+        "hasBeta": false,
+        "writeResidualSum": true,
+        "hasBias": "\"no_bias\" == \"bias\"",
+        "scalar": "\"f16\"",
+        "usesF16": true,
+        "hiddenSize": "dim(shapes.inputT, -1)",
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipSimplifiedLayerNormalization",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": "scalar_no_bias_residual",
+          "dispatch": { "workgroups": "rowCount" }
+        }
+      ]
+    },
+    {
+      "id": "no_bias_output_only_f16",
+      "requires": { "features": ["shader-f16"] },
+      "priority": 0,
+      "when": ["f16_no_bias_output_contract", "normResourcesFit", "rowDispatchFits"],
+      "constants": {
+        "simplified": true,
+        "useSubgroups": false,
+        "hasBeta": false,
+        "writeResidualSum": false,
+        "hasBias": "\"no_bias\" == \"bias\"",
+        "scalar": "\"f16\"",
+        "usesF16": true,
+        "hiddenSize": "dim(shapes.inputT, -1)",
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipSimplifiedLayerNormalization.OutputOnly",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": "scalar_no_bias_output_only",
+          "dispatch": { "workgroups": "rowCount" }
+        }
+      ]
+    },
+    {
+      "id": "no_bias_output_only",
+      "priority": 0,
+      "when": ["f32_no_bias_output_contract", "normResourcesFit", "rowDispatchFits"],
+      "constants": {
+        "simplified": true,
+        "useSubgroups": false,
+        "hasBeta": false,
+        "writeResidualSum": false,
+        "hasBias": "\"no_bias\" == \"bias\"",
+        "scalar": "\"f32\"",
+        "hiddenSize": "dim(shapes.inputT, -1)",
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipSimplifiedLayerNormalization.OutputOnly",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": "scalar_no_bias_output_only",
+          "dispatch": { "workgroups": "rowCount" }
+        }
+      ]
+    },
+    {
+      "id": "bias_vec4_f16",
+      "priority": 21,
+      "when": ["f16_bias_residual_contract", "vec4Aligned", "normResourcesFit", "rowDispatchFits"],
+      "constants": {
+        "scalar": "\"f16\"",
+        "vectorScalar": "\"vec4<f16>\"",
+        "hasBias": "\"bias\" == \"bias\"",
+        "HIDDEN_LEN": "hiddenSize / 4"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipSimplifiedLayerNormalization.Vec4",
+          "source": {
+            "shader": "norm-skip-row-vec4.wgsl.jinja",
+            "inputs": {
+              "simplified": true,
+              "hasBias": "\"bias\" == \"bias\"",
+              "hasBeta": false,
+              "writeResidualSum": true,
+              "usesF16": true,
+              "hidden": "hiddenSize",
+              "hiddenVec": "hiddenSize / 4",
+              "wg": "skipWgVec4",
+              "vecType": "\"vec4<f16>\"",
+              "useSubgroups": "hasSubgroups"
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "vec4_bias_residual",
+          "dispatch": { "workgroups": "rowCount" }
+        }
+      ]
+    },
+    {
+      "id": "bias_vec4",
+      "priority": 20,
+      "when": ["f32_bias_residual_contract", "vec4Aligned", "normResourcesFit", "rowDispatchFits"],
+      "constants": {
+        "scalar": "\"f32\"",
+        "vectorScalar": "\"vec4<f32>\"",
+        "hasBias": "\"bias\" == \"bias\"",
+        "HIDDEN_LEN": "hiddenSize / 4"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipSimplifiedLayerNormalization.Vec4",
+          "source": {
+            "shader": "norm-skip-row-vec4.wgsl.jinja",
+            "inputs": {
+              "simplified": true,
+              "hasBias": "\"bias\" == \"bias\"",
+              "hasBeta": false,
+              "writeResidualSum": true,
+              "usesF16": false,
+              "hidden": "hiddenSize",
+              "hiddenVec": "hiddenSize / 4",
+              "wg": "skipWgVec4",
+              "vecType": "\"vec4<f32>\"",
+              "useSubgroups": "hasSubgroups"
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "vec4_bias_residual",
+          "dispatch": { "workgroups": "rowCount" }
+        }
+      ]
+    },
+    {
+      "id": "bias_output_only_vec4",
+      "priority": 20,
+      "when": ["f32_bias_output_contract", "vec4Aligned", "normResourcesFit", "rowDispatchFits"],
+      "constants": {
+        "scalar": "\"f32\"",
+        "vectorScalar": "\"vec4<f32>\"",
+        "hasBias": "\"bias\" == \"bias\"",
+        "HIDDEN_LEN": "hiddenSize / 4"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipSimplifiedLayerNormalization.Vec4OutputOnly",
+          "source": {
+            "shader": "norm-skip-row-vec4.wgsl.jinja",
+            "inputs": {
+              "simplified": true,
+              "hasBias": "\"bias\" == \"bias\"",
+              "hasBeta": false,
+              "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": "bias_output_only_vec4_f16",
+      "priority": 21,
+      "when": ["f16_bias_output_contract", "vec4Aligned", "normResourcesFit", "rowDispatchFits"],
+      "constants": {
+        "scalar": "\"f16\"",
+        "vectorScalar": "\"vec4<f16>\"",
+        "hasBias": "\"bias\" == \"bias\"",
+        "HIDDEN_LEN": "hiddenSize / 4"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipSimplifiedLayerNormalization.Vec4OutputOnly",
+          "source": {
+            "shader": "norm-skip-row-vec4.wgsl.jinja",
+            "inputs": {
+              "simplified": true,
+              "hasBias": "\"bias\" == \"bias\"",
+              "hasBeta": false,
+              "writeResidualSum": false,
+              "usesF16": true,
+              "hidden": "hiddenSize",
+              "hiddenVec": "hiddenSize / 4",
+              "wg": "skipWgVec4",
+              "vecType": "\"vec4<f16>\"",
+              "useSubgroups": "hasSubgroups"
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "vec4_bias_output_only",
+          "dispatch": { "workgroups": "rowCount" }
+        }
+      ]
+    },
+    {
+      "id": "bias",
+      "priority": 0,
+      "when": ["f32_bias_residual_contract", "normResourcesFit", "rowDispatchFits"],
+      "constants": {
+        "simplified": true,
+        "useSubgroups": false,
+        "hasBeta": false,
+        "writeResidualSum": true,
+        "hasBias": "\"bias\" == \"bias\"",
+        "scalar": "\"f32\"",
+        "hiddenSize": "dim(shapes.inputT, -1)",
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipSimplifiedLayerNormalization",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": "scalar_bias_residual",
+          "dispatch": { "workgroups": "rowCount" }
+        }
+      ]
+    },
+    {
+      "id": "bias_f16",
+      "requires": { "features": ["shader-f16"] },
+      "priority": 0,
+      "when": ["f16_bias_residual_contract", "normResourcesFit", "rowDispatchFits"],
+      "constants": {
+        "simplified": true,
+        "useSubgroups": false,
+        "hasBeta": false,
+        "writeResidualSum": true,
+        "hasBias": "\"bias\" == \"bias\"",
+        "scalar": "\"f16\"",
+        "usesF16": true,
+        "hiddenSize": "dim(shapes.inputT, -1)",
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipSimplifiedLayerNormalization",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": "scalar_bias_residual",
+          "dispatch": { "workgroups": "rowCount" }
+        }
+      ]
+    },
+    {
+      "id": "bias_output_only_f16",
+      "requires": { "features": ["shader-f16"] },
+      "priority": 0,
+      "when": ["f16_bias_output_contract", "normResourcesFit", "rowDispatchFits"],
+      "constants": {
+        "simplified": true,
+        "useSubgroups": false,
+        "hasBeta": false,
+        "writeResidualSum": false,
+        "hasBias": "\"bias\" == \"bias\"",
+        "scalar": "\"f16\"",
+        "usesF16": true,
+        "hiddenSize": "dim(shapes.inputT, -1)",
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipSimplifiedLayerNormalization.OutputOnly",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": "scalar_bias_output_only",
+          "dispatch": { "workgroups": "rowCount" }
+        }
+      ]
+    },
+    {
+      "id": "bias_output_only",
+      "priority": 0,
+      "when": ["f32_bias_output_contract", "normResourcesFit", "rowDispatchFits"],
+      "constants": {
+        "simplified": true,
+        "useSubgroups": false,
+        "hasBeta": false,
+        "writeResidualSum": false,
+        "hasBias": "\"bias\" == \"bias\"",
+        "scalar": "\"f32\"",
+        "hiddenSize": "dim(shapes.inputT, -1)",
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipSimplifiedLayerNormalization.OutputOnly",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": "scalar_bias_output_only",
+          "dispatch": { "workgroups": "rowCount" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,23 @@

+{
+  "name": "com.microsoft.SkipSimplifiedLayerNormalization",
+  "id": "_com_microsoft_skipsimplifiedlayernormalization_webgpu_6026e4e",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "Qjd8vGH/qoXU3rXEQhkz+xTQK2qSM5Kl4BHnWY/nzeQ=",
+      "manifest.json": "0QZ+D26SsY6uIHFcZV326mQMPoEYAClEUEifNI/9jms=",
+      "norm-skip-row-vec4.wgsl.jinja": "L7sH/FpRWVXbyaGm4QTjBAaLMacnNmwb0dbt2pKie5s=",
+      "norm-skip-row.wgsl.jinja": "bo/2mKHvO3Xkb7rbXnHkw/csvMsxID6bhdF8Wv7efKM=",
+      "test.json": "oQayex/aOYAhZXp9ESTPlwVGivSDeuEumSnU6kNNUR0="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "1.0",
+    "specialized": true,
+    "opPath": "ops/com.microsoft.SkipSimplifiedLayerNormalization"
+  }
+}

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

	@@ -0,0 +1,129 @@

+{% 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<f32, WG>;
+fn reduce_scalar(value: f32{% if useSubgroups %}, sg_lane: u32, sg_id: u32, num_sg: u32{% else %}, tid: u32{% endif %}) -> f32 {
+{% if useSubgroups %}
+  let s = subgroupAdd(value);
+  if (num_sg == 1u) {
+    return s;
+  }
+  if (sg_lane == 0u) {
+    sg_partials[sg_id] = s;
+  }
+  workgroupBarrier();
+  var total = 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;
+  let skip_base = base;
+  var acc = 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 %});
+    acc = acc + dot(v, v);
+  }
+  let total = reduce_scalar(acc{% if useSubgroups %}, sg_lane, sg_id, num_sg{% else %}, tid{% endif %});
+  let row_inv = inverseSqrt(total / f32(HIDDEN) + params.epsilon);
+  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 %}
+    output[idx] = {{ source.vecType }}(residual * row_inv * vec4<f32>(gamma[i]));
+  }
+}

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

	@@ -0,0 +1,229 @@

+{% 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 usesF16 %}
+enable f16;
+{% endif %}
+{% 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,837 @@

+{
+  "op": "com.microsoft.SkipSimplifiedLayerNormalization",
+  "cases": [
+    {
+      "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 }
+        }
+      },
+      "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 }
+        }
+      },
+      "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 }
+        },
+        "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_residual_hidden_vector",
+      "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": [-5.0, -5.0, -5.0, -5.0, 3.0, 3.0, 3.0, 3.0] }
+        },
+        "gammaT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [10.0, -2.0, 3.0, 4.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] } },
+        "biasT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.25] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [3, 1], "tolerance": 0.00001 },
+        "residualT": { "dtype": "float32", "shape": [3, 1], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "large_values_scaled_by_rms",
+      "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] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 4], "tolerance": 0.000001 },
+        "residualT": { "dtype": "float32", "shape": [1, 4], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "zero_tokens_null_input",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/skiplayernorm_op_test.cc",
+        "test": "SkipLayerNormTest.SkipLayerNormNullInput",
+        "notes": "Same zero-token lowered shape as ORT's SkipLayerNormalization null-input case, applied to simplified RMS normalization."
+      },
+      "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] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [0, 4], "tolerance": 0 },
+        "residualT": { "dtype": "float32", "shape": [0, 4], "tolerance": 0 }
+      }
+    },
+    {
+      "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 }
+        }
+      },
+      "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 }
+        },
+        "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": "large_hidden_768_no_bias",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 768],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.003, "cosStep": 0.005 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [1, 768],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.007, "cosStep": 0.011 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [768],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "offset": 1.0, "sinStep": 0.013, "cosStep": 0.017 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 768], "tolerance": 0.0002 },
+        "residualT": { "dtype": "float32", "shape": [1, 768], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "large_hidden_768_bias",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 768],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.005, "cosStep": 0.009 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [1, 768],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.011, "cosStep": 0.015 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [768],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "offset": 1.0, "sinStep": 0.017, "cosStep": 0.021 }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [768],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.023, "cosStep": 0.027 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 768], "tolerance": 0.0002 },
+        "residualT": { "dtype": "float32", "shape": [1, 768], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "ort_simplified_batch1_flattened_tokens",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/skiplayernorm_op_test.cc",
+        "test": "SkipLayerNormTest.SkipSimplifiedLayerNormBatch1_Float16",
+        "notes": "ORT fp16 shape [1, 2, 4] is represented as float32 [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.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0] }
+        },
+        "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 },
+        "residualT": { "dtype": "float32", "shape": [2, 4], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "ort_simplified_batch1_bias_flattened_tokens",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/skiplayernorm_op_test.cc",
+        "test": "SkipLayerNormTest.SkipSimplifiedLayerNormBatch1_Bias_Float16",
+        "notes": "ORT fp16 shape [1, 1, 8] is represented as float32 [1, 8] tokens by this lowered kernel."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 8],
+          "data": {
+            "kind": "values",
+            "values": [0.12573242, -0.13208008, 0.640625, 0.10491943, -0.53564453, 0.36157227, 1.3037109, 0.94726562]
+          }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [1, 8],
+          "data": {
+            "kind": "values",
+            "values": [-0.70361328, -1.265625, -0.62304688, 0.041320801, -2.3242188, -0.21875, -1.2460938, -0.73242188]
+          }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": {
+            "kind": "values",
+            "values": [0.94580078, 0.96826172, 1.0410156, 1.1044922, 0.98730469, 1.1367188, 0.93359375, 1.0351562]
+          }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": {
+            "kind": "values",
+            "values": [0.45166016, 0.04699707, -0.37182617, -0.4609375, -0.22888184, 0.11010742, -0.50488281, -0.10461426]
+          }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 8], "tolerance": 0.00002 },
+        "residualT": { "dtype": "float32", "shape": [1, 8], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "output_only_no_bias",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 2.0, -1.0, 0.5, -0.5] }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0.5, -1.0, 1.0, 0.0, 0.0, 1.0, -0.5, 0.5] }
+        },
+        "gammaT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 2.0, -1.0, 0.5] } }
+      },
+      "outputs": {
+        "outputT": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "tolerance": 0.00002,
+          "data": {
+            "kind": "values",
+            "values": [0.5052908659501396, 0.6737211546001862, -1.3474423092003724, 0.6737211546001862, 1.9999900000749995, 0.0, 0.0, 0.0]
+          }
+        }
+      }
+    },
+    {
+      "name": "no_bias_hidden128_compact",
+      "provenance": {
+        "notes": "Compact hidden-size-128 companion for the output-only bench family; keeps the same row width under reference-computed full outputs."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [16, 128],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.017, "cosStep": 0.031 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [16, 128],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.011, "cosStep": 0.023 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "offset": 1.0, "sinStep": 0.007, "cosStep": 0.041 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [16, 128], "tolerance": 0.0002, "relTolerance": 0.0002 },
+        "residualT": { "dtype": "float32", "shape": [16, 128], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "output_only_bias",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0.25, -0.5, 1.0, 2.0, -1.0, 0.5, 1.5, -2.0] }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0.75, 0.25, -0.5, 1.0, 1.0, -1.0, 0.5, 0.25] }
+        },
+        "gammaT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.5, -1.0, 0.5, 2.0] } },
+        "biasT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.1, -0.2, 0.3, 0.4] } }
+      },
+      "outputs": {
+        "outputT": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "tolerance": 0.00002,
+          "data": {
+            "kind": "values",
+            "values": [0.8944258768765443, 0.24393433005723933, 0.21683051560643496, 3.686118765309394, 0.10873177077789489, 0.5074149302968427, 0.8336102426305274, -1.957171874002108]
+          }
+        }
+      }
+    },
+    {
+      "name": "f16_no_bias_hidden768_decode",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float16",
+          "shape": [2, 768],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.029 }
+        },
+        "skipT": {
+          "dtype": "float16",
+          "shape": [2, 768],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.019, "cosStep": 0.037 }
+        },
+        "gammaT": {
+          "dtype": "float16",
+          "shape": [768],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "offset": 1.0, "sinStep": 0.007, "cosStep": 0.041 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 768], "tolerance": 0.01 },
+        "residualT": { "dtype": "float16", "shape": [2, 768], "tolerance": 0.01 }
+      }
+    },
+    {
+      "name": "f16_bias_hidden2048",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float16",
+          "shape": [2, 2048],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.005, "cosStep": 0.009 }
+        },
+        "skipT": {
+          "dtype": "float16",
+          "shape": [2, 2048],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.011, "cosStep": 0.015 }
+        },
+        "gammaT": {
+          "dtype": "float16",
+          "shape": [2048],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "offset": 1.0, "sinStep": 0.017, "cosStep": 0.021 }
+        },
+        "biasT": {
+          "dtype": "float16",
+          "shape": [2048],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.023, "cosStep": 0.027 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 2048], "tolerance": 0.01 },
+        "residualT": { "dtype": "float16", "shape": [2, 2048], "tolerance": 0.01 }
+      }
+    },
+    {
+      "name": "unaligned_hidden770_bias_scalar_fallback",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 770],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.017, "cosStep": 0.031 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [2, 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.013, "cosStep": 0.017 }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [770],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.029, "cosStep": 0.007 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 770], "tolerance": 0.0002 },
+        "residualT": { "dtype": "float32", "shape": [2, 770], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "scalar_fallback_odd_hidden_large_common_offset_residual",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [5, 771],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "offset": 1000.0, "sinStep": 0.017, "cosStep": 0.031 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [5, 771],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.011, "cosStep": 0.023 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [771],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "offset": 1.0, "sinStep": 0.007, "cosStep": 0.041 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [5, 771], "tolerance": 0.0005 },
+        "residualT": { "dtype": "float32", "shape": [5, 771], "tolerance": 0.001 }
+      }
+    },
+    {
+      "name": "empty_rows_vec4_hidden8",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [0, 8], "data": { "kind": "values", "values": [] } },
+        "skipT": { "dtype": "float32", "shape": [0, 8], "data": { "kind": "values", "values": [] } },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0.3, 0.2, 4.0, 2.2, 1.0, -1.0, 0.5, 1.5] }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [0, 8], "tolerance": 0 },
+        "residualT": { "dtype": "float32", "shape": [0, 8], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "rank3_no_bias_residual_2x3x8",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 3, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [2, 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 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 3, 8], "tolerance": 0.00002 },
+        "residualT": { "dtype": "float32", "shape": [2, 3, 8], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "rank3_output_only_bias_2x2x12",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 2, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.27 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [2, 2, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.37, "scale": 0.2 }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.11, "scale": 0.1 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [2, 2, 12], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "f16_rank3_residual_2x2x768",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float16",
+          "shape": [2, 2, 768],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.031 }
+        },
+        "skipT": {
+          "dtype": "float16",
+          "shape": [2, 2, 768],
+          "data": { "kind": "fillFloat32", "sinStep": 0.011, "cosStep": 0.023 }
+        },
+        "gammaT": {
+          "dtype": "float16",
+          "shape": [768],
+          "data": { "kind": "fillFloat32", "sinStep": 0.007, "cosStep": 0.041, "scale": 0.2 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 2, 768], "tolerance": 0.01 },
+        "residualT": { "dtype": "float16", "shape": [2, 2, 768], "tolerance": 0.005 }
+      }
+    },
+    {
+      "name": "f16_no_bias_output_only_768",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float16",
+          "shape": [2, 768],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.029 }
+        },
+        "skipT": {
+          "dtype": "float16",
+          "shape": [2, 768],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.019, "cosStep": 0.037 }
+        },
+        "gammaT": {
+          "dtype": "float16",
+          "shape": [768],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "offset": 1.0, "sinStep": 0.007, "cosStep": 0.041 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float16", "shape": [2, 768], "tolerance": 0.01 } }
+    },
+    {
+      "name": "f16_bias_output_only_2048",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float16",
+          "shape": [2, 2048],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.005, "cosStep": 0.009 }
+        },
+        "skipT": {
+          "dtype": "float16",
+          "shape": [2, 2048],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.011, "cosStep": 0.015 }
+        },
+        "gammaT": {
+          "dtype": "float16",
+          "shape": [2048],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "offset": 1.0, "sinStep": 0.017, "cosStep": 0.021 }
+        },
+        "biasT": {
+          "dtype": "float16",
+          "shape": [2048],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.023, "cosStep": 0.027 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float16", "shape": [2, 2048], "tolerance": 0.01 } }
+    },
+    {
+      "name": "f16_rank3_output_only_unaligned_2x2x767",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float16",
+          "shape": [2, 2, 767],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.031 }
+        },
+        "skipT": {
+          "dtype": "float16",
+          "shape": [2, 2, 767],
+          "data": { "kind": "fillFloat32", "sinStep": 0.011, "cosStep": 0.023 }
+        },
+        "gammaT": {
+          "dtype": "float16",
+          "shape": [767],
+          "data": { "kind": "fillFloat32", "sinStep": 0.007, "cosStep": 0.041, "scale": 0.2 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float16", "shape": [2, 2, 767], "tolerance": 0.01 } }
+    },
+    {
+      "name": "output_only_no_bias_unaligned_hidden6_scalar_row",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 6],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.017, "cosStep": 0.031 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [2, 6],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.011, "cosStep": 0.023 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "offset": 1.0, "sinStep": 0.007, "cosStep": 0.041 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [2, 6], "tolerance": 0.000002 } }
+    },
+    {
+      "name": "output_only_bias_unaligned_hidden6_scalar_row",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 6],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.019, "cosStep": 0.037 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [2, 6],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.029 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "offset": 1.0, "sinStep": 0.007, "cosStep": 0.041 }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.023, "cosStep": 0.027 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [2, 6], "tolerance": 0.000002 } }
+    },
+    {
+      "name": "f16_no_bias_unaligned_hidden6_scalar_row",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float16",
+          "shape": [2, 6],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.017, "cosStep": 0.031 }
+        },
+        "skipT": {
+          "dtype": "float16",
+          "shape": [2, 6],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.011, "cosStep": 0.023 }
+        },
+        "gammaT": {
+          "dtype": "float16",
+          "shape": [6],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "offset": 1.0, "sinStep": 0.007, "cosStep": 0.041 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 6], "tolerance": 0.01 },
+        "residualT": { "dtype": "float16", "shape": [2, 6], "tolerance": 0.005 }
+      }
+    },
+    {
+      "name": "f16_bias_unaligned_hidden6_scalar_row",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float16",
+          "shape": [2, 6],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.019, "cosStep": 0.037 }
+        },
+        "skipT": {
+          "dtype": "float16",
+          "shape": [2, 6],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.029 }
+        },
+        "gammaT": {
+          "dtype": "float16",
+          "shape": [6],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "offset": 1.0, "sinStep": 0.007, "cosStep": 0.041 }
+        },
+        "biasT": {
+          "dtype": "float16",
+          "shape": [6],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.023, "cosStep": 0.027 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 6], "tolerance": 0.01 },
+        "residualT": { "dtype": "float16", "shape": [2, 6], "tolerance": 0.005 }
+      }
+    },
+    {
+      "name": "f16_bias_output_only_unaligned_hidden6_scalar_row",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float16",
+          "shape": [2, 6],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.005, "cosStep": 0.009 }
+        },
+        "skipT": {
+          "dtype": "float16",
+          "shape": [2, 6],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.011, "cosStep": 0.015 }
+        },
+        "gammaT": {
+          "dtype": "float16",
+          "shape": [6],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "offset": 1.0, "sinStep": 0.017, "cosStep": 0.021 }
+        },
+        "biasT": {
+          "dtype": "float16",
+          "shape": [6],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.023, "cosStep": 0.027 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float16", "shape": [2, 6], "tolerance": 0.01 } }
+    }
+  ]
+}