sync 91d990483a17

Browse files

Files changed (12) hide show

README.md +28 -15
build/webgpu/bench.json +725 -17
build/webgpu/chunk-out.wgsl.jinja +1 -2
build/webgpu/chunk-prep.wgsl.jinja +2 -3
build/webgpu/chunk-scan.wgsl.jinja +1 -2
build/webgpu/chunk-ut.wgsl.jinja +1 -2
build/webgpu/linear-attention.scalar.wgsl.jinja +2 -3
build/webgpu/linear-attention.serial.wgsl.jinja +29 -4
build/webgpu/linear-attention.vec4.wgsl.jinja +54 -14
build/webgpu/manifest.json +0 -0
build/webgpu/metadata.json +45 -13
build/webgpu/test.json +840 -34

README.md CHANGED Viewed

@@ -18,21 +18,21 @@ See the [ONNX Runtime `LinearAttention` contrib-operator spec](https://github.co
 ## Inputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
-| `query` | `queryT` | `T` | `3` | — | Query vectors with 3D packed shape `(B, T, H_q * d_k)`; heads are packed into the last dimension. | required |
-| `key` | `keyT` | `T` | `3` | — | Key vectors with 3D packed shape `(B, T, H_k * d_k)`, where positive `H_k` divides `H_kv`; `H_k < H_kv` shares each key head across multiple KV-state heads. Keys should be L2-normalized for `delta`/`gated_delta` modes. | required |
-| `value` | `valueT` | `T` | `3` | — | Value vectors with 3D packed shape `(B, T, H_kv * d_v)`. | required |
-| `past_state` | `pastStateT` | `S` | derived | derived; see description | Recurrent state from the previous step with shape `(B, H_kv, d_k, d_v)`, or `(W, B, H_kv, d_k, d_v)` when `state_window = W > 0`; defaults to zeros if absent. | optional |
-| `decay` | `decayT` | `T` | `3` | — | Exponential decay gate in log-space with shape `(B, T, H_kv * d_k)` or `(B, T, H_kv)`; required for `gated` and `gated_delta` modes. | optional |
-| `beta` | `betaT` | `T` | `3` | — | Update rate (sigmoid output) with shape `(B, T, H_kv)` or `(B, T, 1)`; required for `delta` and `gated_delta` modes. | optional |
 ## Outputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
-| `output` | `outputT` | `T` | `3` | derived; see description | Attention output with 3D packed shape `(B, T, max(H_q, H_kv) * d_v)`. | required |
-| `present_state` | `presentStateT` | `S` | derived | derived; see description | Updated recurrent state with shape `(B, H_kv, d_k, d_v)`, or `(W, B, H_kv, d_k, d_v)` when `state_window = W > 0`. | required |
 ## Attributes
@@ -41,11 +41,11 @@ Attributes and default values (overridable per request):
 | Attribute | Default | Description |
 | --- | --- | --- |
 | `chunk_size` | `64` | Accepted for schema compatibility; does not affect the result. |
 | `scale` | `0` | Scale applied to query-key products. Zero selects `1 / sqrt(d_k)`. |
 | `state_window` | `0` | Number of recent recurrent states retained in `present_state`, in the supported range 0 to 8; zero returns only the current state. |
 | `update_rule` | `"gated_delta"` | Recurrent update rule: `linear`, `gated`, `delta`, or `gated_delta`. |
-| `kv_num_heads` | — | Number of key/value heads. |
-| `q_num_heads` | — | Number of query heads. |
 ## Type constraints
@@ -54,9 +54,18 @@ Attributes and default values (overridable per request):
 | `T` | `float32`, `float16` |
 | `S` | `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
@@ -70,10 +79,14 @@ Attributes and default values (overridable per request):
 ## 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.

 ## Inputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
+| `queryT` | `query` | `T` | `3` | — | Query vectors with 3D packed shape `(B, T, H_q * d_k)`; heads are packed into the last dimension. | required |
+| `keyT` | `key` | `T` | `3` | — | Key vectors with 3D packed shape `(B, T, H_k * d_k)`, where positive `H_k` divides `H_kv`; `H_k < H_kv` shares each key head across multiple KV-state heads. Keys should be L2-normalized for `delta`/`gated_delta` modes. | required |
+| `valueT` | `value` | `T` | `3` | — | Value vectors with 3D packed shape `(B, T, H_kv * d_v)`. | required |
+| `pastStateT` | `past_state` | `S` | derived | derived | Recurrent state from the previous step with shape `(B, H_kv, d_k, d_v)`, or `(W, B, H_kv, d_k, d_v)` when `state_window = W > 0`; defaults to zeros if absent. | optional |
+| `decayT` | `decay` | `T` | `3` | — | Exponential decay gate in log-space with shape `(B, T, H_kv * d_k)` or `(B, T, H_kv)`; required for `gated` and `gated_delta` modes. | optional |
+| `betaT` | `beta` | `T` | `3` | — | Update rate (sigmoid output) with shape `(B, T, H_kv)` or `(B, T, 1)`; required for `delta` and `gated_delta` modes. | optional |
 ## Outputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
+| `outputT` | `output` | `T` | `3` | derived | Attention output with 3D packed shape `(B, T, max(H_q, H_kv) * d_v)`. | required |
+| `presentStateT` | `present_state` | `S` | derived | derived | Updated recurrent state with shape `(B, H_kv, d_k, d_v)`, or `(W, B, H_kv, d_k, d_v)` when `state_window = W > 0`. | required |
 ## Attributes
 | Attribute | Default | Description |
 | --- | --- | --- |
 | `chunk_size` | `64` | Accepted for schema compatibility; does not affect the result. |
+| `kv_num_heads` | — | Number of key/value heads. |
+| `q_num_heads` | — | Number of query heads. |
 | `scale` | `0` | Scale applied to query-key products. Zero selects `1 / sqrt(d_k)`. |
 | `state_window` | `0` | Number of recent recurrent states retained in `present_state`, in the supported range 0 to 8; zero returns only the current state. |
 | `update_rule` | `"gated_delta"` | Recurrent update rule: `linear`, `gated`, `delta`, or `gated_delta`. |
 ## Type constraints
 | `T` | `float32`, `float16` |
 | `S` | `float32`, `float16` |
+## Implementation variants
+One implementation is selected per call from the device capabilities, the request shapes and the dtypes; these notes say what each one covers.
+- `linear_zero_serial_small_dk` — Small key dimensions keep each state column private across the sequence. Linear updates specialize shape-known indexing and loop bounds, while gated-delta updates retain their uniform-driven recurrence.
+- `linear_state_serial_small_dk` — Small key dimensions keep each state column private across the sequence. Linear updates specialize shape-known indexing and loop bounds, while gated-delta updates retain their uniform-driven recurrence.
+- `gated_delta_zero_serial_small_dk` — Small key dimensions keep each state column private across the sequence. Linear updates specialize shape-known indexing and loop bounds, while gated-delta updates retain their uniform-driven recurrence.
+- `gated_delta_state_serial_small_dk` — Small key dimensions keep each state column private across the sequence. Linear updates specialize shape-known indexing and loop bounds, while gated-delta updates retain their uniform-driven recurrence.
 ## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
 - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
 - [`test.json`](build/webgpu/test.json) — correctness cases
 - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
 ## Use with `@huggingface/kernels`
+```sh
+npm install --save-exact @huggingface/kernels@0.0.1-preview.2
+```
+Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
 The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+It follows the `v1` branch as fixes land. To pin exact artifact bytes, pass a 40-character commit `revision` instead of `version`.
 Replace each `*Data` placeholder with a typed array containing the corresponding input data.

build/webgpu/bench.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "com.microsoft.LinearAttention",
   "tunableSpace": {
     "dvGroups": [2, 4, 8],
     "tileV": [4, 8, 16],
@@ -81,7 +80,7 @@
       "name": "linear-attention-linear-state-scalar-f16-seq1536-pathology",
       "preset": "stress",
       "provenance": {
-        "source": "authored for branch coverage",
         "notes": "Long-sequence supplied-state case at head_dim_k 16. It exercises the recurrent small-dk route's serial token recurrence and distinguishes it from the chunked prefill decomposition."
       },
       "vars": { "batch": 4, "seq": 1536, "qHeads": 4, "kvHeads": 2, "dk": 16, "dv": 16 },
@@ -151,8 +150,8 @@
         ]
       },
       "provenance": {
-        "source": "authored for branch coverage",
-        "notes": "Zero-state sibling of the supplied-state long-sequence case, covering the same branch and shape without an entry state."
       }
     },
     {
@@ -213,7 +212,7 @@
       "name": "linear-attention-qwen3next-decode-s1",
       "preset": "model",
       "provenance": {
-        "notes": "Qwen3-Next class defaults (linear_num_value_heads 32, linear_num_key_heads 16, linear_key_head_dim 128, linear_value_head_dim 128) at a decode step, where the recurrence carries the whole cost."
       },
       "vars": { "batch": 1, "seq": 1, "qHeads": 32, "kvHeads": 16, "dk": 128, "dv": 128 },
       "attrs": {
@@ -225,11 +224,33 @@
       },
       "inputs": {
         "queryT": { "shape": [1, 1, 4096], "dtype": "float32", "dist": "normal", "seed": 8100, "scale": 0.3 },
-        "keyT": { "shape": [1, 1, 2048], "dtype": "float32", "dist": "normal", "seed": 8101, "scale": 0.3 },
         "valueT": { "shape": [1, 1, 2048], "dtype": "float32", "dist": "normal", "seed": 8102, "scale": 0.3 },
         "pastStateT": { "shape": [1, 16, 128, 128], "dtype": "float32", "dist": "normal", "seed": 8103, "scale": 0.1 },
-        "decayT": { "shape": [1, 1, 2048], "dtype": "float32", "dist": "uniform", "seed": 8104, "min": 0.9, "max": 1 },
-        "betaT": { "shape": [1, 1, 16], "dtype": "float32", "dist": "uniform", "seed": 8105, "min": 0.1, "max": 0.9 }
       },
       "outputs": {
         "outputT": { "shape": [1, 1, 4096], "dtype": "float32" },
@@ -247,7 +268,9 @@
     {
       "name": "linear-attention-qwen3next-prefill-s512",
       "preset": "model",
-      "provenance": { "notes": "Qwen3-Next class defaults over a 512-token prefill chunk." },
       "vars": { "batch": 1, "seq": 512, "qHeads": 32, "kvHeads": 16, "dk": 128, "dv": 128 },
       "attrs": {
         "q_num_heads": 32,
@@ -258,11 +281,33 @@
       },
       "inputs": {
         "queryT": { "shape": [1, 512, 4096], "dtype": "float32", "dist": "normal", "seed": 8200, "scale": 0.3 },
-        "keyT": { "shape": [1, 512, 2048], "dtype": "float32", "dist": "normal", "seed": 8201, "scale": 0.3 },
         "valueT": { "shape": [1, 512, 2048], "dtype": "float32", "dist": "normal", "seed": 8202, "scale": 0.3 },
         "pastStateT": { "shape": [1, 16, 128, 128], "dtype": "float32", "dist": "normal", "seed": 8203, "scale": 0.1 },
-        "decayT": { "shape": [1, 512, 2048], "dtype": "float32", "dist": "uniform", "seed": 8204, "min": 0.9, "max": 1 },
-        "betaT": { "shape": [1, 512, 16], "dtype": "float32", "dist": "uniform", "seed": 8205, "min": 0.1, "max": 0.9 }
       },
       "outputs": {
         "outputT": { "shape": [1, 512, 4096], "dtype": "float32" },
@@ -281,7 +326,7 @@
       "name": "linear-attention-qwen3next-prefill-s2048",
       "preset": "model",
       "provenance": {
-        "notes": "Qwen3-Next class defaults over a 2048-token prefill chunk, eight chunk_size 64 blocks per workgroup pass."
       },
       "vars": { "batch": 1, "seq": 2048, "qHeads": 32, "kvHeads": 16, "dk": 128, "dv": 128 },
       "attrs": {
@@ -293,7 +338,13 @@
       },
       "inputs": {
         "queryT": { "shape": [1, 2048, 4096], "dtype": "float32", "dist": "normal", "seed": 8300, "scale": 0.3 },
-        "keyT": { "shape": [1, 2048, 2048], "dtype": "float32", "dist": "normal", "seed": 8301, "scale": 0.3 },
         "valueT": { "shape": [1, 2048, 2048], "dtype": "float32", "dist": "normal", "seed": 8302, "scale": 0.3 },
         "pastStateT": { "shape": [1, 16, 128, 128], "dtype": "float32", "dist": "normal", "seed": 8303, "scale": 0.1 },
         "decayT": {
@@ -301,10 +352,19 @@
           "dtype": "float32",
           "dist": "uniform",
           "seed": 8304,
-          "min": 0.9,
-          "max": 1
         },
-        "betaT": { "shape": [1, 2048, 16], "dtype": "float32", "dist": "uniform", "seed": 8305, "min": 0.1, "max": 0.9 }
       },
       "outputs": {
         "outputT": { "shape": [1, 2048, 4096], "dtype": "float32" },
@@ -318,6 +378,654 @@
           }
         ]
       }
     }
   ]
 }

 {
   "tunableSpace": {
     "dvGroups": [2, 4, 8],
     "tileV": [4, 8, 16],
       "name": "linear-attention-linear-state-scalar-f16-seq1536-pathology",
       "preset": "stress",
       "provenance": {
+        "source": "synthetic benchmark",
         "notes": "Long-sequence supplied-state case at head_dim_k 16. It exercises the recurrent small-dk route's serial token recurrence and distinguishes it from the chunked prefill decomposition."
       },
       "vars": { "batch": 4, "seq": 1536, "qHeads": 4, "kvHeads": 2, "dk": 16, "dv": 16 },
         ]
       },
       "provenance": {
+        "source": "synthetic benchmark",
+        "notes": "A long-sequence scalar recurrence uses the zero-state branch with no entry state."
       }
     },
     {
       "name": "linear-attention-qwen3next-decode-s1",
       "preset": "model",
       "provenance": {
+        "notes": "Qwen3-Next class defaults (linear_num_value_heads 32, linear_num_key_heads 16, linear_key_head_dim 128, linear_value_head_dim 128) at a decode step, where the recurrence carries the whole cost. Float uniforms use scale/offset: log-decay spans log(0.9) to 0, beta spans 0.1 to 0.9, and keys are bounded by 1/sqrt(dk) so squared key norms do not exceed one."
       },
       "vars": { "batch": 1, "seq": 1, "qHeads": 32, "kvHeads": 16, "dk": 128, "dv": 128 },
       "attrs": {
       },
       "inputs": {
         "queryT": { "shape": [1, 1, 4096], "dtype": "float32", "dist": "normal", "seed": 8100, "scale": 0.3 },
+        "keyT": {
+          "shape": [1, 1, 2048],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 8101,
+          "scale": 0.08838834764831843
+        },
         "valueT": { "shape": [1, 1, 2048], "dtype": "float32", "dist": "normal", "seed": 8102, "scale": 0.3 },
         "pastStateT": { "shape": [1, 16, 128, 128], "dtype": "float32", "dist": "normal", "seed": 8103, "scale": 0.1 },
+        "decayT": {
+          "shape": [1, 1, 2048],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 8104,
+          "signed": false,
+          "offset": -0.10536051565782628,
+          "scale": 0.10536051565782628
+        },
+        "betaT": {
+          "shape": [1, 1, 16],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 8105,
+          "signed": false,
+          "offset": 0.1,
+          "scale": 0.8
+        }
       },
       "outputs": {
         "outputT": { "shape": [1, 1, 4096], "dtype": "float32" },
     {
       "name": "linear-attention-qwen3next-prefill-s512",
       "preset": "model",
+      "provenance": {
+        "notes": "Qwen3-Next class defaults over a 512-token prefill chunk. Float uniforms use scale/offset: log-decay spans log(0.9) to 0, beta spans 0.1 to 0.9, and keys are bounded by 1/sqrt(dk) so squared key norms do not exceed one."
+      },
       "vars": { "batch": 1, "seq": 512, "qHeads": 32, "kvHeads": 16, "dk": 128, "dv": 128 },
       "attrs": {
         "q_num_heads": 32,
       },
       "inputs": {
         "queryT": { "shape": [1, 512, 4096], "dtype": "float32", "dist": "normal", "seed": 8200, "scale": 0.3 },
+        "keyT": {
+          "shape": [1, 512, 2048],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 8201,
+          "scale": 0.08838834764831843
+        },
         "valueT": { "shape": [1, 512, 2048], "dtype": "float32", "dist": "normal", "seed": 8202, "scale": 0.3 },
         "pastStateT": { "shape": [1, 16, 128, 128], "dtype": "float32", "dist": "normal", "seed": 8203, "scale": 0.1 },
+        "decayT": {
+          "shape": [1, 512, 2048],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 8204,
+          "signed": false,
+          "offset": -0.10536051565782628,
+          "scale": 0.10536051565782628
+        },
+        "betaT": {
+          "shape": [1, 512, 16],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 8205,
+          "signed": false,
+          "offset": 0.1,
+          "scale": 0.8
+        }
       },
       "outputs": {
         "outputT": { "shape": [1, 512, 4096], "dtype": "float32" },
       "name": "linear-attention-qwen3next-prefill-s2048",
       "preset": "model",
       "provenance": {
+        "notes": "Qwen3-Next class defaults over a 2048-token prefill chunk, eight chunk_size 64 blocks per workgroup pass. Float uniforms use scale/offset: log-decay spans log(0.9) to 0, beta spans 0.1 to 0.9, and keys are bounded by 1/sqrt(dk) so squared key norms do not exceed one."
       },
       "vars": { "batch": 1, "seq": 2048, "qHeads": 32, "kvHeads": 16, "dk": 128, "dv": 128 },
       "attrs": {
       },
       "inputs": {
         "queryT": { "shape": [1, 2048, 4096], "dtype": "float32", "dist": "normal", "seed": 8300, "scale": 0.3 },
+        "keyT": {
+          "shape": [1, 2048, 2048],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 8301,
+          "scale": 0.08838834764831843
+        },
         "valueT": { "shape": [1, 2048, 2048], "dtype": "float32", "dist": "normal", "seed": 8302, "scale": 0.3 },
         "pastStateT": { "shape": [1, 16, 128, 128], "dtype": "float32", "dist": "normal", "seed": 8303, "scale": 0.1 },
         "decayT": {
           "dtype": "float32",
           "dist": "uniform",
           "seed": 8304,
+          "signed": false,
+          "offset": -0.10536051565782628,
+          "scale": 0.10536051565782628
         },
+        "betaT": {
+          "shape": [1, 2048, 16],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 8305,
+          "signed": false,
+          "offset": 0.1,
+          "scale": 0.8
+        }
       },
       "outputs": {
         "outputT": { "shape": [1, 2048, 4096], "dtype": "float32" },
           }
         ]
       }
+    },
+    {
+      "name": "bench-linear_geometry_float32_float32_g0_zero",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "linear", "state_window": 3, "scale": 0.375 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float32",
+          "shape": [2, 33, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float32",
+          "shape": [2, 33, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [2, 33, 14],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 33, 28] },
+        "presentStateT": { "dtype": "float32", "shape": [3, 2, 2, 8, 7] }
+      },
+      "preset": "stress",
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "dim(shapes.queryT, 0) * dim(shapes.queryT, 1) * (dim(shapes.valueT, 2) / attrs.kv_num_heads) * (2 * (dim(shapes.queryT, 2) / attrs.q_num_heads) * (attrs.kv_num_heads + max(attrs.q_num_heads, attrs.kv_num_heads)) + max(attrs.q_num_heads, attrs.kv_num_heads))"
+          }
+        ]
+      }
+    },
+    {
+      "name": "bench-linear_geometry_float32_float32_g0_state",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "linear", "state_window": 3, "scale": 0.375 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float32",
+          "shape": [2, 33, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float32",
+          "shape": [2, 33, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [2, 33, 14],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [3, 2, 2, 8, 7],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 33, 28] },
+        "presentStateT": { "dtype": "float32", "shape": [3, 2, 2, 8, 7] }
+      },
+      "preset": "stress",
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "dim(shapes.queryT, 0) * dim(shapes.queryT, 1) * (dim(shapes.valueT, 2) / attrs.kv_num_heads) * (2 * (dim(shapes.queryT, 2) / attrs.q_num_heads) * (attrs.kv_num_heads + max(attrs.q_num_heads, attrs.kv_num_heads)) + max(attrs.q_num_heads, attrs.kv_num_heads))"
+          }
+        ]
+      }
+    },
+    {
+      "name": "bench-linear_geometry_float32_float32_g1_zero",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 2, "kv_num_heads": 4, "update_rule": "linear", "state_window": 0, "scale": 0 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float32",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float32",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [1, 17, 20],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 17, 20] },
+        "presentStateT": { "dtype": "float32", "shape": [1, 4, 6, 5] }
+      },
+      "preset": "stress",
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "dim(shapes.queryT, 0) * dim(shapes.queryT, 1) * (dim(shapes.valueT, 2) / attrs.kv_num_heads) * (2 * (dim(shapes.queryT, 2) / attrs.q_num_heads) * (attrs.kv_num_heads + max(attrs.q_num_heads, attrs.kv_num_heads)) + max(attrs.q_num_heads, attrs.kv_num_heads))"
+          }
+        ]
+      }
+    },
+    {
+      "name": "bench-linear_geometry_float32_float32_g1_state",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 2, "kv_num_heads": 4, "update_rule": "linear", "state_window": 0, "scale": 0 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float32",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float32",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [1, 17, 20],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [1, 4, 6, 5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 17, 20] },
+        "presentStateT": { "dtype": "float32", "shape": [1, 4, 6, 5] }
+      },
+      "preset": "stress",
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "dim(shapes.queryT, 0) * dim(shapes.queryT, 1) * (dim(shapes.valueT, 2) / attrs.kv_num_heads) * (2 * (dim(shapes.queryT, 2) / attrs.q_num_heads) * (attrs.kv_num_heads + max(attrs.q_num_heads, attrs.kv_num_heads)) + max(attrs.q_num_heads, attrs.kv_num_heads))"
+          }
+        ]
+      }
+    },
+    {
+      "name": "bench-linear_geometry_float16_float16_g0_zero",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "linear", "state_window": 3, "scale": 0.375 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float16",
+          "shape": [2, 33, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float16",
+          "shape": [2, 33, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float16",
+          "shape": [2, 33, 14],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 33, 28] },
+        "presentStateT": { "dtype": "float16", "shape": [3, 2, 2, 8, 7] }
+      },
+      "preset": "stress",
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "dim(shapes.queryT, 0) * dim(shapes.queryT, 1) * (dim(shapes.valueT, 2) / attrs.kv_num_heads) * (2 * (dim(shapes.queryT, 2) / attrs.q_num_heads) * (attrs.kv_num_heads + max(attrs.q_num_heads, attrs.kv_num_heads)) + max(attrs.q_num_heads, attrs.kv_num_heads))"
+          }
+        ]
+      }
+    },
+    {
+      "name": "bench-linear_geometry_float16_float16_g0_state",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "linear", "state_window": 3, "scale": 0.375 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float16",
+          "shape": [2, 33, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float16",
+          "shape": [2, 33, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float16",
+          "shape": [2, 33, 14],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "pastStateT": {
+          "dtype": "float16",
+          "shape": [3, 2, 2, 8, 7],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 33, 28] },
+        "presentStateT": { "dtype": "float16", "shape": [3, 2, 2, 8, 7] }
+      },
+      "preset": "stress",
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "dim(shapes.queryT, 0) * dim(shapes.queryT, 1) * (dim(shapes.valueT, 2) / attrs.kv_num_heads) * (2 * (dim(shapes.queryT, 2) / attrs.q_num_heads) * (attrs.kv_num_heads + max(attrs.q_num_heads, attrs.kv_num_heads)) + max(attrs.q_num_heads, attrs.kv_num_heads))"
+          }
+        ]
+      }
+    },
+    {
+      "name": "bench-linear_geometry_float16_float16_g1_zero",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 2, "kv_num_heads": 4, "update_rule": "linear", "state_window": 0, "scale": 0 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float16",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float16",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float16",
+          "shape": [1, 17, 20],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [1, 17, 20] },
+        "presentStateT": { "dtype": "float16", "shape": [1, 4, 6, 5] }
+      },
+      "preset": "stress",
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "dim(shapes.queryT, 0) * dim(shapes.queryT, 1) * (dim(shapes.valueT, 2) / attrs.kv_num_heads) * (2 * (dim(shapes.queryT, 2) / attrs.q_num_heads) * (attrs.kv_num_heads + max(attrs.q_num_heads, attrs.kv_num_heads)) + max(attrs.q_num_heads, attrs.kv_num_heads))"
+          }
+        ]
+      }
+    },
+    {
+      "name": "bench-linear_geometry_float16_float16_g1_state",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 2, "kv_num_heads": 4, "update_rule": "linear", "state_window": 0, "scale": 0 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float16",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float16",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float16",
+          "shape": [1, 17, 20],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "pastStateT": {
+          "dtype": "float16",
+          "shape": [1, 4, 6, 5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [1, 17, 20] },
+        "presentStateT": { "dtype": "float16", "shape": [1, 4, 6, 5] }
+      },
+      "preset": "stress",
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "dim(shapes.queryT, 0) * dim(shapes.queryT, 1) * (dim(shapes.valueT, 2) / attrs.kv_num_heads) * (2 * (dim(shapes.queryT, 2) / attrs.q_num_heads) * (attrs.kv_num_heads + max(attrs.q_num_heads, attrs.kv_num_heads)) + max(attrs.q_num_heads, attrs.kv_num_heads))"
+          }
+        ]
+      }
+    },
+    {
+      "name": "bench-linear_geometry_float16_float32_g0_zero",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "linear", "state_window": 3, "scale": 0.375 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float16",
+          "shape": [2, 33, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float16",
+          "shape": [2, 33, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float16",
+          "shape": [2, 33, 14],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 33, 28] },
+        "presentStateT": { "dtype": "float32", "shape": [3, 2, 2, 8, 7] }
+      },
+      "preset": "stress",
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "dim(shapes.queryT, 0) * dim(shapes.queryT, 1) * (dim(shapes.valueT, 2) / attrs.kv_num_heads) * (2 * (dim(shapes.queryT, 2) / attrs.q_num_heads) * (attrs.kv_num_heads + max(attrs.q_num_heads, attrs.kv_num_heads)) + max(attrs.q_num_heads, attrs.kv_num_heads))"
+          }
+        ]
+      }
+    },
+    {
+      "name": "bench-linear_geometry_float16_float32_g0_state",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "linear", "state_window": 3, "scale": 0.375 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float16",
+          "shape": [2, 33, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float16",
+          "shape": [2, 33, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float16",
+          "shape": [2, 33, 14],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [3, 2, 2, 8, 7],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 33, 28] },
+        "presentStateT": { "dtype": "float32", "shape": [3, 2, 2, 8, 7] }
+      },
+      "preset": "stress",
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "dim(shapes.queryT, 0) * dim(shapes.queryT, 1) * (dim(shapes.valueT, 2) / attrs.kv_num_heads) * (2 * (dim(shapes.queryT, 2) / attrs.q_num_heads) * (attrs.kv_num_heads + max(attrs.q_num_heads, attrs.kv_num_heads)) + max(attrs.q_num_heads, attrs.kv_num_heads))"
+          }
+        ]
+      }
+    },
+    {
+      "name": "bench-linear_geometry_float16_float32_g1_zero",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 2, "kv_num_heads": 4, "update_rule": "linear", "state_window": 0, "scale": 0 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float16",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float16",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float16",
+          "shape": [1, 17, 20],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [1, 17, 20] },
+        "presentStateT": { "dtype": "float32", "shape": [1, 4, 6, 5] }
+      },
+      "preset": "stress",
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "dim(shapes.queryT, 0) * dim(shapes.queryT, 1) * (dim(shapes.valueT, 2) / attrs.kv_num_heads) * (2 * (dim(shapes.queryT, 2) / attrs.q_num_heads) * (attrs.kv_num_heads + max(attrs.q_num_heads, attrs.kv_num_heads)) + max(attrs.q_num_heads, attrs.kv_num_heads))"
+          }
+        ]
+      }
+    },
+    {
+      "name": "bench-linear_geometry_float16_float32_g1_state",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 2, "kv_num_heads": 4, "update_rule": "linear", "state_window": 0, "scale": 0 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float16",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float16",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float16",
+          "shape": [1, 17, 20],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [1, 4, 6, 5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [1, 17, 20] },
+        "presentStateT": { "dtype": "float32", "shape": [1, 4, 6, 5] }
+      },
+      "preset": "stress",
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "dim(shapes.queryT, 0) * dim(shapes.queryT, 1) * (dim(shapes.valueT, 2) / attrs.kv_num_heads) * (2 * (dim(shapes.queryT, 2) / attrs.q_num_heads) * (attrs.kv_num_heads + max(attrs.q_num_heads, attrs.kv_num_heads)) + max(attrs.q_num_heads, attrs.kv_num_heads))"
+          }
+        ]
+      }
+    },
+    {
+      "name": "bench-linear_geometry_float32_float16_g0_zero",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "linear", "state_window": 3, "scale": 0.375 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float32",
+          "shape": [2, 33, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float32",
+          "shape": [2, 33, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [2, 33, 14],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 33, 28] },
+        "presentStateT": { "dtype": "float16", "shape": [3, 2, 2, 8, 7] }
+      },
+      "preset": "stress",
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "dim(shapes.queryT, 0) * dim(shapes.queryT, 1) * (dim(shapes.valueT, 2) / attrs.kv_num_heads) * (2 * (dim(shapes.queryT, 2) / attrs.q_num_heads) * (attrs.kv_num_heads + max(attrs.q_num_heads, attrs.kv_num_heads)) + max(attrs.q_num_heads, attrs.kv_num_heads))"
+          }
+        ]
+      }
+    },
+    {
+      "name": "bench-linear_geometry_float32_float16_g0_state",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "linear", "state_window": 3, "scale": 0.375 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float32",
+          "shape": [2, 33, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float32",
+          "shape": [2, 33, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [2, 33, 14],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "pastStateT": {
+          "dtype": "float16",
+          "shape": [3, 2, 2, 8, 7],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 33, 28] },
+        "presentStateT": { "dtype": "float16", "shape": [3, 2, 2, 8, 7] }
+      },
+      "preset": "stress",
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "dim(shapes.queryT, 0) * dim(shapes.queryT, 1) * (dim(shapes.valueT, 2) / attrs.kv_num_heads) * (2 * (dim(shapes.queryT, 2) / attrs.q_num_heads) * (attrs.kv_num_heads + max(attrs.q_num_heads, attrs.kv_num_heads)) + max(attrs.q_num_heads, attrs.kv_num_heads))"
+          }
+        ]
+      }
+    },
+    {
+      "name": "bench-linear_geometry_float32_float16_g1_zero",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 2, "kv_num_heads": 4, "update_rule": "linear", "state_window": 0, "scale": 0 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float32",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float32",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [1, 17, 20],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 17, 20] },
+        "presentStateT": { "dtype": "float16", "shape": [1, 4, 6, 5] }
+      },
+      "preset": "stress",
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "dim(shapes.queryT, 0) * dim(shapes.queryT, 1) * (dim(shapes.valueT, 2) / attrs.kv_num_heads) * (2 * (dim(shapes.queryT, 2) / attrs.q_num_heads) * (attrs.kv_num_heads + max(attrs.q_num_heads, attrs.kv_num_heads)) + max(attrs.q_num_heads, attrs.kv_num_heads))"
+          }
+        ]
+      }
+    },
+    {
+      "name": "bench-linear_geometry_float32_float16_g1_state",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 2, "kv_num_heads": 4, "update_rule": "linear", "state_window": 0, "scale": 0 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float32",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float32",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [1, 17, 20],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "pastStateT": {
+          "dtype": "float16",
+          "shape": [1, 4, 6, 5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 17, 20] },
+        "presentStateT": { "dtype": "float16", "shape": [1, 4, 6, 5] }
+      },
+      "preset": "stress",
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "dim(shapes.queryT, 0) * dim(shapes.queryT, 1) * (dim(shapes.valueT, 2) / attrs.kv_num_heads) * (2 * (dim(shapes.queryT, 2) / attrs.q_num_heads) * (attrs.kv_num_heads + max(attrs.q_num_heads, attrs.kv_num_heads)) + max(attrs.q_num_heads, attrs.kv_num_heads))"
+          }
+        ]
+      }
     }
   ]
 }

build/webgpu/chunk-out.wgsl.jinja CHANGED Viewed

@@ -89,11 +89,10 @@ var<workgroup> pmrow: array<f32, TOKEN_ROWS * CHUNK>;
 @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>,
 ) {
   let tid = lid.x;
-  let flat = wg.x + wg.y * nwg.x;
 {{ emit_chunk_head_setup(needHeads=true) }}
   let out_heads = max(params.qNumHeads, params.kvNumHeads);
   let chunk = flat % num_chunks;

 @compute @workgroup_size(WG, 1, 1)
 fn main(
   @builtin(workgroup_id) wg: vec3<u32>,
   @builtin(local_invocation_id) lid: vec3<u32>,
 ) {
   let tid = lid.x;
+  let flat = wg.x + wg.y * {{ DISPATCH_FOLD_WIDTH }}u;
 {{ emit_chunk_head_setup(needHeads=true) }}
   let out_heads = max(params.qNumHeads, params.kvNumHeads);
   let chunk = flat % num_chunks;

build/webgpu/chunk-prep.wgsl.jinja CHANGED Viewed

@@ -18,12 +18,11 @@ const WG: u32 = {{ workgroupSize }}u;
 @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 flat (batch * chunk) index: wg.y carries the high bits past the
-  // maxComputeWorkgroupsPerDimension dispatch limit.
-  let flat = wg.x + wg.y * nwg.x;
   let num_chunks = (params.seqLength + CHUNK - 1u) / CHUNK;
   let chunk = flat % num_chunks;
   let batch = flat / num_chunks;

 @compute @workgroup_size(WG, 1, 1)
 fn main(
   @builtin(workgroup_id) wg: vec3<u32>,
   @builtin(local_invocation_id) lid: vec3<u32>,
 ) {
   // 2D-folded flat (batch * chunk) index: wg.y carries the high bits past the
+  // per-axis dispatch fold width.
+  let flat = wg.x + wg.y * {{ DISPATCH_FOLD_WIDTH }}u;
   let num_chunks = (params.seqLength + CHUNK - 1u) / CHUNK;
   let chunk = flat % num_chunks;
   let batch = flat / num_chunks;

build/webgpu/chunk-scan.wgsl.jinja CHANGED Viewed

@@ -94,13 +94,12 @@ var<workgroup> stage: array<f32, TOKEN_TILE * HEAD_DIM_K>;
 @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>,
 ) {
   let tid = lid.x;
   let col = tid % TILE_V;
   let group = tid / TILE_V;
-  let flat = wg.x + wg.y * nwg.x;
 {{ emit_chunk_head_setup() }}
   let v_tiles = HEAD_DIM_V / TILE_V;
   let v_tile = flat % v_tiles;

 @compute @workgroup_size(WG, 1, 1)
 fn main(
   @builtin(workgroup_id) wg: vec3<u32>,
   @builtin(local_invocation_id) lid: vec3<u32>,
 ) {
   let tid = lid.x;
   let col = tid % TILE_V;
   let group = tid / TILE_V;
+  let flat = wg.x + wg.y * {{ DISPATCH_FOLD_WIDTH }}u;
 {{ emit_chunk_head_setup() }}
   let v_tiles = HEAD_DIM_V / TILE_V;
   let v_tile = flat % v_tiles;

build/webgpu/chunk-ut.wgsl.jinja CHANGED Viewed

@@ -89,11 +89,10 @@ var<workgroup> tile_til: array<f32, CHUNK * TK>;
 @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>,
 ) {
   let tid = lid.x;
-  let flat = wg.x + wg.y * nwg.x;
 {{ emit_chunk_head_setup() }}
   let chunk = flat % num_chunks;
   let head = (flat / num_chunks) % params.kvNumHeads;

 @compute @workgroup_size(WG, 1, 1)
 fn main(
   @builtin(workgroup_id) wg: vec3<u32>,
   @builtin(local_invocation_id) lid: vec3<u32>,
 ) {
   let tid = lid.x;
+  let flat = wg.x + wg.y * {{ DISPATCH_FOLD_WIDTH }}u;
 {{ emit_chunk_head_setup() }}
   let chunk = flat % num_chunks;
   let head = (flat / num_chunks) % params.kvNumHeads;

build/webgpu/linear-attention.scalar.wgsl.jinja CHANGED Viewed

@@ -21,9 +21,9 @@ f16({{ expr }}){% else %}{{ expr }}{% endif %}
   let scale = select(inverseSqrt(f32(head_dim_k)), params.scale, params.scale != 0.0);
   // 2D-folded flat (batch*head*dv_tile) index: wg.y carries the high bits past
-  // the maxComputeWorkgroupsPerDimension dispatch limit. Reduces to wg.x when nwg.y == 1; the batch_idx >=
   // params.batchSize guard drops the over-dispatched tail.
-  let workgroup_idx = wg.x + wg.y * nwg.x;
   let dv_tile_idx = workgroup_idx % dv_tiles;
   let bh = workgroup_idx / dv_tiles;
   let head_idx = bh % params.kvNumHeads;
@@ -128,7 +128,6 @@ var<workgroup> broadcast_delta: array<f32, TILE_V>;
 @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>{% if useSubgroups %},
   @builtin(subgroup_invocation_id) sg_lid: u32,
   @builtin(subgroup_size) sg_size: u32{% endif %}

   let scale = select(inverseSqrt(f32(head_dim_k)), params.scale, params.scale != 0.0);
   // 2D-folded flat (batch*head*dv_tile) index: wg.y carries the high bits past
+  // the per-axis dispatch fold width. Reduces to wg.x with no fold; the batch_idx >=
   // params.batchSize guard drops the over-dispatched tail.
+  let workgroup_idx = wg.x + wg.y * {{ DISPATCH_FOLD_WIDTH }}u;
   let dv_tile_idx = workgroup_idx % dv_tiles;
   let bh = workgroup_idx / dv_tiles;
   let head_idx = bh % params.kvNumHeads;
 @compute @workgroup_size(WG, 1, 1)
 fn main(
   @builtin(workgroup_id) wg: vec3<u32>,
   @builtin(local_invocation_id) lid: vec3<u32>{% if useSubgroups %},
   @builtin(subgroup_invocation_id) sg_lid: u32,
   @builtin(subgroup_size) sg_size: u32{% endif %}

build/webgpu/linear-attention.serial.wgsl.jinja CHANGED Viewed

@@ -25,18 +25,43 @@ enable f16;
 {% endif %}
 {{ env.wgsl.resourceDeclarations }}
 const HEAD_DIM_K: u32 = {{ headDimK }}u;
 // Barrier-free small-dk recurrence. One invocation owns one
 // (batch, kv-head, value-dimension) state column and keeps its complete dk
-// slice private across the sequence. This trades dk-lane parallelism for zero
-// workgroup synchronization. Selection limits this route to small dk.
 @compute @workgroup_size(1, 1, 1)
 fn main(
   @builtin(workgroup_id) wg: vec3<u32>,
-  @builtin(num_workgroups) nwg: vec3<u32>,
 ) {
-  let flat_idx = wg.x + wg.y * nwg.x;
   let head_dim_v = params.vPackedDim / params.kvNumHeads;
   let dv_idx = flat_idx % head_dim_v;
   let bh = flat_idx / head_dim_v;

 {% endif %}
 {{ env.wgsl.resourceDeclarations }}
+{% if updateRule == "linear" %}
+{% set geometry = [
+  ["batchSize", "u32", serialBatchSize],
+  ["seqLength", "u32", serialSeqLength],
+  ["qNumHeads", "u32", serialQNumHeads],
+  ["kvNumHeads", "u32", serialKvNumHeads],
+  ["qPackedDim", "u32", serialQPackedDim],
+  ["kPackedDim", "u32", serialKPackedDim],
+  ["vPackedDim", "u32", serialVPackedDim],
+  ["scale", "f32", serialScale]
+] %}
+{% if hasStateWindow %}
+{% set geometry = geometry + [["stateWindow", "u32", serialStateWindow], ["stateSlotStride", "u32", serialStateSlotStride]] %}
+{% endif %}
+struct SerialGeometry {
+{% for field in geometry %}
+  {{ field[0] }}: {{ field[1] }},
+{% endfor %}
+}
+const params = SerialGeometry(
+{% for field in geometry %}
+  {% if field[1] == "f32" %}f32({{ field[2] }}){% else %}{{ field[2] }}u{% endif %},
+{% endfor %}
+);
+{% endif %}
 const HEAD_DIM_K: u32 = {{ headDimK }}u;
 // Barrier-free small-dk recurrence. One invocation owns one
 // (batch, kv-head, value-dimension) state column and keeps its complete dk
+// slice private across the sequence without workgroup synchronization. Small dk
+// bounds the private state retained by each invocation.
 @compute @workgroup_size(1, 1, 1)
 fn main(
   @builtin(workgroup_id) wg: vec3<u32>,
 ) {
+  let flat_idx = wg.x + wg.y * {{ DISPATCH_FOLD_WIDTH }}u;
   let head_dim_v = params.vPackedDim / params.kvNumHeads;
   let dv_idx = flat_idx % head_dim_v;
   let bh = flat_idx / head_dim_v;

build/webgpu/linear-attention.vec4.wgsl.jinja CHANGED Viewed

@@ -21,9 +21,9 @@ f16({{ expr }}){% else %}{{ expr }}{% endif %}
   let scale = select(inverseSqrt(f32(head_dim_k)), params.scale, params.scale != 0.0);
   // 2D-folded flat (batch*head*dv_tile) index: wg.y carries the high bits past
-  // the maxComputeWorkgroupsPerDimension dispatch limit. Reduces to wg.x when nwg.y == 1; the batch_idx >=
   // params.batchSize guard drops the over-dispatched tail.
-  let workgroup_idx = wg.x + wg.y * nwg.x;
   let dv_tile_idx = workgroup_idx % dv_tiles;
   let bh = workgroup_idx / dv_tiles;
   let head_idx = bh % params.kvNumHeads;
@@ -97,6 +97,38 @@ enable f16;
 enable subgroups;
 {% endif %}
 {{ env.wgsl.resourceDeclarations }}
 // LANES threads cooperate on one value tile's reduction axis; DV_GROUPS such groups
 // share a workgroup so the per-token key/query/decay stream is fetched once for all of
@@ -117,10 +149,9 @@ const VEC4_LANES: u32 = 4u;
 // No-subgroup tier: shared-memory linear folds replace subgroupAdd. Every quantity a
 // token reduces is staged before one barrier and read back after it, so the barrier
 // count is a property of the token and not of how many quantities it reduces.
-// A lane's TILE_V components are adjacent, so they travel as one {{ redT }} word and
-// fold with {{ redWidth }}-wide adds -- the same per-component summation order as a
-// scalar fold, at a {{ redWidth }}th of the shared-memory transactions.
-// Folds stay within the caller's own lane group: groups own disjoint value tiles and
 // must not see each other's partials.
 var<workgroup> wg_fold: array<{{ redT }}, WG * {{ redSlots }}u>;
 var<workgroup> wg_fold_out: array<{{ redT }}, DV_GROUPS * {{ redSlots }}u>;
@@ -128,16 +159,26 @@ var<workgroup> wg_fold_out: array<{{ redT }}, DV_GROUPS * {{ redSlots }}u>;
 @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>,
 ) {
   let tid = lid.x;
   let lane = tid % LANES;
-{% if dvGroups > 1 or not useSubgroups %}
   let dv_group = tid / LANES;
 {% endif %}
-  let dk_base = lane * VEC4_LANES;
-{{ emit_tiled_setup(dvGroups=dvGroups) }}
   let lane_active = dk_base < head_dim_k;
   // state[j] holds 4 consecutive dk rows (the 4 components) for dv slot j.
@@ -183,10 +224,9 @@ fn main(
     }
   }
 {% endif %}
-  // A token's key and query do not depend on the recurrent state, so they are fetched one
-  // iteration ahead: the fetch for t+1 is issued before the reductions for t, and its memory
-  // latency overlaps the dependent chain instead of stalling in front of it. The state
-  // recurrence is what serializes this loop, and it leaves the load unit idle otherwise.
 {% macro load_key(token) %}
     if (lane_active) {
       let k_base = ({{ token }} * params.kPackedDim + key_head_idx * head_dim_k) / VEC4_LANES + lane;

   let scale = select(inverseSqrt(f32(head_dim_k)), params.scale, params.scale != 0.0);
   // 2D-folded flat (batch*head*dv_tile) index: wg.y carries the high bits past
+  // the per-axis dispatch fold width. Reduces to wg.x with no fold; the batch_idx >=
   // params.batchSize guard drops the over-dispatched tail.
+  let workgroup_idx = wg.x + wg.y * {{ DISPATCH_FOLD_WIDTH }}u;
   let dv_tile_idx = workgroup_idx % dv_tiles;
   let bh = workgroup_idx / dv_tiles;
   let head_idx = bh % params.kvNumHeads;
 enable subgroups;
 {% endif %}
 {{ env.wgsl.resourceDeclarations }}
+{% if useSubgroups %}
+{% set skipLogicalLaneCount = true %}
+{% set skipLogicalLaneCount = skipLogicalLaneCount is defined and skipLogicalLaneCount %}
+var<workgroup> sgLaneClaims: atomic<u32>;
+struct SubgroupLogicalLanes {
+  ltid: u32,
+  ord: u32,
+  rank: u32,
+  count: u32,
+}
+fn subgroup_logical_lanes() -> SubgroupLogicalLanes {
+  let rank = subgroupExclusiveAdd(1u);
+  let count = subgroupAdd(1u);
+  // Every lane performs the atomic so no collective follows a lane guard (a
+  // subgroup op after a closed `if (rank == 0u)` block is the reconvergence
+  // hazard the render gate flags): only the rank-0 lane adds its subgroup's
+  // claim, every other lane adds 0 and discards its snapshot. The rank-0 lane is
+  // the lowest active lane, which is the lane `subgroupBroadcastFirst` reads.
+  let ticket = atomicAdd(&sgLaneClaims, select(0u, count | (1u << 16u), rank == 0u));
+  let claim = subgroupBroadcastFirst(ticket);
+  return SubgroupLogicalLanes((claim & 0xffffu) + rank, claim >> 16u, rank, count);
+}
+{% if not skipLogicalLaneCount %}
+fn subgroup_logical_count() -> u32 {
+  return atomicLoad(&sgLaneClaims) >> 16u;
+}
+{% endif %}
+{% endif %}
 // LANES threads cooperate on one value tile's reduction axis; DV_GROUPS such groups
 // share a workgroup so the per-token key/query/decay stream is fetched once for all of
 // No-subgroup tier: shared-memory linear folds replace subgroupAdd. Every quantity a
 // token reduces is staged before one barrier and read back after it, so the barrier
 // count is a property of the token and not of how many quantities it reduces.
+// A lane's TILE_V components are adjacent, so they are stored as one {{ redT }} word
+// and fold with {{ redWidth }}-wide adds in the scalar fold's per-component order.
+// Folds stay within each lane group: groups own disjoint value tiles and
 // must not see each other's partials.
 var<workgroup> wg_fold: array<{{ redT }}, WG * {{ redSlots }}u>;
 var<workgroup> wg_fold_out: array<{{ redT }}, DV_GROUPS * {{ redSlots }}u>;
 @compute @workgroup_size(WG, 1, 1)
 fn main(
   @builtin(workgroup_id) wg: vec3<u32>,
+{% if not useSubgroups %}
   @builtin(local_invocation_id) lid: vec3<u32>,
+{% endif %}
 ) {
+{% if useSubgroups %}
+  // Subgroup tier: selection admits this arm only on adapters whose fixed subgroup
+  // width equals LANES, so every subgroup of the workgroup is exactly one lane group
+  // and subgroupAdd over it is the group's reduction. The group and the lane come from
+  // the subgroup itself -- the dense claim order of the subgroup and the invocation's
+  // dense rank among its active lanes -- not from a partition of local_invocation_id,
+  // because WGSL does not promise which invocations share a subgroup.
+  let L = subgroup_logical_lanes();
+  let lane = L.rank;
+  let dv_group = L.ord;
+{% else %}
   let tid = lid.x;
   let lane = tid % LANES;
   let dv_group = tid / LANES;
 {% endif %}
+  let dk_base = lane * VEC4_LANES;{{ emit_tiled_setup(dvGroups=(2 if useSubgroups else dvGroups)) }}
   let lane_active = dk_base < head_dim_k;
   // state[j] holds 4 consecutive dk rows (the 4 components) for dv slot j.
     }
   }
 {% endif %}
+  // A token's key and query do not depend on the recurrent state. Fetch t+1
+  // before completing the reductions for t so the independent load is issued
+  // ahead of its first use.
 {% macro load_key(token) %}
     if (lane_active) {
       let k_base = ({{ token }} * params.kPackedDim + key_head_idx * head_dim_k) / VEC4_LANES + lane;

build/webgpu/manifest.json CHANGED Viewed

The diff for this file is too large to render. See raw diff

build/webgpu/metadata.json CHANGED Viewed

@@ -1,24 +1,56 @@
 {
   "name": "com.microsoft.LinearAttention",
-  "id": "_com_microsoft_linearattention_webgpu_dc21501",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "UnbCbJjzzX17ihI6vVaaXbbgxyNRzghqQ9jFtny3sZs=",
-      "chunk-out.wgsl.jinja": "n2BVpE+dJE8bWsb3Vv9EZalVcE6Bpr0+RYnUxCliSSQ=",
-      "chunk-prep.wgsl.jinja": "U9V7Cld2Zou0sX77ohI+0hmq+UZUsMQohiWrCLFrFm0=",
-      "chunk-scan.wgsl.jinja": "4yrz7CYz19GJAxJ5cBmc4XJt6aXa1ICUi+fg4ma2+mM=",
-      "chunk-ut.wgsl.jinja": "9KqJJCt5PwygvCrBjjy8oglG1ly4HiumH2Im/4Zlz5M=",
-      "linear-attention.scalar.wgsl.jinja": "mc7lY/Xgl1LkvRxiCeaSZ2DwldVlIeoTti2QFqhvCiI=",
-      "linear-attention.serial.wgsl.jinja": "p2AF3fRzazYsMPvGijvV3UyJhxvVxdZJyOhCYLRCJgI=",
-      "linear-attention.vec4.wgsl.jinja": "heG1tNzKu8g1M6PqKtDFED/CzGMDA/qWvlTuZ4Z557M=",
-      "manifest.json": "75pbsCWWrC10tm4NEaVyyoM7nDVgCqxwvj8Bh568M2g=",
-      "test.json": "NtZMMzB3k6UYTuxbvHod7FJlq+HqMV0oYuZi/qhW8sM="
     }
   },
-  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
-  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/com.microsoft.LinearAttention" }
 }

 {
   "name": "com.microsoft.LinearAttention",
+  "id": "_com_microsoft_linearattention_webgpu_13a2aaf",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "SADv8TDNnXFGZdNUCqGMV5D9W9IKy2+gSRRNOZvxh/4=",
+      "chunk-out.wgsl.jinja": "VyrRYpaK7TsAfpKBMZpBiusQfnbDGhS5QIvnA/Vk6Tg=",
+      "chunk-prep.wgsl.jinja": "Iq5ZLzNK/CtE7/i7Tm4pZi5rxMKq6dZ4wEpBqrapUdQ=",
+      "chunk-scan.wgsl.jinja": "aPOMJxL7BhXVa05PwX9Uxl86dRbnuc9DkEwoFZq1Vvo=",
+      "chunk-ut.wgsl.jinja": "MfZk82OOr4oJCJWi4vU/54BLerVSyS31uRsWeGDakw0=",
+      "linear-attention.scalar.wgsl.jinja": "O8NRGtK2wpemLnSk96xiGRl50mqbqHQ3TiEkZpCGlOc=",
+      "linear-attention.serial.wgsl.jinja": "gWOSDZOYAlaSi4wPdGqEq1UWfGxpVDxYTFeUdMDp6lw=",
+      "linear-attention.vec4.wgsl.jinja": "bUlUYYjiumu6uCY/V7+EnbuhmBljoZXUnskKNYS0Hac=",
+      "manifest.json": "7SA3QNshHZfuljHXCA9dVbCwtV4wiiH+4EO78NGq9mU=",
+      "test.json": "NQ2VDXoy+rsZMH/agLHNEpOXHOydLQq+C2veCcWdpyk="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": {
+      "linear_zero_chunked": ["chunk-out.wgsl.jinja", "chunk-scan.wgsl.jinja"],
+      "linear_state_chunked": ["chunk-out.wgsl.jinja", "chunk-scan.wgsl.jinja"],
+      "gated_zero_chunked": ["chunk-out.wgsl.jinja", "chunk-prep.wgsl.jinja", "chunk-scan.wgsl.jinja"],
+      "gated_state_chunked": ["chunk-out.wgsl.jinja", "chunk-prep.wgsl.jinja", "chunk-scan.wgsl.jinja"],
+      "delta_zero_chunked": ["chunk-out.wgsl.jinja", "chunk-scan.wgsl.jinja", "chunk-ut.wgsl.jinja"],
+      "delta_state_chunked": ["chunk-out.wgsl.jinja", "chunk-scan.wgsl.jinja", "chunk-ut.wgsl.jinja"],
+      "gated_delta_zero_chunked": ["chunk-out.wgsl.jinja", "chunk-prep.wgsl.jinja", "chunk-scan.wgsl.jinja", "chunk-ut.wgsl.jinja"],
+      "gated_delta_state_chunked": ["chunk-out.wgsl.jinja", "chunk-prep.wgsl.jinja", "chunk-scan.wgsl.jinja", "chunk-ut.wgsl.jinja"],
+      "linear_zero_serial_small_dk": ["linear-attention.serial.wgsl.jinja"],
+      "linear_state_serial_small_dk": ["linear-attention.serial.wgsl.jinja"],
+      "gated_delta_zero_serial_small_dk": ["linear-attention.serial.wgsl.jinja"],
+      "gated_delta_state_serial_small_dk": ["linear-attention.serial.wgsl.jinja"],
+      "linear_zero_scalar": ["linear-attention.scalar.wgsl.jinja"],
+      "linear_state_scalar": ["linear-attention.scalar.wgsl.jinja"],
+      "gated_delta_zero_scalar": ["linear-attention.scalar.wgsl.jinja"],
+      "gated_delta_state_scalar": ["linear-attention.scalar.wgsl.jinja"],
+      "gated_zero_scalar": ["linear-attention.scalar.wgsl.jinja"],
+      "gated_state_scalar": ["linear-attention.scalar.wgsl.jinja"],
+      "delta_zero_scalar": ["linear-attention.scalar.wgsl.jinja"],
+      "delta_state_scalar": ["linear-attention.scalar.wgsl.jinja"],
+      "linear_zero_vec4": ["linear-attention.vec4.wgsl.jinja"],
+      "linear_state_vec4": ["linear-attention.vec4.wgsl.jinja"],
+      "gated_delta_zero_vec4": ["linear-attention.vec4.wgsl.jinja"],
+      "gated_delta_state_vec4": ["linear-attention.vec4.wgsl.jinja"],
+      "gated_zero_vec4": ["linear-attention.vec4.wgsl.jinja"],
+      "gated_state_vec4": ["linear-attention.vec4.wgsl.jinja"],
+      "delta_zero_vec4": ["linear-attention.vec4.wgsl.jinja"],
+      "delta_state_vec4": ["linear-attention.vec4.wgsl.jinja"]
+    }
+  }
 }

build/webgpu/test.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "com.microsoft.LinearAttention",
   "cases": [
     {
       "name": "linear_zero_state",
@@ -328,7 +327,7 @@
     {
       "name": "gated_delta_scalar_headdim6_seq128_state",
       "provenance": {
-        "notes": "Compact sibling for the long-sequence scalar gated-delta benchmark; preserves qHeads=kvHeads=4, headDimK=6 (non-vec4), past state, per-head decay/beta, and a multi-token recurrence loop."
       },
       "attrs": { "q_num_heads": 4, "kv_num_heads": 4, "update_rule": "gated_delta", "scale": 0.25 },
       "inputs": {
@@ -364,8 +363,8 @@
         }
       },
       "outputs": {
-        "outputT": { "dtype": "float32", "shape": [1, 128, 48], "tolerance": 0.0005, "relTolerance": 0.0005 },
-        "presentStateT": { "dtype": "float32", "shape": [1, 4, 6, 12], "tolerance": 0.0005, "relTolerance": 0.0005 }
       }
     },
     {
@@ -534,7 +533,7 @@
     {
       "name": "linear_zero_scalar_f16_seq128",
       "provenance": {
-        "notes": "Compact sibling for the f16 linear-rule scalar benchmark; f16 query/state force the scalar implementation while seq=128 keeps the zero-state recurrence loop model-shaped."
       },
       "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "linear", "scale": 0.25 },
       "inputs": {
@@ -555,14 +554,14 @@
         }
       },
       "outputs": {
-        "outputT": { "dtype": "float16", "shape": [1, 128, 64], "tolerance": 0.04, "relTolerance": 0.04 },
-        "presentStateT": { "dtype": "float16", "shape": [1, 2, 16, 16], "tolerance": 0.04, "relTolerance": 0.04 }
       }
     },
     {
       "name": "linear_zero_f16_seq128_offset_value_scale_lock",
       "provenance": {
-        "notes": "Scale lock for the zero-past-state linear recurrence over 128 tokens (the \"zero\" in linear_zero names the absent past_state, not the output). Its sibling linear_zero_scalar_f16_seq128 drives q/k/v at amplitude 0.03 about zero, so the state is a 128-term random walk of zero-mean outer products and the output peaks at 1.7e-3 against a 0.04 absolute tolerance - blind to any multiplicative error below 24x, and 13.5x on present_state. Here the key carries a small positive DC and V oscillates about 0.5, so the accumulated state peaks at 1.5 and the output at 2.0; halving either output now scores ~30x the allowed error. Same shape, args and dtypes, so it still covers linear_zero_serial_small_dk, linear_zero_scalar and linear_zero_vec4."
       },
       "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "linear", "scale": 0.25 },
       "inputs": {
@@ -590,7 +589,7 @@
     {
       "name": "linear_state_scalar_f16_seq128",
       "provenance": {
-        "notes": "Compact correctness sibling for the long supplied-state f16 pathology benchmark. It selects linear_state_serial_small_dk and validates that the initial state is incorporated rather than silently treated as zero."
       },
       "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "linear", "scale": 0.25 },
       "inputs": {
@@ -616,14 +615,14 @@
         }
       },
       "outputs": {
-        "outputT": { "dtype": "float16", "shape": [1, 128, 64], "tolerance": 0.04, "relTolerance": 0.04 },
-        "presentStateT": { "dtype": "float16", "shape": [1, 2, 16, 16], "tolerance": 0.04, "relTolerance": 0.04 }
       }
     },
     {
       "name": "linear_state_f16_seq128_offset_value_scale_lock",
       "provenance": {
-        "notes": "Supplied-past-state arm of the same lock. linear_state_scalar_f16_seq128 claims to validate that the initial state is incorporated rather than treated as zero, but at amplitude 0.03 its output peaks at 4.4e-3 against a 0.04 absolute tolerance (9.1x blind) and the past state itself is scaled 0.01, so zeroing past_state entirely stays well inside the tolerance. With the past state at amplitude 0.3, a keyed DC and V about 0.5, the output peaks at 1.3 and the state at 1.5: dropping past_state now scores 11x the allowed error on the output and 35x on present_state, and halving either scores ~28x."
       },
       "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "linear", "scale": 0.25 },
       "inputs": {
@@ -804,7 +803,7 @@
     {
       "name": "gated_delta_default_rule_no_updateRule_arg",
       "provenance": {
-        "notes": "updateRule OMITTED -> manifest derive.effRule defaults to gated_delta (host and reference agree). Requires decay+beta+past_state. headDimK=4 (%4==0) so vec4 path. Verifies the default-rule contract across manifest/host/reference."
       },
       "attrs": { "q_num_heads": 2, "kv_num_heads": 1 },
       "inputs": {
@@ -847,7 +846,7 @@
     {
       "name": "gated_delta_f32_dk128_dv128_compact",
       "provenance": {
-        "notes": "Compact correctness lock for the Bonsai-shaped dK=dV=128 gated-delta benchmark. qHeads=4, kvHeads=2 and seq=2 bound CPU reference work while preserving GQA state sharing and the f32 vec4 path."
       },
       "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "gated_delta", "scale": 0.08838834764831845 },
       "inputs": {
@@ -883,8 +882,13 @@
         }
       },
       "outputs": {
-        "outputT": { "dtype": "float32", "shape": [1, 2, 512], "tolerance": 0.0008, "relTolerance": 0.0008 },
-        "presentStateT": { "dtype": "float32", "shape": [1, 2, 128, 128], "tolerance": 0.0008, "relTolerance": 0.0008 }
       }
     },
     {
@@ -933,7 +937,7 @@
     {
       "name": "gated_delta_f16_dk128_dv128_compact",
       "provenance": {
-        "notes": "Float16 sibling of the dK=dV=128 gated-delta lock. It preserves the model-shaped head dimensions and scalar f16 route while remaining small enough for CPU reference validation."
       },
       "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "gated_delta", "scale": 0.08838834764831845 },
       "inputs": {
@@ -969,8 +973,8 @@
         }
       },
       "outputs": {
-        "outputT": { "dtype": "float16", "shape": [1, 2, 512], "tolerance": 0.08, "relTolerance": 0.08 },
-        "presentStateT": { "dtype": "float16", "shape": [1, 2, 128, 128], "tolerance": 0.08, "relTolerance": 0.08 }
       }
     },
     {
@@ -1097,7 +1101,7 @@
       "provenance": {
         "source": "onnxruntime/contrib_ops/webgpu/bert/linear_attention.cc",
         "test": "head-size-driven workgroup sizing",
-        "notes": "One vec4 lane per four dk rows, so head_dim_k = 2048 needs 512 lanes. The width was pinned to 256 and the head dim to 1024 to match; both are now taken from the device, which is what the second half of the old guard already did. A device that cannot run 512 invocations per workgroup still cannot serve this shape, so the case declares that limit rather than claiming to be universal. Query and key are positive cycles of length 5 and 6 rather than sinusoids: a signed 2048-term dot product cancels down to ~1e-4, where no tolerance the f32 summation order allows can see a scale error. Both cycle lengths are coprime with the four-lane vec4 stride, so every lane still sees every value."
       },
       "attrs": { "q_num_heads": 1, "kv_num_heads": 1, "update_rule": "linear" },
       "inputs": {
@@ -1153,7 +1157,7 @@
     {
       "name": "linear_state_scalar_dk17_above_serial_cap_unaligned",
       "provenance": {
-        "notes": "Supplied-state sibling of the dk=17 scalar lock: headDimK=17 defeats serialHeadDimFits (dk<=16) and the %4 vec4 guard, so the scalar kernel runs with a past state and must fold it into the recurrence instead of starting from zeros."
       },
       "attrs": { "q_num_heads": 1, "kv_num_heads": 1, "update_rule": "linear", "scale": 0.5 },
       "inputs": {
@@ -1186,7 +1190,7 @@
     {
       "name": "linear_state_vec4_dk20_above_serial_cap",
       "provenance": {
-        "notes": "headDimK=20 is %4 (vec4-eligible) but above serialHeadDimFits (dk<=16), so the supplied-state vec4 kernel wins instead of linear_state_serial_small_dk. WG=pow2ceil(ceil(20/4))=8 leaves 3 idle lanes, and head_dim_v=3 exercises a partial dv tile on both output and present_state."
       },
       "attrs": { "q_num_heads": 1, "kv_num_heads": 1, "update_rule": "linear", "scale": 0.5 },
       "inputs": {
@@ -1298,7 +1302,7 @@
     {
       "name": "linear_state_scalar_dk17_window3",
       "provenance": {
-        "notes": "Windowed twin of linear_state_scalar_dk17_above_serial_cap_unaligned, which is the case that selects this variant. The pinned serial window pair anchors the reference; this one checks that the same variant's per-token snapshot and slot indexing agree with it."
       },
       "attrs": { "q_num_heads": 1, "kv_num_heads": 1, "update_rule": "linear", "scale": 0.5, "state_window": 3 },
       "inputs": {
@@ -1331,7 +1335,7 @@
     {
       "name": "linear_state_vec4_dk20_window2",
       "provenance": {
-        "notes": "Windowed twin of linear_state_vec4_dk20_above_serial_cap, which is the case that selects this variant. The pinned serial window pair anchors the reference; this one checks that the same variant's per-token snapshot and slot indexing agree with it."
       },
       "attrs": { "q_num_heads": 1, "kv_num_heads": 1, "update_rule": "linear", "scale": 0.5, "state_window": 2 },
       "inputs": {
@@ -1393,7 +1397,7 @@
     {
       "name": "linear_zero_vec4_window2",
       "provenance": {
-        "notes": "Windowed zero-state coverage for the linear vec4 route, including per-token present-state slots."
       },
       "attrs": { "q_num_heads": 1, "kv_num_heads": 1, "update_rule": "linear", "scale": 0.5, "state_window": 2 },
       "inputs": {
@@ -1420,7 +1424,9 @@
     },
     {
       "name": "gated_zero_window2_all_routes",
-      "provenance": { "notes": "Windowed zero-state coverage for every eligible gated serial, scalar, and vec4 route." },
       "attrs": { "q_num_heads": 1, "kv_num_heads": 1, "update_rule": "gated", "scale": 0.5, "state_window": 2 },
       "inputs": {
         "queryT": {
@@ -1448,7 +1454,7 @@
     {
       "name": "gated_state_window2_all_routes",
       "provenance": {
-        "notes": "Windowed carried-state coverage for every eligible gated serial, scalar, and vec4 route."
       },
       "attrs": { "q_num_heads": 1, "kv_num_heads": 1, "update_rule": "gated", "scale": 0.5, "state_window": 2 },
       "inputs": {
@@ -1481,7 +1487,9 @@
     },
     {
       "name": "delta_zero_window2_all_routes",
-      "provenance": { "notes": "Windowed zero-state coverage for every eligible delta scalar and vec4 route." },
       "attrs": { "q_num_heads": 1, "kv_num_heads": 1, "update_rule": "delta", "scale": 0.5, "state_window": 2 },
       "inputs": {
         "queryT": {
@@ -1508,7 +1516,9 @@
     },
     {
       "name": "delta_state_window2_all_routes",
-      "provenance": { "notes": "Windowed carried-state coverage for every eligible delta scalar and vec4 route." },
       "attrs": { "q_num_heads": 1, "kv_num_heads": 1, "update_rule": "delta", "scale": 0.5, "state_window": 2 },
       "inputs": {
         "queryT": {
@@ -1541,7 +1551,7 @@
     {
       "name": "gated_delta_zero_window2_all_routes",
       "provenance": {
-        "notes": "Windowed zero-state coverage for every eligible gated-delta serial, scalar, and vec4 route."
       },
       "attrs": { "q_num_heads": 1, "kv_num_heads": 1, "update_rule": "gated_delta", "scale": 0.5, "state_window": 2 },
       "inputs": {
@@ -1571,7 +1581,7 @@
     {
       "name": "gated_delta_state_window2_all_routes",
       "provenance": {
-        "notes": "Windowed carried-state coverage for every eligible gated-delta serial, scalar, and vec4 route."
       },
       "attrs": { "q_num_heads": 1, "kv_num_heads": 1, "update_rule": "gated_delta", "scale": 0.5, "state_window": 2 },
       "inputs": {
@@ -1839,7 +1849,7 @@
     {
       "name": "delta_with_initial_state",
       "provenance": {
-        "notes": "Covers the schema-valid delta recurrence with a supplied past_state; providers consume the same optional state for all four update rules."
       },
       "attrs": { "q_num_heads": 1, "kv_num_heads": 1, "update_rule": "delta", "scale": 0.5 },
       "inputs": {
@@ -1939,7 +1949,7 @@
     {
       "name": "gated_delta_zero_scalar_dk17_no_past",
       "provenance": {
-        "notes": "Omitted past_state with headDimK=17 crosses the serial-kernel ceiling and is not divisible by four, selecting the scalar zero-state recurrence rather than a state-bearing or vec4 sibling."
       },
       "attrs": { "q_num_heads": 1, "kv_num_heads": 1, "update_rule": "gated_delta", "scale": 0.25 },
       "inputs": {
@@ -2293,7 +2303,7 @@
     {
       "name": "gated_delta_state_chunked_f16_seq1024",
       "provenance": {
-        "notes": "float16 activations and state through the chunked prefill decomposition, which needs seqLength >= 1024. The chunk passes accumulate in float32 and only their loads and stores are narrowed, so this is the fixture that exercises the f16 enable in every pass."
       },
       "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "gated_delta", "scale": 0.35 },
       "inputs": {
@@ -2332,6 +2342,802 @@
         "outputT": { "dtype": "float16", "shape": [1, 1024, 32], "tolerance": 0.04, "relTolerance": 0.04 },
         "presentStateT": { "dtype": "float16", "shape": [1, 2, 8, 8], "tolerance": 0.04, "relTolerance": 0.04 }
       }
     }
   ]
 }

 {
   "cases": [
     {
       "name": "linear_zero_state",
     {
       "name": "gated_delta_scalar_headdim6_seq128_state",
       "provenance": {
+        "notes": "A compact multi-token gated-delta recurrence uses four query/KV heads, non-four-wide key head size 6, past state, and per-head decay and beta."
       },
       "attrs": { "q_num_heads": 4, "kv_num_heads": 4, "update_rule": "gated_delta", "scale": 0.25 },
       "inputs": {
         }
       },
       "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 128, 48], "tolerance": 0.00005, "relTolerance": 0.00005 },
+        "presentStateT": { "dtype": "float32", "shape": [1, 4, 6, 12], "tolerance": 0.00005, "relTolerance": 0.00005 }
       }
     },
     {
     {
       "name": "linear_zero_scalar_f16_seq128",
       "provenance": {
+        "notes": "Float16 query and state select the scalar linear-rule implementation over a model-shaped 128-token zero-state recurrence."
       },
       "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "linear", "scale": 0.25 },
       "inputs": {
         }
       },
       "outputs": {
+        "outputT": { "dtype": "float16", "shape": [1, 128, 64], "tolerance": 0.00004, "relTolerance": 0.002 },
+        "presentStateT": { "dtype": "float16", "shape": [1, 2, 16, 16], "tolerance": 0.00004, "relTolerance": 0.002 }
       }
     },
     {
       "name": "linear_zero_f16_seq128_offset_value_scale_lock",
       "provenance": {
+        "notes": "A 128-token float16 recurrence omits past state but does not produce zero output. Positive-offset keys and values around 0.5 keep output and present state at order-one magnitude, making multiplicative errors observable on the serial, scalar, and vec4 zero-state routes."
       },
       "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "linear", "scale": 0.25 },
       "inputs": {
     {
       "name": "linear_state_scalar_f16_seq128",
       "provenance": {
+        "notes": "A compact float16 recurrence selects the serial small-key-dimension route and verifies that its supplied initial state is incorporated."
       },
       "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "linear", "scale": 0.25 },
       "inputs": {
         }
       },
       "outputs": {
+        "outputT": { "dtype": "float16", "shape": [1, 128, 64], "tolerance": 0.0002, "relTolerance": 0.004 },
+        "presentStateT": { "dtype": "float16", "shape": [1, 2, 16, 16], "tolerance": 0.0002, "relTolerance": 0.004 }
       }
     },
     {
       "name": "linear_state_f16_seq128_offset_value_scale_lock",
       "provenance": {
+        "notes": "A supplied past state at amplitude 0.3, positive-offset keys, and values around 0.5 keep both outputs at order-one magnitude. Dropping the initial state or uniformly rescaling either output therefore exceeds tolerance."
       },
       "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "linear", "scale": 0.25 },
       "inputs": {
     {
       "name": "gated_delta_default_rule_no_updateRule_arg",
       "provenance": {
+        "notes": "Omitting `updateRule` selects the schema-default gated-delta recurrence with decay, beta, and past state. Key head size 4 exercises the vec4 route."
       },
       "attrs": { "q_num_heads": 2, "kv_num_heads": 1 },
       "inputs": {
     {
       "name": "gated_delta_f32_dk128_dv128_compact",
       "provenance": {
+        "notes": "A compact gated-delta case preserves Bonsai's 128-wide key/value heads, four query heads, two KV heads, shared GQA state, and the float32 vec4 path."
       },
       "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "gated_delta", "scale": 0.08838834764831845 },
       "inputs": {
         }
       },
       "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 2, 512], "tolerance": 0.000005, "relTolerance": 0.0005 },
+        "presentStateT": {
+          "dtype": "float32",
+          "shape": [1, 2, 128, 128],
+          "tolerance": 0.000005,
+          "relTolerance": 0.0005
+        }
       }
     },
     {
     {
       "name": "gated_delta_f16_dk128_dv128_compact",
       "provenance": {
+        "notes": "Float16 tensors with 128-wide key/value heads exercise the scalar gated-delta route at a compact sequence length."
       },
       "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "gated_delta", "scale": 0.08838834764831845 },
       "inputs": {
         }
       },
       "outputs": {
+        "outputT": { "dtype": "float16", "shape": [1, 2, 512], "tolerance": 0.000015, "relTolerance": 0.003 },
+        "presentStateT": { "dtype": "float16", "shape": [1, 2, 128, 128], "tolerance": 0.0002, "relTolerance": 0.003 }
       }
     },
     {
       "provenance": {
         "source": "onnxruntime/contrib_ops/webgpu/bert/linear_attention.cc",
         "test": "head-size-driven workgroup sizing",
+        "notes": "One vec4 lane handles four key rows, so key head size 2048 requires a 512-invocation workgroup and declares that adapter limit. Positive query/key cycles of coprime lengths 5 and 6 avoid cancellation while ensuring every four-wide lane sees every value."
       },
       "attrs": { "q_num_heads": 1, "kv_num_heads": 1, "update_rule": "linear" },
       "inputs": {
     {
       "name": "linear_state_scalar_dk17_above_serial_cap_unaligned",
       "provenance": {
+        "notes": "Key head size 17 exceeds the serial limit and is not divisible by four, selecting the scalar recurrence with a supplied past state."
       },
       "attrs": { "q_num_heads": 1, "kv_num_heads": 1, "update_rule": "linear", "scale": 0.5 },
       "inputs": {
     {
       "name": "linear_state_vec4_dk20_above_serial_cap",
       "provenance": {
+        "notes": "A key head dimension of 20 uses five four-value vectors in an eight-lane workgroup, while value head dimension 3 produces a partial output and present-state tile."
       },
       "attrs": { "q_num_heads": 1, "kv_num_heads": 1, "update_rule": "linear", "scale": 0.5 },
       "inputs": {
     {
       "name": "linear_state_scalar_dk17_window3",
       "provenance": {
+        "notes": "A three-slot window on the scalar, key-head-size-17 route verifies per-token state snapshots and slot indexing with a supplied initial state."
       },
       "attrs": { "q_num_heads": 1, "kv_num_heads": 1, "update_rule": "linear", "scale": 0.5, "state_window": 3 },
       "inputs": {
     {
       "name": "linear_state_vec4_dk20_window2",
       "provenance": {
+        "notes": "A two-slot window on the vec4, key-head-size-20 route verifies per-token state snapshots and slot indexing with a supplied initial state."
       },
       "attrs": { "q_num_heads": 1, "kv_num_heads": 1, "update_rule": "linear", "scale": 0.5, "state_window": 2 },
       "inputs": {
     {
       "name": "linear_zero_vec4_window2",
       "provenance": {
+        "notes": "A zero-state linear recurrence writes per-token present-state slots on the vec4 route."
       },
       "attrs": { "q_num_heads": 1, "kv_num_heads": 1, "update_rule": "linear", "scale": 0.5, "state_window": 2 },
       "inputs": {
     },
     {
       "name": "gated_zero_window2_all_routes",
+      "provenance": {
+        "notes": "A zero-state gated recurrence exercises window publication on its serial, scalar, and vec4 routes."
+      },
       "attrs": { "q_num_heads": 1, "kv_num_heads": 1, "update_rule": "gated", "scale": 0.5, "state_window": 2 },
       "inputs": {
         "queryT": {
     {
       "name": "gated_state_window2_all_routes",
       "provenance": {
+        "notes": "A carried-state gated recurrence exercises window publication on its serial, scalar, and vec4 routes."
       },
       "attrs": { "q_num_heads": 1, "kv_num_heads": 1, "update_rule": "gated", "scale": 0.5, "state_window": 2 },
       "inputs": {
     },
     {
       "name": "delta_zero_window2_all_routes",
+      "provenance": {
+        "notes": "A zero-state delta recurrence exercises window publication on its scalar and vec4 routes."
+      },
       "attrs": { "q_num_heads": 1, "kv_num_heads": 1, "update_rule": "delta", "scale": 0.5, "state_window": 2 },
       "inputs": {
         "queryT": {
     },
     {
       "name": "delta_state_window2_all_routes",
+      "provenance": {
+        "notes": "A carried-state delta recurrence exercises window publication on its scalar and vec4 routes."
+      },
       "attrs": { "q_num_heads": 1, "kv_num_heads": 1, "update_rule": "delta", "scale": 0.5, "state_window": 2 },
       "inputs": {
         "queryT": {
     {
       "name": "gated_delta_zero_window2_all_routes",
       "provenance": {
+        "notes": "A zero-state gated-delta recurrence exercises window publication on its serial, scalar, and vec4 routes."
       },
       "attrs": { "q_num_heads": 1, "kv_num_heads": 1, "update_rule": "gated_delta", "scale": 0.5, "state_window": 2 },
       "inputs": {
     {
       "name": "gated_delta_state_window2_all_routes",
       "provenance": {
+        "notes": "A carried-state gated-delta recurrence exercises window publication on its serial, scalar, and vec4 routes."
       },
       "attrs": { "q_num_heads": 1, "kv_num_heads": 1, "update_rule": "gated_delta", "scale": 0.5, "state_window": 2 },
       "inputs": {
     {
       "name": "delta_with_initial_state",
       "provenance": {
+        "notes": "A schema-valid delta recurrence consumes the optional `past_state`, which has consistent meaning across all four update rules."
       },
       "attrs": { "q_num_heads": 1, "kv_num_heads": 1, "update_rule": "delta", "scale": 0.5 },
       "inputs": {
     {
       "name": "gated_delta_zero_scalar_dk17_no_past",
       "provenance": {
+        "notes": "With past state omitted, key head size 17 exceeds the serial limit and is not divisible by four, selecting the scalar zero-state recurrence."
       },
       "attrs": { "q_num_heads": 1, "kv_num_heads": 1, "update_rule": "gated_delta", "scale": 0.25 },
       "inputs": {
     {
       "name": "gated_delta_state_chunked_f16_seq1024",
       "provenance": {
+        "notes": "A 1,024-token chunked prefill uses float16 activations and state. Every pass accumulates in float32 and narrows only loads and stores."
       },
       "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "gated_delta", "scale": 0.35 },
       "inputs": {
         "outputT": { "dtype": "float16", "shape": [1, 1024, 32], "tolerance": 0.04, "relTolerance": 0.04 },
         "presentStateT": { "dtype": "float16", "shape": [1, 2, 8, 8], "tolerance": 0.04, "relTolerance": 0.04 }
       }
+    },
+    {
+      "name": "gated_delta_chunked_strong_decay_prefix_underflow",
+      "provenance": {
+        "notes": "Per-head decay -8 on every token of a 1024-token prefill: over a 16-token chunk the within-chunk decay prefix reaches exp(-128), below the f32 subnormal floor, so a chunked formulation that divides keys by exp(prefix) produces Inf/NaN while the recurrence itself applies a finite exp(-8) per token. Bounded Q/K/V, beta 0.5, no past state."
+      },
+      "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "gated_delta" },
+      "inputs": {
+        "queryT": {
+          "dtype": "float32",
+          "shape": [1, 1024, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.021, "cosStep": 0.033, "scale": 1.0 }
+        },
+        "keyT": {
+          "dtype": "float32",
+          "shape": [1, 1024, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.027, "scale": 1.0 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [1, 1024, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.009, "cosStep": 0.043, "scale": 1.0 }
+        },
+        "decayT": { "dtype": "float32", "shape": [1, 1024, 2], "data": { "kind": "constant", "value": -8.0 } },
+        "betaT": { "dtype": "float32", "shape": [1, 1024, 2], "data": { "kind": "constant", "value": 0.5 } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 1024, 128], "tolerance": 0.0002, "relTolerance": 0.0005 },
+        "presentStateT": { "dtype": "float32", "shape": [1, 2, 32, 32], "tolerance": 0.0002, "relTolerance": 0.0005 }
+      },
+      "skipGpu": {
+        "category": "todo",
+        "reason": "The chunked prefill decomposition materializes a normalized key as k divided by exp(the within-chunk decay prefix). A sustained decay of -8 drives that prefix to -128 over a 16-token chunk, so the exponential underflows to zero and the division returns Inf/NaN, while every pairwise term the decomposition actually needs carries the bounded ratio exp(p_t - p_s) <= 1. The serial route applies a finite exp(-8) per token and stays finite. Fixing this needs the chunk operands re-anchored to a per-chunk pivot, with the entry-state terms kept in their absolute form."
+      }
+    },
+    {
+      "name": "gated_delta_chunked_moderate_decay_control",
+      "provenance": {
+        "notes": "Decay -4 keeps each 16-token chunk prefix at -64, within the float32 exponent range, so `k / exp(prefix)` remains finite. It is the finite-prefix control for `gated_delta_chunked_strong_decay_prefix_underflow`."
+      },
+      "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "gated_delta" },
+      "inputs": {
+        "queryT": {
+          "dtype": "float32",
+          "shape": [1, 1024, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.021, "cosStep": 0.033, "scale": 1.0 }
+        },
+        "keyT": {
+          "dtype": "float32",
+          "shape": [1, 1024, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.027, "scale": 1.0 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [1, 1024, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.009, "cosStep": 0.043, "scale": 1.0 }
+        },
+        "decayT": { "dtype": "float32", "shape": [1, 1024, 2], "data": { "kind": "constant", "value": -4.0 } },
+        "betaT": { "dtype": "float32", "shape": [1, 1024, 2], "data": { "kind": "constant", "value": 0.5 } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 1024, 128], "tolerance": 0.0002, "relTolerance": 0.0005 },
+        "presentStateT": { "dtype": "float32", "shape": [1, 2, 32, 32], "tolerance": 0.0002, "relTolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "linear_geometry_float32_float32_g0_zero",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "linear", "state_window": 3, "scale": 0.375 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float32",
+          "shape": [2, 33, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float32",
+          "shape": [2, 33, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [2, 33, 14],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 33, 28] },
+        "presentStateT": { "dtype": "float32", "shape": [3, 2, 2, 8, 7] }
+      }
+    },
+    {
+      "name": "linear_geometry_float32_float32_g0_state",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "linear", "state_window": 3, "scale": 0.375 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float32",
+          "shape": [2, 33, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float32",
+          "shape": [2, 33, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [2, 33, 14],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [3, 2, 2, 8, 7],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 33, 28] },
+        "presentStateT": { "dtype": "float32", "shape": [3, 2, 2, 8, 7] }
+      }
+    },
+    {
+      "name": "linear_geometry_float32_float32_g1_zero",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 2, "kv_num_heads": 4, "update_rule": "linear", "state_window": 0, "scale": 0 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float32",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float32",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [1, 17, 20],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 17, 20] },
+        "presentStateT": { "dtype": "float32", "shape": [1, 4, 6, 5] }
+      }
+    },
+    {
+      "name": "linear_geometry_float32_float32_g1_state",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 2, "kv_num_heads": 4, "update_rule": "linear", "state_window": 0, "scale": 0 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float32",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float32",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [1, 17, 20],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [1, 4, 6, 5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 17, 20] },
+        "presentStateT": { "dtype": "float32", "shape": [1, 4, 6, 5] }
+      }
+    },
+    {
+      "name": "linear_geometry_float32_float32_g2_zero",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 3, "kv_num_heads": 3, "update_rule": "linear", "state_window": 2, "scale": 0.375 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float32",
+          "shape": [2, 0, 48],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float32",
+          "shape": [2, 0, 16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [2, 0, 27],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 0, 27] },
+        "presentStateT": { "dtype": "float32", "shape": [2, 2, 3, 16, 9] }
+      }
+    },
+    {
+      "name": "linear_geometry_float32_float32_g2_state",
+      "provenance": {
+        "notes": "A one-token update inside a two-slot state window verifies zero-fill of the leading slot and a nonuniform updated state in the final slot, with shared keys, odd value dimensions, and independent activation/state dtypes."
+      },
+      "attrs": { "q_num_heads": 3, "kv_num_heads": 3, "update_rule": "linear", "state_window": 2, "scale": 0.375 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float32",
+          "shape": [2, 1, 48],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float32",
+          "shape": [2, 1, 16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [2, 1, 27],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [2, 2, 3, 16, 9],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 1, 27] },
+        "presentStateT": { "dtype": "float32", "shape": [2, 2, 3, 16, 9] }
+      }
+    },
+    {
+      "name": "linear_geometry_float16_float16_g0_zero",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "linear", "state_window": 3, "scale": 0.375 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float16",
+          "shape": [2, 33, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float16",
+          "shape": [2, 33, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float16",
+          "shape": [2, 33, 14],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 33, 28], "tolerance": 6e-8, "relTolerance": 0.0005 },
+        "presentStateT": { "dtype": "float16", "shape": [3, 2, 2, 8, 7], "tolerance": 6e-8, "relTolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "linear_geometry_float16_float16_g0_state",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "linear", "state_window": 3, "scale": 0.375 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float16",
+          "shape": [2, 33, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float16",
+          "shape": [2, 33, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float16",
+          "shape": [2, 33, 14],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "pastStateT": {
+          "dtype": "float16",
+          "shape": [3, 2, 2, 8, 7],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 33, 28], "tolerance": 6e-8, "relTolerance": 0.0005 },
+        "presentStateT": { "dtype": "float16", "shape": [3, 2, 2, 8, 7], "tolerance": 6e-8, "relTolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "linear_geometry_float16_float16_g1_zero",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 2, "kv_num_heads": 4, "update_rule": "linear", "state_window": 0, "scale": 0 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float16",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float16",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float16",
+          "shape": [1, 17, 20],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [1, 17, 20], "tolerance": 6e-8, "relTolerance": 0.0005 },
+        "presentStateT": { "dtype": "float16", "shape": [1, 4, 6, 5], "tolerance": 6e-8, "relTolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "linear_geometry_float16_float16_g1_state",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 2, "kv_num_heads": 4, "update_rule": "linear", "state_window": 0, "scale": 0 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float16",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float16",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float16",
+          "shape": [1, 17, 20],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "pastStateT": {
+          "dtype": "float16",
+          "shape": [1, 4, 6, 5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [1, 17, 20], "tolerance": 6e-8, "relTolerance": 0.0005 },
+        "presentStateT": { "dtype": "float16", "shape": [1, 4, 6, 5], "tolerance": 6e-8, "relTolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "linear_geometry_float16_float16_g2_zero",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 3, "kv_num_heads": 3, "update_rule": "linear", "state_window": 2, "scale": 0.375 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float16",
+          "shape": [2, 0, 48],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float16",
+          "shape": [2, 0, 16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float16",
+          "shape": [2, 0, 27],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 0, 27], "tolerance": 6e-8, "relTolerance": 0.0005 },
+        "presentStateT": { "dtype": "float16", "shape": [2, 2, 3, 16, 9], "tolerance": 6e-8, "relTolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "linear_geometry_float16_float16_g2_state",
+      "provenance": {
+        "notes": "A one-token update inside a two-slot state window verifies zero-fill of the leading slot and a nonuniform updated state in the final slot, with shared keys, odd value dimensions, and independent activation/state dtypes."
+      },
+      "attrs": { "q_num_heads": 3, "kv_num_heads": 3, "update_rule": "linear", "state_window": 2, "scale": 0.375 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float16",
+          "shape": [2, 1, 48],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float16",
+          "shape": [2, 1, 16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float16",
+          "shape": [2, 1, 27],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "pastStateT": {
+          "dtype": "float16",
+          "shape": [2, 2, 3, 16, 9],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 1, 27], "tolerance": 6e-8, "relTolerance": 0.0005 },
+        "presentStateT": { "dtype": "float16", "shape": [2, 2, 3, 16, 9], "tolerance": 6e-8, "relTolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "linear_geometry_float16_float32_g0_zero",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "linear", "state_window": 3, "scale": 0.375 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float16",
+          "shape": [2, 33, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float16",
+          "shape": [2, 33, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float16",
+          "shape": [2, 33, 14],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 33, 28], "tolerance": 6e-8, "relTolerance": 0.0005 },
+        "presentStateT": { "dtype": "float32", "shape": [3, 2, 2, 8, 7] }
+      }
+    },
+    {
+      "name": "linear_geometry_float16_float32_g0_state",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "linear", "state_window": 3, "scale": 0.375 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float16",
+          "shape": [2, 33, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float16",
+          "shape": [2, 33, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float16",
+          "shape": [2, 33, 14],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [3, 2, 2, 8, 7],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 33, 28], "tolerance": 6e-8, "relTolerance": 0.0005 },
+        "presentStateT": { "dtype": "float32", "shape": [3, 2, 2, 8, 7] }
+      }
+    },
+    {
+      "name": "linear_geometry_float16_float32_g1_zero",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 2, "kv_num_heads": 4, "update_rule": "linear", "state_window": 0, "scale": 0 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float16",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float16",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float16",
+          "shape": [1, 17, 20],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [1, 17, 20], "tolerance": 6e-8, "relTolerance": 0.0005 },
+        "presentStateT": { "dtype": "float32", "shape": [1, 4, 6, 5] }
+      }
+    },
+    {
+      "name": "linear_geometry_float16_float32_g1_state",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 2, "kv_num_heads": 4, "update_rule": "linear", "state_window": 0, "scale": 0 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float16",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float16",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float16",
+          "shape": [1, 17, 20],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [1, 4, 6, 5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [1, 17, 20], "tolerance": 6e-8, "relTolerance": 0.0005 },
+        "presentStateT": { "dtype": "float32", "shape": [1, 4, 6, 5] }
+      }
+    },
+    {
+      "name": "linear_geometry_float16_float32_g2_zero",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 3, "kv_num_heads": 3, "update_rule": "linear", "state_window": 2, "scale": 0.375 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float16",
+          "shape": [2, 0, 48],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float16",
+          "shape": [2, 0, 16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float16",
+          "shape": [2, 0, 27],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 0, 27], "tolerance": 6e-8, "relTolerance": 0.0005 },
+        "presentStateT": { "dtype": "float32", "shape": [2, 2, 3, 16, 9] }
+      }
+    },
+    {
+      "name": "linear_geometry_float16_float32_g2_state",
+      "provenance": {
+        "notes": "A one-token update inside a two-slot state window verifies zero-fill of the leading slot and a nonuniform updated state in the final slot, with shared keys, odd value dimensions, and independent activation/state dtypes."
+      },
+      "attrs": { "q_num_heads": 3, "kv_num_heads": 3, "update_rule": "linear", "state_window": 2, "scale": 0.375 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float16",
+          "shape": [2, 1, 48],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float16",
+          "shape": [2, 1, 16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float16",
+          "shape": [2, 1, 27],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [2, 2, 3, 16, 9],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 1, 27], "tolerance": 6e-8, "relTolerance": 0.0005 },
+        "presentStateT": { "dtype": "float32", "shape": [2, 2, 3, 16, 9] }
+      }
+    },
+    {
+      "name": "linear_geometry_float32_float16_g0_zero",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "linear", "state_window": 3, "scale": 0.375 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float32",
+          "shape": [2, 33, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float32",
+          "shape": [2, 33, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [2, 33, 14],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 33, 28] },
+        "presentStateT": { "dtype": "float16", "shape": [3, 2, 2, 8, 7], "tolerance": 6e-8, "relTolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "linear_geometry_float32_float16_g0_state",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 4, "kv_num_heads": 2, "update_rule": "linear", "state_window": 3, "scale": 0.375 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float32",
+          "shape": [2, 33, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float32",
+          "shape": [2, 33, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [2, 33, 14],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "pastStateT": {
+          "dtype": "float16",
+          "shape": [3, 2, 2, 8, 7],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 33, 28] },
+        "presentStateT": { "dtype": "float16", "shape": [3, 2, 2, 8, 7], "tolerance": 6e-8, "relTolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "linear_geometry_float32_float16_g1_zero",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 2, "kv_num_heads": 4, "update_rule": "linear", "state_window": 0, "scale": 0 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float32",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float32",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [1, 17, 20],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 17, 20] },
+        "presentStateT": { "dtype": "float16", "shape": [1, 4, 6, 5], "tolerance": 6e-8, "relTolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "linear_geometry_float32_float16_g1_state",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 2, "kv_num_heads": 4, "update_rule": "linear", "state_window": 0, "scale": 0 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float32",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float32",
+          "shape": [1, 17, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [1, 17, 20],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "pastStateT": {
+          "dtype": "float16",
+          "shape": [1, 4, 6, 5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 17, 20] },
+        "presentStateT": { "dtype": "float16", "shape": [1, 4, 6, 5], "tolerance": 6e-8, "relTolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "linear_geometry_float32_float16_g2_zero",
+      "provenance": {
+        "notes": "Small-head linear recurrence with shared keys, standard or inverse query grouping, odd value dimensions, independent state dtype, and retained or empty-sequence state windows."
+      },
+      "attrs": { "q_num_heads": 3, "kv_num_heads": 3, "update_rule": "linear", "state_window": 2, "scale": 0.375 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float32",
+          "shape": [2, 0, 48],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float32",
+          "shape": [2, 0, 16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [2, 0, 27],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 0, 27] },
+        "presentStateT": { "dtype": "float16", "shape": [2, 2, 3, 16, 9], "tolerance": 6e-8, "relTolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "linear_geometry_float32_float16_g2_state",
+      "provenance": {
+        "notes": "A one-token update inside a two-slot state window verifies zero-fill of the leading slot and a nonuniform updated state in the final slot, with shared keys, odd value dimensions, and independent activation/state dtypes."
+      },
+      "attrs": { "q_num_heads": 3, "kv_num_heads": 3, "update_rule": "linear", "state_window": 2, "scale": 0.375 },
+      "inputs": {
+        "queryT": {
+          "dtype": "float32",
+          "shape": [2, 1, 48],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "keyT": {
+          "dtype": "float32",
+          "shape": [2, 1, 16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [2, 1, 27],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.31, "scale": 0.05 }
+        },
+        "pastStateT": {
+          "dtype": "float16",
+          "shape": [2, 2, 3, 16, 9],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 1, 27] },
+        "presentStateT": { "dtype": "float16", "shape": [2, 2, 3, 16, 9], "tolerance": 6e-8, "relTolerance": 0.0005 }
+      }
     }
   ]
 }