sync 6fdf6301e2bb

Browse files

Files changed (6) hide show

README.md +91 -0
build/webgpu/bench.json +924 -0
build/webgpu/engram-gate-rows.wgsl.jinja +172 -0
build/webgpu/manifest.json +196 -0
build/webgpu/metadata.json +26 -0
build/webgpu/test.json +2320 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# com.microsoft.EngramGate
+`com.microsoft`  ·  ONNX Runtime contrib operator  ·  contrib since_version 1
+## Description
+Fuses the Engram gate. For each `(batch, sequence, hc_mult)` row it RMS-normalizes `key` and `query` with their per-stream scales, forms `dot = sum(RMSNorm(key) * RMSNorm(query)) / sqrt(hidden_size)`, and writes `sigmoid(sign(dot) * sqrt(max(abs(dot), 1e-6))) * value`, broadcasting the value row shared by every hyper-connection. That `1e-6` floor is fixed and is not `epsilon`; a zero dot product gives a gate of exactly `0.5`. Both sums of squares and the scaled cross term accumulate in one float32 pass, and only the store narrows. Bfloat16 is not implemented.
+See the [ONNX Runtime `EngramGate` contrib-operator spec](https://github.com/microsoft/onnxruntime/blob/main/docs/ContribOperators.md#com.microsoft.EngramGate) for the reference semantics.
+## Inputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `keyT` | `key` | `T` | `4` | — | Projected Engram keys with shape `(batch_size, sequence_length, hc_mult, hidden_size)`. | required |
+| `queryT` | `query` | `T` | `4` | — | Hidden-state queries, exactly the same shape as `key`. The upstream kernels require equality rather than broadcasting, and so does this one. | required |
+| `valueT` | `value` | `T` | `3` | — | Projected Engram value shared by every hyper-connection, with shape `(batch_size, sequence_length, hidden_size)`. Each row of it is gated `hc_mult` times, once per stream. | required |
+| `keyNormScaleT` | `key_norm_scale` | `T` | `2` | — | RMSNorm scale for keys with shape `(hc_mult, hidden_size)`: one weight row per hyper-connection stream, selected by the row's stream index. | required |
+| `queryNormScaleT` | `query_norm_scale` | `T` | `2` | — | RMSNorm scale for queries with shape `(hc_mult, hidden_size)`, indexed by the same stream index as the key scale. | required |
+## Outputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `outputT` | `output` | `T` | `4` | same as `keyT` | The gated value tensor, with the same shape and dtype as `key`. Every element of a row carries the same gate, so the row's stream index is only observable through the scale lookup. | required |
+## Attributes
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `epsilon` | `0.00001` | Constant added to both mean-of-squares denominators before the reciprocal square root. It does not reach the gate: the `1e-6` floor under the square root is a separate fixed constant. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16` |
+## Implementation variants
+One implementation is selected per call from the device capabilities, the request shapes and the dtypes; these notes say what each one covers.
+- `subgroup_rows_vec4` — One aligned lane group per gated row inside a subgroup, with the row's value words held in registers across the fold. The three accumulators travel as one vector through a segmented butterfly over the low lane bits, so several rows share a subgroup with no workgroup memory and no barrier.
+- `packed_rows_vec4` — Portable vec4 route: the workgroup is split into a power-of-two lane group per row and each group folds its own contiguous slice of one shared vector array, so a short row does not idle the workgroup and no subgroup support is required. It also carries rows too wide to stage in registers, which it walks twice.
+- `subgroup_rows_scalar` — The barrier-free lane-group schedule for a hidden size that is not a multiple of four: the same segmented butterfly over scalar loads.
+- `packed_rows_scalar` — Scalar fallback for a hidden size that is not a multiple of four, and the route every device can select. It keeps the lane-group-per-row schedule so a narrow row still fills the workgroup.
+## Device requirements
+Some implementation variants require `subgroups`. These are route-specific capabilities, not package-wide requirements; availability also depends on the request shape and dtype.
+## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
+- [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
+- [`test.json`](build/webgpu/test.json) — correctness cases
+- [`bench.json`](build/webgpu/bench.json) — benchmark cases
+- [`engram-gate-rows.wgsl.jinja`](build/webgpu/engram-gate-rows.wgsl.jinja)
+## Use with `@huggingface/kernels`
+```sh
+npm install --save-exact @huggingface/kernels@0.0.1-preview.3
+```
+Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
+The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+It follows the `v1` branch as fixes land. To pin exact artifact bytes, pass a 40-character commit `revision` instead of `version`.
+Replace each `*Data` placeholder with a typed array containing the corresponding input data.
+```js
+import { getKernel } from "@huggingface/kernels";
+const kernel = await getKernel("webgpu-kernels/com.microsoft.EngramGate", { version: 1 });
+const { outputT } = await kernel({
+  keyT: { data: keyTData, shape: [1, 1, 1, 2] },
+  queryT: { data: queryTData, shape: [1, 1, 1, 2] },
+  valueT: { data: valueTData, shape: [1, 1, 2] },
+  keyNormScaleT: { data: keyNormScaleTData, shape: [1, 2] },
+  queryNormScaleT: { data: queryNormScaleTData, shape: [1, 2] },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,924 @@

+{
+  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256, 512], "LANE_WORDS": [1, 2, 4, 8] },
+  "cases": [
+    {
+      "name": "engram-nanowhale-decode-h320-f32",
+      "preset": "smoke",
+      "provenance": {
+        "model": "DeepSeek-V4 nanowhale Engram block",
+        "notes": "One decoded token: four hyper-connection streams of 320 channels. The whole operator moves 21 KB, so the measurement is launch latency rather than bandwidth."
+      },
+      "vars": { "dtype": "float32", "rows": 4, "tokens": 1, "hcMult": 4, "hidden": 320 },
+      "inputs": {
+        "keyT": { "shape": [1, 1, 4, 320], "dtype": "float32", "dist": "normal", "seed": 601, "scale": 1 },
+        "queryT": { "shape": [1, 1, 4, 320], "dtype": "float32", "dist": "normal", "seed": 602, "scale": 1 },
+        "valueT": { "shape": [1, 1, 320], "dtype": "float32", "dist": "normal", "seed": 603, "scale": 1 },
+        "keyNormScaleT": {
+          "shape": [4, 320],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 604,
+          "min": 0.2,
+          "max": 1.8
+        },
+        "queryNormScaleT": {
+          "shape": [4, 320],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 605,
+          "min": 0.2,
+          "max": 1.8
+        }
+      },
+      "outputs": { "outputT": { "shape": [1, 1, 4, 320], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(3 * args.rows * args.hidden + args.tokens * args.hidden + 2 * args.hcMult * args.hidden) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "engram-nanowhale-decode-h320-f16",
+      "preset": "smoke",
+      "provenance": {
+        "model": "DeepSeek-V4 nanowhale Engram block",
+        "notes": "The float16 decode step at the same geometry."
+      },
+      "vars": { "dtype": "float16", "rows": 4, "tokens": 1, "hcMult": 4, "hidden": 320 },
+      "inputs": {
+        "keyT": { "shape": [1, 1, 4, 320], "dtype": "float16", "dist": "normal", "seed": 611, "scale": 1 },
+        "queryT": { "shape": [1, 1, 4, 320], "dtype": "float16", "dist": "normal", "seed": 612, "scale": 1 },
+        "valueT": { "shape": [1, 1, 320], "dtype": "float16", "dist": "normal", "seed": 613, "scale": 1 },
+        "keyNormScaleT": {
+          "shape": [4, 320],
+          "dtype": "float16",
+          "dist": "uniform",
+          "seed": 614,
+          "min": 0.2,
+          "max": 1.8
+        },
+        "queryNormScaleT": {
+          "shape": [4, 320],
+          "dtype": "float16",
+          "dist": "uniform",
+          "seed": 615,
+          "min": 0.2,
+          "max": 1.8
+        }
+      },
+      "outputs": { "outputT": { "shape": [1, 1, 4, 320], "dtype": "float16" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(3 * args.rows * args.hidden + args.tokens * args.hidden + 2 * args.hcMult * args.hidden) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "engram-nanowhale-prefill-s64-h320-f32",
+      "preset": "smoke",
+      "provenance": {
+        "model": "DeepSeek-V4 nanowhale Engram block",
+        "notes": "A 64-token prefill: 256 rows of 320, still small enough to sit inside cache."
+      },
+      "vars": { "dtype": "float32", "rows": 256, "tokens": 64, "hcMult": 4, "hidden": 320 },
+      "inputs": {
+        "keyT": { "shape": [1, 64, 4, 320], "dtype": "float32", "dist": "normal", "seed": 621, "scale": 1 },
+        "queryT": { "shape": [1, 64, 4, 320], "dtype": "float32", "dist": "normal", "seed": 622, "scale": 1 },
+        "valueT": { "shape": [1, 64, 320], "dtype": "float32", "dist": "normal", "seed": 623, "scale": 1 },
+        "keyNormScaleT": {
+          "shape": [4, 320],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 624,
+          "min": 0.2,
+          "max": 1.8
+        },
+        "queryNormScaleT": {
+          "shape": [4, 320],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 625,
+          "min": 0.2,
+          "max": 1.8
+        }
+      },
+      "outputs": { "outputT": { "shape": [1, 64, 4, 320], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(3 * args.rows * args.hidden + args.tokens * args.hidden + 2 * args.hcMult * args.hidden) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "engram-nanowhale-prefill-s512-h320-f16",
+      "preset": "smoke",
+      "provenance": {
+        "model": "DeepSeek-V4 nanowhale Engram block",
+        "notes": "Measures the DeepSeek-V4 nanowhale Engram gate over a 512-token prefill: 2048 rows (512 tokens x 4 heads) of width 320 in float16."
+      },
+      "vars": { "dtype": "float16", "rows": 2048, "tokens": 512, "hcMult": 4, "hidden": 320 },
+      "inputs": {
+        "keyT": { "shape": [1, 512, 4, 320], "dtype": "float16", "dist": "normal", "seed": 631, "scale": 1 },
+        "queryT": { "shape": [1, 512, 4, 320], "dtype": "float16", "dist": "normal", "seed": 632, "scale": 1 },
+        "valueT": { "shape": [1, 512, 320], "dtype": "float16", "dist": "normal", "seed": 633, "scale": 1 },
+        "keyNormScaleT": {
+          "shape": [4, 320],
+          "dtype": "float16",
+          "dist": "uniform",
+          "seed": 634,
+          "min": 0.2,
+          "max": 1.8
+        },
+        "queryNormScaleT": {
+          "shape": [4, 320],
+          "dtype": "float16",
+          "dist": "uniform",
+          "seed": 635,
+          "min": 0.2,
+          "max": 1.8
+        }
+      },
+      "outputs": { "outputT": { "shape": [1, 512, 4, 320], "dtype": "float16" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(3 * args.rows * args.hidden + args.tokens * args.hidden + 2 * args.hcMult * args.hidden) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "engram-v4flash-decode-h4096-f32",
+      "preset": "smoke",
+      "provenance": {
+        "model": "DeepSeek-V4 v4flash Engram block",
+        "notes": "One decoded token of the wide configuration: four rows of 4096, the fewest rows this operator ever runs with at a non-trivial row width."
+      },
+      "vars": { "dtype": "float32", "rows": 4, "tokens": 1, "hcMult": 4, "hidden": 4096 },
+      "inputs": {
+        "keyT": { "shape": [1, 1, 4, 4096], "dtype": "float32", "dist": "normal", "seed": 641, "scale": 1 },
+        "queryT": { "shape": [1, 1, 4, 4096], "dtype": "float32", "dist": "normal", "seed": 642, "scale": 1 },
+        "valueT": { "shape": [1, 1, 4096], "dtype": "float32", "dist": "normal", "seed": 643, "scale": 1 },
+        "keyNormScaleT": {
+          "shape": [4, 4096],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 644,
+          "min": 0.2,
+          "max": 1.8
+        },
+        "queryNormScaleT": {
+          "shape": [4, 4096],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 645,
+          "min": 0.2,
+          "max": 1.8
+        }
+      },
+      "outputs": { "outputT": { "shape": [1, 1, 4, 4096], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(3 * args.rows * args.hidden + args.tokens * args.hidden + 2 * args.hcMult * args.hidden) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "engram-v4flash-decode-h4096-f16",
+      "preset": "smoke",
+      "provenance": {
+        "model": "DeepSeek-V4 v4flash Engram block",
+        "notes": "The float16 half of the wide decode step."
+      },
+      "vars": { "dtype": "float16", "rows": 4, "tokens": 1, "hcMult": 4, "hidden": 4096 },
+      "inputs": {
+        "keyT": { "shape": [1, 1, 4, 4096], "dtype": "float16", "dist": "normal", "seed": 651, "scale": 1 },
+        "queryT": { "shape": [1, 1, 4, 4096], "dtype": "float16", "dist": "normal", "seed": 652, "scale": 1 },
+        "valueT": { "shape": [1, 1, 4096], "dtype": "float16", "dist": "normal", "seed": 653, "scale": 1 },
+        "keyNormScaleT": {
+          "shape": [4, 4096],
+          "dtype": "float16",
+          "dist": "uniform",
+          "seed": 654,
+          "min": 0.2,
+          "max": 1.8
+        },
+        "queryNormScaleT": {
+          "shape": [4, 4096],
+          "dtype": "float16",
+          "dist": "uniform",
+          "seed": 655,
+          "min": 0.2,
+          "max": 1.8
+        }
+      },
+      "outputs": { "outputT": { "shape": [1, 1, 4, 4096], "dtype": "float16" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(3 * args.rows * args.hidden + args.tokens * args.hidden + 2 * args.hcMult * args.hidden) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "engram-v4flash-prefill-s512-h4096-f16",
+      "preset": "smoke",
+      "provenance": {
+        "model": "DeepSeek-V4 v4flash Engram block",
+        "notes": "A 512-token prefill of the wide configuration: 8192 rows of 4096, a 55 MB working set well past cache, so the bandwidth figure is a DRAM figure."
+      },
+      "vars": { "dtype": "float16", "rows": 2048, "tokens": 512, "hcMult": 4, "hidden": 4096 },
+      "inputs": {
+        "keyT": { "shape": [1, 512, 4, 4096], "dtype": "float16", "dist": "normal", "seed": 661, "scale": 1 },
+        "queryT": { "shape": [1, 512, 4, 4096], "dtype": "float16", "dist": "normal", "seed": 662, "scale": 1 },
+        "valueT": { "shape": [1, 512, 4096], "dtype": "float16", "dist": "normal", "seed": 663, "scale": 1 },
+        "keyNormScaleT": {
+          "shape": [4, 4096],
+          "dtype": "float16",
+          "dist": "uniform",
+          "seed": 664,
+          "min": 0.2,
+          "max": 1.8
+        },
+        "queryNormScaleT": {
+          "shape": [4, 4096],
+          "dtype": "float16",
+          "dist": "uniform",
+          "seed": 665,
+          "min": 0.2,
+          "max": 1.8
+        }
+      },
+      "outputs": { "outputT": { "shape": [1, 512, 4, 4096], "dtype": "float16" } },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(3 * args.rows * args.hidden + args.tokens * args.hidden + 2 * args.hcMult * args.hidden) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "engram-narrow-h64-f32",
+      "preset": "smoke",
+      "provenance": {
+        "notes": "64 channels per row: 2048 rows each narrower than a subgroup, where a workgroup-per-row schedule would leave most of every workgroup idle."
+      },
+      "vars": { "dtype": "float32", "rows": 2048, "tokens": 512, "hcMult": 4, "hidden": 64 },
+      "inputs": {
+        "keyT": { "shape": [1, 512, 4, 64], "dtype": "float32", "dist": "normal", "seed": 671, "scale": 1 },
+        "queryT": { "shape": [1, 512, 4, 64], "dtype": "float32", "dist": "normal", "seed": 672, "scale": 1 },
+        "valueT": { "shape": [1, 512, 64], "dtype": "float32", "dist": "normal", "seed": 673, "scale": 1 },
+        "keyNormScaleT": {
+          "shape": [4, 64],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 674,
+          "min": 0.2,
+          "max": 1.8
+        },
+        "queryNormScaleT": {
+          "shape": [4, 64],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 675,
+          "min": 0.2,
+          "max": 1.8
+        }
+      },
+      "outputs": { "outputT": { "shape": [1, 512, 4, 64], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(3 * args.rows * args.hidden + args.tokens * args.hidden + 2 * args.hcMult * args.hidden) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "engram-narrow-h128-f16",
+      "preset": "smoke",
+      "provenance": {
+        "notes": "128 channels per row in float16, the other narrow width the hyper-connection family uses."
+      },
+      "vars": { "dtype": "float16", "rows": 2048, "tokens": 512, "hcMult": 4, "hidden": 128 },
+      "inputs": {
+        "keyT": { "shape": [1, 512, 4, 128], "dtype": "float16", "dist": "normal", "seed": 681, "scale": 1 },
+        "queryT": { "shape": [1, 512, 4, 128], "dtype": "float16", "dist": "normal", "seed": 682, "scale": 1 },
+        "valueT": { "shape": [1, 512, 128], "dtype": "float16", "dist": "normal", "seed": 683, "scale": 1 },
+        "keyNormScaleT": {
+          "shape": [4, 128],
+          "dtype": "float16",
+          "dist": "uniform",
+          "seed": 684,
+          "min": 0.2,
+          "max": 1.8
+        },
+        "queryNormScaleT": {
+          "shape": [4, 128],
+          "dtype": "float16",
+          "dist": "uniform",
+          "seed": 685,
+          "min": 0.2,
+          "max": 1.8
+        }
+      },
+      "outputs": { "outputT": { "shape": [1, 512, 4, 128], "dtype": "float16" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(3 * args.rows * args.hidden + args.tokens * args.hidden + 2 * args.hcMult * args.hidden) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "engram-scalar-h130-f32",
+      "preset": "smoke",
+      "provenance": { "notes": "A 130-channel row is not divisible by four and exercises a narrow scalar-width tail." },
+      "vars": { "dtype": "float32", "rows": 2048, "tokens": 512, "hcMult": 4, "hidden": 130 },
+      "inputs": {
+        "keyT": { "shape": [1, 512, 4, 130], "dtype": "float32", "dist": "normal", "seed": 691, "scale": 1 },
+        "queryT": { "shape": [1, 512, 4, 130], "dtype": "float32", "dist": "normal", "seed": 692, "scale": 1 },
+        "valueT": { "shape": [1, 512, 130], "dtype": "float32", "dist": "normal", "seed": 693, "scale": 1 },
+        "keyNormScaleT": {
+          "shape": [4, 130],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 694,
+          "min": 0.2,
+          "max": 1.8
+        },
+        "queryNormScaleT": {
+          "shape": [4, 130],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 695,
+          "min": 0.2,
+          "max": 1.8
+        }
+      },
+      "outputs": { "outputT": { "shape": [1, 512, 4, 130], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(3 * args.rows * args.hidden + args.tokens * args.hidden + 2 * args.hcMult * args.hidden) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "engram-batched-rows-h320-f16",
+      "preset": "scale",
+      "provenance": {
+        "model": "DeepSeek-V4 nanowhale Engram block",
+        "notes": "65536 rows of 320 from a batched prefill, past the point where a one-dimensional workgroup grid would need the folded row index."
+      },
+      "vars": { "dtype": "float16", "rows": 65536, "tokens": 16384, "hcMult": 4, "hidden": 320 },
+      "inputs": {
+        "keyT": { "shape": [32, 512, 4, 320], "dtype": "float16", "dist": "normal", "seed": 701, "scale": 1 },
+        "queryT": { "shape": [32, 512, 4, 320], "dtype": "float16", "dist": "normal", "seed": 702, "scale": 1 },
+        "valueT": { "shape": [32, 512, 320], "dtype": "float16", "dist": "normal", "seed": 703, "scale": 1 },
+        "keyNormScaleT": {
+          "shape": [4, 320],
+          "dtype": "float16",
+          "dist": "uniform",
+          "seed": 704,
+          "min": 0.2,
+          "max": 1.8
+        },
+        "queryNormScaleT": {
+          "shape": [4, 320],
+          "dtype": "float16",
+          "dist": "uniform",
+          "seed": 705,
+          "min": 0.2,
+          "max": 1.8
+        }
+      },
+      "outputs": { "outputT": { "shape": [32, 512, 4, 320], "dtype": "float16" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(3 * args.rows * args.hidden + args.tokens * args.hidden + 2 * args.hcMult * args.hidden) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "engram-v4flash-prefill-s512-h4096-f32",
+      "preset": "model",
+      "provenance": {
+        "model": "DeepSeek-V4 v4flash Engram block",
+        "notes": "The float32 counterpart of the wide prefill, a 110 MB working set."
+      },
+      "vars": { "dtype": "float32", "rows": 2048, "tokens": 512, "hcMult": 4, "hidden": 4096 },
+      "inputs": {
+        "keyT": { "shape": [1, 512, 4, 4096], "dtype": "float32", "dist": "normal", "seed": 711, "scale": 1 },
+        "queryT": { "shape": [1, 512, 4, 4096], "dtype": "float32", "dist": "normal", "seed": 712, "scale": 1 },
+        "valueT": { "shape": [1, 512, 4096], "dtype": "float32", "dist": "normal", "seed": 713, "scale": 1 },
+        "keyNormScaleT": {
+          "shape": [4, 4096],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 714,
+          "min": 0.2,
+          "max": 1.8
+        },
+        "queryNormScaleT": {
+          "shape": [4, 4096],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 715,
+          "min": 0.2,
+          "max": 1.8
+        }
+      },
+      "outputs": { "outputT": { "shape": [1, 512, 4, 4096], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(3 * args.rows * args.hidden + args.tokens * args.hidden + 2 * args.hcMult * args.hidden) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "engram-v4flash-prefill-s1024-h4096-f16",
+      "preset": "model",
+      "provenance": {
+        "model": "DeepSeek-V4 v4flash Engram block",
+        "notes": "A 1024-token prefill of the wide configuration: 110 MB in float16."
+      },
+      "vars": { "dtype": "float16", "rows": 4096, "tokens": 1024, "hcMult": 4, "hidden": 4096 },
+      "inputs": {
+        "keyT": { "shape": [1, 1024, 4, 4096], "dtype": "float16", "dist": "normal", "seed": 721, "scale": 1 },
+        "queryT": { "shape": [1, 1024, 4, 4096], "dtype": "float16", "dist": "normal", "seed": 722, "scale": 1 },
+        "valueT": { "shape": [1, 1024, 4096], "dtype": "float16", "dist": "normal", "seed": 723, "scale": 1 },
+        "keyNormScaleT": {
+          "shape": [4, 4096],
+          "dtype": "float16",
+          "dist": "uniform",
+          "seed": 724,
+          "min": 0.2,
+          "max": 1.8
+        },
+        "queryNormScaleT": {
+          "shape": [4, 4096],
+          "dtype": "float16",
+          "dist": "uniform",
+          "seed": 725,
+          "min": 0.2,
+          "max": 1.8
+        }
+      },
+      "outputs": { "outputT": { "shape": [1, 1024, 4, 4096], "dtype": "float16" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(3 * args.rows * args.hidden + args.tokens * args.hidden + 2 * args.hcMult * args.hidden) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "neighbor_float32_h63_r2048",
+      "preset": "all",
+      "inputs": {
+        "keyT": { "shape": [1, 512, 4, 63], "dtype": "float32", "seed": 691, "dist": "normal", "scale": 1 },
+        "queryT": { "shape": [1, 512, 4, 63], "dtype": "float32", "seed": 692, "dist": "normal", "scale": 1 },
+        "valueT": { "shape": [1, 512, 63], "dtype": "float32", "seed": 693, "dist": "normal", "scale": 1 },
+        "keyNormScaleT": {
+          "shape": [4, 63],
+          "dtype": "float32",
+          "seed": 694,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        },
+        "queryNormScaleT": {
+          "shape": [4, 63],
+          "dtype": "float32",
+          "seed": 695,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        }
+      },
+      "outputs": { "outputT": { "shape": [1, 512, 4, 63], "dtype": "float32" } }
+    },
+    {
+      "name": "neighbor_float32_h65_r2048",
+      "preset": "all",
+      "inputs": {
+        "keyT": { "shape": [1, 512, 4, 65], "dtype": "float32", "seed": 691, "dist": "normal", "scale": 1 },
+        "queryT": { "shape": [1, 512, 4, 65], "dtype": "float32", "seed": 692, "dist": "normal", "scale": 1 },
+        "valueT": { "shape": [1, 512, 65], "dtype": "float32", "seed": 693, "dist": "normal", "scale": 1 },
+        "keyNormScaleT": {
+          "shape": [4, 65],
+          "dtype": "float32",
+          "seed": 694,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        },
+        "queryNormScaleT": {
+          "shape": [4, 65],
+          "dtype": "float32",
+          "seed": 695,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        }
+      },
+      "outputs": { "outputT": { "shape": [1, 512, 4, 65], "dtype": "float32" } }
+    },
+    {
+      "name": "neighbor_float32_h127_r2048",
+      "preset": "all",
+      "inputs": {
+        "keyT": { "shape": [1, 512, 4, 127], "dtype": "float32", "seed": 691, "dist": "normal", "scale": 1 },
+        "queryT": { "shape": [1, 512, 4, 127], "dtype": "float32", "seed": 692, "dist": "normal", "scale": 1 },
+        "valueT": { "shape": [1, 512, 127], "dtype": "float32", "seed": 693, "dist": "normal", "scale": 1 },
+        "keyNormScaleT": {
+          "shape": [4, 127],
+          "dtype": "float32",
+          "seed": 694,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        },
+        "queryNormScaleT": {
+          "shape": [4, 127],
+          "dtype": "float32",
+          "seed": 695,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        }
+      },
+      "outputs": { "outputT": { "shape": [1, 512, 4, 127], "dtype": "float32" } }
+    },
+    {
+      "name": "neighbor_float32_h129_r2048",
+      "preset": "all",
+      "inputs": {
+        "keyT": { "shape": [1, 512, 4, 129], "dtype": "float32", "seed": 691, "dist": "normal", "scale": 1 },
+        "queryT": { "shape": [1, 512, 4, 129], "dtype": "float32", "seed": 692, "dist": "normal", "scale": 1 },
+        "valueT": { "shape": [1, 512, 129], "dtype": "float32", "seed": 693, "dist": "normal", "scale": 1 },
+        "keyNormScaleT": {
+          "shape": [4, 129],
+          "dtype": "float32",
+          "seed": 694,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        },
+        "queryNormScaleT": {
+          "shape": [4, 129],
+          "dtype": "float32",
+          "seed": 695,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        }
+      },
+      "outputs": { "outputT": { "shape": [1, 512, 4, 129], "dtype": "float32" } }
+    },
+    {
+      "name": "neighbor_float32_h131_r2048",
+      "preset": "all",
+      "inputs": {
+        "keyT": { "shape": [1, 512, 4, 131], "dtype": "float32", "seed": 691, "dist": "normal", "scale": 1 },
+        "queryT": { "shape": [1, 512, 4, 131], "dtype": "float32", "seed": 692, "dist": "normal", "scale": 1 },
+        "valueT": { "shape": [1, 512, 131], "dtype": "float32", "seed": 693, "dist": "normal", "scale": 1 },
+        "keyNormScaleT": {
+          "shape": [4, 131],
+          "dtype": "float32",
+          "seed": 694,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        },
+        "queryNormScaleT": {
+          "shape": [4, 131],
+          "dtype": "float32",
+          "seed": 695,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        }
+      },
+      "outputs": { "outputT": { "shape": [1, 512, 4, 131], "dtype": "float32" } }
+    },
+    {
+      "name": "neighbor_float32_h255_r2048",
+      "preset": "all",
+      "inputs": {
+        "keyT": { "shape": [1, 512, 4, 255], "dtype": "float32", "seed": 691, "dist": "normal", "scale": 1 },
+        "queryT": { "shape": [1, 512, 4, 255], "dtype": "float32", "seed": 692, "dist": "normal", "scale": 1 },
+        "valueT": { "shape": [1, 512, 255], "dtype": "float32", "seed": 693, "dist": "normal", "scale": 1 },
+        "keyNormScaleT": {
+          "shape": [4, 255],
+          "dtype": "float32",
+          "seed": 694,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        },
+        "queryNormScaleT": {
+          "shape": [4, 255],
+          "dtype": "float32",
+          "seed": 695,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        }
+      },
+      "outputs": { "outputT": { "shape": [1, 512, 4, 255], "dtype": "float32" } }
+    },
+    {
+      "name": "neighbor_float32_h257_r3",
+      "preset": "all",
+      "inputs": {
+        "keyT": { "shape": [1, 1, 3, 257], "dtype": "float32", "seed": 691, "dist": "normal", "scale": 1 },
+        "queryT": { "shape": [1, 1, 3, 257], "dtype": "float32", "seed": 692, "dist": "normal", "scale": 1 },
+        "valueT": { "shape": [1, 1, 257], "dtype": "float32", "seed": 693, "dist": "normal", "scale": 1 },
+        "keyNormScaleT": {
+          "shape": [3, 257],
+          "dtype": "float32",
+          "seed": 694,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        },
+        "queryNormScaleT": {
+          "shape": [3, 257],
+          "dtype": "float32",
+          "seed": 695,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        }
+      },
+      "outputs": { "outputT": { "shape": [1, 1, 3, 257], "dtype": "float32" } }
+    },
+    {
+      "name": "neighbor_float32_h257_r2048",
+      "preset": "all",
+      "inputs": {
+        "keyT": { "shape": [1, 512, 4, 257], "dtype": "float32", "seed": 691, "dist": "normal", "scale": 1 },
+        "queryT": { "shape": [1, 512, 4, 257], "dtype": "float32", "seed": 692, "dist": "normal", "scale": 1 },
+        "valueT": { "shape": [1, 512, 257], "dtype": "float32", "seed": 693, "dist": "normal", "scale": 1 },
+        "keyNormScaleT": {
+          "shape": [4, 257],
+          "dtype": "float32",
+          "seed": 694,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        },
+        "queryNormScaleT": {
+          "shape": [4, 257],
+          "dtype": "float32",
+          "seed": 695,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        }
+      },
+      "outputs": { "outputT": { "shape": [1, 512, 4, 257], "dtype": "float32" } }
+    },
+    {
+      "name": "neighbor_float16_h63_r2048",
+      "preset": "all",
+      "inputs": {
+        "keyT": { "shape": [1, 512, 4, 63], "dtype": "float16", "seed": 691, "dist": "normal", "scale": 1 },
+        "queryT": { "shape": [1, 512, 4, 63], "dtype": "float16", "seed": 692, "dist": "normal", "scale": 1 },
+        "valueT": { "shape": [1, 512, 63], "dtype": "float16", "seed": 693, "dist": "normal", "scale": 1 },
+        "keyNormScaleT": {
+          "shape": [4, 63],
+          "dtype": "float16",
+          "seed": 694,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        },
+        "queryNormScaleT": {
+          "shape": [4, 63],
+          "dtype": "float16",
+          "seed": 695,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        }
+      },
+      "outputs": { "outputT": { "shape": [1, 512, 4, 63], "dtype": "float16" } }
+    },
+    {
+      "name": "neighbor_float16_h65_r2048",
+      "preset": "all",
+      "inputs": {
+        "keyT": { "shape": [1, 512, 4, 65], "dtype": "float16", "seed": 691, "dist": "normal", "scale": 1 },
+        "queryT": { "shape": [1, 512, 4, 65], "dtype": "float16", "seed": 692, "dist": "normal", "scale": 1 },
+        "valueT": { "shape": [1, 512, 65], "dtype": "float16", "seed": 693, "dist": "normal", "scale": 1 },
+        "keyNormScaleT": {
+          "shape": [4, 65],
+          "dtype": "float16",
+          "seed": 694,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        },
+        "queryNormScaleT": {
+          "shape": [4, 65],
+          "dtype": "float16",
+          "seed": 695,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        }
+      },
+      "outputs": { "outputT": { "shape": [1, 512, 4, 65], "dtype": "float16" } }
+    },
+    {
+      "name": "neighbor_float16_h127_r2048",
+      "preset": "all",
+      "inputs": {
+        "keyT": { "shape": [1, 512, 4, 127], "dtype": "float16", "seed": 691, "dist": "normal", "scale": 1 },
+        "queryT": { "shape": [1, 512, 4, 127], "dtype": "float16", "seed": 692, "dist": "normal", "scale": 1 },
+        "valueT": { "shape": [1, 512, 127], "dtype": "float16", "seed": 693, "dist": "normal", "scale": 1 },
+        "keyNormScaleT": {
+          "shape": [4, 127],
+          "dtype": "float16",
+          "seed": 694,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        },
+        "queryNormScaleT": {
+          "shape": [4, 127],
+          "dtype": "float16",
+          "seed": 695,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        }
+      },
+      "outputs": { "outputT": { "shape": [1, 512, 4, 127], "dtype": "float16" } }
+    },
+    {
+      "name": "neighbor_float16_h129_r2048",
+      "preset": "all",
+      "inputs": {
+        "keyT": { "shape": [1, 512, 4, 129], "dtype": "float16", "seed": 691, "dist": "normal", "scale": 1 },
+        "queryT": { "shape": [1, 512, 4, 129], "dtype": "float16", "seed": 692, "dist": "normal", "scale": 1 },
+        "valueT": { "shape": [1, 512, 129], "dtype": "float16", "seed": 693, "dist": "normal", "scale": 1 },
+        "keyNormScaleT": {
+          "shape": [4, 129],
+          "dtype": "float16",
+          "seed": 694,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        },
+        "queryNormScaleT": {
+          "shape": [4, 129],
+          "dtype": "float16",
+          "seed": 695,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        }
+      },
+      "outputs": { "outputT": { "shape": [1, 512, 4, 129], "dtype": "float16" } }
+    },
+    {
+      "name": "neighbor_float16_h131_r2048",
+      "preset": "all",
+      "inputs": {
+        "keyT": { "shape": [1, 512, 4, 131], "dtype": "float16", "seed": 691, "dist": "normal", "scale": 1 },
+        "queryT": { "shape": [1, 512, 4, 131], "dtype": "float16", "seed": 692, "dist": "normal", "scale": 1 },
+        "valueT": { "shape": [1, 512, 131], "dtype": "float16", "seed": 693, "dist": "normal", "scale": 1 },
+        "keyNormScaleT": {
+          "shape": [4, 131],
+          "dtype": "float16",
+          "seed": 694,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        },
+        "queryNormScaleT": {
+          "shape": [4, 131],
+          "dtype": "float16",
+          "seed": 695,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        }
+      },
+      "outputs": { "outputT": { "shape": [1, 512, 4, 131], "dtype": "float16" } }
+    },
+    {
+      "name": "neighbor_float16_h255_r2048",
+      "preset": "all",
+      "inputs": {
+        "keyT": { "shape": [1, 512, 4, 255], "dtype": "float16", "seed": 691, "dist": "normal", "scale": 1 },
+        "queryT": { "shape": [1, 512, 4, 255], "dtype": "float16", "seed": 692, "dist": "normal", "scale": 1 },
+        "valueT": { "shape": [1, 512, 255], "dtype": "float16", "seed": 693, "dist": "normal", "scale": 1 },
+        "keyNormScaleT": {
+          "shape": [4, 255],
+          "dtype": "float16",
+          "seed": 694,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        },
+        "queryNormScaleT": {
+          "shape": [4, 255],
+          "dtype": "float16",
+          "seed": 695,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        }
+      },
+      "outputs": { "outputT": { "shape": [1, 512, 4, 255], "dtype": "float16" } }
+    },
+    {
+      "name": "neighbor_float16_h257_r3",
+      "preset": "all",
+      "inputs": {
+        "keyT": { "shape": [1, 1, 3, 257], "dtype": "float16", "seed": 691, "dist": "normal", "scale": 1 },
+        "queryT": { "shape": [1, 1, 3, 257], "dtype": "float16", "seed": 692, "dist": "normal", "scale": 1 },
+        "valueT": { "shape": [1, 1, 257], "dtype": "float16", "seed": 693, "dist": "normal", "scale": 1 },
+        "keyNormScaleT": {
+          "shape": [3, 257],
+          "dtype": "float16",
+          "seed": 694,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        },
+        "queryNormScaleT": {
+          "shape": [3, 257],
+          "dtype": "float16",
+          "seed": 695,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        }
+      },
+      "outputs": { "outputT": { "shape": [1, 1, 3, 257], "dtype": "float16" } }
+    },
+    {
+      "name": "neighbor_float16_h257_r2048",
+      "preset": "all",
+      "inputs": {
+        "keyT": { "shape": [1, 512, 4, 257], "dtype": "float16", "seed": 691, "dist": "normal", "scale": 1 },
+        "queryT": { "shape": [1, 512, 4, 257], "dtype": "float16", "seed": 692, "dist": "normal", "scale": 1 },
+        "valueT": { "shape": [1, 512, 257], "dtype": "float16", "seed": 693, "dist": "normal", "scale": 1 },
+        "keyNormScaleT": {
+          "shape": [4, 257],
+          "dtype": "float16",
+          "seed": 694,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        },
+        "queryNormScaleT": {
+          "shape": [4, 257],
+          "dtype": "float16",
+          "seed": 695,
+          "dist": "uniform",
+          "min": 0.2,
+          "max": 1.8
+        }
+      },
+      "outputs": { "outputT": { "shape": [1, 512, 4, 257], "dtype": "float16" } }
+    }
+  ]
+}

build/webgpu/engram-gate-rows.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,172 @@

+{% if foldKind == "subgroup_full" or foldKind == "subgroup_butterfly" %}
+enable subgroups;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+// The Engram gate, fused per (batch, sequence, hc_mult) row:
+//   dot  = sum(RMSNorm(key) * RMSNorm(query)) / sqrt(HIDDEN)
+//   out  = sigmoid(sign(dot) * sqrt(max(abs(dot), 1e-6))) * value
+// Both inverse-RMS factors are row scalars, so they come out of the sum and the
+// two sums of squares and the scaled cross term ride ONE pass as a single vec3
+// accumulator. Normalizing per element first would need two passes and a
+// different rounding.
+//
+// Each row owns a fixed LANE GROUP rather than a whole workgroup. LANES_PER_ROW
+// is the smallest power of two that covers the row, capped by the workgroup
+// (and, on the subgroup routes, by the subgroup) so several rows are resident at
+// once, and widened toward the workgroup when there are too few rows to fill the
+// dispatch. The value words stay in registers across the fold.
+//
+// Every accumulation is float32 whatever the tensor type; only the store narrows.
+const ROW_WORDS: u32 = {{ rowWords }}u;
+const HIDDEN: f32 = {{ hidden }}.0;
+const HC_MULT: u32 = {{ hcMult }}u;
+const EPS: f32 = {{ epsilon }};
+const GATE_FLOOR: f32 = 0.000001;
+const WG: u32 = {{ workgroupSize }}u;
+const LANES_PER_ROW: u32 = {{ lanesPerRow }}u;
+const ROWS_PER_GROUP: u32 = {{ rowsPerGroup }}u;
+{% if foldKind == "workgroup_tree" %}
+// One contiguous LANES_PER_ROW slice per row, carrying all three accumulators as
+// one vector. The fold never crosses a slice boundary, so the rows sharing this
+// array do not interfere.
+var<workgroup> partials: array<vec3<f32>, WG>;
+{% endif %}
+fn sigmoid_safe(x: f32) -> f32 {
+  if (x >= 0.0) {
+    let z = exp(-x);
+    return 1.0 / (1.0 + z);
+  }
+  let z = exp(x);
+  return z / (1.0 + z);
+}
+// sign(0.0) is 0.0 in WGSL, so a zero dot product gives an argument of exactly
+// zero and therefore a gate of exactly 0.5. A copysign spelling would map it to
+// +sqrt(GATE_FLOOR) instead, which is the divergence upstream calls out.
+fn engram_gate(dotValue: f32) -> f32 {
+  return sigmoid_safe(sign(dotValue) * sqrt(max(abs(dotValue), GATE_FLOOR)));
+}
+@compute @workgroup_size(WG)
+fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid: vec3<u32>) {
+  let lane = lid.x % LANES_PER_ROW;
+  let localRow = lid.x / LANES_PER_ROW;
+  // 2D-folded row index: wg.y carries the high bits past the dispatch's per-axis
+  // workgroup fold width.
+  let row = (wg.x + wg.y * {{ DISPATCH_FOLD_WIDTH }}u) * ROWS_PER_GROUP + localRow;
+  let rowValid = row < params.rows;
+  // A lane group past the last row folds an empty row and stores nothing, so the
+  // collectives and barriers below stay in uniform control flow.
+  let limit = select(0u, ROW_WORDS, rowValid);
+  let base = row * ROW_WORDS;
+  // The hyper-connection stream selects the scale rows; the token selects the one
+  // value row that every stream of that token shares.
+  let scaleBase = (row % HC_MULT) * ROW_WORDS;
+  let valueBase = (row / HC_MULT) * ROW_WORDS;
+  // x: sum(key^2), y: sum(query^2), z: sum(key * key_scale * query * query_scale).
+  var acc = vec3<f32>(0.0);
+{% if stageRow %}
+{% for i in range(wordsPerLane) %}
+{% if useVec4 %}
+  var vv{{ i }} = vec4<f32>(0.0);
+{% else %}
+  var vv{{ i }} = 0.0;
+{% endif %}
+  let c{{ i }} = lane + {{ i }}u * LANES_PER_ROW;
+  if (c{{ i }} < limit) {
+{% if useVec4 %}
+    let kv{{ i }} = vec4<f32>(key[base + c{{ i }}]);
+    let qv{{ i }} = vec4<f32>(query[base + c{{ i }}]);
+    let ks{{ i }} = vec4<f32>(kscale[scaleBase + c{{ i }}]);
+    let qs{{ i }} = vec4<f32>(qscale[scaleBase + c{{ i }}]);
+    vv{{ i }} = vec4<f32>(value[valueBase + c{{ i }}]);
+    acc = acc + vec3<f32>(dot(kv{{ i }}, kv{{ i }}), dot(qv{{ i }}, qv{{ i }}), dot(kv{{ i }} * ks{{ i }}, qv{{ i }} * qs{{ i }}));
+{% else %}
+    let kv{{ i }} = f32(key[base + c{{ i }}]);
+    let qv{{ i }} = f32(query[base + c{{ i }}]);
+    let ks{{ i }} = f32(kscale[scaleBase + c{{ i }}]);
+    let qs{{ i }} = f32(qscale[scaleBase + c{{ i }}]);
+    vv{{ i }} = f32(value[valueBase + c{{ i }}]);
+    acc = acc + vec3<f32>(kv{{ i }} * kv{{ i }}, qv{{ i }} * qv{{ i }}, (kv{{ i }} * ks{{ i }}) * (qv{{ i }} * qs{{ i }}));
+{% endif %}
+  }
+{% endfor %}
+{% else %}
+  // The row's value words do not fit in registers at this width, so the statistic
+  // pass and the store pass each walk it.
+  for (var c = lane; c < limit; c = c + LANES_PER_ROW) {
+{% if useVec4 %}
+    let kv = vec4<f32>(key[base + c]);
+    let qv = vec4<f32>(query[base + c]);
+    let ks = vec4<f32>(kscale[scaleBase + c]);
+    let qs = vec4<f32>(qscale[scaleBase + c]);
+    acc = acc + vec3<f32>(dot(kv, kv), dot(qv, qv), dot(kv * ks, qv * qs));
+{% else %}
+    let kv = f32(key[base + c]);
+    let qv = f32(query[base + c]);
+    let ks = f32(kscale[scaleBase + c]);
+    let qs = f32(qscale[scaleBase + c]);
+    acc = acc + vec3<f32>(kv * kv, qv * qv, (kv * ks) * (qv * qs));
+{% endif %}
+  }
+{% endif %}
+{% if foldKind == "subgroup_full" %}
+  let sums = subgroupAdd(acc);
+{% elif foldKind == "subgroup_butterfly" %}
+  // Segmented butterfly: only the low bits of the lane index are exchanged, so each
+  // aligned group of LANES_PER_ROW lanes folds its own row and one subgroup carries
+  // several rows without a barrier or any workgroup memory. All three accumulators
+  // ride one shuffle.
+  var folded = acc;
+{% for step in range(foldSteps) %}
+  folded = folded + subgroupShuffleXor(folded, 1u << {{ step }}u);
+{% endfor %}
+  let sums = folded;
+{% elif foldKind == "workgroup_tree" %}
+  partials[lid.x] = acc;
+  workgroupBarrier();
+  // LANES_PER_ROW is a power of two, so the halving stride is uniform across the
+  // whole workgroup and every lane reaches every barrier.
+  var stride = LANES_PER_ROW / 2u;
+  loop {
+    if (stride == 0u) {
+      break;
+    }
+    if (lane < stride) {
+      partials[lid.x] = partials[lid.x] + partials[lid.x + stride];
+    }
+    workgroupBarrier();
+    stride = stride / 2u;
+  }
+  let sums = partials[localRow * LANES_PER_ROW];
+{% else %}
+  // One lane owns the whole row, so the accumulator already holds the row sums.
+  let sums = acc;
+{% endif %}
+  let keyInv = inverseSqrt(sums.x / HIDDEN + EPS);
+  let queryInv = inverseSqrt(sums.y / HIDDEN + EPS);
+  let gate = engram_gate(sums.z * keyInv * queryInv / sqrt(HIDDEN));
+{% if stageRow %}
+{% for i in range(wordsPerLane) %}
+  if (c{{ i }} < limit) {
+{% if useVec4 %}
+    out[base + c{{ i }}] = {{ vectorScalar }}(vv{{ i }} * gate);
+{% else %}
+    out[base + c{{ i }}] = {{ scalar }}(vv{{ i }} * gate);
+{% endif %}
+  }
+{% endfor %}
+{% else %}
+  for (var c = lane; c < limit; c = c + LANES_PER_ROW) {
+{% if useVec4 %}
+    out[base + c] = {{ vectorScalar }}(vec4<f32>(value[valueBase + c]) * gate);
+{% else %}
+    out[base + c] = {{ scalar }}(f32(value[valueBase + c]) * gate);
+{% endif %}
+  }
+{% endif %}
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,196 @@

+{
+  "domain": "com.microsoft",
+  "name": "EngramGate",
+  "sinceVersion": 1,
+  "inputs": {
+    "keyT": { "onnx": "key", "dtype": "T", "rank": 4 },
+    "queryT": { "onnx": "query", "dtype": "T", "rank": 4 },
+    "valueT": { "onnx": "value", "dtype": "T", "rank": 3 },
+    "keyNormScaleT": { "onnx": "key_norm_scale", "dtype": "T", "rank": 2 },
+    "queryNormScaleT": { "onnx": "query_norm_scale", "dtype": "T", "rank": 2 }
+  },
+  "outputs": { "outputT": { "onnx": "output", "dtype": "T", "rank": 4, "shape": "shapes.keyT" } },
+  "attributes": { "epsilon": { "default": 0.00001 } },
+  "typeConstraints": { "T": ["float32", "float16"] },
+  "tunables": {
+    "WORKGROUP_SIZE": { "default": 256 },
+    "LANE_WORDS": { "default": 8 },
+    "MIN_ROW_GROUPS": { "default": 8 }
+  },
+  "derive": {
+    "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
+    "foldedDispatchCapacity": "min(device.limits.maxComputeWorkgroupsPerDimension, 65535) * min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
+    "hcMult": "dim(shapes.keyT, 2)",
+    "hidden": "dim(shapes.keyT, 3)",
+    "rowCount": "dim(shapes.keyT, 0) * dim(shapes.keyT, 1) * hcMult",
+    "dtypeOk": "tensorDtypes.keyT == tensorDtypes.queryT and tensorDtypes.keyT == tensorDtypes.valueT and tensorDtypes.keyT == tensorDtypes.keyNormScaleT and tensorDtypes.keyT == tensorDtypes.queryNormScaleT and tensorDtypes.keyT == tensorDtypes.outputT and f16Ok(dtypes.T)",
+    "shapeOk": "ranks.keyT == 4 and ranks.valueT == 3 and hidden > 0 and hcMult > 0 and rowCount > 0 and sameShape(shapes.queryT, shapes.keyT) and sameShape(shapes.outputT, shapes.keyT) and sameShape(prefix(shapes.valueT, 2), prefix(shapes.keyT, 2)) and dim(shapes.valueT, 2) == hidden and sameShape(shapes.keyNormScaleT, suffix(shapes.keyT, 2)) and sameShape(shapes.queryNormScaleT, suffix(shapes.keyT, 2))",
+    "contractOk": "dtypeOk and shapeOk",
+    "workgroupOk": "tunables.WORKGROUP_SIZE > 0 and tunables.WORKGROUP_SIZE <= deviceWorkgroupCap and pow2ceil(tunables.WORKGROUP_SIZE) == tunables.WORKGROUP_SIZE and tunables.LANE_WORDS > 0 and tunables.MIN_ROW_GROUPS > 0",
+    "sharedFoldFits": "tunables.WORKGROUP_SIZE * 16 <= device.limits.maxComputeWorkgroupStorageSize",
+    "vec4Ok": "hidden % 4 == 0",
+    "vecWords": "hidden / 4",
+    "vecPreferredLanes": "max(1, min(tunables.WORKGROUP_SIZE / 2, pow2ceil(max(1, ceilDiv(vecWords, max(1, tunables.LANE_WORDS))))))",
+    "vecWideLanes": "max(1, min(tunables.WORKGROUP_SIZE, pow2ceil(max(1, ceilDiv(vecWords, max(1, tunables.LANE_WORDS))))))",
+    "vecBaseLanes": "pick([[ceilDiv(vecWords, vecPreferredLanes) <= 8, vecPreferredLanes]], vecWideLanes)",
+    "vecBaseRowGroups": "ceilDiv(rowCount, max(1, tunables.WORKGROUP_SIZE / max(1, vecBaseLanes)))",
+    "vecFillLanes": "max(1, min(min(tunables.WORKGROUP_SIZE, pow2ceil(max(1, vecWords))), pow2ceil(max(1, ceilDiv(tunables.MIN_ROW_GROUPS * tunables.WORKGROUP_SIZE, max(1, rowCount))))))",
+    "vecLanesPerRow": "pick([[vecBaseRowGroups < tunables.MIN_ROW_GROUPS, max(vecBaseLanes, vecFillLanes)]], vecBaseLanes)",
+    "vecRowsPerGroup": "tunables.WORKGROUP_SIZE / max(1, vecLanesPerRow)",
+    "vecWordsPerLane": "ceilDiv(vecWords, max(1, vecLanesPerRow))",
+    "vecRowGroups": "ceilDiv(rowCount, max(1, vecRowsPerGroup))",
+    "scalarPreferredLanes": "max(1, min(tunables.WORKGROUP_SIZE / 2, pow2ceil(max(1, ceilDiv(hidden, max(8, tunables.LANE_WORDS))))))",
+    "scalarWideLanes": "max(1, min(tunables.WORKGROUP_SIZE, pow2ceil(max(1, ceilDiv(hidden, max(1, tunables.LANE_WORDS))))))",
+    "scalarBaseLanes": "pick([[ceilDiv(hidden, scalarPreferredLanes) <= 8, scalarPreferredLanes]], scalarWideLanes)",
+    "scalarBaseRowGroups": "ceilDiv(rowCount, max(1, tunables.WORKGROUP_SIZE / max(1, scalarBaseLanes)))",
+    "scalarFillLanes": "max(1, min(min(tunables.WORKGROUP_SIZE, pow2ceil(max(1, hidden))), pow2ceil(max(1, ceilDiv(tunables.MIN_ROW_GROUPS * tunables.WORKGROUP_SIZE, max(1, rowCount))))))",
+    "scalarLanesPerRow": "pick([[scalarBaseRowGroups < tunables.MIN_ROW_GROUPS, max(scalarBaseLanes, scalarFillLanes)]], scalarBaseLanes)",
+    "scalarRowsPerGroup": "tunables.WORKGROUP_SIZE / max(1, scalarLanesPerRow)",
+    "scalarWordsPerLane": "ceilDiv(hidden, max(1, scalarLanesPerRow))",
+    "scalarRowGroups": "ceilDiv(rowCount, max(1, scalarRowsPerGroup))",
+    "subgroupShapeKnown": "device.features.has(\"subgroups\") and has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize > 0 and device.adapterInfo.subgroupMinSize <= device.adapterInfo.subgroupMaxSize and device.adapterInfo.subgroupMaxSize <= tunables.WORKGROUP_SIZE and tunables.WORKGROUP_SIZE % device.adapterInfo.subgroupMaxSize == 0",
+    "sgMin": "device.adapterInfo.subgroupMinSize if subgroupShapeKnown else 0",
+    "sgFullWidth": "subgroupShapeKnown and device.adapterInfo.subgroupMinSize == device.adapterInfo.subgroupMaxSize",
+    "sgVecLanesPerRow": "min(sgMin, vecLanesPerRow) if subgroupShapeKnown else 1",
+    "sgVecRowsPerGroup": "tunables.WORKGROUP_SIZE / max(1, sgVecLanesPerRow)",
+    "sgVecWordsPerLane": "ceilDiv(vecWords, max(1, sgVecLanesPerRow))",
+    "sgVecRowGroups": "ceilDiv(rowCount, max(1, sgVecRowsPerGroup))",
+    "sgScalarLanesPerRow": "min(sgMin, scalarLanesPerRow) if subgroupShapeKnown else 1",
+    "sgScalarRowsPerGroup": "tunables.WORKGROUP_SIZE / max(1, sgScalarLanesPerRow)",
+    "sgScalarWordsPerLane": "ceilDiv(hidden, max(1, sgScalarLanesPerRow))",
+    "sgScalarRowGroups": "ceilDiv(rowCount, max(1, sgScalarRowsPerGroup))",
+    "vectorElement": "\"vec4<\" ~ dtypes.T ~ \">\"",
+    "workgroupSize": "tunables.WORKGROUP_SIZE",
+    "epsilon": "attrs.epsilon",
+    "scalar": "dtypes.T",
+    "vectorScalar": "vectorElement"
+  },
+  "when": ["contractOk", "workgroupOk"],
+  "bindings": {
+    "key": { "arg": "keyT", "elementType": "$vectorElement" },
+    "query": { "arg": "queryT", "elementType": "$vectorElement" },
+    "value": { "arg": "valueT", "elementType": "$vectorElement" },
+    "kscale": { "arg": "keyNormScaleT", "elementType": "$vectorElement" },
+    "qscale": { "arg": "queryNormScaleT", "elementType": "$vectorElement" },
+    "out": { "arg": "outputT", "elementType": "$vectorElement" },
+    "key_main": { "arg": "keyT", "name": "key", "elementType": "$scalar" },
+    "query_main": { "arg": "queryT", "name": "query", "elementType": "$scalar" },
+    "value_main": { "arg": "valueT", "name": "value", "elementType": "$scalar" },
+    "kscale_main": { "arg": "keyNormScaleT", "name": "kscale", "elementType": "$scalar" },
+    "qscale_main": { "arg": "queryNormScaleT", "name": "qscale", "elementType": "$scalar" },
+    "out_main": { "arg": "outputT", "name": "out", "elementType": "$scalar" },
+    "params": { "struct": [{ "name": "rows", "type": "u32", "value": "rowCount" }] }
+  },
+  "variants": [
+    {
+      "id": "subgroup_rows_vec4",
+      "priority": 30,
+      "when": ["vec4Ok", "subgroupShapeKnown", "sgVecWordsPerLane >= 1", "sgVecWordsPerLane <= 8", "sgVecRowGroups <= foldedDispatchCapacity"],
+      "requires": { "features": ["subgroups"] },
+      "derive": {
+        "foldKind": "\"none\" if sgVecLanesPerRow <= 1 else (\"subgroup_full\" if (sgFullWidth and sgVecLanesPerRow == sgMin) else \"subgroup_butterfly\")",
+        "foldSteps": "log2ceil(sgVecLanesPerRow)",
+        "useVec4": "true",
+        "stageRow": "true",
+        "rowWords": "vecWords",
+        "wordsPerLane": "sgVecWordsPerLane",
+        "lanesPerRow": "sgVecLanesPerRow",
+        "rowsPerGroup": "sgVecRowsPerGroup"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "EngramGate.SubgroupRowsVec4",
+          "shader": "engram-gate-rows.wgsl.jinja",
+          "bindings": ["key", "query", "value", "kscale", "qscale", "out", "params"],
+          "dispatch": {
+            "x": "min(sgVecRowGroups, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(sgVecRowGroups, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "packed_rows_vec4",
+      "priority": 20,
+      "when": ["vec4Ok", "sharedFoldFits", "vecRowGroups <= foldedDispatchCapacity"],
+      "derive": {
+        "foldKind": "\"none\" if vecLanesPerRow <= 1 else \"workgroup_tree\"",
+        "foldSteps": "log2ceil(vecLanesPerRow)",
+        "useVec4": "true",
+        "stageRow": "vecWordsPerLane <= 8",
+        "rowWords": "vecWords",
+        "wordsPerLane": "vecWordsPerLane",
+        "lanesPerRow": "vecLanesPerRow",
+        "rowsPerGroup": "vecRowsPerGroup"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "EngramGate.PackedRowsVec4",
+          "shader": "engram-gate-rows.wgsl.jinja",
+          "bindings": ["key", "query", "value", "kscale", "qscale", "out", "params"],
+          "dispatch": {
+            "x": "min(vecRowGroups, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(vecRowGroups, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "subgroup_rows_scalar",
+      "priority": 10,
+      "when": ["subgroupShapeKnown", "sgScalarWordsPerLane >= 1", "sgScalarWordsPerLane <= 8", "sgScalarRowGroups <= foldedDispatchCapacity"],
+      "requires": { "features": ["subgroups"] },
+      "derive": {
+        "foldKind": "\"none\" if sgScalarLanesPerRow <= 1 else (\"subgroup_full\" if (sgFullWidth and sgScalarLanesPerRow == sgMin) else \"subgroup_butterfly\")",
+        "foldSteps": "log2ceil(sgScalarLanesPerRow)",
+        "useVec4": "false",
+        "stageRow": "true",
+        "rowWords": "hidden",
+        "wordsPerLane": "sgScalarWordsPerLane",
+        "lanesPerRow": "sgScalarLanesPerRow",
+        "rowsPerGroup": "sgScalarRowsPerGroup"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "EngramGate.SubgroupRowsScalar",
+          "shader": "engram-gate-rows.wgsl.jinja",
+          "bindings": ["key_main", "query_main", "value_main", "kscale_main", "qscale_main", "out_main", "params"],
+          "dispatch": {
+            "x": "min(sgScalarRowGroups, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(sgScalarRowGroups, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "packed_rows_scalar",
+      "priority": 0,
+      "when": ["sharedFoldFits", "scalarRowGroups <= foldedDispatchCapacity"],
+      "derive": {
+        "foldKind": "\"none\" if scalarLanesPerRow <= 1 else \"workgroup_tree\"",
+        "foldSteps": "log2ceil(scalarLanesPerRow)",
+        "useVec4": "false",
+        "stageRow": "scalarWordsPerLane <= 8",
+        "rowWords": "hidden",
+        "wordsPerLane": "scalarWordsPerLane",
+        "lanesPerRow": "scalarLanesPerRow",
+        "rowsPerGroup": "scalarRowsPerGroup"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "EngramGate.PackedRowsScalar",
+          "shader": "engram-gate-rows.wgsl.jinja",
+          "bindings": ["key_main", "query_main", "value_main", "kscale_main", "qscale_main", "out_main", "params"],
+          "dispatch": {
+            "x": "min(scalarRowGroups, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(scalarRowGroups, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,26 @@

+{
+  "name": "com.microsoft.EngramGate",
+  "id": "_com_microsoft_engramgate_webgpu_e6b4c45",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "1AsNY/KLnWqSi3rI1GdriO0bQ5cfQQukYt9WisteVaE=",
+      "engram-gate-rows.wgsl.jinja": "q9Z+7afQ+9LFxjks9tQw2xX93Dsen47WaPBW26breLM=",
+      "manifest.json": "m6y59i8jPJFyMB8AbPpNvfd/t3w9MLtG5/L7rfOQBs8=",
+      "test.json": "OfAVLi7QPYg7OUihAF7HmAdC/xrIq3qziUYSMchd4DQ="
+    }
+  },
+  "provenance": { "kernel": { "sha": "6fdf6301e2bbcc2f03bf1eaf493b7ad55ef33afc", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.1",
+    "variants": {
+      "subgroup_rows_vec4": ["engram-gate-rows.wgsl.jinja"],
+      "packed_rows_vec4": ["engram-gate-rows.wgsl.jinja"],
+      "subgroup_rows_scalar": ["engram-gate-rows.wgsl.jinja"],
+      "packed_rows_scalar": ["engram-gate-rows.wgsl.jinja"]
+    }
+  }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,2320 @@

+{
+  "cases": [
+    {
+      "name": "ort_gate_float",
+      "provenance": {
+        "source": "ORT: onnxruntime/test/contrib_ops/engram_ops_test.cc",
+        "test": "EngramOpsTest.EngramGateFloat",
+        "notes": "The upstream tiny case verbatim: key {0, 2.5}, query {3, 4}, value {2, -1.5}, unit scales, one token and one hyper-connection. Expected values recomputed in float32 from the upstream kernel's own ordering."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "keyT": { "dtype": "float32", "shape": [1, 1, 1, 2], "data": { "kind": "values", "values": [0.0, 2.5] } },
+        "queryT": { "dtype": "float32", "shape": [1, 1, 1, 2], "data": { "kind": "values", "values": [3.0, 4.0] } },
+        "valueT": { "dtype": "float32", "shape": [1, 1, 2], "data": { "kind": "values", "values": [2.0, -1.5] } },
+        "keyNormScaleT": { "dtype": "float32", "shape": [1, 2], "data": { "kind": "values", "values": [1.0, 1.0] } },
+        "queryNormScaleT": { "dtype": "float32", "shape": [1, 2], "data": { "kind": "values", "values": [1.0, 1.0] } }
+      },
+      "outputs": {
+        "outputT": {
+          "dtype": "float32",
+          "shape": [1, 1, 1, 2],
+          "data": { "kind": "values", "values": [1.4867780208587646, -1.1150834560394287] },
+          "tolerance": 0.000002,
+          "relTolerance": 0.00002
+        }
+      }
+    },
+    {
+      "name": "ort_gate_float16",
+      "provenance": {
+        "source": "ORT: onnxruntime/test/contrib_ops/engram_ops_test.cc",
+        "test": "EngramOpsTest.EngramGateFloat16",
+        "notes": "The upstream tiny case verbatim: key {0, 2.5}, query {3, 4}, value {2, -1.5}, unit scales, one token and one hyper-connection. Expected values recomputed in float32 from the upstream kernel's own ordering."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "keyT": { "dtype": "float16", "shape": [1, 1, 1, 2], "data": { "kind": "values", "values": [0.0, 2.5] } },
+        "queryT": { "dtype": "float16", "shape": [1, 1, 1, 2], "data": { "kind": "values", "values": [3.0, 4.0] } },
+        "valueT": { "dtype": "float16", "shape": [1, 1, 2], "data": { "kind": "values", "values": [2.0, -1.5] } },
+        "keyNormScaleT": { "dtype": "float16", "shape": [1, 2], "data": { "kind": "values", "values": [1.0, 1.0] } },
+        "queryNormScaleT": { "dtype": "float16", "shape": [1, 2], "data": { "kind": "values", "values": [1.0, 1.0] } }
+      },
+      "outputs": {
+        "outputT": {
+          "dtype": "float16",
+          "shape": [1, 1, 1, 2],
+          "data": { "kind": "values", "values": [1.486328125, -1.115234375] },
+          "tolerance": 0.0005,
+          "relTolerance": 0.002
+        }
+      }
+    },
+    {
+      "name": "ort_zero_dot_product",
+      "provenance": {
+        "source": "ORT: onnxruntime/test/contrib_ops/engram_ops_test.cc",
+        "test": "EngramOpsTest.EngramGateZeroDotProduct",
+        "notes": "Orthogonal key and query make the dot product vanish, and the gate must be exactly 0.5 rather than sigmoid(sqrt(1e-6)). A copysign spelling of the signed square root fails this case and nothing else."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "keyT": { "dtype": "float32", "shape": [1, 1, 1, 2], "data": { "kind": "values", "values": [1.0, 0.0] } },
+        "queryT": { "dtype": "float32", "shape": [1, 1, 1, 2], "data": { "kind": "values", "values": [0.0, 1.0] } },
+        "valueT": { "dtype": "float32", "shape": [1, 1, 2], "data": { "kind": "values", "values": [1.0, -1.0] } },
+        "keyNormScaleT": { "dtype": "float32", "shape": [1, 2], "data": { "kind": "values", "values": [1.0, 1.0] } },
+        "queryNormScaleT": { "dtype": "float32", "shape": [1, 2], "data": { "kind": "values", "values": [1.0, 1.0] } }
+      },
+      "outputs": {
+        "outputT": {
+          "dtype": "float32",
+          "shape": [1, 1, 1, 2],
+          "data": { "kind": "values", "values": [0.5, -0.5] },
+          "tolerance": 0.00001,
+          "relTolerance": 0.00001
+        }
+      }
+    },
+    {
+      "name": "ort_vectorized_h4_float",
+      "provenance": {
+        "source": "ORT: onnxruntime/test/contrib_ops/engram_ops_test.cc",
+        "test": "EngramOpsTest.EngramGateVectorizedFloat",
+        "notes": "Upstream's vectorized case verbatim, including its bounded sine generator: batch 1, sequence 2, hc_mult 2, hidden 4, with distinct per-stream norm scales so a wrong scale row is visible. Expected values recomputed in float32."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "keyT": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 4],
+          "data": {
+            "kind": "values",
+            "values": [-0.7833269238471985, -0.5646424293518066, -0.295520156621933, 5.960464477539063e-8, 0.29552027583122253, 0.5646424889564514, 0.7833269834518433, 0.9320391416549683, 0.9974949955940247, 0.9738476276397705, 0.8632094264030457, 0.6754631400108337, 0.4273796081542969, 0.14112000167369843, -0.15774588286876678, -0.4425203502178192]
+          }
+        },
+        "queryT": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.9320390820503235, 0.81341552734375, 0.6442177295684814, 0.434965580701828, 0.19866937398910522, -0.04997912049293518, -0.295520156621933, -0.5226871967315674, -0.7173560857772827, -0.8674232363700867, -0.9635581374168396, -0.9997837543487549, -0.9738476276397705, -0.8873624205589294, -0.745705246925354, -0.5576837658882141]
+          }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [1, 2, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.3894183337688446, 0.681638777256012, 0.8912073373794556, 0.9927129745483398, 0.9738476276397705, 0.8368987441062927, 0.5984721779823303, 0.28747788071632385]
+          }
+        },
+        "keyNormScaleT": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.5646424889564514, 0.6442177295684814, 0.7173560857772827, 0.7833269238471985, 0.8414710164070129, 0.8912073373794556, 0.9320390820503235, 0.9635581374168396]
+          }
+        },
+        "queryNormScaleT": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.9854497313499451, 0.9489846229553223, 0.8912073373794556, 0.8134154677391052, 0.7173560857772827, 0.605186402797699, 0.47942548990249634, 0.3428977131843567]
+          }
+        }
+      },
+      "outputs": {
+        "outputT": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.10116856545209885, 0.17708569765090942, 0.23153036832809448, 0.2579009234905243, 0.1250687837600708, 0.21892069280147552, 0.28622743487358093, 0.3188278377056122, 0.2453785240650177, 0.21087177097797394, 0.15079589188098907, 0.07243525236845016, 0.3465985059738159, 0.2978575527667999, 0.21300002932548523, 0.10231518745422363]
+          },
+          "tolerance": 0.000002,
+          "relTolerance": 0.00002
+        }
+      }
+    },
+    {
+      "name": "ort_vectorized_h4_float16",
+      "provenance": {
+        "source": "ORT: onnxruntime/test/contrib_ops/engram_ops_test.cc",
+        "test": "EngramOpsTest.EngramGateVectorizedFloat16",
+        "notes": "Upstream's vectorized case verbatim, including its bounded sine generator: batch 1, sequence 2, hc_mult 2, hidden 4, with distinct per-stream norm scales so a wrong scale row is visible. Expected values recomputed in float32."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "keyT": {
+          "dtype": "float16",
+          "shape": [1, 2, 2, 4],
+          "data": {
+            "kind": "values",
+            "values": [-0.783203125, -0.564453125, -0.29541015625, 5.960464477539063e-8, 0.29541015625, 0.564453125, 0.783203125, 0.93212890625, 0.99755859375, 0.9736328125, 0.86328125, 0.67529296875, 0.427490234375, 0.14111328125, -0.15771484375, -0.442626953125]
+          }
+        },
+        "queryT": {
+          "dtype": "float16",
+          "shape": [1, 2, 2, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.93212890625, 0.8134765625, 0.64404296875, 0.43505859375, 0.1986083984375, -0.04998779296875, -0.29541015625, -0.5224609375, -0.71728515625, -0.8671875, -0.96337890625, -1.0, -0.9736328125, -0.88720703125, -0.74560546875, -0.5576171875]
+          }
+        },
+        "valueT": {
+          "dtype": "float16",
+          "shape": [1, 2, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.389404296875, 0.681640625, 0.89111328125, 0.99267578125, 0.9736328125, 0.8369140625, 0.5986328125, 0.28759765625]
+          }
+        },
+        "keyNormScaleT": {
+          "dtype": "float16",
+          "shape": [2, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.564453125, 0.64404296875, 0.71728515625, 0.783203125, 0.84130859375, 0.89111328125, 0.93212890625, 0.96337890625]
+          }
+        },
+        "queryNormScaleT": {
+          "dtype": "float16",
+          "shape": [2, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.9853515625, 0.94921875, 0.89111328125, 0.8134765625, 0.71728515625, 0.60498046875, 0.4794921875, 0.343017578125]
+          }
+        }
+      },
+      "outputs": {
+        "outputT": {
+          "dtype": "float16",
+          "shape": [1, 2, 2, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.1011962890625, 0.1771240234375, 0.2315673828125, 0.2578125, 0.125, 0.2188720703125, 0.2861328125, 0.31884765625, 0.245361328125, 0.2109375, 0.15087890625, 0.07244873046875, 0.3466796875, 0.2978515625, 0.213134765625, 0.10235595703125]
+          },
+          "tolerance": 0.0005,
+          "relTolerance": 0.002
+        }
+      }
+    },
+    {
+      "name": "ort_multi_iteration_h260_float",
+      "provenance": {
+        "notes": "Upstream's multi-iteration reduction geometry: batch 1, sequence 2, hc_mult 2, hidden 260, the smallest multiple of four above both of the upstream accumulation strides. The inputs use a bounded fill rather than upstream's sine wave, so the expectation comes from the reference."
+      },
+      "inputs": {
+        "keyT": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 260],
+          "data": { "kind": "fillFloat32", "sinStep": 0.12, "cosStep": 0.3, "scale": 1.0 }
+        },
+        "queryT": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 260],
+          "data": { "kind": "fillFloat32", "sinStep": 0.18, "cosStep": 0.34, "scale": 1.0 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [1, 2, 260],
+          "data": { "kind": "fillFloat32", "sinStep": 0.24, "cosStep": 0.39999999999999997, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "dtype": "float32",
+          "shape": [2, 260],
+          "data": { "kind": "fillFloat32", "sinStep": 0.14, "cosStep": 0.42, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "dtype": "float32",
+          "shape": [2, 260],
+          "data": { "kind": "fillFloat32", "sinStep": 0.16, "cosStep": 0.48, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 2, 2, 260], "tolerance": 0.000002, "relTolerance": 0.00002 }
+      }
+    },
+    {
+      "name": "ort_multi_iteration_h260_float16",
+      "provenance": { "notes": "The float16 half of the same multi-iteration reduction geometry." },
+      "inputs": {
+        "keyT": {
+          "dtype": "float16",
+          "shape": [1, 2, 2, 260],
+          "data": { "kind": "fillFloat32", "sinStep": 0.18000000000000002, "cosStep": 0.36, "scale": 1.0 }
+        },
+        "queryT": {
+          "dtype": "float16",
+          "shape": [1, 2, 2, 260],
+          "data": { "kind": "fillFloat32", "sinStep": 0.24000000000000002, "cosStep": 0.4, "scale": 1.0 }
+        },
+        "valueT": {
+          "dtype": "float16",
+          "shape": [1, 2, 260],
+          "data": {
+            "kind": "fillFloat32",
+            "sinStep": 0.30000000000000004,
+            "cosStep": 0.45999999999999996,
+            "scale": 1.0
+          }
+        },
+        "keyNormScaleT": {
+          "dtype": "float16",
+          "shape": [2, 260],
+          "data": { "kind": "fillFloat32", "sinStep": 0.2, "cosStep": 0.48, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "dtype": "float16",
+          "shape": [2, 260],
+          "data": {
+            "kind": "fillFloat32",
+            "sinStep": 0.22000000000000003,
+            "cosStep": 0.54,
+            "scale": 0.5,
+            "offset": 1.0
+          }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [1, 2, 2, 260], "tolerance": 0.0005, "relTolerance": 0.002 }
+      }
+    },
+    {
+      "name": "nanowhale_decode_h320_f32",
+      "provenance": {
+        "notes": "One decoded token of the DeepSeek-V4 nanowhale geometry: four hyper-connection streams of 320 channels. Four rows in the whole dispatch, the launch-bound end of the operator."
+      },
+      "inputs": {
+        "keyT": {
+          "dtype": "float32",
+          "shape": [1, 1, 4, 320],
+          "data": {
+            "kind": "fillFloat32",
+            "sinStep": 0.24000000000000002,
+            "cosStep": 0.42000000000000004,
+            "scale": 1.0
+          }
+        },
+        "queryT": {
+          "dtype": "float32",
+          "shape": [1, 1, 4, 320],
+          "data": { "kind": "fillFloat32", "sinStep": 0.30000000000000004, "cosStep": 0.46, "scale": 1.0 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [1, 1, 320],
+          "data": { "kind": "fillFloat32", "sinStep": 0.36, "cosStep": 0.52, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "dtype": "float32",
+          "shape": [4, 320],
+          "data": { "kind": "fillFloat32", "sinStep": 0.26, "cosStep": 0.54, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "dtype": "float32",
+          "shape": [4, 320],
+          "data": { "kind": "fillFloat32", "sinStep": 0.28, "cosStep": 0.6, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 1, 4, 320], "tolerance": 0.000002, "relTolerance": 0.00002 }
+      }
+    },
+    {
+      "name": "nanowhale_prefill_s5_hc3_h320_f16",
+      "provenance": {
+        "notes": "Fifteen rows of 320 at three streams: neither the row count nor the stream count divides the rows a workgroup carries, so the trailing lane groups fold an empty row."
+      },
+      "inputs": {
+        "keyT": {
+          "dtype": "float16",
+          "shape": [1, 5, 3, 320],
+          "data": { "kind": "fillFloat32", "sinStep": 0.3, "cosStep": 0.48, "scale": 1.0 }
+        },
+        "queryT": {
+          "dtype": "float16",
+          "shape": [1, 5, 3, 320],
+          "data": { "kind": "fillFloat32", "sinStep": 0.36, "cosStep": 0.52, "scale": 1.0 }
+        },
+        "valueT": {
+          "dtype": "float16",
+          "shape": [1, 5, 320],
+          "data": { "kind": "fillFloat32", "sinStep": 0.42, "cosStep": 0.58, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "dtype": "float16",
+          "shape": [3, 320],
+          "data": { "kind": "fillFloat32", "sinStep": 0.31999999999999995, "cosStep": 0.6, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "dtype": "float16",
+          "shape": [3, 320],
+          "data": {
+            "kind": "fillFloat32",
+            "sinStep": 0.33999999999999997,
+            "cosStep": 0.6599999999999999,
+            "scale": 0.5,
+            "offset": 1.0
+          }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [1, 5, 3, 320], "tolerance": 0.0005, "relTolerance": 0.002 }
+      }
+    },
+    {
+      "name": "v4flash_decode_h4096_f32",
+      "provenance": {
+        "notes": "One decoded token of the DeepSeek-V4 v4flash geometry: four streams of 4096 channels, the widest row that still stages in registers."
+      },
+      "inputs": {
+        "keyT": {
+          "dtype": "float32",
+          "shape": [1, 1, 4, 4096],
+          "data": { "kind": "fillFloat32", "sinStep": 0.32, "cosStep": 0.5, "scale": 1.0 }
+        },
+        "queryT": {
+          "dtype": "float32",
+          "shape": [1, 1, 4, 4096],
+          "data": { "kind": "fillFloat32", "sinStep": 0.38, "cosStep": 0.54, "scale": 1.0 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [1, 1, 4096],
+          "data": { "kind": "fillFloat32", "sinStep": 0.44, "cosStep": 0.6, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "dtype": "float32",
+          "shape": [4, 4096],
+          "data": {
+            "kind": "fillFloat32",
+            "sinStep": 0.33999999999999997,
+            "cosStep": 0.62,
+            "scale": 0.5,
+            "offset": 1.0
+          }
+        },
+        "queryNormScaleT": {
+          "dtype": "float32",
+          "shape": [4, 4096],
+          "data": { "kind": "fillFloat32", "sinStep": 0.36, "cosStep": 0.6799999999999999, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 1, 4, 4096], "tolerance": 0.000002, "relTolerance": 0.00002 }
+      }
+    },
+    {
+      "name": "v4flash_decode_h4096_f16",
+      "provenance": { "notes": "The float16 half of the v4flash decode geometry." },
+      "inputs": {
+        "keyT": {
+          "dtype": "float16",
+          "shape": [1, 1, 4, 4096],
+          "data": { "kind": "fillFloat32", "sinStep": 0.38, "cosStep": 0.56, "scale": 1.0 }
+        },
+        "queryT": {
+          "dtype": "float16",
+          "shape": [1, 1, 4, 4096],
+          "data": { "kind": "fillFloat32", "sinStep": 0.44, "cosStep": 0.6, "scale": 1.0 }
+        },
+        "valueT": {
+          "dtype": "float16",
+          "shape": [1, 1, 4096],
+          "data": { "kind": "fillFloat32", "sinStep": 0.5, "cosStep": 0.6599999999999999, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "dtype": "float16",
+          "shape": [4, 4096],
+          "data": { "kind": "fillFloat32", "sinStep": 0.4, "cosStep": 0.6799999999999999, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "dtype": "float16",
+          "shape": [4, 4096],
+          "data": { "kind": "fillFloat32", "sinStep": 0.42, "cosStep": 0.74, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [1, 1, 4, 4096], "tolerance": 0.0005, "relTolerance": 0.002 }
+      }
+    },
+    {
+      "name": "batched_hc2_h64_f32",
+      "provenance": {
+        "notes": "Batch 3, sequence 5, two streams of 64 channels: a row is far narrower than a subgroup here, so several rows fold inside one subgroup and a wrong batch, token or stream index moves the gate."
+      },
+      "inputs": {
+        "keyT": {
+          "dtype": "float32",
+          "shape": [3, 5, 2, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.42, "cosStep": 0.6, "scale": 1.0 }
+        },
+        "queryT": {
+          "dtype": "float32",
+          "shape": [3, 5, 2, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.64, "scale": 1.0 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [3, 5, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.54, "cosStep": 0.7, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "dtype": "float32",
+          "shape": [2, 64],
+          "data": {
+            "kind": "fillFloat32",
+            "sinStep": 0.43999999999999995,
+            "cosStep": 0.72,
+            "scale": 0.5,
+            "offset": 1.0
+          }
+        },
+        "queryNormScaleT": {
+          "dtype": "float32",
+          "shape": [2, 64],
+          "data": {
+            "kind": "fillFloat32",
+            "sinStep": 0.45999999999999996,
+            "cosStep": 0.78,
+            "scale": 0.5,
+            "offset": 1.0
+          }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [3, 5, 2, 64], "tolerance": 0.000002, "relTolerance": 0.00002 }
+      }
+    },
+    {
+      "name": "batched_hc3_h128_f16",
+      "provenance": { "notes": "Batch 2, sequence 3, three streams of 128 channels in float16." },
+      "inputs": {
+        "keyT": {
+          "dtype": "float16",
+          "shape": [2, 3, 3, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.44, "cosStep": 0.62, "scale": 1.0 }
+        },
+        "queryT": {
+          "dtype": "float16",
+          "shape": [2, 3, 3, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.5, "cosStep": 0.66, "scale": 1.0 }
+        },
+        "valueT": {
+          "dtype": "float16",
+          "shape": [2, 3, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.56, "cosStep": 0.72, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "dtype": "float16",
+          "shape": [3, 128],
+          "data": {
+            "kind": "fillFloat32",
+            "sinStep": 0.45999999999999996,
+            "cosStep": 0.74,
+            "scale": 0.5,
+            "offset": 1.0
+          }
+        },
+        "queryNormScaleT": {
+          "dtype": "float16",
+          "shape": [3, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.8, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 3, 3, 128], "tolerance": 0.0005, "relTolerance": 0.002 }
+      }
+    },
+    {
+      "name": "hc_one_h512_f32",
+      "provenance": {
+        "notes": "A single hyper-connection stream: the degenerate case where the per-stream scale lookup and a shared scale agree, kept because upstream's own smallest tests are shaped this way."
+      },
+      "inputs": {
+        "keyT": {
+          "dtype": "float32",
+          "shape": [2, 2, 1, 512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.6599999999999999, "scale": 1.0 }
+        },
+        "queryT": {
+          "dtype": "float32",
+          "shape": [2, 2, 1, 512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.54, "cosStep": 0.7, "scale": 1.0 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [2, 2, 512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.6, "cosStep": 0.76, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "dtype": "float32",
+          "shape": [1, 512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.5, "cosStep": 0.78, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "dtype": "float32",
+          "shape": [1, 512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.52, "cosStep": 0.84, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 2, 1, 512], "tolerance": 0.000002, "relTolerance": 0.00002 }
+      }
+    },
+    {
+      "name": "wide_h8192_f16",
+      "provenance": {
+        "notes": "8192 channels per row: the row fills a whole workgroup and every lane stages the maximum eight vector words."
+      },
+      "inputs": {
+        "keyT": {
+          "dtype": "float16",
+          "shape": [1, 1, 2, 8192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.54, "cosStep": 0.72, "scale": 1.0 }
+        },
+        "queryT": {
+          "dtype": "float16",
+          "shape": [1, 1, 2, 8192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.6000000000000001, "cosStep": 0.76, "scale": 1.0 }
+        },
+        "valueT": {
+          "dtype": "float16",
+          "shape": [1, 1, 8192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.66, "cosStep": 0.8200000000000001, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "dtype": "float16",
+          "shape": [2, 8192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.56, "cosStep": 0.8400000000000001, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "dtype": "float16",
+          "shape": [2, 8192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.5800000000000001, "cosStep": 0.9, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [1, 1, 2, 8192], "tolerance": 0.0005, "relTolerance": 0.002 }
+      }
+    },
+    {
+      "name": "wide_h16384_f32",
+      "provenance": {
+        "notes": "16384 channels per row: too wide to stage, so the statistic pass and the store pass each walk the row."
+      },
+      "inputs": {
+        "keyT": {
+          "dtype": "float32",
+          "shape": [1, 1, 2, 16384],
+          "data": { "kind": "fillFloat32", "sinStep": 0.6, "cosStep": 0.78, "scale": 1.0 }
+        },
+        "queryT": {
+          "dtype": "float32",
+          "shape": [1, 1, 2, 16384],
+          "data": { "kind": "fillFloat32", "sinStep": 0.6599999999999999, "cosStep": 0.82, "scale": 1.0 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [1, 1, 16384],
+          "data": { "kind": "fillFloat32", "sinStep": 0.72, "cosStep": 0.8799999999999999, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "dtype": "float32",
+          "shape": [2, 16384],
+          "data": { "kind": "fillFloat32", "sinStep": 0.62, "cosStep": 0.8999999999999999, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "dtype": "float32",
+          "shape": [2, 16384],
+          "data": { "kind": "fillFloat32", "sinStep": 0.64, "cosStep": 0.96, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 1, 2, 16384], "tolerance": 0.000002, "relTolerance": 0.00005 }
+      }
+    },
+    {
+      "name": "scalar_h6_f32",
+      "provenance": {
+        "notes": "Six channels per row: not a multiple of four, so the scalar route runs, and narrow enough that one lane group covers three rows of a subgroup."
+      },
+      "inputs": {
+        "keyT": {
+          "dtype": "float32",
+          "shape": [1, 5, 3, 6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.62, "cosStep": 0.8, "scale": 1.0 }
+        },
+        "queryT": {
+          "dtype": "float32",
+          "shape": [1, 5, 3, 6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.6799999999999999, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [1, 5, 6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.74, "cosStep": 0.8999999999999999, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "dtype": "float32",
+          "shape": [3, 6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.64, "cosStep": 0.9199999999999999, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "dtype": "float32",
+          "shape": [3, 6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.66, "cosStep": 0.98, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 5, 3, 6], "tolerance": 0.000002, "relTolerance": 0.00002 }
+      }
+    },
+    {
+      "name": "scalar_h130_f16",
+      "provenance": { "notes": "130 channels per row in float16: the scalar route with several words per lane." },
+      "inputs": {
+        "keyT": {
+          "dtype": "float16",
+          "shape": [2, 3, 2, 130],
+          "data": { "kind": "fillFloat32", "sinStep": 0.68, "cosStep": 0.8600000000000001, "scale": 1.0 }
+        },
+        "queryT": {
+          "dtype": "float16",
+          "shape": [2, 3, 2, 130],
+          "data": { "kind": "fillFloat32", "sinStep": 0.74, "cosStep": 0.9, "scale": 1.0 }
+        },
+        "valueT": {
+          "dtype": "float16",
+          "shape": [2, 3, 130],
+          "data": { "kind": "fillFloat32", "sinStep": 0.8, "cosStep": 0.96, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "dtype": "float16",
+          "shape": [2, 130],
+          "data": { "kind": "fillFloat32", "sinStep": 0.7000000000000001, "cosStep": 0.98, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "dtype": "float16",
+          "shape": [2, 130],
+          "data": { "kind": "fillFloat32", "sinStep": 0.7200000000000001, "cosStep": 1.04, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 3, 2, 130], "tolerance": 0.0005, "relTolerance": 0.002 }
+      }
+    },
+    {
+      "name": "scalar_h1_f32",
+      "provenance": {
+        "notes": "One channel per row: a single lane owns the whole row, so no fold runs at all and the statistic is the lane's own accumulator."
+      },
+      "inputs": {
+        "keyT": {
+          "dtype": "float32",
+          "shape": [1, 4, 2, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.72, "cosStep": 0.8999999999999999, "scale": 1.0 }
+        },
+        "queryT": {
+          "dtype": "float32",
+          "shape": [1, 4, 2, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.78, "cosStep": 0.94, "scale": 1.0 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [1, 4, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.84, "cosStep": 1.0, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "dtype": "float32",
+          "shape": [2, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.74, "cosStep": 1.02, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "dtype": "float32",
+          "shape": [2, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.76, "cosStep": 1.08, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 4, 2, 1], "tolerance": 0.000002, "relTolerance": 0.00002 }
+      }
+    },
+    {
+      "name": "scalar_wide_h2050_f32",
+      "provenance": {
+        "notes": "2050 channels per row: an unvectorizable width too wide to stage, the streaming scalar route."
+      },
+      "inputs": {
+        "keyT": {
+          "dtype": "float32",
+          "shape": [1, 1, 2, 2050],
+          "data": { "kind": "fillFloat32", "sinStep": 0.74, "cosStep": 0.9199999999999999, "scale": 1.0 }
+        },
+        "queryT": {
+          "dtype": "float32",
+          "shape": [1, 1, 2, 2050],
+          "data": { "kind": "fillFloat32", "sinStep": 0.8, "cosStep": 0.96, "scale": 1.0 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [1, 1, 2050],
+          "data": { "kind": "fillFloat32", "sinStep": 0.86, "cosStep": 1.02, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "dtype": "float32",
+          "shape": [2, 2050],
+          "data": { "kind": "fillFloat32", "sinStep": 0.76, "cosStep": 1.04, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "dtype": "float32",
+          "shape": [2, 2050],
+          "data": { "kind": "fillFloat32", "sinStep": 0.78, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 1, 2, 2050], "tolerance": 0.000002, "relTolerance": 0.00005 }
+      }
+    },
+    {
+      "name": "epsilon_dominant_h8_f32",
+      "provenance": {
+        "notes": "An epsilon of 1 against key and query rows two orders of magnitude smaller: the epsilon term, not the mean of squares, sets both inverse-RMS factors, which pins epsilon to the normalization denominators and away from the gate's own floor."
+      },
+      "attrs": { "epsilon": 1 },
+      "inputs": {
+        "keyT": {
+          "dtype": "float32",
+          "shape": [1, 3, 2, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.8, "cosStep": 0.98, "scale": 0.01 }
+        },
+        "queryT": {
+          "dtype": "float32",
+          "shape": [1, 3, 2, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.8600000000000001, "cosStep": 1.02, "scale": 0.01 }
+        },
+        "valueT": {
+          "dtype": "float32",
+          "shape": [1, 3, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.92, "cosStep": 1.08, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "dtype": "float32",
+          "shape": [2, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.8200000000000001, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "dtype": "float32",
+          "shape": [2, 8],
+          "data": {
+            "kind": "fillFloat32",
+            "sinStep": 0.8400000000000001,
+            "cosStep": 1.1600000000000001,
+            "scale": 0.5,
+            "offset": 1.0
+          }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 3, 2, 8], "tolerance": 0.000002, "relTolerance": 0.00002 }
+      }
+    },
+    {
+      "name": "packing_float32_h63_r3",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 1, 3, 63],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 1, 3, 63],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 1, 63],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [3, 63],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [3, 63],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 1, 3, 63], "dtype": "float32", "tolerance": 0.000002, "relTolerance": 0.00002 }
+      }
+    },
+    {
+      "name": "packing_float32_h63_r260",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 65, 4, 63],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 65, 4, 63],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 65, 63],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [4, 63],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [4, 63],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 65, 4, 63], "dtype": "float32", "tolerance": 0.000002, "relTolerance": 0.00002 }
+      }
+    },
+    {
+      "name": "packing_float32_h65_r3",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 1, 3, 65],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 1, 3, 65],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 1, 65],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [3, 65],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [3, 65],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 1, 3, 65], "dtype": "float32", "tolerance": 0.000002, "relTolerance": 0.00002 }
+      }
+    },
+    {
+      "name": "packing_float32_h65_r260",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 65, 4, 65],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 65, 4, 65],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 65, 65],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [4, 65],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [4, 65],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 65, 4, 65], "dtype": "float32", "tolerance": 0.000002, "relTolerance": 0.00002 }
+      }
+    },
+    {
+      "name": "packing_float32_h129_r3",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 1, 3, 129],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 1, 3, 129],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 1, 129],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [3, 129],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [3, 129],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 1, 3, 129], "dtype": "float32", "tolerance": 0.000002, "relTolerance": 0.00002 }
+      }
+    },
+    {
+      "name": "packing_float32_h129_r260",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 65, 4, 129],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 65, 4, 129],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 65, 129],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [4, 129],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [4, 129],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 65, 4, 129], "dtype": "float32", "tolerance": 0.000002, "relTolerance": 0.00002 }
+      }
+    },
+    {
+      "name": "packing_float32_h130_r3",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 1, 3, 130],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 1, 3, 130],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 1, 130],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [3, 130],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [3, 130],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 1, 3, 130], "dtype": "float32", "tolerance": 0.000002, "relTolerance": 0.00002 }
+      }
+    },
+    {
+      "name": "packing_float32_h130_r260",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 65, 4, 130],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 65, 4, 130],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 65, 130],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [4, 130],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [4, 130],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 65, 4, 130], "dtype": "float32", "tolerance": 0.000002, "relTolerance": 0.00002 }
+      }
+    },
+    {
+      "name": "packing_float32_h131_r3",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 1, 3, 131],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 1, 3, 131],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 1, 131],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [3, 131],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [3, 131],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 1, 3, 131], "dtype": "float32", "tolerance": 0.000002, "relTolerance": 0.00002 }
+      }
+    },
+    {
+      "name": "packing_float32_h131_r260",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 65, 4, 131],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 65, 4, 131],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 65, 131],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [4, 131],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [4, 131],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 65, 4, 131], "dtype": "float32", "tolerance": 0.000002, "relTolerance": 0.00002 }
+      }
+    },
+    {
+      "name": "packing_float32_h255_r3",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 1, 3, 255],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 1, 3, 255],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 1, 255],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [3, 255],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [3, 255],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 1, 3, 255], "dtype": "float32", "tolerance": 0.000002, "relTolerance": 0.00002 }
+      }
+    },
+    {
+      "name": "packing_float32_h255_r260",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 65, 4, 255],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 65, 4, 255],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 65, 255],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [4, 255],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [4, 255],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 65, 4, 255], "dtype": "float32", "tolerance": 0.000002, "relTolerance": 0.00002 }
+      }
+    },
+    {
+      "name": "packing_float32_h257_r3",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 1, 3, 257],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 1, 3, 257],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 1, 257],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [3, 257],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [3, 257],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 1, 3, 257], "dtype": "float32", "tolerance": 0.000002, "relTolerance": 0.00002 }
+      }
+    },
+    {
+      "name": "packing_float32_h257_r260",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 65, 4, 257],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 65, 4, 257],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 65, 257],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [4, 257],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [4, 257],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 65, 4, 257], "dtype": "float32", "tolerance": 0.000002, "relTolerance": 0.00002 }
+      }
+    },
+    {
+      "name": "packing_float16_h63_r3",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 1, 3, 63],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 1, 3, 63],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 1, 63],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [3, 63],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [3, 63],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 1, 3, 63], "dtype": "float16", "tolerance": 0.0005, "relTolerance": 0.002 }
+      }
+    },
+    {
+      "name": "packing_float16_h63_r260",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 65, 4, 63],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 65, 4, 63],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 65, 63],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [4, 63],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [4, 63],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 65, 4, 63], "dtype": "float16", "tolerance": 0.0005, "relTolerance": 0.002 }
+      }
+    },
+    {
+      "name": "packing_float16_h65_r3",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 1, 3, 65],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 1, 3, 65],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 1, 65],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [3, 65],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [3, 65],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 1, 3, 65], "dtype": "float16", "tolerance": 0.0005, "relTolerance": 0.002 }
+      }
+    },
+    {
+      "name": "packing_float16_h65_r260",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 65, 4, 65],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 65, 4, 65],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 65, 65],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [4, 65],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [4, 65],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 65, 4, 65], "dtype": "float16", "tolerance": 0.0005, "relTolerance": 0.002 }
+      }
+    },
+    {
+      "name": "packing_float16_h129_r3",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 1, 3, 129],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 1, 3, 129],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 1, 129],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [3, 129],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [3, 129],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 1, 3, 129], "dtype": "float16", "tolerance": 0.0005, "relTolerance": 0.002 }
+      }
+    },
+    {
+      "name": "packing_float16_h129_r260",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 65, 4, 129],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 65, 4, 129],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 65, 129],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [4, 129],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [4, 129],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 65, 4, 129], "dtype": "float16", "tolerance": 0.0005, "relTolerance": 0.002 }
+      }
+    },
+    {
+      "name": "packing_float16_h130_r3",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 1, 3, 130],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 1, 3, 130],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 1, 130],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [3, 130],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [3, 130],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 1, 3, 130], "dtype": "float16", "tolerance": 0.0005, "relTolerance": 0.002 }
+      }
+    },
+    {
+      "name": "packing_float16_h130_r260",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 65, 4, 130],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 65, 4, 130],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 65, 130],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [4, 130],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [4, 130],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 65, 4, 130], "dtype": "float16", "tolerance": 0.0005, "relTolerance": 0.002 }
+      }
+    },
+    {
+      "name": "packing_float16_h131_r3",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 1, 3, 131],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 1, 3, 131],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 1, 131],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [3, 131],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [3, 131],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 1, 3, 131], "dtype": "float16", "tolerance": 0.0005, "relTolerance": 0.002 }
+      }
+    },
+    {
+      "name": "packing_float16_h131_r260",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 65, 4, 131],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 65, 4, 131],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 65, 131],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [4, 131],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [4, 131],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 65, 4, 131], "dtype": "float16", "tolerance": 0.0005, "relTolerance": 0.002 }
+      }
+    },
+    {
+      "name": "packing_float16_h255_r3",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 1, 3, 255],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 1, 3, 255],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 1, 255],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [3, 255],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [3, 255],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 1, 3, 255], "dtype": "float16", "tolerance": 0.0005, "relTolerance": 0.002 }
+      }
+    },
+    {
+      "name": "packing_float16_h255_r260",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 65, 4, 255],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 65, 4, 255],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 65, 255],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [4, 255],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [4, 255],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 65, 4, 255], "dtype": "float16", "tolerance": 0.0005, "relTolerance": 0.002 }
+      }
+    },
+    {
+      "name": "packing_float16_h257_r3",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 1, 3, 257],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 1, 3, 257],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 1, 257],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [3, 257],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [3, 257],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 1, 3, 257], "dtype": "float16", "tolerance": 0.0005, "relTolerance": 0.002 }
+      }
+    },
+    {
+      "name": "packing_float16_h257_r260",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 65, 4, 257],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 65, 4, 257],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 65, 257],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [4, 257],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [4, 257],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 65, 4, 257], "dtype": "float16", "tolerance": 0.0005, "relTolerance": 0.002 }
+      }
+    },
+    {
+      "name": "packing_override_float32_h65_wg64_depth1",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 65, 4, 65],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 65, 4, 65],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 65, 65],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [4, 65],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [4, 65],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 65, 4, 65], "dtype": "float32", "tolerance": 0.000002, "relTolerance": 0.00002 }
+      },
+      "tunables": { "WORKGROUP_SIZE": 64, "LANE_WORDS": 1 }
+    },
+    {
+      "name": "packing_override_float32_h131_wg128_depth16",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 65, 4, 131],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 65, 4, 131],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 65, 131],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [4, 131],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [4, 131],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 65, 4, 131], "dtype": "float32", "tolerance": 0.000002, "relTolerance": 0.00002 }
+      },
+      "tunables": { "WORKGROUP_SIZE": 128, "LANE_WORDS": 16 }
+    },
+    {
+      "name": "packing_override_float16_h65_wg64_depth1",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 65, 4, 65],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 65, 4, 65],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 65, 65],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [4, 65],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [4, 65],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 65, 4, 65], "dtype": "float16", "tolerance": 0.0005, "relTolerance": 0.002 }
+      },
+      "tunables": { "WORKGROUP_SIZE": 64, "LANE_WORDS": 1 }
+    },
+    {
+      "name": "packing_override_float16_h131_wg128_depth16",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 65, 4, 131],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 65, 4, 131],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 65, 131],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [4, 131],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [4, 131],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 65, 4, 131], "dtype": "float16", "tolerance": 0.0005, "relTolerance": 0.002 }
+      },
+      "tunables": { "WORKGROUP_SIZE": 128, "LANE_WORDS": 16 }
+    },
+    {
+      "name": "vector_depth_float32_h252_r195_wg256",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 65, 3, 252],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 65, 3, 252],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 65, 252],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [3, 252],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [3, 252],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 65, 3, 252], "dtype": "float32", "tolerance": 0.000002, "relTolerance": 0.00002 }
+      },
+      "provenance": {
+        "notes": "Lane-depth boundary with nonuniform affine scales and hyper-connection value reuse; row and vector tails exercise shared row groups at the default dtype tolerance."
+      }
+    },
+    {
+      "name": "vector_depth_float32_h260_r195_wg256",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 65, 3, 260],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 65, 3, 260],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 65, 260],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [3, 260],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [3, 260],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 65, 3, 260], "dtype": "float32", "tolerance": 0.000002, "relTolerance": 0.00002 }
+      },
+      "provenance": {
+        "notes": "Lane-depth boundary with nonuniform affine scales and hyper-connection value reuse; row and vector tails exercise shared row groups at the default dtype tolerance."
+      }
+    },
+    {
+      "name": "vector_depth_float32_h320_r129_wg256",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 43, 3, 320],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 43, 3, 320],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 43, 320],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [3, 320],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [3, 320],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 43, 3, 320], "dtype": "float32", "tolerance": 0.000002, "relTolerance": 0.00002 }
+      },
+      "provenance": {
+        "notes": "Lane-depth boundary with nonuniform affine scales and hyper-connection value reuse; row and vector tails exercise shared row groups at the default dtype tolerance."
+      }
+    },
+    {
+      "name": "vector_depth_float32_h516_r129_wg256",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 43, 3, 516],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 43, 3, 516],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 43, 516],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [3, 516],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [3, 516],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 43, 3, 516], "dtype": "float32", "tolerance": 0.000002, "relTolerance": 0.00002 }
+      },
+      "provenance": {
+        "notes": "Lane-depth boundary with nonuniform affine scales and hyper-connection value reuse; row and vector tails exercise shared row groups at the default dtype tolerance."
+      }
+    },
+    {
+      "name": "vector_depth_float32_h1028_r195_wg64",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 65, 3, 1028],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 65, 3, 1028],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 65, 1028],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [3, 1028],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [3, 1028],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 65, 3, 1028], "dtype": "float32", "tolerance": 0.000002, "relTolerance": 0.00002 }
+      },
+      "tunables": { "WORKGROUP_SIZE": 64 },
+      "provenance": {
+        "notes": "Lane-depth boundary with nonuniform affine scales and hyper-connection value reuse; row and vector tails exercise shared row groups at the default dtype tolerance."
+      }
+    },
+    {
+      "name": "vector_depth_float32_h320_r3_wg256",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 1, 3, 320],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 1, 3, 320],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 1, 320],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [3, 320],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [3, 320],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 1, 3, 320], "dtype": "float32", "tolerance": 0.000002, "relTolerance": 0.00002 }
+      },
+      "provenance": {
+        "notes": "Lane-depth boundary with nonuniform affine scales and hyper-connection value reuse; row and vector tails exercise shared row groups at the default dtype tolerance."
+      }
+    },
+    {
+      "name": "vector_depth_float16_h252_r195_wg256",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 65, 3, 252],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 65, 3, 252],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 65, 252],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [3, 252],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [3, 252],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 65, 3, 252], "dtype": "float16", "tolerance": 0.0005, "relTolerance": 0.002 }
+      },
+      "provenance": {
+        "notes": "Lane-depth boundary with nonuniform affine scales and hyper-connection value reuse; row and vector tails exercise shared row groups at the default dtype tolerance."
+      }
+    },
+    {
+      "name": "vector_depth_float16_h260_r195_wg256",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 65, 3, 260],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 65, 3, 260],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 65, 260],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [3, 260],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [3, 260],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 65, 3, 260], "dtype": "float16", "tolerance": 0.0005, "relTolerance": 0.002 }
+      },
+      "provenance": {
+        "notes": "Lane-depth boundary with nonuniform affine scales and hyper-connection value reuse; row and vector tails exercise shared row groups at the default dtype tolerance."
+      }
+    },
+    {
+      "name": "vector_depth_float16_h320_r129_wg256",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 43, 3, 320],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 43, 3, 320],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 43, 320],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [3, 320],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [3, 320],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 43, 3, 320], "dtype": "float16", "tolerance": 0.0005, "relTolerance": 0.002 }
+      },
+      "provenance": {
+        "notes": "Lane-depth boundary with nonuniform affine scales and hyper-connection value reuse; row and vector tails exercise shared row groups at the default dtype tolerance."
+      }
+    },
+    {
+      "name": "vector_depth_float16_h516_r129_wg256",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 43, 3, 516],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 43, 3, 516],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 43, 516],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [3, 516],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [3, 516],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 43, 3, 516], "dtype": "float16", "tolerance": 0.0005, "relTolerance": 0.002 }
+      },
+      "provenance": {
+        "notes": "Lane-depth boundary with nonuniform affine scales and hyper-connection value reuse; row and vector tails exercise shared row groups at the default dtype tolerance."
+      }
+    },
+    {
+      "name": "vector_depth_float16_h1028_r195_wg64",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 65, 3, 1028],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 65, 3, 1028],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 65, 1028],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [3, 1028],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [3, 1028],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 65, 3, 1028], "dtype": "float16", "tolerance": 0.0005, "relTolerance": 0.002 }
+      },
+      "tunables": { "WORKGROUP_SIZE": 64 },
+      "provenance": {
+        "notes": "Lane-depth boundary with nonuniform affine scales and hyper-connection value reuse; row and vector tails exercise shared row groups at the default dtype tolerance."
+      }
+    },
+    {
+      "name": "vector_depth_float16_h320_r3_wg256",
+      "inputs": {
+        "keyT": {
+          "shape": [1, 1, 3, 320],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.71, "scale": 1.0 }
+        },
+        "queryT": {
+          "shape": [1, 1, 3, 320],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.84, "scale": 1.0 }
+        },
+        "valueT": {
+          "shape": [1, 1, 320],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.97, "scale": 1.0 }
+        },
+        "keyNormScaleT": {
+          "shape": [3, 320],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.7, "cosStep": 1.1, "scale": 0.5, "offset": 1.0 }
+        },
+        "queryNormScaleT": {
+          "shape": [3, 320],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.81, "cosStep": 1.23, "scale": 0.5, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 1, 3, 320], "dtype": "float16", "tolerance": 0.0005, "relTolerance": 0.002 }
+      },
+      "provenance": {
+        "notes": "Lane-depth boundary with nonuniform affine scales and hyper-connection value reuse; row and vector tails exercise shared row groups at the default dtype tolerance."
+      }
+    }
+  ]
+}