diff --git a/CLAUDE.md b/CLAUDE.md
index 781c6698..8966e86d 100644
--- a/CLAUDE.md
+++ b/CLAUDE.md
@@ -69,6 +69,7 @@ implementation of the same op on that workload.
 | Device time | The compared quantity is `device_busy_ms`, never wall-clock span. |
 | Two questions | `Ratio` says whether someone else's kernel is faster; `SOL` how much faster the hardware allows anyone to go, with the binding resource (`mem`/`comp`/`lat`) in its own `Bound` column. The SOL arithmetic and thresholds are imported from the checkout's roofline tool (M5) — never re-derived here. |
 | Order follows the API Reference | `DATA_PAGES` takes the API nav's order over the same families — a page per family except `Conv & Pool` (two) and `Other` (Top-k, FFT, mHC, Engram, the rest). Within a page, ops sit in the order `docs/api/` names them, read by `api_op_order()`; an op no API page names comes last, ranked by verdict. `_BENCH_ORDER` in `hooks.py` repeats the page order for the nav: change one, change the other. |
+| Rows follow the manifest | One row group per manifest label, one row per dtype under it in a `dtype` column. Labels keep the order the snapshot lists them in, which is the manifest's; the key above the table lists them in the same order. A row no manifest describes takes its id, trailing dtype names split off, as its label. |
 | Workload shapes | The snapshot names a workload but does not carry its shapes. `scripts/workload_shape.py` reads them from the TileOPs spec manifest at the commit the benchmark ran on, joined by the `<label>-<dtype>` the benchmark id is built from. A workload the manifest does not declare keeps its id and gets no shapes — never a guessed one. |
 
 ## Bilingual pages (en / zh)
diff --git a/docs/api/index.md b/docs/api/index.md
index a42e169b..045c76ac 100644
--- a/docs/api/index.md
+++ b/docs/api/index.md
@@ -17,7 +17,8 @@ The pages are ordered by how much an op composes: the pointwise transforms first
 the axis reductions and the normalizations built on them, then the matmul and the
 expert routing over it, then the windowed and spectral transforms, then the
 sequence-model kernels built on all of the above. It is the order `tileops` declares
-its op families in.
+its op families in. The exception is Top-k, whose one op is exported from
+`tileops.attention`.
 
 | Page | What it covers |
 | --- | --- |
diff --git a/docs/assets/extra.css b/docs/assets/extra.css
index 62d8f88f..b323dcd1 100644
--- a/docs/assets/extra.css
+++ b/docs/assets/extra.css
@@ -634,7 +634,11 @@ html[lang="zh"] .md-typeset .keystone {
   text-align: center;
 }
 
-.md-typeset .datatable table:not([class]) td.colsep {
+/* The label spans its dtype rows, and wraps at a hyphen rather than widening
+ * the column: an attention label runs past seventy characters.
+ */
+.md-typeset .datatable table:not([class]) td.wl-name {
+  white-space: normal;
   text-align: left;
 }
 
@@ -669,8 +673,8 @@ html[lang="zh"] .md-typeset .keystone {
 }
 
 /* Two header rows, still one middle rule: it belongs under the whole header
- * block, not between its halves. The `Workload` cell spans both rows, so its
- * own bottom edge already sits on that line.
+ * block, not between its halves. The `Workload` and `dtype` cells span both
+ * rows, so their own bottom edges already sit on that line.
  */
 /* Each header word underlined thin and grey — under `Alternatives` that also
  * says how far a header spanning two columns reaches. The heavy rule stays
@@ -698,9 +702,15 @@ html[lang="zh"] .md-typeset .keystone {
 
 /* A cool cast on the names, so the teal rule divides a label column from ink
  * measurements rather than sitting between them. The same in every row: a
- * colour that never varies cannot be mistaken for a verdict.
- */
-.md-typeset .datatable table:not([class]) td.colsep code {
+ * colour that never varies cannot be mistaken for a verdict. As wide as the
+ * label up to 18em and no narrower: a table squeezed to the page would
+ * otherwise break the label at every hyphen.
+ */
+.md-typeset .datatable table:not([class]) td.wl-name code {
+  display: inline-block;
+  width: max-content;
+  max-width: 18em;
+  white-space: normal;
   color: #316c6f;
 }
 
@@ -1191,12 +1201,10 @@ html[lang="zh"] .md-nav__title {
   text-align: left;
 }
 
-/* The workload key carries the shapes at the width of the page, so the table
- * holds only `W1`, `W2`, … and the measurements.
- */
 /* The workload key: what an op's workloads have in common, stated once, and
- * what varies, one line per workload. Three columns — code, what varies, the
- * benchmark's own id — so a reader compares down a column.
+ * what varies, one line per label. Two columns — the label, what varies — so a
+ * reader compares down a column. It carries the shapes at the width of the
+ * page, so the table holds only the labels and the measurements.
  */
 .md-typeset .wl-key {
   margin: 0.35em 0 0.8em;
@@ -1224,7 +1232,7 @@ html[lang="zh"] .md-nav__title {
 .md-typeset .wl-key .wl-shared {
   display: flex;
   flex-wrap: wrap;
-  gap: 0.15em 0.32em;
+  row-gap: 0.15em;
   margin: 0 0 0.15em;
   font-size: inherit;
   line-height: inherit;
@@ -1235,17 +1243,26 @@ html[lang="zh"] .md-nav__title {
   margin-bottom: 0.2em;
 }
 
+/* A middot between entries, not a second rule: punctuation separates without
+ * claiming to align. Each entry holds its middot in its left padding; the line
+ * is shifted left by that padding and clipped, so no line starts with one.
+ */
+.md-typeset .wl-key .wl-shared,
+.md-typeset .wl-key .wl-flow {
+  margin-left: -1.1em;
+  clip-path: inset(0 0 0 1.1em);
+}
+
 .md-typeset .wl-key .wl-cell {
+  position: relative;
+  padding-left: 1.1em;
   white-space: nowrap;
 }
 
-/* A middot between entries, not a second rule: the rule after the code falls in
- * the same place on every row, a rule between entries would land wherever the
- * previous name ended. Punctuation separates without claiming to align.
- */
-.md-typeset .wl-key .wl-cell + .wl-cell::before {
+.md-typeset .wl-key .wl-cell::before {
   content: "·";
-  margin-right: 0.32em;
+  position: absolute;
+  left: 0.25em;
   color: rgba(26, 23, 32, 0.5);
   font-weight: 400;
 }
@@ -1280,90 +1297,79 @@ html[lang="zh"] .md-nav__title {
   }
 }
 
+/* Labels in one column, what varies in the next. A label past the column's
+ * width wraps at its hyphens, so a long one never pushes the values off.
+ */
 .md-typeset .wl-key ul.wl-rows {
+  display: grid;
+  grid-template-columns: fit-content(24em) minmax(0, 1fr);
+  gap: 0.1em 0.9em;
+  align-items: baseline;
   margin: 0;
   padding: 0;
   list-style: none;
 }
 
-/* The id ranges right, off the same edge on every row, so the middle column is
- * free to be as long as one workload needs without ragging the rest.
- */
 .md-typeset .wl-key ul.wl-rows li {
-  display: flex;
-  flex-wrap: wrap;
-  align-items: baseline;
-  gap: 0 0.9em;
-  margin: 0;
+  display: contents;
 }
 
-.md-typeset .wl-key ul.wl-rows b {
-  flex: 0 0 2.1em;
-  color: var(--tf-violet);
-  font-weight: 700;
+/* No width for two columns: each label on its own line, what varies under it. */
+@media (max-width: 50rem) {
+  .md-typeset .wl-key ul.wl-rows {
+    grid-template-columns: minmax(0, 1fr);
+  }
 }
 
-/* A hairline between the code and what it ran on — the same rule the table
- * draws between the workload column and its measurements.
+/* A hairline between the label and what it ran on — the same rule the table
+ * draws between the workload columns and their measurements.
  */
 .md-typeset .wl-key .wl-delta {
-  display: flex;
-  flex: 1 1 auto;
-  flex-wrap: wrap;
-  gap: 0.15em 0.32em;
   padding-left: 0.9em;
   border-left: 1px solid rgba(26, 23, 32, 0.2);
 }
 
-.md-typeset .wl-key .wl-delta:empty {
-  border-left: 0;
+/* Tensors, then scalars: one line when both fit, else the scalars wrap whole. */
+.md-typeset .wl-key .wl-flow,
+.md-typeset .wl-key .wl-part {
+  display: flex;
+  flex-wrap: wrap;
+  row-gap: 0.15em;
 }
 
-/* A tensor list wider than the column: one tensor per line, and no middots —
- * stacked entries are already separated by the line break.
- */
-.md-typeset .wl-key .wl-shared.wl-stack {
-  flex-direction: column;
-  align-items: flex-start;
-  gap: 0;
+.md-typeset .wl-key .wl-part {
+  min-width: 0;
 }
 
-.md-typeset .wl-key .wl-shared.wl-stack .wl-cell + .wl-cell::before {
-  content: none;
-  margin-right: 0;
+/* Labels too long for the label column: each on its own line, values under it. */
+.md-typeset .wl-key ul.wl-rows.wl-long {
+  grid-template-columns: minmax(0, 1fr);
+  row-gap: 0;
 }
 
-/* Scalars wider than the column: the id takes the first line beside the code,
- * the scalars the second, with the rule running down both.
- */
-.md-typeset .wl-key ul.wl-long li {
-  margin-bottom: 0.3em;
+.md-typeset .wl-key ul.wl-rows.wl-long .wl-id {
+  margin-top: 0.4em;
 }
 
-.md-typeset .wl-key ul.wl-long li:last-child {
-  margin-bottom: 0;
+.md-typeset .wl-key ul.wl-rows.wl-long li:first-child .wl-id {
+  margin-top: 0;
 }
 
-.md-typeset .wl-key ul.wl-long .wl-id {
-  order: 1;
-  flex: 1 1 auto;
-  margin-left: 0;
-  padding-left: 0.9em;
-  border-left: 1px solid rgba(26, 23, 32, 0.2);
+.md-typeset .wl-key .wl-delta:empty {
+  border-left: 0;
 }
 
-.md-typeset .wl-key ul.wl-long .wl-delta {
-  order: 2;
-  flex: 0 0 calc(100% - 3em);
-  margin-left: 3em;
+/* A tensor list wider than the column: one tensor per line. */
+.md-typeset .wl-key .wl-shared.wl-stack {
+  flex-direction: column;
+  align-items: flex-start;
+  gap: 0;
 }
 
 .md-typeset .wl-key .wl-id {
-  flex: 0 0 auto;
-  margin-left: auto;
   padding: 0;
   background: transparent;
-  color: var(--tf-muted);
+  color: #316c6f;
 }
 
 /* A tensor's name is not its shape, and a symbol is not its value: teal
@@ -1380,16 +1386,6 @@ html[lang="zh"] .md-nav__title {
   color: var(--tf-muted);
 }
 
-/* The workload column is now two characters wide, so it centres like the
- * numbers rather than ranging left as a label did.
- */
-.md-typeset .datatable table:not([class]) td.colsep b {
-  color: var(--tf-violet);
-  font-weight: 700;
-}
-
-
-
 /* Roofline figure (performance-guides/memory-bound). Log-log, so the two decades
    between an elementwise kernel and the ridge stay legible. Teal is the roof,
    violet the ridge, rose a kernel below it — the tables' rose for "behind". */
diff --git a/docs/backends.md b/docs/backends.md
index 5e3f57f6..95358271 100644
--- a/docs/backends.md
+++ b/docs/backends.md
@@ -276,48 +276,49 @@ constructing its first op, imports the module named there, and the registration
 module top level fill the registry. There is no base class to inherit and no interface to
 implement.
 
-**Step 2, a target name and a `detect`.** This `detect` is a single line — it claims every
-CPU device:
+**Step 2, a target name and a `detect`.** Both live in `target.py`, and `detect` is a
+single line — it claims every CPU device:
 
 ```python
-from tileops.backend import TensorSpec, register_detector, register_kernel_builder
+TARGET = "torch_cpu"
 
-register_detector(
-    target="torch_cpu",
-    detect=lambda device: device.type == "cpu",
-)
+
+def detect(device: torch.device) -> bool:
+    return device.type == "cpu"
 ```
 
-The two names mean different things: `target="torch_cpu"` names this set of kernels and is
+The two names mean different things: `TARGET = "torch_cpu"` names this set of kernels and is
 the backend author's to choose, while `device.type == "cpu"` is the device type it claims,
 defined by torch.
 
 **Step 3, a `build_kernel` written to the manifest signature.** `RMSNormFwdOp`'s spec
 declares two inputs, `x` and an optional `weight`, and two params, `normalized_shape` and `eps`; the
-function's parameters follow that declaration:
+function's parameters follow that declaration. It sits in `ops/rms_norm.py` with the
+kernel class `CpuRMSNorm`:
 
 ```python
-from .kernels import CpuRMSNorm
-
-
-def build_rms_norm(
-    x: TensorSpec,
-    weight: TensorSpec | None,
-    *,
-    normalized_shape,
-    eps,
-):
+def build_rms_norm(x: TensorSpec, weight: TensorSpec | None, *, normalized_shape, eps):
+    if eps is None:                     # a null manifest default: the ref API's meaning
+        eps = torch.finfo(torch.float32).eps
     return CpuRMSNorm(normalized_shape, eps, x.dtype)
 ```
 
-**Step 4, the registrations at module top level.** One per op taken over:
+**Step 4, the registrations at module top level.** `BUILDERS` in `ops/__init__.py` lists
+every op the target takes over, with each key spelled as in the manifest. The package's
+`__init__.py` registers one detector, then one builder per entry:
 
 ```python
-register_kernel_builder(
-    op="RMSNormFwdOp",
-    target="torch_cpu",
-    build_kernel=build_rms_norm,
-)
+BUILDERS = {
+    "RMSNormFwdOp": build_rms_norm,
+    "GemmFwdOp": build_gemm,
+}
+```
+
+```python
+register_detector(target=TARGET, detect=detect)
+
+for _op, _build_kernel in BUILDERS.items():
+    register_kernel_builder(op=_op, target=TARGET, build_kernel=_build_kernel)
 ```
 
 ## The template project: layout, tests, and retargeting it
@@ -330,11 +331,12 @@ Each file in the example covers one part of the work:
 | --- | --- |
 | `pyproject.toml` | the entry point declaration, which is the whole install mechanism |
 | `src/tileops_cpu/__init__.py` | all registration code |
-| `src/tileops_cpu/kernels.py` | the kernel implementation; a real backend compiles here |
-| `src/tileops_cpu/pending.py` | a builder registered under `op="GemmOp"` — the manifest key is `GemmFwdOp`, and a name that does not match is never called |
+| `src/tileops_cpu/target.py` | the target name and `detect` |
+| `src/tileops_cpu/ops/__init__.py` | `BUILDERS`, the table of every op the target takes over. Its keys are spelled as in the manifest; a builder under a misspelt key is never called |
+| `src/tileops_cpu/ops/rms_norm.py`, `ops/gemm.py` | one module per op, holding the kernel and the builder that constructs it. A real backend compiles in the kernel's constructor |
 | `tests/test_takeover.py` | numerics, validation, normalisation, outputs |
 | `tests/test_discovery.py` | entry point and registration |
-| `tests/test_errors.py` | the four error paths |
+| `tests/test_errors.py` | the three error paths: an unregistered op raises rather than falling back, an unknown target raises, and a failed call binds the op to no target |
 | `tests/test_memoization.py` | when `build_kernel` is called again |
 
 `CpuRMSNorm` does not receive the row count when it is constructed, which is
@@ -346,7 +348,7 @@ The tests need an environment with `tileops` installed:
 
 ```bash
 pip install -e .          # add --no-deps when tileops is already installed
-python -m pytest -q       # with a GPU: 22 passed; without: 20 passed, 2 skipped
+python -m pytest -q       # 24 passed, both with two H200s visible and with CUDA_VISIBLE_DEVICES=""
 ```
 
 The same holds inside the TileOPs dev image, again without modifying TileOPs:
@@ -368,9 +370,9 @@ that TileOPs is too old.
 ### Turning it into a backend for real hardware
 
 1. Copy the repository, rename `tileops_cpu` to `tileops_<hardware>`, and the target name with it.
-2. Change `_detect` to claim the corresponding device type.
-3. Replace `kernels.py` with real kernels — compile on construction, launch on `__call__`.
-4. Pick the first op to take over and write its `build_kernel` against the op's manifest signature.
+2. Change `detect` in `target.py` to claim the corresponding device type.
+3. Replace the kernels in the `ops/` modules with real ones, which compile on construction and launch on `__call__`.
+4. Pick the first op to take over, write its `build_kernel` against the op's manifest signature, and add it to `BUILDERS`.
 5. The four files under [`tests/`](https://github.com/lcy-seso/tileops-backend-example/tree/main/tests) carry over largely as they are; substitute the op and target names.
 6. Add a `build_kernel` per op from there, until every op the target model uses is covered.
 
@@ -496,8 +498,8 @@ If such a failure does show up on your hardware, the traceback stops inside Tile
 than in the backend's `build_kernel`. That is a regression on the TileOPs side: file an
 issue with the traceback.
 
-The example marks two tests `requires_cuda_runtime`, and they skip on a machine with no
-GPU — what they check is this premise, not the backend.
+Every test in the example also passes where no GPU is visible, and none is skipped; that
+run is itself the check of this premise.
 
 ## Error messages and what to do
 
@@ -578,7 +580,7 @@ author meets them:
 | --- | --- | --- |
 | 1 | The public torch-side API and the meaning of each parameter | How the op is called, the parameter names and their semantics are settled; a backend neither defines nor changes them |
 | 2 | Manifest validation | A call whose dtype or shape does not conform is rejected at the op layer and never reaches the backend |
-| 3 | Parameters by name | Parameters arrive under the manifest's `params` names, with the values the op instance holds; a parameter the manifest defaults to null and the caller did not give arrives as `None` |
+| 3 | Parameters by name | Parameters arrive under the manifest's `params` names, with the values the op instance holds. For a parameter the manifest defaults to null, that is the value the op settled on, either `None` or a number |
 | 4 | Input contiguity | Every input the call does not write arrives contiguous; an input it writes arrives as the caller passed it, unless the manifest declares it `contiguous: true` |
 | 5 | Memoisation and reuse of kernels | The builder is called once per specialization: a later call with the same device and input signature reuses the previous return value. A builder may therefore compile, and the op layer guarantees it is not called again |
 | 6 | The `torch.compile` and CUDA-graph boundary | The op layer wraps a call as an opaque operator and registers a fake alongside, so the compiler can infer the output's shape and dtype without executing. **A backend's kernels do nothing for compilation**; see [Bringing an op into torch.compile](torch-compile.md) |
@@ -613,8 +615,9 @@ reaches `build_kernel`](#from-op-layer).
 
 Two consequences:
 
-- **The memo table is bounded and an entry can be evicted at any time.** A backend
-  must not assume the callable it returned stays alive; whatever resources it
+- **An entry may not last.** When a call fails, the op revokes its target decision and
+  drops its memo table. A backend must not assume the callable it returned stays
+  alive; whatever resources it
   depends on, it holds references to itself.
 - **A finer or a coarser grain is resolved on the backend side.** Finer
   distinctions happen inside the backend; to rebuild less often, add a cache inside
diff --git a/docs/backends.zh.md b/docs/backends.zh.md
index 9b6a1b97..c9a8f832 100644
--- a/docs/backends.zh.md
+++ b/docs/backends.zh.md
@@ -218,43 +218,41 @@ torch_cpu = "tileops_cpu"
 
 `pip install` 之后不需要任何初始化：TileOPs 在构造第一个算子时枚举这个组、import 其中声明的模块，模块顶层的注册调用就把注册表填好。既没有需要继承的基类，也没有需要实现的接口。
 
-**第二步，起 target 名，写 `detect`。** 这里的 `detect` 只有一行 —— 认领所有 CPU 设备：
+**第二步，起 target 名，写 `detect`。** 两者都在 `target.py` 里，`detect` 只有一行 —— 认领所有 CPU 设备：
 
 ```python
-from tileops.backend import TensorSpec, register_detector, register_kernel_builder
+TARGET = "torch_cpu"
 
-register_detector(
-    target="torch_cpu",
-    detect=lambda device: device.type == "cpu",
-)
+
+def detect(device: torch.device) -> bool:
+    return device.type == "cpu"
 ```
 
-两个名字含义不同：`target="torch_cpu"` 是这一套 kernel 的名字，由后端作者决定；`device.type == "cpu"` 是它认领的设备类型，由 torch 定义。
+两个名字含义不同：`TARGET = "torch_cpu"` 是这一套 kernel 的名字，由后端作者决定；`device.type == "cpu"` 是它认领的设备类型，由 torch 定义。
 
-**第三步，照 manifest 签名写 `build_kernel`。** `RMSNormFwdOp` 的 spec 声明了两个输入 `x`、`weight`（可选）与两个参数 `normalized_shape`、`eps`，函数的形参照抄这份声明：
+**第三步，照 manifest 签名写 `build_kernel`。** `RMSNormFwdOp` 的 spec 声明了两个输入 `x`、`weight`（可选）与两个参数 `normalized_shape`、`eps`，函数的形参照抄这份声明。它和 kernel 类 `CpuRMSNorm` 一起放在 `ops/rms_norm.py` 里：
 
 ```python
-from .kernels import CpuRMSNorm
-
-
-def build_rms_norm(
-    x: TensorSpec,
-    weight: TensorSpec | None,
-    *,
-    normalized_shape,
-    eps,
-):
+def build_rms_norm(x: TensorSpec, weight: TensorSpec | None, *, normalized_shape, eps):
+    if eps is None:                     # manifest 默认为 null：取参考 API 的含义
+        eps = torch.finfo(torch.float32).eps
     return CpuRMSNorm(normalized_shape, eps, x.dtype)
 ```
 
-**第四步，在模块顶层注册。** 每个要接管的算子登记一次：
+**第四步，在模块顶层注册。** `ops/__init__.py` 里的 `BUILDERS` 列出这个 target 接管的全部算子，每个键都按 manifest 中的拼写书写；包的 `__init__.py` 登记一个 detector，再给表里每个算子登记一次 builder：
 
 ```python
-register_kernel_builder(
-    op="RMSNormFwdOp",
-    target="torch_cpu",
-    build_kernel=build_rms_norm,
-)
+BUILDERS = {
+    "RMSNormFwdOp": build_rms_norm,
+    "GemmFwdOp": build_gemm,
+}
+```
+
+```python
+register_detector(target=TARGET, detect=detect)
+
+for _op, _build_kernel in BUILDERS.items():
+    register_kernel_builder(op=_op, target=TARGET, build_kernel=_build_kernel)
 ```
 
 ## 模板项目：结构、测试与改造
@@ -267,11 +265,12 @@ register_kernel_builder(
 | --- | --- |
 | `pyproject.toml` | entry point 声明，也就是全部安装机制 |
 | `src/tileops_cpu/__init__.py` | 全部注册代码 |
-| `src/tileops_cpu/kernels.py` | kernel 实现，真实后端在此处编译 |
-| `src/tileops_cpu/pending.py` | 一个用 `op="GemmOp"` 注册的 builder —— manifest 的键是 `GemmFwdOp`，名字不一致就永远不会被调到 |
+| `src/tileops_cpu/target.py` | target 名与 `detect` |
+| `src/tileops_cpu/ops/__init__.py` | `BUILDERS`，即这个 target 接管的全部算子。键按 manifest 中的拼写书写，键写错的 builder 永远不会被调用 |
+| `src/tileops_cpu/ops/rms_norm.py`、`ops/gemm.py` | 每个算子一个模块，存放 kernel 实现和构造它的 builder。真实后端在 kernel 的构造函数里编译 |
 | `tests/test_takeover.py` | 数值、校验、归一与输出 |
 | `tests/test_discovery.py` | entry point 与注册 |
-| `tests/test_errors.py` | 四条错误路径 |
+| `tests/test_errors.py` | 三条错误路径：未登记的算子报错而不回退，未知的 target 报错，调用失败后算子不固定到任何 target |
 | `tests/test_memoization.py` | `build_kernel` 在什么条件下被重新调用 |
 
 其中 `CpuRMSNorm` 在构造时**得不到行数**，正是「构造函数只接收编译期参数」这一条的体现。
@@ -282,7 +281,7 @@ register_kernel_builder(
 
 ```bash
 pip install -e .          # tileops 已安装时加 --no-deps
-python -m pytest -q       # 有 GPU：22 passed；无 GPU：20 passed, 2 skipped
+python -m pytest -q       # 两块 H200 可见与 CUDA_VISIBLE_DEVICES="" 两种情况下都是 24 passed
 ```
 
 在 TileOPs 的 dev 镜像中运行同样不需要修改 TileOPs：
@@ -300,9 +299,9 @@ docker run --rm --gpus all -v "$PWD/..":/work -w /work \
 ### 改造成面向真实硬件的后端
 
 1. 复制该仓库，把 `tileops_cpu` 改为 `tileops_<硬件名>`，target 名同样改写。
-2. 修改 `_detect`，认领对应的设备类型。
-3. 把 `kernels.py` 替换为真实 kernel，构造时编译，`__call__` 时启动。
-4. 选定第一个要接管的算子，照它的 manifest 签名编写 `build_kernel`。
+2. 修改 `target.py` 里的 `detect`，认领对应的设备类型。
+3. 把 `ops/` 下各模块的 kernel 替换为真实 kernel：构造时编译，调用 `__call__` 时启动。
+4. 选定第一个要接管的算子，照它的 manifest 签名编写 `build_kernel`，并在 `BUILDERS` 里加一行。
 5. [`tests/`](https://github.com/lcy-seso/tileops-backend-example/tree/main/tests) 中的四个文件大体可以直接沿用，替换其中的算子名与 target 名即可。
 6. 之后逐个算子增加 `build_kernel`，直到覆盖目标模型用到的全部算子。
 
@@ -397,7 +396,7 @@ target 定下来之前，算子层不查询与特定硬件绑定的信息 ——
 
 万一在自己的硬件上撞到这种失败，调用栈会停在 TileOPs 内部、而不是后端的 `build_kernel` 里。那是 TileOPs 一侧的回退，提 issue 并附上调用栈。
 
-示例仓库为两个测试加了 `requires_cuda_runtime` 标记，在没有 GPU 的机器上自动跳过 —— 它们检验的正是这条前提，而不是后端本身。
+在看不到 GPU 的环境里，示例仓库的全部测试照常通过，没有跳过项；这次运行本身就检查了这条前提。
 
 ## 错误信息与处理
 
@@ -462,7 +461,7 @@ TileOPs 不解析 `torch.device`，而是把它原样传给 `detect`。这样做
 | --- | --- | --- |
 | 1 | torch 侧的公开 API 与参数语义 | 这个算子如何被调用、参数名与各参数的含义均已确定，后端既不定义也不能改动 |
 | 2 | manifest 校验 | dtype 或形状不合规的调用被算子层拒绝，不会到达后端 |
-| 3 | 参数按名字传入 | 后端收到的参数名是 manifest `params` 的名字，值是算子实例保存的值；manifest 默认为 null 且调用方没给的参数，收到的是 `None` |
+| 3 | 参数按名字传入 | 后端收到的参数名是 manifest `params` 的名字，值是算子实例保存的值。manifest 中默认为 null 的参数，传入的是算子选定的值，可能是 `None`，也可能是一个具体数值 |
 | 4 | 输入的连续性归一 | 本次调用不写入的输入都转成连续张量；被写入的输入按调用方传入的原样交给后端，除非 manifest 声明它 `contiguous: true` |
 | 5 | kernel 的记忆与重用 | 构造函数按特化调用一次：设备与输入签名相同的后续调用直接使用上一次的返回值。因此构造函数内部可以编译，算子层保证它不会被重复调用 |
 | 6 | `torch.compile` 与 CUDA graph 的边界 | 算子层把一次调用包成不透明算子并另配一个 fake，使编译器在不执行的前提下也能推出输出的形状与 dtype。**后端的 kernel 不为编译做任何事**，细节见[接入 torch.compile](torch-compile.md) |
@@ -486,7 +485,7 @@ TileOPs 按**设备加输入签名**记住 `build_kernel` 的返回值。这个
 
 两点由此而来：
 
-- **记忆表有上限，条目可以在任何时候被淘汰。** 后端不得假设自己返回的可调用对象一直存活；它所依赖的资源应由它自己持有引用。
+- **条目不一定一直保留。** 一次调用失败时，算子会撤销它的 target 判定，并清空记忆表。后端不得假设自己返回的可调用对象一直存活；它所依赖的资源应由它自己持有引用。
 - **需要更细或更粗的粒度，都在后端一侧解决。** 更细的区分在后端内部处理；希望减少重建次数，可以在 `build_kernel` 内部另加一层缓存。
 
 ## 调用方可用的接口
diff --git a/docs/index.md b/docs/index.md
index 43d98da4..07a2ed68 100644
--- a/docs/index.md
+++ b/docs/index.md
@@ -13,7 +13,7 @@ judgement call. An implementation can therefore be regenerated from its spec,
 while the reverse does not hold.
 
 To a caller it is simply a set of operators: shapes and dtype come from the call,
-the specialized kernel is auto-tuned and cached on first use and works under CUDA
+the specialized kernel is built and cached on first use and works under CUDA
 graphs afterwards, and each op declares whether it supports
 `torch.compile(fullgraph=True)`.
 
diff --git a/docs/index.zh.md b/docs/index.zh.md
index fee0ee4b..2803ff7b 100644
--- a/docs/index.zh.md
+++ b/docs/index.zh.md
@@ -4,7 +4,7 @@ TileOPs 是一个面向大模型推理的算子库，构建在 [TileLang](https:
 
 它与手写算子库的不同之处在于组织方式：每个算子先以一份 spec 声明，再由 agent 依据这份 spec 生成实现。spec 既是代码生成的唯一依据，也是验收的标准 —— 正确性对照 spec 指定的参考实现，性能对照 roofline 模型给出的上界，两项都不依赖人的判断。因此一个实现可以随时从 spec 重新生成，而反过来做不到。
 
-对使用者而言，它就是一批可以直接调用的算子：形状与 dtype 在调用时确定，特化后的 kernel 在首次使用时自动调优并缓存，随后可以与 CUDA graph 配合使用；每个算子各自声明是否支持 `torch.compile(fullgraph=True)`。
+对使用者而言，它就是一批可以直接调用的算子：形状与 dtype 在调用时确定，特化后的 kernel 在首次使用时构造并缓存，随后可以与 CUDA graph 配合使用；每个算子各自声明是否支持 `torch.compile(fullgraph=True)`。
 
 ## 安装
 
diff --git a/docs/manifest.md b/docs/manifest.md
index 8227f51e..230ec6ce 100644
--- a/docs/manifest.md
+++ b/docs/manifest.md
@@ -91,7 +91,7 @@ RMSNormFwdOp:
        dtype_cases: [{T: float16}, {T: bfloat16}], label: llama-8b-prefill}
     - {B: [4, 2048], normalized_shape: [128], dtype_cases: [{T: bfloat16}], label: qk-norm-head}
   roofline:
-    flops: "4 * prod(B) * prod(normalized_shape)"
+    flops: "(4 if present(weight) else 3) * prod(B) * prod(normalized_shape)"
 ```
 
 Read it in five steps:
@@ -196,11 +196,11 @@ and the roofline reads presence with `present`:
       - "num_groups > 0 and C % num_groups == 0"
       - "B * (C // num_groups) * prod(L) != 1"
   workloads:
-    - {B: 8, C: 128, L: [32, 32], num_groups: 32, dtype_cases: [{T: float16}], label: image-g32}
+    - {B: 8, C: 128, L: [32, 32], num_groups: 32, dtype_cases: [{T: float16}], label: image}
     - {B: 8, C: 128, L: [32, 32], num_groups: 32, some: [weight, bias],
-       dtype_cases: [{T: float16}], label: image-g32-affine}
+       dtype_cases: [{T: float16}], label: image-affine}
   roofline:
-    flops: "(5 if present(weight) or present(bias) else 3) * B * C * prod(L)"
+    flops: "(5 + (1 if present(weight) else 0) + (1 if present(bias) else 0)) * B * C * prod(L)"
 ```
 
 (Excerpt: the entry has more rows and dtypes.) An op may dispatch on whether an optional
diff --git a/docs/manifest.zh.md b/docs/manifest.zh.md
index 15522c5a..cf4cbb4b 100644
--- a/docs/manifest.zh.md
+++ b/docs/manifest.zh.md
@@ -75,7 +75,7 @@ RMSNormFwdOp:
        dtype_cases: [{T: float16}, {T: bfloat16}], label: llama-8b-prefill}
     - {B: [4, 2048], normalized_shape: [128], dtype_cases: [{T: bfloat16}], label: qk-norm-head}
   roofline:
-    flops: "4 * prod(B) * prod(normalized_shape)"
+    flops: "(4 if present(weight) else 3) * prod(B) * prod(normalized_shape)"
 ```
 
 分五步读：
@@ -154,11 +154,11 @@ load_workloads("RMSNormFwdOp")             # that op's workload rows
       - "num_groups > 0 and C % num_groups == 0"
       - "B * (C // num_groups) * prod(L) != 1"
   workloads:
-    - {B: 8, C: 128, L: [32, 32], num_groups: 32, dtype_cases: [{T: float16}], label: image-g32}
+    - {B: 8, C: 128, L: [32, 32], num_groups: 32, dtype_cases: [{T: float16}], label: image}
     - {B: 8, C: 128, L: [32, 32], num_groups: 32, some: [weight, bias],
-       dtype_cases: [{T: float16}], label: image-g32-affine}
+       dtype_cases: [{T: float16}], label: image-affine}
   roofline:
-    flops: "(5 if present(weight) or present(bias) else 3) * B * C * prod(L)"
+    flops: "(5 + (1 if present(weight) else 0) + (1 if present(bias) else 0)) * B * C * prod(L)"
 ```
 
 （节选：条目里还有更多行和 dtype。）算子可以按可选输入是否传入来分派，不能读张量内容来决定。
diff --git a/docs/new-op.md b/docs/new-op.md
index d612cfc1..7ffc39e4 100644
--- a/docs/new-op.md
+++ b/docs/new-op.md
@@ -50,7 +50,7 @@ GemmFwdOp:
       d: {dtype: T, shape: "[M, N]"}
   workloads:
     - {M: 4096, N: 4096, K: 7168, trans_a: false, trans_b: true,
-       dtype_cases: [{T: float16}, {T: bfloat16}], label: ds-v3-prefill-attn-proj}
+       dtype_cases: [{T: float16}, {T: bfloat16}], label: ds-v3-prefill-mlp-up}
   roofline:
     flops: "2 * M * N * K"
 ```
@@ -165,7 +165,8 @@ reference cannot collect it.
 ## Step 3: write the kernel
 
 A kernel class subclasses [`Kernel`](https://github.com/tile-ai/TileOPs/blob/main/src/tileops/kernels/kernel_base.py), lives under [`src/tileops/kernels/`](https://github.com/tile-ai/TileOPs/tree/main/src/tileops/kernels), is written in
-TileLang, compiles at construction and launches on `__call__`. Its constructor is what its
+TileLang, compiles at construction and implements `forward`, which the base class's
+`__call__` runs. Its constructor is what its
 `entry_for` builder calls, and its call signature is the `kernel(a, b)` of step 2.
 
 This is the one place of the six the spec does not constrain: a kernel neither reads the
@@ -181,13 +182,13 @@ class GemmTmaKernel(Kernel):
         self.kernel = _gemm_kernel(m, n, k, trans_a, trans_b, self.dtype_str, ...)  # compiles
         self.init_config(config, tune)      # tile sizes and pipeline depth
 
-    def __call__(self, a, b):               # a call passes tensors, nothing else
+    def forward(self, a, b):                # a call passes tensors, nothing else
         ...
 ```
 
 `m`, `n`, `k`, the dtype and the two layout flags are constructor arguments because the
 generated code treats them as constants: loop bounds, TMA descriptors and the WGMMA shape
-all unroll from them, as do the tile sizes. The tensors belong to `__call__`, where each
+all unroll from them, as do the tile sizes. The tensors belong to `forward`, where each
 call swaps pointers.
 
 Dividing them wrong costs a recompile. A decode step advances one token at a time, so
@@ -209,8 +210,9 @@ misses, every step compiles, and decode goes nowhere.
 
 Tests live in [`tests/ops/`](https://github.com/tile-ai/TileOPs/tree/main/tests/ops) and compare against `ref_program`, the reference the
 workload (or the test class) defines, over shapes the test chooses to reach the kernel's
-branches — small shapes marked `smoke` for the PR checks, large ones `full` for the
-nightly. The workload rows are not unit-test coverage; the
+branches. Cases marked `smoke` run on every PR. Cases marked `full` run on any PR that
+changes their test file, and on the nightly. Long-running cases are marked `nightly` and
+run only on the nightly. The workload rows are not unit-test coverage; the
 contract tests already run each of them through the op.
 
 The scaffolding is `TestBase` and `FixtureBase` from
diff --git a/docs/new-op.zh.md b/docs/new-op.zh.md
index 0dd15b0b..9ff07bd0 100644
--- a/docs/new-op.zh.md
+++ b/docs/new-op.zh.md
@@ -45,7 +45,7 @@ GemmFwdOp:
       d: {dtype: T, shape: "[M, N]"}
   workloads:
     - {M: 4096, N: 4096, K: 7168, trans_a: false, trans_b: true,
-       dtype_cases: [{T: float16}, {T: bfloat16}], label: ds-v3-prefill-attn-proj}
+       dtype_cases: [{T: float16}, {T: bfloat16}], label: ds-v3-prefill-mlp-up}
   roofline:
     flops: "2 * M * N * K"
 ```
@@ -130,7 +130,7 @@ def entry_for(self, role, call):                    # call is the input dtype
 
 ## 第三步：写 kernel
 
-kernel 类继承 [`Kernel`](https://github.com/tile-ai/TileOPs/blob/main/src/tileops/kernels/kernel_base.py)，放在 [`src/tileops/kernels/`](https://github.com/tile-ai/TileOPs/tree/main/src/tileops/kernels) 下，用 TileLang 写，构造时编译、`__call__` 时启动。构造函数由它的 `entry_for` 构造方法调用，调用签名就是第二步里的 `kernel(a, b)`。
+kernel 类继承 [`Kernel`](https://github.com/tile-ai/TileOPs/blob/main/src/tileops/kernels/kernel_base.py)，放在 [`src/tileops/kernels/`](https://github.com/tile-ai/TileOPs/tree/main/src/tileops/kernels) 下，用 TileLang 编写，在构造时编译。它实现 `forward`，基类的 `__call__` 会调用它。构造函数由它的 `entry_for` 构造方法调用，调用签名就是第二步里的 `kernel(a, b)`。
 
 它是这六处里唯一不受 spec 约束的一处：kernel 不读 spec，也不对照 spec 检查。
 
@@ -142,11 +142,11 @@ class GemmTmaKernel(Kernel):
         self.kernel = _gemm_kernel(m, n, k, trans_a, trans_b, self.dtype_str, ...)  # compiles
         self.init_config(config, tune)      # tile sizes and pipeline depth
 
-    def __call__(self, a, b):               # a call passes tensors, nothing else
+    def forward(self, a, b):                # a call passes tensors, nothing else
         ...
 ```
 
-`m`、`n`、`k`、dtype 与两个布局标志进了构造函数，因为生成的代码里这些值是常量：循环边界、TMA 描述符、WGMMA 的形状都按它们展开，tile 尺寸同理。张量本身留给 `__call__`，每次调用只换指针。
+`m`、`n`、`k`、dtype 与两个布局标志进了构造函数，因为生成的代码里这些值是常量：循环边界、TMA 描述符、WGMMA 的形状都按它们展开，tile 尺寸同理。张量本身留给 `forward`，每次调用只换指针。
 
 分错的代价是重新编译。decode 一步一步往前走，`seq_len` 每步 +1，batch 随 running set 变化：
 
@@ -163,7 +163,7 @@ out = kernel(q, k, v)                       # seq_len 从张量形状里读
 
 ## 第四步：写测试
 
-测试放在 [`tests/ops/`](https://github.com/tile-ai/TileOPs/tree/main/tests/ops)，比对对象是 workload（或测试类）定义的参考实现 `ref_program`，形状由测试自己挑，以覆盖 kernel 的各个分支；小形状标 `smoke` 进 PR 检查，大形状标 `full` 留给 nightly。workload 行不是单元测试的覆盖面，契约测试已经把每一行都交给算子跑过。
+测试放在 [`tests/ops/`](https://github.com/tile-ai/TileOPs/tree/main/tests/ops)，比对对象是 workload（或测试类）定义的参考实现 `ref_program`，形状由测试自己挑，以覆盖 kernel 的各个分支；`smoke` 用例每个 PR 都运行；`full` 用例在改动其测试文件的 PR 和 nightly 中运行；耗时长的用例标 `nightly`，只在 nightly 运行。workload 行不是单元测试的覆盖面，契约测试已经把每一行都交给算子跑过。
 
 骨架用 [`tests/test_base.py`](https://github.com/tile-ai/TileOPs/blob/main/tests/test_base.py) 里的 `TestBase` 与 `FixtureBase`，用例写在 `PARAMS` 里。
 
diff --git a/docs/performance-guides/memory-bound/index.md b/docs/performance-guides/memory-bound/index.md
index c662860c..c6409f1a 100644
--- a/docs/performance-guides/memory-bound/index.md
+++ b/docs/performance-guides/memory-bound/index.md
@@ -9,16 +9,16 @@ with a [macro benchmark](https://github.com/tile-ai/TileOPs/tree/main/benchmarks
 the hardware peak from the specification, multiplied by that factor, gives the
 effective peak a kernel can actually reach. That effective peak is the reference
 point for tuning. On an H200, the measured values are
-[**57.27** TFLOP/s for fp32 FMA and **4.07** TB/s of memory bandwidth](https://github.com/tile-ai/TileOPs/blob/main/src/tileops/perf/profiles/h200.yaml).
+[**57.27** TFLOP/s for fp32 FMA and **4.50** TB/s of memory bandwidth](https://github.com/tile-ai/TileOPs/blob/main/src/tileops/perf/profiles/h200.yaml).
 The **ridge point** of the roofline is where the bandwidth slope meets the
 compute ceiling. Dividing compute by bandwidth places it at an arithmetic
-intensity of **14.07 flop/byte**: the number of floating-point operations per
+intensity of **12.72 flop/byte**: the number of floating-point operations per
 byte moved when compute and bandwidth are saturated at the same time:
 
 <figure class="roofline" markdown="1">
 
-<svg class="tf-roofline" viewBox="0 0 520 306" role="img" aria-label="Roofline for an H200: the 4.07 TB/s bandwidth slope meets the 57.27 TFLOP/s compute ceiling at 14 flop per byte; silu sits far down the slope, reaching 9% of the compute ceiling.">
-<path class="tf-rl-region" d="M 52.0 250.0 L 52.0 217.9 L 357.4 67.1 L 357.4 250.0 Z"/>
+<svg class="tf-roofline" viewBox="0 0 520 306" role="img" aria-label="Roofline for an H200: the 4.50 TB/s bandwidth slope meets the 57.27 TFLOP/s compute ceiling at 13 flop per byte; silu sits far down the slope, reaching 10% of the compute ceiling.">
+<path class="tf-rl-region" d="M 52.0 250.0 L 52.0 213.3 L 348.1 67.1 L 348.1 250.0 Z"/>
 <line class="tf-rl-grid" x1="115.4" y1="52" x2="115.4" y2="250"/>
 <text class="tf-rl-tick" x="115.4" y="270" text-anchor="middle">1</text>
 <line class="tf-rl-grid" x1="178.9" y1="52" x2="178.9" y2="250"/>
@@ -51,19 +51,19 @@ byte moved when compute and bandwidth are saturated at the same time:
 <line class="tf-rl-axis" x1="52" y1="52" x2="52" y2="250"/>
 <text class="tf-rl-axis-title" x="4" y="34">Attainable TFLOP/s (log)</text>
 <text class="tf-rl-axis-title" x="496" y="294" text-anchor="end">Arithmetic intensity: flop per byte (log)</text>
-<polyline class="tf-rl-roof" points="52.0,217.9 357.4,67.1 496.0,67.1"/>
-<line class="tf-rl-drop" x1="357.4" y1="67.1" x2="357.4" y2="250"/>
-<circle class="tf-rl-ridge" cx="357.4" cy="67.1" r="6"/>
-<text class="tf-rl-label tf-rl-label--ridge" x="370.4" y="89.1">Ridge point</text>
-<text class="tf-rl-sub" x="370.4" y="107.1">14 flop/byte</text>
-<circle class="tf-rl-point" cx="135.8" cy="176.5" r="5.5"/>
-<text class="tf-rl-label tf-rl-label--point" x="148.8" y="196.5">silu (fp16)</text>
-<text class="tf-rl-sub" x="148.8" y="214.5">5 flop / 4 bytes — ceiling 5.1</text>
-<text class="tf-rl-roof-label" x="152.5" y="153.3" transform="rotate(-30 152.5 153.3)">bandwidth roof 4.07 TB/s</text>
+<polyline class="tf-rl-roof" points="52.0,213.3 348.1,67.1 496.0,67.1"/>
+<line class="tf-rl-drop" x1="348.1" y1="67.1" x2="348.1" y2="250"/>
+<circle class="tf-rl-ridge" cx="348.1" cy="67.1" r="6"/>
+<text class="tf-rl-label tf-rl-label--ridge" x="361.1" y="89.1">Ridge point</text>
+<text class="tf-rl-sub" x="361.1" y="107.1">13 flop/byte</text>
+<circle class="tf-rl-point" cx="135.8" cy="171.9" r="5.5"/>
+<text class="tf-rl-label tf-rl-label--point" x="148.8" y="191.9">silu (fp16)</text>
+<text class="tf-rl-sub" x="148.8" y="209.9">5 flop / 4 bytes — ceiling 5.6</text>
+<text class="tf-rl-roof-label" x="152.5" y="148.7" transform="rotate(-30 152.5 148.7)">bandwidth roof 4.50 TB/s</text>
 <text class="tf-rl-roof-label" x="492.0" y="54.1" text-anchor="end">compute roof 57.3 TFLOP/s</text>
 </svg>
 
-<figcaption>The bent line is the roofline. Every kernel's performance point lies below it. To the left of the ridge, the ceiling is arithmetic intensity times bandwidth, so attainable compute rises linearly with intensity. To the right, the ceiling is the compute peak, and higher intensity no longer raises it. <code>silu</code> performs 5 operations per 4 bytes moved, for an arithmetic intensity of 1.25 flop/byte, one eleventh of the ridge. Even with bandwidth fully saturated, it can reach only 9% of the compute ceiling.</figcaption>
+<figcaption>The bent line is the roofline. Every kernel's performance point lies below it. To the left of the ridge, the ceiling is arithmetic intensity times bandwidth, so attainable compute rises linearly with intensity. To the right, the ceiling is the compute peak, and higher intensity no longer raises it. <code>silu</code> performs 5 operations per 4 bytes moved, for an arithmetic intensity of 1.25 flop/byte, one tenth of the ridge. Even with bandwidth fully saturated, it can reach only 10% of the compute ceiling.</figcaption>
 
 </figure>
 
diff --git a/docs/performance-guides/memory-bound/index.zh.md b/docs/performance-guides/memory-bound/index.zh.md
index 404f4915..e9d83131 100644
--- a/docs/performance-guides/memory-bound/index.zh.md
+++ b/docs/performance-guides/memory-bound/index.zh.md
@@ -2,12 +2,12 @@
 
 ## 什么是访存受限
 
-在 GPU 上，一个 kernel 跑多快，取决于算力与带宽哪一个先成为瓶颈。TileOPs 用 [macro benchmark](https://github.com/tile-ai/TileOPs/tree/main/benchmarks/hardware) 测算出一个**[校准系数](https://github.com/tile-ai/TileOPs/blob/main/src/tileops/perf/profiles/h200.yaml)**：硬件 spec 给出的理论峰值乘以校准系数，得到实际可达的有效值，以此作为性能优化的指导标准。我们在 H200 上实测出 [fp32 FMA 的算力为 **57.27** TFLOP/s，访存带宽为 **4.07** TB/s](https://github.com/tile-ai/TileOPs/blob/main/src/tileops/perf/profiles/h200.yaml)。roofline 的**拐点**（ridge point）是带宽斜线与算力上限这两段的交点，两者相除给出它的横坐标，也就是拐点处的算术强度 **14.07 flop/byte** —— 算力与带宽同时用满时，每搬运一个字节对应的浮点运算次数：
+在 GPU 上，一个 kernel 跑多快，取决于算力与带宽哪一个先成为瓶颈。TileOPs 用 [macro benchmark](https://github.com/tile-ai/TileOPs/tree/main/benchmarks/hardware) 测算出一个**[校准系数](https://github.com/tile-ai/TileOPs/blob/main/src/tileops/perf/profiles/h200.yaml)**：硬件 spec 给出的理论峰值乘以校准系数，得到实际可达的有效值，以此作为性能优化的指导标准。我们在 H200 上实测出 [fp32 FMA 的算力为 **57.27** TFLOP/s，访存带宽为 **4.50** TB/s](https://github.com/tile-ai/TileOPs/blob/main/src/tileops/perf/profiles/h200.yaml)。roofline 的**拐点**（ridge point）是带宽斜线与算力上限这两段的交点，两者相除给出它的横坐标，也就是拐点处的算术强度 **12.72 flop/byte** —— 算力与带宽同时用满时，每搬运一个字节对应的浮点运算次数：
 
 <figure class="roofline" markdown="1">
 
-<svg class="tf-roofline" viewBox="0 0 520 306" role="img" aria-label="H200 的 roofline：带宽上限 4.07 TB/s 的斜线在每字节 14 次浮点运算处与 57.27 TFLOP/s 的算力上限相交；silu 位于斜线左端，可达算力为算力上限的 9%。">
-<path class="tf-rl-region" d="M 52.0 250.0 L 52.0 217.9 L 357.4 67.1 L 357.4 250.0 Z"/>
+<svg class="tf-roofline" viewBox="0 0 520 306" role="img" aria-label="H200 的 roofline：带宽上限 4.50 TB/s 的斜线在每字节 13 次浮点运算处与 57.27 TFLOP/s 的算力上限相交；silu 位于斜线左端，可达算力为算力上限的 10%。">
+<path class="tf-rl-region" d="M 52.0 250.0 L 52.0 213.3 L 348.1 67.1 L 348.1 250.0 Z"/>
 <line class="tf-rl-grid" x1="115.4" y1="52" x2="115.4" y2="250"/>
 <text class="tf-rl-tick" x="115.4" y="270" text-anchor="middle">1</text>
 <line class="tf-rl-grid" x1="178.9" y1="52" x2="178.9" y2="250"/>
@@ -40,19 +40,19 @@
 <line class="tf-rl-axis" x1="52" y1="52" x2="52" y2="250"/>
 <text class="tf-rl-axis-title" x="4" y="34">可达算力 TFLOP/s（对数轴）</text>
 <text class="tf-rl-axis-title" x="496" y="294" text-anchor="end">算术强度：每字节的浮点运算数（对数轴）</text>
-<polyline class="tf-rl-roof" points="52.0,217.9 357.4,67.1 496.0,67.1"/>
-<line class="tf-rl-drop" x1="357.4" y1="67.1" x2="357.4" y2="250"/>
-<circle class="tf-rl-ridge" cx="357.4" cy="67.1" r="6"/>
-<text class="tf-rl-label tf-rl-label--ridge" x="370.4" y="89.1">拐点</text>
-<text class="tf-rl-sub" x="370.4" y="107.1">每字节 14 次</text>
-<circle class="tf-rl-point" cx="135.8" cy="176.5" r="5.5"/>
-<text class="tf-rl-label tf-rl-label--point" x="148.8" y="196.5">silu (fp16)</text>
-<text class="tf-rl-sub" x="148.8" y="214.5">5 次 / 4 字节，上限 5.1</text>
-<text class="tf-rl-roof-label" x="152.5" y="153.3" transform="rotate(-30 152.5 153.3)">带宽上限 4.07 TB/s</text>
+<polyline class="tf-rl-roof" points="52.0,213.3 348.1,67.1 496.0,67.1"/>
+<line class="tf-rl-drop" x1="348.1" y1="67.1" x2="348.1" y2="250"/>
+<circle class="tf-rl-ridge" cx="348.1" cy="67.1" r="6"/>
+<text class="tf-rl-label tf-rl-label--ridge" x="361.1" y="89.1">拐点</text>
+<text class="tf-rl-sub" x="361.1" y="107.1">每字节 13 次</text>
+<circle class="tf-rl-point" cx="135.8" cy="171.9" r="5.5"/>
+<text class="tf-rl-label tf-rl-label--point" x="148.8" y="191.9">silu (fp16)</text>
+<text class="tf-rl-sub" x="148.8" y="209.9">5 次 / 4 字节，上限 5.6</text>
+<text class="tf-rl-roof-label" x="152.5" y="148.7" transform="rotate(-30 152.5 148.7)">带宽上限 4.50 TB/s</text>
 <text class="tf-rl-roof-label" x="492.0" y="54.1" text-anchor="end">算力上限 57.3 TFLOP/s</text>
 </svg>
 
-<figcaption>图上这条折线就是 roofline，任何 kernel 的性能点都落在它以下。拐点以左，上限是「算术强度 × 带宽」，可达算力随算术强度线性上升；拐点以右，上限就是算力峰值，不再随算术强度变化。<code>silu</code> 每搬运 4 个字节做 5 次运算，算术强度 1.25 flop/byte，只有拐点的 1/11，所以即便带宽完全用满，也只能达到算力上限的 9%。</figcaption>
+<figcaption>图上这条折线就是 roofline，任何 kernel 的性能点都落在它以下。拐点以左，上限是「算术强度 × 带宽」，可达算力随算术强度线性上升；拐点以右，上限就是算力峰值，不再随算术强度变化。<code>silu</code> 每搬运 4 个字节做 5 次运算，算术强度 1.25 flop/byte，只有拐点的 1/10，所以即便带宽完全用满，也只能达到算力上限的 10%。</figcaption>
 
 </figure>
 
diff --git a/docs/timing.md b/docs/timing.md
index cab6aeb6..c9eb2767 100644
--- a/docs/timing.md
+++ b/docs/timing.md
@@ -40,7 +40,7 @@ Inside, `bench_kernel` is three stages — collect, attribute, measure:
 
 ```python
 # Collect: each call runs under its own iteration number
-with _phase_session():                          # kernels + mappings + launch APIs
+with _phase_session():                          # kernels + copies + mappings + launch APIs
     for i in range(n_repeat):
         with _labelled(_PREPARE_ID):            # the L2 flush and other preparation
             prepare_one(i)
@@ -87,7 +87,7 @@ Five choices in it, each for a reason:
 
 | Case | Raises | Meaning |
 | --- | --- | --- |
-| CUPTI discarded records | `_CUPTIRecordsLostError` | the reading is gone though the iteration did run — the whole phase is measured again, up to 3 times, asking for a 4× larger buffer each time |
+| CUPTI discarded records | `_CUPTIRecordsLostError` | the reading is gone though the iteration did run — the whole phase is measured again, up to 3 attempts in all, asking for a 4× larger buffer each time |
 | Nothing discarded, but a kernel carries no iteration number | `_OffThreadLaunchError` | a thread that never pushed an id launched it |
 | Nothing discarded, and one iteration has no kernels at all | `_CUPTIAttributionError` | that call never reached the device |
 
@@ -163,7 +163,7 @@ call for a stop:
 | --- | --- | --- |
 | The timed closure contains `Tensor.backward` or `torch.autograd.grad` | the backward kernels come from the autograd engine's own thread, carry no iteration number, and the case raises instead of producing a figure | Drive a single fused node with `backward_of(out)`; for a chain, set `torch.autograd.set_multithreading_enabled(False)` |
 | Another thread in the process uses the GPU, or the timed closure uses CUPTI's `CUSTOM0` external id | those kernels carry no iteration number, or the closure overwrites the one the timer set, and it raises either way | Have the timed call launch its own work; use `CUSTOM1` / `CUSTOM2` instead |
-| The op produces its result through `copy_` — in-place elementwise, MoE's write-back | that copy goes through `cudaMemcpyAsync`, which a kernel-only collection never sees, so the reading is short and nothing raises | Collect `MEMCPY` / `MEMSET` as a check and confirm the window has none |
+| The op produces its result through `copy_` — in-place elementwise, MoE's write-back | the timer collects the copy but leaves it out of `device_busy_ms` and reports it as `uncounted_copy_ms`, so the reading is short | Pass `count_copies=True` for the case; every tag's reading then includes the copies |
 | One call launches several kernels | the gaps between kernels land in `latency_ms`, so comparing by it against a fused implementation charges them to your side | Conclude from `device_busy_ms` only; `latency_ms` compares between rows of equal `n_kernels` |
 | One call takes more than 10 ms | the iteration count hits the floor of 10, the wall-clock far exceeds the 100 ms budget, and p10/p90 over 10 samples are coarse | Accept the longer wall-clock, or state an iteration count and the sample size |
 | You want a kernel-level benchmark | the op has no spec, so shapes and roofline have to be written by hand and the spec validator cannot see them | Measure through the Op interface and write a [spec](manifest.md) |
diff --git a/docs/timing.zh.md b/docs/timing.zh.md
index 7a1f5e11..dffa3d8b 100644
--- a/docs/timing.zh.md
+++ b/docs/timing.zh.md
@@ -27,7 +27,7 @@ samples = bench_kernel(op, args=(x, weight))   # 每次迭代一个 Sample
 
 ```python
 # 采集：每次调用在自己的迭代号下执行
-with _phase_session():                          # kernel + 映射 + launch 类 API
+with _phase_session():                          # kernel + 拷贝 + 映射 + launch 类 API
     for i in range(n_repeat):
         with _labelled(_PREPARE_ID):            # 准备工作用专用 id 标记
             prepare_one(i)
@@ -64,7 +64,7 @@ for i in range(n_repeat):
 
 | 情形 | 报什么 | 含义 |
 | --- | --- | --- |
-| CUPTI 丢了记录 | `_CUPTIRecordsLostError` | 那次迭代跑过，只是读数丢了 —— 整个 phase 重测，最多 3 次，每次把缓冲要大 4 倍 |
+| CUPTI 丢了记录 | `_CUPTIRecordsLostError` | 那次迭代跑过，只是读数丢了 —— 整个 phase 重测，总共最多测 3 次，每次把缓冲要大 4 倍 |
 | 没有丢弃，但有 kernel 带不上迭代号 | `_OffThreadLaunchError` | 某个没有标记迭代号的线程发起了它 |
 | 没有丢弃，某次迭代一个 kernel 都没有 | `_CUPTIAttributionError` | 这次调用没上设备 |
 
@@ -112,7 +112,7 @@ decode 尺度上，算子的执行时间可能短于发起它的那次 Python 
 | --- | --- | --- |
 | 被测闭包里包了 `Tensor.backward` 或 `torch.autograd.grad` | 反向的 kernel 由 autograd 引擎的线程发出，带不上迭代号，整个用例报错，不产出数字 | 单个融合节点用 `backward_of(out)` 直接驱动；多节点链改用 `torch.autograd.set_multithreading_enabled(False)` |
 | 进程里有别的线程在用 GPU，或被测闭包自己用了 CUPTI 的 `CUSTOM0` external id | 那些 kernel 带不上迭代号，或者迭代号被闭包盖掉，同样报错 | 让被计时的调用自己启动它的工作；external id 改用 `CUSTOM1` / `CUSTOM2` |
-| 算子靠 `copy_` 回写才产出结果，例如原地 elementwise 与 MoE 的写回 | 那次拷贝走 `cudaMemcpyAsync`，只采集 kernel 活动的计时看不见它，读数偏小，而且不报错 | 采集 `MEMCPY` / `MEMSET` 当检查用，确认窗口内没有 |
+| 算子靠 `copy_` 回写才产出结果，例如原地 elementwise 与 MoE 的写回 | 计时会采集这次拷贝，但默认不计入 `device_busy_ms`，另记在 `uncounted_copy_ms`，读数因此偏小 | 给这个用例传 `count_copies=True`，所有 tag 的读数都会计入拷贝 |
 | 一次调用发多个 kernel | kernel 之间的空隙落在 `latency_ms` 里，用它与融合实现比较，空隙算在多 kernel 这一方 | 结论只用 `device_busy_ms`；`latency_ms` 仅在两行 `n_kernels` 相同时可比 |
 | 单次调用超过 10 ms | 迭代次数被下限 10 顶住，墙钟远超 100 ms 的预算，10 个样本给出的 p10/p90 很粗 | 接受更长的墙钟，或显式指定迭代次数并写明样本量 |
 | 想加一个 kernel 级的 benchmark | 这个算子没有 spec，形状与 roofline 只能手写，spec 校验器也查不到它 | 经 Op 接口测，并补一份 [spec](manifest.md) |
diff --git a/docs/torch-compile.md b/docs/torch-compile.md
index 53f6a814..2ba0f197 100644
--- a/docs/torch-compile.md
+++ b/docs/torch-compile.md
@@ -68,8 +68,10 @@ compiled path behave differently from the eager one.
 - **Warm up before a CUDA-graph capture.** Call once with real tensors at the same
   shape: building a kernel may compile, while capture allows only a memo hit and the call
   that follows it. See [what each phase may do](backends.md#phase-limits).
-- **Every device builds its own kernel.** The device is part of the kernel memo key, so
-  the same instance builds again on a second card. A `target=` named in the constructor
+- **A second card may need its own build.** For a call a target serves, the device is
+  part of the kernel's memo key, so the same instance builds again on a second card. An
+  in-tree kernel is keyed by what its `entry_for` returns, which includes the device only
+  when the build depends on it. A `target=` named in the constructor
   is honoured on the first compiled call too, and a failed build pins the op to no
   target.
 
@@ -116,8 +118,9 @@ entry, one operator per effect branch: its tensor arguments are `signature.input
 order, what it returns is `signature.outputs`, the arguments it writes are exactly the
 inputs marked `mutated`, and each output's shape and dtype come from the signature. Its
 name is `tileops::<family>_<snake(class)>`, with the family written once where the class
-name already opens with it — here `tileops::norm_rms_norm_fwd`; where the entry declares
-a written input or `buffer: out`, that branch's operator adds `_writes_<input>` or `_out`.
+name already opens with it — here `tileops::norm_rms_norm_fwd`; when a branch writes
+an input, fills a `buffer: out` or omits an output, its operator's name gains
+`_writes_<input>`, `_out` and `_without_<output>` respectively, in that order.
 No op chooses its own name, so `compile_op_names` cannot disagree with what was
 registered.
 
diff --git a/docs/torch-compile.zh.md b/docs/torch-compile.zh.md
index 3df47968..a35941a9 100644
--- a/docs/torch-compile.zh.md
+++ b/docs/torch-compile.zh.md
@@ -51,7 +51,7 @@ block(x, w)
 - **不要依赖 stride 原样传递。** 算子不写入的非连续输入在节点内部连续化，算子自己分配的输出恒为连续张量；后续计算需要别的布局，在算子之外自行转换。输出若就是被写入的输入（`alias`）或调用方给的 `out`，沿用那个张量的存储。
 - **不能用 meta 张量预热。** 有了边界，传入 meta 或 fake 张量的调用就在 fake 处返回，走不到构造 kernel 那一步。
 - **CUDA graph 捕获之前先行预热。** 用真实张量、相同形状至少调用一次：构造 kernel 允许编译，捕获期间只允许查表命中后直接调用。各阶段分别允许执行哪些操作，见[各阶段允许做什么](backends.md#phase-limits)。
-- **每个设备各自构造 kernel。** 设备是 kernel 记忆键的一部分，同一个实例换一块卡就重新构造一次。构造函数里指名的 `target=` 在首次编译调用中同样生效；构造失败不会把算子固定到任何 target。
+- **换一块卡可能要重新构造。** 由 target 服务的调用，设备是 kernel 记忆键的一部分，同一个实例换到另一块卡上会重新构造一次。自带 kernel 的记忆键是它的 `entry_for` 的返回值，只有构造结果与设备有关时才包含设备。构造函数里指名的 `target=` 在首次编译调用中同样生效；构造失败不会把算子固定到任何 target。
 
 ### 接入之后成立的三项保证
 
@@ -83,7 +83,7 @@ class RMSNormFwdOp(Op):
         return kernel(x, weight)
 ```
 
-声明就这么多。operator 与它的 fake 都从条目生成，每个副作用分支一个 operator：张量参数是 `signature.inputs` 的顺序，返回什么看 `signature.outputs`，写哪些参数恰好是标了 `mutated` 的输入，每个输出的形状与 dtype 取自签名。名字是 `tileops::<family>_<snake(class)>`（类名本身以 family 开头时只写一次），这里就是 `tileops::norm_rms_norm_fwd`；条目声明了写入输入或 `buffer: out` 时，对应分支的 operator 在名字后面加上 `_writes_<输入>` 或 `_out`。没有算子自己起名字，`compile_op_names` 也就不可能和注册的名字对不上。
+声明就这么多。operator 与它的 fake 都从条目生成，每个副作用分支一个 operator：张量参数是 `signature.inputs` 的顺序，返回什么看 `signature.outputs`，写哪些参数恰好是标了 `mutated` 的输入，每个输出的形状与 dtype 取自签名。名字是 `tileops::<family>_<snake(class)>`（类名本身以 family 开头时只写一次），这里就是 `tileops::norm_rms_norm_fwd`；一个分支的 operator 若写入某个输入、填写 `buffer: out` 或不产出某个输出，名字后面就按这个顺序分别加上 `_writes_<输入>`、`_out`、`_without_<输出>`。没有算子自己起名字，`compile_op_names` 也就不可能和注册的名字对不上。
 
 一次调用经过的各层，以及边界落在哪里：
 
diff --git a/mkdocs.yml b/mkdocs.yml
index 1403d9ca..718f7bcf 100644
--- a/mkdocs.yml
+++ b/mkdocs.yml
@@ -42,7 +42,8 @@ nav:
   # the sequence-model kernels built on all of the above. Pages are not one per
   # family. The FP8 lightning indexer is a section of the Attention page, where its
   # single op is used, and the sequence-modeling family is two pages, mHC and Engram,
-  # being two unrelated algorithms. Trace comes last — a tool rather than an op.
+  # being two unrelated algorithms. Top-k keeps its page before GEMM though its op is
+  # exported from `tileops.attention`. Trace comes last — a tool rather than an op.
   - API Reference:
     - api/index.md
     - Elementwise: api/elementwise.md
diff --git a/scripts/gen_bench_pages.py b/scripts/gen_bench_pages.py
index 038c1cf0..5528c8cd 100644
--- a/scripts/gen_bench_pages.py
+++ b/scripts/gen_bench_pages.py
@@ -607,9 +607,10 @@ def _bound_cell(sol: dict | None) -> str:
     "<table>",
     "<thead>",
     "<tr>",
-    # These three say what the row is; the rule after them divides that from
+    # These two say what the row is; the rule after them divides that from
     # what it measured. Each spans both header rows, so the rule has no gap.
-    '<th rowspan="2" class="colsep">Workload</th>',
+    '<th rowspan="2">Workload</th>',
+    '<th rowspan="2" class="colsep">dtype</th>',
     "<th>Ratio</th>",
     "<th>Device time</th>",
     '<th colspan="2">Alternatives</th>',
@@ -653,27 +654,60 @@ def _stack(cells: list[str]) -> str:
     return f"{head}<br>{tail}"
 
 
-WORKLOAD_CODE = "W"
+def _label_of(w: dict, m: dict) -> tuple[str, list[str]]:
+    """The manifest label a row ran under, and the dtypes its case id names.
+
+    A case id is the label followed by one dtype per dtype index the case
+    assigns, so `square-b4-float8_e4m3fn-bfloat16` is the label `square-b4` at
+    `fp8e4m3/bf16`. A row with no manifest entry strips the trailing dtype names
+    off its own id; an id that ends in none is its own label.
+    """
+    config, spec = w["config"], w.get("spec")
+    label = spec.label if spec and spec.label else None
+    if label and config.startswith(label + "-"):
+        tokens = config[len(label) + 1:].split("-")
+    elif label:
+        tokens = [spec.dtype] if spec.dtype else []
+    else:
+        parts = config.split("-")
+        tokens = []
+        while len(parts) > 1 and parts[-1] in workload_shape.DTYPE_ABBR:
+            tokens.insert(0, parts.pop())
+        label = "-".join(parts)
+        if not tokens and m.get("dtype"):
+            tokens = [m["dtype"]]
+    return label, tokens
+
+
+def _dtype_text(tokens: list[str]) -> str:
+    return "/".join(workload_shape.abbr_dtype(t) for t in tokens) or EMPTY
 
 
-def _facts(spec) -> OrderedDict:
-    """Every scalar a workload sets, as name -> (value, kind).
+def _label_html(label: str) -> str:
+    """A label that may wrap after any hyphen, and nowhere else."""
+    return html.escape(label).replace("-", "-<wbr>")
 
-    The symbols of the tensor templates come first, since those are the numbers
-    the shapes are made of; then the dimensions the manifest names outright,
-    then the parameters, which are how the op was called rather than how big it
-    is and stay dimmed wherever they are printed.
+
+def _facts(spec, dtypes=()) -> OrderedDict:
+    """Every scalar a workload sets but its dtype, as name -> (value, kind).
+
+    The symbols of the tensor templates come first, marked `sym` since a row's
+    concrete shapes already state them; then the dimensions the manifest names
+    outright, then the parameters, dimmed wherever they are printed. A parameter set to one of
+    the row's `dtypes` — an `out_dtype` — is part of the case id, and the
+    dtype column states it.
     """
     facts = OrderedDict()
-    if spec.dtype:
-        facts["dtype"] = (workload_shape.abbr_dtype(spec.dtype), "dim")
+    if not spec:
+        return facts
     if spec.symbolic:
         for name, value in spec.bindings.items():
-            facts[name] = (str(value), "dim")
+            facts[name] = (str(value), "sym")
     for name, value in spec.dims:
         facts.setdefault(name, (value, "dim"))
     for name, value in spec.params:
-        facts[name] = (value, "param")
+        if value not in dtypes:
+            facts[name] = (value, "param")
     return facts
 
 
@@ -709,10 +743,13 @@ def _cells(entries) -> str:
 
 _TAG = __import__("re").compile(r"<[^>]+>")
 
-# How wide a line of entries may be before it is set differently, in characters
-# of the key's monospace against the content column.
-WRAP_TENSORS = 96   # a tensor list wraps -> one tensor per line
-WRAP_DELTA = 74     # a row's own scalars wrap -> id and scalars on two lines
+# How wide a tensor list may be before it is set one tensor per line, in
+# characters of the key's monospace against the content column.
+WRAP_TENSORS = 96
+
+
+# Past this a label wraps in the key's label column; its cluster stacks instead.
+LONG_LABEL = 32
 
 
 def _width(entries) -> int:
@@ -720,87 +757,131 @@ def _width(entries) -> int:
     return sum(len(_TAG.sub("", body)) + 3 for body, _ in entries)
 
 
-def _cluster(rows: list) -> list:
-    """An op's workloads, grouped by what they have in common.
+def _tensors_of(specs: list, bases: list, symbolic: bool = True):
+    """One tensor list for every dtype a label ran at, or None if none exists.
 
-    The key is the description with every size taken out: which tensors, named
-    together how, in the manifest's symbols. The dtype is not part of it — a
-    workload in `f16` beside one in `bf16` carries that as one more thing that
-    varies. A tensor in a dtype of its own does split the group, since then the
-    tensor list itself differs.
+    Returns `(kind, [(names, shape, dtype)])`: `kind` is `"sym"` where the
+    shapes are the manifest's symbols, `"con"` where they are concrete. A
+    tensor states a dtype where it is the same in every row and not the row's
+    own — an `f32` bias beside `bf16` and `f16` inputs — and none where it
+    follows the row. A row's own dtype is `bases[i]`, the first one its dtype
+    column shows, so a case named by an `out_dtype` marks every input not in
+    it. Any other difference between the rows leaves no one list true for all
+    of them. `symbolic=False` asks for the concrete list alone.
     """
-    groups = OrderedDict()
-    for code, w in rows:
-        spec = w.get("spec")
-        if not spec:
-            key = None
-        elif spec.symbolic:
-            key = ("sym", *(c for c, _ in _tensors_html(spec.symbolic,
-                                                        spec.dtype)))
+    if not all(specs):
+        return None
+    lists = [("sym", s.symbolic) if symbolic and s.symbolic else ("con", s.tensors)
+             for s in specs]
+    kind, first = lists[0]
+    if any(k != kind or [e[:2] for e in es] != [e[:2] for e in first]
+           for k, es in lists):
+        return None
+    out = []
+    for i, (names, shape, _) in enumerate(first):
+        dtypes = [es[i][2] or s.dtype for s, (_, es) in zip(specs, lists, strict=True)]
+        if all(d == b for d, b in zip(dtypes, bases, strict=True)):
+            out.append((names, shape, None))
+        elif len(set(dtypes)) == 1:
+            out.append((names, shape, dtypes[0]))
         else:
-            # No template to write the shapes in: the row prints its own, so
-            # the group is only there to hold workloads taking the same tensors.
-            key = ("con", *(n for n, _, _ in spec.tensors))
-        groups.setdefault(key, []).append((code, w))
-    return list(groups.values())
+            return None
+    return kind, out
+
 
+def label_groups(workloads: list, metrics: list) -> list:
+    """An op's rows, one group per label, groups clustered by shape.
 
-def workload_key(rows: list) -> list:
-    """What each row of an op's table ran on, listed above it.
+    Returns clusters of groups, each group a dict of `label`, `facts` (the
+    scalars it was set to), `tensors` (see `_tensors_of`), `concrete` (the same
+    tensors with the symbols substituted) and `rows`, a list of
+    `(dtype, metrics)`. The snapshot lists an op's rows in the manifest's
+    order, and every level keeps the order it first meets its members in, so a
+    size sweep reads in the order the manifest declares it.
 
-    An op's workloads are nearly the same run at several sizes, so what they
-    share is stated once — in the manifest's own symbols, `q k [B, H, DK]` — and
-    each row carries only the symbols that vary on it. Inside the table's first
-    column the same shapes would squeeze the measurements off the window.
+    A label whose dtypes differ in anything but the dtype — another shape,
+    another symbol value — is one group per description, each under the same
+    label: folding them would state one row's shapes for another's numbers.
 
-    Order follows the table's rows, so `W3` is the third row.
+    A cluster holds the groups written in the same symbols: which tensors,
+    named together how. The dtype is not part of it. Groups with concrete
+    shapes cluster by tensor names alone, since each prints its own shapes.
     """
-    if not rows:
+    buckets = OrderedDict()
+    for w, m in zip(workloads, metrics, strict=True):
+        label, dtypes = _label_of(w, m)
+        spec = w.get("spec")
+        facts = _facts(spec, dtypes)
+        base = dtypes[0] if dtypes else spec and spec.dtype
+        buckets.setdefault((label, tuple(facts.items())), []).append(
+            (spec, facts, _dtype_text(dtypes), m, base))
+
+    groups = []
+    for (label, _), members in buckets.items():
+        parts = [members]
+        if _tensors_of([mb[0] for mb in members], [mb[4] for mb in members]) is None:
+            split = OrderedDict()
+            for member in members:
+                spec = member[0]
+                key = spec and (spec.symbolic, spec.tensors, spec.dtype, member[4])
+                split.setdefault(repr(key), []).append(member)
+            parts = list(split.values())
+        for part in parts:
+            specs, bases = [mb[0] for mb in part], [mb[4] for mb in part]
+            groups.append({"label": label, "facts": part[0][1],
+                           "tensors": _tensors_of(specs, bases),
+                           "concrete": _tensors_of(specs, bases, symbolic=False),
+                           "rows": [(mb[2], mb[3]) for mb in part]})
+
+    clusters = OrderedDict()
+    for group in groups:
+        t = group["tensors"]
+        key = (None if t is None else
+               (t[0], *t[1]) if t[0] == "sym" else (t[0], *(n for n, _, _ in t[1])))
+        clusters.setdefault(key, []).append(group)
+    return list(clusters.values())
+
+
+def workload_key(clusters: list) -> list:
+    """What each label of an op's table ran on, listed above it.
+
+    Labels written in the same symbols share one template, `q, k: [B, H, DK]`,
+    stated once; each label prints its shapes substituted, `q, k: [1, 8, 128]`.
+    What every label of a cluster shares goes above them, only when there are
+    two or more. Order follows the table's row groups.
+    """
+    if not clusters:
         return []
     blocks = []
-    for group in _cluster(rows):
-        specs = [w.get("spec") for _, w in group]
-        if not specs[0]:
-            # No manifest entry: the id is all there is to say. Still a group
-            # and still a list, so the code column lines up with every other
-            # op's and the styling is the one thing that does not vary.
-            bare = "".join(
-                f'<li><b>{code}</b><span class="wl-delta"></span>'
-                f'<code class="wl-id">{html.escape(w["config"])}</code></li>'
-                for code, w in group)
-            blocks.append('<div class="wl-group"><ul class="wl-rows">'
-                          + bare + "</ul></div>")
-            continue
-        facts = [_facts(s) for s in specs]
-        # A symbolic template describes the whole group only where every
-        # workload in it resolved to the same symbols.
-        symbolic = specs[0].symbolic
-        if not all(s.symbolic == symbolic for s in specs):
-            symbolic = None
-        shared = [n for n in facts[0]
-                  if all(f.get(n) == facts[0][n] for f in facts)]
-
-        head = _tensors_html(symbolic, specs[0].dtype) if symbolic else []
+    for cluster in clusters:
+        facts = [OrderedDict((n, v) for n, v in g["facts"].items()
+                             if not (g["concrete"] and v[1] == "sym"))
+                 for g in cluster]
+        for f, g in zip(facts, cluster, strict=True):
+            dtypes = list(OrderedDict.fromkeys(d for d, _ in g["rows"]))
+            f["dtype"] = (", ".join(dtypes), "dim")
+        shared = ([n for n in facts[0]
+                   if all(f.get(n) == facts[0][n] for f in facts)]
+                  if len(cluster) > 1 else [])
+        # One cluster is one symbolic template, stated once above its labels.
+        tensors = cluster[0]["tensors"]
+        symbolic = tensors is not None and tensors[0] == "sym"
+        head = _tensors_html(tensors[1], None) if symbolic else []
         scalars = [_fact_html(n, *facts[0][n]) for n in shared]
 
-        rows_html, deltas = [], []
-        for (code, _), spec, fact in zip(group, specs, facts, strict=True):
-            varies = [_fact_html(n, *v) for n, v in fact.items()
-                      if n not in shared]
-            # Without a template the group has nothing to hold in common, so
-            # the row states its own tensors outright.
-            if not symbolic:
-                varies = _tensors_html(spec.tensors, spec.dtype) + varies
-            deltas.append(varies)
+        rows_html = []
+        for group, fact in zip(cluster, facts, strict=True):
+            # A long row breaks between tensors and scalars before inside either.
+            tensor_cells = (_tensors_html(group["concrete"][1], None)
+                            if group["concrete"] else [])
+            scalar_cells = [_fact_html(n, *v) for n, v in fact.items()
+                            if n not in shared]
+            parts = "".join(f'<span class="wl-part">{_cells(c)}</span>'
+                            for c in (tensor_cells, scalar_cells) if c)
+            flow = f'<span class="wl-flow">{parts}</span>' if parts else ""
             rows_html.append(
-                f'<li><b>{code}</b><span class="wl-delta">'
-                + _cells(varies) +
-                f'</span><code class="wl-id">{html.escape(spec.label)}</code></li>')
-        # Where one row's scalars are wider than the column, the whole group is
-        # set on two lines — id under the code, scalars under that — so every
-        # row in it breaks the same way.
-        long = (" wl-long" if any(_width(d) > WRAP_DELTA for d in deltas)
-                else "")
+                f'<li><code class="wl-id">{_label_html(group["label"])}</code>'
+                f'<span class="wl-delta">{flow}</span></li>')
 
         block = []
         # Tensors and scalars on lines of their own: what the op takes, then how
@@ -811,12 +892,17 @@ def workload_key(rows: list) -> list:
             block.append(f'<p class="wl-shared{stack}">{_cells(head)}</p>')
         if scalars:
             block.append(f'<p class="wl-shared">{_cells(scalars)}</p>')
-        block.append(f'<ul class="wl-rows{long}">' + "".join(rows_html) + "</ul>")
+        long = any(len(g["label"]) > LONG_LABEL for g in cluster)
+        block.append(f'<ul class="wl-rows{" wl-long" if long else ""}">'
+                     + "".join(rows_html) + "</ul>")
         blocks.append(f'<div class="wl-group">{"".join(block)}</div>')
     return ['<div class="wl-key">', *blocks, "</div>", ""]
 
 
-def detail_row(code: str, m: dict) -> str:
+def detail_row(dtype: str, m: dict, label: str | None = None,
+               span: int = 1) -> str:
+    """One measured row. The first row of a label group carries the label,
+    spanning the group's dtype rows."""
     ordered = sorted(m["rivals"].items(), key=lambda kv: kv[1]["busy_ms"])
     names = _stack([f"<code>{html.escape(t)}</code>" for t, _ in ordered])
     times = _stack([_sig_ms(r["busy_ms"]) for _, r in ordered])
@@ -827,9 +913,15 @@ def detail_row(code: str, m: dict) -> str:
     gap = _ratio_cell(real[0]["speedup"] if real else
                       weak[0]["speedup"] if weak else None,
                       rated=bool(real))
+    name = ""
+    if label is not None:
+        rowspan = f' rowspan="{span}"' if span > 1 else ""
+        name = (f'<td class="wl-name"{rowspan}>'
+                f"<code>{_label_html(label)}</code></td>")
     return (
         "<tr>"
-        f'<td class="colsep"><b>{code}</b></td>'
+        f"{name}"
+        f'<td class="colsep">{html.escape(dtype)}</td>'
         f"<td>{gap}</td>"
         f"<td>{_sig_ms(m['busy_ms'])}</td>"
         f"<td>{names}</td>"
@@ -994,15 +1086,20 @@ def reading_page(sol_engine=(None, None)) -> str:
         "## Columns", "",
         "| Column | Meaning |",
         "| --- | --- |",
-        "| **Workload** | `W1`, `W2`, … — the key above each table spells each "
-        "one out: the benchmark's own id for it, the dtype it ran at, and every "
-        "input tensor as `name: shape, dtype`. Tensors sharing a shape are "
-        "named together, and each carries its own dtype, so a `mask` in `bool` "
-        "says so where it is read. After the tensors come the dimensions the op "
-        "is sized by rather than shaped by (`m`, `n`, `k` for a GEMM, "
-        "`num_experts` for MoE routing), then dimmed, the parameters the call "
-        "did not leave at the signature's default. A quantity the others "
-        "already fix — `max_seqlen_q` is `max(q_lens)` — is not repeated. |",
+        "| **Workload** | The label the TileOPs spec manifest gives the "
+        "workload, once per group of rows that ran it at several dtypes. The "
+        "key above each table spells each label out: every input tensor as "
+        "`name: shape, dtype`, tensors sharing a shape named together, and a "
+        "tensor in a dtype of its own — a `mask` in `bool` — saying so where it "
+        "is read, and a template shared by several labels, `q, k: [B, H, DK]`, "
+        "stated once above them. After the tensors come the dimensions the op is sized by "
+        "rather than shaped by (`m`, `n`, `k` for a GEMM, `num_experts` for MoE "
+        "routing), then dimmed, the parameters the call did not leave at the "
+        "signature's default, then the dtypes the label ran at. A quantity the "
+        "others already fix — `max_seqlen_q` is `max(q_lens)` — is not "
+        "repeated. Labels and rows keep the manifest's order. |",
+        "| **dtype** | The dtype the row ran at; a case with two dtype indices "
+        "names both, `fp8e4m3/bf16`. |",
         "| **Ratio** | `alt / ours` — the fastest alternative's device time "
         "divided by ours, the one number the colour grades. |",
         "| **Device time** | Milliseconds the device spent executing the call's "
@@ -1144,17 +1241,17 @@ def data_page(title: str, fams: list[str], rows_by_fam: dict,
             # count is the length of the list under it, and a tick on every op
             # says nothing. Only a mark that warns survives.
             warn = f" <small>{tmark}</small>" if tmark in ("❌", "⏭️") else ""
-            ordered = sorted(zip(workloads_of[op], metrics_by_op[op], strict=True),
-                             key=lambda z: z[0]["config"])
-            coded = [(f"{WORKLOAD_CODE}{i}", w)
-                     for i, (w, _) in enumerate(ordered, 1)]
+            clusters = label_groups(workloads_of[op], metrics_by_op[op])
             lines += [f"{op_h} {_op_cell(op, module, ref)}{warn}",
-                      "", *workload_key(coded),
+                      "", *workload_key(clusters),
                       # No `markdown="1"`, and no blank line until `</div>`: a
                       # blank line would end the raw-HTML block mid-table.
                       '<div class="datatable">', *DETAIL_HEADER]
-            for (code, _), (_, m) in zip(coded, ordered, strict=True):
-                lines.append(detail_row(code, m))
+            for group in (g for c in clusters for g in c):
+                rows = group["rows"]
+                for i, (dtype, m) in enumerate(rows):
+                    lines.append(detail_row(dtype, m, group["label"] if i == 0 else None,
+                                            len(rows)))
             lines += [*DETAIL_FOOTER, "</div>", ""]
     return "\n".join(lines) + "\n"
 
diff --git a/scripts/workload_shape.py b/scripts/workload_shape.py
index 1af6ea4c..b8837d10 100644
--- a/scripts/workload_shape.py
+++ b/scripts/workload_shape.py
@@ -41,7 +41,7 @@
     "float16": "f16", "bfloat16": "bf16", "float32": "f32", "float64": "f64",
     "float8_e4m3fn": "fp8e4m3", "float8_e5m2": "fp8e5m2",
     "int8": "i8", "int16": "i16", "int32": "i32", "int64": "i64",
-    "uint8": "u8", "bool": "bool",
+    "uint8": "u8", "bool": "bool", "complex64": "c64", "complex128": "c128",
 }
 
 
diff --git a/tests/fixtures/bench_results.xml b/tests/fixtures/bench_results.xml
index ca9fe698..1cfd1852 100644
--- a/tests/fixtures/bench_results.xml
+++ b/tests/fixtures/bench_results.xml
@@ -23,6 +23,23 @@
         <property name="torch-ref_ratio" value="11.39"/>
       </properties>
     </testcase>
+    <testcase classname="bench" name="test_bench[decode-b1-h8-float16]">
+      <properties>
+        <property name="op" value="DeltaDecodeFwdOp"/>
+        <property name="op_module" value="tileops.ops.linear_attention.delta"/>
+        <property name="tileops_device_busy_ms" value="0.0033"/>
+        <property name="tileops_latency_ms" value="0.0090"/>
+        <property name="tileops_tflops" value="0.508"/>
+        <property name="tileops_flops" value="1572864"/>
+        <property name="tileops_bytes" value="8388608"/>
+        <property name="tileops_timing" value="cupti"/>
+        <property name="tileops_n_kernels" value="1"/>
+        <property name="fla_device_busy_ms" value="0.0142"/>
+        <property name="fla_ratio" value="4.30"/>
+        <property name="torch-ref_device_busy_ms" value="0.0353"/>
+        <property name="torch-ref_ratio" value="10.70"/>
+      </properties>
+    </testcase>
 
     <!-- Behind the same library kernel. -->
     <testcase classname="bench" name="test_bench[decode-b8-h8-bfloat16]">
diff --git a/tests/fixtures/manifest/ops.yaml b/tests/fixtures/manifest/ops.yaml
index b4449129..c9748bca 100644
--- a/tests/fixtures/manifest/ops.yaml
+++ b/tests/fixtures/manifest/ops.yaml
@@ -17,7 +17,7 @@ DeltaDecodeFwdOp:
       - "q.shape == (B, H, DK)"
   workloads:
     - {q_shape: [1, 8, 128], k_shape: [1, 8, 128], v_shape: [1, 8, 128],
-       state_shape: [1, 8, 128, 128], dtypes: [bfloat16], label: decode-b1-h8}
+       state_shape: [1, 8, 128, 128], dtypes: [bfloat16, float16], label: decode-b1-h8}
     - {q_shape: [8, 8, 128], k_shape: [8, 8, 128], v_shape: [8, 8, 128],
        state_shape: [8, 8, 128, 128], dtypes: [bfloat16], label: decode-b8-h8}
 
diff --git a/tests/golden/elementwise.md b/tests/golden/elementwise.md
index 7a221287..f606eb83 100644
--- a/tests/golden/elementwise.md
+++ b/tests/golden/elementwise.md
@@ -7,14 +7,15 @@ One table per op, one row per workload. `Ratio` is the fastest other implementat
 ## [MysteryFwd](https://github.com/tile-ai/TileOPs/search?q=repo%3Atile-ai%2FTileOPs+MysteryFwdOp&type=code) <small>⏭️</small>
 
 <div class="wl-key">
-<div class="wl-group"><ul class="wl-rows"><li><b>W1</b><span class="wl-delta"></span><code class="wl-id">undeclared-op-case-float16</code></li></ul></div>
+<div class="wl-group"><ul class="wl-rows"><li><code class="wl-id">undeclared-<wbr>op-<wbr>case</code><span class="wl-delta"><span class="wl-flow"><span class="wl-part"><span class="wl-cell wl-scalar"><span class="wl-k">dtype</span>=<span class="wl-v">f16</span></span></span></span></span></li></ul></div>
 </div>
 
 <div class="datatable">
 <table>
 <thead>
 <tr>
-<th rowspan="2" class="colsep">Workload</th>
+<th rowspan="2">Workload</th>
+<th rowspan="2" class="colsep">dtype</th>
 <th>Ratio</th>
 <th>Device time</th>
 <th colspan="2">Alternatives</th>
@@ -33,7 +34,7 @@ One table per op, one row per workload. `Ratio` is the fastest other implementat
 </tr>
 </thead>
 <tbody>
-<tr><td class="colsep"><b>W1</b></td><td><span class="perf-ahead">2.00×</span></td><td>0.004</td><td><code>triton</code><br><span class="alt-slow"><code>brand-new-lib</code></span></td><td>0.008<br><span class="alt-slow">0.0120</span></td><td>·</td><td>·</td><td>·</td></tr>
+<tr><td class="wl-name"><code>undeclared-<wbr>op-<wbr>case</code></td><td class="colsep">f16</td><td><span class="perf-ahead">2.00×</span></td><td>0.004</td><td><code>triton</code><br><span class="alt-slow"><code>brand-new-lib</code></span></td><td>0.008<br><span class="alt-slow">0.0120</span></td><td>·</td><td>·</td><td>·</td></tr>
 </tbody>
 </table>
 </div>
@@ -41,14 +42,15 @@ One table per op, one row per workload. `Ratio` is the fastest other implementat
 ## [SquareFwd](https://github.com/tile-ai/TileOPs/search?q=repo%3Atile-ai%2FTileOPs+SquareFwdOp&type=code)
 
 <div class="wl-key">
-<div class="wl-group"><p class="wl-shared"><span class="wl-cell wl-scalar"><span class="wl-k">dtype</span>=<span class="wl-v">f16</span></span></p><ul class="wl-rows"><li><b>W1</b><span class="wl-delta"><span class="wl-cell wl-tensor"><span class="wl-k">a</span>: [64, 32]</span></span><code class="wl-id">oblong</code></li></ul></div>
+<div class="wl-group"><ul class="wl-rows"><li><code class="wl-id">oblong</code><span class="wl-delta"><span class="wl-flow"><span class="wl-part"><span class="wl-cell wl-tensor"><span class="wl-k">a</span>: [64, 32]</span></span><span class="wl-part"><span class="wl-cell wl-scalar"><span class="wl-k">dtype</span>=<span class="wl-v">f16</span></span></span></span></span></li></ul></div>
 </div>
 
 <div class="datatable">
 <table>
 <thead>
 <tr>
-<th rowspan="2" class="colsep">Workload</th>
+<th rowspan="2">Workload</th>
+<th rowspan="2" class="colsep">dtype</th>
 <th>Ratio</th>
 <th>Device time</th>
 <th colspan="2">Alternatives</th>
@@ -67,7 +69,7 @@ One table per op, one row per workload. `Ratio` is the fastest other implementat
 </tr>
 </thead>
 <tbody>
-<tr><td class="colsep"><b>W1</b></td><td><span class="perf-ahead">1.50×</span></td><td>0.002</td><td><code>torch</code></td><td>0.003</td><td>·</td><td>·</td><td>·</td></tr>
+<tr><td class="wl-name"><code>oblong</code></td><td class="colsep">f16</td><td><span class="perf-ahead">1.50×</span></td><td>0.002</td><td><code>torch</code></td><td>0.003</td><td>·</td><td>·</td><td>·</td></tr>
 </tbody>
 </table>
 </div>
@@ -75,14 +77,15 @@ One table per op, one row per workload. `Ratio` is the fastest other implementat
 ## [TemplatedFwd](https://github.com/tile-ai/TileOPs/search?q=repo%3Atile-ai%2FTileOPs+TemplatedFwdOp&type=code)
 
 <div class="wl-key">
-<div class="wl-group"><p class="wl-shared"><span class="wl-cell wl-tensor"><span class="wl-k">x</span>: [rows, cols]</span><span class="wl-cell wl-tensor"><span class="wl-k">mask</span>: [rows], <span class="wl-dt">bool</span></span></p><p class="wl-shared"><span class="wl-cell wl-scalar"><span class="wl-k">dtype</span>=<span class="wl-v">f16</span></span><span class="wl-cell wl-scalar"><span class="wl-k">cols</span>=<span class="wl-v">256</span></span></p><ul class="wl-rows"><li><b>W1</b><span class="wl-delta"><span class="wl-cell wl-scalar"><span class="wl-k">rows</span>=<span class="wl-v">128</span></span></span><code class="wl-id">templated-128x256</code></li><li><b>W2</b><span class="wl-delta"><span class="wl-cell wl-scalar"><span class="wl-k">rows</span>=<span class="wl-v">64</span></span></span><code class="wl-id">templated-64x256</code></li></ul></div>
+<div class="wl-group"><p class="wl-shared"><span class="wl-cell wl-tensor"><span class="wl-k">x</span>: [rows, cols]</span><span class="wl-cell wl-tensor"><span class="wl-k">mask</span>: [rows], <span class="wl-dt">bool</span></span></p><p class="wl-shared"><span class="wl-cell wl-scalar"><span class="wl-k">dtype</span>=<span class="wl-v">f16</span></span></p><ul class="wl-rows"><li><code class="wl-id">templated-<wbr>64x256</code><span class="wl-delta"><span class="wl-flow"><span class="wl-part"><span class="wl-cell wl-tensor"><span class="wl-k">x</span>: [64, 256]</span><span class="wl-cell wl-tensor"><span class="wl-k">mask</span>: [64], <span class="wl-dt">bool</span></span></span></span></span></li><li><code class="wl-id">templated-<wbr>128x256</code><span class="wl-delta"><span class="wl-flow"><span class="wl-part"><span class="wl-cell wl-tensor"><span class="wl-k">x</span>: [128, 256]</span><span class="wl-cell wl-tensor"><span class="wl-k">mask</span>: [128], <span class="wl-dt">bool</span></span></span></span></span></li></ul></div>
 </div>
 
 <div class="datatable">
 <table>
 <thead>
 <tr>
-<th rowspan="2" class="colsep">Workload</th>
+<th rowspan="2">Workload</th>
+<th rowspan="2" class="colsep">dtype</th>
 <th>Ratio</th>
 <th>Device time</th>
 <th colspan="2">Alternatives</th>
@@ -101,8 +104,8 @@ One table per op, one row per workload. `Ratio` is the fastest other implementat
 </tr>
 </thead>
 <tbody>
-<tr><td class="colsep"><b>W1</b></td><td><span class="perf-ahead">1.43×</span></td><td>0.007</td><td><code>torch</code></td><td>0.0100</td><td>·</td><td>·</td><td>·</td></tr>
-<tr><td class="colsep"><b>W2</b></td><td><span class="perf-ahead">1.50×</span></td><td>0.006</td><td><code>torch</code></td><td>0.009</td><td>·</td><td>·</td><td>·</td></tr>
+<tr><td class="wl-name"><code>templated-<wbr>64x256</code></td><td class="colsep">f16</td><td><span class="perf-ahead">1.50×</span></td><td>0.006</td><td><code>torch</code></td><td>0.009</td><td>·</td><td>·</td><td>·</td></tr>
+<tr><td class="wl-name"><code>templated-<wbr>128x256</code></td><td class="colsep">f16</td><td><span class="perf-ahead">1.43×</span></td><td>0.007</td><td><code>torch</code></td><td>0.0100</td><td>·</td><td>·</td><td>·</td></tr>
 </tbody>
 </table>
 </div>
diff --git a/tests/golden/index.md b/tests/golden/index.md
index cbbc83f6..46ec1f0e 100644
--- a/tests/golden/index.md
+++ b/tests/golden/index.md
@@ -2,7 +2,7 @@
 
 !!! info "Nightly snapshot"
 
-    **GPU** NVIDIA H200 · **commit** [`0123456789ab`](https://github.com/tile-ai/TileOPs/commit/0123456789abcdef0123456789abcdef01234567) · **run date** 2026-01-01 · **5 ops**, 8 workloads
+    **GPU** NVIDIA H200 · **commit** [`0123456789ab`](https://github.com/tile-ai/TileOPs/commit/0123456789abcdef0123456789abcdef01234567) · **run date** 2026-01-01 · **5 ops**, 9 workloads
     · [nightly run](https://github.com/tile-ai/TileOPs/actions/runs/1234567890)
 
 ## Environment
@@ -32,5 +32,5 @@ Not published by this run: `image`, `driver`, `cuda`, `torch`, `tilelang`.
 | Page | Ops | Workloads |
 | --- | --- | --- |
 | [Elementwise](elementwise.md) | 3 | 4 |
-| [Linear Attention](linear-attention.md) | 1 | 2 |
+| [Linear Attention](linear-attention.md) | 1 | 3 |
 | [SSM](ssm.md) | 1 | 2 |
diff --git a/tests/golden/linear-attention.md b/tests/golden/linear-attention.md
index bb6823a6..9cf9a938 100644
--- a/tests/golden/linear-attention.md
+++ b/tests/golden/linear-attention.md
@@ -1,20 +1,21 @@
 # Linear Attention
 
-**1 ops, 2 workloads.**
+**1 ops, 3 workloads.**
 
 One table per op, one row per workload. `Ratio` is the fastest other implementation's device time divided by ours, so <span class="perf-ahead">green</span> is faster than it, <span class="perf-par">plain</span> is level with it, <span class="perf-behind">red</span> is slower. Times are in ms. [How these numbers are taken](reading.md).
 
 ## [DeltaDecodeFwd](https://github.com/tile-ai/TileOPs/search?q=repo%3Atile-ai%2FTileOPs+DeltaDecodeFwdOp&type=code)
 
 <div class="wl-key">
-<div class="wl-group"><p class="wl-shared"><span class="wl-cell wl-tensor"><span class="wl-k">q, k</span>: [B, H, DK]</span><span class="wl-cell wl-tensor"><span class="wl-k">v</span>: [B, H, DV]</span><span class="wl-cell wl-tensor"><span class="wl-k">state</span>: [B, H, DK, DV]</span></p><p class="wl-shared"><span class="wl-cell wl-scalar"><span class="wl-k">dtype</span>=<span class="wl-v">bf16</span></span><span class="wl-cell wl-scalar"><span class="wl-k">H</span>=<span class="wl-v">8</span></span><span class="wl-cell wl-scalar"><span class="wl-k">DK</span>=<span class="wl-v">128</span></span><span class="wl-cell wl-scalar"><span class="wl-k">DV</span>=<span class="wl-v">128</span></span></p><ul class="wl-rows"><li><b>W1</b><span class="wl-delta"><span class="wl-cell wl-scalar"><span class="wl-k">B</span>=<span class="wl-v">1</span></span></span><code class="wl-id">decode-b1-h8</code></li><li><b>W2</b><span class="wl-delta"><span class="wl-cell wl-scalar"><span class="wl-k">B</span>=<span class="wl-v">8</span></span></span><code class="wl-id">decode-b8-h8</code></li></ul></div>
+<div class="wl-group"><p class="wl-shared"><span class="wl-cell wl-tensor"><span class="wl-k">q, k</span>: [B, H, DK]</span><span class="wl-cell wl-tensor"><span class="wl-k">v</span>: [B, H, DV]</span><span class="wl-cell wl-tensor"><span class="wl-k">state</span>: [B, H, DK, DV]</span></p><ul class="wl-rows"><li><code class="wl-id">decode-<wbr>b1-<wbr>h8</code><span class="wl-delta"><span class="wl-flow"><span class="wl-part"><span class="wl-cell wl-tensor"><span class="wl-k">q, k, v</span>: [1, 8, 128]</span><span class="wl-cell wl-tensor"><span class="wl-k">state</span>: [1, 8, 128, 128]</span></span><span class="wl-part"><span class="wl-cell wl-scalar"><span class="wl-k">dtype</span>=<span class="wl-v">bf16, f16</span></span></span></span></span></li><li><code class="wl-id">decode-<wbr>b8-<wbr>h8</code><span class="wl-delta"><span class="wl-flow"><span class="wl-part"><span class="wl-cell wl-tensor"><span class="wl-k">q, k, v</span>: [8, 8, 128]</span><span class="wl-cell wl-tensor"><span class="wl-k">state</span>: [8, 8, 128, 128]</span></span><span class="wl-part"><span class="wl-cell wl-scalar"><span class="wl-k">dtype</span>=<span class="wl-v">bf16</span></span></span></span></span></li></ul></div>
 </div>
 
 <div class="datatable">
 <table>
 <thead>
 <tr>
-<th rowspan="2" class="colsep">Workload</th>
+<th rowspan="2">Workload</th>
+<th rowspan="2" class="colsep">dtype</th>
 <th>Ratio</th>
 <th>Device time</th>
 <th colspan="2">Alternatives</th>
@@ -33,8 +34,9 @@ One table per op, one row per workload. `Ratio` is the fastest other implementat
 </tr>
 </thead>
 <tbody>
-<tr><td class="colsep"><b>W1</b></td><td><span class="perf-ahead">4.58×</span></td><td>0.0031</td><td><code>fla</code><br><span class="alt-slow"><code>torch-ref</code></span></td><td>0.0142<br><span class="alt-slow">0.0353</span></td><td>0.508</td><td>·</td><td>·</td></tr>
-<tr><td class="colsep"><b>W2</b></td><td><span class="perf-behind">0.50×</span></td><td>0.0200</td><td><code>fla</code></td><td>0.0100</td><td>0.629</td><td>·</td><td>·</td></tr>
+<tr><td class="wl-name" rowspan="2"><code>decode-<wbr>b1-<wbr>h8</code></td><td class="colsep">bf16</td><td><span class="perf-ahead">4.58×</span></td><td>0.0031</td><td><code>fla</code><br><span class="alt-slow"><code>torch-ref</code></span></td><td>0.0142<br><span class="alt-slow">0.0353</span></td><td>0.508</td><td>·</td><td>·</td></tr>
+<tr><td class="colsep">f16</td><td><span class="perf-ahead">4.30×</span></td><td>0.0033</td><td><code>fla</code><br><span class="alt-slow"><code>torch-ref</code></span></td><td>0.0142<br><span class="alt-slow">0.0353</span></td><td>0.508</td><td>·</td><td>·</td></tr>
+<tr><td class="wl-name"><code>decode-<wbr>b8-<wbr>h8</code></td><td class="colsep">bf16</td><td><span class="perf-behind">0.50×</span></td><td>0.0200</td><td><code>fla</code></td><td>0.0100</td><td>0.629</td><td>·</td><td>·</td></tr>
 </tbody>
 </table>
 </div>
diff --git a/tests/golden/ssm.md b/tests/golden/ssm.md
index bc9f7a41..3af9b0ec 100644
--- a/tests/golden/ssm.md
+++ b/tests/golden/ssm.md
@@ -7,14 +7,15 @@ One table per op, one row per workload. `Ratio` is the fastest other implementat
 ## [ChunkScanFwd](https://github.com/tile-ai/TileOPs/search?q=repo%3Atile-ai%2FTileOPs+ChunkScanFwdOp&type=code) <small>❌</small>
 
 <div class="wl-key">
-<div class="wl-group"><p class="wl-shared"><span class="wl-cell wl-scalar"><span class="wl-k">dtype</span>=<span class="wl-v">bf16</span></span><span class="wl-cell wl-scalar"><span class="wl-k">num_chunks</span>=<span class="wl-v">4</span></span><span class="wl-cell wl-scalar"><span class="wl-k">chunk_len</span>=<span class="wl-v">64</span></span><span class="wl-cell wl-scalar"><span class="wl-k">N</span>=<span class="wl-v">128</span></span></p><ul class="wl-rows"><li><b>W1</b><span class="wl-delta"><span class="wl-cell wl-scalar"><span class="wl-k">batch</span>=<span class="wl-v">2</span></span></span><code class="wl-id">scan-b2</code></li><li><b>W2</b><span class="wl-delta"><span class="wl-cell wl-scalar"><span class="wl-k">batch</span>=<span class="wl-v">4</span></span><span class="wl-cell wl-scalar"><span class="wl-dim"><span class="wl-k">is_causal</span>=<span class="wl-v">false</span></span></span></span><code class="wl-id">scan-b4</code></li></ul></div>
+<div class="wl-group"><p class="wl-shared"><span class="wl-cell wl-scalar"><span class="wl-k">num_chunks</span>=<span class="wl-v">4</span></span><span class="wl-cell wl-scalar"><span class="wl-k">chunk_len</span>=<span class="wl-v">64</span></span><span class="wl-cell wl-scalar"><span class="wl-k">N</span>=<span class="wl-v">128</span></span><span class="wl-cell wl-scalar"><span class="wl-k">dtype</span>=<span class="wl-v">bf16</span></span></p><ul class="wl-rows"><li><code class="wl-id">scan-<wbr>b2</code><span class="wl-delta"><span class="wl-flow"><span class="wl-part"><span class="wl-cell wl-scalar"><span class="wl-k">batch</span>=<span class="wl-v">2</span></span></span></span></span></li><li><code class="wl-id">scan-<wbr>b4</code><span class="wl-delta"><span class="wl-flow"><span class="wl-part"><span class="wl-cell wl-scalar"><span class="wl-k">batch</span>=<span class="wl-v">4</span></span><span class="wl-cell wl-scalar"><span class="wl-dim"><span class="wl-k">is_causal</span>=<span class="wl-v">false</span></span></span></span></span></span></li></ul></div>
 </div>
 
 <div class="datatable">
 <table>
 <thead>
 <tr>
-<th rowspan="2" class="colsep">Workload</th>
+<th rowspan="2">Workload</th>
+<th rowspan="2" class="colsep">dtype</th>
 <th>Ratio</th>
 <th>Device time</th>
 <th colspan="2">Alternatives</th>
@@ -33,8 +34,8 @@ One table per op, one row per workload. `Ratio` is the fastest other implementat
 </tr>
 </thead>
 <tbody>
-<tr><td class="colsep"><b>W1</b></td><td><span class="perf-unrated">4.00×</span></td><td>0.0500</td><td><code>torch-ref</code></td><td>0.2000</td><td>·</td><td>·</td><td>·</td></tr>
-<tr><td class="colsep"><b>W2</b></td><td><span class="perf-ahead">1.60×</span></td><td>0.1000</td><td><code>torch</code></td><td>0.1000</td><td>·</td><td>·</td><td>·</td></tr>
+<tr><td class="wl-name"><code>scan-<wbr>b2</code></td><td class="colsep">bf16</td><td><span class="perf-unrated">4.00×</span></td><td>0.0500</td><td><code>torch-ref</code></td><td>0.2000</td><td>·</td><td>·</td><td>·</td></tr>
+<tr><td class="wl-name"><code>scan-<wbr>b4</code></td><td class="colsep">bf16</td><td><span class="perf-ahead">1.60×</span></td><td>0.1000</td><td><code>torch</code></td><td>0.1000</td><td>·</td><td>·</td><td>·</td></tr>
 </tbody>
 </table>
 </div>
diff --git a/tests/test_renderer.py b/tests/test_renderer.py
index 595a765a..e64e5bef 100644
--- a/tests/test_renderer.py
+++ b/tests/test_renderer.py
@@ -79,9 +79,11 @@ def test_a_run_reports_what_it_could_not_describe(rendered):
 def test_without_a_manifest_the_pages_still_render(tmp_path):
     pages, _ = render(str(tmp_path / "bare"), manifest_dir=str(tmp_path / "none"))
     assert pages, "a missing manifest must not stop the deploy"
-    # No shapes to state, so every workload is named by its benchmark id alone.
+    # No shapes to state, so every workload is named by its benchmark id alone,
+    # the trailing dtype split off so the label's dtypes still share one group.
     assert "wl-tensor" not in "".join(pages.values())
-    assert "decode-b1-h8-bfloat16" in "".join(pages.values())
+    assert ('<td class="wl-name" rowspan="2"><code>decode-<wbr>b1-<wbr>h8</code></td>'
+            '<td class="colsep">bf16</td>') in "".join(pages.values())
 
 
 def test_a_manifest_in_a_subdirectory_renders_the_same_pages(tmp_path, rendered):