sync 6fdf6301e2bb

Browse files

Files changed (6) hide show

README.md +89 -0
build/webgpu/bench.json +283 -0
build/webgpu/gated-rms-norm-rows.wgsl.jinja +161 -0
build/webgpu/manifest.json +238 -0
build/webgpu/metadata.json +28 -0
build/webgpu/test.json +881 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# com.microsoft.GatedRMSNorm
+`com.microsoft`  ·  ONNX Runtime contrib operator  ·  contrib since_version 1
+## Description
+Gated RMS normalization as used by Mamba2 and gated DeltaNet attention outputs: `Y = X * rsqrt(mean(X^2) + epsilon) * scale * SiLU(gate)`. The mean of squares is taken over each contiguous group of `C` elements, where `C` is the element count of `scale`, so a per-head norm runs on a packed `(..., H * C)` tensor with no surrounding reshape. Rows are counted flat as the element count of `X` divided by `C` and any rank of at least 1 is accepted. Every intermediate including SiLU is float32 whatever the tensor type; only the store narrows. Bfloat16 is not implemented.
+See the [ONNX Runtime `GatedRMSNorm` contrib-operator spec](https://github.com/microsoft/onnxruntime/blob/main/docs/ContribOperators.md#com.microsoft.GatedRMSNorm) for the reference semantics.
+## Inputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `xT` | `X` | `T` | — | — | Input activations with shape `(..., H * C)`. Any rank of at least 1 is accepted and the leading axes are folded into the row count. | required |
+| `scaleT` | `scale` | `T` | — | — | Normalization weight. Its element count is the normalization span `C`; the standard spelling is `(C)`, and the reshaped `(1, C)` and `(C, 1)` spellings of the same vector are accepted because upstream measures the span as an element count. | required |
+| `gateT` | `gate` | `T` | — | — | Gate activations with exactly the same shape as `X`; SiLU is applied to them internally. | required |
+## Outputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `yT` | `Y` | `T` | same as `xT` | same as `xT` | Normalized, scaled and gated activations with the same shape and dtype as `X`. | required |
+## Attributes
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `epsilon` | `0.00001` | Constant added to the mean of squares before the reciprocal square root. The standard default is `1e-5`; note that the Hugging Face gated-RMS archetype defaults to `1e-6` instead, so a value must be carried explicitly when porting between the two. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16` |
+## Implementation variants
+One implementation is selected per call from the device capabilities, the request shapes and the dtypes; these notes say what each one covers.
+- `subgroup_rows_vec4` — One aligned lane group per normalization row inside a subgroup, with the row held in registers. The fold is a segmented butterfly over the low lane bits, so several rows share a subgroup with no workgroup memory, no barrier and no idle lanes beyond the row's own power-of-two padding.
+- `packed_rows_vec4` — Portable vec4 route: the workgroup is split into a power-of-two lane group per row and each group folds its own contiguous slice of the shared array, so a short row does not idle the workgroup and no subgroup support is required. It also carries rows too wide to stage in registers, which it walks twice.
+- `subgroup_rows_scalar` — The barrier-free lane-group schedule for a normalization span that is not a multiple of four: the same segmented butterfly over scalar loads.
+- `packed_rows_scalar` — Scalar fallback for a span that is not a multiple of four, and the route every device can select. It keeps the lane-group-per-row schedule so a narrow span still fills the workgroup.
+- `subgroup_wide_rows_vec4` — Subgroup partials for a row spanning several whole subgroups; every lane reads the short row sum after one barrier.
+- `subgroup_wide_rows_scalar` — Subgroup partials for a row spanning several whole subgroups; every lane reads the short row sum after one barrier.
+## Device requirements
+Some implementation variants require `subgroups`. 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, per-variant templates, 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 cases
+- [`gated-rms-norm-rows.wgsl.jinja`](build/webgpu/gated-rms-norm-rows.wgsl.jinja)
+## Use with `@huggingface/kernels`
+```sh
+npm install --save-exact @huggingface/kernels@0.0.1-preview.3
+```
+Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
+The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+It follows the `v1` branch as fixes land. To pin exact artifact bytes, pass a 40-character commit `revision` instead of `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.GatedRMSNorm", { version: 1 });
+const { yT } = await kernel({
+  xT: { data: xTData, shape: [3, 4] },
+  scaleT: { data: scaleTData, shape: [1] },
+  gateT: { data: gateTData, shape: [3, 4] },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,283 @@

+{
+  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256, 512], "LANE_WORDS": [1, 2, 4, 8] },
+  "cases": [
+    {
+      "name": "gated-rmsnorm-decode-h16-d128-f32",
+      "preset": "smoke",
+      "provenance": {
+        "model": "Qwen3.5 / Qwen3-Next gated linear attention output norm",
+        "notes": "One decoded token: 16 value heads of 128. The whole operator is 24 KB of traffic, so the measurement is launch latency rather than bandwidth."
+      },
+      "vars": { "dtype": "float32", "count": 2048, "span": 128 },
+      "inputs": {
+        "xT": { "shape": [1, 1, 2048], "dtype": "float32", "dist": "normal", "seed": 501, "scale": 1 },
+        "scaleT": { "shape": [128], "dtype": "float32", "dist": "uniform", "seed": 502, "min": 0.2, "max": 1.8 },
+        "gateT": { "shape": [1, 1, 2048], "dtype": "float32", "dist": "normal", "seed": 503, "scale": 2 }
+      },
+      "outputs": { "yT": { "shape": [1, 1, 2048], "dtype": "float32" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(args.count * 3 + args.span) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "gated-rmsnorm-decode-h16-d128-f16",
+      "preset": "smoke",
+      "provenance": {
+        "model": "Qwen3.5 / Qwen3-Next gated linear attention output norm",
+        "notes": "The float16 decode step, the dtype the shipped model graphs use."
+      },
+      "vars": { "dtype": "float16", "count": 2048, "span": 128 },
+      "inputs": {
+        "xT": { "shape": [1, 1, 2048], "dtype": "float16", "dist": "normal", "seed": 501, "scale": 1 },
+        "scaleT": { "shape": [128], "dtype": "float16", "dist": "uniform", "seed": 502, "min": 0.2, "max": 1.8 },
+        "gateT": { "shape": [1, 1, 2048], "dtype": "float16", "dist": "normal", "seed": 503, "scale": 2 }
+      },
+      "outputs": { "yT": { "shape": [1, 1, 2048], "dtype": "float16" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(args.count * 3 + args.span) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "gated-rmsnorm-prefill-s512-h16-d128-f32",
+      "preset": "smoke",
+      "provenance": {
+        "model": "Qwen3.5 / Qwen3-Next gated linear attention output norm",
+        "notes": "Measures gated RMS normalization over a 512-token prefill: 8192 rows (512 tokens x 16 heads) of 128 float32 channels each."
+      },
+      "vars": { "dtype": "float32", "count": 1048576, "span": 128 },
+      "inputs": {
+        "xT": { "shape": [1, 512, 2048], "dtype": "float32", "dist": "normal", "seed": 511, "scale": 1 },
+        "scaleT": { "shape": [128], "dtype": "float32", "dist": "uniform", "seed": 512, "min": 0.2, "max": 1.8 },
+        "gateT": { "shape": [1, 512, 2048], "dtype": "float32", "dist": "normal", "seed": 513, "scale": 2 }
+      },
+      "outputs": { "yT": { "shape": [1, 512, 2048], "dtype": "float32" } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "(args.count * 3 + args.span) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "gated-rmsnorm-prefill-s512-h16-d128-f16",
+      "preset": "smoke",
+      "provenance": {
+        "model": "Qwen3.5 / Qwen3-Next gated linear attention output norm",
+        "notes": "The float16 prefill of the same 8192 spans."
+      },
+      "vars": { "dtype": "float16", "count": 1048576, "span": 128 },
+      "inputs": {
+        "xT": { "shape": [1, 512, 2048], "dtype": "float16", "dist": "normal", "seed": 511, "scale": 1 },
+        "scaleT": { "shape": [128], "dtype": "float16", "dist": "uniform", "seed": 512, "min": 0.2, "max": 1.8 },
+        "gateT": { "shape": [1, 512, 2048], "dtype": "float16", "dist": "normal", "seed": 513, "scale": 2 }
+      },
+      "outputs": { "yT": { "shape": [1, 512, 2048], "dtype": "float16" } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "(args.count * 3 + args.span) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "gated-rmsnorm-single-group-d1536-f16",
+      "preset": "smoke",
+      "provenance": {
+        "notes": "Measures gated RMS normalization over 512 rows of 1536 float16 channels each, normalized as a single group (H=1) rather than split across multiple heads."
+      },
+      "vars": { "dtype": "float16", "count": 786432, "span": 1536 },
+      "inputs": {
+        "xT": { "shape": [1, 512, 1536], "dtype": "float16", "dist": "normal", "seed": 521, "scale": 1 },
+        "scaleT": { "shape": [1536], "dtype": "float16", "dist": "uniform", "seed": 522, "min": 0.2, "max": 1.8 },
+        "gateT": { "shape": [1, 512, 1536], "dtype": "float16", "dist": "normal", "seed": 523, "scale": 2 }
+      },
+      "outputs": { "yT": { "shape": [1, 512, 1536], "dtype": "float16" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(args.count * 3 + args.span) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "gated-rmsnorm-single-group-d5120-f32",
+      "preset": "smoke",
+      "provenance": {
+        "notes": "A 5120-wide span: wide enough that the row no longer fits in registers at this workgroup width, so the statistic and the store each walk it."
+      },
+      "vars": { "dtype": "float32", "count": 1310720, "span": 5120 },
+      "inputs": {
+        "xT": { "shape": [1, 256, 5120], "dtype": "float32", "dist": "normal", "seed": 531, "scale": 1 },
+        "scaleT": { "shape": [5120], "dtype": "float32", "dist": "uniform", "seed": 532, "min": 0.2, "max": 1.8 },
+        "gateT": { "shape": [1, 256, 5120], "dtype": "float32", "dist": "normal", "seed": 533, "scale": 2 }
+      },
+      "outputs": { "yT": { "shape": [1, 256, 5120], "dtype": "float32" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(args.count * 3 + args.span) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "gated-rmsnorm-small-span-d48-f16",
+      "preset": "smoke",
+      "provenance": {
+        "notes": "A 48-element span over the same tensor: 16384 very short rows, where a workgroup-per-row schedule would leave most of every workgroup idle."
+      },
+      "vars": { "dtype": "float16", "count": 786432, "span": 48 },
+      "inputs": {
+        "xT": { "shape": [1, 512, 1536], "dtype": "float16", "dist": "normal", "seed": 541, "scale": 1 },
+        "scaleT": { "shape": [48], "dtype": "float16", "dist": "uniform", "seed": 542, "min": 0.2, "max": 1.8 },
+        "gateT": { "shape": [1, 512, 1536], "dtype": "float16", "dist": "normal", "seed": 543, "scale": 2 }
+      },
+      "outputs": { "yT": { "shape": [1, 512, 1536], "dtype": "float16" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(args.count * 3 + args.span) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "gated-rmsnorm-span-six-scalar-f32",
+      "preset": "smoke",
+      "provenance": { "notes": "A normalization span of six values leaves a two-element tail after four-wide groups." },
+      "vars": { "dtype": "float32", "count": 786432, "span": 6 },
+      "inputs": {
+        "xT": { "shape": [1, 512, 1536], "dtype": "float32", "dist": "normal", "seed": 551, "scale": 1 },
+        "scaleT": { "shape": [6], "dtype": "float32", "dist": "uniform", "seed": 552, "min": 0.2, "max": 1.8 },
+        "gateT": { "shape": [1, 512, 1536], "dtype": "float32", "dist": "normal", "seed": 553, "scale": 2 }
+      },
+      "outputs": { "yT": { "shape": [1, 512, 1536], "dtype": "float32" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(args.count * 3 + args.span) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "gated-rmsnorm-prefill-s4096-h16-d128-f16",
+      "preset": "model",
+      "provenance": {
+        "model": "Qwen3.5 / Qwen3-Next gated linear attention output norm",
+        "notes": "A 4096-token prefill: 65536 spans, a working set far past cache so the bandwidth figure is a DRAM figure."
+      },
+      "vars": { "dtype": "float16", "count": 8388608, "span": 128 },
+      "inputs": {
+        "xT": { "shape": [1, 4096, 2048], "dtype": "float16", "dist": "normal", "seed": 561, "scale": 1 },
+        "scaleT": { "shape": [128], "dtype": "float16", "dist": "uniform", "seed": 562, "min": 0.2, "max": 1.8 },
+        "gateT": { "shape": [1, 4096, 2048], "dtype": "float16", "dist": "normal", "seed": 563, "scale": 2 }
+      },
+      "outputs": { "yT": { "shape": [1, 4096, 2048], "dtype": "float16" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(args.count * 3 + args.span) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "gated-rmsnorm-prefill-s2048-h16-d128-f32",
+      "preset": "model",
+      "provenance": {
+        "model": "Qwen3.5 / Qwen3-Next gated linear attention output norm",
+        "notes": "The float32 counterpart at the same 48 MB working set."
+      },
+      "vars": { "dtype": "float32", "count": 4194304, "span": 128 },
+      "inputs": {
+        "xT": { "shape": [1, 2048, 2048], "dtype": "float32", "dist": "normal", "seed": 571, "scale": 1 },
+        "scaleT": { "shape": [128], "dtype": "float32", "dist": "uniform", "seed": 572, "min": 0.2, "max": 1.8 },
+        "gateT": { "shape": [1, 2048, 2048], "dtype": "float32", "dist": "normal", "seed": 573, "scale": 2 }
+      },
+      "outputs": { "yT": { "shape": [1, 2048, 2048], "dtype": "float32" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(args.count * 3 + args.span) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "gated-rmsnorm-short-c5-r8193-float32",
+      "preset": "smoke",
+      "provenance": {
+        "notes": "Measures gated RMS normalization over 8193 rows of 5 channels in float32, one row past a full 256-row group. Its smaller row count sets an overhead baseline against the float32 131,073-row sibling."
+      },
+      "vars": { "dtype": "float32", "count": 40965, "span": 5 },
+      "inputs": {
+        "xT": { "shape": [1, 8193, 5], "dtype": "float32", "dist": "normal", "seed": 551, "scale": 1 },
+        "scaleT": { "shape": [5], "dtype": "float32", "dist": "uniform", "seed": 552, "min": 0.2, "max": 1.8 },
+        "gateT": { "shape": [1, 8193, 5], "dtype": "float32", "dist": "normal", "seed": 553, "scale": 2 }
+      },
+      "outputs": { "yT": { "shape": [1, 8193, 5], "dtype": "float32" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(args.count * 3 + args.span) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "gated-rmsnorm-short-c5-r131073-float32",
+      "preset": "smoke",
+      "provenance": {
+        "notes": "Measures gated RMS normalization over 131,073 rows of 5 channels in float32, one row past a full 256-row group. Its larger row count isolates per-row cost from the 8193-row float32 sibling's per-dispatch overhead."
+      },
+      "vars": { "dtype": "float32", "count": 655365, "span": 5 },
+      "inputs": {
+        "xT": { "shape": [1, 131073, 5], "dtype": "float32", "dist": "normal", "seed": 551, "scale": 1 },
+        "scaleT": { "shape": [5], "dtype": "float32", "dist": "uniform", "seed": 552, "min": 0.2, "max": 1.8 },
+        "gateT": { "shape": [1, 131073, 5], "dtype": "float32", "dist": "normal", "seed": 553, "scale": 2 }
+      },
+      "outputs": { "yT": { "shape": [1, 131073, 5], "dtype": "float32" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(args.count * 3 + args.span) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "gated-rmsnorm-short-c5-r131073-float16",
+      "preset": "smoke",
+      "provenance": {
+        "notes": "Measures gated RMS normalization over 131,073 rows of 5 channels in float16, one row past a full 256-row group; compares against its float32 counterpart at the same shape."
+      },
+      "vars": { "dtype": "float16", "count": 655365, "span": 5 },
+      "inputs": {
+        "xT": { "shape": [1, 131073, 5], "dtype": "float16", "dist": "normal", "seed": 551, "scale": 1 },
+        "scaleT": { "shape": [5], "dtype": "float16", "dist": "uniform", "seed": 552, "min": 0.2, "max": 1.8 },
+        "gateT": { "shape": [1, 131073, 5], "dtype": "float16", "dist": "normal", "seed": 553, "scale": 2 }
+      },
+      "outputs": { "yT": { "shape": [1, 131073, 5], "dtype": "float16" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(args.count * 3 + args.span) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "gated-rmsnorm-short-c6-r131073-float16",
+      "preset": "smoke",
+      "provenance": {
+        "notes": "Measures gated RMS normalization over 131,073 rows of 6 channels in float16, one row past a full 256-row group."
+      },
+      "vars": { "dtype": "float16", "count": 786438, "span": 6 },
+      "inputs": {
+        "xT": { "shape": [1, 131073, 6], "dtype": "float16", "dist": "normal", "seed": 551, "scale": 1 },
+        "scaleT": { "shape": [6], "dtype": "float16", "dist": "uniform", "seed": 552, "min": 0.2, "max": 1.8 },
+        "gateT": { "shape": [1, 131073, 6], "dtype": "float16", "dist": "normal", "seed": 553, "scale": 2 }
+      },
+      "outputs": { "yT": { "shape": [1, 131073, 6], "dtype": "float16" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(args.count * 3 + args.span) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "gated-rmsnorm-short-c7-r131073-float32",
+      "preset": "smoke",
+      "provenance": {
+        "notes": "Measures gated RMS normalization over 131,073 rows of 7 channels in float32, one row past a full 256-row group; compares against its float16 counterpart at the same shape."
+      },
+      "vars": { "dtype": "float32", "count": 917511, "span": 7 },
+      "inputs": {
+        "xT": { "shape": [1, 131073, 7], "dtype": "float32", "dist": "normal", "seed": 551, "scale": 1 },
+        "scaleT": { "shape": [7], "dtype": "float32", "dist": "uniform", "seed": 552, "min": 0.2, "max": 1.8 },
+        "gateT": { "shape": [1, 131073, 7], "dtype": "float32", "dist": "normal", "seed": 553, "scale": 2 }
+      },
+      "outputs": { "yT": { "shape": [1, 131073, 7], "dtype": "float32" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(args.count * 3 + args.span) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "gated-rmsnorm-short-c7-r131073-float16",
+      "preset": "smoke",
+      "provenance": {
+        "notes": "Measures gated RMS normalization over 131,073 rows of 7 channels in float16, one row past a full 256-row group; compares against its float32 counterpart at the same shape."
+      },
+      "vars": { "dtype": "float16", "count": 917511, "span": 7 },
+      "inputs": {
+        "xT": { "shape": [1, 131073, 7], "dtype": "float16", "dist": "normal", "seed": 551, "scale": 1 },
+        "scaleT": { "shape": [7], "dtype": "float16", "dist": "uniform", "seed": 552, "min": 0.2, "max": 1.8 },
+        "gateT": { "shape": [1, 131073, 7], "dtype": "float16", "dist": "normal", "seed": 553, "scale": 2 }
+      },
+      "outputs": { "yT": { "shape": [1, 131073, 7], "dtype": "float16" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(args.count * 3 + args.span) * dtypeBytes(args.dtype)" }]
+      }
+    }
+  ]
+}

build/webgpu/gated-rms-norm-rows.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,161 @@

+{% if foldKind == "subgroup_full" or foldKind == "subgroup_butterfly" or foldKind == "subgroup_tree" %}
+enable subgroups;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+// Gated RMS normalization over each contiguous group of C elements of a row:
+//   Y = X * inverseSqrt(mean(X^2) + epsilon) * scale * SiLU(gate)
+// C is the scale length, so a per-head norm runs on a packed (..., H * C) tensor
+// with no surrounding reshape. Every intermediate is float32 and only the store
+// narrows, so an f16 tensor keeps an f32 statistic.
+//
+// Each normalization group owns a fixed LANE GROUP rather than a whole workgroup:
+// a 128-element group would otherwise leave most of a workgroup idle and still pay
+// a full-width fold. LANES_PER_ROW is the smallest power of two that covers the
+// row, capped by the workgroup (and, on the subgroup routes, by the subgroup) so
+// several rows are resident at once.
+const ROW_WORDS: u32 = {{ rowWords }}u;
+const NORM_SIZE: f32 = {{ normSize }}.0;
+const EPS: f32 = {{ epsilon }};
+const WG: u32 = {{ workgroupSize }}u;
+const LANES_PER_ROW: u32 = {{ lanesPerRow }}u;
+const ROWS_PER_GROUP: u32 = {{ rowsPerGroup }}u;
+{% if foldKind == "subgroup_tree" %}
+var<workgroup> partials: array<f32, {{ (workgroupSize / subgroupWidth) | int }}>;
+{% endif %}
+{% if foldKind == "workgroup_tree" %}
+// One contiguous LANES_PER_ROW slice per row. The fold never crosses a slice
+// boundary, so the rows sharing this array do not interfere.
+var<workgroup> partials: array<f32, WG>;
+{% endif %}
+{% if useVec4 %}
+fn silu4(v: vec4<f32>) -> vec4<f32> {
+  return v / (vec4<f32>(1.0) + exp(-v));
+}
+{% else %}
+fn silu(x: f32) -> f32 {
+  return x / (1.0 + exp(-x));
+}
+{% endif %}
+@compute @workgroup_size(WG)
+{% if foldKind == "subgroup_tree" %}
+fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(subgroup_invocation_id) sg_lane: u32, @builtin(subgroup_id) sg_id: u32) {
+  let tid = sg_id * {{ subgroupWidth }}u + sg_lane;
+{% else %}
+fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid: vec3<u32>) {
+{% endif %}
+  let lane = {{ "tid" if foldKind == "subgroup_tree" else "lid.x" }} % LANES_PER_ROW;
+  let localRow = {{ "tid" if foldKind == "subgroup_tree" else "lid.x" }} / LANES_PER_ROW;
+  // 2D-folded row index: wg.y carries the high bits past the dispatch's per-axis
+  // workgroup fold width.
+  let row = (wg.x + wg.y * {{ DISPATCH_FOLD_WIDTH }}u) * ROWS_PER_GROUP + localRow;
+  let rowValid = row < params.rows;
+  // A lane group past the last row folds an empty row and stores nothing, so the
+  // collectives and barriers below stay in uniform control flow.
+  let limit = select(0u, ROW_WORDS, rowValid);
+  let base = row * ROW_WORDS;
+  var acc = 0.0;
+{% if stageRow %}
+{% for i in range(wordsPerLane) %}
+{% if useVec4 %}
+  var xv{{ i }} = vec4<f32>(0.0);
+  var gv{{ i }} = vec4<f32>(0.0);
+{% else %}
+  var xv{{ i }} = 0.0;
+  var gv{{ i }} = 0.0;
+{% endif %}
+  let c{{ i }} = lane + {{ i }}u * LANES_PER_ROW;
+  if (c{{ i }} < limit) {
+{% if useVec4 %}
+    xv{{ i }} = vec4<f32>(x[base + c{{ i }}]);
+    gv{{ i }} = vec4<f32>(gate[base + c{{ i }}]);
+    acc = acc + dot(xv{{ i }}, xv{{ i }});
+{% else %}
+    xv{{ i }} = f32(x[base + c{{ i }}]);
+    gv{{ i }} = f32(gate[base + c{{ i }}]);
+    acc = fma(xv{{ i }}, xv{{ i }}, acc);
+{% endif %}
+  }
+{% endfor %}
+{% else %}
+  // The row does not fit in registers at this width, so the statistic pass and the
+  // store pass each walk it.
+  for (var c = lane; c < limit; c = c + LANES_PER_ROW) {
+{% if useVec4 %}
+    let v = vec4<f32>(x[base + c]);
+    acc = acc + dot(v, v);
+{% else %}
+    let v = f32(x[base + c]);
+    acc = fma(v, v, acc);
+{% endif %}
+  }
+{% endif %}
+{% if foldKind == "subgroup_full" %}
+  let sumSq = subgroupAdd(acc);
+{% elif foldKind == "subgroup_butterfly" %}
+  // Segmented butterfly: only the low bits of the lane index are exchanged, so each
+  // aligned group of LANES_PER_ROW lanes folds its own row and one subgroup carries
+  // several rows without a barrier or any workgroup memory.
+  var folded = acc;
+{% for step in range(foldSteps) %}
+  folded = folded + subgroupShuffleXor(folded, 1u << {{ step }}u);
+{% endfor %}
+  let sumSq = folded;
+{% elif foldKind == "subgroup_tree" %}
+  // Subgroup IDs define the row coordinates independently of local invocation
+  // ordering. Only subgroup leaders publish; the row's short final sum is read
+  // by every lane, avoiding a second broadcast barrier.
+  let subgroupPartial = subgroupAdd(acc);
+  if (sg_lane == 0u) {
+    partials[sg_id] = subgroupPartial;
+  }
+  workgroupBarrier();
+  let sumSq = {% for i in range(subgroupsPerRow) %}{{ "+ " if not loop.first else "" }}partials[localRow * {{ subgroupsPerRow }}u + {{ i }}u]{% endfor %};
+{% elif foldKind == "workgroup_tree" %}
+  partials[lid.x] = acc;
+  workgroupBarrier();
+  // LANES_PER_ROW is a power of two, so the halving stride is uniform across the
+  // whole workgroup and every lane reaches every barrier.
+  var stride = LANES_PER_ROW / 2u;
+  loop {
+    if (stride == 0u) {
+      break;
+    }
+    if (lane < stride) {
+      partials[lid.x] = partials[lid.x] + partials[lid.x + stride];
+    }
+    workgroupBarrier();
+    stride = stride / 2u;
+  }
+  let sumSq = partials[localRow * LANES_PER_ROW];
+{% else %}
+  // One lane owns the whole row, so the accumulator is already the row statistic.
+  let sumSq = acc;
+{% endif %}
+  let inv = inverseSqrt(sumSq / NORM_SIZE + EPS);
+{% if stageRow %}
+{% for i in range(wordsPerLane) %}
+  if (c{{ i }} < limit) {
+{% if useVec4 %}
+    y[base + c{{ i }}] = {{ vectorScalar }}(xv{{ i }} * inv * vec4<f32>(scale[c{{ i }}]) * silu4(gv{{ i }}));
+{% else %}
+    y[base + c{{ i }}] = {{ scalar }}(xv{{ i }} * inv * f32(scale[c{{ i }}]) * silu(gv{{ i }}));
+{% endif %}
+  }
+{% endfor %}
+{% else %}
+  for (var c = lane; c < limit; c = c + LANES_PER_ROW) {
+    let idx = base + c;
+{% if useVec4 %}
+    y[idx] = {{ vectorScalar }}(vec4<f32>(x[idx]) * inv * vec4<f32>(scale[c]) * silu4(vec4<f32>(gate[idx])));
+{% else %}
+    y[idx] = {{ scalar }}(f32(x[idx]) * inv * f32(scale[c]) * silu(f32(gate[idx])));
+{% endif %}
+  }
+{% endif %}
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,238 @@

+{
+  "domain": "com.microsoft",
+  "name": "GatedRMSNorm",
+  "sinceVersion": 1,
+  "inputs": {
+    "xT": { "onnx": "X", "dtype": "T" },
+    "scaleT": { "onnx": "scale", "dtype": "T" },
+    "gateT": { "onnx": "gate", "dtype": "T" }
+  },
+  "outputs": { "yT": { "onnx": "Y", "dtype": "T", "rank": "ranks.xT", "shape": "shapes.xT" } },
+  "attributes": { "epsilon": { "default": 0.00001 } },
+  "typeConstraints": { "T": ["float32", "float16"] },
+  "tunables": { "WORKGROUP_SIZE": { "default": 256 }, "LANE_WORDS": { "default": 2 } },
+  "derive": {
+    "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
+    "foldedDispatchCapacity": "min(device.limits.maxComputeWorkgroupsPerDimension, 65535) * min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
+    "normSize": "numel(shapes.scaleT)",
+    "lastDim": "dim(shapes.xT, ranks.xT - 1)",
+    "elementCount": "numel(shapes.xT)",
+    "rowCount": "elementCount / max(1, normSize)",
+    "dtypeOk": "tensorDtypes.xT == tensorDtypes.scaleT and tensorDtypes.xT == tensorDtypes.gateT and tensorDtypes.xT == tensorDtypes.yT and f16Ok(dtypes.T)",
+    "shapeOk": "ranks.xT >= 1 and ranks.scaleT >= 1 and normSize > 0 and lastDim % normSize == 0 and sameShape(shapes.gateT, shapes.xT) and sameShape(shapes.yT, shapes.xT)",
+    "contractOk": "dtypeOk and shapeOk",
+    "workgroupOk": "tunables.WORKGROUP_SIZE > 0 and tunables.WORKGROUP_SIZE <= deviceWorkgroupCap and pow2ceil(tunables.WORKGROUP_SIZE) == tunables.WORKGROUP_SIZE and tunables.WORKGROUP_SIZE * 4 <= device.limits.maxComputeWorkgroupStorageSize and tunables.LANE_WORDS > 0",
+    "vec4Ok": "normSize % 4 == 0",
+    "vecWords": "normSize / 4",
+    "vecPreferredLanes": "max(1, min(tunables.WORKGROUP_SIZE / 2, pow2ceil(max(1, ceilDiv(vecWords, max(1, tunables.LANE_WORDS))))))",
+    "vecWideLanes": "max(1, min(tunables.WORKGROUP_SIZE, pow2ceil(max(1, ceilDiv(vecWords, max(1, tunables.LANE_WORDS))))))",
+    "vecLanesPerRow": "pick([[ceilDiv(vecWords, vecPreferredLanes) <= 8, vecPreferredLanes]], vecWideLanes)",
+    "vecRowsPerGroup": "tunables.WORKGROUP_SIZE / max(1, vecLanesPerRow)",
+    "vecWordsPerLane": "ceilDiv(vecWords, max(1, vecLanesPerRow))",
+    "vecRowGroups": "ceilDiv(rowCount, max(1, vecRowsPerGroup))",
+    "scalarPreferredLanes": "1 if normSize <= 8 else (max(1, min(tunables.WORKGROUP_SIZE / 2, pow2ceil(max(1, ceilDiv(normSize, max(1, tunables.LANE_WORDS)))))))",
+    "scalarWideLanes": "max(1, min(tunables.WORKGROUP_SIZE, pow2ceil(max(1, ceilDiv(normSize, max(1, tunables.LANE_WORDS))))))",
+    "scalarLanesPerRow": "pick([[ceilDiv(normSize, scalarPreferredLanes) <= 8, scalarPreferredLanes]], scalarWideLanes)",
+    "scalarRowsPerGroup": "tunables.WORKGROUP_SIZE / max(1, scalarLanesPerRow)",
+    "scalarWordsPerLane": "ceilDiv(normSize, max(1, scalarLanesPerRow))",
+    "scalarRowGroups": "ceilDiv(rowCount, max(1, scalarRowsPerGroup))",
+    "subgroupShapeKnown": "device.features.has(\"subgroups\") and has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize > 0 and device.adapterInfo.subgroupMinSize <= device.adapterInfo.subgroupMaxSize and device.adapterInfo.subgroupMaxSize <= tunables.WORKGROUP_SIZE and tunables.WORKGROUP_SIZE % device.adapterInfo.subgroupMaxSize == 0",
+    "sgMin": "device.adapterInfo.subgroupMinSize if subgroupShapeKnown else 0",
+    "sgFullWidth": "subgroupShapeKnown and device.adapterInfo.subgroupMinSize == device.adapterInfo.subgroupMaxSize",
+    "sgVecLanesPerRow": "min(sgMin, vecLanesPerRow) if subgroupShapeKnown else 1",
+    "sgVecRowsPerGroup": "tunables.WORKGROUP_SIZE / max(1, sgVecLanesPerRow)",
+    "sgVecWordsPerLane": "ceilDiv(vecWords, max(1, sgVecLanesPerRow))",
+    "sgVecRowGroups": "ceilDiv(rowCount, max(1, sgVecRowsPerGroup))",
+    "sgScalarLanesPerRow": "min(sgMin, scalarLanesPerRow) if subgroupShapeKnown else 1",
+    "sgScalarRowsPerGroup": "tunables.WORKGROUP_SIZE / max(1, sgScalarLanesPerRow)",
+    "sgScalarWordsPerLane": "ceilDiv(normSize, max(1, sgScalarLanesPerRow))",
+    "sgScalarRowGroups": "ceilDiv(rowCount, max(1, sgScalarRowsPerGroup))",
+    "vectorElement": "\"vec4<\" ~ dtypes.T ~ \">\"",
+    "workgroupSize": "tunables.WORKGROUP_SIZE",
+    "epsilon": "attrs.epsilon",
+    "scalar": "dtypes.T",
+    "vectorScalar": "vectorElement"
+  },
+  "when": ["contractOk", "workgroupOk"],
+  "bindings": {
+    "x": { "arg": "xT", "elementType": "$vectorElement" },
+    "scale": { "arg": "scaleT", "elementType": "$vectorElement" },
+    "gate": { "arg": "gateT", "elementType": "$vectorElement" },
+    "y": { "arg": "yT", "elementType": "$vectorElement" },
+    "x_x_t": { "arg": "xT", "name": "x", "elementType": "$scalar" },
+    "scale_main": { "arg": "scaleT", "name": "scale", "elementType": "$scalar" },
+    "gate_main": { "arg": "gateT", "name": "gate", "elementType": "$scalar" },
+    "y_y_t": { "arg": "yT", "name": "y", "elementType": "$scalar" },
+    "params": { "struct": [{ "name": "rows", "type": "u32", "value": "rowCount" }] }
+  },
+  "variants": [
+    {
+      "id": "subgroup_rows_vec4",
+      "priority": 30,
+      "when": ["vec4Ok", "subgroupShapeKnown", "sgVecWordsPerLane >= 1", "sgVecWordsPerLane <= 8", "sgVecRowGroups <= foldedDispatchCapacity"],
+      "requires": { "features": ["subgroups"] },
+      "derive": {
+        "foldKind": "\"none\" if sgVecLanesPerRow <= 1 else (\"subgroup_full\" if (sgFullWidth and sgVecLanesPerRow == sgMin) else \"subgroup_butterfly\")",
+        "foldSteps": "log2ceil(sgVecLanesPerRow)",
+        "useVec4": "true",
+        "stageRow": "true",
+        "rowWords": "vecWords",
+        "wordsPerLane": "sgVecWordsPerLane",
+        "lanesPerRow": "sgVecLanesPerRow",
+        "rowsPerGroup": "sgVecRowsPerGroup"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "GatedRMSNorm.SubgroupRowsVec4",
+          "shader": "gated-rms-norm-rows.wgsl.jinja",
+          "bindings": ["x", "scale", "gate", "y", "params"],
+          "dispatch": {
+            "x": "min(sgVecRowGroups, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(sgVecRowGroups, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "packed_rows_vec4",
+      "priority": 20,
+      "when": ["vec4Ok", "vecRowGroups <= foldedDispatchCapacity"],
+      "derive": {
+        "foldKind": "\"none\" if vecLanesPerRow <= 1 else \"workgroup_tree\"",
+        "foldSteps": "log2ceil(vecLanesPerRow)",
+        "useVec4": "true",
+        "stageRow": "vecWordsPerLane <= 8",
+        "rowWords": "vecWords",
+        "wordsPerLane": "vecWordsPerLane",
+        "lanesPerRow": "vecLanesPerRow",
+        "rowsPerGroup": "vecRowsPerGroup"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "GatedRMSNorm.PackedRowsVec4",
+          "shader": "gated-rms-norm-rows.wgsl.jinja",
+          "bindings": ["x", "scale", "gate", "y", "params"],
+          "dispatch": {
+            "x": "min(vecRowGroups, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(vecRowGroups, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "subgroup_rows_scalar",
+      "priority": 10,
+      "when": ["subgroupShapeKnown", "sgScalarWordsPerLane >= 1", "sgScalarWordsPerLane <= 8", "sgScalarRowGroups <= foldedDispatchCapacity"],
+      "requires": { "features": ["subgroups"] },
+      "derive": {
+        "foldKind": "\"none\" if sgScalarLanesPerRow <= 1 else (\"subgroup_full\" if (sgFullWidth and sgScalarLanesPerRow == sgMin) else \"subgroup_butterfly\")",
+        "foldSteps": "log2ceil(sgScalarLanesPerRow)",
+        "useVec4": "false",
+        "stageRow": "true",
+        "rowWords": "normSize",
+        "wordsPerLane": "sgScalarWordsPerLane",
+        "lanesPerRow": "sgScalarLanesPerRow",
+        "rowsPerGroup": "sgScalarRowsPerGroup"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "GatedRMSNorm.SubgroupRowsScalar",
+          "shader": "gated-rms-norm-rows.wgsl.jinja",
+          "bindings": ["x_x_t", "scale_main", "gate_main", "y_y_t", "params"],
+          "dispatch": {
+            "x": "min(sgScalarRowGroups, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(sgScalarRowGroups, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "packed_rows_scalar",
+      "priority": 0,
+      "when": ["scalarRowGroups <= foldedDispatchCapacity"],
+      "derive": {
+        "foldKind": "\"none\" if scalarLanesPerRow <= 1 else \"workgroup_tree\"",
+        "foldSteps": "log2ceil(scalarLanesPerRow)",
+        "useVec4": "false",
+        "stageRow": "scalarWordsPerLane <= 8",
+        "rowWords": "normSize",
+        "wordsPerLane": "scalarWordsPerLane",
+        "lanesPerRow": "scalarLanesPerRow",
+        "rowsPerGroup": "scalarRowsPerGroup"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "GatedRMSNorm.PackedRowsScalar",
+          "shader": "gated-rms-norm-rows.wgsl.jinja",
+          "bindings": ["x_x_t", "scale_main", "gate_main", "y_y_t", "params"],
+          "dispatch": {
+            "x": "min(scalarRowGroups, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(scalarRowGroups, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "subgroup_wide_rows_vec4",
+      "priority": 25,
+      "when": ["vec4Ok", "vecRowGroups <= foldedDispatchCapacity", "dtypes.T == \"f32\"", "sgFullWidth", "device.wgslLanguageFeatures.has(\"subgroup_id\")", "vecLanesPerRow > sgMin", "vecLanesPerRow % max(1, sgMin) == 0"],
+      "requires": { "features": ["subgroups"] },
+      "derive": {
+        "foldKind": "\"subgroup_tree\"",
+        "useVec4": "true",
+        "stageRow": "vecWordsPerLane <= 8",
+        "rowWords": "vecWords",
+        "wordsPerLane": "vecWordsPerLane",
+        "lanesPerRow": "vecLanesPerRow",
+        "rowsPerGroup": "vecRowsPerGroup",
+        "subgroupWidth": "sgMin",
+        "subgroupsPerRow": "vecLanesPerRow / max(1, sgMin)"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "GatedRMSNorm.SubgroupWideRowsVec4",
+          "shader": "gated-rms-norm-rows.wgsl.jinja",
+          "bindings": ["x", "scale", "gate", "y", "params"],
+          "dispatch": {
+            "x": "min(vecRowGroups, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(vecRowGroups, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "subgroup_wide_rows_scalar",
+      "priority": 5,
+      "when": ["scalarRowGroups <= foldedDispatchCapacity", "dtypes.T == \"f32\"", "sgFullWidth", "device.wgslLanguageFeatures.has(\"subgroup_id\")", "scalarLanesPerRow > sgMin", "scalarLanesPerRow % max(1, sgMin) == 0"],
+      "requires": { "features": ["subgroups"] },
+      "derive": {
+        "foldKind": "\"subgroup_tree\"",
+        "useVec4": "false",
+        "stageRow": "scalarWordsPerLane <= 8",
+        "rowWords": "normSize",
+        "wordsPerLane": "scalarWordsPerLane",
+        "lanesPerRow": "scalarLanesPerRow",
+        "rowsPerGroup": "scalarRowsPerGroup",
+        "subgroupWidth": "sgMin",
+        "subgroupsPerRow": "scalarLanesPerRow / max(1, sgMin)"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "GatedRMSNorm.SubgroupWideRowsScalar",
+          "shader": "gated-rms-norm-rows.wgsl.jinja",
+          "bindings": ["x_x_t", "scale_main", "gate_main", "y_y_t", "params"],
+          "dispatch": {
+            "x": "min(scalarRowGroups, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(scalarRowGroups, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,28 @@

+{
+  "name": "com.microsoft.GatedRMSNorm",
+  "id": "_com_microsoft_gatedrmsnorm_webgpu_3fd919f",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "E5LQpx8I5PNcqbNTX4KyCxcXvNNp3cHMKCbDTk1JlFI=",
+      "gated-rms-norm-rows.wgsl.jinja": "FNlhlwGawl2E/vwWNng/ZX+vThDezoBS43E/hxT1/9k=",
+      "manifest.json": "/o1FvJuHt1tdlx9h75CH1UA2wgYdt880zc3UQmntRrg=",
+      "test.json": "o0Og3R+dIxLG/ZwxG7CgIS7tbiRjxtApNu9xFzFxcZg="
+    }
+  },
+  "provenance": { "kernel": { "sha": "6fdf6301e2bbcc2f03bf1eaf493b7ad55ef33afc", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.1",
+    "variants": {
+      "subgroup_rows_vec4": ["gated-rms-norm-rows.wgsl.jinja"],
+      "packed_rows_vec4": ["gated-rms-norm-rows.wgsl.jinja"],
+      "subgroup_rows_scalar": ["gated-rms-norm-rows.wgsl.jinja"],
+      "packed_rows_scalar": ["gated-rms-norm-rows.wgsl.jinja"],
+      "subgroup_wide_rows_vec4": ["gated-rms-norm-rows.wgsl.jinja"],
+      "subgroup_wide_rows_scalar": ["gated-rms-norm-rows.wgsl.jinja"]
+    }
+  }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,881 @@

+{
+  "cases": [
+    {
+      "name": "ort_float_per_head",
+      "provenance": {
+        "source": "ORT: onnxruntime/test/contrib_ops/linear_attention_gates_op_test.cc",
+        "test": "ContribOpGatedRMSNormTest.Float_PerHead",
+        "notes": "Upstream draws its tensors from a seeded mt19937 and recomputes the expectation in the test body, so the geometry and epsilon transfer but the values do not. Batch 1, sequence 4, 32 heads of 128."
+      },
+      "attrs": { "epsilon": 0.000001 },
+      "inputs": {
+        "xT": {
+          "dtype": "float32",
+          "shape": [1, 4, 4096],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29 }
+        },
+        "scaleT": {
+          "dtype": "float32",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.5, "offset": 1.0 }
+        },
+        "gateT": {
+          "dtype": "float32",
+          "shape": [1, 4, 4096],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.19, "scale": 2.0 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [1, 4, 4096], "tolerance": 0.000002, "relTolerance": 0.00002 } }
+    },
+    {
+      "name": "ort_float_single_group",
+      "provenance": {
+        "source": "ORT: onnxruntime/test/contrib_ops/linear_attention_gates_op_test.cc",
+        "test": "ContribOpGatedRMSNormTest.Float_SingleGroup",
+        "notes": "One normalization group per row (H = 1), the case where a per-channel and a shared-group weight index agree, so it cannot distinguish them on its own."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "xT": {
+          "dtype": "float32",
+          "shape": [2, 3, 512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41 }
+        },
+        "scaleT": {
+          "dtype": "float32",
+          "shape": [512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.5, "offset": 1.0 }
+        },
+        "gateT": {
+          "dtype": "float32",
+          "shape": [2, 3, 512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.05, "cosStep": 0.43, "scale": 2.0 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [2, 3, 512], "tolerance": 0.000002, "relTolerance": 0.00002 } }
+    },
+    {
+      "name": "ort_float16_per_head",
+      "provenance": {
+        "source": "ORT: onnxruntime/test/contrib_ops/linear_attention_gates_op_test.cc",
+        "test": "ContribOpGatedRMSNormTest.Float16_PerHead"
+      },
+      "attrs": { "epsilon": 0.000001 },
+      "inputs": {
+        "xT": {
+          "dtype": "float16",
+          "shape": [1, 4, 4096],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29 }
+        },
+        "scaleT": {
+          "dtype": "float16",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.5, "offset": 1.0 }
+        },
+        "gateT": {
+          "dtype": "float16",
+          "shape": [1, 4, 4096],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.19, "scale": 2.0 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [1, 4, 4096], "tolerance": 0.0005, "relTolerance": 0.002 } }
+    },
+    {
+      "name": "ort_float16_prefill",
+      "provenance": {
+        "source": "ORT: onnxruntime/test/contrib_ops/linear_attention_gates_op_test.cc",
+        "test": "ContribOpGatedRMSNormTest.Float16_Prefill",
+        "notes": "Batch 2, sequence 17: an odd token count so the row count is not a multiple of the rows a workgroup carries."
+      },
+      "attrs": { "epsilon": 0.000001 },
+      "inputs": {
+        "xT": {
+          "dtype": "float16",
+          "shape": [2, 17, 4096],
+          "data": { "kind": "fillFloat32", "sinStep": 0.061, "cosStep": 0.257 }
+        },
+        "scaleT": {
+          "dtype": "float16",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.5, "offset": 1.0 }
+        },
+        "gateT": {
+          "dtype": "float16",
+          "shape": [2, 17, 4096],
+          "data": { "kind": "fillFloat32", "sinStep": 0.09, "cosStep": 0.23, "scale": 2.0 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [2, 17, 4096], "tolerance": 0.0005, "relTolerance": 0.002 } }
+    },
+    {
+      "name": "ort_float16_small_norm_size",
+      "provenance": {
+        "source": "ORT: onnxruntime/test/contrib_ops/linear_attention_gates_op_test.cc",
+        "test": "ContribOpGatedRMSNormTest.Float16_SmallNormSize",
+        "notes": "A 48-element span: a multiple of four but not of the lane group a wide row would use, so the trailing lanes of the group fold nothing."
+      },
+      "attrs": { "epsilon": 0.000001 },
+      "inputs": {
+        "xT": {
+          "dtype": "float16",
+          "shape": [1, 4, 384],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.33 }
+        },
+        "scaleT": {
+          "dtype": "float16",
+          "shape": [48],
+          "data": { "kind": "fillFloat32", "sinStep": 0.41, "cosStep": 0.07, "scale": 0.5, "offset": 1.0 }
+        },
+        "gateT": {
+          "dtype": "float16",
+          "shape": [1, 4, 384],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.27, "scale": 2.0 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [1, 4, 384], "tolerance": 0.0005, "relTolerance": 0.002 } }
+    },
+    {
+      "name": "ort_float16_large_norm_size",
+      "provenance": {
+        "source": "ORT: onnxruntime/test/contrib_ops/linear_attention_gates_op_test.cc",
+        "test": "ContribOpGatedRMSNormTest.Float16_LargeNormSize"
+      },
+      "attrs": { "epsilon": 0.000001 },
+      "inputs": {
+        "xT": {
+          "dtype": "float16",
+          "shape": [1, 2, 3072],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.053 }
+        },
+        "scaleT": {
+          "dtype": "float16",
+          "shape": [1536],
+          "data": { "kind": "fillFloat32", "sinStep": 0.029, "cosStep": 0.013, "scale": 0.5, "offset": 1.0 }
+        },
+        "gateT": {
+          "dtype": "float16",
+          "shape": [1, 2, 3072],
+          "data": { "kind": "fillFloat32", "sinStep": 0.037, "cosStep": 0.071, "scale": 2.0 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [1, 2, 3072], "tolerance": 0.0008, "relTolerance": 0.003 } }
+    },
+    {
+      "name": "rank1_decode_row",
+      "inputs": {
+        "xT": {
+          "dtype": "float32",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.21, "cosStep": 0.39 }
+        },
+        "scaleT": {
+          "dtype": "float32",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.5, "offset": 1.0 }
+        },
+        "gateT": {
+          "dtype": "float32",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.47, "scale": 2.0 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [128], "tolerance": 0.000002, "relTolerance": 0.00002 } }
+    },
+    {
+      "name": "rank2_many_rows",
+      "inputs": {
+        "xT": {
+          "dtype": "float32",
+          "shape": [512, 256],
+          "data": { "kind": "fillFloat32", "sinStep": 0.043, "cosStep": 0.121 }
+        },
+        "scaleT": {
+          "dtype": "float32",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.5, "offset": 1.0 }
+        },
+        "gateT": {
+          "dtype": "float32",
+          "shape": [512, 256],
+          "data": { "kind": "fillFloat32", "sinStep": 0.083, "cosStep": 0.151, "scale": 2.0 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [512, 256], "tolerance": 0.000002, "relTolerance": 0.00002 } }
+    },
+    {
+      "name": "rank4_batched_multi_group",
+      "inputs": {
+        "xT": {
+          "dtype": "float16",
+          "shape": [2, 3, 4, 256],
+          "data": { "kind": "fillFloat32", "sinStep": 0.067, "cosStep": 0.139 }
+        },
+        "scaleT": {
+          "dtype": "float16",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.17, "scale": 0.5, "offset": 1.0 }
+        },
+        "gateT": {
+          "dtype": "float16",
+          "shape": [2, 3, 4, 256],
+          "data": { "kind": "fillFloat32", "sinStep": 0.101, "cosStep": 0.211, "scale": 2.0 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [2, 3, 4, 256], "tolerance": 0.0005, "relTolerance": 0.002 } }
+    },
+    {
+      "name": "scale_spelled_leading_one",
+      "inputs": {
+        "xT": {
+          "dtype": "float32",
+          "shape": [8, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.31, "cosStep": 0.09 }
+        },
+        "scaleT": {
+          "dtype": "float32",
+          "shape": [1, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.5, "offset": 1.0 }
+        },
+        "gateT": {
+          "dtype": "float32",
+          "shape": [8, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.15, "cosStep": 0.53, "scale": 2.0 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [8, 128], "tolerance": 0.000002, "relTolerance": 0.00002 } }
+    },
+    {
+      "name": "scale_spelled_trailing_one",
+      "inputs": {
+        "xT": {
+          "dtype": "float32",
+          "shape": [8, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.31, "cosStep": 0.09 }
+        },
+        "scaleT": {
+          "dtype": "float32",
+          "shape": [128, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.5, "offset": 1.0 }
+        },
+        "gateT": {
+          "dtype": "float32",
+          "shape": [8, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.15, "cosStep": 0.53, "scale": 2.0 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [8, 128], "tolerance": 0.000002, "relTolerance": 0.00002 } }
+    },
+    {
+      "name": "span_six_scalar_route",
+      "inputs": {
+        "xT": {
+          "dtype": "float32",
+          "shape": [5, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.27, "cosStep": 0.61 }
+        },
+        "scaleT": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.44, "cosStep": 0.22, "scale": 0.5, "offset": 1.0 }
+        },
+        "gateT": {
+          "dtype": "float32",
+          "shape": [5, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.35, "cosStep": 0.14, "scale": 2.0 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [5, 12], "tolerance": 0.000002, "relTolerance": 0.00002 } }
+    },
+    {
+      "name": "span_one_single_lane_fold",
+      "inputs": {
+        "xT": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.51, "cosStep": 0.13 }
+        },
+        "scaleT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.75] } },
+        "gateT": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.63, "scale": 2.0 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [3, 4], "tolerance": 0.000002, "relTolerance": 0.00002 } }
+    },
+    {
+      "name": "span_four_single_word",
+      "inputs": {
+        "xT": {
+          "dtype": "float32",
+          "shape": [16, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.47, "cosStep": 0.23 }
+        },
+        "scaleT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.5, 1.25, -0.75, 2.0] } },
+        "gateT": {
+          "dtype": "float32",
+          "shape": [16, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.71, "scale": 2.0 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [16, 4], "tolerance": 0.000002, "relTolerance": 0.00002 } }
+    },
+    {
+      "name": "span_four_single_word_f16",
+      "inputs": {
+        "xT": {
+          "dtype": "float16",
+          "shape": [16, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.47, "cosStep": 0.23 }
+        },
+        "scaleT": { "dtype": "float16", "shape": [4], "data": { "kind": "values", "values": [0.5, 1.25, -0.75, 2.0] } },
+        "gateT": {
+          "dtype": "float16",
+          "shape": [16, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.71, "scale": 2.0 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [16, 4], "tolerance": 0.0005, "relTolerance": 0.002 } }
+    },
+    {
+      "name": "epsilon_dominates_statistic",
+      "attrs": { "epsilon": 100 },
+      "inputs": {
+        "xT": {
+          "dtype": "float32",
+          "shape": [4, 16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.33, "cosStep": 0.77 }
+        },
+        "scaleT": {
+          "dtype": "float32",
+          "shape": [16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.21, "cosStep": 0.49, "scale": 0.5, "offset": 1.0 }
+        },
+        "gateT": {
+          "dtype": "float32",
+          "shape": [4, 16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.57, "cosStep": 0.17, "scale": 2.0 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [4, 16], "tolerance": 0.000002, "relTolerance": 0.00002 } }
+    },
+    {
+      "name": "gate_saturating_tails",
+      "inputs": {
+        "xT": {
+          "dtype": "float32",
+          "shape": [4, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.25, "cosStep": 0.55 }
+        },
+        "scaleT": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.15, "cosStep": 0.35, "scale": 0.5, "offset": 1.0 }
+        },
+        "gateT": {
+          "dtype": "float32",
+          "shape": [4, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.45, "cosStep": 0.85, "scale": 12.0 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [4, 64], "tolerance": 0.000002, "relTolerance": 0.00002 } }
+    },
+    {
+      "name": "span_2048_two_words_per_lane",
+      "inputs": {
+        "xT": {
+          "dtype": "float32",
+          "shape": [4, 2048],
+          "data": { "kind": "fillFloat32", "sinStep": 0.011, "cosStep": 0.047 }
+        },
+        "scaleT": {
+          "dtype": "float32",
+          "shape": [2048],
+          "data": { "kind": "fillFloat32", "sinStep": 0.019, "cosStep": 0.007, "scale": 0.5, "offset": 1.0 }
+        },
+        "gateT": {
+          "dtype": "float32",
+          "shape": [4, 2048],
+          "data": { "kind": "fillFloat32", "sinStep": 0.023, "cosStep": 0.059, "scale": 2.0 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [4, 2048], "tolerance": 0.000004, "relTolerance": 0.00004 } }
+    },
+    {
+      "name": "span_16384_streaming_route",
+      "inputs": {
+        "xT": {
+          "dtype": "float32",
+          "shape": [2, 16384],
+          "data": { "kind": "fillFloat32", "sinStep": 0.0031, "cosStep": 0.0079 }
+        },
+        "scaleT": {
+          "dtype": "float32",
+          "shape": [16384],
+          "data": { "kind": "fillFloat32", "sinStep": 0.0011, "cosStep": 0.0023, "scale": 0.5, "offset": 1.0 }
+        },
+        "gateT": {
+          "dtype": "float32",
+          "shape": [2, 16384],
+          "data": { "kind": "fillFloat32", "sinStep": 0.0043, "cosStep": 0.0017, "scale": 2.0 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [2, 16384], "tolerance": 0.00001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "span_16384_streaming_route_f16",
+      "inputs": {
+        "xT": {
+          "dtype": "float16",
+          "shape": [2, 16384],
+          "data": { "kind": "fillFloat32", "sinStep": 0.0031, "cosStep": 0.0079 }
+        },
+        "scaleT": {
+          "dtype": "float16",
+          "shape": [16384],
+          "data": { "kind": "fillFloat32", "sinStep": 0.0011, "cosStep": 0.0023, "scale": 0.5, "offset": 1.0 }
+        },
+        "gateT": {
+          "dtype": "float16",
+          "shape": [2, 16384],
+          "data": { "kind": "fillFloat32", "sinStep": 0.0043, "cosStep": 0.0017, "scale": 2.0 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [2, 16384], "tolerance": 0.0008, "relTolerance": 0.003 } }
+    },
+    {
+      "name": "wide_subgroup_span_1026_rows_3_wg_256",
+      "inputs": {
+        "xT": {
+          "dtype": "float32",
+          "shape": [3, 1026],
+          "data": { "kind": "fillFloat32", "sinStep": 0.011, "cosStep": 0.047 }
+        },
+        "scaleT": {
+          "dtype": "float32",
+          "shape": [1026],
+          "data": { "kind": "fillFloat32", "sinStep": 0.019, "cosStep": 0.007, "scale": 0.5, "offset": 1.0 }
+        },
+        "gateT": {
+          "dtype": "float32",
+          "shape": [3, 1026],
+          "data": { "kind": "fillFloat32", "sinStep": 0.023, "cosStep": 0.059, "scale": 2.0 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [3, 1026], "tolerance": 0.000004, "relTolerance": 0.00004 } },
+      "tunables": { "WORKGROUP_SIZE": 256 }
+    },
+    {
+      "name": "wide_subgroup_span_1028_rows_3_wg_256",
+      "inputs": {
+        "xT": {
+          "dtype": "float32",
+          "shape": [3, 1028],
+          "data": { "kind": "fillFloat32", "sinStep": 0.011, "cosStep": 0.047 }
+        },
+        "scaleT": {
+          "dtype": "float32",
+          "shape": [1028],
+          "data": { "kind": "fillFloat32", "sinStep": 0.019, "cosStep": 0.007, "scale": 0.5, "offset": 1.0 }
+        },
+        "gateT": {
+          "dtype": "float32",
+          "shape": [3, 1028],
+          "data": { "kind": "fillFloat32", "sinStep": 0.023, "cosStep": 0.059, "scale": 2.0 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [3, 1028], "tolerance": 0.000004, "relTolerance": 0.00004 } },
+      "tunables": { "WORKGROUP_SIZE": 256 }
+    },
+    {
+      "name": "wide_subgroup_span_2050_rows_3_wg_64",
+      "inputs": {
+        "xT": {
+          "dtype": "float32",
+          "shape": [3, 2050],
+          "data": { "kind": "fillFloat32", "sinStep": 0.011, "cosStep": 0.047 }
+        },
+        "scaleT": {
+          "dtype": "float32",
+          "shape": [2050],
+          "data": { "kind": "fillFloat32", "sinStep": 0.019, "cosStep": 0.007, "scale": 0.5, "offset": 1.0 }
+        },
+        "gateT": {
+          "dtype": "float32",
+          "shape": [3, 2050],
+          "data": { "kind": "fillFloat32", "sinStep": 0.023, "cosStep": 0.059, "scale": 2.0 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [3, 2050], "tolerance": 0.000004, "relTolerance": 0.00004 } },
+      "tunables": { "WORKGROUP_SIZE": 64 }
+    },
+    {
+      "name": "short_scalar_float32_c1_r257",
+      "inputs": {
+        "xT": {
+          "dtype": "float32",
+          "shape": [257, 1],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_x" } }
+        },
+        "scaleT": {
+          "dtype": "float32",
+          "shape": [1],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_scale" } }
+        },
+        "gateT": {
+          "dtype": "float32",
+          "shape": [257, 1],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_gate" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [257, 1], "tolerance": 0.000002, "relTolerance": 0.00002 } },
+      "provenance": {
+        "notes": "Short scalar row schedule: mixed magnitudes, signed and zero scale, SiLU tails, and rows on either side of a workgroup boundary, with float32 intermediates at the default dtype tolerances; spans eight/nine cover the staging boundary."
+      }
+    },
+    {
+      "name": "short_scalar_float32_c2_r257",
+      "inputs": {
+        "xT": {
+          "dtype": "float32",
+          "shape": [257, 2],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_x" } }
+        },
+        "scaleT": {
+          "dtype": "float32",
+          "shape": [2],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_scale" } }
+        },
+        "gateT": {
+          "dtype": "float32",
+          "shape": [257, 2],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_gate" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [257, 2], "tolerance": 0.000002, "relTolerance": 0.00002 } },
+      "provenance": {
+        "notes": "Short scalar row schedule: mixed magnitudes, signed and zero scale, SiLU tails, and rows on either side of a workgroup boundary, with float32 intermediates at the default dtype tolerances; spans eight/nine cover the staging boundary."
+      }
+    },
+    {
+      "name": "short_scalar_float32_c3_r257",
+      "inputs": {
+        "xT": {
+          "dtype": "float32",
+          "shape": [257, 3],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_x" } }
+        },
+        "scaleT": {
+          "dtype": "float32",
+          "shape": [3],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_scale" } }
+        },
+        "gateT": {
+          "dtype": "float32",
+          "shape": [257, 3],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_gate" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [257, 3], "tolerance": 0.000002, "relTolerance": 0.00002 } },
+      "provenance": {
+        "notes": "Short scalar row schedule: mixed magnitudes, signed and zero scale, SiLU tails, and rows on either side of a workgroup boundary, with float32 intermediates at the default dtype tolerances; spans eight/nine cover the staging boundary."
+      }
+    },
+    {
+      "name": "short_scalar_float32_c5_r257",
+      "inputs": {
+        "xT": {
+          "dtype": "float32",
+          "shape": [257, 5],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_x" } }
+        },
+        "scaleT": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_scale" } }
+        },
+        "gateT": {
+          "dtype": "float32",
+          "shape": [257, 5],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_gate" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [257, 5], "tolerance": 0.000002, "relTolerance": 0.00002 } },
+      "provenance": {
+        "notes": "Short scalar row schedule: mixed magnitudes, signed and zero scale, SiLU tails, and rows on either side of a workgroup boundary, with float32 intermediates at the default dtype tolerances; spans eight/nine cover the staging boundary."
+      }
+    },
+    {
+      "name": "short_scalar_float32_c6_r257",
+      "inputs": {
+        "xT": {
+          "dtype": "float32",
+          "shape": [257, 6],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_x" } }
+        },
+        "scaleT": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_scale" } }
+        },
+        "gateT": {
+          "dtype": "float32",
+          "shape": [257, 6],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_gate" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [257, 6], "tolerance": 0.000002, "relTolerance": 0.00002 } },
+      "provenance": {
+        "notes": "Short scalar row schedule: mixed magnitudes, signed and zero scale, SiLU tails, and rows on either side of a workgroup boundary, with float32 intermediates at the default dtype tolerances; spans eight/nine cover the staging boundary."
+      }
+    },
+    {
+      "name": "short_scalar_float32_c7_r257",
+      "inputs": {
+        "xT": {
+          "dtype": "float32",
+          "shape": [257, 7],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_x" } }
+        },
+        "scaleT": {
+          "dtype": "float32",
+          "shape": [7],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_scale" } }
+        },
+        "gateT": {
+          "dtype": "float32",
+          "shape": [257, 7],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_gate" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [257, 7], "tolerance": 0.000002, "relTolerance": 0.00002 } },
+      "provenance": {
+        "notes": "Short scalar row schedule: mixed magnitudes, signed and zero scale, SiLU tails, and rows on either side of a workgroup boundary, with float32 intermediates at the default dtype tolerances; spans eight/nine cover the staging boundary."
+      }
+    },
+    {
+      "name": "short_scalar_float32_c8_r257",
+      "inputs": {
+        "xT": {
+          "dtype": "float32",
+          "shape": [257, 8],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_x" } }
+        },
+        "scaleT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_scale" } }
+        },
+        "gateT": {
+          "dtype": "float32",
+          "shape": [257, 8],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_gate" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [257, 8], "tolerance": 0.000002, "relTolerance": 0.00002 } },
+      "provenance": {
+        "notes": "Short scalar row schedule: mixed magnitudes, signed and zero scale, SiLU tails, and rows on either side of a workgroup boundary, with float32 intermediates at the default dtype tolerances; spans eight/nine cover the staging boundary."
+      }
+    },
+    {
+      "name": "short_scalar_float32_c9_r257",
+      "inputs": {
+        "xT": {
+          "dtype": "float32",
+          "shape": [257, 9],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_x" } }
+        },
+        "scaleT": {
+          "dtype": "float32",
+          "shape": [9],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_scale" } }
+        },
+        "gateT": {
+          "dtype": "float32",
+          "shape": [257, 9],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_gate" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [257, 9], "tolerance": 0.000002, "relTolerance": 0.00002 } },
+      "provenance": {
+        "notes": "Short scalar row schedule: mixed magnitudes, signed and zero scale, SiLU tails, and rows on either side of a workgroup boundary, with float32 intermediates at the default dtype tolerances; spans eight/nine cover the staging boundary."
+      }
+    },
+    {
+      "name": "short_scalar_float16_c3_r257",
+      "inputs": {
+        "xT": {
+          "dtype": "float16",
+          "shape": [257, 3],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_x" } }
+        },
+        "scaleT": {
+          "dtype": "float16",
+          "shape": [3],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_scale" } }
+        },
+        "gateT": {
+          "dtype": "float16",
+          "shape": [257, 3],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_gate" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [257, 3], "tolerance": 0.0005, "relTolerance": 0.002 } },
+      "provenance": {
+        "notes": "Short scalar row schedule: mixed magnitudes, signed and zero scale, SiLU tails, and rows on either side of a workgroup boundary, with float32 intermediates at the default dtype tolerances; spans eight/nine cover the staging boundary."
+      }
+    },
+    {
+      "name": "short_scalar_float16_c5_r257",
+      "inputs": {
+        "xT": {
+          "dtype": "float16",
+          "shape": [257, 5],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_x" } }
+        },
+        "scaleT": {
+          "dtype": "float16",
+          "shape": [5],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_scale" } }
+        },
+        "gateT": {
+          "dtype": "float16",
+          "shape": [257, 5],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_gate" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [257, 5], "tolerance": 0.0005, "relTolerance": 0.002 } },
+      "provenance": {
+        "notes": "Short scalar row schedule: mixed magnitudes, signed and zero scale, SiLU tails, and rows on either side of a workgroup boundary, with float32 intermediates at the default dtype tolerances; spans eight/nine cover the staging boundary."
+      }
+    },
+    {
+      "name": "short_scalar_float16_c6_r257",
+      "inputs": {
+        "xT": {
+          "dtype": "float16",
+          "shape": [257, 6],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_x" } }
+        },
+        "scaleT": {
+          "dtype": "float16",
+          "shape": [6],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_scale" } }
+        },
+        "gateT": {
+          "dtype": "float16",
+          "shape": [257, 6],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_gate" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [257, 6], "tolerance": 0.0005, "relTolerance": 0.002 } },
+      "provenance": {
+        "notes": "Short scalar row schedule: mixed magnitudes, signed and zero scale, SiLU tails, and rows on either side of a workgroup boundary, with float32 intermediates at the default dtype tolerances; spans eight/nine cover the staging boundary."
+      }
+    },
+    {
+      "name": "short_scalar_float16_c7_r257",
+      "inputs": {
+        "xT": {
+          "dtype": "float16",
+          "shape": [257, 7],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_x" } }
+        },
+        "scaleT": {
+          "dtype": "float16",
+          "shape": [7],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_scale" } }
+        },
+        "gateT": {
+          "dtype": "float16",
+          "shape": [257, 7],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_gate" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [257, 7], "tolerance": 0.0005, "relTolerance": 0.002 } },
+      "provenance": {
+        "notes": "Short scalar row schedule: mixed magnitudes, signed and zero scale, SiLU tails, and rows on either side of a workgroup boundary, with float32 intermediates at the default dtype tolerances; spans eight/nine cover the staging boundary."
+      }
+    },
+    {
+      "name": "short_scalar_float32_c7_r65_wg64_words4",
+      "inputs": {
+        "xT": {
+          "dtype": "float32",
+          "shape": [65, 7],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_x" } }
+        },
+        "scaleT": {
+          "dtype": "float32",
+          "shape": [7],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_scale" } }
+        },
+        "gateT": {
+          "dtype": "float32",
+          "shape": [65, 7],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_gate" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [65, 7], "tolerance": 0.000002, "relTolerance": 0.00002 } },
+      "tunables": { "WORKGROUP_SIZE": 64, "LANE_WORDS": 4 },
+      "provenance": {
+        "notes": "Short scalar row schedule: mixed magnitudes, signed and zero scale, SiLU tails, and rows on either side of a workgroup boundary, with float32 intermediates at the default dtype tolerances; spans eight/nine cover the staging boundary."
+      }
+    },
+    {
+      "name": "short_scalar_float16_c7_r65_wg64_words4",
+      "inputs": {
+        "xT": {
+          "dtype": "float16",
+          "shape": [65, 7],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_x" } }
+        },
+        "scaleT": {
+          "dtype": "float16",
+          "shape": [7],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_scale" } }
+        },
+        "gateT": {
+          "dtype": "float16",
+          "shape": [65, 7],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_gate" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [65, 7], "tolerance": 0.0005, "relTolerance": 0.002 } },
+      "tunables": { "WORKGROUP_SIZE": 64, "LANE_WORDS": 4 },
+      "provenance": {
+        "notes": "Short scalar row schedule: mixed magnitudes, signed and zero scale, SiLU tails, and rows on either side of a workgroup boundary, with float32 intermediates at the default dtype tolerances; spans eight/nine cover the staging boundary."
+      }
+    },
+    {
+      "name": "short_scalar_float32_c6_r3_wg1_words1",
+      "inputs": {
+        "xT": {
+          "dtype": "float32",
+          "shape": [3, 6],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_x" } }
+        },
+        "scaleT": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_scale" } }
+        },
+        "gateT": {
+          "dtype": "float32",
+          "shape": [3, 6],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/short_scalar_gate" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [3, 6], "tolerance": 0.000002, "relTolerance": 0.00002 } },
+      "tunables": { "WORKGROUP_SIZE": 1, "LANE_WORDS": 1 },
+      "provenance": {
+        "notes": "Three rows of 6 float32 channels each, processed one lane per row; checks correctness at the smallest possible per-row lane allocation, below the width where multiple lanes would cooperate on one row."
+      }
+    }
+  ],
+  "fixtureArrays": {
+    "short_scalar_x": [-0.0009765625, 0.5, -1, 2, 8, -16, 64],
+    "short_scalar_scale": [-1, 0, 0.125, 1.5, -2],
+    "short_scalar_gate": [-80, -4, -1, 0, 1, 4, 80]
+  }
+}