sync 2e7068faf55e

Browse files

Files changed (7) hide show

README.md +93 -0
build/webgpu/bench.json +38 -0
build/webgpu/manifest.json +1948 -0
build/webgpu/matmul-nbits-fused-rms-norm.wgsl.jinja +117 -0
build/webgpu/metadata.json +19 -0
build/webgpu/mlp-gate-up.wgsl.jinja +290 -0
build/webgpu/test.json +1703 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# com.microsoft.MatMulNBitsMlp
+`com.microsoft`  ·  ONNX Runtime contrib operator  ·  contrib since_version 1
+## Description
+Fuses a gated MLP over two block-quantized projections that share one activation: `Y = silu(A_norm @ gate + gate_bias) * (A_norm @ up + up_bias)`, using the `MatMulNBits` weight packing with no zero-point input. `A_norm` is `A`, `SimplifiedLayerNormalization(A, norm_scale)`, or `SkipSimplifiedLayerNormalization(A, skip, norm_scale)`, whose residual sum may be returned as a second output. Only `silu` and the default `accuracy_level = 0` are implemented; bfloat16 is not implemented.
+See the [ONNX Runtime `MatMulNBitsMlp` contrib-operator spec](https://github.com/microsoft/onnxruntime/blob/main/docs/ContribOperators.md#com.microsoft.MatMulNBitsMlp) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `A` | `aT` | `T1` | — | — | Shared activation of rank 2 `(M, K)` or rank 3 `(batch, sequence, K)`; only the last axis is the reduction axis. | required |
+| `skip` | `skipT` | `T1` | — | — | Residual added to `A` before normalization, with `A`'s shape. Requires `norm_scale`. | optional |
+| `norm_scale` | `normScaleT` | `T1` | `1` | — | Simplified-layer-normalization (RMS) gain of shape `[K]`. Absent means the projections read `A` unnormalized. | optional |
+| `gate_B` | `gateBT` | `uint8` | `3` | — | Bit-packed uint8 gate weights of shape `(N, k_blocks, blob_size)`. | required |
+| `gate_scales` | `gateScalesT` | `T1` | `2` | — | Per-block gate scales of shape `(N, k_blocks)`, with the same dtype as `A`. Quantization is symmetric: this operator has no zero-point input, so codes are offset by the midpoint `2^(bits - 1)`. | required |
+| `gate_bias` | `gateBiasT` | `T1` | `1` | — | Optional gate bias of shape `[N]`, added before the activation. | optional |
+| `up_B` | `upBT` | `uint8` | `3` | — | Bit-packed up weights, same shape and packing as gate_B. | required |
+| `up_scales` | `upScalesT` | `T1` | `2` | — | Per-block up scales of shape `(N, k_blocks)`. | required |
+| `up_bias` | `upBiasT` | `T1` | `1` | — | Optional up bias of shape `[N]`, added before the product. | optional |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `Y` | `yT` | `T1` | same as `A` | derived; see description | Gated MLP output: A's leading axes with a trailing N. | required |
+| `input_skip_bias_sum` | `residualT` | `T1` | same as `A` | same as `A` | The residual sum A + skip, with A's shape. Requires the skip input. | optional |
+## Attributes
+Attributes and default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `accuracy_level` | `0` | Minimum internal accuracy level: 0 (unset), 1 (float32), 2 (float16), 3 (bfloat16), or 4 (int8). |
+| `bits` | `4` | Bit width used to quantize both weight matrices; this implementation supports 2, 4, and 8. |
+| `epsilon` | `0.00001` | Epsilon used by the optional fused RMS normalization. |
+| `K` | — | Input feature dimension shared by both quantized weight matrices. |
+| `N` | — | Output feature dimension shared by both quantized weight matrices. |
+| `activation` | — | Activation applied to the gate projection; this implementation supports `silu`. |
+| `block_size` | — | Size of each quantization block along K. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T1` | `float32`, `float16` |
+## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
+- [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
+- [`test.json`](build/webgpu/test.json) — correctness cases
+- [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
+- [`matmul-nbits-fused-rms-norm.wgsl.jinja`](build/webgpu/matmul-nbits-fused-rms-norm.wgsl.jinja)
+- [`mlp-gate-up.wgsl.jinja`](build/webgpu/mlp-gate-up.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.MatMulNBitsMlp", { version: 1 });
+const { yT } = await kernel({
+  aT: { data: aTData, shape: [2, 16] },
+  gateBT: { data: gateBTData, shape: [4, 2, 4] },
+  gateScalesT: { data: gateScalesTData, shape: [4, 2] },
+  upBT: { data: upBTData, shape: [4, 2, 4] },
+  upScalesT: { data: upScalesTData, shape: [4, 2] },
+}, {
+  attrs: {
+    K: 16,
+    N: 4,
+    block_size: 8,
+    activation: "silu",
+  },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,38 @@

+{
+  "op": "com.microsoft.MatMulNBitsMlp",
+  "tunableSpace": { "TILE_N": [4, 8, 16], "LANES": [8, 16, 32] },
+  "cases": [
+    {
+      "name": "mlp-q4-decode-k2048-n5632",
+      "preset": "smoke",
+      "vars": { "dtype": "float32" },
+      "attrs": { "K": 2048, "N": 5632, "bits": 4, "block_size": 32, "activation": "silu" },
+      "inputs": {
+        "aT": { "shape": [1, 2048], "dtype": "float32", "dist": "normal", "seed": 8101, "scale": 1 },
+        "normScaleT": { "shape": [2048], "dtype": "float32", "dist": "normal", "seed": 8102, "scale": 1 },
+        "gateBT": { "shape": [5632, 64, 16], "dtype": "uint8", "dist": "uniform", "seed": 8103, "scale": 255 },
+        "gateScalesT": { "shape": [5632, 64], "dtype": "float32", "dist": "normal", "seed": 8104, "scale": 0.05 },
+        "upBT": { "shape": [5632, 64, 16], "dtype": "uint8", "dist": "uniform", "seed": 8105, "scale": 255 },
+        "upScalesT": { "shape": [5632, 64], "dtype": "float32", "dist": "normal", "seed": 8106, "scale": 0.05 }
+      },
+      "outputs": { "yT": { "shape": [1, 5632], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 5632 * 64 * 16 * 4" }] }
+    },
+    {
+      "name": "mlp-q4-prefill-m64-k2048-n5632",
+      "preset": "model",
+      "vars": { "dtype": "float32" },
+      "attrs": { "K": 2048, "N": 5632, "bits": 4, "block_size": 32, "activation": "silu" },
+      "inputs": {
+        "aT": { "shape": [64, 2048], "dtype": "float32", "dist": "normal", "seed": 8107, "scale": 1 },
+        "normScaleT": { "shape": [2048], "dtype": "float32", "dist": "normal", "seed": 8108, "scale": 1 },
+        "gateBT": { "shape": [5632, 64, 16], "dtype": "uint8", "dist": "uniform", "seed": 8109, "scale": 255 },
+        "gateScalesT": { "shape": [5632, 64], "dtype": "float32", "dist": "normal", "seed": 8110, "scale": 0.05 },
+        "upBT": { "shape": [5632, 64, 16], "dtype": "uint8", "dist": "uniform", "seed": 8111, "scale": 255 },
+        "upScalesT": { "shape": [5632, 64], "dtype": "float32", "dist": "normal", "seed": 8112, "scale": 0.05 }
+      },
+      "outputs": { "yT": { "shape": [64, 5632], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * 2 * 64 * 2048 * 5632" }] }
+    }
+  ]
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,1948 @@

+{
+  "domain": "com.microsoft",
+  "name": "MatMulNBitsMlp",
+  "sinceVersion": 1,
+  "description": "Fuses a gated MLP over two block-quantized projections that share one activation: `Y = silu(A_norm @ gate + gate_bias) * (A_norm @ up + up_bias)`, using the `MatMulNBits` weight packing with no zero-point input. `A_norm` is `A`, `SimplifiedLayerNormalization(A, norm_scale)`, or `SkipSimplifiedLayerNormalization(A, skip, norm_scale)`, whose residual sum may be returned as a second output. Only `silu` and the default `accuracy_level = 0` are implemented; bfloat16 is not implemented.",
+  "inputs": [
+    {
+      "role": "A",
+      "dtype": "T1",
+      "description": "Shared activation of rank 2 `(M, K)` or rank 3 `(batch, sequence, K)`; only the last axis is the reduction axis."
+    },
+    {
+      "role": "skip",
+      "dtype": "T1",
+      "optional": true,
+      "description": "Residual added to `A` before normalization, with `A`'s shape. Requires `norm_scale`."
+    },
+    {
+      "role": "norm_scale",
+      "dtype": "T1",
+      "rank": 1,
+      "optional": true,
+      "description": "Simplified-layer-normalization (RMS) gain of shape `[K]`. Absent means the projections read `A` unnormalized."
+    },
+    {
+      "role": "gate_B",
+      "dtype": "uint8",
+      "rank": 3,
+      "description": "Bit-packed uint8 gate weights of shape `(N, k_blocks, blob_size)`."
+    },
+    {
+      "role": "gate_scales",
+      "dtype": "T1",
+      "rank": 2,
+      "description": "Per-block gate scales of shape `(N, k_blocks)`, with the same dtype as `A`. Quantization is symmetric: this operator has no zero-point input, so codes are offset by the midpoint `2^(bits - 1)`."
+    },
+    {
+      "role": "gate_bias",
+      "dtype": "T1",
+      "rank": 1,
+      "optional": true,
+      "description": "Optional gate bias of shape `[N]`, added before the activation."
+    },
+    {
+      "role": "up_B",
+      "dtype": "uint8",
+      "rank": 3,
+      "description": "Bit-packed up weights, same shape and packing as gate_B."
+    },
+    { "role": "up_scales", "dtype": "T1", "rank": 2, "description": "Per-block up scales of shape `(N, k_blocks)`." },
+    {
+      "role": "up_bias",
+      "dtype": "T1",
+      "rank": 1,
+      "optional": true,
+      "description": "Optional up bias of shape `[N]`, added before the product."
+    }
+  ],
+  "outputs": [
+    {
+      "role": "Y",
+      "dtype": "T1",
+      "rank": "ranks.aT",
+      "shape": "shapes.aT[:-1] + [attrs.N]",
+      "description": "Gated MLP output: A's leading axes with a trailing N."
+    },
+    {
+      "role": "input_skip_bias_sum",
+      "dtype": "T1",
+      "rank": "ranks.aT",
+      "optional": true,
+      "shape": "shapes.aT",
+      "description": "The residual sum A + skip, with A's shape. Requires the skip input."
+    }
+  ],
+  "attributes": { "accuracy_level": 0, "bits": 4, "epsilon": 0.00001 },
+  "attributeDescriptions": {
+    "K": "Input feature dimension shared by both quantized weight matrices.",
+    "N": "Output feature dimension shared by both quantized weight matrices.",
+    "accuracy_level": "Minimum internal accuracy level: 0 (unset), 1 (float32), 2 (float16), 3 (bfloat16), or 4 (int8).",
+    "activation": "Activation applied to the gate projection; this implementation supports `silu`.",
+    "bits": "Bit width used to quantize both weight matrices; this implementation supports 2, 4, and 8.",
+    "block_size": "Size of each quantization block along K.",
+    "epsilon": "Epsilon used by the optional fused RMS normalization."
+  },
+  "attributeConstraints": {
+    "K": { "required": true },
+    "N": { "required": true },
+    "accuracy_level": { "values": [0] },
+    "activation": { "required": true, "values": ["silu"] },
+    "bits": { "values": [2, 4, 8] },
+    "block_size": { "required": true }
+  },
+  "typeConstraints": { "T1": ["float32", "float16"] },
+  "args": {
+    "aT": { "kind": "tensor", "semantic": "A", "role": "input" },
+    "skipT": { "kind": "tensor", "semantic": "skip", "role": "input", "required": false },
+    "normScaleT": { "kind": "tensor", "semantic": "norm_scale", "role": "weights", "required": false },
+    "gateBT": { "kind": "tensor", "semantic": "gate_B", "role": "weights" },
+    "gateScalesT": { "kind": "tensor", "semantic": "gate_scales", "role": "weights" },
+    "gateBiasT": { "kind": "tensor", "semantic": "gate_bias", "role": "weights", "required": false },
+    "upBT": { "kind": "tensor", "semantic": "up_B", "role": "weights" },
+    "upScalesT": { "kind": "tensor", "semantic": "up_scales", "role": "weights" },
+    "upBiasT": { "kind": "tensor", "semantic": "up_bias", "role": "weights", "required": false },
+    "yT": { "kind": "tensor", "semantic": "Y", "role": "output" },
+    "residualT": { "kind": "tensor", "semantic": "input_skip_bias_sum", "role": "output", "required": false }
+  },
+  "tunables": { "TILE_N": 8, "LANES": 8, "NORM_WORKGROUP_SIZE": 128, "ROW_TILE": 8 },
+  "derive": {
+    "aRows": "numel(shapes.aT) / max(1, attrs.K)",
+    "rowTile": "1 if aRows <= 1 else min(aRows, tunables.ROW_TILE)",
+    "rowGroups": "ceilDiv(aRows, rowTile)",
+    "kBlocks": "dim(shapes.gateBT, 1)",
+    "blobSize": "dim(shapes.gateBT, 2)",
+    "codesPerByte": "8 / attrs.bits",
+    "codeMask": "3 if attrs.bits == 2 else (15 if attrs.bits == 4 else 255)",
+    "epsilonValue": "attrs.epsilon",
+    "bitsSupported": "attrs.bits == 2 or attrs.bits == 4 or attrs.bits == 8",
+    "weightShapeOk": "ranks.gateBT == 3 and ranks.upBT == 3 and dim(shapes.gateBT, 0) == attrs.N and dim(shapes.upBT, 0) == attrs.N and dim(shapes.upBT, 1) == kBlocks and dim(shapes.upBT, 2) == blobSize and kBlocks == ceilDiv(attrs.K, attrs.block_size) and blobSize * 8 == attrs.block_size * attrs.bits",
+    "scaleShapeOk": "ranks.gateScalesT == 2 and ranks.upScalesT == 2 and dim(shapes.gateScalesT, 0) == attrs.N and dim(shapes.gateScalesT, 1) == kBlocks and dim(shapes.upScalesT, 0) == attrs.N and dim(shapes.upScalesT, 1) == kBlocks",
+    "ioShapeOk": "(ranks.aT == 2 or ranks.aT == 3) and dim(shapes.aT, ranks.aT - 1) == attrs.K and ranks.yT == ranks.aT and dim(shapes.yT, ranks.yT - 1) == attrs.N and sameShape(prefix(shapes.yT, ranks.yT - 1), prefix(shapes.aT, ranks.aT - 1))",
+    "biasShapeOk": "(ranks.gateBiasT == 1 and dim(shapes.gateBiasT, 0) == attrs.N if present.gateBiasT else true) and (ranks.upBiasT == 1 and dim(shapes.upBiasT, 0) == attrs.N if present.upBiasT else true)",
+    "dtypeOk": "tensorDtypes.gateScalesT == tensorDtypes.aT and tensorDtypes.upScalesT == tensorDtypes.aT and tensorDtypes.yT == tensorDtypes.aT and f16Ok(tensorDtypes.aT)",
+    "lanesPow2": "tunables.LANES == pow2ceil(tunables.LANES)",
+    "mlpShapeOk": "bitsSupported and weightShapeOk and scaleShapeOk and ioShapeOk and biasShapeOk and dtypeOk and lanesPow2 and attrs.K > 0 and attrs.N > 0 and attrs.block_size > 0",
+    "normContractOk": "present.normScaleT and ranks.normScaleT == 1 and dim(shapes.normScaleT, 0) == attrs.K and tensorDtypes.normScaleT == tensorDtypes.aT and (sameShape(shapes.skipT, shapes.aT) and tensorDtypes.skipT == tensorDtypes.aT if present.skipT else true) and (sameShape(shapes.residualT, shapes.aT) and tensorDtypes.residualT == tensorDtypes.aT and present.skipT if present.residualT else true)",
+    "gateUpDispatchFits": "ceilDiv(attrs.N, tunables.TILE_N) <= device.limits.maxComputeWorkgroupsPerDimension and aRows <= device.limits.maxComputeWorkgroupsPerDimension and tunables.TILE_N * tunables.LANES <= device.limits.maxComputeInvocationsPerWorkgroup and tunables.TILE_N * tunables.LANES <= device.limits.maxComputeWorkgroupSizeX",
+    "normDispatchFits": "tunables.NORM_WORKGROUP_SIZE <= device.limits.maxComputeInvocationsPerWorkgroup and tunables.NORM_WORKGROUP_SIZE <= device.limits.maxComputeWorkgroupSizeX",
+    "biasPresence_nogb_noub": "not present.gateBiasT and not present.upBiasT",
+    "biasPresence_nogb_ub": "not present.gateBiasT and present.upBiasT",
+    "biasPresence_gb_noub": "present.gateBiasT and not present.upBiasT",
+    "biasPresence_gb_ub": "present.gateBiasT and present.upBiasT"
+  },
+  "constants": {
+    "aScalar": "\"f16\" if tensorDtypes.aT == \"float16\" else \"f32\"",
+    "scalar": "\"f16\" if tensorDtypes.aT == \"float16\" else \"f32\"",
+    "usesF16": "tensorDtypes.aT == \"float16\"",
+    "K": "attrs.K",
+    "N": "attrs.N",
+    "blockSize": "attrs.block_size",
+    "kBlocks": "kBlocks",
+    "blobSize": "blobSize",
+    "bits": "attrs.bits",
+    "codesPerByte": "codesPerByte",
+    "codeMask": "codeMask",
+    "defaultZero": "\"2.0\" if attrs.bits == 2 else (\"8.0\" if attrs.bits == 4 else \"128.0\")",
+    "tileN": "tunables.TILE_N",
+    "lanes": "tunables.LANES",
+    "rowTile": "rowTile",
+    "rows": "aRows",
+    "hidden": "attrs.K",
+    "workgroupSize": "tunables.NORM_WORKGROUP_SIZE",
+    "epsilon": "epsilonValue",
+    "hasGateBias": "present.gateBiasT",
+    "hasUpBias": "present.upBiasT",
+    "hasSkip": "present.skipT",
+    "writeResidual": "present.residualT",
+    "K_LEN": "attrs.K",
+    "N_LEN": "attrs.N"
+  },
+  "bindingSets": {
+    "normFull": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "skip",
+        "arg": "skipT",
+        "semantic": "skip",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "norm_scale",
+        "arg": "normScaleT",
+        "semantic": "norm_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$K_LEN"
+      },
+      { "name": "normed", "semantic": "normedA", "buffer": { "type": "storage" }, "elementType": "f32" },
+      {
+        "name": "residual",
+        "arg": "residualT",
+        "semantic": "input_skip_bias_sum",
+        "buffer": { "type": "storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "rows", "type": "u32", "value": "aRows" }] }
+      }
+    ],
+    "normSkip": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "skip",
+        "arg": "skipT",
+        "semantic": "skip",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "norm_scale",
+        "arg": "normScaleT",
+        "semantic": "norm_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$K_LEN"
+      },
+      { "name": "normed", "semantic": "normedA", "buffer": { "type": "storage" }, "elementType": "f32" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "rows", "type": "u32", "value": "aRows" }] }
+      }
+    ],
+    "normOnly": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "norm_scale",
+        "arg": "normScaleT",
+        "semantic": "norm_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$K_LEN"
+      },
+      { "name": "normed", "semantic": "normedA", "buffer": { "type": "storage" }, "elementType": "f32" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "rows", "type": "u32", "value": "aRows" }] }
+      }
+    ],
+    "gateUpAFull": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "gate_b",
+        "arg": "gateBT",
+        "semantic": "gate_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "gate_scales",
+        "arg": "gateScalesT",
+        "semantic": "gate_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "gate_bias",
+        "arg": "gateBiasT",
+        "semantic": "gate_bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$N_LEN"
+      },
+      {
+        "name": "up_b",
+        "arg": "upBT",
+        "semantic": "up_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "up_scales",
+        "arg": "upScalesT",
+        "semantic": "up_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "up_bias",
+        "arg": "upBiasT",
+        "semantic": "up_bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$N_LEN"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$aScalar" }
+    ],
+    "gateUpNormedFull": [
+      { "name": "normed", "semantic": "normedA", "buffer": { "type": "read-only-storage" }, "elementType": "f32" },
+      {
+        "name": "gate_b",
+        "arg": "gateBT",
+        "semantic": "gate_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "gate_scales",
+        "arg": "gateScalesT",
+        "semantic": "gate_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "gate_bias",
+        "arg": "gateBiasT",
+        "semantic": "gate_bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$N_LEN"
+      },
+      {
+        "name": "up_b",
+        "arg": "upBT",
+        "semantic": "up_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "up_scales",
+        "arg": "upScalesT",
+        "semantic": "up_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "up_bias",
+        "arg": "upBiasT",
+        "semantic": "up_bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$N_LEN"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$aScalar" }
+    ],
+    "fusedFull": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "skip",
+        "arg": "skipT",
+        "semantic": "skip",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "norm_scale",
+        "arg": "normScaleT",
+        "semantic": "norm_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$K_LEN"
+      },
+      {
+        "name": "gate_b",
+        "arg": "gateBT",
+        "semantic": "gate_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "gate_scales",
+        "arg": "gateScalesT",
+        "semantic": "gate_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "gate_bias",
+        "arg": "gateBiasT",
+        "semantic": "gate_bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$N_LEN"
+      },
+      {
+        "name": "up_b",
+        "arg": "upBT",
+        "semantic": "up_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "up_scales",
+        "arg": "upScalesT",
+        "semantic": "up_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "up_bias",
+        "arg": "upBiasT",
+        "semantic": "up_bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$N_LEN"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$aScalar" },
+      {
+        "name": "residual",
+        "arg": "residualT",
+        "semantic": "input_skip_bias_sum",
+        "buffer": { "type": "storage" },
+        "elementType": "$aScalar"
+      }
+    ],
+    "gateUpA_nogb_noub": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "gate_b",
+        "arg": "gateBT",
+        "semantic": "gate_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "gate_scales",
+        "arg": "gateScalesT",
+        "semantic": "gate_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "up_b",
+        "arg": "upBT",
+        "semantic": "up_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "up_scales",
+        "arg": "upScalesT",
+        "semantic": "up_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$aScalar" }
+    ],
+    "gateUpNormed_nogb_noub": [
+      { "name": "normed", "semantic": "normedA", "buffer": { "type": "read-only-storage" }, "elementType": "f32" },
+      {
+        "name": "gate_b",
+        "arg": "gateBT",
+        "semantic": "gate_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "gate_scales",
+        "arg": "gateScalesT",
+        "semantic": "gate_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "up_b",
+        "arg": "upBT",
+        "semantic": "up_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "up_scales",
+        "arg": "upScalesT",
+        "semantic": "up_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$aScalar" }
+    ],
+    "fused_norm_nogb_noub": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "norm_scale",
+        "arg": "normScaleT",
+        "semantic": "norm_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$K_LEN"
+      },
+      {
+        "name": "gate_b",
+        "arg": "gateBT",
+        "semantic": "gate_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "gate_scales",
+        "arg": "gateScalesT",
+        "semantic": "gate_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "up_b",
+        "arg": "upBT",
+        "semantic": "up_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "up_scales",
+        "arg": "upScalesT",
+        "semantic": "up_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$aScalar" }
+    ],
+    "fused_skip_nogb_noub": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "skip",
+        "arg": "skipT",
+        "semantic": "skip",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "norm_scale",
+        "arg": "normScaleT",
+        "semantic": "norm_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$K_LEN"
+      },
+      {
+        "name": "gate_b",
+        "arg": "gateBT",
+        "semantic": "gate_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "gate_scales",
+        "arg": "gateScalesT",
+        "semantic": "gate_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "up_b",
+        "arg": "upBT",
+        "semantic": "up_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "up_scales",
+        "arg": "upScalesT",
+        "semantic": "up_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$aScalar" }
+    ],
+    "fused_skipsum_nogb_noub": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "skip",
+        "arg": "skipT",
+        "semantic": "skip",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "norm_scale",
+        "arg": "normScaleT",
+        "semantic": "norm_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$K_LEN"
+      },
+      {
+        "name": "gate_b",
+        "arg": "gateBT",
+        "semantic": "gate_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "gate_scales",
+        "arg": "gateScalesT",
+        "semantic": "gate_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "up_b",
+        "arg": "upBT",
+        "semantic": "up_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "up_scales",
+        "arg": "upScalesT",
+        "semantic": "up_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$aScalar" },
+      {
+        "name": "residual",
+        "arg": "residualT",
+        "semantic": "input_skip_bias_sum",
+        "buffer": { "type": "storage" },
+        "elementType": "$aScalar"
+      }
+    ],
+    "gateUpA_nogb_ub": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "gate_b",
+        "arg": "gateBT",
+        "semantic": "gate_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "gate_scales",
+        "arg": "gateScalesT",
+        "semantic": "gate_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "up_b",
+        "arg": "upBT",
+        "semantic": "up_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "up_scales",
+        "arg": "upScalesT",
+        "semantic": "up_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "up_bias",
+        "arg": "upBiasT",
+        "semantic": "up_bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$N_LEN"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$aScalar" }
+    ],
+    "gateUpNormed_nogb_ub": [
+      { "name": "normed", "semantic": "normedA", "buffer": { "type": "read-only-storage" }, "elementType": "f32" },
+      {
+        "name": "gate_b",
+        "arg": "gateBT",
+        "semantic": "gate_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "gate_scales",
+        "arg": "gateScalesT",
+        "semantic": "gate_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "up_b",
+        "arg": "upBT",
+        "semantic": "up_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "up_scales",
+        "arg": "upScalesT",
+        "semantic": "up_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "up_bias",
+        "arg": "upBiasT",
+        "semantic": "up_bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$N_LEN"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$aScalar" }
+    ],
+    "fused_norm_nogb_ub": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "norm_scale",
+        "arg": "normScaleT",
+        "semantic": "norm_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$K_LEN"
+      },
+      {
+        "name": "gate_b",
+        "arg": "gateBT",
+        "semantic": "gate_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "gate_scales",
+        "arg": "gateScalesT",
+        "semantic": "gate_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "up_b",
+        "arg": "upBT",
+        "semantic": "up_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "up_scales",
+        "arg": "upScalesT",
+        "semantic": "up_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "up_bias",
+        "arg": "upBiasT",
+        "semantic": "up_bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$N_LEN"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$aScalar" }
+    ],
+    "fused_skip_nogb_ub": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "skip",
+        "arg": "skipT",
+        "semantic": "skip",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "norm_scale",
+        "arg": "normScaleT",
+        "semantic": "norm_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$K_LEN"
+      },
+      {
+        "name": "gate_b",
+        "arg": "gateBT",
+        "semantic": "gate_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "gate_scales",
+        "arg": "gateScalesT",
+        "semantic": "gate_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "up_b",
+        "arg": "upBT",
+        "semantic": "up_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "up_scales",
+        "arg": "upScalesT",
+        "semantic": "up_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "up_bias",
+        "arg": "upBiasT",
+        "semantic": "up_bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$N_LEN"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$aScalar" }
+    ],
+    "fused_skipsum_nogb_ub": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "skip",
+        "arg": "skipT",
+        "semantic": "skip",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "norm_scale",
+        "arg": "normScaleT",
+        "semantic": "norm_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$K_LEN"
+      },
+      {
+        "name": "gate_b",
+        "arg": "gateBT",
+        "semantic": "gate_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "gate_scales",
+        "arg": "gateScalesT",
+        "semantic": "gate_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "up_b",
+        "arg": "upBT",
+        "semantic": "up_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "up_scales",
+        "arg": "upScalesT",
+        "semantic": "up_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "up_bias",
+        "arg": "upBiasT",
+        "semantic": "up_bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$N_LEN"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$aScalar" },
+      {
+        "name": "residual",
+        "arg": "residualT",
+        "semantic": "input_skip_bias_sum",
+        "buffer": { "type": "storage" },
+        "elementType": "$aScalar"
+      }
+    ],
+    "gateUpA_gb_noub": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "gate_b",
+        "arg": "gateBT",
+        "semantic": "gate_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "gate_scales",
+        "arg": "gateScalesT",
+        "semantic": "gate_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "gate_bias",
+        "arg": "gateBiasT",
+        "semantic": "gate_bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$N_LEN"
+      },
+      {
+        "name": "up_b",
+        "arg": "upBT",
+        "semantic": "up_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "up_scales",
+        "arg": "upScalesT",
+        "semantic": "up_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$aScalar" }
+    ],
+    "gateUpNormed_gb_noub": [
+      { "name": "normed", "semantic": "normedA", "buffer": { "type": "read-only-storage" }, "elementType": "f32" },
+      {
+        "name": "gate_b",
+        "arg": "gateBT",
+        "semantic": "gate_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "gate_scales",
+        "arg": "gateScalesT",
+        "semantic": "gate_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "gate_bias",
+        "arg": "gateBiasT",
+        "semantic": "gate_bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$N_LEN"
+      },
+      {
+        "name": "up_b",
+        "arg": "upBT",
+        "semantic": "up_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "up_scales",
+        "arg": "upScalesT",
+        "semantic": "up_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$aScalar" }
+    ],
+    "fused_norm_gb_noub": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "norm_scale",
+        "arg": "normScaleT",
+        "semantic": "norm_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$K_LEN"
+      },
+      {
+        "name": "gate_b",
+        "arg": "gateBT",
+        "semantic": "gate_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "gate_scales",
+        "arg": "gateScalesT",
+        "semantic": "gate_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "gate_bias",
+        "arg": "gateBiasT",
+        "semantic": "gate_bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$N_LEN"
+      },
+      {
+        "name": "up_b",
+        "arg": "upBT",
+        "semantic": "up_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "up_scales",
+        "arg": "upScalesT",
+        "semantic": "up_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$aScalar" }
+    ],
+    "fused_skip_gb_noub": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "skip",
+        "arg": "skipT",
+        "semantic": "skip",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "norm_scale",
+        "arg": "normScaleT",
+        "semantic": "norm_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$K_LEN"
+      },
+      {
+        "name": "gate_b",
+        "arg": "gateBT",
+        "semantic": "gate_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "gate_scales",
+        "arg": "gateScalesT",
+        "semantic": "gate_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "gate_bias",
+        "arg": "gateBiasT",
+        "semantic": "gate_bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$N_LEN"
+      },
+      {
+        "name": "up_b",
+        "arg": "upBT",
+        "semantic": "up_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "up_scales",
+        "arg": "upScalesT",
+        "semantic": "up_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$aScalar" }
+    ],
+    "fused_skipsum_gb_noub": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "skip",
+        "arg": "skipT",
+        "semantic": "skip",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "norm_scale",
+        "arg": "normScaleT",
+        "semantic": "norm_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$K_LEN"
+      },
+      {
+        "name": "gate_b",
+        "arg": "gateBT",
+        "semantic": "gate_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "gate_scales",
+        "arg": "gateScalesT",
+        "semantic": "gate_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "gate_bias",
+        "arg": "gateBiasT",
+        "semantic": "gate_bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$N_LEN"
+      },
+      {
+        "name": "up_b",
+        "arg": "upBT",
+        "semantic": "up_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "up_scales",
+        "arg": "upScalesT",
+        "semantic": "up_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$aScalar" },
+      {
+        "name": "residual",
+        "arg": "residualT",
+        "semantic": "input_skip_bias_sum",
+        "buffer": { "type": "storage" },
+        "elementType": "$aScalar"
+      }
+    ],
+    "gateUpA_gb_ub": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "gate_b",
+        "arg": "gateBT",
+        "semantic": "gate_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "gate_scales",
+        "arg": "gateScalesT",
+        "semantic": "gate_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "gate_bias",
+        "arg": "gateBiasT",
+        "semantic": "gate_bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$N_LEN"
+      },
+      {
+        "name": "up_b",
+        "arg": "upBT",
+        "semantic": "up_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "up_scales",
+        "arg": "upScalesT",
+        "semantic": "up_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "up_bias",
+        "arg": "upBiasT",
+        "semantic": "up_bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$N_LEN"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$aScalar" }
+    ],
+    "gateUpNormed_gb_ub": [
+      { "name": "normed", "semantic": "normedA", "buffer": { "type": "read-only-storage" }, "elementType": "f32" },
+      {
+        "name": "gate_b",
+        "arg": "gateBT",
+        "semantic": "gate_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "gate_scales",
+        "arg": "gateScalesT",
+        "semantic": "gate_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "gate_bias",
+        "arg": "gateBiasT",
+        "semantic": "gate_bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$N_LEN"
+      },
+      {
+        "name": "up_b",
+        "arg": "upBT",
+        "semantic": "up_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "up_scales",
+        "arg": "upScalesT",
+        "semantic": "up_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "up_bias",
+        "arg": "upBiasT",
+        "semantic": "up_bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$N_LEN"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$aScalar" }
+    ],
+    "fused_norm_gb_ub": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "norm_scale",
+        "arg": "normScaleT",
+        "semantic": "norm_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$K_LEN"
+      },
+      {
+        "name": "gate_b",
+        "arg": "gateBT",
+        "semantic": "gate_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "gate_scales",
+        "arg": "gateScalesT",
+        "semantic": "gate_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "gate_bias",
+        "arg": "gateBiasT",
+        "semantic": "gate_bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$N_LEN"
+      },
+      {
+        "name": "up_b",
+        "arg": "upBT",
+        "semantic": "up_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "up_scales",
+        "arg": "upScalesT",
+        "semantic": "up_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "up_bias",
+        "arg": "upBiasT",
+        "semantic": "up_bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$N_LEN"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$aScalar" }
+    ],
+    "fused_skip_gb_ub": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "skip",
+        "arg": "skipT",
+        "semantic": "skip",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "norm_scale",
+        "arg": "normScaleT",
+        "semantic": "norm_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$K_LEN"
+      },
+      {
+        "name": "gate_b",
+        "arg": "gateBT",
+        "semantic": "gate_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "gate_scales",
+        "arg": "gateScalesT",
+        "semantic": "gate_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "gate_bias",
+        "arg": "gateBiasT",
+        "semantic": "gate_bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$N_LEN"
+      },
+      {
+        "name": "up_b",
+        "arg": "upBT",
+        "semantic": "up_B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "up_scales",
+        "arg": "upScalesT",
+        "semantic": "up_scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "up_bias",
+        "arg": "upBiasT",
+        "semantic": "up_bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$N_LEN"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$aScalar" }
+    ]
+  },
+  "variants": [
+    {
+      "id": "plain_nogb_noub",
+      "priority": 10,
+      "when": ["mlpShapeOk", "not present.normScaleT", "not present.skipT", "not present.residualT", "biasPresence_nogb_noub", "gateUpDispatchFits"],
+      "constants": { "inlineNorm": "0", "fromNormed": "0" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "MatMulNBitsMlp.GateUp",
+          "shader": "mlp-gate-up.wgsl.jinja",
+          "bindings": "gateUpA_nogb_noub",
+          "dispatch": { "x": "ceilDiv(attrs.N, tunables.TILE_N)", "y": "rowGroups" }
+        }
+      ]
+    },
+    {
+      "id": "plain_nogb_ub",
+      "priority": 10,
+      "when": ["mlpShapeOk", "not present.normScaleT", "not present.skipT", "not present.residualT", "biasPresence_nogb_ub", "gateUpDispatchFits"],
+      "constants": { "inlineNorm": "0", "fromNormed": "0" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "MatMulNBitsMlp.GateUp",
+          "shader": "mlp-gate-up.wgsl.jinja",
+          "bindings": "gateUpA_nogb_ub",
+          "dispatch": { "x": "ceilDiv(attrs.N, tunables.TILE_N)", "y": "rowGroups" }
+        }
+      ]
+    },
+    {
+      "id": "plain_gb_noub",
+      "priority": 10,
+      "when": ["mlpShapeOk", "not present.normScaleT", "not present.skipT", "not present.residualT", "biasPresence_gb_noub", "gateUpDispatchFits"],
+      "constants": { "inlineNorm": "0", "fromNormed": "0" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "MatMulNBitsMlp.GateUp",
+          "shader": "mlp-gate-up.wgsl.jinja",
+          "bindings": "gateUpA_gb_noub",
+          "dispatch": { "x": "ceilDiv(attrs.N, tunables.TILE_N)", "y": "rowGroups" }
+        }
+      ]
+    },
+    {
+      "id": "plain_gb_ub",
+      "priority": 10,
+      "when": ["mlpShapeOk", "not present.normScaleT", "not present.skipT", "not present.residualT", "biasPresence_gb_ub", "gateUpDispatchFits"],
+      "constants": { "inlineNorm": "0", "fromNormed": "0" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "MatMulNBitsMlp.GateUp",
+          "shader": "mlp-gate-up.wgsl.jinja",
+          "bindings": "gateUpA_gb_ub",
+          "dispatch": { "x": "ceilDiv(attrs.N, tunables.TILE_N)", "y": "rowGroups" }
+        }
+      ]
+    },
+    {
+      "id": "staged_norm_nogb_noub",
+      "priority": 10,
+      "when": ["mlpShapeOk", "normContractOk", "biasPresence_nogb_noub", "gateUpDispatchFits", "normDispatchFits", "not present.skipT", "not present.residualT"],
+      "constants": { "inlineNorm": "0", "fromNormed": "1" },
+      "intermediates": [{ "id": "normedA", "dtype": "float32", "shape": "[numel(shapes.aT)]" }],
+      "passes": [
+        {
+          "id": "norm",
+          "name": "MatMulNBitsMlp.RmsNorm",
+          "shader": "matmul-nbits-fused-rms-norm.wgsl.jinja",
+          "bindings": "normOnly",
+          "dispatch": { "workgroups": "aRows" }
+        },
+        {
+          "id": "main",
+          "name": "MatMulNBitsMlp.GateUp",
+          "shader": "mlp-gate-up.wgsl.jinja",
+          "bindings": "gateUpNormed_nogb_noub",
+          "dispatch": { "x": "ceilDiv(attrs.N, tunables.TILE_N)", "y": "rowGroups" }
+        }
+      ]
+    },
+    {
+      "id": "staged_skip_nogb_noub",
+      "priority": 10,
+      "when": ["mlpShapeOk", "normContractOk", "biasPresence_nogb_noub", "gateUpDispatchFits", "normDispatchFits", "present.skipT", "not present.residualT"],
+      "constants": { "inlineNorm": "0", "fromNormed": "1" },
+      "intermediates": [{ "id": "normedA", "dtype": "float32", "shape": "[numel(shapes.aT)]" }],
+      "passes": [
+        {
+          "id": "norm",
+          "name": "MatMulNBitsMlp.RmsNorm",
+          "shader": "matmul-nbits-fused-rms-norm.wgsl.jinja",
+          "bindings": "normSkip",
+          "dispatch": { "workgroups": "aRows" }
+        },
+        {
+          "id": "main",
+          "name": "MatMulNBitsMlp.GateUp",
+          "shader": "mlp-gate-up.wgsl.jinja",
+          "bindings": "gateUpNormed_nogb_noub",
+          "dispatch": { "x": "ceilDiv(attrs.N, tunables.TILE_N)", "y": "rowGroups" }
+        }
+      ]
+    },
+    {
+      "id": "staged_skipsum_nogb_noub",
+      "priority": 10,
+      "when": ["mlpShapeOk", "normContractOk", "biasPresence_nogb_noub", "gateUpDispatchFits", "normDispatchFits", "present.skipT", "present.residualT"],
+      "constants": { "inlineNorm": "0", "fromNormed": "1" },
+      "intermediates": [{ "id": "normedA", "dtype": "float32", "shape": "[numel(shapes.aT)]" }],
+      "passes": [
+        {
+          "id": "norm",
+          "name": "MatMulNBitsMlp.RmsNorm",
+          "shader": "matmul-nbits-fused-rms-norm.wgsl.jinja",
+          "bindings": "normFull",
+          "dispatch": { "workgroups": "aRows" }
+        },
+        {
+          "id": "main",
+          "name": "MatMulNBitsMlp.GateUp",
+          "shader": "mlp-gate-up.wgsl.jinja",
+          "bindings": "gateUpNormed_nogb_noub",
+          "dispatch": { "x": "ceilDiv(attrs.N, tunables.TILE_N)", "y": "rowGroups" }
+        }
+      ]
+    },
+    {
+      "id": "staged_norm_nogb_ub",
+      "priority": 10,
+      "when": ["mlpShapeOk", "normContractOk", "biasPresence_nogb_ub", "gateUpDispatchFits", "normDispatchFits", "not present.skipT", "not present.residualT"],
+      "constants": { "inlineNorm": "0", "fromNormed": "1" },
+      "intermediates": [{ "id": "normedA", "dtype": "float32", "shape": "[numel(shapes.aT)]" }],
+      "passes": [
+        {
+          "id": "norm",
+          "name": "MatMulNBitsMlp.RmsNorm",
+          "shader": "matmul-nbits-fused-rms-norm.wgsl.jinja",
+          "bindings": "normOnly",
+          "dispatch": { "workgroups": "aRows" }
+        },
+        {
+          "id": "main",
+          "name": "MatMulNBitsMlp.GateUp",
+          "shader": "mlp-gate-up.wgsl.jinja",
+          "bindings": "gateUpNormed_nogb_ub",
+          "dispatch": { "x": "ceilDiv(attrs.N, tunables.TILE_N)", "y": "rowGroups" }
+        }
+      ]
+    },
+    {
+      "id": "staged_skip_nogb_ub",
+      "priority": 10,
+      "when": ["mlpShapeOk", "normContractOk", "biasPresence_nogb_ub", "gateUpDispatchFits", "normDispatchFits", "present.skipT", "not present.residualT"],
+      "constants": { "inlineNorm": "0", "fromNormed": "1" },
+      "intermediates": [{ "id": "normedA", "dtype": "float32", "shape": "[numel(shapes.aT)]" }],
+      "passes": [
+        {
+          "id": "norm",
+          "name": "MatMulNBitsMlp.RmsNorm",
+          "shader": "matmul-nbits-fused-rms-norm.wgsl.jinja",
+          "bindings": "normSkip",
+          "dispatch": { "workgroups": "aRows" }
+        },
+        {
+          "id": "main",
+          "name": "MatMulNBitsMlp.GateUp",
+          "shader": "mlp-gate-up.wgsl.jinja",
+          "bindings": "gateUpNormed_nogb_ub",
+          "dispatch": { "x": "ceilDiv(attrs.N, tunables.TILE_N)", "y": "rowGroups" }
+        }
+      ]
+    },
+    {
+      "id": "staged_skipsum_nogb_ub",
+      "priority": 10,
+      "when": ["mlpShapeOk", "normContractOk", "biasPresence_nogb_ub", "gateUpDispatchFits", "normDispatchFits", "present.skipT", "present.residualT"],
+      "constants": { "inlineNorm": "0", "fromNormed": "1" },
+      "intermediates": [{ "id": "normedA", "dtype": "float32", "shape": "[numel(shapes.aT)]" }],
+      "passes": [
+        {
+          "id": "norm",
+          "name": "MatMulNBitsMlp.RmsNorm",
+          "shader": "matmul-nbits-fused-rms-norm.wgsl.jinja",
+          "bindings": "normFull",
+          "dispatch": { "workgroups": "aRows" }
+        },
+        {
+          "id": "main",
+          "name": "MatMulNBitsMlp.GateUp",
+          "shader": "mlp-gate-up.wgsl.jinja",
+          "bindings": "gateUpNormed_nogb_ub",
+          "dispatch": { "x": "ceilDiv(attrs.N, tunables.TILE_N)", "y": "rowGroups" }
+        }
+      ]
+    },
+    {
+      "id": "staged_norm_gb_noub",
+      "priority": 10,
+      "when": ["mlpShapeOk", "normContractOk", "biasPresence_gb_noub", "gateUpDispatchFits", "normDispatchFits", "not present.skipT", "not present.residualT"],
+      "constants": { "inlineNorm": "0", "fromNormed": "1" },
+      "intermediates": [{ "id": "normedA", "dtype": "float32", "shape": "[numel(shapes.aT)]" }],
+      "passes": [
+        {
+          "id": "norm",
+          "name": "MatMulNBitsMlp.RmsNorm",
+          "shader": "matmul-nbits-fused-rms-norm.wgsl.jinja",
+          "bindings": "normOnly",
+          "dispatch": { "workgroups": "aRows" }
+        },
+        {
+          "id": "main",
+          "name": "MatMulNBitsMlp.GateUp",
+          "shader": "mlp-gate-up.wgsl.jinja",
+          "bindings": "gateUpNormed_gb_noub",
+          "dispatch": { "x": "ceilDiv(attrs.N, tunables.TILE_N)", "y": "rowGroups" }
+        }
+      ]
+    },
+    {
+      "id": "staged_skip_gb_noub",
+      "priority": 10,
+      "when": ["mlpShapeOk", "normContractOk", "biasPresence_gb_noub", "gateUpDispatchFits", "normDispatchFits", "present.skipT", "not present.residualT"],
+      "constants": { "inlineNorm": "0", "fromNormed": "1" },
+      "intermediates": [{ "id": "normedA", "dtype": "float32", "shape": "[numel(shapes.aT)]" }],
+      "passes": [
+        {
+          "id": "norm",
+          "name": "MatMulNBitsMlp.RmsNorm",
+          "shader": "matmul-nbits-fused-rms-norm.wgsl.jinja",
+          "bindings": "normSkip",
+          "dispatch": { "workgroups": "aRows" }
+        },
+        {
+          "id": "main",
+          "name": "MatMulNBitsMlp.GateUp",
+          "shader": "mlp-gate-up.wgsl.jinja",
+          "bindings": "gateUpNormed_gb_noub",
+          "dispatch": { "x": "ceilDiv(attrs.N, tunables.TILE_N)", "y": "rowGroups" }
+        }
+      ]
+    },
+    {
+      "id": "staged_skipsum_gb_noub",
+      "priority": 10,
+      "when": ["mlpShapeOk", "normContractOk", "biasPresence_gb_noub", "gateUpDispatchFits", "normDispatchFits", "present.skipT", "present.residualT"],
+      "constants": { "inlineNorm": "0", "fromNormed": "1" },
+      "intermediates": [{ "id": "normedA", "dtype": "float32", "shape": "[numel(shapes.aT)]" }],
+      "passes": [
+        {
+          "id": "norm",
+          "name": "MatMulNBitsMlp.RmsNorm",
+          "shader": "matmul-nbits-fused-rms-norm.wgsl.jinja",
+          "bindings": "normFull",
+          "dispatch": { "workgroups": "aRows" }
+        },
+        {
+          "id": "main",
+          "name": "MatMulNBitsMlp.GateUp",
+          "shader": "mlp-gate-up.wgsl.jinja",
+          "bindings": "gateUpNormed_gb_noub",
+          "dispatch": { "x": "ceilDiv(attrs.N, tunables.TILE_N)", "y": "rowGroups" }
+        }
+      ]
+    },
+    {
+      "id": "staged_norm_gb_ub",
+      "priority": 10,
+      "when": ["mlpShapeOk", "normContractOk", "biasPresence_gb_ub", "gateUpDispatchFits", "normDispatchFits", "not present.skipT", "not present.residualT"],
+      "constants": { "inlineNorm": "0", "fromNormed": "1" },
+      "intermediates": [{ "id": "normedA", "dtype": "float32", "shape": "[numel(shapes.aT)]" }],
+      "passes": [
+        {
+          "id": "norm",
+          "name": "MatMulNBitsMlp.RmsNorm",
+          "shader": "matmul-nbits-fused-rms-norm.wgsl.jinja",
+          "bindings": "normOnly",
+          "dispatch": { "workgroups": "aRows" }
+        },
+        {
+          "id": "main",
+          "name": "MatMulNBitsMlp.GateUp",
+          "shader": "mlp-gate-up.wgsl.jinja",
+          "bindings": "gateUpNormed_gb_ub",
+          "dispatch": { "x": "ceilDiv(attrs.N, tunables.TILE_N)", "y": "rowGroups" }
+        }
+      ]
+    },
+    {
+      "id": "staged_skip_gb_ub",
+      "priority": 10,
+      "when": ["mlpShapeOk", "normContractOk", "biasPresence_gb_ub", "gateUpDispatchFits", "normDispatchFits", "present.skipT", "not present.residualT"],
+      "constants": { "inlineNorm": "0", "fromNormed": "1" },
+      "intermediates": [{ "id": "normedA", "dtype": "float32", "shape": "[numel(shapes.aT)]" }],
+      "passes": [
+        {
+          "id": "norm",
+          "name": "MatMulNBitsMlp.RmsNorm",
+          "shader": "matmul-nbits-fused-rms-norm.wgsl.jinja",
+          "bindings": "normSkip",
+          "dispatch": { "workgroups": "aRows" }
+        },
+        {
+          "id": "main",
+          "name": "MatMulNBitsMlp.GateUp",
+          "shader": "mlp-gate-up.wgsl.jinja",
+          "bindings": "gateUpNormed_gb_ub",
+          "dispatch": { "x": "ceilDiv(attrs.N, tunables.TILE_N)", "y": "rowGroups" }
+        }
+      ]
+    },
+    {
+      "id": "staged_skipsum_gb_ub",
+      "priority": 10,
+      "when": ["mlpShapeOk", "normContractOk", "biasPresence_gb_ub", "gateUpDispatchFits", "normDispatchFits", "present.skipT", "present.residualT"],
+      "constants": { "inlineNorm": "0", "fromNormed": "1" },
+      "intermediates": [{ "id": "normedA", "dtype": "float32", "shape": "[numel(shapes.aT)]" }],
+      "passes": [
+        {
+          "id": "norm",
+          "name": "MatMulNBitsMlp.RmsNorm",
+          "shader": "matmul-nbits-fused-rms-norm.wgsl.jinja",
+          "bindings": "normFull",
+          "dispatch": { "workgroups": "aRows" }
+        },
+        {
+          "id": "main",
+          "name": "MatMulNBitsMlp.GateUp",
+          "shader": "mlp-gate-up.wgsl.jinja",
+          "bindings": "gateUpNormed_gb_ub",
+          "dispatch": { "x": "ceilDiv(attrs.N, tunables.TILE_N)", "y": "rowGroups" }
+        }
+      ]
+    },
+    {
+      "id": "fused_norm_nogb_noub",
+      "priority": 30,
+      "requires": { "limits": { "maxStorageBuffersPerShaderStage": 7 } },
+      "when": ["mlpShapeOk", "normContractOk", "biasPresence_nogb_noub", "gateUpDispatchFits", "aRows == 1", "not present.skipT", "not present.residualT"],
+      "constants": { "inlineNorm": "1", "fromNormed": "0", "rowTile": "1" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "MatMulNBitsMlp.FusedDecode",
+          "shader": "mlp-gate-up.wgsl.jinja",
+          "bindings": "fused_norm_nogb_noub",
+          "dispatch": { "x": "ceilDiv(attrs.N, tunables.TILE_N)", "y": "aRows" }
+        }
+      ]
+    },
+    {
+      "id": "fused_skip_nogb_noub",
+      "priority": 30,
+      "requires": { "limits": { "maxStorageBuffersPerShaderStage": 8 } },
+      "when": ["mlpShapeOk", "normContractOk", "biasPresence_nogb_noub", "gateUpDispatchFits", "aRows == 1", "present.skipT", "not present.residualT"],
+      "constants": { "inlineNorm": "1", "fromNormed": "0", "rowTile": "1" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "MatMulNBitsMlp.FusedDecode",
+          "shader": "mlp-gate-up.wgsl.jinja",
+          "bindings": "fused_skip_nogb_noub",
+          "dispatch": { "x": "ceilDiv(attrs.N, tunables.TILE_N)", "y": "aRows" }
+        }
+      ]
+    },
+    {
+      "id": "fused_skipsum_nogb_noub",
+      "priority": 30,
+      "requires": { "limits": { "maxStorageBuffersPerShaderStage": 9 } },
+      "when": ["mlpShapeOk", "normContractOk", "biasPresence_nogb_noub", "gateUpDispatchFits", "aRows == 1", "present.skipT", "present.residualT"],
+      "constants": { "inlineNorm": "1", "fromNormed": "0", "rowTile": "1" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "MatMulNBitsMlp.FusedDecode",
+          "shader": "mlp-gate-up.wgsl.jinja",
+          "bindings": "fused_skipsum_nogb_noub",
+          "dispatch": { "x": "ceilDiv(attrs.N, tunables.TILE_N)", "y": "aRows" }
+        }
+      ]
+    },
+    {
+      "id": "fused_norm_nogb_ub",
+      "priority": 30,
+      "requires": { "limits": { "maxStorageBuffersPerShaderStage": 8 } },
+      "when": ["mlpShapeOk", "normContractOk", "biasPresence_nogb_ub", "gateUpDispatchFits", "aRows == 1", "not present.skipT", "not present.residualT"],
+      "constants": { "inlineNorm": "1", "fromNormed": "0", "rowTile": "1" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "MatMulNBitsMlp.FusedDecode",
+          "shader": "mlp-gate-up.wgsl.jinja",
+          "bindings": "fused_norm_nogb_ub",
+          "dispatch": { "x": "ceilDiv(attrs.N, tunables.TILE_N)", "y": "aRows" }
+        }
+      ]
+    },
+    {
+      "id": "fused_skip_nogb_ub",
+      "priority": 30,
+      "requires": { "limits": { "maxStorageBuffersPerShaderStage": 9 } },
+      "when": ["mlpShapeOk", "normContractOk", "biasPresence_nogb_ub", "gateUpDispatchFits", "aRows == 1", "present.skipT", "not present.residualT"],
+      "constants": { "inlineNorm": "1", "fromNormed": "0", "rowTile": "1" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "MatMulNBitsMlp.FusedDecode",
+          "shader": "mlp-gate-up.wgsl.jinja",
+          "bindings": "fused_skip_nogb_ub",
+          "dispatch": { "x": "ceilDiv(attrs.N, tunables.TILE_N)", "y": "aRows" }
+        }
+      ]
+    },
+    {
+      "id": "fused_skipsum_nogb_ub",
+      "priority": 30,
+      "requires": { "limits": { "maxStorageBuffersPerShaderStage": 10 } },
+      "when": ["mlpShapeOk", "normContractOk", "biasPresence_nogb_ub", "gateUpDispatchFits", "aRows == 1", "present.skipT", "present.residualT"],
+      "constants": { "inlineNorm": "1", "fromNormed": "0", "rowTile": "1" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "MatMulNBitsMlp.FusedDecode",
+          "shader": "mlp-gate-up.wgsl.jinja",
+          "bindings": "fused_skipsum_nogb_ub",
+          "dispatch": { "x": "ceilDiv(attrs.N, tunables.TILE_N)", "y": "aRows" }
+        }
+      ]
+    },
+    {
+      "id": "fused_norm_gb_noub",
+      "priority": 30,
+      "requires": { "limits": { "maxStorageBuffersPerShaderStage": 8 } },
+      "when": ["mlpShapeOk", "normContractOk", "biasPresence_gb_noub", "gateUpDispatchFits", "aRows == 1", "not present.skipT", "not present.residualT"],
+      "constants": { "inlineNorm": "1", "fromNormed": "0", "rowTile": "1" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "MatMulNBitsMlp.FusedDecode",
+          "shader": "mlp-gate-up.wgsl.jinja",
+          "bindings": "fused_norm_gb_noub",
+          "dispatch": { "x": "ceilDiv(attrs.N, tunables.TILE_N)", "y": "aRows" }
+        }
+      ]
+    },
+    {
+      "id": "fused_skip_gb_noub",
+      "priority": 30,
+      "requires": { "limits": { "maxStorageBuffersPerShaderStage": 9 } },
+      "when": ["mlpShapeOk", "normContractOk", "biasPresence_gb_noub", "gateUpDispatchFits", "aRows == 1", "present.skipT", "not present.residualT"],
+      "constants": { "inlineNorm": "1", "fromNormed": "0", "rowTile": "1" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "MatMulNBitsMlp.FusedDecode",
+          "shader": "mlp-gate-up.wgsl.jinja",
+          "bindings": "fused_skip_gb_noub",
+          "dispatch": { "x": "ceilDiv(attrs.N, tunables.TILE_N)", "y": "aRows" }
+        }
+      ]
+    },
+    {
+      "id": "fused_skipsum_gb_noub",
+      "priority": 30,
+      "requires": { "limits": { "maxStorageBuffersPerShaderStage": 10 } },
+      "when": ["mlpShapeOk", "normContractOk", "biasPresence_gb_noub", "gateUpDispatchFits", "aRows == 1", "present.skipT", "present.residualT"],
+      "constants": { "inlineNorm": "1", "fromNormed": "0", "rowTile": "1" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "MatMulNBitsMlp.FusedDecode",
+          "shader": "mlp-gate-up.wgsl.jinja",
+          "bindings": "fused_skipsum_gb_noub",
+          "dispatch": { "x": "ceilDiv(attrs.N, tunables.TILE_N)", "y": "aRows" }
+        }
+      ]
+    },
+    {
+      "id": "fused_norm_gb_ub",
+      "priority": 30,
+      "requires": { "limits": { "maxStorageBuffersPerShaderStage": 9 } },
+      "when": ["mlpShapeOk", "normContractOk", "biasPresence_gb_ub", "gateUpDispatchFits", "aRows == 1", "not present.skipT", "not present.residualT"],
+      "constants": { "inlineNorm": "1", "fromNormed": "0", "rowTile": "1" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "MatMulNBitsMlp.FusedDecode",
+          "shader": "mlp-gate-up.wgsl.jinja",
+          "bindings": "fused_norm_gb_ub",
+          "dispatch": { "x": "ceilDiv(attrs.N, tunables.TILE_N)", "y": "aRows" }
+        }
+      ]
+    },
+    {
+      "id": "fused_skip_gb_ub",
+      "priority": 30,
+      "requires": { "limits": { "maxStorageBuffersPerShaderStage": 10 } },
+      "when": ["mlpShapeOk", "normContractOk", "biasPresence_gb_ub", "gateUpDispatchFits", "aRows == 1", "present.skipT", "not present.residualT"],
+      "constants": { "inlineNorm": "1", "fromNormed": "0", "rowTile": "1" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "MatMulNBitsMlp.FusedDecode",
+          "shader": "mlp-gate-up.wgsl.jinja",
+          "bindings": "fused_skip_gb_ub",
+          "dispatch": { "x": "ceilDiv(attrs.N, tunables.TILE_N)", "y": "aRows" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/matmul-nbits-fused-rms-norm.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,117 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+// Shared MatMulNBits MLP/QKV normalization pass.
+//   normed[row, d] = (A + skip)[row, d] * inverseSqrt(mean_d((A + skip)^2) + eps) * norm_scale[d]
+// One workgroup owns one row. Every intermediate stays in f32 and `normed` is an
+// f32 scratch tensor, so the projection pass reads exactly what the fused
+// single-dispatch kernel keeps in registers -- the two paths agree for float16
+// inputs instead of differing by one narrowing.
+const HIDDEN: u32 = {{ hidden }}u;
+const WG: u32 = {{ workgroupSize }}u;
+const EPSILON: f32 = {{ epsilon }};
+var<workgroup> partial: array<f32, WG>;
+{% 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 %}
+// Reusing partial after this reduction requires a barrier between the read of
+// partial[0] and the next write, or the next round can race the prior readers.
+{% set trailingBarrier = trailingBarrier is defined and trailingBarrier %}
+fn reduce_sum(value: f32, tid: u32) -> f32 {
+  partial[tid] = value;
+  workgroupBarrier();
+{{ wgsl_tree_fold(["partial"], idx="tid", wg="WG", form="head") }}
+{% if trailingBarrier %}
+  let total = partial[0];
+  workgroupBarrier();
+  return total;
+{% else %}
+  return partial[0];
+{% endif %}
+}
+fn row_value(index: u32) -> f32 {
+{% if hasSkip %}
+  return f32(a[index]) + f32(skip[index]);
+{% else %}
+  return f32(a[index]);
+{% endif %}
+}
+@compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(workgroup_id) wg: vec3<u32>,
+        @builtin(num_workgroups) nwg: vec3<u32>,
+        @builtin(local_invocation_id) lid: vec3<u32>) {
+  // 2D-folded row index: wg.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit. Reduces to wg.x when nwg.y == 1.
+  let row = wg.x + wg.y * nwg.x;
+  if (row >= params.rows) {
+    return;
+  }
+  let tid = lid.x;
+  let base = row * HIDDEN;
+  var local_sq = 0.0;
+  for (var d = tid; d < HIDDEN; d = d + WG) {
+    let value = row_value(base + d);
+    local_sq = local_sq + value * value;
+  }
+  let inv = inverseSqrt(reduce_sum(local_sq, tid) / f32(HIDDEN) + EPSILON);
+  for (var d = tid; d < HIDDEN; d = d + WG) {
+    let index = base + d;
+    let value = row_value(index);
+{% if writeResidual %}
+    residual[index] = {{ scalar }}(value);
+{% endif %}
+    normed[index] = value * inv * f32(norm_scale[d]);
+  }
+}

build/webgpu/metadata.json ADDED Viewed

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

+{
+  "name": "com.microsoft.MatMulNBitsMlp",
+  "id": "_com_microsoft_matmulnbitsmlp_webgpu_5f0af61",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "FmfZaHEP/sBLMvGjNKkLAHUkkW9+LWNCTLnfAniZXZY=",
+      "manifest.json": "2PUEf3quWni+CsZC4X4XMrR/GVu+mhWWWPqOjzAbbyU=",
+      "matmul-nbits-fused-rms-norm.wgsl.jinja": "4TI9Mc/RZWVgMl6NVG2iwe5OpvqQmswhe6xtGvctNmI=",
+      "mlp-gate-up.wgsl.jinja": "Zncc9Fsa1Vja8Dz8384/+MSAlJXnpb8OwIoGO1ht0rs=",
+      "test.json": "sKKYxHcXolfGhck7PQNdGq3iSM8B58IYf/w5hl6JPLg="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/com.microsoft.MatMulNBitsMlp" }
+}

build/webgpu/mlp-gate-up.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,290 @@

+{% macro matmul_nbits_packed_code(fn="packed_weight", buffer="b", kBlocks="params.kBlocks", blobSize="params.blobSize", bits=4) %}
+fn {{ fn }}(n: u32, block: u32, offset: u32) -> u32 {
+{% if bits == 2 %}
+  let byte_index = offset / 4u;
+  let shift = (offset % 4u) * 2u;
+  let packed_index = (n * {{ kBlocks }} + block) * {{ blobSize }} + byte_index;
+  return ({{ buffer }}[packed_index] >> shift) & 3u;
+{% elif bits == 4 %}
+  let byte_index = offset / 2u;
+  let shift = (offset % 2u) * 4u;
+  let packed_index = (n * {{ kBlocks }} + block) * {{ blobSize }} + byte_index;
+  return ({{ buffer }}[packed_index] >> shift) & 15u;
+{% else %}
+  let packed_index = (n * {{ kBlocks }} + block) * {{ blobSize }} + offset;
+  return {{ buffer }}[packed_index] & 255u;
+{% endif %}
+}
+{%- endmacro %}
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+// com.microsoft.MatMulNBitsMlp, gate/up projection pass.
+//   Y[row, n] = silu(dot(A_norm[row], gate[n]) + gate_bias[n])
+//               * (dot(A_norm[row], up[n]) + up_bias[n])
+// Both projections share the activation row, so one workgroup walks the row once
+// and feeds TILE_N output columns of each projection from it. LANES threads
+// cooperate on one column: thread (column, lane) strides the reduction axis by
+// LANES, and the per-column partials are folded within the column's own lane
+// group at the end. Codes are dequantized as (code - ZERO) * block_scale, with
+// the block scale hoisted out of the inner loop; this operator has no
+// zero-point input, so ZERO is the symmetric midpoint for the bit width.
+// A workgroup covers ROW_TILE activation rows, reusing each unpacked weight code
+// across their accumulators.
+const K: u32 = {{ K }}u;
+const N: u32 = {{ N }}u;
+const BLOCK_SIZE: u32 = {{ blockSize }}u;
+const KBLOCKS: u32 = {{ kBlocks }}u;
+const BLOB_SIZE: u32 = {{ blobSize }}u;
+const TILE_N: u32 = {{ tileN }}u;
+const LANES: u32 = {{ lanes }}u;
+const ROW_TILE: u32 = {{ rowTile }}u;
+const ROWS: u32 = {{ rows }}u;
+const WG: u32 = TILE_N * LANES;
+const ZERO: f32 = {{ defaultZero }};
+// Codes per logical byte and the mask for one code. Physical uint8 storage uses
+// a u32 slot, so packing increases the number of codes returned by each load.
+const BITS: u32 = {{ bits }}u;
+const CODES_PER_BYTE: u32 = {{ codesPerByte }}u;
+const CODE_MASK: u32 = {{ codeMask }}u;
+{% if inlineNorm %}
+const EPSILON: f32 = {{ epsilon }};
+{% endif %}
+{% for stream in ["gate", "up"] %}
+{{ matmul_nbits_packed_code(fn=stream ~ "_code", buffer=stream ~ "_b", kBlocks="KBLOCKS", blobSize="BLOB_SIZE", bits=bits) }}
+// Decode two consecutive reduction-axis codes. Below 8 bits an even offset and
+// its successor share one stored byte; at 8 bits they occupy adjacent slots. An
+// odd offset would straddle bytes, so callers advance by two from an even start.
+fn {{ stream }}_code_pair(n: u32, block: u32, offset: u32) -> vec2<u32> {
+  let base = (n * KBLOCKS + block) * BLOB_SIZE;
+  let shift = (offset % CODES_PER_BYTE) * BITS;
+  let lo = {{ stream }}_b[base + offset / CODES_PER_BYTE];
+  let hi = {{ "lo" if codesPerByte > 1 else (stream ~ "_b[base + offset + 1u]") }};
+  return vec2<u32>((lo >> shift) & CODE_MASK,
+                   (hi >> {{ "(shift + BITS)" if codesPerByte > 1 else "0u" }}) & CODE_MASK);
+}
+{% endfor %}
+fn silu(x: f32) -> f32 {
+  return x / (1.0 + exp(-x));
+}
+var<workgroup> red_gate: array<f32, WG * ROW_TILE>;
+var<workgroup> red_up: array<f32, WG * ROW_TILE>;
+{% if inlineNorm %}
+var<workgroup> partial: array<f32, WG>;
+var<workgroup> row_inv: f32;
+{% 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 %}
+// Reusing partial after this reduction requires a barrier between the read of
+// partial[0] and the next write, or the next round can race the prior readers.
+{% set trailingBarrier = trailingBarrier is defined and trailingBarrier %}
+fn reduce_sum(value: f32, tid: u32) -> f32 {
+  partial[tid] = value;
+  workgroupBarrier();
+{{ wgsl_tree_fold(["partial"], idx="tid", wg="WG", form="head") }}
+{% if trailingBarrier %}
+  let total = partial[0];
+  workgroupBarrier();
+  return total;
+{% else %}
+  return partial[0];
+{% endif %}
+}
+fn row_value(index: u32) -> f32 {
+{% if hasSkip %}
+  return f32(a[index]) + f32(skip[index]);
+{% else %}
+  return f32(a[index]);
+{% endif %}
+}
+{% endif %}
+{% macro act(b, k) %}{% if inlineNorm %}row_value({{ b }} + {{ k }}) * row_inv * f32(norm_scale[{{ k }}]){% elif fromNormed %}normed[{{ b }} + {{ k }}]{% else %}f32(a[{{ b }} + {{ k }}]){% endif %}{%- endmacro %}
+{% macro walk_block(guarded) %}
+        for (var offset = lane * 2u; offset + 1u < BLOCK_SIZE; offset = offset + LANES * 2u) {
+          let k = k_base + offset;
+{% if guarded %}
+          if (k + 1u < K) {
+{% endif %}
+            let gate_codes = gate_code_pair(n, block, offset);
+            let up_codes = up_code_pair(n, block, offset);
+            let gate_lo = f32(gate_codes.x) - ZERO;
+            let gate_hi = f32(gate_codes.y) - ZERO;
+            let up_lo = f32(up_codes.x) - ZERO;
+            let up_hi = f32(up_codes.y) - ZERO;
+{% for r in range(rowTile) %}
+            {
+              let v0 = {{ act("base_" ~ r, "k") }};
+              let v1 = {{ act("base_" ~ r, "k + 1u") }};
+              block_gate_{{ r }} = block_gate_{{ r }} + v0 * gate_lo + v1 * gate_hi;
+              block_up_{{ r }} = block_up_{{ r }} + v0 * up_lo + v1 * up_hi;
+            }
+{% endfor %}
+{% if guarded %}
+          } else if (k < K) {
+            let gate_value = f32(gate_code(n, block, offset)) - ZERO;
+            let up_value = f32(up_code(n, block, offset)) - ZERO;
+{% for r in range(rowTile) %}
+            {
+              let v0 = {{ act("base_" ~ r, "k") }};
+              block_gate_{{ r }} = block_gate_{{ r }} + v0 * gate_value;
+              block_up_{{ r }} = block_up_{{ r }} + v0 * up_value;
+            }
+{% endfor %}
+          }
+{% endif %}
+        }
+{%- endmacro %}
+@compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid: vec3<u32>) {
+  let row0 = wg.y * ROW_TILE;
+  let tid = lid.x;
+  let column = tid / LANES;
+  let lane = tid % LANES;
+  let n = wg.x * TILE_N + column;
+{% for r in range(rowTile) %}
+  // Rows past the end of the batch clamp onto the last real row; their
+  // accumulators stay finite and the store guard drops them.
+  let base_{{ r }} = min(row0 + {{ r }}u, ROWS - 1u) * K;
+{% endfor %}
+{% if inlineNorm %}
+  var local_sq = 0.0;
+  for (var d = tid; d < K; d = d + WG) {
+    let value = row_value(base_0 + d);
+    local_sq = local_sq + value * value;
+  }
+  let inv = inverseSqrt(reduce_sum(local_sq, tid) / f32(K) + EPSILON);
+  if (tid == 0u) {
+    row_inv = inv;
+  }
+  // Separates the reduction's readers of partial[0] from the projection's
+  // reuse of the same workgroup array below.
+  workgroupBarrier();
+{% if writeResidual %}
+  // Every N tile computes the same residual row; only the first one stores it,
+  // so the tiles never write the same location.
+  if (wg.x == 0u) {
+    for (var d = tid; d < K; d = d + WG) {
+      residual[base_0 + d] = {{ scalar }}(row_value(base_0 + d));
+    }
+  }
+{% endif %}
+{% endif %}
+{% for r in range(rowTile) %}
+  var acc_gate_{{ r }} = 0.0;
+  var acc_up_{{ r }} = 0.0;
+{% endfor %}
+  if (n < N) {
+    for (var block = 0u; block < KBLOCKS; block = block + 1u) {
+      let gate_scale = f32(gate_scales[n * KBLOCKS + block]);
+      let up_scale = f32(up_scales[n * KBLOCKS + block]);
+      let k_base = block * BLOCK_SIZE;
+{% for r in range(rowTile) %}
+      var block_gate_{{ r }} = 0.0;
+      var block_up_{{ r }} = 0.0;
+{% endfor %}
+      // Each trip handles two codes. BLOCK_SIZE is even for every admitted
+      // packing, so paired trips cover a full block. Only a final partial block
+      // needs bounds checks; the branch is workgroup-uniform.
+      if (k_base + BLOCK_SIZE <= K) {
+{{ walk_block(false) }}
+      } else {
+{{ walk_block(true) }}
+      }
+{% for r in range(rowTile) %}
+      acc_gate_{{ r }} = acc_gate_{{ r }} + block_gate_{{ r }} * gate_scale;
+      acc_up_{{ r }} = acc_up_{{ r }} + block_up_{{ r }} * up_scale;
+{% endfor %}
+    }
+  }
+{% for r in range(rowTile) %}
+  red_gate[{{ r }}u * WG + tid] = acc_gate_{{ r }};
+  red_up[{{ r }}u * WG + tid] = acc_up_{{ r }};
+{% endfor %}
+  workgroupBarrier();
+  // Fold within each column's own LANES-wide group, one group per staged row.
+  // LANES is a power of two, so lane + stride never leaves the group and no
+  // column can read another's tail.
+  for (var stride = LANES / 2u; stride > 0u; stride = stride / 2u) {
+    if (lane < stride) {
+{% for r in range(rowTile) %}
+      red_gate[{{ r }}u * WG + tid] = red_gate[{{ r }}u * WG + tid] + red_gate[{{ r }}u * WG + tid + stride];
+      red_up[{{ r }}u * WG + tid] = red_up[{{ r }}u * WG + tid] + red_up[{{ r }}u * WG + tid + stride];
+{% endfor %}
+    }
+    workgroupBarrier();
+  }
+  if (lane == 0u && n < N) {
+{% for r in range(rowTile) %}
+    if (row0 + {{ r }}u < ROWS) {
+      var gate_value_{{ r }} = red_gate[{{ r }}u * WG + tid];
+      var up_value_{{ r }} = red_up[{{ r }}u * WG + tid];
+{% if hasGateBias %}
+      gate_value_{{ r }} = gate_value_{{ r }} + f32(gate_bias[n]);
+{% endif %}
+{% if hasUpBias %}
+      up_value_{{ r }} = up_value_{{ r }} + f32(up_bias[n]);
+{% endif %}
+      y[(row0 + {{ r }}u) * N + n] = {{ scalar }}(silu(gate_value_{{ r }}) * up_value_{{ r }});
+    }
+{% endfor %}
+  }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,1703 @@

+{
+  "op": "com.microsoft.MatMulNBitsMlp",
+  "fixtureArrays": {
+    "pinned_plain_gb_ub_input_aT": [0.9181, 1.1011, 1.115, 0.9522, 0.6333, 0.2043, -0.2696, -0.7143, -1.0586, -1.2453, -1.2411, -1.0421, -0.6753, -0.1947, 0.3263, 0.8068, 1.1707, 1.3587, 1.3385, 1.1098, 0.7051, 0.1852, -0.3707, -0.8762, -1.2521, -1.4384, -1.4042, -1.1526, -0.7208, -0.1744, 0.4029, 0.9218],
+    "pinned_plain_gb_ub_input_gateBT": [52, 93, 210, 163, 160, 89, 30, 255, 204, 21, 42, 27, 184, 145, 246, 247, 100, 205, 130, 147, 208, 201, 206, 239, 252, 133, 218, 11, 232, 1, 166, 231],
+    "pinned_plain_gb_ub_input_upBT": [248, 209, 54, 55, 164, 13, 194, 211, 16, 9, 14, 47, 60, 197, 26, 75, 40, 65, 230, 39, 212, 125, 114, 195, 64, 121, 190, 31, 108, 53, 202, 59],
+    "pinned_norm_nogb_noub_input_normScaleT": [1.1719, 1.2817, 1.29, 1.1936, 1.0274, 0.8514, 0.7289, 0.7042, 0.7862, 0.9454, 1.1242, 1.2582, 1.2991, 1.2321, 1.0814, 0.9013]
+  },
+  "cases": [
+    {
+      "name": "plain_nogb_noub",
+      "attrs": { "K": 32, "N": 8, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [3, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [3, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "plain_nogb_ub",
+      "attrs": { "K": 32, "N": 8, "bits": 4, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [3, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        },
+        "upBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [3, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "plain_gb_noub",
+      "attrs": { "K": 32, "N": 8, "bits": 4, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [3, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "gateBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.53, "scale": 0.3 }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [3, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "plain_gb_ub",
+      "attrs": { "K": 32, "N": 8, "bits": 4, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [3, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "gateBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.53, "scale": 0.3 }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        },
+        "upBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [3, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "decode_norm_nogb_noub",
+      "attrs": { "K": 32, "N": 8, "bits": 4, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [1, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "normScaleT": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.4, "offset": 1.0 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [1, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "decode_skip_nogb_noub",
+      "attrs": { "K": 32, "N": 8, "bits": 4, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [1, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [1, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.5 }
+        },
+        "normScaleT": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.4, "offset": 1.0 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [1, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "decode_skipsum_nogb_noub",
+      "attrs": { "K": 32, "N": 8, "bits": 4, "block_size": 16, "activation": "silu", "epsilon": 0.001 },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [1, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [1, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.5 }
+        },
+        "normScaleT": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.4, "offset": 1.0 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        }
+      },
+      "outputs": {
+        "yT": { "dtype": "float32", "shape": [1, 8], "tolerance": 0.0001, "relTolerance": 0.0001 },
+        "residualT": { "dtype": "float32", "shape": [1, 32], "tolerance": 0.000001, "relTolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "decode_norm_nogb_ub",
+      "attrs": { "K": 32, "N": 8, "bits": 4, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [1, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "normScaleT": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.4, "offset": 1.0 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        },
+        "upBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [1, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "decode_skip_nogb_ub",
+      "attrs": { "K": 32, "N": 8, "bits": 4, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [1, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [1, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.5 }
+        },
+        "normScaleT": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.4, "offset": 1.0 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        },
+        "upBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [1, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "decode_skipsum_nogb_ub",
+      "attrs": { "K": 32, "N": 8, "bits": 4, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [1, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [1, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.5 }
+        },
+        "normScaleT": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.4, "offset": 1.0 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        },
+        "upBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": {
+        "yT": { "dtype": "float32", "shape": [1, 8], "tolerance": 0.0001, "relTolerance": 0.0001 },
+        "residualT": { "dtype": "float32", "shape": [1, 32], "tolerance": 0.000001, "relTolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "decode_norm_gb_noub",
+      "attrs": { "K": 32, "N": 8, "bits": 4, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [1, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "normScaleT": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.4, "offset": 1.0 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "gateBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.53, "scale": 0.3 }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [1, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "decode_skip_gb_noub",
+      "attrs": { "K": 32, "N": 8, "bits": 4, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [1, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [1, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.5 }
+        },
+        "normScaleT": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.4, "offset": 1.0 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "gateBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.53, "scale": 0.3 }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [1, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "decode_skipsum_gb_noub",
+      "attrs": { "K": 32, "N": 8, "bits": 4, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [1, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [1, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.5 }
+        },
+        "normScaleT": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.4, "offset": 1.0 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "gateBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.53, "scale": 0.3 }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        }
+      },
+      "outputs": {
+        "yT": { "dtype": "float32", "shape": [1, 8], "tolerance": 0.0001, "relTolerance": 0.0001 },
+        "residualT": { "dtype": "float32", "shape": [1, 32], "tolerance": 0.000001, "relTolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "decode_norm_gb_ub",
+      "attrs": { "K": 32, "N": 8, "bits": 4, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [1, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "normScaleT": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.4, "offset": 1.0 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "gateBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.53, "scale": 0.3 }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        },
+        "upBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [1, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "decode_skip_gb_ub",
+      "attrs": { "K": 32, "N": 8, "bits": 4, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [1, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [1, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.5 }
+        },
+        "normScaleT": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.4, "offset": 1.0 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "gateBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.53, "scale": 0.3 }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        },
+        "upBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [1, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "decode_skipsum_gb_ub",
+      "attrs": { "K": 32, "N": 8, "bits": 4, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [1, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [1, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.5 }
+        },
+        "normScaleT": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.4, "offset": 1.0 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "gateBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.53, "scale": 0.3 }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        },
+        "upBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": {
+        "yT": { "dtype": "float32", "shape": [1, 8], "tolerance": 0.0001, "relTolerance": 0.0001 },
+        "residualT": { "dtype": "float32", "shape": [1, 32], "tolerance": 0.000001, "relTolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "prefill_norm_nogb_noub",
+      "attrs": { "K": 32, "N": 8, "bits": 4, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [5, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "normScaleT": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.4, "offset": 1.0 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [5, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "prefill_norm_nogb_ub",
+      "attrs": { "K": 32, "N": 8, "bits": 4, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [5, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "normScaleT": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.4, "offset": 1.0 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        },
+        "upBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [5, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "prefill_norm_gb_noub",
+      "attrs": { "K": 32, "N": 8, "bits": 4, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [5, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "normScaleT": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.4, "offset": 1.0 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "gateBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.53, "scale": 0.3 }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [5, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "prefill_norm_gb_ub",
+      "attrs": { "K": 32, "N": 8, "bits": 4, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [5, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "normScaleT": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.4, "offset": 1.0 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "gateBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.53, "scale": 0.3 }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        },
+        "upBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [5, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "prefill_skip_gb_ub",
+      "attrs": { "K": 32, "N": 8, "bits": 4, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [5, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [5, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.5 }
+        },
+        "normScaleT": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.4, "offset": 1.0 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "gateBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.53, "scale": 0.3 }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        },
+        "upBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [5, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "prefill_skipsum_gb_ub",
+      "attrs": { "K": 32, "N": 8, "bits": 4, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [5, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [5, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.5 }
+        },
+        "normScaleT": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.4, "offset": 1.0 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "gateBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.53, "scale": 0.3 }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        },
+        "upBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": {
+        "yT": { "dtype": "float32", "shape": [5, 8], "tolerance": 0.0001, "relTolerance": 0.0001 },
+        "residualT": { "dtype": "float32", "shape": [5, 32], "tolerance": 0.000001, "relTolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "tail_column_not_tile_aligned",
+      "provenance": {
+        "notes": "N is not a multiple of the column tile, so the last tile's out-of-range lanes must contribute nothing."
+      },
+      "attrs": { "K": 32, "N": 5, "bits": 4, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [2, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "normScaleT": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.4, "offset": 1.0 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [5, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [5, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "gateBiasT": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.53, "scale": 0.3 }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [5, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [5, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        },
+        "upBiasT": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [2, 5], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "tail_block_partial_k",
+      "provenance": { "notes": "K is not a multiple of block_size, so the final block is partly out of range." },
+      "attrs": { "K": 20, "N": 8, "bits": 4, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [2, 20],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "normScaleT": {
+          "dtype": "float32",
+          "shape": [20],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.4, "offset": 1.0 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "gateBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.53, "scale": 0.3 }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        },
+        "upBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [2, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "rank3_activation",
+      "provenance": { "notes": "Rank-3 activation: the leading axes fold into the row count." },
+      "attrs": { "K": 32, "N": 8, "bits": 4, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [1, 4, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "normScaleT": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.4, "offset": 1.0 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "gateBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.53, "scale": 0.3 }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        },
+        "upBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [1, 4, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "bits2_decode",
+      "provenance": { "notes": "2-bit codes pack four per stored byte with a default zero point of 2." },
+      "attrs": { "K": 32, "N": 8, "bits": 2, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [1, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "normScaleT": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.4, "offset": 1.0 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 4],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "gateBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.53, "scale": 0.3 }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 4],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        },
+        "upBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [1, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "bits8_prefill",
+      "provenance": { "notes": "8-bit codes are one byte per weight with a default zero point of 128." },
+      "attrs": { "K": 32, "N": 8, "bits": 8, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [3, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "normScaleT": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.4, "offset": 1.0 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 16],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "gateBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.53, "scale": 0.3 }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 16],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        },
+        "upBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [3, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "block32_decode",
+      "provenance": { "notes": "block_size 32, the size ONNX Runtime's fused decode kernel is specialized for." },
+      "attrs": { "K": 64, "N": 8, "bits": 4, "block_size": 32, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [1, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [1, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.5 }
+        },
+        "normScaleT": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.4, "offset": 1.0 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 16],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "gateBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.53, "scale": 0.3 }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 16],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        },
+        "upBiasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": {
+        "yT": { "dtype": "float32", "shape": [1, 8], "tolerance": 0.0001, "relTolerance": 0.0001 },
+        "residualT": { "dtype": "float32", "shape": [1, 64], "tolerance": 0.000001, "relTolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "f16_decode_skipsum",
+      "attrs": { "K": 32, "N": 8, "bits": 4, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [1, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "skipT": {
+          "dtype": "float16",
+          "shape": [1, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.5 }
+        },
+        "normScaleT": {
+          "dtype": "float16",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.4, "offset": 1.0 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float16",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "gateBiasT": {
+          "dtype": "float16",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.53, "scale": 0.3 }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float16",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        },
+        "upBiasT": {
+          "dtype": "float16",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": {
+        "yT": { "dtype": "float16", "shape": [1, 8], "tolerance": 0.002, "relTolerance": 0.01 },
+        "residualT": { "dtype": "float16", "shape": [1, 32], "tolerance": 0.002, "relTolerance": 0.002 }
+      }
+    },
+    {
+      "name": "f16_prefill_norm",
+      "attrs": { "K": 32, "N": 8, "bits": 4, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [4, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "normScaleT": {
+          "dtype": "float16",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.4, "offset": 1.0 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float16",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "gateBiasT": {
+          "dtype": "float16",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.53, "scale": 0.3 }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float16",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        },
+        "upBiasT": {
+          "dtype": "float16",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [4, 8], "tolerance": 0.002, "relTolerance": 0.01 } }
+    },
+    {
+      "name": "f16_plain",
+      "attrs": { "K": 32, "N": 8, "bits": 4, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [3, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float16",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "gateBiasT": {
+          "dtype": "float16",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.53, "scale": 0.3 }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float16",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        },
+        "upBiasT": {
+          "dtype": "float16",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [3, 8], "tolerance": 0.002, "relTolerance": 0.01 } }
+    },
+    {
+      "name": "pinned_plain_gb_ub",
+      "provenance": {
+        "notes": "Expected values computed by an independent implementation written from the ONNX Runtime schema text alone, so this case checks the trusted reference as well as the kernels. No normalization: the projections read A directly."
+      },
+      "attrs": { "K": 16, "N": 4, "bits": 4, "block_size": 8, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [2, 16],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/pinned_plain_gb_ub_input_aT" } }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [4, 2, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/pinned_plain_gb_ub_input_gateBT" } }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [4, 2],
+          "data": { "kind": "values", "values": [0.03, 0.04, 0.05, 0.06, 0.07, 0.08, 0.09, 0.1] }
+        },
+        "gateBiasT": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [0.1782, 0.1617, -0.0315, -0.1903] }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [4, 2, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/pinned_plain_gb_ub_input_upBT" } }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [4, 2],
+          "data": { "kind": "values", "values": [0.05, 0.046, 0.042, 0.038, 0.034, 0.03, 0.026, 0.022] }
+        },
+        "upBiasT": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [0.0932, -0.0341, -0.1356, -0.1345] }
+        }
+      },
+      "outputs": {
+        "yT": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.2154822, -0.4471286, 0.0238219, -0.4921483, 0.272635, -0.6797003, -0.0049887, -0.738616]
+          },
+          "tolerance": 0.00001,
+          "relTolerance": 0.0001
+        }
+      }
+    },
+    {
+      "name": "pinned_norm_nogb_noub",
+      "provenance": {
+        "notes": "Expected values computed by an independent implementation written from the ONNX Runtime schema text alone, so this case checks the trusted reference as well as the kernels. SimplifiedLayerNormalization with no biases."
+      },
+      "attrs": { "K": 16, "N": 4, "bits": 4, "block_size": 8, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [2, 16],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/pinned_plain_gb_ub_input_aT" } }
+        },
+        "normScaleT": {
+          "dtype": "float32",
+          "shape": [16],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/pinned_norm_nogb_noub_input_normScaleT" } }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [4, 2, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/pinned_plain_gb_ub_input_gateBT" } }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [4, 2],
+          "data": { "kind": "values", "values": [0.03, 0.04, 0.05, 0.06, 0.07, 0.08, 0.09, 0.1] }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [4, 2, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/pinned_plain_gb_ub_input_upBT" } }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [4, 2],
+          "data": { "kind": "values", "values": [0.05, 0.046, 0.042, 0.038, 0.034, 0.03, 0.026, 0.022] }
+        }
+      },
+      "outputs": {
+        "yT": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.0651757, -0.5406582, 0.1351256, -0.4018507, 0.0572953, -0.6108609, 0.1088786, -0.5550907]
+          },
+          "tolerance": 0.00001,
+          "relTolerance": 0.0001
+        }
+      }
+    },
+    {
+      "name": "pinned_skipsum_gb_ub",
+      "provenance": {
+        "notes": "Expected values computed by an independent implementation written from the ONNX Runtime schema text alone, so this case checks the trusted reference as well as the kernels. SkipSimplifiedLayerNormalization with both biases and the residual-sum output."
+      },
+      "attrs": { "K": 16, "N": 4, "bits": 4, "block_size": 8, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [2, 16],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/pinned_plain_gb_ub_input_aT" } }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [2, 16],
+          "data": {
+            "kind": "values",
+            "values": [0.0537, 0.2754, 0.2738, 0.0827, -0.1607, -0.2913, -0.2055, 0.0807, 0.4361, 0.6796, 0.672, 0.3892, -0.0618, -0.4922, -0.7224, -0.6678, -0.3779, -0.0076, 0.2621, 0.3158, 0.1593, -0.0878, -0.2585, -0.2289, 0.0133, 0.3593, 0.6324, 0.678, 0.4447, 0.0142, -0.4352, -0.7141]
+          }
+        },
+        "normScaleT": {
+          "dtype": "float32",
+          "shape": [16],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/pinned_norm_nogb_noub_input_normScaleT" } }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [4, 2, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/pinned_plain_gb_ub_input_gateBT" } }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [4, 2],
+          "data": { "kind": "values", "values": [0.03, 0.04, 0.05, 0.06, 0.07, 0.08, 0.09, 0.1] }
+        },
+        "gateBiasT": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [0.1782, 0.1617, -0.0315, -0.1903] }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [4, 2, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/pinned_plain_gb_ub_input_upBT" } }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [4, 2],
+          "data": { "kind": "values", "values": [0.05, 0.046, 0.042, 0.038, 0.034, 0.03, 0.026, 0.022] }
+        },
+        "upBiasT": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [0.0932, -0.0341, -0.1356, -0.1345] }
+        }
+      },
+      "outputs": {
+        "yT": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": {
+            "kind": "values",
+            "values": [-0.052766, -0.1362556, 0.2952684, -1.5736477, 0.0464634, -0.1608242, -0.1441075, -0.5477401]
+          },
+          "tolerance": 0.00001,
+          "relTolerance": 0.0001
+        },
+        "residualT": {
+          "dtype": "float32",
+          "shape": [2, 16],
+          "data": {
+            "kind": "values",
+            "values": [0.9718, 1.3765, 1.3888, 1.0349, 0.4726, -0.087, -0.4751, -0.6336, -0.6225, -0.5657, -0.5691, -0.6529, -0.7371, -0.6869, -0.3961, 0.139, 0.7928, 1.3511, 1.6006, 1.4256, 0.8644, 0.0974, -0.6292, -1.1051, -1.2388, -1.0791, -0.7718, -0.4746, -0.2761, -0.1602, -0.0323, 0.2077]
+          },
+          "tolerance": 0.000001,
+          "relTolerance": 0.000001
+        }
+      }
+    },
+    {
+      "name": "prefill_skip_nogb_noub",
+      "provenance": {
+        "notes": "Five activation rows force the two-pass prefill schedule. Supplying skip and norm_scale while omitting both projection biases exercises staged SkipSimplifiedLayerNormalization without the residual-sum output."
+      },
+      "attrs": { "K": 32, "N": 8, "bits": 4, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [5, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "skipT": {
+          "dtype": "float32",
+          "shape": [5, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.5 }
+        },
+        "normScaleT": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.4, "offset": 1.0 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [5, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "norm_rows_past_one_tile",
+      "provenance": {
+        "notes": "Ten activation rows against a ROW_TILE of eight, so the row axis dispatches two groups and the second holds two real rows and six that clamp onto the last one. It is the only case where the store guard has anything to drop; every other multi-row case fits one group, where the guard cannot fire."
+      },
+      "attrs": { "K": 32, "N": 8, "bits": 4, "block_size": 16, "activation": "silu" },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [10, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.8 }
+        },
+        "normScaleT": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.4, "offset": 1.0 }
+        },
+        "gateBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195] }
+        },
+        "gateScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "upBT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 8],
+          "data": { "kind": "cycle", "values": [211, 44, 137, 98, 165, 20, 233, 121, 78, 190, 15, 252, 87, 143, 61] }
+        },
+        "upScalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": [0.045, 0.03, 0.07, 0.05, 0.08, 0.035] }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [10, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    }
+  ]
+}