diff --git a/benchmark/amd/catalog.py b/benchmark/amd/catalog.py
index 9ce6d1a3ae..8c52ce6427 100644
--- a/benchmark/amd/catalog.py
+++ b/benchmark/amd/catalog.py
@@ -89,6 +89,8 @@ def _qwen3tts(name: str, size: str, runner_model: str, sglang_omni_config: str,
     _qwen("4B"),
     _qwen("9B"),
     _qwen("9B", gpus=2, tp=2, name="qwen3_5_9b_tp2"),
+    # speculative decoding with the checkpoint's MTP head drafting k tokens (configs/qwen3_5_<size>_mtp<k>.yaml)
+    *(_qwen(size, name=f"qwen3_5_{size.lower()}_mtp{k}") for size in ("0.8B", "9B") for k in (2, 3)),
     Row("kokoro", "tts", "hexgrad/Kokoro-82M", "kokoro"),
     Row("orpheus", "tts", "canopylabs/orpheus-3b-0.1-ft", "orpheus", gpus=2),
     Row("orpheus_colocated", "tts", "canopylabs/orpheus-3b-0.1-ft", "orpheus"),
diff --git a/benchmark/spec_decode/__init__.py b/benchmark/spec_decode/__init__.py
new file mode 100644
index 0000000000..35995edeca
--- /dev/null
+++ b/benchmark/spec_decode/__init__.py
@@ -0,0 +1,2 @@
+"""Speculative decoding measurements that run without a server: the acceptance rate of a draft model against its
+target over the benchmark prompt sets, and the cost of a draft token."""
diff --git a/benchmark/spec_decode/draft_cost.py b/benchmark/spec_decode/draft_cost.py
new file mode 100644
index 0000000000..730d77dfd3
--- /dev/null
+++ b/benchmark/spec_decode/draft_cost.py
@@ -0,0 +1,183 @@
+"""What one draft token costs on the GPU: the Qwen3.5 MTP head at batch 1 plus the shared lm_head.
+
+    python -m benchmark.spec_decode.draft_cost --model Qwen/Qwen3.5-9B --context 1024
+
+Times, with CUDA events over repeated runs: the lm_head GEMM at 1..6 rows (a verify step's rows at concurrency 1),
+each GEMM of the head at one row, and one complete draft step (fc, the full-attention decoder layer over a
+fixed-length KV cache, final norm, lm_head, argmax) both eager and replayed from a HIP graph. The graph number is the
+launch-free cost a captured engine step would pay; the eager number is what the acceptance harness pays.
+"""
+from __future__ import annotations
+
+import argparse
+import sys
+
+import torch
+import torch.nn.functional as F
+from transformers.models.qwen3_5.modeling_qwen3_5 import apply_rotary_pos_emb, repeat_kv
+
+from benchmark.spec_decode.qwen3_5_mtp import BaseModel, MtpHead, load_base
+
+
+def time_fn(fn, iters: int, repeats: int = 3) -> float:
+    """Mean milliseconds per call, best of `repeats` blocks of `iters` calls."""
+    for _ in range(3):
+        fn()
+    torch.cuda.synchronize()
+    best = float("inf")
+    for _ in range(repeats):
+        start, end = torch.cuda.Event(enable_timing=True), torch.cuda.Event(enable_timing=True)
+        start.record()
+        for _ in range(iters):
+            fn()
+        end.record()
+        torch.cuda.synchronize()
+        best = min(best, start.elapsed_time(end) / iters)
+    return best
+
+
+def gemm_row(name: str, weight: torch.Tensor, rows: int, iters: int) -> dict:
+    x = torch.randn(rows, weight.shape[1], device=weight.device, dtype=weight.dtype)
+    ms = time_fn(lambda: F.linear(x, weight), iters)
+    byte_count = weight.numel() * weight.element_size()
+    return {"name": name, "rows": rows, "shape": list(weight.shape), "ms": ms, "mb": byte_count / 1e6,
+            "tb_s": byte_count / (ms * 1e-3) / 1e12}
+
+
+class DraftStep:
+    """One draft token through the head with static buffers, so it can be captured into a HIP graph."""
+
+    def __init__(self, head: MtpHead, base: BaseModel, context: int):
+        self.head, self.base = head, base
+        layer = head.layers[0]
+        attn = layer.self_attn
+        self.attn, self.layer = attn, layer
+        cfg = base.config
+        device, dtype = base.device, base.dtype
+        self.kv_heads = cfg.num_key_value_heads
+        self.groups = attn.num_key_value_groups
+        self.head_dim = attn.head_dim
+        hidden = cfg.hidden_size
+        self.h = torch.randn(1, hidden, device=device, dtype=dtype)
+        self.tok = torch.tensor([1], device=device)
+        # Keys and values of `context` earlier positions; the new token lands in the last slot.
+        self.K = torch.randn(1, self.kv_heads, context, self.head_dim, device=device, dtype=dtype)
+        self.V = torch.randn(1, self.kv_heads, context, self.head_dim, device=device, dtype=dtype)
+        pos = torch.tensor([[[context - 1]]], device=device).expand(3, 1, 1)
+        self.cos, self.sin = base.rotary(self.h[None], pos)
+        self.gqa = True
+        self.out = None
+
+    def attention(self, q, k, v):
+        if self.gqa:
+            try:
+                return F.scaled_dot_product_attention(q, k, v, scale=self.attn.scaling, enable_gqa=True)
+            except (RuntimeError, TypeError):
+                self.gqa = False
+        return F.scaled_dot_product_attention(
+            q, repeat_kv(k, self.groups), repeat_kv(v, self.groups), scale=self.attn.scaling
+        )
+
+    def layer_forward(self) -> torch.Tensor:
+        head, layer, attn = self.head, self.layer, self.attn
+        emb = head.pre_fc_norm_embedding(self.base.embed(self.tok))
+        x = head.fc(torch.cat((emb, head.pre_fc_norm_hidden(self.h)), dim=-1))[None]  # (1, 1, hidden)
+        residual = x
+        y = layer.input_layernorm(x)
+        q, gate = torch.chunk(attn.q_proj(y).view(1, 1, -1, self.head_dim * 2), 2, dim=-1)
+        gate = gate.reshape(1, 1, -1)
+        q = attn.q_norm(q.reshape(1, 1, -1, self.head_dim)).transpose(1, 2)
+        k = attn.k_norm(attn.k_proj(y).view(1, 1, -1, self.head_dim)).transpose(1, 2)
+        v = attn.v_proj(y).view(1, 1, -1, self.head_dim).transpose(1, 2)
+        q, k = apply_rotary_pos_emb(q, k, self.cos, self.sin)
+        self.K[:, :, -1:] = k
+        self.V[:, :, -1:] = v
+        o = self.attention(q, self.K, self.V).transpose(1, 2).reshape(1, 1, -1)
+        x = residual + attn.o_proj(o * torch.sigmoid(gate))
+        residual = x
+        x = residual + layer.mlp(layer.post_attention_layernorm(x))
+        return head.norm(x)[0, -1]
+
+    def full(self) -> None:
+        self.out = self.base.lm_head(self.layer_forward()).argmax()
+
+    def layer_only(self) -> None:
+        self.out = self.layer_forward()
+
+
+def capture(fn) -> torch.cuda.CUDAGraph:
+    stream = torch.cuda.Stream()
+    stream.wait_stream(torch.cuda.current_stream())
+    with torch.cuda.stream(stream):
+        for _ in range(3):
+            fn()
+    torch.cuda.current_stream().wait_stream(stream)
+    graph = torch.cuda.CUDAGraph()
+    with torch.cuda.graph(graph):
+        fn()
+    return graph
+
+
+def main(argv=None) -> int:
+    ap = argparse.ArgumentParser(description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter)
+    ap.add_argument("--model", default="Qwen/Qwen3.5-9B")
+    ap.add_argument("--device", default="cuda")
+    ap.add_argument("--context", type=int, nargs="+", default=[256, 1024])
+    ap.add_argument("--iters", type=int, default=50)
+    args = ap.parse_args(argv)
+
+    base, snapshot = load_base(args.model, args.device)
+    head = MtpHead.from_snapshot(snapshot, base.config, base.device, base.dtype)
+    print(f"{args.model} on {torch.cuda.get_device_name(0)}, torch {torch.__version__}, {base.dtype}")
+    lm_head = base.lm_head.weight
+    mb = lambda w: w.numel() * w.element_size() / 1e6  # noqa: E731
+    layer_bytes = sum(mb(p) for p in head.parameters())
+    print(f"mtp head weights {layer_bytes:.1f} MB, lm_head {mb(lm_head):.1f} MB")
+
+    with torch.inference_mode():
+        print("\nlm_head GEMM (bf16, torch F.linear):")
+        for rows in range(1, 7):
+            r = gemm_row("lm_head", lm_head, rows, args.iters)
+            print(f"  rows={rows}: {r['ms']:.3f} ms  {r['tb_s']:.2f} TB/s")
+
+        print("\nmtp head GEMMs at one row:")
+        layer = head.layers[0]
+        attn, mlp = layer.self_attn, layer.mlp
+        weights = [
+            ("fc", head.fc.weight), ("q_proj", attn.q_proj.weight), ("k_proj", attn.k_proj.weight),
+            ("v_proj", attn.v_proj.weight), ("o_proj", attn.o_proj.weight), ("gate_proj", mlp.gate_proj.weight),
+            ("up_proj", mlp.up_proj.weight), ("down_proj", mlp.down_proj.weight),
+        ]
+        total_ms = 0.0
+        for name, w in weights:
+            r = gemm_row(name, w, 1, args.iters)
+            total_ms += r["ms"]
+            print(f"  {name:10s} {str(r['shape']):16s} {r['ms']:.3f} ms  {r['mb']:7.1f} MB  {r['tb_s']:.2f} TB/s")
+        lm = gemm_row("lm_head", lm_head, 1, args.iters)
+        print(f"  GEMM sum: head {total_ms:.3f} ms + lm_head {lm['ms']:.3f} ms = {total_ms + lm['ms']:.3f} ms")
+
+        for context in args.context:
+            step = DraftStep(head, base, context)
+            eager_full = time_fn(step.full, args.iters)
+            eager_layer = time_fn(step.layer_only, args.iters)
+            print(
+                f"\ndraft step, context {context} (gqa sdpa={step.gqa}):\n"
+                f"  eager: full {eager_full:.3f} ms, layer only {eager_layer:.3f} ms"
+            )
+            try:
+                g_full = capture(step.full)
+                graph_full = time_fn(g_full.replay, args.iters)
+                g_layer = capture(step.layer_only)
+                graph_layer = time_fn(g_layer.replay, args.iters)
+            except RuntimeError as exc:
+                print(f"  graph capture failed: {exc!r}")
+                continue
+            print(
+                f"  graph: full {graph_full:.3f} ms, layer only {graph_layer:.3f} ms, "
+                f"lm_head share {(graph_full - graph_layer) / graph_full * 100:.0f}%"
+            )
+    return 0
+
+
+if __name__ == "__main__":
+    sys.exit(main())
diff --git a/benchmark/spec_decode/draft_head_variants.py b/benchmark/spec_decode/draft_head_variants.py
new file mode 100644
index 0000000000..a447c98f40
--- /dev/null
+++ b/benchmark/spec_decode/draft_head_variants.py
@@ -0,0 +1,229 @@
+"""What a cheaper lm_head for the draft costs in acceptance, measured without a server.
+
+    python -m benchmark.spec_decode.draft_head_variants --model Qwen/Qwen3.5-9B --k 3 --out variants.json
+
+Runs greedy speculative decoding once with the full bf16 head and, at every draft position, also reads the head's
+output through each variant (a quantized copy, or a subset of the vocabulary). A variant's accepted run in a step is
+the longest prefix of its tokens equal to what the target emitted. The chain itself follows the bf16 drafts, so a
+variant that is right where bf16 was wrong stops counting there: the figure is a slight lower bound on what the variant
+reaches on its own chain. ``bf16`` is the control: its figure is the run's own tokens per step, give or take the last
+step of each prompt.
+"""
+from __future__ import annotations
+
+import argparse
+import json
+import sys
+import time
+from dataclasses import dataclass, field
+from pathlib import Path
+
+import torch
+import torch.nn.functional as F
+
+from benchmark.spec_decode.mtp_acceptance import (
+    CHAT_PROMPTS,
+    LONG_PROMPTS,
+    chat_ids,
+    read_prompts,
+    stop_ids_of,
+)
+from benchmark.spec_decode.qwen3_5_mtp import HFDraft, HFTarget, MtpHead, load_base, load_tokenizer
+from benchmark.spec_decode.spec_loop import run_speculative
+from benchmark.spec_decode.verify import SamplingParams
+
+
+def fp8_round(x: torch.Tensor, scale: torch.Tensor | float) -> torch.Tensor:
+    """``x`` rounded through float8_e4m3fnuz at ``scale`` (the format holds up to 240), back in the original units."""
+    q = (x.float() * scale).to(torch.float8_e4m3fnuz)
+    return q.float() / scale
+
+
+def int_round(x: torch.Tensor, scale: torch.Tensor | float, bits: int) -> torch.Tensor:
+    """``x`` rounded to a signed ``bits``-bit integer at ``scale``, clipped, back in the original units."""
+    top = 2 ** (bits - 1) - 1
+    return ((x.float() * scale).round().clamp(-top, top)) / scale
+
+
+@dataclass
+class Variant:
+    """One way of reading the draft head: ``weight`` [ids, hidden] with ``ids`` naming its rows (None: every row), and
+    ``x_round`` applied to the head's input first."""
+    name: str
+    weight: torch.Tensor
+    ids: torch.Tensor | None = None
+    x_round: object = None
+    agree: int = 0
+    total: int = 0
+    accepted: int = 0
+    per_position: list[int] = field(default_factory=list)
+
+    def token(self, hidden: torch.Tensor) -> int:
+        x = hidden.float()
+        if self.x_round is not None:
+            x = self.x_round(x)
+        logits = F.linear(x, self.weight.float())
+        best = int(logits.argmax())
+        return best if self.ids is None else int(self.ids[best])
+
+
+class RecordingDraft(HFDraft):
+    """The MTP head, keeping per step the bf16 draft tokens and each variant's: the steps' lists line up with the
+    positions they predict, which the generated sequence settles afterwards."""
+
+    def __init__(self, head: MtpHead, base, variants: list[Variant]):
+        super().__init__(head, base)
+        self.variants = variants
+        # one entry per draft position: (predicted position, {variant name: token})
+        self.steps: list[list[tuple[int, dict[str, int]]]] = []
+
+    def reset(self) -> None:
+        super().reset()
+        self.steps = []
+
+    def extend(self, hiddens: torch.Tensor, tokens: torch.Tensor) -> torch.Tensor:
+        logits = super().extend(hiddens, tokens)
+        self.steps.append([self._record()])
+        return logits
+
+    def chain(self, token: int) -> torch.Tensor:
+        logits = super().chain(token)
+        self.steps[-1].append(self._record())
+        return logits
+
+    def _record(self) -> tuple[int, dict[str, int]]:
+        # position p of the draft predicts the token at p + 2; its cache now holds positions up to length - 1
+        return self.length + 1, {v.name: v.token(self.last_hidden[None]) for v in self.variants}
+
+
+def score(variants: list[Variant], steps, prompt: list[int], generated: list[int], k: int) -> None:
+    seq = prompt + generated
+    for step in steps:
+        if len(step) < k:
+            continue  # a step the generation ended before verifying
+        if step[0][0] >= len(seq):
+            continue
+        for v in variants:
+            run = 0
+            alive = True
+            for j, (position, tokens) in enumerate(step):
+                if position >= len(seq):
+                    break
+                if alive and tokens[v.name] == seq[position]:
+                    run += 1
+                    v.per_position[j] += 1
+                else:
+                    alive = False
+                v.agree += tokens[v.name] == tokens["bf16"]
+                v.total += 1
+            v.accepted += run
+
+
+def build_variants(names: list[str], weight: torch.Tensor, stop_ids: frozenset[int], k: int) -> list[Variant]:
+    device = weight.device
+    w = weight.float()
+    variants = [Variant("bf16", w)]
+    row_max = w.abs().amax(dim=1, keepdim=True)
+    row_norm = w.norm(dim=1)
+    specials = torch.tensor(sorted(stop_ids), device=device)
+    for name in names:
+        if name == "fp8":  # one scale for the head: argmax needs no per-channel scale back
+            variants.append(Variant(name, fp8_round(w, 240.0 / float(row_max.max()))))
+        elif name == "fp8-rowscale":
+            variants.append(Variant(name, fp8_round(w, 240.0 / row_max)))
+        elif name == "fp8-x":  # the input too, as an fp8 x fp8 kernel reads it
+            variants.append(Variant(name, fp8_round(w, 240.0 / float(row_max.max())),
+                                    x_round=lambda x: fp8_round(x, 240.0 / x.abs().amax(dim=-1, keepdim=True))))
+        elif name.startswith("int"):
+            # three of the head's standard deviations reach full scale: the far tail clips, the bulk keeps its bits
+            bits, _, rowscale = name[3:].partition("-")
+            top = 2 ** (int(bits) - 1) - 1
+            scale = top / (3 * w.std(dim=1, keepdim=True)) if rowscale else top / (3 * float(w.std()))
+            variants.append(Variant(name, int_round(w, scale, int(bits))))
+        elif name.startswith("prefix-"):
+            n = int(name.split("-")[1])
+            ids = torch.cat((torch.arange(n, device=device), specials[specials >= n]))
+            variants.append(Variant(name, w[ids], ids=ids))
+        elif name.startswith("norm-"):
+            n = int(name.split("-")[1])
+            ids = torch.unique(torch.cat((row_norm.topk(n).indices, specials)))
+            variants.append(Variant(name, w[ids], ids=ids))
+        else:
+            raise SystemExit(f"unknown variant {name}")
+    for v in variants:
+        v.per_position = [0] * k
+    return variants
+
+
+def main(argv=None) -> int:
+    ap = argparse.ArgumentParser(description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter)
+    ap.add_argument("--model", default="Qwen/Qwen3.5-9B")
+    ap.add_argument("--device", default="cuda")
+    ap.add_argument("--k", type=int, default=3)
+    ap.add_argument("--n-chat", type=int, default=32)
+    ap.add_argument("--n-long", type=int, default=16)
+    ap.add_argument("--max-new-chat", type=int, default=256)
+    ap.add_argument("--max-new-long", type=int, default=128)
+    ap.add_argument("--extra-prompts", help="a file of further prompts, one a line, scored as their own set")
+    ap.add_argument("--variants", nargs="+", default=[
+        "fp8", "fp8-rowscale", "fp8-x", "int8", "int8-rowscale", "int6-rowscale", "int4-rowscale",
+        "prefix-64000", "prefix-96000", "prefix-128000", "prefix-151000",
+        "norm-64000", "norm-124160",
+    ])
+    ap.add_argument("--seed", type=int, default=0)
+    ap.add_argument("--out", required=True)
+    args = ap.parse_args(argv)
+
+    base, snapshot = load_base(args.model, args.device)
+    tokenizer = load_tokenizer(snapshot)
+    head = MtpHead.from_snapshot(snapshot, base.config, base.device, base.dtype)
+    stop_ids = stop_ids_of(tokenizer, base.config)
+    variants = build_variants(args.variants, base.lm_head.weight.detach(), stop_ids, args.k)
+    target = HFTarget(base, "post")
+    draft = RecordingDraft(head, base, variants)
+    sets = {
+        "chat": ([chat_ids(tokenizer, p) for p in read_prompts(CHAT_PROMPTS, args.n_chat)], args.max_new_chat),
+        "long": ([chat_ids(tokenizer, p) for p in read_prompts(LONG_PROMPTS, args.n_long)], args.max_new_long),
+    }
+    if args.extra_prompts:
+        extra = read_prompts(Path(args.extra_prompts), 1 << 30)
+        sets["extra"] = ([chat_ids(tokenizer, p) for p in extra], args.max_new_chat)
+    out: dict = {"model": args.model, "k": args.k, "sets": {}}
+    for name, (prompts, max_new) in sets.items():
+        for v in variants:
+            v.agree = v.total = v.accepted = 0
+            v.per_position = [0] * args.k
+        run_steps = run_tokens = accepted = 0
+        t0 = time.perf_counter()
+        for i, prompt in enumerate(prompts):
+            generator = torch.Generator(device=base.device)
+            generator.manual_seed(args.seed * 1000 + i)
+            generated, stats = run_speculative(
+                target, draft, prompt, args.k, SamplingParams(), max_new, stop_ids, generator,
+            )
+            run_steps += stats.steps
+            run_tokens += stats.tokens
+            accepted += stats.accepted_total
+            score(variants, draft.steps, prompt, generated, args.k)
+        wall = time.perf_counter() - t0
+        rows = {}
+        control = max(next(v for v in variants if v.name == "bf16").accepted, 1)
+        for v in variants:
+            rows[v.name] = {
+                "tokens_per_step": 1.0 + v.accepted / max(run_steps, 1),
+                "accepted_vs_bf16": v.accepted / control,
+                "agree_with_bf16": v.agree / max(v.total, 1),
+                "accepted_per_position": [c / max(run_steps, 1) for c in v.per_position],
+            }
+            print(f"{name} {v.name:16s} tok/step {rows[v.name]['tokens_per_step']:.3f} accepted/bf16 "
+                  f"{rows[v.name]['accepted_vs_bf16']:.4f} agree {rows[v.name]['agree_with_bf16']:.4f}", flush=True)
+        out["sets"][name] = {"steps": run_steps, "tokens": run_tokens, "wall_s": wall, "variants": rows,
+                             "run_tokens_per_step": 1 + accepted / max(run_steps, 1)}
+        print(f"{name}: {len(prompts)} prompts, {run_steps} steps, {run_tokens} tokens, run tok/step "
+              f"{1 + accepted / max(run_steps, 1):.3f}, {wall:.0f} s", flush=True)
+        Path(args.out).write_text(json.dumps(out, indent=1))
+    return 0
+
+
+if __name__ == "__main__":
+    sys.exit(main())
diff --git a/benchmark/spec_decode/mtp_acceptance.py b/benchmark/spec_decode/mtp_acceptance.py
new file mode 100644
index 0000000000..cda6e75cc4
--- /dev/null
+++ b/benchmark/spec_decode/mtp_acceptance.py
@@ -0,0 +1,282 @@
+"""Acceptance rate of the Qwen3.5 MTP head against its own base model, measured without a server.
+
+    python -m benchmark.spec_decode.mtp_acceptance --model Qwen/Qwen3.5-9B --calibrate --out acceptance.json
+
+Loads the base model and the ``mtp.*`` head from the cached snapshot (offline), builds the chat and long-prompt sets
+the sweep uses (benchmark/assets, chat template with thinking off), then runs real speculative decoding for every
+(k, sampling mode) cell: k drafts chained from the head, verified token by token against the base model (spec_loop).
+Reports per-position acceptance, mean accepted drafts and tokens per step, and writes everything to a JSON file
+after each cell so a cut-off run still leaves its finished cells behind.
+
+``--calibrate`` first compares the four input conventions the head could have been trained with (hidden state before
+or after the base's final norm, embedding or hidden state first in the fc input) on a few prompts and keeps the one
+with the highest greedy acceptance; a wrong convention shows up as near-zero acceptance, not as an error.
+"""
+from __future__ import annotations
+
+import argparse
+import json
+import sys
+import time
+from pathlib import Path
+
+import torch
+
+from benchmark.spec_decode.qwen3_5_mtp import HFDraft, HFTarget, MtpHead, load_base, load_tokenizer
+from benchmark.spec_decode.spec_loop import RunStats, run_speculative
+from benchmark.spec_decode.verify import SamplingParams
+
+REPO = Path(__file__).resolve().parents[2]
+CHAT_PROMPTS = REPO / "benchmark/assets/simple_text_queries.txt"
+LONG_PROMPTS = REPO / "benchmark/assets/long_text_queries_1k.txt"
+
+
+def read_prompts(path: Path, n: int) -> list[str]:
+    lines = [line.strip() for line in path.read_text(encoding="utf-8").splitlines() if line.strip()]
+    return lines[:n]
+
+
+def chat_ids(tokenizer, prompt: str) -> list[int]:
+    text = tokenizer.apply_chat_template(
+        [{"role": "user", "content": prompt}], tokenize=False, add_generation_prompt=True, enable_thinking=False
+    )
+    return tokenizer(text, add_special_tokens=False)["input_ids"]
+
+
+def stop_ids_of(tokenizer, config) -> frozenset[int]:
+    """The chat turn's end and the document end, as the serving stack stops on both."""
+    ids = set()
+    for value in (tokenizer.eos_token_id, getattr(config, "eos_token_id", None)):
+        if isinstance(value, int):
+            ids.add(value)
+        elif isinstance(value, (list, tuple)):
+            ids.update(int(v) for v in value)
+    for token in ("<|im_end|>", "<|endoftext|>"):
+        token_id = tokenizer.convert_tokens_to_ids(token)
+        if isinstance(token_id, int) and token_id >= 0 and token_id != tokenizer.unk_token_id:
+            ids.add(token_id)
+    return frozenset(ids)
+
+
+def fmt_rates(rates: list[float]) -> str:
+    return "[" + " ".join(f"{r:.3f}" for r in rates) + "]"
+
+
+def run_set(
+    target: HFTarget,
+    draft: HFDraft,
+    prompts: list[list[int]],
+    k: int,
+    params: SamplingParams,
+    max_new: int,
+    stop_ids: frozenset[int],
+    seed: int,
+    deadline: float | None,
+) -> tuple[RunStats, list[list[int]], float, bool]:
+    """Run one prompt set. Returns the merged stats, the generations, the wall time and whether the deadline cut it."""
+    stats = RunStats(k=k)
+    outputs: list[list[int]] = []
+    t0 = time.perf_counter()
+    cut = False
+    for i, prompt in enumerate(prompts):
+        if deadline is not None and time.time() > deadline:
+            cut = True
+            break
+        generator = torch.Generator(device=target.base.device)
+        generator.manual_seed(seed * 1000 + i)
+        tokens, run_stats = run_speculative(target, draft, prompt, k, params, max_new, stop_ids, generator)
+        stats.merge(run_stats)
+        outputs.append(tokens)
+    return stats, outputs, time.perf_counter() - t0, cut
+
+
+def main(argv=None) -> int:
+    ap = argparse.ArgumentParser(description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter)
+    ap.add_argument("--model", default="Qwen/Qwen3.5-0.8B")
+    ap.add_argument("--device", default="cuda")
+    ap.add_argument("--chat-prompts", default=str(CHAT_PROMPTS))
+    ap.add_argument("--long-prompts", default=str(LONG_PROMPTS))
+    ap.add_argument("--n-chat", type=int, default=32)
+    ap.add_argument("--n-long", type=int, default=16)
+    ap.add_argument("--max-new-chat", type=int, default=256)
+    ap.add_argument("--max-new-long", type=int, default=128)
+    ap.add_argument("--k", type=int, nargs="+", default=[1, 2, 3, 4])
+    ap.add_argument("--modes", nargs="+", default=["greedy", "sampling"], choices=["greedy", "sampling"])
+    ap.add_argument("--temperature", type=float, default=0.7)
+    ap.add_argument("--top-p", type=float, default=0.95)
+    ap.add_argument("--hidden", default="post", choices=["post", "pre"], help="which base hidden state the head sees")
+    ap.add_argument("--fc-order", default="emb_first", choices=["emb_first", "hidden_first"])
+    ap.add_argument("--calibrate", action="store_true", help="pick --hidden/--fc-order by greedy acceptance first")
+    ap.add_argument("--calibrate-prompts", type=int, default=6)
+    ap.add_argument("--calibrate-tokens", type=int, default=48)
+    ap.add_argument("--seed", type=int, default=0)
+    ap.add_argument("--deadline", type=float, default=None, help="epoch seconds; stop starting prompts after it")
+    ap.add_argument("--out", required=True)
+    args = ap.parse_args(argv)
+
+    t0 = time.perf_counter()
+    base, snapshot = load_base(args.model, args.device)
+    tokenizer = load_tokenizer(snapshot)
+    head = MtpHead.from_snapshot(snapshot, base.config, base.device, base.dtype)
+    load_s = time.perf_counter() - t0
+    lm_head_bytes = base.lm_head.weight.numel() * base.lm_head.weight.element_size()
+    print(
+        f"{args.model}: snapshot {snapshot.name}, {type(base.model).__name__}, {base.dtype}, "
+        f"{base.config.num_hidden_layers} layers, hidden {base.config.hidden_size}, vocab {base.config.vocab_size}; "
+        f"mtp head {head.weight_bytes() / 1e6:.1f} MB (+ lm_head {lm_head_bytes / 1e6:.1f} MB shared); "
+        f"loaded in {load_s:.0f} s",
+        flush=True,
+    )
+    stop_ids = stop_ids_of(tokenizer, base.config)
+
+    chat_prompts = read_prompts(Path(args.chat_prompts), args.n_chat)
+    long_prompts = read_prompts(Path(args.long_prompts), args.n_long)
+    sets = {
+        "chat": ([chat_ids(tokenizer, p) for p in chat_prompts], args.max_new_chat),
+        "long": ([chat_ids(tokenizer, p) for p in long_prompts], args.max_new_long),
+    }
+    for name, (prompts, max_new) in sets.items():
+        lengths = [len(p) for p in prompts]
+        print(f"{name}: {len(prompts)} prompts, {min(lengths)}-{max(lengths)} tokens, up to {max_new} new", flush=True)
+
+    result = {
+        "model": args.model,
+        "snapshot": str(snapshot),
+        "torch": torch.__version__,
+        "device": torch.cuda.get_device_name(0) if torch.cuda.is_available() else "cpu",
+        "config": {
+            "num_hidden_layers": base.config.num_hidden_layers,
+            "hidden_size": base.config.hidden_size,
+            "vocab_size": base.config.vocab_size,
+            "mtp_num_hidden_layers": getattr(base.config, "mtp_num_hidden_layers", None),
+        },
+        "mtp_weight_bytes": head.weight_bytes(),
+        "lm_head_bytes": lm_head_bytes,
+        "prompts": {
+            name: {"file": args.chat_prompts if name == "chat" else args.long_prompts, "n": len(prompts),
+                   "max_new": max_new, "token_counts": [len(p) for p in prompts]}
+            for name, (prompts, max_new) in sets.items()
+        },
+        "sampling": {"temperature": args.temperature, "top_p": args.top_p},
+        "calibration": [],
+        "cells": [],
+        "samples": [],
+    }
+    out = Path(args.out)
+
+    def save() -> None:
+        out.write_text(json.dumps(result, indent=1), encoding="utf-8")
+
+    import transformers
+
+    result["transformers"] = transformers.__version__
+
+    hidden, fc_order = args.hidden, args.fc_order
+    if args.calibrate:
+        cal_prompts = sets["chat"][0][: args.calibrate_prompts]
+        best = None
+        for cand_hidden in ("post", "pre"):
+            for cand_order in ("emb_first", "hidden_first"):
+                target = HFTarget(base, cand_hidden)
+                draft = HFDraft(head, base, cand_order)
+                stats, _, wall, _ = run_set(
+                    target, draft, cal_prompts, 4, SamplingParams(), args.calibrate_tokens, stop_ids, args.seed, None
+                )
+                target.close()
+                row = {"hidden": cand_hidden, "fc_order": cand_order, "mean_accepted": stats.mean_accepted,
+                       "position_rates": stats.position_rates(), "steps": stats.steps}
+                result["calibration"].append(row)
+                print(
+                    f"calibrate hidden={cand_hidden} fc_order={cand_order}: accepted/step {stats.mean_accepted:.3f} "
+                    f"p={fmt_rates(stats.position_rates())} ({stats.steps} steps, {wall:.0f} s)",
+                    flush=True,
+                )
+                if best is None or stats.mean_accepted > best[0]:
+                    best = (stats.mean_accepted, cand_hidden, cand_order)
+        _, hidden, fc_order = best
+        print(f"calibrate: using hidden={hidden} fc_order={fc_order}", flush=True)
+    result["chosen"] = {"hidden": hidden, "fc_order": fc_order}
+    save()
+
+    target = HFTarget(base, hidden)
+    draft = HFDraft(head, base, fc_order)
+    greedy_outputs: dict[tuple[str, int], list[list[int]]] = {}
+    cut = False
+    for mode in args.modes:
+        params = SamplingParams() if mode == "greedy" else SamplingParams(args.temperature, args.top_p)
+        for k in args.k:
+            if cut:
+                break
+            cell = {"mode": mode, "temperature": params.temperature, "top_p": params.top_p, "k": k, "sets": {}}
+            total = RunStats(k=k)
+            wall_total = 0.0
+            for name, (prompts, max_new) in sets.items():
+                stats, outputs, wall, cut = run_set(
+                    target, draft, prompts, k, params, max_new, stop_ids, args.seed, args.deadline
+                )
+                wall_total += wall
+                summary = stats.summary()
+                summary["wall_s"] = wall
+                summary["hf_ms_per_token"] = 1e3 * wall / max(stats.tokens, 1)
+                summary["complete"] = not cut
+                cell["sets"][name] = summary
+                total.merge(stats)
+                if mode == "greedy":
+                    greedy_outputs[(name, k)] = outputs
+                if outputs and k == args.k[0]:
+                    for i in range(min(2, len(outputs))):
+                        result["samples"].append({
+                            "set": name, "index": i, "mode": mode, "k": k,
+                            "text": tokenizer.decode(outputs[i], skip_special_tokens=False)[:600],
+                        })
+                print(
+                    f"{args.model} {mode} k={k} {name}: {stats.prompts} prompts, {stats.steps} steps, "
+                    f"{stats.tokens} tokens, accepted/step {stats.mean_accepted:.3f}, tokens/step "
+                    f"{stats.tokens_per_step:.3f}, p={fmt_rates(stats.position_rates())} "
+                    f"E[p]={fmt_rates(stats.position_expected_rates())} reach={fmt_rates(stats.reach_rates())} "
+                    f"({wall:.0f} s, {summary['hf_ms_per_token']:.1f} ms/token){' CUT' if cut else ''}",
+                    flush=True,
+                )
+                if cut:
+                    break
+            cell["sets"]["all"] = total.summary()
+            cell["wall_s"] = wall_total
+            result["cells"].append(cell)
+            save()
+        if cut:
+            print("deadline reached; remaining cells skipped", flush=True)
+            break
+
+    # Greedy speculative decoding must reproduce greedy decoding whatever k is: the target is driven one token at
+    # a time either way, so the continuations should agree across k token for token.
+    consistency = {"prompts_compared": 0, "prompts_differing": 0}
+    for name in sets:
+        runs = [greedy_outputs[(name, k)] for k in args.k if (name, k) in greedy_outputs]
+        if len(runs) < 2:
+            continue
+        for i in range(min(len(r) for r in runs)):
+            consistency["prompts_compared"] += 1
+            if any(r[i] != runs[0][i] for r in runs[1:]):
+                consistency["prompts_differing"] += 1
+    result["greedy_consistency"] = consistency
+    save()
+
+    print(f"\n{args.model} (hidden={hidden}, fc_order={fc_order}); greedy outputs differing across k: "
+          f"{consistency['prompts_differing']}/{consistency['prompts_compared']}")
+    print(f"{'mode':9s} {'k':>2s} {'set':5s} {'steps':>6s} {'tok/step':>8s} {'acc/step':>8s} {'p1..pk':30s} E[p1..pk]")
+    for cell in result["cells"]:
+        for name in ("chat", "long", "all"):
+            s = cell["sets"].get(name)
+            if s is None:
+                continue
+            print(
+                f"{cell['mode']:9s} {cell['k']:2d} {name:5s} {s['steps']:6d} {s['tokens_per_step']:8.3f} "
+                f"{s['mean_accepted']:8.3f} {fmt_rates(s['position_rates']):30s} "
+                f"{fmt_rates(s['position_expected_rates']):30s}"
+            )
+    print(f"written {out}")
+    return 0
+
+
+if __name__ == "__main__":
+    sys.exit(main())
diff --git a/benchmark/spec_decode/qwen3_5_mtp.py b/benchmark/spec_decode/qwen3_5_mtp.py
new file mode 100644
index 0000000000..e90bf1e089
--- /dev/null
+++ b/benchmark/spec_decode/qwen3_5_mtp.py
@@ -0,0 +1,217 @@
+"""Qwen3.5 on Hugging Face transformers as the target, and the checkpoint's MTP head as the draft.
+
+transformers' Qwen3.5 model ignores the ``mtp.*`` tensors, so the head is assembled here from the same building blocks
+(one full-attention decoder layer, the fc projection and its two input norms, a final norm) and loaded straight from
+the safetensors shards of the cached snapshot. It shares the embedding and the lm_head with the base model: the
+checkpoints set ``mtp_use_dedicated_embeddings: false`` and ship no ``mtp.lm_head``.
+
+Draft input at position p is fc(cat(norm_e(embed(x_{p+1})), norm_h(h_p))), with h_p the target's hidden state after its
+final norm (the tensor its lm_head reads) at position p; a chained draft feeds the head's own output hidden state
+back as h. Both conventions are the vLLM / SGLang ones for this head; the CLI's calibration compares them against the
+pre-norm hidden state and the other concatenation order on a few prompts before trusting them.
+"""
+from __future__ import annotations
+
+import json
+import os
+from pathlib import Path
+
+import torch
+from safetensors import safe_open
+from torch import nn
+from transformers import AutoModelForCausalLM, AutoTokenizer, DynamicCache
+from transformers.models.qwen3_5.modeling_qwen3_5 import Qwen3_5DecoderLayer, Qwen3_5RMSNorm
+
+MTP_PREFIX = "mtp."
+
+
+def snapshot_dir(repo_id: str) -> Path:
+    """The cached snapshot of a Hub repo, found without touching the network."""
+    hub = Path(os.environ.get("HF_HOME", Path.home() / ".cache" / "huggingface")) / "hub"
+    repo = hub / ("models--" + repo_id.replace("/", "--")) / "snapshots"
+    candidates = [p for p in repo.iterdir() if (p / "config.json").exists()] if repo.exists() else []
+    if not candidates:
+        raise FileNotFoundError(f"no cached snapshot with config.json for {repo_id} under {hub}")
+    return sorted(candidates, key=lambda p: p.stat().st_mtime)[-1]
+
+
+class BaseModel:
+    """The text stack of a loaded checkpoint: the decoder, its lm_head, embedding and rotary module."""
+
+    def __init__(self, model: nn.Module):
+        self.model = model
+        if hasattr(model.model, "language_model"):  # Qwen3_5ForConditionalGeneration
+            self.text = model.model.language_model
+        else:  # Qwen3_5ForCausalLM
+            self.text = model.model
+        self.lm_head = model.lm_head
+        self.embed = self.text.embed_tokens
+        self.rotary = self.text.rotary_emb
+        self.config = self.text.config
+        self.device = next(self.text.parameters()).device
+        self.dtype = next(self.text.parameters()).dtype
+
+
+def load_base(repo_id: str, device: str | torch.device, dtype: torch.dtype = torch.bfloat16) -> tuple[BaseModel, Path]:
+    snapshot = snapshot_dir(repo_id)
+    try:
+        model = AutoModelForCausalLM.from_pretrained(snapshot, dtype=dtype, attn_implementation="sdpa")
+    except Exception as exc:  # the text-only wrapper is not available for every transformers build
+        from transformers import Qwen3_5ForConditionalGeneration
+
+        print(f"AutoModelForCausalLM failed ({exc!r}); loading Qwen3_5ForConditionalGeneration")
+        model = Qwen3_5ForConditionalGeneration.from_pretrained(snapshot, dtype=dtype, attn_implementation="sdpa")
+    model.to(device).eval()
+    return BaseModel(model), snapshot
+
+
+def load_tokenizer(snapshot: Path):
+    return AutoTokenizer.from_pretrained(snapshot)
+
+
+class MtpHead(nn.Module):
+    """The one-layer multi-token-prediction head of the Qwen3.5 checkpoints."""
+
+    def __init__(self, text_config):
+        super().__init__()
+        import copy
+
+        config = copy.deepcopy(text_config)
+        config.layer_types = ["full_attention"]
+        config.num_hidden_layers = 1
+        config._attn_implementation = "sdpa"
+        hidden, eps = config.hidden_size, config.rms_norm_eps
+        self.fc = nn.Linear(2 * hidden, hidden, bias=False)
+        self.pre_fc_norm_hidden = Qwen3_5RMSNorm(hidden, eps=eps)
+        self.pre_fc_norm_embedding = Qwen3_5RMSNorm(hidden, eps=eps)
+        self.layers = nn.ModuleList([Qwen3_5DecoderLayer(config, 0)])
+        self.norm = Qwen3_5RMSNorm(hidden, eps=eps)
+
+    @classmethod
+    def from_snapshot(cls, snapshot: Path, text_config, device, dtype: torch.dtype) -> MtpHead:
+        weight_map = json.loads((snapshot / "model.safetensors.index.json").read_text())["weight_map"]
+        by_shard: dict[str, list[str]] = {}
+        for name, shard in weight_map.items():
+            if name.startswith(MTP_PREFIX):
+                by_shard.setdefault(shard, []).append(name)
+        if not by_shard:
+            raise ValueError(f"{snapshot} has no {MTP_PREFIX}* tensors")
+        state = {}
+        for shard, names in by_shard.items():
+            with safe_open(str(snapshot / shard), framework="pt", device="cpu") as f:
+                for name in names:
+                    state[name[len(MTP_PREFIX):]] = f.get_tensor(name)
+        head = cls(text_config)
+        head.load_state_dict(state, strict=True)
+        return head.to(device=device, dtype=dtype).eval()
+
+    def weight_bytes(self) -> int:
+        return sum(p.numel() * p.element_size() for p in self.parameters())
+
+
+class HFTarget:
+    """Qwen3.5 driven one token at a time through transformers' cache. hidden_source selects which hidden state
+    the draft sees: "post" (after the final norm) or "pre" (the last decoder layer's output)."""
+
+    def __init__(self, base: BaseModel, hidden_source: str = "post"):
+        if hidden_source not in ("post", "pre"):
+            raise ValueError(hidden_source)
+        self.base = base
+        self.hidden_source = hidden_source
+        self.cache: DynamicCache | None = None
+        self._pre_norm: torch.Tensor | None = None
+        self._hook = None
+        if hidden_source == "pre":
+            self._hook = base.text.norm.register_forward_pre_hook(self._capture)
+
+    def _capture(self, module, args):
+        self._pre_norm = args[0]
+
+    def close(self) -> None:
+        if self._hook is not None:
+            self._hook.remove()
+            self._hook = None
+
+    def _forward(self, ids: torch.Tensor) -> tuple[torch.Tensor, torch.Tensor]:
+        out = self.base.text(input_ids=ids, past_key_values=self.cache, use_cache=True)
+        post = out.last_hidden_state[0]
+        hidden = self._pre_norm[0] if self.hidden_source == "pre" else post
+        logits = self.base.lm_head(post[-1]).float()
+        return logits, hidden
+
+    @torch.inference_mode()
+    def prefill(self, tokens: list[int]) -> tuple[torch.Tensor, torch.Tensor]:
+        self.cache = DynamicCache(config=self.base.config)
+        ids = torch.tensor([tokens], dtype=torch.long, device=self.base.device)
+        return self._forward(ids)
+
+    @torch.inference_mode()
+    def step(self, token: int) -> tuple[torch.Tensor, torch.Tensor]:
+        ids = torch.tensor([[token]], dtype=torch.long, device=self.base.device)
+        logits, hidden = self._forward(ids)
+        return logits, hidden[-1]
+
+
+class HFDraft:
+    """The MTP head as a draft model with its own KV cache. fc_order "emb_first" concatenates the normed embedding
+    before the normed hidden state; "hidden_first" the other way round."""
+
+    def __init__(self, head: MtpHead, base: BaseModel, fc_order: str = "emb_first"):
+        if fc_order not in ("emb_first", "hidden_first"):
+            raise ValueError(fc_order)
+        self.head = head
+        self.base = base
+        self.fc_order = fc_order
+        self.cache = DynamicCache()
+        self.length = 0
+        self.last_hidden: torch.Tensor | None = None
+
+    def reset(self) -> None:
+        self.cache = DynamicCache()
+        self.length = 0
+        self.last_hidden = None
+
+    @torch.inference_mode()
+    def _run(self, hiddens: torch.Tensor, tokens: torch.Tensor) -> torch.Tensor:
+        n = tokens.shape[0]
+        emb = self.head.pre_fc_norm_embedding(self.base.embed(tokens))
+        hid = self.head.pre_fc_norm_hidden(hiddens.to(self.base.dtype))
+        parts = (emb, hid) if self.fc_order == "emb_first" else (hid, emb)
+        x = self.head.fc(torch.cat(parts, dim=-1))[None]  # (1, n, hidden)
+        positions = torch.arange(self.length, self.length + n, device=x.device)
+        cos, sin = self.base.rotary(x, positions.view(1, 1, n).expand(3, 1, n))
+        mask = None
+        if n > 1 and self.length > 0:
+            # Several new queries over an existing cache: sdpa's is_causal flag aligns top-left, so spell the mask
+            # out, as an additive mask so that the eager attention path reads it too.
+            total = self.length + n
+            allowed = (torch.arange(total, device=x.device)[None, :] <= positions[:, None]).view(1, 1, n, total)
+            mask = torch.zeros(allowed.shape, device=x.device, dtype=x.dtype)
+            mask.masked_fill_(~allowed, torch.finfo(x.dtype).min)
+        y = self.head.layers[0](
+            x,
+            position_embeddings=(cos, sin),
+            attention_mask=mask,
+            position_ids=positions[None],
+            past_key_values=self.cache,
+            use_cache=True,
+        )
+        g = self.head.norm(y)[0]
+        self.length += n
+        self.last_hidden = g[-1]
+        return self.base.lm_head(g[-1]).float()
+
+    def extend(self, hiddens: torch.Tensor, tokens: torch.Tensor) -> torch.Tensor:
+        return self._run(hiddens.to(self.base.device), tokens.to(self.base.device))
+
+    def chain(self, token: int) -> torch.Tensor:
+        tokens = torch.tensor([token], dtype=torch.long, device=self.base.device)
+        return self._run(self.last_hidden[None], tokens)
+
+    def rollback(self, n: int) -> None:
+        if n <= 0:
+            return
+        if n > self.length:
+            raise ValueError(f"rollback of {n} positions from a cache of {self.length}")
+        self.cache.crop(-n)
+        self.length -= n
diff --git a/benchmark/spec_decode/sampled_dist.py b/benchmark/spec_decode/sampled_dist.py
new file mode 100644
index 0000000000..4aa6c75544
--- /dev/null
+++ b/benchmark/spec_decode/sampled_dist.py
@@ -0,0 +1,149 @@
+"""Whether a server samples text from the same distribution as another, for one that cannot reproduce the other's
+draws at a seed (speculative decoding consumes its random numbers differently from plain decoding).
+
+    python -m benchmark.spec_decode.sampled_dist save --url http://127.0.0.1:8100 --out spec.json
+    python -m benchmark.spec_decode.sampled_dist compare base.json spec.json [--tokenizer Qwen/Qwen3.5-0.8B]
+
+``save`` sends each of the first ``--prompts`` chat prompts ``--samples`` times (one seed per sample) at temperature
+0.7, top-p 0.9 and top-k 20 for ``--max-tokens`` tokens. ``compare`` re-tokenizes every completion and, per prompt and
+position 1..N-1 (position 0 is the prompt's own token, sampled the same way on both paths), tests the two token
+marginals with a permutation test of their total-variation distance. Under the null each p-value is uniform, so the
+fraction below 0.01 should be near 1%. The same test between two halves of the first run (``--null``) shows the
+level the harness reaches on identical servers.
+"""
+
+from __future__ import annotations
+
+import argparse
+import collections
+import concurrent.futures as cf
+import json
+import random
+import sys
+from pathlib import Path
+
+import requests
+
+from benchmark.spec_decode.serve_parity import prompt_set
+
+
+def one(url: str, model: str, prompt: str, seed: int, args) -> str:
+    body = {
+        "model": model, "messages": [{"role": "user", "content": prompt}], "max_tokens": args.max_tokens,
+        "temperature": args.temperature, "top_p": args.top_p, "top_k": args.top_k, "seed": seed,
+        "stream": False, "enable_thinking": False,
+    }
+    r = requests.post(f"{url}/v1/chat/completions", json=body, timeout=args.timeout)
+    r.raise_for_status()
+    return r.json()["choices"][0]["message"]["content"]
+
+
+def save(args) -> int:
+    prompts = [p for _, p, _ in prompt_set()][: args.prompts]
+    jobs = [(i, s) for i in range(len(prompts)) for s in range(args.samples)]
+    with cf.ThreadPoolExecutor(args.concurrency) as pool:
+        texts = list(pool.map(lambda js: one(args.url, args.model, prompts[js[0]], 1000 * js[0] + js[1], args), jobs))
+    out = {"settings": {k: getattr(args, k) for k in ("temperature", "top_p", "top_k", "max_tokens", "samples")},
+           "texts": [[texts[i * args.samples + s] for s in range(args.samples)] for i in range(len(prompts))]}
+    Path(args.out).write_text(json.dumps(out))
+    print(f"{len(jobs)} samples over {len(prompts)} prompts")
+    return 0
+
+
+def tv(a: list[int], b: list[int]) -> float:
+    ca, cb = collections.Counter(a), collections.Counter(b)
+    return 0.5 * sum(abs(ca[t] / len(a) - cb[t] / len(b)) for t in set(ca) | set(cb))
+
+
+def perm_p(a: list[int], b: list[int], rounds: int, rng: random.Random) -> float:
+    observed, pooled, n = tv(a, b), a + b, len(a)
+    hits = 0
+    for _ in range(rounds):
+        rng.shuffle(pooled)
+        hits += tv(pooled[:n], pooled[n:]) >= observed - 1e-12
+    return (hits + 1) / (rounds + 1)
+
+
+def tests(xs: list[list[list[int]]], ys: list[list[list[int]]], positions: int, rounds: int, seed: int):
+    rng = random.Random(seed)
+    rows = []
+    for i, (x, y) in enumerate(zip(xs, ys, strict=True)):
+        for p in range(1, positions):
+            a = [s[p] for s in x if len(s) > p]
+            b = [s[p] for s in y if len(s) > p]
+            if len(a) < 10 or len(b) < 10:
+                continue
+            rows.append((i, p, tv(a, b), perm_p(a, b, rounds, rng)))
+    return rows
+
+
+def fisher(ps: list[float]) -> float:
+    """Fisher's combined p-value: -2 sum(ln p) is chi-squared with 2m degrees of freedom."""
+    import math
+
+    if not ps:
+        return 1.0
+    x, m = -2 * sum(math.log(p) for p in ps), len(ps)
+    # chi-squared survival at even degrees of freedom, in closed form
+    term = total = math.exp(-x / 2)
+    for i in range(1, m):
+        term *= x / (2 * i)
+        total += term
+    return min(total, 1.0)
+
+
+def summary(name: str, rows) -> str:
+    ps = [r[3] for r in rows]
+    small = sum(p < 0.01 for p in ps)
+    share = 100 * small / max(len(rows), 1)
+    return (f"{name}: {len(rows)} (prompt, position) tests, p<0.01 in {small} ({share:.1f}%), "
+            f"p<0.05 in {sum(p < 0.05 for p in ps)}, mean TV {sum(r[2] for r in rows) / max(len(rows), 1):.3f}, "
+            f"min p {min(ps, default=1):.4f}, Fisher combined p {fisher(ps):.3g}")
+
+
+def compare(args) -> int:
+    from transformers import AutoTokenizer
+
+    tok = AutoTokenizer.from_pretrained(args.tokenizer)
+    a, b = (json.loads(Path(p).read_text()) for p in (args.a, args.b))
+    enc = [[[tok.encode(t, add_special_tokens=False) for t in prompt] for prompt in run["texts"]] for run in (a, b)]
+    rows = tests(enc[0], enc[1], args.positions, args.rounds, 0)
+    print(summary("a vs b", rows))
+    for r in rows:
+        if r[3] < 0.01:
+            print(f"  prompt {r[0]} position {r[1]}: TV {r[2]:.3f} p {r[3]:.4f}")
+    if args.null:
+        half = len(enc[0][0]) // 2
+        print(summary("a vs a (halves)", tests([p[:half] for p in enc[0]], [p[half:] for p in enc[0]],
+                                              args.positions, args.rounds, 1)))
+    return 0
+
+
+def main(argv=None) -> int:
+    ap = argparse.ArgumentParser(description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter)
+    sub = ap.add_subparsers(dest="cmd", required=True)
+    s = sub.add_parser("save")
+    s.add_argument("--url", required=True)
+    s.add_argument("--out", required=True)
+    s.add_argument("--model", default="qwen3_5")
+    s.add_argument("--prompts", type=int, default=8)
+    s.add_argument("--samples", type=int, default=200)
+    s.add_argument("--max-tokens", type=int, default=6)
+    s.add_argument("--temperature", type=float, default=0.7)
+    s.add_argument("--top-p", type=float, default=0.9)
+    s.add_argument("--top-k", type=int, default=20)
+    s.add_argument("--concurrency", type=int, default=32)
+    s.add_argument("--timeout", type=float, default=300)
+    c = sub.add_parser("compare")
+    c.add_argument("a")
+    c.add_argument("b")
+    c.add_argument("--tokenizer", required=True)
+    c.add_argument("--positions", type=int, default=5)
+    c.add_argument("--rounds", type=int, default=2000)
+    c.add_argument("--null", action="store_true")
+    args = ap.parse_args(argv)
+    return save(args) if args.cmd == "save" else compare(args)
+
+
+if __name__ == "__main__":
+    sys.exit(main())
diff --git a/benchmark/spec_decode/serve_parity.py b/benchmark/spec_decode/serve_parity.py
new file mode 100644
index 0000000000..167ff91657
--- /dev/null
+++ b/benchmark/spec_decode/serve_parity.py
@@ -0,0 +1,148 @@
+"""Greedy outputs of a running server on T2's prompt set, and the comparison of two such runs.
+
+    python -m benchmark.spec_decode.serve_parity save --url http://127.0.0.1:8100 --out spec.json \\
+        --tokenizer Qwen/Qwen3.5-9B [--concurrency 1]
+    python -m benchmark.spec_decode.serve_parity compare base.json spec.json
+
+The set is the one T2's acceptance harness used: the first 32 prompts of the sweep's chat set (256 new tokens)
+and the 16 ~1k-token prompts (128 new), through the chat template with thinking off, greedy, streamed, so a
+speculative server's multi-token chunks go through the same SSE path a client reads. ``save`` also records each
+request's decode rate (generated tokens after the first over the time from the first chunk to the last), which
+at concurrency 1 is the per-token latency a client sees.
+"""
+
+from __future__ import annotations
+
+import argparse
+import concurrent.futures as cf
+import json
+import sys
+import time
+from pathlib import Path
+
+import requests
+
+REPO = Path(__file__).resolve().parents[2]
+CHAT = REPO / "benchmark/assets/simple_text_queries.txt"
+LONG = REPO / "benchmark/assets/long_text_queries_1k.txt"
+
+
+def prompt_set(n_chat: int = 32, n_long: int = 16, chat_new: int = 256, long_new: int = 128):
+    chat = [line.strip() for line in CHAT.read_text().splitlines() if line.strip()][:n_chat]
+    long = [line.strip() for line in LONG.read_text().splitlines() if line.strip()][:n_long]
+    return [(f"chat{i}", p, chat_new) for i, p in enumerate(chat)] + [
+        (f"long{i}", p, long_new) for i, p in enumerate(long)
+    ]
+
+
+def one(url: str, model: str, pid: str, prompt: str, max_new: int, timeout: float,
+        sampling: dict | None = None) -> dict:
+    body = {
+        "model": model, "messages": [{"role": "user", "content": prompt}], "max_tokens": max_new,
+        "temperature": 0.0, "stream": True, "enable_thinking": False, **(sampling or {}),
+    }
+    t0 = time.monotonic()
+    arrivals, parts = [], []
+    with requests.post(f"{url}/v1/chat/completions", json=body, stream=True, timeout=timeout) as r:
+        r.raise_for_status()
+        for raw in r.iter_lines():
+            if not raw or not raw.startswith(b"data: "):
+                continue
+            data = raw[6:]
+            if data == b"[DONE]":
+                break
+            doc = json.loads(data)
+            if "error" in doc:
+                raise RuntimeError(doc["error"])
+            for choice in doc.get("choices", []):
+                content = (choice.get("delta") or {}).get("content")
+                if content:
+                    parts.append(content)
+                    arrivals.append(time.monotonic())
+    return {"id": pid, "text": "".join(parts), "chunks": len(parts), "e2e_s": time.monotonic() - t0,
+            "ttft_s": arrivals[0] - t0 if arrivals else None,
+            "decode_s": arrivals[-1] - arrivals[0] if len(arrivals) > 1 else 0.0}
+
+
+def save(args) -> int:
+    from transformers import AutoTokenizer
+
+    tok = AutoTokenizer.from_pretrained(args.tokenizer)
+    items = prompt_set()
+    t0 = time.monotonic()
+    with cf.ThreadPoolExecutor(args.concurrency) as pool:
+        futures = [pool.submit(one, args.url, args.tokenizer, pid, p, n, args.timeout) for pid, p, n in items]
+        results = [f.result() for f in futures]
+    wall = time.monotonic() - t0
+    for r in results:
+        r["tokens"] = len(tok.encode(r["text"], add_special_tokens=False))
+        r["ms_per_token"] = 1000 * r["decode_s"] / (r["tokens"] - 1) if r["tokens"] > 1 else None
+    rates = sorted(r["ms_per_token"] for r in results if r["ms_per_token"])
+    sampled = []
+    for pid, p, _ in items[:args.sampled]:
+        # a smoke check of sampled decoding: it must answer, and answer differently each time
+        try:
+            texts = [one(args.url, args.tokenizer, pid, p, 64, args.timeout,
+                         {"temperature": 0.7})["text"] for _ in range(2)]
+            sampled.append({"id": pid, "texts": texts, "distinct": texts[0] != texts[1]})
+        except Exception as e:  # noqa: BLE001 - recorded, the greedy pass stands on its own
+            sampled.append({"id": pid, "error": f"{type(e).__name__}: {e}"})
+    out = {"url": args.url, "concurrency": args.concurrency, "wall_s": wall, "results": results,
+           "sampled": sampled,
+           "ms_per_token_p50": rates[len(rates) // 2] if rates else None,
+           "tokens_per_chunk": sum(r["tokens"] for r in results) / max(sum(r["chunks"] for r in results), 1)}
+    Path(args.out).write_text(json.dumps(out, indent=1))
+    print(f"{len(results)} prompts, wall {wall:.1f} s, decode p50 {out['ms_per_token_p50']} ms/token, "
+          f"{out['tokens_per_chunk']:.2f} tokens/chunk"
+          + (f", sampled {sum(1 for x in sampled if 'texts' in x)}/{len(sampled)} answered, "
+             f"{sum(1 for x in sampled if x.get('distinct'))} distinct" if sampled else ""))
+    return 0
+
+
+def compare(args) -> int:
+    from transformers import AutoTokenizer
+
+    a, b = (json.loads(Path(p).read_text()) for p in (args.a, args.b))
+    tok = AutoTokenizer.from_pretrained(args.tokenizer) if args.tokenizer else None
+    ra = {r["id"]: r for r in a["results"]}
+    same, diffs = 0, []
+    for r in b["results"]:
+        x, y = ra[r["id"]]["text"], r["text"]
+        if x == y:
+            same += 1
+            continue
+        div = next((i for i, (c1, c2) in enumerate(zip(x, y, strict=False)) if c1 != c2), min(len(x), len(y)))
+        entry = {"id": r["id"], "char": div, "a": x[max(div - 40, 0):div + 40], "b": y[max(div - 40, 0):div + 40]}
+        if tok is not None:
+            ia, ib = tok.encode(x, add_special_tokens=False), tok.encode(y, add_special_tokens=False)
+            pairs = enumerate(zip(ia, ib, strict=False))
+            entry["token"] = next((i for i, (t1, t2) in pairs if t1 != t2), min(len(ia), len(ib)))
+            entry["len_a"], entry["len_b"] = len(ia), len(ib)
+        diffs.append(entry)
+    print(f"identical {same}/{len(b['results'])}")
+    for d in diffs:
+        print(json.dumps(d))
+    return 0 if not diffs else 1
+
+
+def main(argv=None) -> int:
+    ap = argparse.ArgumentParser(description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter)
+    sub = ap.add_subparsers(dest="cmd", required=True)
+    s = sub.add_parser("save")
+    s.add_argument("--url", required=True)
+    s.add_argument("--out", required=True)
+    s.add_argument("--tokenizer", required=True, help="the served checkpoint's HF id")
+    s.add_argument("--concurrency", type=int, default=1)
+    s.add_argument("--timeout", type=float, default=600)
+    s.add_argument("--sampled", type=int, default=0,
+                   help="also send this many prompts twice at temperature 0.7 (64 tokens)")
+    c = sub.add_parser("compare")
+    c.add_argument("a")
+    c.add_argument("b")
+    c.add_argument("--tokenizer", default=None)
+    args = ap.parse_args(argv)
+    return save(args) if args.cmd == "save" else compare(args)
+
+
+if __name__ == "__main__":
+    sys.exit(main())
diff --git a/benchmark/spec_decode/spec_loop.py b/benchmark/spec_decode/spec_loop.py
new file mode 100644
index 0000000000..271789eb8b
--- /dev/null
+++ b/benchmark/spec_decode/spec_loop.py
@@ -0,0 +1,183 @@
+"""Speculative decoding with sequential verification, over abstract target and draft models.
+
+A step drafts k tokens, then feeds the target one token at a time: the target's distribution at a position is known
+before the draft there is judged, so a rejected draft is never fed to the target and the target's state is never
+rolled back. The target does exactly one forward per token it emits, which keeps every (k, sampling mode) cell cheap
+enough to run as the real process rather than a replay of one continuation. Only the draft's cache is rolled back:
+the positions it extended from its own hidden states are dropped and recomputed from the target's hidden states once
+the step's tokens are known, as a serving engine's draft pass does.
+"""
+from __future__ import annotations
+
+from dataclasses import dataclass, field
+from typing import Protocol
+
+import torch
+
+from benchmark.spec_decode.verify import Decision, SamplingParams, accept_draft, sample, sampling_probs
+
+
+class Target(Protocol):
+    def prefill(self, tokens: list[int]) -> tuple[torch.Tensor, torch.Tensor]:
+        """Process the prompt. Returns the next-token logits and the hidden state of every prompt position."""
+
+    def step(self, token: int) -> tuple[torch.Tensor, torch.Tensor]:
+        """Append one token. Returns the next-token logits and that position's hidden state."""
+
+
+class Draft(Protocol):
+    def reset(self) -> None: ...
+
+    def extend(self, hiddens: torch.Tensor, tokens: torch.Tensor) -> torch.Tensor:
+        """Append positions computed from the target's hidden states, each paired with the token that followed it.
+        Returns the logits of the draft for the token after the last one."""
+
+    def chain(self, token: int) -> torch.Tensor:
+        """Append one position computed from the draft's own last hidden state and a drafted token; returns the
+        logits for the next draft."""
+
+    def rollback(self, n: int) -> None:
+        """Drop the last n positions (a finished step's chain positions)."""
+
+
+@dataclass
+class RunStats:
+    """Acceptance accounting for one k, summed over prompts."""
+
+    k: int
+    prompts: int = 0
+    steps: int = 0
+    tokens: int = 0  # generated tokens, the first sampled token included
+    accepted_total: int = 0
+    # Per draft position j: steps that reached it (every earlier draft accepted), how many accepted it, and the sum
+    # of min(1, p/q) there.
+    evaluated: list[int] = field(default_factory=list)
+    accepted: list[int] = field(default_factory=list)
+    accept_prob_sum: list[float] = field(default_factory=list)
+
+    def __post_init__(self) -> None:
+        if not self.evaluated:
+            self.evaluated = [0] * self.k
+            self.accepted = [0] * self.k
+            self.accept_prob_sum = [0.0] * self.k
+
+    def record(self, position: int, decision: Decision) -> None:
+        self.evaluated[position] += 1
+        self.accepted[position] += int(decision.accepted)
+        self.accept_prob_sum[position] += decision.accept_prob
+
+    def merge(self, other: RunStats) -> None:
+        assert other.k == self.k
+        self.prompts += other.prompts
+        self.steps += other.steps
+        self.tokens += other.tokens
+        self.accepted_total += other.accepted_total
+        for j in range(self.k):
+            self.evaluated[j] += other.evaluated[j]
+            self.accepted[j] += other.accepted[j]
+            self.accept_prob_sum[j] += other.accept_prob_sum[j]
+
+    @property
+    def mean_accepted(self) -> float:
+        return self.accepted_total / max(self.steps, 1)
+
+    @property
+    def tokens_per_step(self) -> float:
+        """Tokens a step emits: the accepted drafts plus the target's own token (bonus or correction)."""
+        return 1.0 + self.mean_accepted
+
+    def position_rates(self) -> list[float]:
+        """Acceptance at position j given the step reached it."""
+        return [self.accepted[j] / self.evaluated[j] if self.evaluated[j] else float("nan") for j in range(self.k)]
+
+    def position_expected_rates(self) -> list[float]:
+        """The same, estimated from min(1, p/q) instead of the coin flips."""
+        return [
+            self.accept_prob_sum[j] / self.evaluated[j] if self.evaluated[j] else float("nan") for j in range(self.k)
+        ]
+
+    def reach_rates(self) -> list[float]:
+        """Fraction of steps that accepted at least j+1 drafts."""
+        return [self.accepted[j] / max(self.steps, 1) for j in range(self.k)]
+
+    def summary(self) -> dict:
+        return {
+            "k": self.k,
+            "prompts": self.prompts,
+            "steps": self.steps,
+            "tokens": self.tokens,
+            "mean_accepted": self.mean_accepted,
+            "tokens_per_step": self.tokens_per_step,
+            "position_rates": self.position_rates(),
+            "position_expected_rates": self.position_expected_rates(),
+            "reach_rates": self.reach_rates(),
+        }
+
+
+def run_speculative(
+    target: Target,
+    draft: Draft,
+    prompt: list[int],
+    k: int,
+    params: SamplingParams,
+    max_new_tokens: int,
+    stop_ids: frozenset[int] = frozenset(),
+    generator: torch.Generator | None = None,
+) -> tuple[list[int], RunStats]:
+    """Generate up to max_new_tokens after the prompt with k drafts per step. Returns the generated tokens (the stop
+    token included when one was reached) and the step statistics."""
+    stats = RunStats(k=k, prompts=1)
+    logits, hiddens = target.prefill(prompt)
+    first = sample(sampling_probs(logits, params), params, generator)
+    generated = [first]
+    draft.reset()
+    # Draft position p pairs the target's hidden state at p with the token at p+1.
+    q_logits = draft.extend(hiddens, torch.tensor(prompt[1:] + [first], dtype=torch.long))
+    pending = first  # emitted, not yet fed to the target
+
+    def finished() -> bool:
+        return len(generated) >= max_new_tokens or generated[-1] in stop_ids
+
+    while not finished():
+        drafts: list[int] = []
+        draft_dists: list[torch.Tensor] = []
+        for j in range(k):
+            q = sampling_probs(q_logits, params)
+            drafts.append(sample(q, params, generator))
+            draft_dists.append(q)
+            if j + 1 < k:
+                q_logits = draft.chain(drafts[j])
+
+        logits, h = target.step(pending)
+        new_hiddens = [h]
+        n_accepted = 0
+        next_token: int | None = None
+        for j in range(k):
+            decision = accept_draft(drafts[j], draft_dists[j], sampling_probs(logits, params), generator)
+            stats.record(j, decision)
+            if not decision.accepted:
+                next_token = decision.replacement
+                break
+            n_accepted += 1
+            generated.append(drafts[j])
+            if finished():
+                break
+            logits, h = target.step(drafts[j])
+            new_hiddens.append(h)
+        else:
+            # Every draft accepted: the target's distribution after the last one gives a bonus token.
+            next_token = sample(sampling_probs(logits, params), params, generator)
+        stats.steps += 1
+        stats.accepted_total += n_accepted
+        if next_token is None:
+            break  # a stop token or the length cap landed on an accepted draft
+        generated.append(next_token)
+        if finished():
+            break
+        draft.rollback(k - 1)
+        new_tokens = generated[-(n_accepted + 1):]
+        q_logits = draft.extend(torch.stack(new_hiddens), torch.tensor(new_tokens, dtype=torch.long))
+        pending = next_token
+
+    stats.tokens = len(generated)
+    return generated, stats
diff --git a/benchmark/spec_decode/verify.py b/benchmark/spec_decode/verify.py
new file mode 100644
index 0000000000..3359b20726
--- /dev/null
+++ b/benchmark/spec_decode/verify.py
@@ -0,0 +1,74 @@
+"""Draft verification for speculative decoding: the sampler's distribution and the per-position accept decision
+(speculative sampling, Leviathan et al. 2023 / Chen et al. 2023).
+
+Greedy decoding is the temperature-0 case. Both distributions collapse to one-hot vectors, so a draft is accepted
+exactly when it is the target's argmax and the replacement on rejection is that argmax; one code path serves both
+modes and the greedy decision does not depend on what the generator draws.
+"""
+from __future__ import annotations
+
+from dataclasses import dataclass
+
+import torch
+
+
+@dataclass(frozen=True)
+class SamplingParams:
+    temperature: float = 0.0
+    top_p: float = 1.0
+
+    @property
+    def greedy(self) -> bool:
+        return self.temperature <= 0.0
+
+
+def sampling_probs(logits: torch.Tensor, params: SamplingParams) -> torch.Tensor:
+    """The distribution a sampler with these parameters draws from, as float32 probabilities over the last dim."""
+    logits = logits.float()
+    if params.greedy:
+        return torch.nn.functional.one_hot(logits.argmax(-1), logits.shape[-1]).to(torch.float32)
+    probs = torch.softmax(logits / params.temperature, dim=-1)
+    if params.top_p < 1.0:
+        sorted_probs, order = probs.sort(dim=-1, descending=True)
+        cumulative = sorted_probs.cumsum(-1)
+        # Keep the smallest prefix whose mass reaches top_p, including the token that crosses it.
+        sorted_probs = sorted_probs.masked_fill(cumulative - sorted_probs > params.top_p, 0.0)
+        probs = torch.zeros_like(probs).scatter_(-1, order, sorted_probs)
+        probs = probs / probs.sum(-1, keepdim=True)
+    return probs
+
+
+def sample(probs: torch.Tensor, params: SamplingParams, generator: torch.Generator | None) -> int:
+    """One token from a probability vector."""
+    if params.greedy:
+        return int(probs.argmax().item())
+    return int(torch.multinomial(probs, 1, generator=generator).item())
+
+
+@dataclass(frozen=True)
+class Decision:
+    accepted: bool
+    # On rejection, the token drawn from the residual distribution norm(max(target - draft, 0)).
+    replacement: int | None
+    # min(1, target(token) / draft(token)): the probability of accepting this draft given the two distributions,
+    # which is the acceptance rate's low-variance estimator.
+    accept_prob: float
+
+
+def accept_draft(
+    token: int, draft_probs: torch.Tensor, target_probs: torch.Tensor, generator: torch.Generator | None
+) -> Decision:
+    """Judge one draft token against the target's distribution at the same position."""
+    q = float(draft_probs[token])
+    p = float(target_probs[token])
+    accept_prob = 1.0 if q <= 0.0 else min(1.0, p / q)
+    u = torch.rand((), device=draft_probs.device, generator=generator).item()
+    if u < accept_prob:
+        return Decision(True, None, accept_prob)
+    residual = (target_probs - draft_probs).clamp_(min=0.0)
+    total = residual.sum()
+    if total <= 0.0:
+        # Identical distributions can still reject on rounding; the target itself is the correct fallback then.
+        residual, total = target_probs, target_probs.sum()
+    replacement = int(torch.multinomial(residual / total, 1, generator=generator).item())
+    return Decision(False, replacement, accept_prob)
diff --git a/configs/qwen3_5_0.8b_mtp2.yaml b/configs/qwen3_5_0.8b_mtp2.yaml
new file mode 100644
index 0000000000..9debe22d4c
--- /dev/null
+++ b/configs/qwen3_5_0.8b_mtp2.yaml
@@ -0,0 +1,20 @@
+model: "qwen3_5_0.8b"
+max_seq_len: 32768
+node_groups:
+  # Speculative decoding is checked on text prompts; images still prefill
+  # through the encoder, unwarmed.
+  - node_names: [LLM, vision_encoder]
+    ranks: [0]
+model_kwargs:
+  # Speculative decoding: the checkpoint's MTP head drafts 2 tokens a step
+  # and each decode step verifies them (see SpecDecodeResource).
+  num_draft_tokens: 2
+  # drafts read an fp8 copy of lm_head (see Fp8DraftHead), the verify the checkpoint's
+  draft_head: fp8
+  # verify steps may sample the target too (a second decode capture)
+  spec_sampling: true
+resources:
+  gdn_state:
+    # A request holds num_draft_tokens + 1 state slots: 64 requests, plus the sink.
+    max_slots: 193
+    # state_dtype: bfloat16  # default: the checkpoint's mamba_ssm_dtype (float32)
diff --git a/configs/qwen3_5_0.8b_mtp3.yaml b/configs/qwen3_5_0.8b_mtp3.yaml
new file mode 100644
index 0000000000..67ec4d6b46
--- /dev/null
+++ b/configs/qwen3_5_0.8b_mtp3.yaml
@@ -0,0 +1,18 @@
+model: "qwen3_5_0.8b"
+max_seq_len: 32768
+node_groups:
+  # Speculative decoding is checked on text prompts; images still prefill
+  # through the encoder, unwarmed.
+  - node_names: [LLM, vision_encoder]
+    ranks: [0]
+model_kwargs:
+  # Speculative decoding: the checkpoint's MTP head drafts 3 tokens a step
+  # and each decode step verifies them (see SpecDecodeResource).
+  num_draft_tokens: 3
+  # drafts read an fp8 copy of lm_head (see Fp8DraftHead), the verify the checkpoint's
+  draft_head: fp8
+resources:
+  gdn_state:
+    # A request holds num_draft_tokens + 1 state slots: 64 requests, plus the sink.
+    max_slots: 257
+    # state_dtype: bfloat16  # default: the checkpoint's mamba_ssm_dtype (float32)
diff --git a/configs/qwen3_5_9b_mtp2.yaml b/configs/qwen3_5_9b_mtp2.yaml
new file mode 100644
index 0000000000..c9114e4e48
--- /dev/null
+++ b/configs/qwen3_5_9b_mtp2.yaml
@@ -0,0 +1,20 @@
+model: "qwen3_5_9b"
+max_seq_len: 32768
+node_groups:
+  # Speculative decoding is checked on text prompts; images still prefill
+  # through the encoder, unwarmed.
+  - node_names: [LLM, vision_encoder]
+    ranks: [0]
+model_kwargs:
+  # Speculative decoding: the checkpoint's MTP head drafts 2 tokens a step
+  # and each decode step verifies them (see SpecDecodeResource).
+  num_draft_tokens: 2
+  # drafts read an fp8 copy of lm_head (see Fp8DraftHead), the verify the checkpoint's
+  draft_head: fp8
+  # verify steps may sample the target too (a second decode capture)
+  spec_sampling: true
+resources:
+  gdn_state:
+    # A request holds num_draft_tokens + 1 state slots: 64 requests, plus the sink.
+    max_slots: 193
+    # state_dtype: bfloat16  # default: the checkpoint's mamba_ssm_dtype (float32)
diff --git a/configs/qwen3_5_9b_mtp3.yaml b/configs/qwen3_5_9b_mtp3.yaml
new file mode 100644
index 0000000000..1eab9c8b62
--- /dev/null
+++ b/configs/qwen3_5_9b_mtp3.yaml
@@ -0,0 +1,18 @@
+model: "qwen3_5_9b"
+max_seq_len: 32768
+node_groups:
+  # Speculative decoding is checked on text prompts; images still prefill
+  # through the encoder, unwarmed.
+  - node_names: [LLM, vision_encoder]
+    ranks: [0]
+model_kwargs:
+  # Speculative decoding: the checkpoint's MTP head drafts 3 tokens a step
+  # and each decode step verifies them (see SpecDecodeResource).
+  num_draft_tokens: 3
+  # drafts read an fp8 copy of lm_head (see Fp8DraftHead), the verify the checkpoint's
+  draft_head: fp8
+resources:
+  gdn_state:
+    # A request holds num_draft_tokens + 1 state slots: 64 requests, plus the sink.
+    max_slots: 257
+    # state_dtype: bfloat16  # default: the checkpoint's mamba_ssm_dtype (float32)
diff --git a/docs/adding_models.rst b/docs/adding_models.rst
index 24fdf1d795..a5fdc23536 100644
--- a/docs/adding_models.rst
+++ b/docs/adding_models.rst
@@ -792,6 +792,12 @@ The spec types are:
        so a prefill capture bucket larger than ``max_captured_prefill_tokens`` (512
        by default) runs eagerly; smaller ones capture a sequential kernel whose cost
        grows with the longest row.
+   * - ``SpecDecodeSpec(config=SpecDecodeConfig(num_draft_tokens=k), kv_cache=..., attn=...)``
+     - Speculative decoding: one row of device registers per request (its committed
+       length, how much its last verify step accepted, the next step's tokens, its
+       budget) and the in-graph operations a verify step and its draft passes call. See
+       `Speculative decoding`_. Declare it only on a deployment that drafts; nothing
+       else changes when it is absent.
    * - ``SamplerSpec(vocab_size=..., enable_repetion_penalty=...)``
      - A sampler with its own per-request parameter buffers, philox stream, and optional
        seen-token mask. ``enable_repetion_penalty`` declares a capability, not an
@@ -2425,6 +2431,64 @@ Run it with:
 
    mstar-serve --config configs/your_model.yaml --host 0.0.0.0 --port 8000
 
+Speculative decoding
+~~~~~~~~~~~~~~~~~~~~
+
+A model whose checkpoint carries a draft (Qwen3.5's one-layer MTP head) can decode
+speculatively: each decode step verifies ``k`` drafted tokens with the target and drafts
+the next ``k``, inside one captured graph and with no host sync.
+``SpecDecodeResource`` (``mstar/engine/resources/spec_decode/``) has the design in its
+module docstring. What a model declares and calls:
+
+- **Resources.** A ``SpecDecodeSpec`` over the node's paged KV and attention
+  resources; the recurrent pool with ``RecurrentStateConfig(spec_tokens=k)`` (a request
+  holds ``k + 1`` state slots, and conv windows grow by ``k``) and the linear-attention
+  spec with ``LinearAttnConfig(spec_decode=<key>)``; one more KV layer for the draft
+  head's attention; no combined walk mixing prompts with decode rows. ``k`` is a
+  model kwarg (``model_kwargs: {num_draft_tokens: k}`` in the serving config), since
+  these shapes are sized from it.
+- **Decode steps.** Every row is ``k + 1`` tokens. ``declare_step`` gives each row a
+  ``SpecDecodeStep(phase=VERIFY)`` with ``retract`` set from
+  ``pending_retract(rid)``, the same amounts negated on ``KVStep.commit_adjust``, and
+  ``AttentionStep(row_queries=k + 1)``; no sampler or position step. The forward calls
+  ``verify_inputs()`` (tokens and positions), runs the target, ``accept(logits)``,
+  then for each draft pass the draft head and ``record_draft(j, logits)``. Pass 0 runs
+  over the verify's tokens and the others one token a row, ``j`` positions past each
+  row's last accepted slot (``draft_slots()``). Every pass attends one query a row:
+  ``attend_draft(j, qkv, ...)`` takes the layer's packed projection and does its split,
+  norms, rotation and KV write in one kernel (``mstar::qkv_norm_rope_cache``) before
+  the attention, so the head's work after attention, and its vocabulary projection, run
+  at ``[rows]``. ``record_draft`` writes the token to the next step's register in the
+  same kernels that take the argmax, and returns it as the next pass's input, so nothing
+  between passes reads the host.
+- **Capture.** ``decode_rows=True`` on the decode ``BatchedAcceleratorGraphConfig``,
+  so a platform compiles the verify bucket as decode, by its rows.
+- **The draft's head.** A draft pass needs one argmax, and the target checks it, so
+  ``Fp8DraftHead`` (``mstar/engine/resources/spec_decode/draft_head.py``) can hold an fp8
+  copy of the vocabulary head for the draft alone: half the bytes of the pass's largest
+  kernel, for a little acceptance and no correctness. Verification always reads the
+  checkpoint's head. A vocabulary subset is not offered: it saves as much on Latin-script
+  text and loses a third to a half of the accepted drafts on other scripts.
+- **The prompt's last step.** It declares ``SpecDecodeStep(phase=SEED)`` (budgets
+  and ``ignore_eos`` rows) and ``KVStep.commit_adjust`` of ``seed_lookahead`` tokens,
+  runs the draft head over the prompt, and calls ``seed(first_token, first_draft)``.
+- **Stops.** ``accept`` returns each row's emitted tokens (-1 past the last), how
+  many positions the target accepted and a ``done`` flag (a stop id or the budget);
+  ``check_stop_batched`` reads ``done`` and hands the accepted counts to ``learn``,
+  which is how the host gives back the KV it reserved optimistically. A step emits a
+  varying number of tokens, so it takes the per-request tensor path to the client,
+  not the inline one, and the API side drops the -1s.
+
+The async worker plans step N+1 while N runs, so a verify step is planned against the
+length it would have if N accepted every draft, and its KV write slots, attention
+lengths, rotary positions and recurrent start slots are corrected from the registers
+on the device. That needs an attention backend whose plan reads per-row KV lengths and
+a block table on the device: AITER's unified attention on AMD GPUs. Verification is
+greedy, or with ``SpecDecodeConfig(sampling=True)`` a draw of the target's token at each
+position from the request's top-k/top-p distribution at its temperature, by the
+sampler's captured kernel (``draw_tokens``; with argmax drafts, accepting while the draw
+equals the next draft is rejection sampling). No ``repetition_penalty``.
+
 Tuning resources per deployment
 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
 
diff --git a/docs/compile_cache.rst b/docs/compile_cache.rst
new file mode 100644
index 0000000000..2b54f27482
--- /dev/null
+++ b/docs/compile_cache.rst
@@ -0,0 +1,109 @@
+Compile-cache bundles
+=====================
+
+A first start compiles every captured forward with Inductor (tuning it at ``-O1`` and
+``-O2``), Triton builds the kernels, and on ROCm MIOpen searches its convolution solvers.
+All three cache what they produce (:ref:`serving:Restarts and the compile cache`), so a
+start on the same host skips that work. A *bundle* carries those caches to another host or
+container with the same software stack, so its first start is warm too.
+
+.. code-block:: bash
+
+   # on a host that has served the configs you want to ship warm
+   mstar cache pack /shared/mstar-bundles --note "qwen3_5_9b, whisper_large"
+
+   # on the new host, before serving
+   mstar cache unpack /shared/mstar-bundles
+   # or in one step
+   mstar serve qwen3_5_9b --compile-bundle /shared/mstar-bundles
+
+The stack key
+-------------
+
+A bundle is only useful on the stack that compiled it, so bundles are kept per key:
+
+.. code-block:: text
+
+   <product>/torch-<torch version>-<hip|cuda>-<version>-triton-<triton version>
+   mi325x-vf/torch-2.12.1+rocm7.2-hip-7.2.53211-triton-3.7.1
+
+``mstar cache key`` prints this host's. The product is the GPU's name, lower-cased with
+``AMD Instinct`` dropped. The manifest also records the GPU arch (``gfx942``, ``sm_90``).
+
+- The arch, torch, HIP or CUDA, and Triton versions must match. Inductor and Triton key
+  their own entries on them, so a bundle from another stack would unpack without error and
+  never be read. ``mstar cache unpack`` refuses it (``--force`` unpacks it anyway).
+- The product doesn't have to match. Autotuning picks are timings taken on the product
+  that compiled them: an MI300X bundle loads on an MI325X (both gfx942) and starts warm, but
+  its kernels were tuned on the MI300X. The startup line says so.
+- Nothing in the key covers the driver, ``ptxas`` or the M* version. A newer M* that changes
+  a captured forward's graph compiles that graph again on first start and keeps the rest.
+
+Layout
+------
+
+A bundle root holds one directory per key; each is a bundle:
+
+.. code-block:: text
+
+   <root>/<product>/<torch-...-triton-...>/
+       manifest.json        the key, member count and size per tar, and one history entry per pack
+       inductor.tar         TORCHINDUCTOR_CACHE_DIR, less a Triton dir inside it
+       triton.tar           TRITON_CACHE_DIR
+       miopen.tar           MIOPEN_USER_DB_PATH (ROCm)
+       miopen-kernels.tar   MIOPEN_CUSTOM_CACHE_DIR, only when it is a different directory
+       .append.lock         taken by every pack
+
+Each tar is uncompressed and holds its directory's files as ``./<path>``, so it can be
+read and merged with plain ``tar``. ``unpack`` and ``pack`` find these directories the way
+a server started from the same environment does: an exported variable, else M*'s default
+for Inductor and Triton (``~/.cache/mstar/torchinductor`` and its ``triton``
+subdirectory, see :doc:`environment_variables`), else the library's own (MIOpen's
+``~/.config/miopen`` and ``~/.cache/miopen``). ``mstar cache dirs`` prints them.
+
+Not in a bundle: FlashInfer's JIT builds (``FLASHINFER_WORKSPACE_BASE``), AITER's and
+FlyDSL's on ROCm (``AITER_JIT_DIR``, ``FLYDSL_RUNTIME_CACHE_DIR``), which are keyed by their
+own package versions rather than by this key, and the weights.
+
+Building a bundle
+-----------------
+
+#. Point the caches at empty directories (or start from a bundle for the same key), for
+   example ``export XDG_CACHE_HOME=$(mktemp -d)``.
+#. Start each config you want warm and send it a request of every shape that compiles on
+   first use, such as each image size a diffusion model will serve. Most of the work
+   happens before ``/health`` answers: every capture bucket is compiled at startup.
+#. Stop the servers, then ``mstar cache pack ROOT``. Packing must not run while a server
+   is writing to the caches.
+
+``pack`` adds to the bundle for this host's key under ``ROOT``, creating it if needed. It
+appends the Inductor and Triton files the bundle doesn't have and those whose size or
+modification time differ from its last copy; their entries never change once written.
+MIOpen keeps each database in one file (the find and perf databases as ``key=value``
+text, compiled kernels in SQLite), so a changed one is merged with the bundle's copy,
+this host's record winning where both have one: caches warmed separately, each from
+empty, add up rather than replace each other. Each tar is copied, appended to and renamed
+into place, so a concurrent ``unpack`` never reads a partial tar, and the lock serializes
+packs, so several hosts with the same key can add to one bundle in turn.
+
+Unpacking
+---------
+
+``mstar cache unpack SRC`` takes a bundle or a root of bundles, picks the one for this
+host's key, and extracts each tar into its cache directory over what is there; where both
+have a file, the bundle's copy wins, and with repeated copies in a tar the last one does.
+It exits with status 3 when ``SRC`` has no bundle for this stack, and ``mstar serve
+--compile-bundle`` then starts cold with a warning. ``unpack`` leaves
+``mstar-bundle.json`` in the Inductor directory, naming the bundle and its key.
+
+At startup each worker logs its stack key and what it found:
+
+.. code-block:: text
+
+   Worker audio_encoder: compile cache: stack mi325x-vf/torch-...-triton-3.7.1; bundle /shared/mstar-bundles/mi325x-vf/torch-... unpacked 2026-10-29T21:04:11+00:00 matches
+   Worker decoder: compile cache: stack mi325x-vf/...; no bundle unpacked into /root/.cache/mstar/torchinductor
+
+A bundle built for another arch, torch, HIP, CUDA or Triton logs a warning instead, since
+its entries won't be used.
+
+``mstar cache show ROOT`` lists the bundles under a root with their sizes and last packs.
diff --git a/docs/index.rst b/docs/index.rst
index 6d962ea388..d1f4779bfb 100644
--- a/docs/index.rst
+++ b/docs/index.rst
@@ -30,6 +30,7 @@ vision-language-action policies, and world models — through a **Python SDK**,
    :caption: User Guide
 
    serving
+   compile_cache
    metrics
    clients
    evaluation
diff --git a/docs/models.rst b/docs/models.rst
index 2214c84758..d56d904b45 100644
--- a/docs/models.rst
+++ b/docs/models.rst
@@ -544,6 +544,21 @@ supported yet. On any deployment,
 ``"return_token_ids": true`` adds each text chunk's sampled ids
 (``{"token_ids": [...]}``), stop tokens included, so a trainer need not
 re-tokenize the text.
+Speculative decoding with the checkpoints' MTP head (AMD GPUs): serve
+``configs/qwen3_5_{0.8b,9b}_mtp{2,3}.yaml``, which set ``model_kwargs:
+{num_draft_tokens: k}``. Each decode step verifies ``k`` drafts of the head and
+drafts the next ``k``, in one captured graph, and streams the 1 to ``k + 1``
+tokens it accepts as one chunk. A request holds ``k + 1`` recurrent-state slots,
+so ``gdn_state.max_slots`` is ``1 + concurrency * (k + 1)`` (the configs: 64
+requests); the KV cache gains one layer for the head. The ``mtp3`` configs verify
+greedily and refuse ``temperature > 0``; the ``mtp2`` ones set ``spec_sampling:
+true`` and also sample (temperature, ``top_p``, ``top_k``; a second decode capture),
+seeded by the request but not reproducing the plain path's draws. Neither takes
+``repetition_penalty``. The configs set ``draft_head: fp8``: drafts read an fp8 copy
+of ``lm_head`` (one rank; ``draft_head: bf16`` keeps one head), which gives up about
+1 % of the accepted drafts for a step about 6 % cheaper at 9B c=1 on MI325X, and the
+verify always reads the full head. They need
+AITER's attention (device-side KV lengths) and are checked on text prompts. See "Speculative decoding" in :doc:`adding_models` for the design.
 
 Qwen3.8-27B (``qwen3_8_27b``) is the Qwen3.5 architecture and runs on the same
 code, served from ``configs/qwen3_8_27b_tp4.yaml`` (TP4)::
diff --git a/docs/serving.rst b/docs/serving.rst
index d03b2cdbd7..2ff8d9b54f 100644
--- a/docs/serving.rst
+++ b/docs/serving.rst
@@ -813,6 +813,9 @@ next start's compile. On a quota-limited or network-mounted home (NFS), or a vol
 concurrent jobs on several hosts share, export ``TORCHINDUCTOR_CACHE_DIR`` to a node-local
 path and copy a warm cache in at job start.
 
+To give a new host or container these caches warm, pack them into a bundle for the stack
+on a host that has served the configs, and unpack it on the new one before serving
+(:doc:`compile_cache`).
 .. _host-idle:
 
 Host idle states
diff --git a/mstar/cli/cache.py b/mstar/cli/cache.py
new file mode 100644
index 0000000000..a98fdd83f1
--- /dev/null
+++ b/mstar/cli/cache.py
@@ -0,0 +1,101 @@
+"""``mstar cache``: compile-cache bundles (docs/compile_cache.rst).
+
+    mstar cache key                  # this host's stack key, the bundle path a pack writes
+    mstar cache dirs                 # the cache dirs a server started here reads and writes
+    mstar cache pack ROOT            # after a warm-up: add this host's caches to ROOT/<key>/
+    mstar cache unpack SRC           # before serving: the bundle for this stack, into the cache dirs
+    mstar cache show ROOT            # the bundles under ROOT
+"""
+
+from __future__ import annotations
+
+import argparse
+import sys
+import time
+from pathlib import Path
+
+# unpack's exit status when SRC has no bundle for this stack, so callers can serve cold
+EXIT_NO_MATCH = 3
+
+
+def _key(args: argparse.Namespace) -> None:
+    from mstar.utils import compile_bundle
+
+    key = compile_bundle.stack_key()
+    print(key.path)
+    print(f"  product {key.product}, arch {key.arch}, torch {key.torch}, {key.runtime}, triton {key.triton}")
+
+
+def _dirs(args: argparse.Namespace) -> None:
+    from mstar.utils import compile_bundle
+
+    for kind, path in compile_bundle.cache_dirs().items():
+        print(f"{kind:15s} {path}")
+
+
+def _pack(args: argparse.Namespace) -> None:
+    from mstar.utils import compile_bundle
+
+    dirs = compile_bundle.cache_dirs()
+    kinds = args.kinds.split(",") if args.kinds else None
+    unknown = sorted(set(kinds or ()) - set(dirs))
+    if unknown:
+        sys.exit(f"error: no {', '.join(unknown)} cache here; kinds here: {', '.join(dirs)}")
+    t0 = time.monotonic()
+    bundle, _ = compile_bundle.pack(Path(args.root), compile_bundle.stack_key(), dirs, kinds=kinds,
+                                    note=args.note)
+    print(f"packed into {bundle} in {time.monotonic() - t0:.1f} s")
+
+
+def _unpack(args: argparse.Namespace) -> None:
+    from mstar.utils import compile_bundle
+
+    t0 = time.monotonic()
+    try:
+        done = compile_bundle.unpack(Path(args.src), compile_bundle.stack_key(), compile_bundle.cache_dirs(),
+                                     force=args.force)
+    except compile_bundle.BundleMismatchError as exc:
+        print(f"not unpacked: {exc}", file=sys.stderr)
+        sys.exit(EXIT_NO_MATCH)
+    kinds = ", ".join(f"{k} -> {d}" for k, d in done.kinds.items())
+    print(f"unpacked {done.bundle} in {time.monotonic() - t0:.1f} s: {kinds}")
+
+
+def _show(args: argparse.Namespace) -> None:
+    from mstar.utils import compile_bundle
+
+    root = Path(args.root)
+    bundles = compile_bundle.find_bundles(root)
+    if not bundles:
+        sys.exit(f"no bundles under {root}")
+    for b in bundles:
+        manifest = compile_bundle.read_manifest(b) or {}
+        tars = sorted(b.glob("*.tar"))
+        size = sum(t.stat().st_size for t in tars) / 2**30
+        print(f"{b}: {size:.2f} GiB in {', '.join(t.stem for t in tars)}; "
+              f"updated {manifest.get('updated', 'unknown (no manifest)')}")
+        for h in manifest.get("history", [])[-3:]:
+            print(f"  {h['time']} {h['host']} mstar {h.get('mstar')}: added {h['added']}"
+                  + (f" ({h['note']})" if h.get("note") else ""))
+
+
+def add_parser(sub: argparse._SubParsersAction) -> None:
+    cache = sub.add_parser(
+        "cache", help="pack or unpack compile-cache bundles, so a new host or container starts warm",
+        description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter,
+    )
+    csub = cache.add_subparsers(dest="cache_command", required=True)
+    csub.add_parser("key", help="this host's stack key (GPU, torch, HIP or CUDA, Triton)").set_defaults(func=_key)
+    csub.add_parser("dirs", help="the cache directories a server started here uses").set_defaults(func=_dirs)
+    pack = csub.add_parser("pack", help="add this host's caches to the bundle for its stack under ROOT")
+    pack.add_argument("root", help="bundle root; the bundle is ROOT/<stack key>/")
+    pack.add_argument("--note", default=None, help="recorded in the manifest, e.g. which configs were warmed")
+    pack.add_argument("--kinds", default=None, help="comma list of caches to pack (default: all here)")
+    pack.set_defaults(func=_pack)
+    unpack = csub.add_parser("unpack", help="extract the bundle for this stack into the cache dirs")
+    unpack.add_argument("src", help="a bundle, or a root holding one per stack")
+    unpack.add_argument("--force", action="store_true", help="unpack a bundle built for another stack")
+    unpack.set_defaults(func=_unpack)
+    show = csub.add_parser("show", help="list the bundles under ROOT")
+    show.add_argument("root")
+    show.set_defaults(func=_show)
diff --git a/mstar/cli/main.py b/mstar/cli/main.py
index 37746f0703..1d468926b0 100644
--- a/mstar/cli/main.py
+++ b/mstar/cli/main.py
@@ -4,6 +4,8 @@
     mstar serve qwen3_omni --gpus 0,1,2
     mstar serve orpheus --port 9000
 
+``mstar cache`` packs and unpacks compile-cache bundles (mstar/cli/cache.py).
+
 ``serve <model>`` resolves a sensible default config, fills in the
 plumbing that the low-level ``mstar-serve`` requires (socket/upload dirs, a
 single-node-safe tensor protocol, HF cache), and then delegates to the same
@@ -17,11 +19,13 @@
 import getpass
 import os
 import shutil
+import subprocess
 import sys
 import tempfile
 from pathlib import Path
 
 import mstar
+from mstar.cli import cache
 from mstar.utils.socket_prefix import per_port_socket_prefix
 
 # Model name -> default config (relative to the repo's configs/).
@@ -248,6 +252,14 @@ def _serve(args: argparse.Namespace) -> None:
     apply_visible_devices(args.gpus)
     if args.compile_level is not None:
         os.environ["MSTAR_COMPILE_LEVEL"] = str(args.compile_level)
+    if args.compile_bundle:
+        # A child process, so this one never touches the GPU before the server's own processes do.
+        p = subprocess.run([sys.executable, "-m", "mstar.cli.main", "cache", "unpack", args.compile_bundle],
+                           check=False)
+        if p.returncode == cache.EXIT_NO_MATCH:
+            print("warning: no compile-cache bundle for this stack; starting cold", file=sys.stderr)
+        elif p.returncode != 0:
+            sys.exit(f"error: unpacking {args.compile_bundle} failed (exit {p.returncode})")
 
     who = getpass.getuser() if hasattr(getpass, "getuser") else "mstar"
     socket_prefix = args.socket_path_prefix or per_port_socket_prefix(args.port)
@@ -322,6 +334,11 @@ def build_parser() -> argparse.ArgumentParser:
              "fastest start; -O1 (default) coordinate descent; -O2 full max-autotune",
     )
     serve.add_argument("--config", default=None, help="override the default config (path to YAML)")
+    serve.add_argument(
+        "--compile-bundle", default=None, metavar="SRC",
+        help="unpack this compile-cache bundle (or the one for this stack in a root of bundles) into the "
+             "compile caches before starting; see `mstar cache`",
+    )
     serve.add_argument("--cache-dir", default=None, help="HuggingFace weight cache directory")
     serve.add_argument(
         "--socket-path-prefix", default=None,
@@ -362,6 +379,7 @@ def build_parser() -> argparse.ArgumentParser:
              "rust/server/target/release)",
     )
     serve.set_defaults(func=_serve)
+    cache.add_parser(sub)
     return parser
 
 
diff --git a/mstar/engine/accelerator_graph_config.py b/mstar/engine/accelerator_graph_config.py
index 81c9b6d419..cc21a31ac5 100644
--- a/mstar/engine/accelerator_graph_config.py
+++ b/mstar/engine/accelerator_graph_config.py
@@ -74,6 +74,11 @@ def __init__(
         # which can pick the wrong dim when an unrelated axis happens to
         # match the bucket's token count.
         input_seq_dims: dict[str, int] | None = None,
+        # Compile one graph whose bucket-varying dims are dynamic, rather than
+        # one per bucket (size 1 still gets its own). For forwards whose
+        # kernels gain little from a fixed size, e.g. an encoder over
+        # fixed-size windows; each bucket is still captured on its own.
+        compile_dynamic: bool = False,
     ):
         self.capture_graph_walk = capture_graph_walk
         self.replay_graph_walks = replay_graph_walks or [capture_graph_walk]
@@ -90,6 +95,7 @@ def __init__(
         # scales by this. Subclasses that support it set it.
         self.total_tokens_multiplier = 1
         self.input_seq_dims = input_seq_dims
+        self.compile_dynamic = compile_dynamic
 
     @abstractmethod
     def get_config_type(self) -> AcceleratorGraphConfigType:
@@ -118,6 +124,8 @@ def __init__(
         total_tokens_multiplier: int = 1,
         compile_mode: str | None = None,
         input_seq_dims: dict[str, int] | None = None,
+        compile_dynamic: bool = False,
+        decode_rows: bool = False,
     ):
         super().__init__(
             capture_graph_walk=capture_graph_walk,
@@ -129,8 +137,12 @@ def __init__(
             capture_forward_method=capture_forward_method,
             caps_eager_batch_size=caps_eager_batch_size,
             input_seq_dims=input_seq_dims,
+            compile_dynamic=compile_dynamic,
         )
         self.single_request_inputs = single_request_inputs
+        # A few tokens a row that are decode all the same (a speculative verify
+        # row): the platform compiles the bucket by its rows, as decode.
+        self.decode_rows = decode_rows
         # ``single_request_inputs.input_seq_len`` is also read per-label by the
         # submodule's own ``declare_step`` (e.g. replicated across cond/uncond
         # labels for a combined guidance step), so it must stay a per-label
diff --git a/mstar/engine/accelerator_graph_runner.py b/mstar/engine/accelerator_graph_runner.py
index f8b6dff102..4bd4f428a7 100644
--- a/mstar/engine/accelerator_graph_runner.py
+++ b/mstar/engine/accelerator_graph_runner.py
@@ -119,6 +119,20 @@ def compile_kwargs(config_mode: str | None) -> dict[str, Any]:
     return kwargs
 
 
+def mark_bucket_dims_dynamic(
+    static_inputs: Mapping[str, Any], seq_dims: Mapping[tuple[int, str], int],
+    config_idx: int, num_tokens: int,
+) -> None:
+    """Hint dynamo that each static input's bucket-varying dim (the one its
+    shared buffer is sliced along) is dynamic, so one compiled graph serves
+    every bucket. Only a hint: a forward that needs the size fixed (a reshape
+    to a constant, say) still compiles per bucket, and size 1 always does."""
+    for key, value in static_inputs.items():
+        dim = seq_dims.get((config_idx, key))
+        if isinstance(value, torch.Tensor) and dim is not None and value.shape[dim] == num_tokens:
+            torch._dynamo.maybe_mark_dynamic(value, dim)
+
+
 def blocked_capture_resources(step, resources: Mapping[str, Resource]) -> list[str]:
     """Return resource keys that prevent capture of the declared step."""
     if step is None:
@@ -1249,6 +1263,10 @@ def prepare() -> dict[str, Any]:
                 key for key, value in static_inputs.items()
                 if isinstance(value, torch.Tensor)
             )
+            if config.compile and config.compile_dynamic:
+                mark_bucket_dims_dynamic(
+                    static_inputs, self._static_buffer_seq_dims, spec.config_idx, spec.num_tokens,
+                )
 
             forward = self._forward_for(spec)
             planner = self._resident
@@ -1310,7 +1328,8 @@ def _forward_for(self, spec: CGSlotSpec):
 
         The platform's inductor overrides wrap that shared callable per bucket,
         so each bucket compiles on its first call under the overrides for its
-        own token count.
+        own token count. A ``compile_dynamic`` config's graph serves every
+        bucket, so it takes the overrides for any count.
         """
         forward = getattr(self._submodule, spec.config.capture_forward_method)
         if not spec.config.compile:
@@ -1319,8 +1338,13 @@ def _forward_for(self, spec: CGSlotSpec):
             self._compiled_forwards[spec.config_idx] = torch.compile(
                 forward, **compile_kwargs(spec.config.compile_mode),
             )
+        if spec.config.compile_dynamic:
+            return with_recompile_limit(
+                with_inductor_overrides(self._compiled_forwards[spec.config_idx]), self._submodule,
+            )
         return with_recompile_limit(with_inductor_overrides(
             self._compiled_forwards[spec.config_idx], spec.num_tokens, spec.bs,
+            getattr(spec.config, "decode_rows", False),
         ), self._submodule)
 
     def _dummy_engine_inputs(
diff --git a/mstar/engine/resources/attn/aiter.py b/mstar/engine/resources/attn/aiter.py
index 662a1ae6ed..7197cb6c22 100644
--- a/mstar/engine/resources/attn/aiter.py
+++ b/mstar/engine/resources/attn/aiter.py
@@ -418,6 +418,13 @@ def _unified(
 _DELEGATED = object()
 
 
+# The context a captured decode step's kernel variant is chosen for. The cache's whole
+# context (32768) launches 128 segments a row for keys that rarely fill two. MI325X, 4
+# K/V heads: 1 row 13.2 -> 11.0 us at 200 keys, 19.4 -> 16.2 at 4k, 105 -> 98 at 32k, 16
+# rows 20.5 -> 17.2, 64 rows level, 4 rows at 16k the one loss (177 -> 189).
+DECODE_SEGMENT_PIN = 4096
+
+
 class AiterPagedAttentionManager(_UnifiedPlans, AttentionManager):
     """Paged attention on AITER's unified attention (see the module docstring).
 
@@ -457,6 +464,10 @@ def __init__(
         self._preplan_delegated = False
         # FlashInfer manager for non-causal steps with multi-query rows
         self._fallback: AttentionManager | None = None
+        # [0, 1, ..., rows], for a step that attends one query a row off a plan
+        # staged with more (`attend_rows`); grown when planning, never in a graph
+        self._row_starts: torch.Tensor | None = None
+        self._retired_row_starts: list[torch.Tensor | None] = []
 
     def depends_on(self):
         return {self._kv_cache_name}
@@ -512,7 +523,7 @@ def plan(self, step: AttentionStep, ctx: StepContext):
             self._preplan_plans = {}
             self._preplanned = False
             self._preplan_delegated = False
-            return
+            return self._current_plans
 
         plan_outputs: KVPlanOutputs = ctx.plan_results.get(self._kv_cache_name)
         assert plan_outputs is not None, (
@@ -534,7 +545,9 @@ def plan(self, step: AttentionStep, ctx: StepContext):
             plans = dict.fromkeys(rows, _DELEGATED)
         elif lease is not None:
             plans = {
-                label: self._stage_captured(label, r, lease, ctx.force_prefill)
+                label: self._stage_captured(
+                    label, r, lease, ctx.force_prefill, step.row_queries,
+                )
                 for label, r in rows.items()
             }
         else:
@@ -546,6 +559,9 @@ def plan(self, step: AttentionStep, ctx: StepContext):
             self._preplanned = True
         else:
             self._current_plans = plans
+        # label -> `_Staged` (or the fallback's marker): the device views a
+        # speculative step rewrites in-graph (`SpecDecodeResource`)
+        return plans
 
     def _check_kind(self, lease: SlotLease, label: str, kind: object) -> None:
         key = CGSlotKey(bucket=lease.bucket, slot=lease.slot, label=label)
@@ -559,6 +575,7 @@ def _check_kind(self, lease: SlotLease, label: str, kind: object) -> None:
 
     def _stage_captured(
         self, label: str, rows: _Rows, lease: SlotLease, force_prefill: bool,
+        row_queries: int | None = None,
     ) -> _Staged:
         num_rows = rows.qo_indptr.shape[0] - 1
         # a decode bucket always has one token per row, while one row of a
@@ -566,6 +583,15 @@ def _stage_captured(
         # as many rows as tokens and replays mixed steps (`force_prefill`)
         num_tokens = lease.bucket.num_tokens
         max_q = 1 if num_rows == num_tokens and not force_prefill else num_tokens
+        if row_queries is not None and not force_prefill:
+            # every row declared the same query count (a verify step)
+            max_q = row_queries
+            self._reserve_row_starts(num_rows)
+            if num_rows and int(np.diff(rows.qo_indptr).max()) > row_queries:
+                raise RuntimeError(
+                    f"AITER attention: bucket {lease.bucket} planned a row with "
+                    f"more than the {row_queries} queries its step declared"
+                )
         if max_q == 1 and not single_query_rows(rows.qo_indptr):
             raise RuntimeError(
                 f"AITER attention: decode bucket {lease.bucket} planned a row "
@@ -573,13 +599,24 @@ def _stage_captured(
             )
         self._check_kind(lease, label, (num_rows, max_q))
         image = self._cg_image(lease, label, _image_numel(num_rows, self._cg_width))
+        # A decode row's pin only picks how many segments split its keys, and the
+        # kernel covers a longer row by widening them, so a row past this still reads
+        # every key.
+        max_k = self._kv_config.max_seq_len
         return stage_plan(
             image, rows, self._cg_width, self._pad_cols,
-            max_q, self._kv_config.max_seq_len,
+            max_q, min(max_k, DECODE_SEGMENT_PIN) if max_q == 1 else max_k,
         )
 
+    def _reserve_row_starts(self, rows: int) -> None:
+        if self._row_starts is None or self._row_starts.numel() <= rows:
+            # captured graphs keep reading the buffer they recorded
+            self._retired_row_starts.append(self._row_starts)
+            self._row_starts = torch.arange(max(rows, 128) + 1, dtype=torch.int32, device=self._device)
+
     def _stage_eager(self, label: str, rows: _Rows, slot: int) -> _Staged:
         num_rows = rows.qo_indptr.shape[0] - 1
+        self._reserve_row_starts(num_rows)
         counts = np.diff(rows.kv_indptr)
         width = (int(counts.max()) if num_rows else 0) + self._pad_cols
         image = self._eager_image(label, slot, _image_numel(num_rows, width))
@@ -636,6 +673,19 @@ def _attend(
             return self._fallback.run(q, label, kv_cache_layer)
         return self._unified(q, plan, kv_cache_layer)
 
+    @torch.compiler.disable
+    def attend_rows(
+        self, q: torch.Tensor, kv_cache_layer: torch.Tensor, kv_lens: torch.Tensor,
+        label: str = "main",
+    ) -> torch.Tensor:
+        """One query a row over this step's block table, each row reading
+        ``kv_lens`` keys, whatever query count the step was planned for."""
+        rows = q.shape[0]
+        plan = self._current_plans[label]
+        return self._unified(q, plan._replace(
+            cu_seqlens_q=self._row_starts[: rows + 1], seqused_k=kv_lens, max_seqlen_q=1, num_tokens=rows,
+        ), kv_cache_layer)
+
 
 class AiterCrossManager(_UnifiedPlans, CrossAttentionManager):
     """Non-causal attention over a write-once context, on AITER's unified
diff --git a/mstar/engine/resources/attn/config.py b/mstar/engine/resources/attn/config.py
index cb088c7ddd..e750c9eb13 100644
--- a/mstar/engine/resources/attn/config.py
+++ b/mstar/engine/resources/attn/config.py
@@ -258,6 +258,10 @@ class AttentionStep(ResourceStep):
     # `segments` then carry the query counts, the cache's segments the appended counts.
     # Only the MLA backend reads it; the others refuse it (`refuse_context_only`).
     context_only: bool = False
+    # every row of the step carries this many queries (a speculative verify
+    # step's k + 1), so a captured bucket may size its kernel for that rather
+    # than for one row holding all of the bucket's tokens
+    row_queries: int | None = None
 
 
 def refuse_context_only(step: AttentionStep, manager: str) -> None:
diff --git a/mstar/engine/resources/attn/prologue.py b/mstar/engine/resources/attn/prologue.py
new file mode 100644
index 0000000000..2c5b2b13ed
--- /dev/null
+++ b/mstar/engine/resources/attn/prologue.py
@@ -0,0 +1,121 @@
+"""The attention prologue of a decode step as one kernel.
+
+A packed ``[tokens, q(+gate) | k | v]`` row is split, q and k are RMS-normalised
+per head (gemma style: scaled by ``1 + weight``), the leading rotary dims of
+each are rotated (NeoX halves, read from cos/sin tables at the step's
+positions), and k and v are written into the paged cache at one flat slot per
+token. A decode step otherwise runs a norm-and-rope kernel per tensor and a
+cache scatter for each attention layer.
+
+``mstar::qkv_norm_rope_cache`` runs AITER's ``fused_qkv_split_qk_norm_rope_cache``
+on AMD GPUs, the only place a caller reaches it: speculative decoding needs
+AITER's attention.
+"""
+
+from __future__ import annotations
+
+import torch
+
+from mstar.kernels.numerics import Tolerance
+from mstar.kernels.ops import KernelOp, ShapePolicy
+
+
+def _split(qkv, heads, kv_heads, head_dim, gated):
+    tokens = qkv.shape[0]
+    q_size = heads * head_dim * (2 if gated else 1)
+    kv_size = kv_heads * head_dim
+    q, k, v = qkv.split([q_size, kv_size, kv_size], dim=-1)
+    if gated:
+        # the checkpoint interleaves query and gate per head
+        q, gate = q.reshape(tokens, heads, 2 * head_dim).chunk(2, dim=-1)
+        gate = gate.contiguous()
+    else:
+        q, gate = q.reshape(tokens, heads, head_dim), qkv.new_empty(0)
+    return q, gate, k.reshape(tokens, kv_heads, head_dim), v.reshape(tokens, kv_heads, head_dim)
+
+
+def _reference(qkv, q_weight, k_weight, cos, sin, positions, key_cache, value_cache, slots,
+               heads, kv_heads, head_dim, eps, gated):
+    q, gate, k, v = _split(qkv, heads, kv_heads, head_dim, gated)
+    half = cos.shape[-1]
+    c, s = cos[positions].float()[:, None, :], sin[positions].float()[:, None, :]
+
+    def norm_rope(x, weight):
+        xf = x.float()
+        xf = xf * torch.rsqrt(xf.pow(2).mean(-1, keepdim=True) + eps) * (1.0 + weight.float())
+        x1, x2 = xf[..., :half], xf[..., half:2 * half]
+        return torch.cat((x1 * c - x2 * s, x2 * c + x1 * s, xf[..., 2 * half:]), dim=-1).to(x.dtype)
+
+    q, k = norm_rope(q, q_weight), norm_rope(k, k_weight)
+    block = key_cache.shape[1]
+    pages, offsets = slots // block, slots % block
+    key_cache[pages, offsets] = k.to(key_cache.dtype)
+    value_cache[pages, offsets] = v.to(value_cache.dtype)
+    # fresh outputs: a one-token row's v slice is contiguous, and would alias qkv
+    return q, gate, k, v.clone(memory_format=torch.contiguous_format)
+
+
+def _fake(qkv, q_weight, k_weight, cos, sin, positions, key_cache, value_cache, slots,
+          heads, kv_heads, head_dim, eps, gated):
+    tokens = qkv.shape[0]
+    gate = qkv.new_empty((tokens, heads, head_dim)) if gated else qkv.new_empty(0)
+    return (qkv.new_empty((tokens, heads, head_dim)), gate,
+            qkv.new_empty((tokens, kv_heads, head_dim)), qkv.new_empty((tokens, kv_heads, head_dim)))
+
+
+def _aiter():
+    from aiter.ops.triton.rope.fused_qkv_split_qk_norm_rope_cache import fused_qkv_split_qk_norm_rope_cache
+
+    def qkv_norm_rope_cache(qkv, q_weight, k_weight, cos, sin, positions, key_cache, value_cache, slots,
+                            heads, kv_heads, head_dim, eps, gated):
+        out = fused_qkv_split_qk_norm_rope_cache(
+            qkv, q_weight, k_weight, cos, sin, positions, key_cache, value_cache, slots,
+            heads, kv_heads, head_dim, is_neox=True, reuse_freqs_front_part=True,
+            attn_output_gate=gated, eps=eps, gated_qkv_layout="interleaved", kv_cache_layout="NHD",
+        )
+        if gated:
+            return out
+        q, k, v = out
+        return q, qkv.new_empty(0), k, v
+
+    return qkv_norm_rope_cache
+
+
+def _inputs(shape, dtype, device, generator=None):
+    tokens, heads, kv_heads, head_dim, half = (shape[d] for d in ("tokens", "heads", "kv_heads", "head_dim", "half"))
+    page_size, pages = 16, 8
+
+    def rand(*size):
+        return torch.randn(*size, dtype=dtype, device=device, generator=generator)
+
+    qkv = rand(tokens, (2 * heads + 2 * kv_heads) * head_dim)
+    steps = torch.arange(1, half + 1, device=device)
+    angle = torch.arange(64, device=device, dtype=torch.float32)[:, None] * 0.01 * steps
+    cache = torch.zeros(pages, 2, page_size, kv_heads, head_dim, dtype=dtype, device=device)
+    slots = torch.randperm(pages * page_size, device=device, generator=generator)[:tokens]
+    positions = torch.randint(0, 64, (tokens,), device=device, generator=generator)
+    return (qkv, rand(head_dim) * 0.1, rand(head_dim) * 0.1, angle.cos().to(dtype), angle.sin().to(dtype),
+            positions, cache[:, 0], cache[:, 1], slots, heads, kv_heads, head_dim, 1e-6, True)
+
+
+QKV_NORM_ROPE_CACHE = KernelOp(
+    "qkv_norm_rope_cache",
+    "(Tensor qkv, Tensor q_weight, Tensor k_weight, Tensor cos, Tensor sin, Tensor positions, "
+    "Tensor(a!) key_cache, Tensor(b!) value_cache, Tensor slots, int heads, int kv_heads, int head_dim, "
+    "float eps, bool gated) -> (Tensor, Tensor, Tensor, Tensor)",
+    fake=_fake,
+    reference=_reference,
+    impls={"rocm": _aiter},
+    shapes=ShapePolicy(
+        dims=("tokens", "heads", "kv_heads", "head_dim", "half"),
+        key=lambda qkv, q_weight, k_weight, cos, *rest: {
+            "tokens": qkv.shape[0], "heads": rest[-5], "kv_heads": rest[-4], "head_dim": rest[-3],
+            "half": cos.shape[-1]},
+        inputs=_inputs,
+        warm=({"tokens": 1, "heads": 16, "kv_heads": 4, "head_dim": 256, "half": 32},
+              {"tokens": 64, "heads": 16, "kv_heads": 4, "head_dim": 256, "half": 32}),
+        dtypes=(torch.bfloat16,),
+    ),
+    # fp32 norm and rotation, one rounding (or two: the kernel rounds the normalised q before rotating)
+    tolerance=Tolerance.accumulated(rtol=1.6e-2, atol_scale=1e-2),
+)
diff --git a/mstar/engine/resources/kv/config.py b/mstar/engine/resources/kv/config.py
index cf7eda44e6..edcfb18ce7 100644
--- a/mstar/engine/resources/kv/config.py
+++ b/mstar/engine/resources/kv/config.py
@@ -765,6 +765,12 @@ class KVStep(ResourceStep):
     # Set by the manager when the step commits, while copies to the host are on: its count of commits so far.
     # The submodule doesn't declare it, as it doesn't declare `SubmoduleStep`'s context.
     commits: int = field(default=0, init=False, compare=False, repr=False)
+    # (request_id, label) -> tokens this commit adds beyond the segment's span:
+    # positive reserves room past it (admit sizes the pages for it too),
+    # negative gives back tokens an earlier commit over-counted. Speculative
+    # decoding commits the most a verify step could keep and settles the
+    # difference here once the host has learned it.
+    commit_adjust: Mapping[tuple[str, str], int] = field(default_factory=dict)
 
 
 @dataclass(frozen=True, kw_only=True)
diff --git a/mstar/engine/resources/kv/manager.py b/mstar/engine/resources/kv/manager.py
index ddb5b0e17e..bf6ec0d5ad 100644
--- a/mstar/engine/resources/kv/manager.py
+++ b/mstar/engine/resources/kv/manager.py
@@ -2074,6 +2074,7 @@ def admit(self, step: KVStep, ctx: StepContext) -> AdmitOutcome:
             # the counts are in tokens; the setting names the rows whose KV
             # grows by them, so they are a span this label can reserve
             reserve_ahead = ctx.reserve_whole_prefill
+            adjust = step.commit_adjust
             ahead = ctx.chunk_tokens_after \
                 if reserve_ahead.value is not SettingValue.UNSET else None
             # One pass for the reservations and the streams the in-flight mark
@@ -2086,7 +2087,7 @@ def admit(self, step: KVStep, ctx: StepContext) -> AdmitOutcome:
             # chunked row reserves ahead, so those steps take every row here.
             ready: list[CacheStream] = []
             rows = step.segments
-            if not ctx.capture and not ahead and len(self._groups) == 1:
+            if not ctx.capture and not ahead and not adjust and len(self._groups) == 1:
                 ready, slow = self._arena.admit_rows(streams, rows, ctx.real_rows(), sink or None, page_size)
                 rows = [rows[i] for i in slow]
             for segment in rows:
@@ -2109,6 +2110,8 @@ def admit(self, step: KVStep, ctx: StepContext) -> AdmitOutcome:
                 if ahead and reserve_ahead.holds_for(rid):
                     span += ahead.get(rid, 0)
                 seq_len = span + stream.stored_len
+                if adjust:
+                    seq_len += max(adjust.get((rid, label), 0), 0)
                 # the common decode row: its pages already cover the token it
                 # appends, so there is nothing to acquire and nothing to record
                 if (
@@ -2459,6 +2462,7 @@ def _setup_plan_states(
                     pages=indptrs.decode_pages, offsets=indptrs.decode_offsets,
                 )
                 self._staged_outputs[(lease.slot, label)] = indptrs
+                indptrs.write_state = plan_state
                 if ctx.is_preplan:
                     self._preplan_states[label] = plan_state
                 else:
@@ -2473,6 +2477,7 @@ def _setup_plan_states(
                     lease.bucket.num_tokens, page_size,
                 )
                 self._staged_outputs[(lease.slot, label)] = indptrs
+                indptrs.write_state = plan_state
                 if ctx.is_preplan:
                     self._preplan_states[label] = plan_state
                 else:
@@ -2495,6 +2500,7 @@ def _setup_plan_states(
             if eager_write and ctx.is_preplan and plan_state.token_to_page.is_cuda:
                 plan_state.ready = torch.cuda.Event()
                 plan_state.ready.record()
+            indptrs.write_state = plan_state
             if ctx.is_preplan:
                 self._preplan_states[label] = plan_state
             else:
@@ -3276,11 +3282,13 @@ def commit(self, step: KVStep, ctx: StepContext):
                 # another batch that is as big (surely sharing rows) or that
                 # names a cached row: the cache no longer describes the streams
                 self._plan_epoch += 1
-        if not step.commit:
+        if not step.commit or step.commit_adjust:
             self._plan_epoch += 1
         rows = self._readonly_rows
         if rows and any(seg.span and seg.request_id in rows for seg in step.segments):
             self._drop_readonly_plans()
+
+        adjust = step.commit_adjust
         # the index only sees real rows of a non-capture step; checked once
         index_on = self._index is not None and not ctx.capture
         windows = self._windows
@@ -3297,8 +3305,8 @@ def commit(self, step: KVStep, ctx: StepContext):
             is_pad = ctx.is_padding_row
             sink = step.sink
             rows = () if ctx.capture else step.segments
-            if rows and self._index is None and not retention and not windows and not _DEBUG_ASSERTS:
-                # nothing to index, adopt, retain or check: the arena moves
+            if rows and self._index is None and not retention and not windows and not adjust and not _DEBUG_ASSERTS:
+                # nothing to index, adopt, retain, adjust or check: the arena moves
                 # every row's counts in one call, and hands back the ones it can't
                 slow = self._arena.commit_rows(streams, rows, ctx.real_rows(), sink or None, step.commit)
                 rows = [rows[i] for i in slow]
@@ -3344,6 +3352,14 @@ def commit(self, step: KVStep, ctx: StepContext):
                             self._assert_layout(segment, stream, touched, ctx, policy)
                     if index_on and stream.chain is not None:
                         self._index_filled_pages(segment, stream, ctx, touched)
+                    if adjust:
+                        # after the layout and index saw the span; pages are a
+                        # high-water mark, so a tail given back keeps them
+                        stream.stored_len = max(
+                            stream.stored_len
+                            + adjust.get((segment.request_id, segment.label), 0),
+                            0,
+                        )
                     # so a claim taken in a window the mark misses still fails
                     # `_commit_offload`'s generation guard
                     stream.generation += 1
diff --git a/mstar/engine/resources/kv/plan.py b/mstar/engine/resources/kv/plan.py
index 9b555ff04c..992ccd8fe8 100644
--- a/mstar/engine/resources/kv/plan.py
+++ b/mstar/engine/resources/kv/plan.py
@@ -233,7 +233,7 @@ class KVPlanOutput:
 
     __slots__ = (
         "cpu_indptrs", "_views", "_view_builder", "decode_pages", "decode_offsets",
-        "rows", "_decode", "columns",
+        "rows", "_decode", "columns", "write_state",
     )
 
     def __init__(
@@ -261,6 +261,9 @@ def __init__(
         # columns. Unlike ``cpu_indptrs``, which may sit in a pinned ring a
         # later plan rewrites, they stay valid as long as the output does.
         self.columns = columns
+        # the device slots this step's writes go to (a `KVPlanState`), set by
+        # `KVManager.plan`; a speculative step rewrites them in-graph
+        self.write_state = None
 
     @property
     def views(self) -> "list[SequenceView]":
diff --git a/mstar/engine/resources/linear_attn/config.py b/mstar/engine/resources/linear_attn/config.py
index 3498cfc00d..0d29bcf5af 100644
--- a/mstar/engine/resources/linear_attn/config.py
+++ b/mstar/engine/resources/linear_attn/config.py
@@ -66,13 +66,20 @@ class LinearAttnConfig:
     # chunked prefill (FlashInfer's) has no bound and ignores this.
     max_captured_prefill_tokens: int | None = None
 
+    # GDN only: the speculative-decoding resource whose verify steps this one
+    # runs, k + 1 tokens a row from each request's last accepted state
+    spec_decode: str | None = None
+
 
 @dataclass
 class LinearAttnSpec(NodeResourceSpec):
     config: LinearAttnConfig
 
     def depends_on(self) -> set[str]:
-        return {self.config.recurrent_state}
+        keys = {self.config.recurrent_state}
+        if self.config.spec_decode is not None:
+            keys.add(self.config.spec_decode)
+        return keys
 
     def combined_splits(self) -> tuple[CombinedSplit, ...]:
         return self.config.splits
diff --git a/mstar/engine/resources/linear_attn/gdn.py b/mstar/engine/resources/linear_attn/gdn.py
index 56ee90b03f..1b17876bc6 100644
--- a/mstar/engine/resources/linear_attn/gdn.py
+++ b/mstar/engine/resources/linear_attn/gdn.py
@@ -22,6 +22,7 @@
     GDNDecodeWrapper,
     GDNMixedWrapper,
     GDNPrefillWrapper,
+    GDNVerifyWrapper,
     GDNWrapper,
 )
 from mstar.engine.resources.recurrent.config import DeltaNetGeometry
@@ -71,6 +72,7 @@ def __init__(
         self.state_dtype = state_dtype
         self._device = device
         self._pool_key = config.recurrent_state
+        self._spec_key = config.spec_decode
         self._has_sink = has_sink
         # chosen once, by `LinearAttnManager.build`; every wrapper runs these
         self._kernels = kernels
@@ -97,7 +99,10 @@ def __init__(
         self._cached_plan_output: dict[str, GDNWrapper] | None = None
 
     def depends_on(self) -> set[str]:
-        # the pool plans first; its addressing arrives via `ctx.plan_results`
+        # the pool plans first; its addressing arrives via `ctx.plan_results`,
+        # and a speculative step's start slots via the speculative resource's
+        if self._spec_key is not None:
+            return {self._pool_key, self._spec_key}
         return {self._pool_key}
 
     def resident_refusal(self) -> str | None:
@@ -132,9 +137,13 @@ def plan(self, step: LinearAttnStep, ctx: StepContext):
             return self._current
 
         addressing: dict[str, RecurrentAddressing] = ctx.plan_results[self._pool_key]
+        verify = self._verify_plan(ctx)
 
         self._current = {}
         for label, segments in self._group_by_label(step.segments or ()).items():
+            if verify is not None:
+                self._current[label] = self._build_verify(label, addressing[label], verify, ctx)
+                continue
             self._current[label] = self._build_plan(
                 label, segments, addressing[label], ctx
             )
@@ -156,6 +165,46 @@ def _group_by_label(segments) -> dict[str, list[Segment]]:
             out.setdefault(seg.label, []).append(seg)
         return out
 
+    def _verify_plan(self, ctx: StepContext):
+        """The speculative resource's plan, if this step verifies drafts."""
+        if self._spec_key is None:
+            return None
+        from mstar.engine.resources.spec_decode.config import SpecPhase
+
+        plan = ctx.plan_results.get(self._spec_key)
+        if plan is None or plan.phase is not SpecPhase.VERIFY:
+            return None
+        return plan
+
+    def _build_verify(
+        self, label: str, addressing: RecurrentAddressing, verify, ctx: StepContext,
+    ) -> GDNWrapper:
+        groups = addressing.host_slot_groups
+        if groups is None:
+            raise RuntimeError(
+                "a speculative verify step needs the pool's slot groups; declare "
+                "the pool with RecurrentStateConfig.spec_tokens"
+            )
+        lease = ctx.slot_lease
+        if lease is None:
+            key, store, bs = (label, "verify"), self._eager_wrappers, None
+        else:
+            key, store = CGSlotKey(lease.bucket, lease.slot, label), self._cg_wrappers
+            bs = max(self._cg_max_bs, lease.bucket.bs)
+        wrapper = store.get(key)
+        if wrapper is None:
+            wrapper = store[key] = GDNVerifyWrapper(
+                device=self._device,
+                pad_slot_id=PAD_SLOT_ID,
+                tokens_per_row=len(groups[0]) if groups else 1,
+                bs=bs,
+                accelerator_graph=lease is not None,
+                null_slot_id=SINK_SLOT if self._has_sink else NO_SLOT,
+                kernels=self._kernels,
+            )
+        wrapper.plan(groups[: addressing.num_rows], verify.num_accepted)
+        return wrapper
+
     def _get_wrapper(
         self, label: str, is_decode: bool, lease: SlotLease | None = None,
         split: CombinedSplit | None = None,
diff --git a/mstar/engine/resources/linear_attn/gdn_kernels.py b/mstar/engine/resources/linear_attn/gdn_kernels.py
index 207f11f702..97c2c584e5 100644
--- a/mstar/engine/resources/linear_attn/gdn_kernels.py
+++ b/mstar/engine/resources/linear_attn/gdn_kernels.py
@@ -66,6 +66,24 @@ class GDNDecodePlan:
     num_rows: int
 
 
+@dataclass(frozen=True)
+class GDNVerifyPlan:
+    """A speculative verify step: every row is ``tokens_per_row`` tokens.
+
+    Row ``i`` starts from the state in ``slot_map[i, num_accepted[i] - 1]``
+    and leaves the state after its token ``t`` in ``slot_map[i, t]``; -1 rows
+    are skipped. ``num_accepted`` is filled on the device by the speculative
+    resource before the layers run.
+    """
+
+    slot_map: torch.Tensor      # [n, tokens_per_row] int32
+    conv_slots: torch.Tensor    # [n] int32, each row's conv window
+    cu_seqlens: torch.Tensor    # [n + 1] int32
+    num_accepted: torch.Tensor  # [n] int32
+    num_rows: int
+    tokens_per_row: int
+
+
 def _with_room(t: torch.Tensor, rows: int) -> torch.Tensor:
     """``t`` with ``rows`` zeroed rows appended along dim 0."""
     return torch.cat((t, t.new_zeros((rows, *t.shape[1:]))), dim=0)
@@ -179,6 +197,15 @@ def captured_prefill(self, bs: int, num_tokens: int) -> GDNPrefillKernel:
     def decode(self, bs: int | None) -> GDNDecodeKernel:
         """For a decode wrapper: a capture bucket's ``bs``, None for eager."""
 
+    def verify(self) -> GDNPrefillKernel:
+        """For a speculative verify step: rows of k + 1 tokens, each starting
+        from the state its last step accepted and keeping the state after
+        every token (``GDNVerifyWrapper``)."""
+        raise NotImplementedError(
+            f"{type(self).__name__} has no speculative verify kernel; the 'aiter' "
+            "GDN backend has one"
+        )
+
 
 # What each NVIDIA arch's chunked prefill kernel will read as its packed state.
 # SM90 is fp32 only; Blackwell takes the narrower ones, so a bf16 pool is not
diff --git a/mstar/engine/resources/linear_attn/gdn_kernels_aiter.py b/mstar/engine/resources/linear_attn/gdn_kernels_aiter.py
index adfbdf65ff..52124b8dde 100644
--- a/mstar/engine/resources/linear_attn/gdn_kernels_aiter.py
+++ b/mstar/engine/resources/linear_attn/gdn_kernels_aiter.py
@@ -122,6 +122,8 @@ def __init__(
         self._dt_bias: dict[int, tuple[torch.Tensor, int, torch.Tensor]] = {}
         # [0, 1, ..., n] for one-token rows, grown outside capture
         self._row_offsets: torch.Tensor | None = None
+        # verify outputs by shape and dtype; see `verify_output`
+        self._verify_out: dict[tuple, torch.Tensor] = {}
 
     @classmethod
     def build(
@@ -219,6 +221,11 @@ def decode(self, bs: int | None) -> GDNDecodeKernel:
             return FlyDSLDecode(self)
         return AiterSequentialDecode(self)
 
+    def verify(self) -> GDNPrefillKernel:
+        if self.flydsl_decode:
+            return FlyDSLVerify(self)
+        return AiterVerify(self)
+
     # Shared state
 
     def dt_bias_as(self, dt_bias: torch.Tensor, dtype: torch.dtype) -> torch.Tensor:
@@ -257,6 +264,30 @@ def row_offsets(self, rows: int, device: torch.device) -> torch.Tensor:
             )
         return self._row_offsets[: rows + 1]
 
+    def verify_output(
+        self, rows: int, tokens: int, dtype: torch.dtype, device: torch.device,
+    ) -> torch.Tensor:
+        """``[rows, tokens, V heads, V]`` for FlyDSL's verify kernel to write.
+
+        The kernel skips padding rows without writing them, so a zeroed buffer
+        kept across steps leaves them finite with no fill per layer. Layers run
+        one after another on one stream, so they all share it. Allocated
+        outside capture, like ``dt_bias_as``.
+        """
+        key = (rows, tokens, dtype, device)
+        out = self._verify_out.get(key)
+        if out is None:
+            if torch.cuda.is_current_stream_capturing():
+                raise RuntimeError(
+                    "gdn verify first saw this shape inside graph capture; "
+                    "capture must follow an eager warmup"
+                )
+            geometry = self.geometry
+            out = self._verify_out[key] = torch.zeros(
+                (rows, tokens, geometry.num_v_heads, geometry.head_v_dim), dtype=dtype, device=device,
+            )
+        return out
+
 
 class AiterChunkedPrefill(GDNPrefillKernel):
     """``chunk_gated_delta_rule_opt_vk`` with in-place pool slots, for eager
@@ -470,3 +501,71 @@ def run(self, plan, q, k, v, a, b, state, a_log, dt_bias):
             core_attn_out=out,
         )
         return out
+
+
+class FlyDSLVerify(GDNPrefillKernel):
+    """FlyDSL's delta rule over a linear draft chain, for a speculative verify
+    step: each row's k + 1 tokens from the state its last step accepted, the
+    state after every token kept in its own slot (see ``GDNVerifyPlan``).
+    q/k/v/a/b are read as the projection's column views, so unlike the
+    sequential kernel nothing is packed or copied per layer.
+
+    Compiled on first call per launch shape, in each bucket's eager warmup.
+    """
+
+    takes_strided_inputs = True
+
+    def __init__(self, kernels: AiterGDNKernels):
+        self._kernels = kernels
+        self.qk_l2norm_in_kernel = kernels.qk_l2norm
+
+    def run(self, plan, q, k, v, a, b, state, a_log, dt_bias):
+        from aiter.ops.flydsl.linear_attention_kernels import flydsl_gdr_mtp
+
+        chain = (plan.num_rows, plan.tokens_per_row)
+        out = self._kernels.verify_output(*chain, q.dtype, q.device)
+        flydsl_gdr_mtp(
+            q.unflatten(0, chain), k.unflatten(0, chain), v.unflatten(0, chain),
+            a.unflatten(0, chain), b.unflatten(0, chain),
+            self._kernels.dt_bias_as(dt_bias, q.dtype), a_log, state, out,
+            plan.slot_map, plan.num_accepted,
+            use_qk_l2norm=self._kernels.qk_l2norm,
+        )
+        return out.flatten(0, 1)
+
+
+class AiterVerify(GDNPrefillKernel):
+    """The sequential kernel over a speculative verify step: each row's k + 1
+    tokens from the state its last step accepted (``num_accepted_tokens``
+    picks the column of the row's slot map), the state after every token kept
+    in its own slot (see ``GDNVerifyPlan``)."""
+
+    takes_strided_inputs = True
+
+    def __init__(self, kernels: AiterGDNKernels):
+        self._kernels = kernels
+        self.qk_l2norm_in_kernel = kernels.qk_l2norm
+
+    def run(self, plan, q, k, v, a, b, state, a_log, dt_bias):
+        from aiter.ops.triton.gated_delta_net import (
+            fused_rearrange_sigmoid_gated_delta_rule,
+        )
+
+        n = q.shape[0]
+        qkv = torch.cat((q.reshape(n, -1), k.reshape(n, -1), v.reshape(n, -1)), dim=-1)
+        # skipped (padding) rows are never written; zeros keep them finite
+        out = torch.zeros((n, v.shape[1], v.shape[2]), dtype=q.dtype, device=q.device)
+        fused_rearrange_sigmoid_gated_delta_rule(
+            A_log=a_log, a=a, b=b, dt_bias=dt_bias, qkv=qkv,
+            key_dim=q.shape[1] * q.shape[2], value_dim=v.shape[1] * v.shape[2],
+            head_k_dim=q.shape[2], head_v_dim=v.shape[2],
+            scale=self._kernels.scale,
+            initial_state=state,
+            inplace_final_state=True,
+            cu_seqlens=plan.cu_seqlens,
+            ssm_state_indices=plan.slot_map,
+            num_accepted_tokens=plan.num_accepted,
+            use_qk_l2norm_in_kernel=self._kernels.qk_l2norm,
+            core_attn_out=out,
+        )
+        return out
diff --git a/mstar/engine/resources/linear_attn/wrappers.py b/mstar/engine/resources/linear_attn/wrappers.py
index 8bbff86de6..daf99d9db0 100644
--- a/mstar/engine/resources/linear_attn/wrappers.py
+++ b/mstar/engine/resources/linear_attn/wrappers.py
@@ -20,8 +20,11 @@
     GDNKernels,
     GDNPrefillKernel,
     GDNPrefillPlan,
+    GDNVerifyPlan,
 )
+from mstar.platforms import current_platform
 from mstar.platforms.compile import compiled
+from mstar.utils.h2d import H2DMirror, PinnedStager
 
 # Must match `BLOCK_M` in mstar/utils/causal_conv1d/kernels.py, which passes it
 # as an explicit constexpr rather than autotuning it.
@@ -447,8 +450,9 @@ def run_conv(
         bias: torch.Tensor | None = None,
         activation: str | None = "silu",
     ):
-        from mstar.utils.causal_conv1d import causal_conv1d_update
-        return causal_conv1d_update(
+        from mstar.utils.causal_conv1d import causal_conv1d_decode, causal_conv1d_update
+        conv = causal_conv1d_decode if current_platform.is_rocm() else causal_conv1d_update
+        return conv(
             x=x,
             conv_state=conv_layer,
             weight=weight,
@@ -717,3 +721,114 @@ def run(self, q, k, v, a, b, state, a_log, dt_bias):
             f = self._front
             out = self._prefill.run(q[f:], k[f:], v[f:], a[f:], b[f:], state, a_log, dt_bias)
         return out if rows is None else self._merge(out, rows)
+
+
+class GDNVerifyWrapper(GDNWrapper):
+    """A speculative verify step: every row ``tokens_per_row`` tokens (a
+    request's last emitted token and its k drafts).
+
+    A request holds one slot per token (the pool's slot group). Its first token
+    starts from the slot its previous step accepted up to, which only the
+    device knows (``num_accepted``, written by the speculative resource), and
+    the state after token ``t`` lands in slot ``t`` of the group, so whatever
+    the step accepts, the next one finds that state without a copy. The conv
+    window is one slot wide by k + width - 1 and slides by the accepted count
+    the same way. Padding rows (a group on the pool's null slot) are skipped
+    by both kernels.
+    """
+
+    def __init__(
+        self,
+        device: torch.device,
+        pad_slot_id: int,
+        tokens_per_row: int,
+        bs: int | None = None,
+        accelerator_graph: bool = False,
+        null_slot_id: int = 0,
+        kernels: GDNKernels | None = None,
+    ):
+        super().__init__(
+            device=device, pad_slot_id=pad_slot_id,
+            num_tokens=None if bs is None else bs * tokens_per_row, bs=bs,
+            accelerator_graph=accelerator_graph, null_slot_id=null_slot_id,
+        )
+        if kernels is None:
+            raise ValueError("GDNVerifyWrapper needs the resource's GDNKernels")
+        self._kernel: GDNPrefillKernel = kernels.verify()
+        self._tokens = tokens_per_row
+        self._stager = PinnedStager(torch.int32)
+        self._mirror = H2DMirror() if accelerator_graph else None
+        # [slot_map | conv_slots | cu_seqlens], one H2D a step, sized once to
+        # the bucket under capture
+        self._image: torch.Tensor | None = None
+        self._rows_capacity = 0
+        self._plan_state: GDNVerifyPlan | None = None
+
+    def _views(self, rows: int):
+        t = self._tokens
+        cap = self._rows_capacity
+        image = self._image
+        slot_map = image[: cap * t].view(cap, t)
+        conv = image[cap * t: cap * t + cap]
+        cu = image[cap * (t + 1): cap * (t + 1) + cap + 1]
+        return slot_map[:rows], conv[:rows], cu[: rows + 1]
+
+    def plan(self, groups, num_accepted: torch.Tensor):
+        rows, t = len(groups), self._tokens
+        need = max(rows, self._bs or 0) if self._accelerator_graph else rows
+        if need > self._rows_capacity:
+            if self._image is not None and self._accelerator_graph:
+                raise RuntimeError(
+                    f"GDNVerifyWrapper: {rows} rows past the {self._rows_capacity} "
+                    "its captured graph holds"
+                )
+            self._rows_capacity = max(need, 1)
+            self._image = torch.zeros(
+                self._rows_capacity * (t + 2) + 1, dtype=torch.int32, device=self._device,
+            )
+            if self._mirror is not None:
+                self._mirror.invalidate()
+        cap = self._rows_capacity
+        null = self._null_slot_id
+        host = np.full(cap * (t + 2) + 1, -1, dtype=np.int32)
+        slot_map = host[: cap * t].reshape(cap, t)
+        conv = host[cap * t: cap * t + cap]
+        conv[:] = null
+        for i, group in enumerate(groups):
+            if group[0] != null:
+                slot_map[i] = group
+                conv[i] = group[0]
+        host[cap * (t + 1): cap * (t + 1) + cap + 1] = np.arange(cap + 1, dtype=np.int32) * t
+        self._stager.copy_(self._image, host, mirror=self._mirror)
+        slot_map_d, conv_d, cu_d = self._views(rows)
+        self._plan_state = GDNVerifyPlan(
+            slot_map=slot_map_d, conv_slots=conv_d, cu_seqlens=cu_d,
+            num_accepted=num_accepted[:rows], num_rows=rows, tokens_per_row=t,
+        )
+
+    def run_conv(
+        self,
+        x: torch.Tensor,
+        conv_layer: torch.Tensor,
+        weight: torch.Tensor,
+        bias: torch.Tensor | None = None,
+        activation: str | None = "silu",
+    ):
+        from mstar.utils.causal_conv1d import causal_conv1d_update
+
+        plan = self._plan_state
+        return causal_conv1d_update(
+            x=x,
+            conv_state=conv_layer,
+            weight=weight,
+            bias=bias,
+            activation=activation,
+            conv_state_indices=plan.conv_slots,
+            num_accepted_tokens=plan.num_accepted,
+            query_start_loc=plan.cu_seqlens,
+            max_query_len=plan.tokens_per_row,
+            null_block_id=self._null_slot_id,
+        )
+
+    def run(self, q, k, v, a, b, state, a_log, dt_bias):
+        return self._kernel.run(self._plan_state, q, k, v, a, b, state, a_log, dt_bias)
diff --git a/mstar/engine/resources/recurrent/config.py b/mstar/engine/resources/recurrent/config.py
index 684e18e1eb..490bb47465 100644
--- a/mstar/engine/resources/recurrent/config.py
+++ b/mstar/engine/resources/recurrent/config.py
@@ -34,6 +34,11 @@ class RecurrentBlockConfig:
     dtype: torch.dtype | None
     shard_dims: tuple[int, ...] = ()
     per_layer: bool = True
+    # A sliding window over recent tokens along the last axis (a conv tail).
+    # Under speculative decoding a verify step appends k + 1 tokens and the
+    # next one starts after however many were accepted, so the pool keeps k
+    # more entries there (`RecurrentStateConfig.spec_tokens`).
+    speculative_window: bool = False
 
     def __post_init__(self):
         self.shape = tuple(self.shape)
@@ -52,6 +57,13 @@ def shard(self, num_shards: int) -> None:
             shape[dim] = divide(self._unsharded_shape[dim], num_shards)
         self.shape = tuple(shape)
 
+    def storage_shape(self, spec_tokens: int = 0) -> tuple[int, ...]:
+        """The shape a slot stores: ``shape``, with a speculative window's
+        last axis ``spec_tokens`` longer."""
+        if not (self.speculative_window and spec_tokens):
+            return self.shape
+        return (*self.shape[:-1], self.shape[-1] + spec_tokens)
+
 
 class RecurrentGeometry(ABC):
     @abstractmethod
@@ -115,6 +127,7 @@ def to_blocks(
                 shape=(self.conv_dim, self.conv_kernel_size - 1),
                 dtype=conv_dtype,
                 shard_dims=(0,),
+                speculative_window=True,
             ),
         }
         if speculative_tokens > 0:
@@ -303,6 +316,12 @@ class RecurrentStateConfig:
     prefix_slots: int = 0
     kv_cache: str | None = None
 
+    # Drafts per speculative verify step (k), 0 without speculation. A request
+    # then holds k + 1 slots per label, one per token of a verify row, so the
+    # step can keep the state after each token and the next one resume from the
+    # last accepted; speculative windows grow by k (`RecurrentBlockConfig`).
+    spec_tokens: int = 0
+
     def __post_init__(self):
         if not self.blocks:
             raise ValueError("a recurrent state pool must declare a block")
@@ -319,6 +338,12 @@ def check_slots(self) -> None:
                 f"label(s)" + ("." if self.disable_sink_slot else
                                ": the sink takes one. Raise it or set disable_sink_slot.")
             )
+        if self.spec_tokens < 0:
+            raise ValueError(f"spec_tokens={self.spec_tokens} must be >= 0")
+
+    @property
+    def slots_per_request(self) -> int:
+        return self.spec_tokens + 1
 
     @property
     def usable_slots(self) -> int:
diff --git a/mstar/engine/resources/recurrent/pool.py b/mstar/engine/resources/recurrent/pool.py
index 91eb1d0308..d1792332fc 100644
--- a/mstar/engine/resources/recurrent/pool.py
+++ b/mstar/engine/resources/recurrent/pool.py
@@ -79,6 +79,13 @@ class SlotState:
     generation: int = 0
     # tokens the state has taken in, which places its page boundaries
     tokens: int = 0
+    # under speculative decoding, the slots after `index` that hold the state
+    # after each further token of a verify row
+    extra: tuple[int, ...] = ()
+
+    @property
+    def group(self) -> tuple[int, ...]:
+        return (self.index, *self.extra)
 
 
 class SnapshotLease(NamedTuple):
@@ -137,6 +144,9 @@ class RecurrentAddressing:
     snapshot_every: int = 0
     # scatter_'s int64 copy of slot_indices, keyed by whether a graph capture made it
     scatter_index: dict[bool, torch.Tensor] = field(default_factory=dict, compare=False, repr=False)
+    # Under speculative decoding, each row's slots in token order (its
+    # `slot_indices` entry first), `pad_index` throughout for a row with none
+    host_slot_groups: Sequence[tuple[int, ...]] | None = None
 
 
 class RecurrentStatePool(Resource):
@@ -184,7 +194,7 @@ def __init__(
         self._blocks: dict[str, torch.Tensor] = {
             name: torch.zeros(
                 (*((config.num_layers,) if block.per_layer else ()),
-                 config.max_slots + config.prefix_slots, *block.shape),
+                 config.max_slots + config.prefix_slots, *block.storage_shape(config.spec_tokens)),
                 dtype=block.resolved_dtype(kv_dtype),
                 device=device,
             )
@@ -212,7 +222,8 @@ def __init__(
         # Admission, as `KVManager` decides it: only at one rank, since a
         # reservation held on rank 0 alone would make the ranks refuse different
         # allocations. `_reserved` holds, per admitted request, the slots kept
-        # for labels it hasn't taken yet, and `_kept` their sum.
+        # for labels it hasn't taken yet (`slots_per_request` each), and `_kept`
+        # their sum.
         self._decides = world_size == 1
         self._reserved: dict[str, int] = {}
         self._kept = 0
@@ -405,9 +416,10 @@ def reset_request(self, rid: str, free: bool = False):
             if not labels:
                 return
             for slot in labels.values():
-                self._zero_slot(slot.index)
-                if free:
-                    self._free.append(slot.index)
+                for index in slot.group:
+                    self._zero_slot(index)
+                    if free:
+                        self._free.append(index)
             if free:
                 labels.clear()
             else:
@@ -441,19 +453,22 @@ def _alloc(self, rid: str, label: str) -> SlotState | None:
             if slot is not None:
                 return slot
             kept = self._reserved.get(rid, 0)
-            if label not in self.config.labels or len(self._free) - self._kept + bool(kept) <= 0:
+            need = self.config.slots_per_request
+            if label not in self.config.labels or len(self._free) - self._kept + (need if kept else 0) < need:
                 # an undeclared label, or every free slot is kept for an admitted request
                 return None
             if kept:
-                self._kept -= 1
-                if kept == 1:
+                self._kept -= need
+                if kept == need:
                     del self._reserved[rid]
                 else:
-                    self._reserved[rid] = kept - 1
+                    self._reserved[rid] = kept - need
             # Slot ids must be unique across the batch: kernels assume it rather
             # than check, and a scatter with two rows on one slot is
             # nondeterministic. The free list guarantees it.
-            slot = labels[label] = SlotState(index=self._free.pop())
+            index = self._free.pop()
+            extra = tuple(self._free.pop() for _ in range(need - 1))
+            slot = labels[label] = SlotState(index=index, extra=extra)
             self._admit_taken[(rid, label)] = (slot.index, bool(kept))
             self._version += 1
             return slot
@@ -478,7 +493,8 @@ def ask(self, rid: str, node_name: str, graph_walk: str) -> AdmitOutcome | None:
 
     def reserve(self, rid: str, node_name: str, graph_walk: str, published: None) -> int | None:
         with self._lock:
-            need, room = len(self.config.labels), len(self._free) - self._kept
+            need = len(self.config.labels) * self.config.slots_per_request
+            room = len(self._free) - self._kept
             if need > room:
                 return room
             self._reserved[rid] = need
@@ -570,11 +586,13 @@ def _release(self, keys: list[tuple[str, str]]) -> None:
             if slot is None or slot.index != index:
                 continue  # reset or removed under us; the slot went with it
             del labels[label]
-            self._zero_slot(index)
-            self._free.append(index)
+            for taken_index in slot.group:
+                self._zero_slot(taken_index)
+                self._free.append(taken_index)
             if reserved:
-                self._reserved[rid] = self._reserved.get(rid, 0) + 1
-                self._kept += 1
+                need = self.config.slots_per_request
+                self._reserved[rid] = self._reserved.get(rid, 0) + need
+                self._kept += need
             self._version += 1
 
     def _refused(self, rid: str, label: str) -> AdmitOutcome:
@@ -1098,10 +1116,13 @@ def _build_addressing(
             )
             memo = self._addressing_memo.get(label)
             if memo is not None and memo[0] == memo_key:
-                return self._stage_addressing(label, segments, ctx, memo[1], memo[2])
+                return self._stage_addressing(label, segments, ctx, memo[1], memo[2], memo[3])
         pending_forks = pending_forks or {}
         indices = []
         has_state = []
+        spec = self.config.spec_tokens > 0
+        groups = [] if spec else None
+        pad_group = (pad,) * self.config.slots_per_request
         for seg in segments:
             slot = (
                 None if ctx.capture or ctx.is_padding_row(seg.request_id)
@@ -1112,18 +1133,22 @@ def _build_addressing(
                 # sink (or the sentinel), which nothing reads back
                 indices.append(pad)
                 has_state.append(False)
+                if spec:
+                    groups.append(pad_group)
             else:
                 indices.append(slot.index)
                 has_state.append(
                     pending_forks.get((seg.request_id, label), slot.has_state)
                 )
+                if spec:
+                    groups.append(slot.group)
         if memo_key is not None:
-            self._addressing_memo[label] = (memo_key, indices, has_state)
-        return self._stage_addressing(label, segments, ctx, indices, has_state)
+            self._addressing_memo[label] = (memo_key, indices, has_state, groups)
+        return self._stage_addressing(label, segments, ctx, indices, has_state, groups)
 
     def _stage_addressing(
         self, label: str, segments: list[Segment], ctx: StepContext,
-        indices: list[int], has_state: list[bool],
+        indices: list[int], has_state: list[bool], groups: list[tuple[int, ...]] | None = None,
     ) -> RecurrentAddressing:
         pad = self.pad_index
         snapshots = ()
@@ -1163,6 +1188,7 @@ def _stage_addressing(
                 self._checkpoint_every(segments, ctx) if snapshots
                 else self._kv.config.tokens_per_page if self._kv is not None else 0
             ),
+            host_slot_groups=groups,
         )
 
     def _addressing_buffers(
diff --git a/mstar/engine/resources/sampler/utils.py b/mstar/engine/resources/sampler/utils.py
index 2072cf8ab5..76fdb1296c 100644
--- a/mstar/engine/resources/sampler/utils.py
+++ b/mstar/engine/resources/sampler/utils.py
@@ -1622,6 +1622,25 @@ def _sample_cuda(
         return Draw(result, probs) if return_probs else result
 
 
+def draw_tokens(
+    logits: torch.Tensor,
+    temperature: torch.Tensor,
+    top_k: torch.Tensor,
+    top_p: torch.Tensor,
+    seed: torch.Tensor,
+    offset: torch.Tensor,
+) -> torch.Tensor:
+    """One token per row of ``logits`` from its top-k, then top-p, distribution
+    at its temperature (argmax where the temperature is 0), drawn with per-row
+    device ``seed``/``offset`` so a captured graph draws afresh on each replay
+    with the offsets it is handed. For a caller that samples rows of its own
+    (a speculative verify step); the sampler resource's per-request state is
+    not involved. Accelerator devices only."""
+    return _sample_cuda(
+        logits, temperature, top_k, top_p, 1.0, None, True, None, seed, offset,
+    )
+
+
 def _needs_explicit_filter(history: HistoryRows | None, order: FilterOrder | None) -> bool:
     """Whether the draw needs the filtered distribution in hand: FlashInfer's
     fused sampler has neither HF's order nor a floor applied after filtering."""
diff --git a/mstar/engine/resources/spec_decode/__init__.py b/mstar/engine/resources/spec_decode/__init__.py
new file mode 100644
index 0000000000..c30b860475
--- /dev/null
+++ b/mstar/engine/resources/spec_decode/__init__.py
@@ -0,0 +1,20 @@
+"""Speculative decoding: what a model declares, kept import-light; the resource
+is reached by its own path (``spec_decode.resource``)."""
+
+from mstar.engine.resources.spec_decode.config import (
+    AcceptedTokens,
+    DraftKind,
+    SpecDecodeConfig,
+    SpecDecodeSpec,
+    SpecDecodeStep,
+    SpecPhase,
+)
+
+__all__ = [
+    "AcceptedTokens",
+    "DraftKind",
+    "SpecDecodeConfig",
+    "SpecDecodeSpec",
+    "SpecDecodeStep",
+    "SpecPhase",
+]
diff --git a/mstar/engine/resources/spec_decode/config.py b/mstar/engine/resources/spec_decode/config.py
new file mode 100644
index 0000000000..55e9683e17
--- /dev/null
+++ b/mstar/engine/resources/spec_decode/config.py
@@ -0,0 +1,132 @@
+"""What a model declares to decode speculatively: the spec, and its step.
+
+Kept free of the resource so a submodule can declare a step without importing
+the device-side bookkeeping.
+"""
+
+from __future__ import annotations
+
+from collections.abc import Mapping
+from dataclasses import dataclass, field
+from enum import Enum
+from typing import TYPE_CHECKING, NamedTuple
+
+import torch
+
+from mstar.engine.resources.spec import NodeResourceSpec
+from mstar.engine.resources.step import ResourceStep
+
+if TYPE_CHECKING:
+    from mstar.engine.resources.base import Resource
+
+
+class DraftKind(Enum):
+    # the checkpoint's own multi-token-prediction head, chained k times
+    MTP = "mtp"
+
+
+@dataclass
+class SpecDecodeConfig:
+    # drafts per verify step (k): a verify row carries k + 1 tokens
+    num_draft_tokens: int
+    draft: DraftKind = DraftKind.MTP
+    # requests this node can hold at once, each with one row of device registers
+    max_requests: int = 256
+    # Whether verify steps can sample the target (temperature, top-k, top-p),
+    # not only take its argmax. A capability: the model captures a sampled
+    # decode variant beside the greedy one and routes a batch by its rows.
+    sampling: bool = False
+
+    def __post_init__(self):
+        if self.num_draft_tokens < 1:
+            raise ValueError(
+                f"num_draft_tokens={self.num_draft_tokens}: a speculative node "
+                "drafts at least one token; leave the resource undeclared to "
+                "decode one token a step"
+            )
+        if self.max_requests < 1:
+            raise ValueError(f"max_requests={self.max_requests} must be >= 1")
+
+    @property
+    def tokens_per_row(self) -> int:
+        return self.num_draft_tokens + 1
+
+
+@dataclass
+class SpecDecodeSpec(NodeResourceSpec):
+    config: SpecDecodeConfig
+    # the paged KV the verify rows write and the attention resource that reads
+    # it; this resource rewrites their device-side addressing per pass
+    kv_cache: str = "kv_cache"
+    attn: str = "attn"
+    # ids that end generation, masked in-graph so a step never emits past one
+    stop_token_ids: tuple[int, ...] = ()
+
+    def depends_on(self) -> set[str]:
+        return {self.kv_cache, self.attn}
+
+    @property
+    def resource_class(self) -> "type[Resource]":
+        from mstar.engine.resources.spec_decode.resource import SpecDecodeResource
+
+        return SpecDecodeResource
+
+    def apply_yaml_overrides(self, max_requests: int | None = None):
+        """The register count is a deployment's concurrency. ``k`` is not here:
+        the KV layers, recurrent slots and capture buckets are sized from it at
+        declaration, so it is a model kwarg."""
+        if max_requests is not None:
+            self.config.max_requests = max_requests
+
+
+class SpecPhase(Enum):
+    # prefill rows: seed each request's registers from the prompt's last token
+    SEED = "seed"
+    # decode rows: verify k drafts, accept a prefix, draft the next k
+    VERIFY = "verify"
+
+
+@dataclass(frozen=True)
+class SpecDecodeStep(ResourceStep):
+    phase: SpecPhase = SpecPhase.VERIFY
+    # rid -> tokens the host over-counted in this request's KV stream and has
+    # since learned it can give back; KV drops them in the same commit
+    # (`KVStep.commit_adjust`), and this resource forgets them then
+    retract: Mapping[int, int] = field(default_factory=dict)
+    # seed rows: rid -> decode tokens the request may emit (`max_tokens`), and
+    # the requests that run past stop ids (`ignore_eos`)
+    budgets: Mapping[int, int] = field(default_factory=dict)
+    ignore_stop: frozenset[int] = frozenset()
+    # seed rows of a sampling node: rid -> (temperature, top_p, top_k, seed); a row
+    # left out (or at temperature 0) verifies greedily
+    sampling: Mapping[int, tuple[float, float, int]] = field(default_factory=dict)
+
+
+class DraftAddress(NamedTuple):
+    """Where a draft pass writes and reads, on the device: its tokens' flat KV
+    slots and rotary positions (all k + 1 slots of a row on the first pass,
+    one a row on the others) and the keys each row's one real query reads."""
+
+    slots: torch.Tensor      # [tokens] int64: page * page_size + offset
+    positions: torch.Tensor  # [tokens] int64
+    kv_lens: torch.Tensor    # [rows] int32
+
+
+class AcceptedTokens(NamedTuple):
+    """What a verify step hands back, one row per request. A NamedTuple, so a
+    compiled forward takes it back from the (uncompiled) resource call as a
+    plain tuple of tensors.
+
+    ``tokens`` holds the tokens the step emits in order, -1 past the last one;
+    ``accepted`` how many of the row's k + 1 positions the target took (the
+    stream's real growth, which the host needs to give back the rest of its
+    reservation); ``done`` whether the row emitted a stop id or reached its
+    budget, so the host can end the loop without reading the tokens.
+    """
+
+    tokens: torch.Tensor    # [rows, k + 1] int64
+    accepted: torch.Tensor  # [rows] int32
+    done: torch.Tensor      # [rows] bool
+    # the target's own token at each verify position, unmasked: the draft's
+    # first pass reads them
+    target: torch.Tensor    # [rows, k + 1] int64
diff --git a/mstar/engine/resources/spec_decode/draft_head.py b/mstar/engine/resources/spec_decode/draft_head.py
new file mode 100644
index 0000000000..739dce9ee4
--- /dev/null
+++ b/mstar/engine/resources/spec_decode/draft_head.py
@@ -0,0 +1,122 @@
+"""A draft's vocabulary head in fp8.
+
+A draft pass reads the head for one argmax, and the target verifies the result,
+so rounding the weights to 8 bits costs a little acceptance and no correctness.
+The head is the largest kernel of a draft pass and all of it is the weight
+stream, so one byte a weight in place of two halves that kernel.
+
+``Fp8DraftHead`` keeps a copy of the head at one scale. Its logits are only
+good for an argmax over a row: each row carries its own scale.
+
+``mstar::fp8_head`` quantises each row of ``x`` and reads the head in fp8. At one
+or two rows AITER's split-K kernel reads the weight once across every CU, 1.7-1.9x
+the bf16 kernel on a 248k x 4K head on MI325X; from three rows on that kernel is
+an order of magnitude slower and hipBLASLt's fp8 GEMM takes over, 1.2-1.6x.
+"""
+
+from __future__ import annotations
+
+import functools
+
+import torch
+from torch import nn
+
+from mstar.kernels.fp8_gemm import quantize_rows
+from mstar.kernels.numerics import Tolerance
+from mstar.kernels.ops import KernelOp, ShapePolicy
+from mstar.platforms import current_platform
+
+# where AITER's fp8 split-K kernel beats hipBLASLt's
+SKINNY_MAX_ROWS = 2
+
+
+@functools.cache
+def _cu_count(device_index: int) -> int:
+    return torch.cuda.get_device_properties(device_index).multi_processor_count
+
+
+@functools.cache
+def _unit_scales(shape: tuple[int, ...], device: torch.device) -> torch.Tensor:
+    """Scales of 1, for GEMMs that take them. Cached, so the first call of a
+    shape, an eager warmup's, allocates them outside any captured graph."""
+    return torch.ones(shape, dtype=torch.float32, device=device)
+
+
+def _aiter_impl():
+    from aiter.ops.custom import wvSplitKQ
+    from aiter.ops.triton.quant import dynamic_per_token_quant_fp8_i8
+
+    def fp8_head(x: torch.Tensor, weight: torch.Tensor) -> torch.Tensor:
+        rows, vocab, device = x.shape[0], weight.shape[0], x.device
+        xq = torch.empty(x.shape, dtype=weight.dtype, device=device)
+        dynamic_per_token_quant_fp8_i8(xq, x, torch.empty(rows, dtype=torch.float32, device=device))
+        if rows > SKINNY_MAX_ROWS:
+            return torch._scaled_mm(
+                xq, weight.T, scale_a=_unit_scales((rows, 1), device),
+                scale_b=_unit_scales((1, vocab), device), out_dtype=x.dtype,
+            )
+        out = torch.empty((rows, vocab), dtype=x.dtype, device=device)
+        one = _unit_scales((1,), device)
+        wvSplitKQ(weight, xq, out, one, one, _cu_count(device.index or 0))
+        return out
+
+    return fp8_head
+
+
+def _reference(x, weight):
+    compute = torch.float64 if x.dtype == torch.float64 else torch.float32
+    xq = quantize_rows(x, current_platform.fp8_dtype() or torch.float8_e4m3fn)[0]
+    return (xq.to(compute) @ weight.to(compute).T).to(x.dtype)
+
+
+def _fake(x, weight):
+    return x.new_empty((x.shape[0], weight.shape[0]))
+
+
+def _shape(x, weight) -> dict[str, int]:
+    return {"rows": x.shape[0], "n": weight.shape[0], "k": x.shape[1]}
+
+
+def _inputs(shape, dtype, device, generator=None):
+    fp8 = current_platform.fp8_dtype() or torch.float8_e4m3fn
+    x = torch.randn(shape["rows"], shape["k"], device=device, generator=generator).to(dtype)
+    weight = torch.randn(shape["n"], shape["k"], device=device, generator=generator) * shape["k"] ** -0.5
+    return x, quantize_rows(weight, fp8)[0]
+
+
+FP8_HEAD = KernelOp(
+    "fp8_head",
+    "(Tensor x, Tensor weight) -> Tensor",
+    fake=_fake,
+    reference=_reference,
+    impls={f"rocm:{arch}": _aiter_impl for arch in ("gfx942", "gfx950")},
+    shapes=ShapePolicy(
+        dims=("rows", "n", "k"),
+        key=_shape,
+        inputs=_inputs,
+        warm=({"rows": 1, "n": 2048, "k": 1024}, {"rows": 2, "n": 2048, "k": 1024},
+              {"rows": 8, "n": 2048, "k": 1024}),
+    ),
+    # one rounding of a long fp32 sum of exact fp8 products, plus summation order
+    tolerance=Tolerance.accumulated(),
+)
+
+
+class Fp8DraftHead(nn.Module):
+    """``weight`` ``[vocab, hidden]`` rounded to fp8, for argmaxes over its rows."""
+
+    def __init__(self, weight: torch.Tensor):
+        super().__init__()
+        fp8 = current_platform.fp8_dtype()
+        if fp8 is None:
+            raise ValueError(f"{current_platform.name} has no fp8 GEMM for a draft head")
+        scale = weight.float().abs().amax() / torch.finfo(fp8).max
+        self.register_buffer("weight", (weight.float() / scale).to(fp8), persistent=False)
+
+    def _apply(self, fn, recurse=True):
+        # the engine moves a whole submodule to its dtype, which would widen the copy
+        fp8 = self.weight.dtype
+        return super()._apply(lambda t: fn(t).to(fp8), recurse)
+
+    def forward(self, x: torch.Tensor) -> torch.Tensor:
+        return FP8_HEAD(x, self.weight)
diff --git a/mstar/engine/resources/spec_decode/kernels.py b/mstar/engine/resources/spec_decode/kernels.py
new file mode 100644
index 0000000000..aa57615fb0
--- /dev/null
+++ b/mstar/engine/resources/spec_decode/kernels.py
@@ -0,0 +1,300 @@
+"""A speculative step's bookkeeping as one kernel per phase.
+
+What a verify step does around its layers is integer work on a few rows: read
+the registers, address the step's writes, compare drafts, cut stop ids, address
+the draft passes. Spelled in torch it is 60-odd small kernels at c=1, each paying
+the launch floor of a replayed graph, so a Triton program per row does each
+phase in one launch.
+
+``mstar::spec_verify_inputs`` runs before the layers, ``mstar::spec_accept``
+after the target's tokens are known. The torch references are the semantics,
+and serve CPU tensors and devices without an implementation.
+"""
+
+from __future__ import annotations
+
+import torch
+import triton
+import triton.language as tl
+
+from mstar.kernels.numerics import Tolerance
+from mstar.kernels.ops import KernelOp, ShapePolicy
+
+
+@triton.jit
+def _verify_inputs_kernel(
+    rows_ptr, real_ptr, seq_len_ptr, num_acc_ptr, tokens_ptr, block_table_ptr,
+    ids_ptr, positions_ptr, slots_ptr, pages_ptr, offsets_ptr, kv_lens_ptr, num_accepted_ptr,
+    registers, table_stride, page_size, T,
+    BLOCK: tl.constexpr,
+):
+    r = tl.program_id(0)
+    row = tl.load(rows_ptr + r)
+    real = tl.load(real_ptr + r) != 0
+    base = tl.where(real, tl.load(seq_len_ptr + row), 0)
+    tl.store(num_accepted_ptr + r, tl.where(real, tl.load(num_acc_ptr + row), 1))
+    tl.store(kv_lens_ptr + r, base + T)
+
+    slot = tl.arange(0, BLOCK)
+    mask = slot < T
+    position = base + slot
+    entry = tl.load(block_table_ptr + r * table_stride + position // page_size, mask=mask, other=0)
+    offset = position % page_size
+    at = r * T + slot
+    tl.store(ids_ptr + at, tl.load(tokens_ptr + slot * registers + row, mask=mask, other=0), mask=mask)
+    tl.store(positions_ptr + at, position, mask=mask)
+    tl.store(pages_ptr + at, entry, mask=mask)
+    tl.store(offsets_ptr + at, offset, mask=mask)
+    tl.store(slots_ptr + at, entry.to(tl.int64) * page_size + offset, mask=mask)
+
+
+def _verify_inputs(
+    rows, real, seq_len, num_acc, tokens, block_table, page_size, ids, positions, slots,
+    write_pages, write_offsets, kv_lens, num_accepted,
+) -> None:
+    b, t = rows.shape[0], tokens.shape[0]
+    _verify_inputs_kernel[(b,)](
+        rows, real, seq_len, num_acc, tokens, block_table,
+        ids, positions, slots, write_pages, write_offsets, kv_lens, num_accepted,
+        tokens.shape[1], block_table.stride(0), page_size, t,
+        BLOCK=triton.next_power_of_2(t), num_warps=1,
+    )
+
+
+def _verify_inputs_reference(
+    rows, real, seq_len, num_acc, tokens, block_table, page_size, ids, positions, slots,
+    write_pages, write_offsets, kv_lens, num_accepted,
+) -> None:
+    b, t = rows.shape[0], tokens.shape[0]
+    rows_l, real_b = rows.long(), real.bool()
+    base = torch.where(real_b, seq_len[rows_l], 0)
+    num_accepted.copy_(torch.where(real_b, num_acc[rows_l], 1))
+    kv_lens.copy_(base + t)
+    position = base[:, None] + torch.arange(t, dtype=base.dtype, device=base.device)[None, :]
+    page = torch.div(position, page_size, rounding_mode="floor")
+    entry = block_table[:b].gather(1, page.long()).long()
+    offset = (position - page * page_size).long()
+    ids.copy_(tokens[:, rows_l].T.reshape(-1))
+    positions.copy_(position.reshape(-1).long())
+    write_pages.copy_(entry.reshape(-1))
+    write_offsets.copy_(offset.reshape(-1))
+    slots.copy_((entry * page_size + offset).reshape(-1))
+
+
+def _verify_inputs_fake(*args) -> None:
+    return None
+
+
+def _verify_inputs_shape(rows, real, seq_len, num_acc, tokens, *rest):
+    return {"rows": rows.shape[0], "tokens": tokens.shape[0]}
+
+
+def _verify_inputs_args(shape, dtype, device, generator=None):
+    b, t = shape["rows"], shape["tokens"]
+    registers, page_size, width = b + 1, 8, 32
+
+    def i32(values):
+        return values.to(torch.int32).to(device)
+
+    rows = i32(torch.arange(b, 0, -1))
+    real = i32(torch.ones(b).long())
+    # a padding row: it reads the sink row's registers as a fresh request
+    real[-1] = 0
+    seq_len = i32(torch.arange(registers) * 11 + 3)
+    num_acc = i32(torch.arange(registers) % t + 1)
+    tokens = (torch.arange(t * registers).view(t, registers) * 7).to(device)
+    block_table = i32(1 + torch.arange(b * width).view(b, width))
+    long = dict(dtype=torch.long, device=device)
+    return (rows, real, seq_len, num_acc, tokens, block_table, page_size,
+            torch.zeros(b * t, **long), torch.zeros(b * t, **long), torch.zeros(b * t, **long),
+            torch.zeros(b * t, **long), torch.zeros(b * t, **long),
+            torch.zeros(b, dtype=torch.int32, device=device), torch.zeros(b, dtype=torch.int32, device=device))
+
+
+SPEC_VERIFY_INPUTS = KernelOp(
+    "spec_verify_inputs",
+    "(Tensor rows, Tensor real, Tensor seq_len, Tensor num_acc, Tensor tokens, Tensor block_table, int page_size, "
+    "Tensor(a!) ids, Tensor(b!) positions, Tensor(c!) slots, Tensor(d!) write_pages, Tensor(e!) write_offsets, "
+    "Tensor(f!) kv_lens, Tensor(g!) num_accepted) -> ()",
+    fake=_verify_inputs_fake,
+    reference=_verify_inputs_reference,
+    impls={"rocm": lambda: _verify_inputs},
+    fallback_to_reference=True,
+    shapes=ShapePolicy(
+        dims=("rows", "tokens"),
+        key=_verify_inputs_shape,
+        inputs=_verify_inputs_args,
+        warm=({"rows": 4, "tokens": 3},),
+        dtypes=(torch.int32,),
+    ),
+    tolerance=Tolerance.exact(),
+)
+
+
+@triton.jit
+def _accept_kernel(
+    rows_ptr, ids_ptr, targets_ptr, positions_ptr, block_table_ptr, stop_ids_ptr,
+    seq_len_ptr, num_acc_ptr, tokens_ptr, emitted_ptr, budget_ptr, stop_on_ptr,
+    out_tokens_ptr, accepted_ptr, done_ptr, keep_ptr,
+    draft_slots_ptr, draft_positions_ptr, draft_lens_ptr,
+    table_stride, page_size, num_stops, num_rows, T, K,
+    BLOCK: tl.constexpr, STOPS: tl.constexpr, PASSES: tl.constexpr,
+):
+    r = tl.program_id(0)
+    row = tl.load(rows_ptr + r)
+    slot = tl.arange(0, BLOCK)
+    mask = slot < T
+    target = tl.load(targets_ptr + r * T + slot, mask=mask, other=-1)
+    drafted = tl.load(ids_ptr + r * T + slot + 1, mask=slot + 1 < T, other=-2)
+    # the first slot whose next draft the target refused ends the step: never
+    # one past the last, which has no draft after it
+    last = tl.argmin((drafted == target).to(tl.int32), axis=0)
+    accepted = last + 1
+    base = tl.load(positions_ptr + r * T).to(tl.int32)
+
+    stop = tl.arange(0, STOPS)
+    stop_ids = tl.load(stop_ids_ptr + stop, mask=stop < num_stops, other=-2)
+    is_stop = tl.max((target[:, None] == stop_ids[None, :]).to(tl.int32), axis=1) > 0
+    emit = slot <= last
+    stops = is_stop & emit & (tl.load(stop_on_ptr + row) > 0)
+    any_stop = tl.max(stops.to(tl.int32), axis=0) > 0
+    first_stop = tl.where(any_stop, tl.argmax(stops.to(tl.int32), axis=0), T)
+    emitted = tl.load(emitted_ptr + row)
+    budget = tl.load(budget_ptr + row)
+    emit = emit & (slot <= first_stop) & (slot < budget - emitted)
+    emitted = emitted + tl.sum(emit.to(tl.int32), axis=0)
+
+    tl.store(out_tokens_ptr + r * T + slot, tl.where(emit, target, -1), mask=mask)
+    tl.store(accepted_ptr + r, accepted)
+    tl.store(done_ptr + r, any_stop | (emitted >= budget))
+    tl.store(keep_ptr + r, r * T + last)
+    tl.store(seq_len_ptr + row, base + accepted)
+    tl.store(num_acc_ptr + row, accepted)
+    tl.store(emitted_ptr + row, emitted)
+    tl.store(tokens_ptr + row, tl.load(targets_ptr + r * T + last))
+
+    # draft pass j's one real query: the last accepted slot, j positions on
+    j = tl.arange(0, PASSES)
+    in_pass = j < K
+    position = base + last + j
+    entry = tl.load(block_table_ptr + r * table_stride + position // page_size, mask=in_pass, other=0)
+    at = j * num_rows + r
+    tl.store(draft_slots_ptr + at, entry.to(tl.int64) * page_size + position % page_size, mask=in_pass)
+    tl.store(draft_positions_ptr + at, position, mask=in_pass)
+    tl.store(draft_lens_ptr + at, position + 1, mask=in_pass)
+
+
+def _accept(
+    rows, ids, targets, positions, block_table, page_size, stop_ids,
+    seq_len, num_acc, tokens, emitted, budget, stop_on,
+    out_tokens, accepted, done, keep, draft_slots, draft_positions, draft_lens,
+) -> None:
+    b, k = rows.shape[0], draft_slots.shape[0]
+    t = out_tokens.shape[1]
+    _accept_kernel[(b,)](
+        rows, ids, targets, positions, block_table, stop_ids,
+        seq_len, num_acc, tokens, emitted, budget, stop_on,
+        out_tokens, accepted, done, keep, draft_slots, draft_positions, draft_lens,
+        block_table.stride(0), page_size, stop_ids.shape[0], b, t, k,
+        BLOCK=triton.next_power_of_2(t), STOPS=triton.next_power_of_2(stop_ids.shape[0]),
+        PASSES=triton.next_power_of_2(k), num_warps=1,
+    )
+
+
+def _accept_reference(
+    rows, ids, targets, positions, block_table, page_size, stop_ids,
+    seq_len, num_acc, tokens, emitted, budget, stop_on,
+    out_tokens, accepted, done, keep, draft_slots, draft_positions, draft_lens,
+) -> None:
+    b, t, k = rows.shape[0], out_tokens.shape[1], draft_slots.shape[0]
+    rows_l = rows.long()
+    y = targets.view(b, t).long()
+    x = ids.view(b, t)
+    base = positions.view(b, t)[:, 0].int()
+    count = (x[:, 1:] == y[:, :-1]).int().cumprod(dim=1).sum(dim=1) + 1
+    last = (count - 1)[:, None].long()
+    accepted.copy_(count.int())
+    seq_len.index_copy_(0, rows_l, base + count.int())
+    num_acc.index_copy_(0, rows_l, count.int())
+    tokens[0].index_copy_(0, rows_l, y.gather(1, last)[:, 0])
+
+    slots = torch.arange(t, dtype=torch.int32, device=rows.device)[None, :]
+    emit = slots < count[:, None]
+    stops = torch.isin(y, stop_ids) & emit & (stop_on[rows_l] > 0)[:, None]
+    hit = stops.any(dim=1)
+    first_stop = torch.where(hit, stops.int().argmax(dim=1), t)
+    emit &= slots <= first_stop[:, None]
+    so_far = emitted[rows_l]
+    emit &= slots < (budget[rows_l] - so_far)[:, None]
+    so_far = so_far + emit.sum(dim=1, dtype=torch.int32)
+    emitted.index_copy_(0, rows_l, so_far)
+    out_tokens.copy_(torch.where(emit, y, -1))
+    done.copy_(hit | (so_far >= budget[rows_l]))
+    keep.copy_(torch.arange(b, device=rows.device) * t + last[:, 0])
+
+    shift = torch.arange(k, dtype=torch.int32, device=rows.device)[:, None]
+    position = (base + count.int() - 1)[None, :] + shift
+    page = torch.div(position, page_size, rounding_mode="floor")
+    entry = block_table[:b].T.gather(0, page.long()).long()
+    draft_slots.copy_(entry * page_size + (position - page * page_size).long())
+    draft_positions.copy_(position.long())
+    draft_lens.copy_(position + 1)
+
+
+def _accept_fake(*args) -> None:
+    return None
+
+
+def _accept_shape(rows, *rest):
+    return {"rows": rows.shape[0], "tokens": rest[-7].shape[1]}
+
+
+def _accept_args(shape, dtype, device, generator=None):
+    b, t = shape["rows"], shape["tokens"]
+    k = t - 1
+    registers, page_size, width = b + 1, 8, 32
+
+    def i32(values):
+        return values.to(torch.int32).to(device)
+
+    long = dict(dtype=torch.long, device=device)
+    rows = i32(1 + torch.arange(b))
+    ids = (torch.arange(b * t) % 5).to(device)
+    # each row agrees with its drafts for a row-dependent stretch; every third row meets a stop id
+    targets = ids.roll(-1).clone()
+    for r in range(b):
+        targets[r * t + r % t] = 99 if r % 3 else 7
+    positions = (torch.arange(b)[:, None] * 13 + torch.arange(t)[None, :] + 20).reshape(-1).to(device)
+    block_table = i32(1 + torch.arange(b * width).view(b, width))
+    return (rows, ids, targets, positions, block_table, page_size, torch.tensor([7, 8], **long),
+            torch.zeros(registers, dtype=torch.int32, device=device),
+            torch.ones(registers, dtype=torch.int32, device=device),
+            torch.zeros(t, registers, **long),
+            i32(torch.arange(registers)),
+            torch.full((registers,), 5, dtype=torch.int32, device=device),
+            i32(torch.arange(registers) % 2),
+            torch.zeros(b, t, **long), torch.zeros(b, dtype=torch.int32, device=device),
+            torch.zeros(b, dtype=torch.bool, device=device), torch.zeros(b, **long),
+            torch.zeros(k, b, **long), torch.zeros(k, b, **long),
+            torch.zeros(k, b, dtype=torch.int32, device=device))
+
+
+SPEC_ACCEPT = KernelOp(
+    "spec_accept",
+    "(Tensor rows, Tensor ids, Tensor targets, Tensor positions, Tensor block_table, int page_size, Tensor stop_ids, "
+    "Tensor(a!) seq_len, Tensor(b!) num_acc, Tensor(c!) tokens, Tensor(d!) emitted, Tensor budget, Tensor stop_on, "
+    "Tensor(e!) out_tokens, Tensor(f!) accepted, Tensor(g!) done, Tensor(h!) keep, Tensor(i!) draft_slots, "
+    "Tensor(j!) draft_positions, Tensor(k!) draft_lens) -> ()",
+    fake=_accept_fake,
+    reference=_accept_reference,
+    impls={"rocm": lambda: _accept},
+    fallback_to_reference=True,
+    shapes=ShapePolicy(
+        dims=("rows", "tokens"),
+        key=_accept_shape,
+        inputs=_accept_args,
+        warm=({"rows": 6, "tokens": 3}, {"rows": 6, "tokens": 4}),
+        dtypes=(torch.int32,),
+    ),
+    tolerance=Tolerance.exact(),
+)
diff --git a/mstar/engine/resources/spec_decode/resource.py b/mstar/engine/resources/spec_decode/resource.py
new file mode 100644
index 0000000000..fcc7d97710
--- /dev/null
+++ b/mstar/engine/resources/spec_decode/resource.py
@@ -0,0 +1,600 @@
+"""Speculative decoding: device-side registers for draft-and-verify steps.
+
+A verify row carries k + 1 tokens ``[x_0, d_1 .. d_k]`` at positions ``P ..
+P + k``: ``x_0`` is the last token the request emitted, the ``d_i`` the drafts,
+``P`` the request's committed length. The target runs them as one causal
+extend, and greedy acceptance keeps the longest run of drafts that equal the
+target's own next token (``n`` of them), so the step emits ``n + 1`` tokens and
+the stream grows by as many. The draft then runs ``k`` passes: pass 0 over the
+verify's ``(rows, k + 1)`` tokens, fed the target's hidden state at ``P + i`` and
+the target's token after it, which writes the draft's K/V for every position
+the step may have kept; pass ``j > 0`` one token a row at ``P + n + j``, the
+draft's own output from the pass before. Only slot ``n`` of pass 0 is a real
+query, so every pass attends one query a row, over the verify's block table,
+and its work after attention runs on ``rows`` rows. A pass's token is draft ``j
++ 1`` of the next step, written straight into that step's registers.
+
+``P``, ``n`` and the next step's tokens live here, one register row per
+request, and the host never reads them before the next step runs: the async
+worker plans step N+1 while N is on the GPU, so N+1 is planned against an
+optimistic length (all of N accepted) and corrected on the device. The KV write
+slots, the attention lengths, the rotary positions and the recurrent state's
+start slot all derive from ``P`` in-graph, through buffers the KV and attention
+resources hand over in their plan results. The host learns ``n`` from the D2H
+``check_stop`` already reads, and gives back what it over-counted in a later
+commit (``KVStep.commit_adjust``).
+
+What keeps the optimistic reservation sufficient: at every admit the host's
+stream length is at least ``P + k - 1``. A verify step writes target KV at ``P ..
+P + k`` and draft KV at up to ``P + 2k - 1``, and its view covers ``stored + k +
+1 >= P + 2k``. The seed commit (the prompt's last prefill) adds ``k - 1`` past
+the prompt to establish it, and each learned step gives back exactly what it
+over-counted, ``(k + 1) - (n + 1)``.
+
+A row also stops itself: tokens past a stop id or past the request's budget
+are emitted as -1, and ``done`` tells the host, so a step never needs the host
+to cut its output.
+"""
+
+from __future__ import annotations
+
+import logging
+import threading
+from collections.abc import Mapping, Sequence
+from typing import NamedTuple
+
+import numpy as np
+import torch
+
+from mstar.engine.resources.attn.prologue import QKV_NORM_ROPE_CACHE
+from mstar.engine.resources.base import CGSlotKey, CGSlotSpec, EngineResourceInfo, Resource
+from mstar.engine.resources.sampler import utils as sampling
+from mstar.engine.resources.sampler.unfiltered import SAMPLE_UNFILTERED
+from mstar.engine.resources.spec import ResourceReqConfig
+from mstar.engine.resources.spec_decode.config import (
+    AcceptedTokens,
+    DraftAddress,
+    SpecDecodeSpec,
+    SpecDecodeStep,
+    SpecPhase,
+)
+from mstar.engine.resources.spec_decode.kernels import SPEC_ACCEPT, SPEC_VERIFY_INPUTS
+from mstar.engine.resources.step import (
+    ADMIT_OK,
+    AdmitOutcome,
+    AllocationFailed,
+    StepContext,
+)
+from mstar.utils.h2d import H2DMirror, PinnedStager
+
+logger = logging.getLogger(__name__)
+
+# The register row padding rows read and write; never handed to a request.
+SINK_ROW = 0
+# Fields of the per-step image staged in one H2D copy, each `rows` long.
+_FIELDS = ("rows", "real", "base", "budget", "stop_on", "top_k")
+
+
+class _StepImage(NamedTuple):
+    """One step's host-known per-row values on the device, as views of one
+    int32 buffer filled by one copy, plus a buffer the step fills in-graph."""
+    image: torch.Tensor          # [len(_FIELDS) * rows] int32
+    rows: torch.Tensor           # register row per step row, SINK_ROW for padding
+    real: torch.Tensor           # 1 for a real request, 0 for padding
+    base: torch.Tensor           # seed steps: the stream's length after the prompt
+    budget: torch.Tensor         # seed steps: decode tokens the request may emit
+    stop_on: torch.Tensor        # seed steps: 0 where the request ignores stop ids
+    top_k: torch.Tensor          # seed steps: the request's top_k (0: off)
+    # seed steps: [temperature | top_p] per row, float32, its own copy
+    params: torch.Tensor
+    seeds: torch.Tensor          # seed steps: each request's sampling seed, int64
+    num_accepted: torch.Tensor   # verify steps: the row's last accepted count, for the recurrent kernels
+    # What an argmax through the unfiltered draw reads, [rows * (k + 1)]:
+    # temperature 0, and offsets it stores back unchanged. Never read for a token.
+    zero_temperature: torch.Tensor
+    zero_seed: torch.Tensor
+    zero_offset: torch.Tensor
+
+
+class SpecDecodePlan(NamedTuple):
+    """What ``plan`` hands its dependents: the phase, the rows, and the device
+    buffers of the other resources a verify step rewrites per pass."""
+    phase: SpecPhase
+    num_rows: int
+    image: _StepImage
+    # attention's block table and per-row KV lengths, and KV's write slots
+    block_table: torch.Tensor | None = None
+    kv_lens: torch.Tensor | None = None
+    write_pages: torch.Tensor | None = None
+    write_offsets: torch.Tensor | None = None
+
+    @property
+    def num_accepted(self) -> torch.Tensor:
+        """Each row's accepted count from its previous verify (1 after a seed),
+        filled by ``verify_inputs`` before the layers run."""
+        return self.image.num_accepted[: self.num_rows]
+
+
+class SpecDecodeResource(Resource):
+    @classmethod
+    def build(cls, spec: SpecDecodeSpec, info: EngineResourceInfo) -> SpecDecodeResource:
+        kv = info.dependency(spec.kv_cache)
+        page_size = getattr(kv.config, "tokens_per_page", None)
+        if page_size is None:
+            raise ValueError(
+                f"speculative decoding writes paged KV; {spec.kv_cache!r} is "
+                f"{type(kv.config).__name__}"
+            )
+        return cls(
+            device=info.device,
+            num_draft_tokens=spec.config.num_draft_tokens,
+            sampling=spec.config.sampling,
+            max_requests=spec.config.max_requests,
+            page_size=page_size,
+            kv_key=spec.kv_cache,
+            attn_key=spec.attn,
+            stop_token_ids=spec.stop_token_ids,
+        )
+
+    def __init__(
+        self,
+        device: torch.device,
+        num_draft_tokens: int,
+        max_requests: int,
+        page_size: int,
+        kv_key: str,
+        attn_key: str,
+        stop_token_ids: Sequence[int] = (),
+        sampling: bool = False,
+    ):
+        self._device = device
+        self.num_draft_tokens = num_draft_tokens
+        self.tokens_per_row = num_draft_tokens + 1
+        self._page_size = page_size
+        self._kv_key = kv_key
+        self._attn_key = attn_key
+        rows = max_requests + 1
+        i32 = dict(dtype=torch.int32, device=device)
+        # Per request: committed length, last accepted count, the next verify's
+        # tokens (by slot, so a draft pass's token master is a plain column),
+        # decode tokens emitted, the budget, and whether stop ids apply.
+        self._seq_len = torch.zeros(rows, **i32)
+        self._num_acc = torch.ones(rows, **i32)
+        self._tokens = torch.zeros(self.tokens_per_row, rows, dtype=torch.long, device=device)
+        self._emitted = torch.zeros(rows, **i32)
+        self._budget = torch.zeros(rows, **i32)
+        self._stop_on = torch.zeros(rows, **i32)
+        # per request, for a sampling node: temperature (0 greedy), top-p, top-k
+        self.sampling = sampling
+        self._temperature = torch.zeros(rows, dtype=torch.float32, device=device)
+        self._top_p = torch.ones(rows, dtype=torch.float32, device=device)
+        self._top_k = torch.zeros(rows, **i32)
+        # the request's seed, and the Philox offset its next verify step draws at
+        self._seed = torch.zeros(rows, dtype=torch.long, device=device)
+        self._rng_offset = torch.zeros(rows, dtype=torch.long, device=device)
+        self._stop_ids = torch.tensor(
+            sorted(set(stop_token_ids)) or [-2], dtype=torch.long, device=device,
+        )
+        self._zero_master = torch.zeros(rows, dtype=torch.long, device=device)
+
+        self._lock = threading.Lock()
+        # rid -> register row; reversed so `pop` hands out ascending rows
+        self._row_of: dict[int, int] = {}
+        self._free = list(reversed(range(1, rows)))
+        # rid -> KV tokens over-counted by verify steps the host has learned
+        # about but not yet given back
+        self._retract: dict[int, int] = {}
+
+        self._stager = PinnedStager(torch.int32)
+        self._param_stager = PinnedStager(torch.float32)
+        self._seed_stager = PinnedStager(torch.int64)
+        self._cg_max_bs = 0
+        self._cg_images: dict[CGSlotKey, tuple[_StepImage, H2DMirror]] = {}
+        self._eager_images: dict[int, _StepImage] = {}
+
+        self._current: SpecDecodePlan | None = None
+        self._preplanned = False
+        self._staged: SpecDecodePlan | None = None
+        # set by `verify_inputs` / `accept_tokens` within one forward
+        self._ids: torch.Tensor | None = None
+        self._positions: torch.Tensor | None = None
+        self._slots: torch.Tensor | None = None
+        self._keep: torch.Tensor | None = None
+        # [k, rows]: each draft pass's KV slot, position and keys, computed
+        # together once the step knows what it accepted
+        self._draft_slots: torch.Tensor | None = None
+        self._draft_positions: torch.Tensor | None = None
+        self._draft_lens: torch.Tensor | None = None
+        # bound by the engine; a draft pass writes and attends through them
+        self._kv = None
+        self._attn = None
+
+    def depends_on(self) -> set[str]:
+        # their plans are what this one rewrites on the device
+        return {self._kv_key, self._attn_key}
+
+    def bind_dependencies(self, resources: Mapping[str, Resource]) -> None:
+        self._kv = resources.get(self._kv_key)
+        self._attn = resources.get(self._attn_key)
+
+    # Request lifecycle
+
+    def ingest_request(self, rid: int, overrides: ResourceReqConfig | None = None):
+        with self._lock:
+            if rid in self._row_of or not self._free:
+                return
+            self._row_of[rid] = self._free.pop()
+
+    def ingest_capture_row(self, rid: int):
+        # capture rows address the sink row
+        return
+
+    def remove_request(self, rid: int):
+        with self._lock:
+            row = self._row_of.pop(rid, None)
+            self._retract.pop(rid, None)
+            if row is not None:
+                self._free.append(row)
+
+    # Host bookkeeping
+
+    def learn(self, request_ids: Sequence[int], accepted: Sequence[int]) -> None:
+        """Record what verify steps accepted, read off the host copy
+        ``check_stop`` already has. Each over-counted ``(k + 1) - accepted``
+        KV tokens; the next step that declares this request gives them back."""
+        full = self.tokens_per_row
+        with self._lock:
+            for rid, count in zip(request_ids, accepted, strict=True):
+                if rid in self._row_of:
+                    self._retract[rid] = self._retract.get(rid, 0) + full - int(count)
+
+    def pending_retract(self, rid: int) -> int:
+        """KV tokens a step declaring ``rid`` now may give back."""
+        return self._retract.get(rid, 0)
+
+    @property
+    def seed_lookahead(self) -> int:
+        """KV tokens a seed commit reserves past the prompt; see the module
+        docstring for why a verify step needs them."""
+        return self.num_draft_tokens - 1
+
+    # Step lifecycle
+
+    def admit(self, step: SpecDecodeStep, ctx: StepContext) -> AdmitOutcome:
+        if ctx.capture:
+            return ADMIT_OK
+        for seg in step.segments or ():
+            rid = seg.request_id
+            if ctx.is_padding_row(rid) or rid in self._row_of:
+                continue
+            self.ingest_request(rid)
+            if rid not in self._row_of:
+                return AdmitOutcome(ok=False, reason=AllocationFailed(
+                    message=(
+                        f"speculative decoding holds {len(self._row_of)} requests, "
+                        "its max_requests; none free"
+                    ),
+                    pages_short=1, label=seg.label, request_id=rid,
+                ))
+        return ADMIT_OK
+
+    @property
+    def supports_preplan(self):
+        return True
+
+    @property
+    def force_double_buffer(self):
+        # a step's rows are staged into a slot's buffer behind an async H2D
+        return True
+
+    def clear_preplan(self):
+        self._preplanned = False
+        self._staged = None
+
+    def plan(self, step: SpecDecodeStep, ctx: StepContext) -> SpecDecodePlan:
+        if self._preplanned and not ctx.is_preplan:
+            self._current = self._staged
+            self.clear_preplan()
+            return self._current
+        segments = step.segments or ()
+        num_rows = len(segments)
+        image, mirror = self._image_for(ctx, num_rows)
+        rows, real, base, budget, stop_on, top_k = (
+            np.zeros(num_rows, dtype=np.int32) for _ in _FIELDS
+        )
+        temperature = np.zeros(num_rows, dtype=np.float32)
+        top_p = np.ones(num_rows, dtype=np.float32)
+        seeds = np.zeros(num_rows, dtype=np.int64)
+        views = None
+        if step.phase is SpecPhase.SEED:
+            kv_plan = ctx.plan_results[self._kv_key].get("main")
+            views = kv_plan.views if kv_plan is not None else None
+        for i, seg in enumerate(segments):
+            rid = seg.request_id
+            if ctx.capture or ctx.is_padding_row(rid):
+                continue
+            rows[i] = self._row_of[rid]
+            real[i] = 1
+            if views is not None:
+                base[i] = views[i].length
+                budget[i] = step.budgets.get(rid, 0)
+                stop_on[i] = rid not in step.ignore_stop
+                if rid in step.sampling:
+                    temperature[i], top_p[i], top_k[i], seeds[i] = step.sampling[rid]
+        width = image.rows.numel()
+        packed = np.zeros(len(_FIELDS) * width, dtype=np.int32)
+        for f, values in enumerate((rows, real, base, budget, stop_on, top_k)):
+            packed[f * width: f * width + num_rows] = values
+        self._stager.copy_(image.image, packed, mirror=mirror)
+        if self.sampling and step.phase is SpecPhase.SEED:
+            params = np.ones(2 * width, dtype=np.float32)
+            params[:num_rows] = temperature
+            params[width: width + num_rows] = top_p
+            self._param_stager.copy_(image.params, params)
+            self._seed_stager.copy_(image.seeds, seeds, pad_value=0)
+
+        plan = SpecDecodePlan(phase=step.phase, num_rows=num_rows, image=image)
+        if step.phase is SpecPhase.VERIFY:
+            plan = self._verify_plan(plan, ctx)
+        if ctx.is_preplan:
+            self._preplanned = True
+            self._staged = plan
+        else:
+            self._current = plan
+        return plan
+
+    def _verify_plan(self, plan: SpecDecodePlan, ctx: StepContext) -> SpecDecodePlan:
+        attn = (ctx.plan_results.get(self._attn_key) or {}).get("main")
+        kv = (ctx.plan_results.get(self._kv_key) or {}).get("main")
+        block_table = getattr(attn, "block_table", None)
+        kv_lens = getattr(attn, "seqused_k", None)
+        write_state = getattr(kv, "write_state", None)
+        if block_table is None or kv_lens is None or write_state is None:
+            raise RuntimeError(
+                "speculative verify needs an attention backend whose plan reads "
+                "per-row KV lengths and a block table on the device (AITER's "
+                f"unified attention); {self._attn_key!r} planned {type(attn).__name__}"
+            )
+        n = plan.num_rows * self.tokens_per_row
+        return plan._replace(
+            block_table=block_table[: plan.num_rows],
+            kv_lens=kv_lens[: plan.num_rows],
+            write_pages=write_state.token_to_page[:n],
+            write_offsets=write_state.token_to_cache[:n],
+        )
+
+    def commit(self, step: SpecDecodeStep, ctx: StepContext) -> None:
+        if not step.retract:
+            return
+        with self._lock:
+            for rid, given in step.retract.items():
+                left = self._retract.get(rid, 0) - given
+                if left > 0:
+                    self._retract[rid] = left
+                else:
+                    self._retract.pop(rid, None)
+
+    def _image_for(self, ctx: StepContext, num_rows: int) -> tuple[_StepImage, H2DMirror | None]:
+        lease = ctx.slot_lease
+        if lease is None:
+            image = self._eager_images.get(ctx.slot)
+            if image is None or image.rows.numel() < num_rows:
+                image = self._eager_images[ctx.slot] = self._new_image(max(num_rows, 8))
+            return image, None
+        key = CGSlotKey(bucket=lease.bucket, slot=lease.slot, label="main")
+        hit = self._cg_images.get(key)
+        if hit is None:
+            size = max(self._cg_max_bs, lease.bucket.bs, num_rows)
+            hit = self._cg_images[key] = (self._new_image(size), H2DMirror())
+        return hit
+
+    def _new_image(self, rows: int) -> _StepImage:
+        image = torch.zeros(len(_FIELDS) * rows, dtype=torch.int32, device=self._device)
+        views = image.view(len(_FIELDS), rows)
+        tokens = rows * self.tokens_per_row
+        return _StepImage(
+            image=image, rows=views[0], real=views[1], base=views[2],
+            budget=views[3], stop_on=views[4], top_k=views[5],
+            params=torch.ones(2 * rows, dtype=torch.float32, device=self._device),
+            seeds=torch.zeros(rows, dtype=torch.long, device=self._device),
+            num_accepted=torch.ones(rows, dtype=torch.int32, device=self._device),
+            zero_temperature=torch.zeros(tokens, dtype=torch.float32, device=self._device),
+            zero_seed=torch.zeros(tokens, dtype=torch.long, device=self._device),
+            zero_offset=torch.zeros(tokens, dtype=torch.long, device=self._device),
+        )
+
+    def build_accelerator_graph_buffers(
+        self, slots: list[CGSlotSpec], max_bs: int, max_seq_len: int,
+    ) -> None:
+        del slots, max_seq_len
+        self._cg_max_bs = max(self._cg_max_bs, max_bs)
+
+    def cleanup(self):
+        self._cg_images.clear()
+        self._eager_images.clear()
+
+    # In-graph operations, called by the forward in this order: `seed` on a
+    # prompt's step; `verify_inputs`, `accept`, then `attend_draft` and
+    # `record_draft` once per draft pass on a decode step.
+
+    def _plan_now(self, phase: SpecPhase) -> SpecDecodePlan:
+        plan = self._current
+        if plan is None or plan.phase is not phase:
+            raise RuntimeError(
+                f"speculative decoding has no {phase.value} plan for this step; "
+                "the step must declare a SpecDecodeStep of that phase"
+            )
+        return plan
+
+    def _rows(self, plan: SpecDecodePlan) -> torch.Tensor:
+        return plan.image.rows[: plan.num_rows].long()
+
+    @torch.compiler.disable
+    def seed(self, first_token: torch.Tensor, first_draft: torch.Tensor) -> None:
+        """Start each prompt row's decode: its first token (sampled off the
+        prompt), the draft head's guess after it, and its committed length."""
+        plan = self._plan_now(SpecPhase.SEED)
+        rows = self._rows(plan)
+        b = plan.num_rows
+        image = plan.image
+        tokens = torch.cat(
+            (first_token.view(1, b).long(),
+             first_draft.view(1, b).long().expand(self.num_draft_tokens, b)),
+            dim=0,
+        )
+        self._tokens.index_copy_(1, rows, tokens)
+        self._seq_len.index_copy_(0, rows, image.base[:b])
+        self._num_acc.index_fill_(0, rows, 1)
+        self._emitted.index_fill_(0, rows, 0)
+        self._budget.index_copy_(0, rows, image.budget[:b])
+        self._stop_on.index_copy_(0, rows, image.stop_on[:b])
+        if self.sampling:
+            width = image.rows.numel()
+            self._temperature.index_copy_(0, rows, image.params[:b])
+            self._top_p.index_copy_(0, rows, image.params[width: width + b])
+            self._top_k.index_copy_(0, rows, image.top_k[:b])
+            self._seed.index_copy_(0, rows, image.seeds[:b])
+            self._rng_offset.index_fill_(0, rows, 0)
+
+    @torch.compiler.disable
+    def verify_inputs(self) -> tuple[torch.Tensor, torch.Tensor]:
+        """The verify tokens and their positions, ``[rows * (k + 1)]`` each;
+        points this step's KV writes, attention lengths and recurrent start
+        slots at the request's real length."""
+        plan = self._plan_now(SpecPhase.VERIFY)
+        image = plan.image
+        b, n = plan.num_rows, plan.num_rows * self.tokens_per_row
+        long = dict(dtype=torch.long, device=self._device)
+        self._ids = torch.empty(n, **long)
+        self._positions = torch.empty(n, **long)
+        self._slots = torch.empty(n, **long)
+        SPEC_VERIFY_INPUTS(
+            image.rows[:b], image.real[:b], self._seq_len, self._num_acc, self._tokens,
+            plan.block_table, self._page_size, self._ids, self._positions, self._slots,
+            plan.write_pages, plan.write_offsets, plan.kv_lens, plan.num_accepted,
+        )
+        return self._ids, self._positions
+
+    @torch.compiler.disable
+    def accept(self, logits: torch.Tensor, sample: bool = False) -> AcceptedTokens:
+        """Verification of this step's drafts against ``logits`` (``[rows *
+        (k + 1), vocab]``, the target's at each verify position): its argmax,
+        or with ``sample`` a draw from each row's top-k/top-p distribution at
+        its temperature (greedy where the row's temperature is 0).
+
+        The drafts are the draft head's argmax, a point mass, for which
+        rejection sampling is: draw the target's token at every position
+        independently and accept while it equals the next draft. A rejected
+        position then holds a draw from the target's distribution with the
+        draft taken out, which is the residual the rule prescribes, so the
+        emitted tokens follow the target exactly. The draws use the
+        request's seed (see `_sample`), so they repeat for it, but they are not
+        the draws a non-speculative step would make with that seed.
+        """
+        if not sample:
+            return self.accept_tokens(self._argmax(logits))
+        if not self.sampling:
+            raise RuntimeError("this speculative node was declared without sampling")
+        return self.accept_tokens(self._sample(logits))
+
+    def _argmax(self, logits: torch.Tensor) -> torch.Tensor:
+        """Each row's argmax, through the unfiltered draw's two kernels."""
+        n, image = logits.shape[0], self._plan_now(SpecPhase.VERIFY).image
+        out = torch.empty(n, dtype=torch.long, device=logits.device)
+        SAMPLE_UNFILTERED(logits, image.zero_temperature[:n], image.zero_seed[:n], image.zero_offset[:n], out, 0)
+        return out
+
+    def _sample(self, logits: torch.Tensor) -> torch.Tensor:
+        """One draw per row of ``logits``: top-k, then top-p over what is
+        left, at the row's temperature (argmax at temperature 0), by the
+        sampler's captured kernel (`draw_tokens`). Each verify position draws
+        from its own stream (the request's seed, offset per position), so the
+        positions are independent; the offset advances once a step, so steps
+        are too, and a request's draws repeat for its seed."""
+        plan = self._plan_now(SpecPhase.VERIFY)
+        rows = self._rows(plan)
+        t, vocab = self.tokens_per_row, logits.shape[-1]
+        temperature = self._temperature[rows].repeat_interleave(t)
+        top_p = self._top_p[rows].repeat_interleave(t)
+        # 0 is off for the request, which the kernel would read as k = 0
+        top_k = self._top_k[rows].repeat_interleave(t)
+        top_k = torch.where(top_k > 0, top_k, vocab)
+        offset = self._rng_offset[rows]
+        position = torch.arange(t, dtype=torch.long, device=logits.device)
+        seed = self._seed[rows].repeat_interleave(t)
+        offsets = (offset[:, None] * t + position[None, :]).reshape(-1)
+        self._rng_offset.index_copy_(0, rows, offset + 1)
+        return sampling.draw_tokens(logits, temperature, top_k, top_p, seed, offsets).long()
+
+    @torch.compiler.disable
+    def accept_tokens(self, target: torch.Tensor) -> AcceptedTokens:
+        """Verification against the target's own token at each verify position
+        (``[rows * (k + 1)]``): its argmax, or what a sampler drew there."""
+        plan = self._plan_now(SpecPhase.VERIFY)
+        image = plan.image
+        b, t, k = plan.num_rows, self.tokens_per_row, self.num_draft_tokens
+        device = self._device
+        long = dict(dtype=torch.long, device=device)
+        tokens = torch.empty(b, t, **long)
+        accepted = torch.empty(b, dtype=torch.int32, device=device)
+        done = torch.empty(b, dtype=torch.bool, device=device)
+        self._keep = torch.empty(b, **long)
+        self._draft_slots = torch.empty(k, b, **long)
+        self._draft_positions = torch.empty(k, b, **long)
+        self._draft_lens = torch.empty(k, b, dtype=torch.int32, device=device)
+        SPEC_ACCEPT(
+            image.rows[:b], self._ids, target, self._positions, plan.block_table, self._page_size,
+            self._stop_ids, self._seq_len, self._num_acc, self._tokens, self._emitted, self._budget,
+            self._stop_on, tokens, accepted, done, self._keep,
+            self._draft_slots, self._draft_positions, self._draft_lens,
+        )
+        return AcceptedTokens(tokens=tokens, accepted=accepted, done=done, target=target.view(b, t))
+
+    def draft_address(self, j: int) -> DraftAddress:
+        """Where draft pass ``j`` writes and reads: pass 0 writes every verify
+        position, the others one position a row."""
+        if j == 0:
+            return DraftAddress(self._slots, self._positions, self._draft_lens[0])
+        return DraftAddress(self._draft_slots[j], self._draft_positions[j], self._draft_lens[j])
+
+    @torch.compiler.disable
+    def draft_slots(self) -> torch.Tensor:
+        """Each row's slot ``n``, the one real query of draft pass 0, as flat
+        indices into ``[rows * (k + 1)]``."""
+        return self._keep
+
+    @torch.compiler.disable
+    def attend_draft(
+        self, j: int, qkv: torch.Tensor, q_weight: torch.Tensor, k_weight: torch.Tensor,
+        cos: torch.Tensor, sin: torch.Tensor, heads: int, kv_heads: int, head_dim: int,
+        eps: float, gated: bool, layer_idx: int,
+    ) -> tuple[torch.Tensor, torch.Tensor]:
+        """Draft pass ``j``'s attention over its packed ``[q | k | v]`` tokens:
+        split, q/k norm, rotation (``cos``/``sin`` are tables over positions) and
+        the KV write as one kernel, then one query a row. Returns the attention
+        output and the gate, ``[rows, heads, head_dim]`` each."""
+        where = self.draft_address(j)
+        layer = self._kv.layer_view(layer_idx)
+        q, gate, _, _ = QKV_NORM_ROPE_CACHE(
+            qkv, q_weight, k_weight, cos, sin, where.positions, layer[:, 0], layer[:, 1], where.slots,
+            heads, kv_heads, head_dim, eps, gated,
+        )
+        if j == 0:
+            q = q.index_select(0, self._keep)
+            if gated:
+                gate = gate.index_select(0, self._keep)
+        return self._attn.attend_rows(q, layer, where.kv_lens), gate
+
+    @torch.compiler.disable
+    def record_draft(self, j: int, logits: torch.Tensor) -> torch.Tensor:
+        """Pass ``j``'s token of each row, the argmax of ``logits`` (``[rows,
+        vocab]``): draft ``j + 1`` of the next step, written to its register by
+        the same kernel, and the next pass's input."""
+        plan = self._plan_now(SpecPhase.VERIFY)
+        b, image = plan.num_rows, plan.image
+        drafted = torch.empty(b, dtype=torch.long, device=logits.device)
+        SAMPLE_UNFILTERED(
+            logits, image.zero_temperature[:b], image.zero_seed[:b], image.zero_offset[:b], drafted, 0,
+            image.rows[:b], self._zero_master, self._tokens[j + 1],
+        )
+        return drafted
diff --git a/mstar/kernels/configs/gfx942/mi325x/decode_gemv.json b/mstar/kernels/configs/gfx942/mi325x/decode_gemv.json
index 767f61a616..b0e2d760cd 100644
--- a/mstar/kernels/configs/gfx942/mi325x/decode_gemv.json
+++ b/mstar/kernels/configs/gfx942/mi325x/decode_gemv.json
@@ -23,7 +23,8 @@
   "harness": "benchmark.kernels.search, one MI325X VF GPU, one run per cell",
   "jobs": [
    "45",
-   "83"
+   "83",
+   "390"
   ],
   "margin": 0.02,
   "note": "Triton dot/FMA families vs hipBLASLt and AITER wvSpltK; jobs 45 and 83; measured shapes only, classes keep AITER for 1-2 rows and hipBLASLt above",
@@ -31,6 +32,28 @@
   "store": "amd/results/innovation/f2-kernels/store/gfx942/mi325x"
  },
  "shapes": {
+  "k12288-n4096-rows12": {
+   "block_k": 256,
+   "block_n": 32,
+   "impl": "triton",
+   "num_stages": 1,
+   "num_warps": 2,
+   "split_k": 8,
+   "use_dot": true,
+   "waves_per_eu": 2,
+   "xcd_remap": false
+  },
+  "k12288-n4096-rows3": {
+   "block_k": 256,
+   "block_n": 32,
+   "impl": "triton",
+   "num_stages": 1,
+   "num_warps": 2,
+   "split_k": 8,
+   "use_dot": true,
+   "waves_per_eu": 2,
+   "xcd_remap": false
+  },
   "k12288-n4096-rows4": {
    "block_k": 256,
    "block_n": 32,
@@ -42,6 +65,17 @@
    "waves_per_eu": 2,
    "xcd_remap": false
   },
+  "k12288-n4096-rows6": {
+   "block_k": 256,
+   "block_n": 32,
+   "impl": "triton",
+   "num_stages": 1,
+   "num_warps": 2,
+   "split_k": 8,
+   "use_dot": true,
+   "waves_per_eu": 2,
+   "xcd_remap": false
+  },
   "k2048-n12288-rows16": {
    "block_k": 256,
    "block_n": 32,
@@ -128,6 +162,17 @@
    "waves_per_eu": 2,
    "xcd_remap": false
   },
+  "k4096-n12352-rows3": {
+   "block_k": 256,
+   "block_n": 32,
+   "impl": "triton",
+   "num_stages": 1,
+   "num_warps": 2,
+   "split_k": 8,
+   "use_dot": true,
+   "waves_per_eu": 2,
+   "xcd_remap": false
+  },
   "k4096-n12352-rows4": {
    "block_k": 256,
    "block_n": 32,
@@ -172,6 +217,17 @@
    "waves_per_eu": 2,
    "xcd_remap": false
   },
+  "k4096-n248320-rows12": {
+   "block_k": 256,
+   "block_n": 32,
+   "impl": "triton",
+   "num_stages": 2,
+   "num_warps": 2,
+   "split_k": 1,
+   "use_dot": true,
+   "waves_per_eu": 0,
+   "xcd_remap": true
+  },
   "k4096-n248320-rows16": {
    "block_k": 128,
    "block_n": 32,
@@ -183,9 +239,23 @@
    "waves_per_eu": 0,
    "xcd_remap": true
   },
+  "k4096-n248320-rows3": {
+   "impl": "aiter"
+  },
   "k4096-n248320-rows4": {
    "impl": "aiter"
   },
+  "k4096-n248320-rows6": {
+   "block_k": 256,
+   "block_n": 32,
+   "impl": "triton",
+   "num_stages": 2,
+   "num_warps": 2,
+   "split_k": 1,
+   "use_dot": true,
+   "waves_per_eu": 0,
+   "xcd_remap": true
+  },
   "k4096-n248320-rows8": {
    "block_k": 256,
    "block_n": 32,
@@ -197,6 +267,17 @@
    "waves_per_eu": 0,
    "xcd_remap": true
   },
+  "k4096-n4096-rows12": {
+   "block_k": 256,
+   "block_n": 32,
+   "impl": "triton",
+   "num_stages": 1,
+   "num_warps": 2,
+   "split_k": 8,
+   "use_dot": true,
+   "waves_per_eu": 2,
+   "xcd_remap": false
+  },
   "k4096-n4096-rows16": {
    "block_k": 256,
    "block_n": 32,
@@ -208,6 +289,17 @@
    "waves_per_eu": 2,
    "xcd_remap": false
   },
+  "k4096-n4096-rows3": {
+   "block_k": 256,
+   "block_n": 32,
+   "impl": "triton",
+   "num_stages": 1,
+   "num_warps": 2,
+   "split_k": 8,
+   "use_dot": true,
+   "waves_per_eu": 2,
+   "xcd_remap": false
+  },
   "k4096-n4096-rows4": {
    "block_k": 256,
    "block_n": 32,
@@ -219,6 +311,17 @@
    "waves_per_eu": 2,
    "xcd_remap": false
   },
+  "k4096-n4096-rows6": {
+   "block_k": 256,
+   "block_n": 32,
+   "impl": "triton",
+   "num_stages": 1,
+   "num_warps": 2,
+   "split_k": 8,
+   "use_dot": true,
+   "waves_per_eu": 2,
+   "xcd_remap": false
+  },
   "k4096-n4096-rows8": {
    "block_k": 256,
    "block_n": 32,
@@ -251,6 +354,31 @@
    "use_dot": true,
    "waves_per_eu": 2,
    "xcd_remap": false
+  },
+  "k8192-n4096-rows12": {
+   "block_k": 256,
+   "block_n": 32,
+   "impl": "triton",
+   "num_stages": 1,
+   "num_warps": 2,
+   "split_k": 8,
+   "use_dot": true,
+   "waves_per_eu": 2,
+   "xcd_remap": false
+  },
+  "k8192-n4096-rows16": {
+   "block_k": 256,
+   "block_n": 32,
+   "impl": "triton",
+   "num_stages": 1,
+   "num_warps": 2,
+   "split_k": 8,
+   "use_dot": true,
+   "waves_per_eu": 2,
+   "xcd_remap": false
+  },
+  "k8192-n4096-rows3": {
+   "impl": "aiter"
   }
  }
 }
diff --git a/mstar/kernels/configs/gfx942/mi325x/decode_gemv_silu_mul.json b/mstar/kernels/configs/gfx942/mi325x/decode_gemv_silu_mul.json
index 24107a951d..31335abca4 100644
--- a/mstar/kernels/configs/gfx942/mi325x/decode_gemv_silu_mul.json
+++ b/mstar/kernels/configs/gfx942/mi325x/decode_gemv_silu_mul.json
@@ -23,7 +23,8 @@
   "harness": "benchmark.kernels.search, one MI325X VF GPU, one run per cell",
   "jobs": [
    "45",
-   "83"
+   "83",
+   "390"
   ],
   "margin": 0.02,
   "note": "fused gate/up + silu*mul vs the unfused pair (AITER or hipBLASLt, then inductor's silu*mul); jobs 45 and 83; measured shapes only",
@@ -85,6 +86,17 @@
    "waves_per_eu": 2,
    "xcd_remap": false
   },
+  "k4096-n24576-rows12": {
+   "block_k": 256,
+   "block_n": 32,
+   "impl": "triton",
+   "num_stages": 2,
+   "num_warps": 2,
+   "split_k": 1,
+   "use_dot": true,
+   "waves_per_eu": 0,
+   "xcd_remap": true
+  },
   "k4096-n24576-rows16": {
    "block_k": 256,
    "block_n": 32,
@@ -107,6 +119,17 @@
    "waves_per_eu": 2,
    "xcd_remap": false
   },
+  "k4096-n24576-rows3": {
+   "block_k": 256,
+   "block_n": 32,
+   "impl": "triton",
+   "num_stages": 2,
+   "num_warps": 2,
+   "split_k": 8,
+   "use_dot": true,
+   "waves_per_eu": 2,
+   "xcd_remap": false
+  },
   "k4096-n24576-rows4": {
    "block_k": 256,
    "block_n": 32,
@@ -118,6 +141,17 @@
    "waves_per_eu": 2,
    "xcd_remap": false
   },
+  "k4096-n24576-rows6": {
+   "block_k": 256,
+   "block_n": 32,
+   "impl": "triton",
+   "num_stages": 2,
+   "num_warps": 2,
+   "split_k": 8,
+   "use_dot": true,
+   "waves_per_eu": 2,
+   "xcd_remap": false
+  },
   "k4096-n24576-rows8": {
    "block_k": 256,
    "block_n": 32,
diff --git a/mstar/kernels/ops.py b/mstar/kernels/ops.py
index f4d462de72..635ea8fe49 100644
--- a/mstar/kernels/ops.py
+++ b/mstar/kernels/ops.py
@@ -45,8 +45,11 @@
 
 # Modules that declare ops, imported by `load_all` so a listing sees them all.
 DECLARING_MODULES = (
+    "mstar.engine.resources.attn.prologue",
     "mstar.engine.resources.rms_norm.rms_norm",
     "mstar.engine.resources.sampler.unfiltered",
+    "mstar.engine.resources.spec_decode.draft_head",
+    "mstar.engine.resources.spec_decode.kernels",
     "mstar.kernels.fp8_gemm",
     "mstar.platforms.rocm_gemm",
 )
diff --git a/mstar/model/qwen3_5/components/language_model.py b/mstar/model/qwen3_5/components/language_model.py
index 1688eca03a..de41de9f10 100644
--- a/mstar/model/qwen3_5/components/language_model.py
+++ b/mstar/model/qwen3_5/components/language_model.py
@@ -136,6 +136,22 @@ def forward(self, hidden_states: torch.Tensor) -> torch.Tensor:
             out = out * torch.sigmoid(gate.reshape(num_tokens, -1))
         return self.o_proj(out)
 
+    def forward_draft(self, hidden_states: torch.Tensor, attend) -> torch.Tensor:
+        """``forward`` for a speculative draft pass. ``attend(qkv, q_weight,
+        k_weight, heads, kv_heads, head_dim, eps, gated)`` takes the packed
+        projection and does the rest up to the attention, one query a row, so
+        this returns ``[rows, hidden]`` however many tokens came in."""
+        out, gate = attend(
+            self.qkv_proj(hidden_states), self.q_norm.weight, self.k_norm.weight,
+            self.num_heads, self.num_kv_heads, self.head_dim, self.q_norm.variance_epsilon,
+            self.output_gate,
+        )
+        rows = out.shape[0]
+        out = out.reshape(rows, self.num_heads * self.head_dim)
+        if self.output_gate:
+            out = out * torch.sigmoid(gate.reshape(rows, -1))
+        return self.o_proj(out)
+
 
 def _norm(config: Qwen3_5Config) -> RMSNorm:
     """Gemma-style RMSNorm (scales by ``1 + weight``); the gated norm is not."""
@@ -225,6 +241,7 @@ def __init__(
             persistent=False,
         )
         self.norm = _norm(config)
+        self._rope_tables: tuple[torch.Tensor, torch.Tensor] | None = None
 
         # precomputed: which cursor each layer advances, and to what index
         self._cursor_attr = [
@@ -270,6 +287,20 @@ def build_cos_sin(
             target_dtype=dtype,
         )
 
+    # out of dynamo: built once, before any capture, read by every replay
+    @torch.compiler.disable
+    def rope_tables(self, dtype: torch.dtype) -> tuple[torch.Tensor, torch.Tensor]:
+        """cos/sin over every position, ``[positions, rotary_dim // 2]``, for
+        a kernel that rotates from a position index. Text only: the three MRoPE
+        grids agree there, so this is the 1D table, rounded as ``build_cos_sin``
+        rounds."""
+        tables = self._rope_tables
+        if tables is None or tables[0].dtype != dtype or tables[0].device != self.inv_freq.device:
+            pos = torch.arange(self.config.max_position_embeddings, device=self.inv_freq.device).float()
+            freqs = self.inv_freq[None, :].float() * pos[:, None]
+            tables = self._rope_tables = (freqs.cos().to(dtype), freqs.sin().to(dtype))
+        return tables
+
     def forward(
         self, query_sequence: torch.Tensor, *, label: str,
         cos_sin: tuple[torch.Tensor, torch.Tensor] | None = None,
@@ -287,13 +318,90 @@ def _set_cursor(self, i: int, mixer: nn.Module) -> None:
         getattr(mixer, self._cursor_attr[i]).set_layer_idx(self._resource_idx[i])
 
 
+class Qwen3_5MTPHead(nn.Module):
+    """The checkpoints' one-layer multi-token-prediction head (``mtp.*``).
+
+    Position ``p`` takes the hidden state the lm_head reads at ``p`` (after the
+    final norm) and the embedding of the token at ``p + 1``, and predicts the
+    token at ``p + 2``: ``fc([norm_e(embed) | norm_h(hidden)])``, one
+    full-attention layer with its own KV (the cache's last layer), and a final
+    norm whose output feeds the shared lm_head and, chained, the next position.
+    """
+
+    def __init__(self, config: Qwen3_5Config, kv_layer_idx: int, comm_group: CommGroup):
+        super().__init__()
+        self.config = config
+        self.kv_layer_idx = kv_layer_idx
+        self.comm_group = comm_group
+        hidden = config.hidden_size
+        # replicated: its input is the replicated hidden state
+        self.fc = nn.Linear(2 * hidden, hidden, bias=False)
+        self.pre_fc_norm_embedding = _norm(config)
+        self.pre_fc_norm_hidden = _norm(config)
+        self.rope_cache = RopeCache()
+        self.layers = nn.ModuleList([_build_full_attn_layer(config, self.rope_cache, comm_group)])
+        self.norm = _norm(config)
+        layer = self.layers[0]
+        layer.self_attn.o_proj.reduce_results = False
+        layer.mlp.down_proj.reduce_results = False
+
+    def _project(self, hidden: torch.Tensor, next_embeds: torch.Tensor) -> torch.Tensor:
+        return self.fc(torch.cat(
+            (self.pre_fc_norm_embedding(next_embeds), self.pre_fc_norm_hidden(hidden)), dim=-1,
+        ))
+
+    def _tail(self, h: torch.Tensor, x: torch.Tensor) -> torch.Tensor:
+        """The layer after attention, and the final norm."""
+        cg = self.comm_group if self.comm_group.world_size > 1 else None
+        layer = self.layers[0]
+        h, residual = layer.post_attention_layernorm.forward_residual(h, x, cg)
+        h = layer.mlp(h)
+        normed, _ = self.norm.forward_residual(h, residual, cg)
+        return normed
+
+    def forward(
+        self, hidden: torch.Tensor, next_embeds: torch.Tensor, *, label: str,
+        cos_sin: tuple[torch.Tensor, torch.Tensor],
+    ) -> torch.Tensor:
+        """The head over a prompt's tokens, its KV written for each."""
+        x = self._project(hidden, next_embeds)
+        self.rope_cache.set(*cos_sin)
+        layer = self.layers[0]
+        attend = layer.self_attn.attend
+        attend.bind_step(label)
+        attend.set_layer_idx(self.kv_layer_idx)
+        return self._tail(layer.self_attn(layer.input_layernorm(x)), x)
+
+    def draft(
+        self, hidden: torch.Tensor, next_embeds: torch.Tensor, attend,
+        keep: torch.Tensor | None = None,
+    ) -> torch.Tensor:
+        """A draft pass: ``attend`` (see ``Qwen3_5Attention.forward_draft``)
+        runs one query a row whatever ``hidden`` holds, so the residual and
+        everything after run at ``[rows]``, ``keep`` picking that row where
+        ``hidden`` has several tokens a row."""
+        x = self._project(hidden, next_embeds)
+        layer = self.layers[0]
+        h = layer.self_attn.forward_draft(layer.input_layernorm(x), attend)
+        return self._tail(h, x if keep is None else x.index_select(0, keep))
+
+
 class Qwen3_5ForCausalLM(nn.Module):
     def __init__(
         self, config: Qwen3_5Config, comm_group: CommGroup | None = None,
+        with_mtp: bool = False,
     ):
         super().__init__()
         self.config = config
         self.model = Qwen3_5LanguageModel(config, comm_group)
+        # the draft for speculative decoding; its KV is one layer past the stack's
+        self.mtp = (
+            Qwen3_5MTPHead(
+                config, kv_layer_idx=len(config.full_layer_indices),
+                comm_group=self.model.comm_group,
+            )
+            if with_mtp else None
+        )
         # Gathers the vocab so the sampler is TP-oblivious. Its `[vocab/tp,
         # hidden]` shard matches the embedding's, so the tie works per rank.
         self.lm_head = ColumnParallelLinear(
@@ -305,6 +413,8 @@ def __init__(
         )
         if config.tie_word_embeddings:
             self.lm_head.weight = self.model.embed_tokens.weight
+        # what a draft pass reads in place of the lm_head, if not None
+        self.draft_head: nn.Module | None = None
 
     def forward(
         self, input_ids: torch.Tensor, *, label: str,
diff --git a/mstar/model/qwen3_5/config.py b/mstar/model/qwen3_5/config.py
index 9655099829..53448ef3ba 100644
--- a/mstar/model/qwen3_5/config.py
+++ b/mstar/model/qwen3_5/config.py
@@ -25,6 +25,8 @@
 LINEAR_ATTN = "linear_attn"
 ROPE = "rope"
 SAMPLER = "sampler"
+# speculative decoding's device registers, declared when the deployment drafts
+SPEC_DECODE = "spec_decode"
 # the ViT tower's cacheless attention, on the `vision_encoder` node
 VISION_ATTN = "vision_attn"
 
diff --git a/mstar/model/qwen3_5/qwen3_5_model.py b/mstar/model/qwen3_5/qwen3_5_model.py
index 27749aedef..3ff33cee5b 100644
--- a/mstar/model/qwen3_5/qwen3_5_model.py
+++ b/mstar/model/qwen3_5/qwen3_5_model.py
@@ -28,6 +28,7 @@
 from mstar.engine.resources.recurrent.config import DeltaNetGeometry, RecurrentStateConfig, RecurrentStateSpec
 from mstar.engine.resources.sampler.config import SamplerSpec, SamplingReqConfig
 from mstar.engine.resources.spec import CombinedSplit, NodeResourceSpec, ResourceReqConfig
+from mstar.engine.resources.spec_decode.config import SpecDecodeConfig, SpecDecodeSpec
 from mstar.graph.base import (
     GraphEdge,
     GraphNode,
@@ -70,6 +71,7 @@
     LLM_MIXED,
     ROPE,
     SAMPLER,
+    SPEC_DECODE,
     VISION_ATTN,
     Qwen3_5Config,
     Qwen3_5VisionConfig,
@@ -165,12 +167,18 @@ class Qwen3_5DenseModel(Model):
     # thread is steadier but slower (13 ms).
     PREPROCESS_TORCH_THREADS = 4
     INTERNAL_KWARGS = (IMAGE_GRIDS_KEY, VIDEO_GRIDS_KEY)
+    # drafts a speculative decode step verifies (see `__init__`); 0 decodes one
+    # token a step
+    num_draft_tokens: int = 0
 
     def __init__(
         self,
         model_path_hf: str,
         cache_dir: str | None = None,
         prefill_token_buckets: list[int] | None = None,
+        num_draft_tokens: int = 0,
+        spec_sampling: bool = False,
+        draft_head: str = "bf16",
         **kwargs: Any,
     ):
         self.model_path_hf = model_path_hf
@@ -180,6 +188,28 @@ def __init__(
 
         self.local_dir = _resolve_model_metadata(model_path_hf, cache_dir)
         self.config = Qwen3_5Config.from_hf(self.local_dir)
+        # Speculative decoding with the checkpoint's MTP head drafting this
+        # many tokens a step (`model_kwargs: {num_draft_tokens: k}` in the
+        # serving yaml); 0 decodes one token a step. Every process builds the
+        # model from the same kwargs, so the walks, resources and capture
+        # buckets sized from it agree everywhere.
+        if num_draft_tokens < 0:
+            raise ValueError(f"num_draft_tokens={num_draft_tokens} must be >= 0")
+        if num_draft_tokens and self.config.mtp_num_hidden_layers < 1:
+            raise ValueError(
+                f"{model_path_hf} has no MTP head (mtp_num_hidden_layers="
+                f"{self.config.mtp_num_hidden_layers}) to draft with; set "
+                "num_draft_tokens to 0"
+            )
+        self.num_draft_tokens = num_draft_tokens
+        # verify steps can also sample the target (`spec_sampling: true`),
+        # at the cost of a second decode capture
+        self.spec_sampling = bool(spec_sampling and num_draft_tokens)
+        # the vocabulary head a draft pass reads (`draft_head: fp8` keeps a copy
+        # in fp8, see `Fp8DraftHead`); verification always reads the checkpoint's
+        if draft_head not in ("bf16", "fp8"):
+            raise ValueError(f"draft_head={draft_head!r}: expected 'bf16' or 'fp8'")
+        self.draft_head = draft_head
         # None on a text-only checkpoint (`_create_submodule` then refuses a
         # vision_encoder node)
         self.vision_config = Qwen3_5VisionConfig.from_hf_or_none(self.local_dir)
@@ -220,7 +250,9 @@ def _gdn_state_dtype(self) -> torch.dtype:
         return torch.float32
 
     def get_node_resources(self) -> list[NodeResourceSpec]:
-        num_kv_layers = len(self.config.full_layer_indices)
+        k = self.num_draft_tokens
+        # the MTP head's attention keeps its KV as the cache's last layer
+        num_kv_layers = len(self.config.full_layer_indices) + (1 if k else 0)
         num_gdn_layers = len(self.config.linear_layer_indices)
         return [
             KVSpec(
@@ -256,7 +288,9 @@ def get_node_resources(self) -> list[NodeResourceSpec]:
                         head_v_dim=self.config.linear_value_head_dim,
                         conv_kernel_size=self.config.linear_conv_kernel_dim
                     ).to_blocks(state_dtype=self._gdn_state_dtype()),
-                    kv_cache=KV_CACHE,
+                    # no snapshots under speculation, where the label isn't keyed (`kv_labels`)
+                    kv_cache=None if k else KV_CACHE,
+                    spec_tokens=k,
                 )
             ),
             LinearAttnSpec(
@@ -264,7 +298,8 @@ def get_node_resources(self) -> list[NodeResourceSpec]:
                 config=LinearAttnConfig(
                     recurrent_state=GDN_STATE,
                     variant=LinearAttnVariant.GDN,
-                    splits=(_MIXED_SPLIT,),
+                    splits=() if k else (_MIXED_SPLIT,),
+                    spec_decode=SPEC_DECODE if k else None,
                 ),
             ),
             SamplerSpec(
@@ -272,9 +307,25 @@ def get_node_resources(self) -> list[NodeResourceSpec]:
                 vocab_size=self.config.vocab_size,
                 enable_repetion_penalty=True,
             ),
+            *self._spec_resources(),
             *self._vision_resources(),
         ]
 
+    def _spec_resources(self) -> list[NodeResourceSpec]:
+        """Speculative decoding's registers, only on a drafting deployment."""
+        if not self.num_draft_tokens:
+            return []
+        return [
+            SpecDecodeSpec(
+                resource_key=SPEC_DECODE, nodes={"LLM"},
+                config=SpecDecodeConfig(
+                    num_draft_tokens=self.num_draft_tokens, sampling=self.spec_sampling,
+                ),
+                kv_cache=KV_CACHE, attn=ATTN,
+                stop_token_ids=tuple(self.config.stop_token_ids),
+            ),
+        ]
+
     def _vision_resources(self) -> list[NodeResourceSpec]:
         """The ViT tower's attention, its only resource; none when text-only."""
         if self.vision_config is None:
@@ -308,17 +359,22 @@ def _vision_resources(self) -> list[NodeResourceSpec]:
     # Model ABC: walk graph declaration
     # -----------------------------------------------------------------------
     def get_combined_graph_walks(self) -> dict[str, dict[str, set[str]]]:
+        if self.num_draft_tokens:
+            # a verify row is k + 1 tokens, which the mixed step's split
+            # (one-token decode rows beside prompts) does not take
+            return {}
         return {"LLM": {LLM_MIXED: {"prefill_text", "decode"}}}
 
     def get_chunking_policies(self) -> dict[str, tuple[ChunkingPolicy, ...]]:
         from mstar.model.qwen3_5.submodules import LLMSubmodule
 
         budget = LLMSubmodule.MAX_BATCH_TOKENS
-        return {"LLM": (
-            ChunkingPolicy(walks=("prefill_text",), chunkable=True, max_batch_tokens=budget),
-            # carries the budget for the steps that mix a chunk with decode rows
-            ChunkingPolicy(walks=(LLM_MIXED,), max_batch_tokens=budget),
-        )}
+        policies = [ChunkingPolicy(walks=("prefill_text",), chunkable=True, max_batch_tokens=budget)]
+        if not self.num_draft_tokens:
+            # carries the budget for the steps that mix a chunk with decode rows;
+            # a speculative node has no mixed walk (`get_combined_graph_walks`)
+            policies.append(ChunkingPolicy(walks=(LLM_MIXED,), max_batch_tokens=budget))
+        return {"LLM": tuple(policies)}
 
     def get_graph_walk_graphs(self) -> dict[str, GraphSection]:
         prefill_text = GraphNode(
@@ -634,11 +690,14 @@ def media(step: PromptPart) -> tuple[Span]:
             return (Span("rows" if video else "digest", walk, rows, vision_position_advance(grid, merge) + 2, source),)
 
         plan = prefill_plan(expand_video_parts(parts, groups))
+        # a verify step writes up to 2k - 1 past the committed length, and the
+        # host counts up to k more until it learns what a step accepted
         layout = LabelLayout(
             spans=tuple(prompt_spans(plan, segments, walk, media)), growth=(("decode_loop", 1),),
             # The keys name `prefill_text` alone, and the preprocess worker refuses
             # a keyed span in any other walk.
             keyed_spans=0 if walk == "prefill_vision" else None,
+            scratch=3 * self.num_draft_tokens,
         )
 
         # always set, so a `video_groups` the client sent never reaches the
@@ -655,10 +714,10 @@ def get_loop_caps(self, **model_kwargs) -> dict[str, int]:
         return {"decode_loop": max(1, self.get_max_output_tokens(**model_kwargs) - 1)}
 
     def kv_labels(self) -> dict[str, dict[str, KVLabel]]:
-        return {KV_CACHE: {"main": KVLabel(
-            growth=(("decode", "decode_loop"),),
-            keys=PrefixKeys("text_inputs", ("prefill_text",), growth=True),
-        )}}
+        # Not keyed under speculation: the host's length runs ahead of what a
+        # verify step kept, so its pages can't be filed under a prompt's keys.
+        keys = None if self.num_draft_tokens else PrefixKeys("text_inputs", ("prefill_text",), growth=True)
+        return {KV_CACHE: {"main": KVLabel(growth=(("decode", "decode_loop"),), keys=keys)}}
 
     def checkpoint_path(self) -> str:
         return self._resolve_weights()
@@ -674,6 +733,9 @@ def postprocess(
         request_kwargs: dict | None = None,
     ) -> bytes:
         if modality == "text":
+            if self.num_draft_tokens:
+                # a verify step's row holds -1 past the tokens it emitted
+                output = output[output >= 0]
             return super().postprocess(output, modality, request_kwargs)
         raise ValueError(f"Unsupported modality for Qwen 3.5: {modality!r}")
 
@@ -907,11 +969,19 @@ def get_request_resource_configs(
         keys = [
             "temperature", "top_p", "top_k", "repetition_penalty", "ignore_eos", "return_logprobs", "logprobs_mode",
         ]
-        return {
-            SAMPLER: SamplingReqConfig(
-                **{k: model_kwargs[k] for k in keys if k in model_kwargs}
+        sampling = SamplingReqConfig(**{k: model_kwargs[k] for k in keys if k in model_kwargs})
+        if self.num_draft_tokens and (
+            (sampling.temperature and not self.spec_sampling)
+            or sampling.repetition_penalty not in (None, 1)
+        ):
+            # Verify steps take the target's argmax unless the deployment
+            # samples (`spec_sampling`), and track no tokens for a penalty.
+            raise ValueError(
+                "this deployment decodes speculatively; it takes temperature 0 "
+                "(or any, with spec_sampling) and no repetition_penalty. Serve it "
+                "without num_draft_tokens for those"
             )
-        }
+        return {SAMPLER: sampling}
 
     # ------------------------------------------------------------------
     # Model ABC: submodule loading
@@ -934,6 +1004,16 @@ def get_submodule(
         self._submodule_cache[node_name] = submodule
         return submodule
 
+    def validate_config_yaml(self, config: dict, config_path: str) -> None:
+        kv = (config.get("resources") or {}).get(KV_CACHE) or {}
+        limit = self.config.max_position_embeddings
+        if self.num_draft_tokens and kv.get("max_seq_len", 0) > limit:
+            # a draft pass rotates from a table of the checkpoint's positions
+            raise ValueError(
+                f"{config_path}: resources.{KV_CACHE}.max_seq_len={kv['max_seq_len']} is past the "
+                f"{limit} positions speculative decoding can address"
+            )
+
     def _create_submodule(
         self, node_name: str, device: str,
         autocast_dtype: torch.dtype | None = None,
@@ -973,17 +1053,27 @@ def _create_submodule(
         from mstar.model.qwen3_5.weight_loader import load_qwen3_5_weights
 
         model = self._build(
-            lambda: Qwen3_5ForCausalLM(self.config, tp_group), dtype, device,
+            lambda: Qwen3_5ForCausalLM(
+                self.config, tp_group, with_mtp=bool(self.num_draft_tokens),
+            ),
+            dtype, device,
         )
         load_qwen3_5_weights(model, weights_dir, device=device)
         model.requires_grad_(False).eval()
         if tp_group is not None and tp_group.world_size > 1:
             # collective over the TP group; every rank builds the LLM here
             tp_group.init_allreduce_fusion(self.config.hidden_size, dtype)
+        if self.num_draft_tokens and self.draft_head == "fp8":
+            from mstar.engine.resources.spec_decode.draft_head import Fp8DraftHead
+
+            if tp_group is not None and tp_group.world_size > 1:
+                raise ValueError("draft_head: fp8 reads one rank's whole head; serve it without TP")
+            model.draft_head = Fp8DraftHead(model.lm_head.weight)
         logger.info("Loaded Qwen3.5 LLM submodule onto %s", device)
         return LLMSubmodule(
             model, self.config, self.vision_config,
             prefill_token_buckets=self.prefill_token_buckets,
+            num_draft_tokens=self.num_draft_tokens, spec_sampling=self.spec_sampling,
         )
 
     @staticmethod
diff --git a/mstar/model/qwen3_5/submodules.py b/mstar/model/qwen3_5/submodules.py
index c7c7f84d84..bbf02d20dd 100644
--- a/mstar/model/qwen3_5/submodules.py
+++ b/mstar/model/qwen3_5/submodules.py
@@ -6,9 +6,11 @@
 counter, and the 3D ids are built here and passed in as cos/sin.
 """
 
+import functools
 import logging
 import os
 from collections.abc import Mapping
+from dataclasses import replace
 from time import perf_counter
 from typing import Any
 
@@ -37,6 +39,7 @@
 from mstar.engine.resources.sampler.config import SamplerStep
 from mstar.engine.resources.sampler.resource import SamplerResource
 from mstar.engine.resources.sampler.unfiltered import with_unfiltered_sampling
+from mstar.engine.resources.spec_decode.config import SpecDecodeStep, SpecPhase
 from mstar.engine.resources.step import Segment, SlotLease, SubmoduleStep
 from mstar.model.components.mrope import mrope_position_ids
 from mstar.model.components.qwen3_vit import (
@@ -63,6 +66,7 @@
     MIXED_MAX_PREFILL_ROWS,
     ROPE,
     SAMPLER,
+    SPEC_DECODE,
     VISION_ATTN,
     Qwen3_5Config,
     Qwen3_5VisionConfig,
@@ -73,6 +77,7 @@
     ARNodeSubmodule,
     BatchedModelOutput,
     DeviceLoopback,
+    HostRows,
     InputSeqLenInfo,
     ModelInputsFromEngine,
     NodeInputs,
@@ -102,6 +107,11 @@ def _mixed_token_buckets(prefill_buckets) -> list[int]:
 
 
 class LLMSubmodule(ARNodeSubmodule):
+    # drafts a speculative decode step verifies (see `__init__`); 0 decodes one
+    # token a step
+    num_draft_tokens: int = 0
+    spec_sampling: bool = False
+
     # Total tokens of a packed text-prefill (and mixed) step, one captured
     # graph per (bucket, batch size). A prompt past the top bucket prefills
     # eagerly, and an eager 8k prefill is launch-bound (~1,100 kernel
@@ -157,6 +167,8 @@ def __init__(
         config: Qwen3_5Config,
         vision_config: Qwen3_5VisionConfig | None = None,
         prefill_token_buckets: list[int] | None = None,
+        num_draft_tokens: int = 0,
+        spec_sampling: bool = False,
     ):
         super().__init__(prefill_token_buckets)
         if "MSTAR_QWEN35_PREFILL_TOKEN_BUCKETS" in os.environ:
@@ -169,12 +181,37 @@ def __init__(
         self.model = model
         self.config = config
         self.vision_config = vision_config
+        self._vision_sentinels: tuple[torch.Tensor, torch.Tensor] | None = None
+        # Speculative decoding: a decode step verifies this many drafts of the
+        # MTP head (`model.mtp`) and drafts the next ones, each row k + 1
+        # tokens; the drafts and lengths live in the SPEC_DECODE resource.
+        self.num_draft_tokens = num_draft_tokens
+        # verify steps may sample (a second decode capture, keyed by
+        # `cg_key_info`); off, they take the target's argmax
+        self.spec_sampling = spec_sampling and num_draft_tokens > 0
+        if num_draft_tokens and model.mtp is None:
+            raise ValueError("a speculative LLM submodule needs the model's MTP head")
         self.token_stop = token_stop(config.stop_token_ids)
-        # the decode walk reads its input ids off the sampler's slot masters
-        self.device_loopback = DeviceLoopback(
+        if self._speculative:
+            # a verify step emits a varying number of tokens a row: its stops
+            # come off its own flags (`check_stop_batched`), not the runtime's
+            # read of one token
+            self.token_stop = replace(self.token_stop, runtime_check=False)
+        # The decode walk reads its input ids off the sampler's slot masters.
+        # Not under speculation: the verify row's tokens are the speculative
+        # resource's, and inline emission carries one token a row.
+        self.device_loopback = None if self._speculative else DeviceLoopback(
             walk="decode", signal="text_inputs", sampler=SAMPLER, buffer="new_token",
         )
-        self._vision_sentinels: tuple[torch.Tensor, torch.Tensor] | None = None
+
+    @property
+    def _speculative(self) -> bool:
+        return self.num_draft_tokens > 0
+
+    @property
+    def _decode_tokens(self) -> int:
+        """Tokens of one decode row: 1, or a verify row's k + 1."""
+        return self.num_draft_tokens + 1
 
     # ------------------------------------------------------------------
     # Engine lifecycle
@@ -184,6 +221,10 @@ def get_accelerator_graph_configs(
         self, device: torch.device, tp_world_size: int = 1,
     ) -> list[AcceleratorGraphConfig]:
         """Decode, text prefill and vision prefill all capture."""
+        if self._speculative:
+            # built here, before any capture, so no replay rebuilds them
+            self.model.model.rope_tables(self.model.model.embed_tokens.weight.dtype)
+
         def dummy(n: int) -> ARNodeInputs:
             return ARNodeInputs(
                 input_ids=torch.zeros(n, dtype=torch.long, device=device),
@@ -215,14 +256,29 @@ def vision_dummy(n: int) -> ARNodeInputs:
         # the cost and that guarantee sit on this graph.
         decode_compile_mode = DETERMINISTIC_COMPILE_MODE if compile_deterministic() else None
 
-        # decode also captures once more for batches that filter nothing
-        return with_unfiltered_sampling([
-            BatchedAcceleratorGraphConfig(
+        mixed = [] if self._speculative else [
+            PackedAcceleratorGraphConfig(
+                capture_graph_walk=LLM_MIXED,
+                capture_token_lengths=self.MIXED_TOKEN_BUCKETS,
+                make_node_input=dummy,
+                capture_batch_sizes=self.MIXED_CAPTURE_BATCH_SIZES,
+                # the layout its GDN splits; a replay's padding rows stay empty
+                token_layout=decode_rows_then_prompt,
+                compile_mode=DETERMINISTIC_COMPILE_MODE,
+            ),
+        ]
+        # a speculative node that samples captures decode twice: greedy and sampled
+        decode_keys = (False, True) if self.spec_sampling else (None,)
+        configs = [
+            *(BatchedAcceleratorGraphConfig(
                 capture_graph_walk="decode",
-                single_request_inputs=dummy(1),
+                # a verify row under speculation
+                single_request_inputs=dummy(self._decode_tokens),
                 capture_batch_sizes=self.DECODE_CAPTURE_BATCH_SIZES,
                 compile_mode=decode_compile_mode,
-            ),
+                additional_key_info=key,
+                decode_rows=self._speculative,
+            ) for key in decode_keys),
             PackedAcceleratorGraphConfig(
                 capture_graph_walk="prefill_text",
                 capture_token_lengths=self.PREFILL_TOKEN_BUCKETS,
@@ -230,15 +286,7 @@ def vision_dummy(n: int) -> ARNodeInputs:
                 capture_batch_sizes=self.PREFILL_CAPTURE_BATCH_SIZES,
                 compile_mode=DETERMINISTIC_COMPILE_MODE,
             ),
-            PackedAcceleratorGraphConfig(
-                capture_graph_walk=LLM_MIXED,
-                capture_token_lengths=self.MIXED_TOKEN_BUCKETS,
-                make_node_input=dummy,
-                capture_batch_sizes=self.MIXED_CAPTURE_BATCH_SIZES,
-                # the layout its GDN splits; a replay's padding rows stay empty
-                token_layout=decode_rows_then_prompt,
-                compile_mode=DETERMINISTIC_COMPILE_MODE,
-            ),
+            *mixed,
             PackedAcceleratorGraphConfig(
                 capture_graph_walk="prefill_vision",
                 capture_token_lengths=self.PREFILL_VISION_TOKEN_BUCKETS,
@@ -246,17 +294,10 @@ def vision_dummy(n: int) -> ARNodeInputs:
                 capture_batch_sizes=self.PREFILL_VISION_CAPTURE_BATCH_SIZES,
                 compile_mode=DETERMINISTIC_COMPILE_MODE,
             ),
-        ], walks=("decode",))
-
-    def cg_key_info(
-        self, graph_walk: str,
-        per_request_info: dict[str, CurrentForwardPassInfo],
-        per_request_input_metadata=None,
-        **kwargs: Any,
-    ) -> Any:
-        if graph_walk != "decode":
-            return None
-        return self.node_resources[SAMPLER].unfiltered_key(graph_walk, per_request_info.keys())
+        ]
+        # a plain decode also captures once more for batches that filter nothing;
+        # a verify step doesn't draw through the sampler
+        return configs if self._speculative else with_unfiltered_sampling(configs, walks=("decode",))
 
     # ------------------------------------------------------------------
     # Step declaration
@@ -404,8 +445,18 @@ def prepare_inputs(
 
     def resident_walks(self) -> frozenset[str]:
         # a captured decode step builds its rope inside the graph on every config
-        # (`_rope_in_graph`), so its `preprocess` reads nothing the tables fill
-        return frozenset({"decode"})
+        # (`_rope_in_graph`), so its `preprocess` reads nothing the tables fill;
+        # a verify row is k + 1 tokens, not the one a resident row takes
+        return frozenset() if self._speculative else frozenset({"decode"})
+
+    def uniform_row_inputs(
+        self, graph_walk: str, has_loopback: bool = False,
+    ) -> ARNodeInputs | None:
+        if graph_walk == "decode" and self._speculative:
+            # the verify row's tokens are the speculative resource's, on the
+            # device; the routed signal is the previous step's output, unread
+            return ARNodeInputs(input_seq_len=self._decode_tokens)
+        return super().uniform_row_inputs(graph_walk, has_loopback)
 
     @staticmethod
     def _rope_in_graph(engine_inputs: ModelInputsFromEngine, graph_walk: str) -> bool:
@@ -471,6 +522,9 @@ def preprocess(
         inputs: list[ARNodeInputs],
     ) -> dict[str, torch.Tensor | Any]:
         out: dict[str, torch.Tensor | Any] = {}
+        if graph_walk == "decode" and self._speculative:
+            # everything a verify step reads is on the device already
+            return out
         in_graph = self._rope_in_graph(engine_inputs, graph_walk)
         sampler: SamplerResource = engine_inputs.resources[SAMPLER]
         if in_graph and graph_walk == "decode" and sampler.loopback_in_graph:
@@ -541,6 +595,13 @@ def declare_step(
         per_request_info: Mapping[int, CurrentForwardPassInfo] | None = None,
         **kwargs,
     ) -> SubmoduleStep:
+        if graph_walk == "decode" and self._speculative:
+            lease = kwargs.get("slot_lease")
+            key = (
+                lease.bucket.cg_key_info if lease is not None
+                else self.cg_key_info(graph_walk, per_request_info or {})
+            )
+            return self._declare_verify(request_ids, key)
         prefill_tokens = {}
         if graph_walk in ("prefill_text", LLM_MIXED):
             # a mixed step's decode rows track their token through the sampler as usual
@@ -576,6 +637,20 @@ def declare_step(
                 int(inp.tensor_inputs["mrope_advance"]) for inp in inputs
             )
 
+        steps = {
+            KV_CACHE: KVStep(),
+            ATTN: AttentionStep(causal=True),
+            GDN_STATE: RecurrentStep(),
+            LINEAR_ATTN: LinearAttnStep(),
+            SAMPLER: SamplerStep(
+                apply_penalty=True,
+                prefill_tracked_tokens=prefill_tokens,
+                kept_rids=kept_rids,
+            ),
+            ROPE: PositionStep(advance=advance),
+        }
+        if self._speculative:
+            self._declare_seed(steps, request_ids, per_request_info or {})
         # the leased bucket's key, which `cg_key_info` chose (or the capture's)
         lease = kwargs.get("slot_lease")
         return SubmoduleStep(
@@ -588,20 +663,87 @@ def declare_step(
                 )
                 for rid, inp in zip(request_ids, inputs, strict=True)
             ],
+            steps=steps,
+        )
+
+    def cg_key_info(
+        self, graph_walk: str,
+        per_request_info: Mapping[int, CurrentForwardPassInfo],
+        per_request_input_metadata=None,
+        **kwargs,
+    ):
+        """Which capture a decode batch takes. Plain decode: the unfiltered
+        draw's, when no row filters. Under speculation with sampling, whether
+        the batch samples: any row at a temperature above 0 takes the sampled
+        capture (its greedy rows still take the argmax there)."""
+        del per_request_input_metadata, kwargs
+        if graph_walk != "decode":
+            return None
+        if not self._speculative:
+            return self.node_resources[SAMPLER].unfiltered_key(graph_walk, per_request_info.keys())
+        if not self.spec_sampling:
+            return None
+        return any(
+            (cfg := info.resource_configs.get(SAMPLER)) is not None and cfg.temperature > 0
+            for info in per_request_info.values()
+        )
+
+    def _declare_verify(self, request_ids: list[int], key=None) -> SubmoduleStep:
+        """A speculative decode step: k + 1 tokens a row, positions and
+        lengths settled on the device. Each row gives back the KV its earlier
+        steps over-counted, as far as the host has learned."""
+        spec = self.node_resources[SPEC_DECODE]
+        span = self._decode_tokens
+        retract = {}
+        for rid in request_ids:
+            given = spec.pending_retract(rid)
+            if given:
+                retract[rid] = given
+        return SubmoduleStep(
+            segments=[Segment(request_id=rid, label="main", span=span) for rid in request_ids],
+            cg_key_info=key,
             steps={
-                KV_CACHE: KVStep(),
-                ATTN: AttentionStep(causal=True),
+                KV_CACHE: KVStep(commit_adjust={
+                    (rid, "main"): -given for rid, given in retract.items()
+                }),
+                ATTN: AttentionStep(causal=True, row_queries=span),
                 GDN_STATE: RecurrentStep(),
                 LINEAR_ATTN: LinearAttnStep(),
-                SAMPLER: SamplerStep(
-                    apply_penalty=True,
-                    prefill_tracked_tokens=prefill_tokens,
-                    kept_rids=kept_rids,
-                ),
-                ROPE: PositionStep(advance=advance),
+                SPEC_DECODE: SpecDecodeStep(phase=SpecPhase.VERIFY, retract=retract),
             },
         )
 
+    def _declare_seed(
+        self, steps: dict, request_ids: list[int],
+        info: Mapping[int, CurrentForwardPassInfo],
+    ) -> None:
+        """A prompt step's share of speculation: seed each row's registers,
+        and on the last prefill reserve the KV a verify step writes past the
+        committed length (`SpecDecodeResource.seed_lookahead`)."""
+        spec = self.node_resources[SPEC_DECODE]
+        budgets, ignore, adjust, sampling = {}, set(), {}, {}
+        for rid in request_ids:
+            row = info.get(rid)
+            if row is None:
+                continue
+            # the prompt's step emits the first token; decode the rest
+            budgets[rid] = max(row.max_tokens - 1, 0)
+            # a capture or padding row carries no sampling config
+            cfg = row.resource_configs.get(SAMPLER)
+            if cfg is not None and cfg.ignore_eos:
+                ignore.add(rid)
+            if cfg is not None and self.spec_sampling and cfg.temperature > 0:
+                sampling[rid] = (
+                    float(cfg.temperature), float(cfg.top_p), int(cfg.top_k or 0), int(cfg.seed),
+                )
+            if row.step_metadata.get("last_prefill", False) and spec.seed_lookahead:
+                adjust[(rid, "main")] = spec.seed_lookahead
+        steps[KV_CACHE] = KVStep(commit_adjust=adjust)
+        steps[SPEC_DECODE] = SpecDecodeStep(
+            phase=SpecPhase.SEED, budgets=budgets, ignore_stop=frozenset(ignore),
+            sampling=sampling,
+        )
+
     # ------------------------------------------------------------------
     # Forward
     # ------------------------------------------------------------------
@@ -615,6 +757,8 @@ def _forward(
         input_ids: torch.Tensor | None = None,
         input_embeds: torch.Tensor | None = None,
     ) -> torch.Tensor:
+        if graph_walk == "decode" and self._speculative:
+            return self._verify(engine_inputs)
         attn: AttentionManager = engine_inputs.resources[ATTN]
         sampler: SamplerResource = engine_inputs.resources[SAMPLER]
         if input_ids is None and input_embeds is None:
@@ -633,11 +777,62 @@ def _forward(
         )
         hidden = self.model.model(embeds, label="main", cos_sin=(cos_3d, sin_3d))
 
-        if graph_walk != "decode":
-            # only the last token of each prefill row feeds the head
-            hidden = attn.select_last_hidden(hidden)
-        logits = self.model.lm_head(hidden)
-        return sampler.sample(logits)
+        if graph_walk == "decode":
+            return sampler.sample(self.model.lm_head(hidden))
+        # only the last token of each prefill row feeds the head
+        token = sampler.sample(self.model.lm_head(attn.select_last_hidden(hidden)))
+        if self._speculative:
+            self._seed_drafts(engine_inputs, hidden, embeds, token, (cos_3d, sin_3d))
+        return token
+
+    def _cos_sin_1d(self, positions: torch.Tensor) -> tuple[torch.Tensor, torch.Tensor]:
+        """Rotary tables for text positions (the three MRoPE grids agree)."""
+        lm = self.model.model
+        return lm.build_cos_sin(
+            positions.unsqueeze(0).expand(3, -1), dtype=lm.embed_tokens.weight.dtype,
+        )
+
+    def _seed_drafts(
+        self, engine_inputs: ModelInputsFromEngine, hidden: torch.Tensor,
+        embeds: torch.Tensor, token: torch.Tensor,
+        cos_sin: tuple[torch.Tensor, torch.Tensor],
+    ) -> None:
+        """Run the MTP head over the prompt, which fills its KV for every
+        prompt position, and seed decode with the first token and the head's
+        draft after it. Position p reads the next token's embedding: the
+        prompt's own, and at each row's last the token just sampled."""
+        attn: AttentionManager = engine_inputs.resources[ATTN]
+        embed = self.model.model.embed_tokens
+        last = (attn.qo_indptr_buf("main")[1:] - 1).long()
+        following = embeds.roll(-1, 0).index_copy(0, last, embed(token).to(embeds.dtype))
+        drafted = self.model.mtp(hidden, following, label="main", cos_sin=cos_sin)
+        draft = self.model.lm_head(drafted.index_select(0, last)).argmax(dim=-1)
+        engine_inputs.resources[SPEC_DECODE].seed(token, draft)
+
+    def _verify(self, engine_inputs: ModelInputsFromEngine):
+        """A speculative decode step: verify each row's k drafts, then draft
+        the next k with the MTP head (see `SpecDecodeResource`)."""
+        spec = engine_inputs.resources[SPEC_DECODE]
+        lm, head = self.model.model, self.model.lm_head
+        ids, positions = spec.verify_inputs()
+        hidden = lm(lm.embed_tokens(ids), label="main", cos_sin=self._cos_sin_1d(positions))
+        step = engine_inputs.step
+        accepted = spec.accept(head(hidden), sample=bool(step is not None and step.cg_key_info))
+        ids = accepted.target.reshape(-1)
+        keep = spec.draft_slots()
+        draft_head = head if self.model.draft_head is None else self.model.draft_head
+        for j in range(self.num_draft_tokens):
+            attend = functools.partial(self._attend_draft, spec, j)
+            out = self.model.mtp.draft(hidden, lm.embed_tokens(ids), attend, keep=keep if j == 0 else None)
+            ids = spec.record_draft(j, draft_head(out))
+            hidden = out
+        return accepted
+
+    def _attend_draft(self, spec, j: int, qkv, q_weight, k_weight, *layout) -> tuple[torch.Tensor, torch.Tensor]:
+        """Draft pass ``j``'s attention, from the MTP layer's packed projection."""
+        lm = self.model.model
+        cos, sin = lm.rope_tables(lm.embed_tokens.weight.dtype)
+        return spec.attend_draft(j, qkv, q_weight, k_weight, cos, sin, *layout, self.model.mtp.kv_layer_idx)
 
     def forward(
         self,
@@ -649,12 +844,13 @@ def forward(
         input_embeds: torch.Tensor | None = None,
         **kwargs,
     ) -> NameToTensorList:
-        return {
-            "new_token": self._forward(
-                graph_walk, engine_inputs, cos_3d, sin_3d,
-                input_ids=input_ids, input_embeds=input_embeds,
-            )
-        }
+        out = self._forward(
+            graph_walk, engine_inputs, cos_3d, sin_3d,
+            input_ids=input_ids, input_embeds=input_embeds,
+        )
+        if graph_walk == "decode" and self._speculative:
+            out = out.tokens
+        return {"new_token": out}
 
     def can_batch(
         self, batch: ExecutingBatch, model_inputs: list[NodeInputs],
@@ -684,6 +880,17 @@ def forward_batched(
             graph_walk, engine_inputs, cos_3d, sin_3d,
             input_ids=input_ids, input_embeds=input_embeds,
         )
+        if graph_walk == "decode" and self._speculative:
+            # a row's tokens, -1 past what it emits; the stop check reads the
+            # accepted count (for the KV the host gives back) and `done`
+            return BatchedModelOutput(
+                row_outputs={"new_token": new_tokens.tokens},
+                check_stop_buffers={
+                    "new_token": new_tokens.tokens,
+                    "accepted": new_tokens.accepted,
+                    "done": new_tokens.done,
+                },
+            )
         # Row i is request i: one clone and one D2H copy per step instead of
         # one per request.
         return BatchedModelOutput(
@@ -720,6 +927,22 @@ def _token_is_real(request_info: CurrentForwardPassInfo) -> bool:
             "last_prefill", False,
         )
 
+    def check_stop(
+        self,
+        request_id: str,
+        request_info: CurrentForwardPassInfo,
+        outputs: dict[str, list[torch.Tensor]],
+    ):
+        if self.num_draft_tokens and request_info.graph_walk == "decode":
+            # The batched check reads the step's own `done`; this fallback
+            # sees only the tokens, so it stops on a stop id and leaves the
+            # budget to the loop's cap.
+            if "new_token" not in outputs or request_info.resource_configs[SAMPLER].ignore_eos:
+                return set()
+            tokens = outputs["new_token"][0].reshape(-1).tolist()
+            stop_ids = self.config.stop_token_ids
+            return {"decode_loop"} if any(t in stop_ids for t in tokens) else set()
+        return super().check_stop(request_id, request_info, outputs)
     def step_is_input_free(self, graph_walk: str) -> bool:
         # decode rows are one token each, like the template; the prefill walks
         # feed `prefill_tokens` from the inputs
@@ -727,9 +950,51 @@ def step_is_input_free(self, graph_walk: str) -> bool:
 
     def speculative_input_seq_len(self, graph_walk, fwd_info, prev_seq_len):
         if graph_walk == "decode":
-            return 1
+            # a verify row under speculation
+            return self._decode_tokens
         return super().speculative_input_seq_len(graph_walk, fwd_info, prev_seq_len)
 
+    def check_stop_batched(
+        self,
+        request_ids: list[str],
+        request_infos: dict[str, CurrentForwardPassInfo],
+        host_rows: HostRows,
+    ):
+        # a verify step's stops come off its own flags (no mixed step runs
+        # under speculation, so a step with a decode row has only those)
+        if self.num_draft_tokens and any(request_infos[rid].graph_walk == "decode" for rid in request_ids):
+            return self._check_stop_verify(request_ids, request_infos, host_rows)
+        return super().check_stop_batched(request_ids, request_infos, host_rows)
+
+
+    def _check_stop_verify(
+        self,
+        request_ids: list[int],
+        request_infos: dict[int, CurrentForwardPassInfo],
+        host_rows: HostRows,
+    ) -> dict[int, set[str]] | None:
+        """A verify step's stops, off the flags it computed itself; and what
+        each row accepted, which tells the speculative resource how much of
+        the KV the step reserved it can give back."""
+        accepted = host_rows.buffers.get("accepted")
+        done = host_rows.buffers.get("done")
+        if not (torch.is_tensor(accepted) and torch.is_tensor(done)):
+            return None
+        accepted, done = accepted.reshape(-1).tolist(), done.reshape(-1).tolist()
+        row_of = {rid: i for i, rid in enumerate(host_rows.request_ids)}
+        learned, counts = [], []
+        stops: dict[int, set[str]] = {}
+        for rid in request_ids:
+            i = row_of.get(rid)
+            if i is None or i >= len(accepted) or request_infos[rid].graph_walk != "decode":
+                continue
+            learned.append(rid)
+            counts.append(accepted[i])
+            if done[i]:
+                stops[rid] = {"decode_loop"}
+        self.node_resources[SPEC_DECODE].learn(learned, counts)
+        return stops
+
 
 # The vision tower's block loop, as a piecewise capture region
 VISION_BLOCK_LOOP = "vision_block_loop"
diff --git a/mstar/model/qwen3_5/weight_loader.py b/mstar/model/qwen3_5/weight_loader.py
index 6a2bc148ee..750e4bea76 100644
--- a/mstar/model/qwen3_5/weight_loader.py
+++ b/mstar/model/qwen3_5/weight_loader.py
@@ -4,7 +4,8 @@
 net mixer is ``linear_attn`` (ours: ``self_attn``); vision tensors lose their
 ``model.visual.`` prefix. ``_STACKED_PARAMS`` routes the separate q/k/v, gate/up
 and delta net input projections by shard id into their fused parameters. The two
-towers load separately, as separate nodes. The MTP head (``mtp.*``) is skipped.
+towers load separately, as separate nodes. The MTP head (``mtp.*``) loads only
+into a model built with it (a speculative deployment).
 """
 from __future__ import annotations
 
@@ -34,6 +35,8 @@
 _VISION_PREFIX = "model.visual."
 # at the checkpoint's root, and only when untied (9B and 27B)
 _LM_HEAD = "lm_head.weight"
+# the multi-token-prediction head, loaded only for a speculative deployment
+_MTP_PREFIX = "mtp."
 
 
 def qwen3_5_name_remapper(name: str) -> str | None:
@@ -47,6 +50,14 @@ def qwen3_5_name_remapper(name: str) -> str | None:
     )
 
 
+def qwen3_5_mtp_name_remapper(name: str) -> str | None:
+    """The text tower's names, and the MTP head's (``mtp.*``, the same in the
+    checkpoint and in ours) for a model that drafts with it."""
+    if name.startswith(_MTP_PREFIX):
+        return name
+    return qwen3_5_name_remapper(name)
+
+
 def qwen3_5_vision_name_remapper(name: str) -> str | None:
     if not name.startswith(_VISION_PREFIX):
         return None
@@ -95,8 +106,12 @@ def load_qwen3_5_weights(
     if not tied:
         # an untied head sits at the root, outside the text prefix
         selectors.append({"keys": {_LM_HEAD}})
+    remapper = qwen3_5_name_remapper
+    if getattr(model, "mtp", None) is not None:
+        selectors.append({"prefix": _MTP_PREFIX})
+        remapper = qwen3_5_mtp_name_remapper
     return _load(
-        model, path, device, qwen3_5_name_remapper, selectors,
+        model, path, device, remapper, selectors,
         expected={
             n for n, _ in model.named_parameters()
             # tied: the checkpoint carries no tensor for it
diff --git a/mstar/platforms/compile.py b/mstar/platforms/compile.py
index a2de694090..a3bc7d1ec0 100644
--- a/mstar/platforms/compile.py
+++ b/mstar/platforms/compile.py
@@ -13,9 +13,10 @@
 
 def with_inductor_overrides(
     fn: Callable[..., Any], num_tokens: int | None = None, rows: int | None = None,
+    decode_rows: bool = False,
 ) -> Callable[..., Any]:
     """Wrap ``fn``, a ``torch.compile``'d callable, so it runs under the
-    platform's ``inductor_overrides(num_tokens, rows)``.
+    platform's ``inductor_overrides(num_tokens, rows, decode_rows)``.
 
     torch.compile compiles on call, not when wrapped, so the config is patched
     around every call. With no overrides ``fn`` is returned as is and costs
@@ -24,7 +25,7 @@ def with_inductor_overrides(
     This resolves ``current_platform``, so call it once the process has its
     device, not at import time.
     """
-    overrides = current_platform.inductor_overrides(num_tokens, rows)
+    overrides = current_platform.inductor_overrides(num_tokens, rows, decode_rows)
     if not overrides:
         return fn
 
diff --git a/mstar/platforms/interface.py b/mstar/platforms/interface.py
index f1ce5b5369..1ffc3c1331 100644
--- a/mstar/platforms/interface.py
+++ b/mstar/platforms/interface.py
@@ -114,10 +114,13 @@ def device_facts(self, device=None) -> DeviceFacts:
 
     def inductor_overrides(
         self, num_tokens: int | None = None, rows: int | None = None,
+        decode_rows: bool = False,
     ) -> dict[str, object]:
         """``torch._inductor.config`` values to patch in while the engine
         compiles a forward for ``num_tokens`` tokens in ``rows`` rows, a
         capture bucket's counts, or ``None`` when the forward serves any count.
+        ``decode_rows`` marks a bucket of several tokens a row that still
+        compiles as decode (a speculative verify step).
 
         Meant for inductor codegen settings on this stack (say, which GEMM
         backends to autotune at small sizes), not for picking kernels. Empty
@@ -281,6 +284,7 @@ def fp8_dtype(self, device=None) -> torch.dtype | None:
 
     def inductor_overrides(
         self, num_tokens: int | None = None, rows: int | None = None,
+        decode_rows: bool = False,
     ) -> dict[str, object]:
         # At decode sizes inductor's autotuner can't reliably pick between its
         # Triton GEMM templates and hipBLASLt on MI300X. Its timings (12-17 us)
@@ -294,9 +298,12 @@ def inductor_overrides(
         # templates, which win there.
         if num_tokens is None:
             return {}
-        decode = rows is not None and rows == num_tokens
+        decode = rows is not None and (rows == num_tokens or decode_rows)
         limit = self.ATEN_GEMM_MAX_DECODE_ROWS if decode else self.ATEN_GEMM_MAX_TOKENS
-        if num_tokens > limit:
+        # a verify bucket's draft GEMMs run at its rows, its target's at a few
+        # times that, so its rows bound the shapes these choices are about
+        size = rows if decode_rows else num_tokens
+        if size > limit:
             return {}
         overrides: dict[str, object] = {"max_autotune_gemm_backends": "ATEN"}
         from mstar.platforms.rocm_gemm import (
@@ -309,12 +316,12 @@ def inductor_overrides(
         )
 
         post = None
-        if num_tokens <= DECODE_GEMV_MAX_ROWS and (passes := skinny_gemm_passes()) is not None:
+        if size <= DECODE_GEMV_MAX_ROWS and (passes := skinny_gemm_passes()) is not None:
             # linears of up to 16 rows on AITER's skinny GEMM or the Triton
             # weight stream, as the device's tuned config picks per shape
             overrides["pre_grad_custom_pass"], post = passes
         if (
-            decode and SKINNY_MAX_ROWS < num_tokens <= DECODE_LINEAR_MAX_ROWS
+            decode and SKINNY_MAX_ROWS < num_tokens and size <= DECODE_LINEAR_MAX_ROWS
             and (tuned := decode_linear_pass()) is not None
         ):
             # decode linears still on hipBLASLt take its tuned solution, where the device has one
diff --git a/mstar/utils/causal_conv1d/__init__.py b/mstar/utils/causal_conv1d/__init__.py
index a54a14a477..6ee919bb7d 100644
--- a/mstar/utils/causal_conv1d/__init__.py
+++ b/mstar/utils/causal_conv1d/__init__.py
@@ -7,8 +7,9 @@
 """
 from mstar.utils.causal_conv1d.kernels import (
     PAD_SLOT_ID,
+    causal_conv1d_decode,
     causal_conv1d_fn,
     causal_conv1d_update,
 )
 
-__all__ = ["PAD_SLOT_ID", "causal_conv1d_fn", "causal_conv1d_update"]
+__all__ = ["PAD_SLOT_ID", "causal_conv1d_decode", "causal_conv1d_fn", "causal_conv1d_update"]
diff --git a/mstar/utils/causal_conv1d/kernels.py b/mstar/utils/causal_conv1d/kernels.py
index 9129f868a0..e6e50c10a3 100644
--- a/mstar/utils/causal_conv1d/kernels.py
+++ b/mstar/utils/causal_conv1d/kernels.py
@@ -1131,6 +1131,84 @@ def _causal_conv1d_update_kernel(
         tl.store(o_ptrs, acc, mask=mask_1d)
 
 
+@triton.jit
+def _causal_conv1d_decode_kernel(
+    x_ptr, w_ptr, bias_ptr, state_ptr, slots_ptr,
+    dim, stride_x_row, stride_w_dim, stride_w_width, stride_state_slot, stride_state_dim, stride_state_tok,
+    stride_slots, null_slot,
+    HAS_BIAS: tl.constexpr, SILU_ACTIVATION: tl.constexpr, BLOCK_N: tl.constexpr,
+):
+    row = tl.program_id(0)
+    cols = tl.program_id(1) * BLOCK_N + tl.arange(0, BLOCK_N)
+    slot = tl.load(slots_ptr + row * stride_slots).to(tl.int64)
+    if slot == null_slot:
+        return
+    mask = cols < dim
+    state = state_ptr + slot * stride_state_slot + cols * stride_state_dim
+    col0 = tl.load(state, mask=mask, other=0.0)
+    col1 = tl.load(state + stride_state_tok, mask=mask, other=0.0)
+    col2 = tl.load(state + 2 * stride_state_tok, mask=mask, other=0.0)
+    x_ptrs = x_ptr + row * stride_x_row + cols
+    x = tl.load(x_ptrs, mask=mask, other=0.0)
+    w = w_ptr + cols * stride_w_dim
+    if HAS_BIAS:
+        acc = tl.load(bias_ptr + cols, mask=mask, other=0.0).to(tl.float32)
+    else:
+        acc = tl.zeros((BLOCK_N,), dtype=tl.float32)
+    acc += col0 * tl.load(w, mask=mask, other=0.0)
+    acc += col1 * tl.load(w + stride_w_width, mask=mask, other=0.0)
+    acc += col2 * tl.load(w + 2 * stride_w_width, mask=mask, other=0.0)
+    acc += x * tl.load(w + 3 * stride_w_width, mask=mask, other=0.0)
+    tl.store(state, col1, mask=mask)
+    tl.store(state + stride_state_tok, col2, mask=mask)
+    tl.store(state + 2 * stride_state_tok, x, mask=mask)
+    if SILU_ACTIVATION:
+        acc = acc / (1 + tl.exp(-acc))
+    tl.store(x_ptrs, acc, mask=mask)
+
+
+def decode_kernel_fits(x: torch.Tensor, conv_state: torch.Tensor, weight: torch.Tensor) -> bool:
+    """Whether `causal_conv1d_decode` has its own kernel for these operands."""
+    return (
+        x.dim() == 2 and x.stride(1) == 1 and x.dtype == conv_state.dtype
+        and weight.shape[1] == 4 and conv_state.shape[2] == 3
+    )
+
+
+def causal_conv1d_decode(
+    x: torch.Tensor,
+    conv_state: torch.Tensor,
+    weight: torch.Tensor,
+    bias: torch.Tensor | None = None,
+    activation: bool | str | None = None,
+    conv_state_indices: torch.Tensor | None = None,
+    null_block_id: int = NULL_BLOCK_ID,
+) -> torch.Tensor:
+    """`causal_conv1d_update` for a decode step: one token a row, a width-4
+    filter, the output in place over ``x``, and bit-equal results.
+
+    The general kernel loads the state twice and branches on what a decode step
+    never has, which on MI325X costs more than a launch per call at 64 rows.
+    Other operands go to the general kernel.
+    """
+    if not decode_kernel_fits(x, conv_state, weight):
+        return causal_conv1d_update(
+            x, conv_state, weight, bias=bias, activation=activation,
+            conv_state_indices=conv_state_indices, null_block_id=null_block_id,
+        )
+    rows, dim = x.shape
+    block_n = 256
+    _causal_conv1d_decode_kernel[(rows, triton.cdiv(dim, block_n))](
+        x, weight, bias, conv_state, conv_state_indices,
+        dim, x.stride(0), weight.stride(0), weight.stride(1),
+        conv_state.stride(0), conv_state.stride(1), conv_state.stride(2),
+        conv_state_indices.stride(0), null_block_id,
+        HAS_BIAS=bias is not None, SILU_ACTIVATION=activation in (True, "silu", "swish"), BLOCK_N=block_n,
+        num_warps=4,
+    )
+    return x
+
+
 def causal_conv1d_update(
     x: torch.Tensor,
     conv_state: torch.Tensor,
diff --git a/mstar/utils/compile_bundle.py b/mstar/utils/compile_bundle.py
new file mode 100644
index 0000000000..fcc537371e
--- /dev/null
+++ b/mstar/utils/compile_bundle.py
@@ -0,0 +1,469 @@
+"""Compile-cache bundles: the Inductor, Triton and MIOpen caches a warmed-up
+server leaves behind, packed per software stack so that a new host or
+container with the same stack starts warm instead of compiling every capture
+bucket again.
+
+A bundle is a directory named by its stack key, under a root that may hold
+bundles for several stacks::
+
+    <root>/<product>/torch-<torch>-<hip|cuda>-<version>-triton-<triton>/
+        manifest.json        the key, what each tar holds, and who added to it
+        inductor.tar         TORCHINDUCTOR_CACHE_DIR, less a Triton dir inside it
+        triton.tar           TRITON_CACHE_DIR
+        miopen.tar           MIOPEN_USER_DB_PATH (ROCm)
+        miopen-kernels.tar   MIOPEN_CUSTOM_CACHE_DIR, when it is not the same dir
+
+Each tar holds its directory's files as ``./<path>`` and is uncompressed, so
+packing into an existing bundle appends only the files it lacks or holds an
+older copy of, and unpacking keeps the last copy of each file. The layout and
+the matching rules are in docs/compile_cache.rst.
+"""
+
+from __future__ import annotations
+
+import contextlib
+import datetime
+import fcntl
+import importlib
+import json
+import logging
+import os
+import re
+import shutil
+import socket
+import sqlite3
+import subprocess
+import tarfile
+import tempfile
+from collections.abc import Callable, Iterator
+from dataclasses import asdict, dataclass
+from pathlib import Path
+
+logger = logging.getLogger(__name__)
+
+MANIFEST = "manifest.json"
+# written into the Inductor dir by unpack and read at server startup
+MARKER = "mstar-bundle.json"
+LOCK = ".append.lock"
+KINDS = ("inductor", "triton", "miopen", "miopen-kernels")
+_KEY_DIR = re.compile(r"torch-(?P<torch>.+)-(?P<runtime>(?:hip|cuda)-[^-]+)-triton-(?P<triton>[^/]+)")
+
+
+class BundleMismatchError(RuntimeError):
+    """No bundle under the given path was built for this host's stack."""
+
+
+@dataclass(frozen=True)
+class StackKey:
+    """The stack a bundle's entries were compiled on.
+
+    Inductor and Triton key their own entries on the GPU arch and the torch,
+    HIP or CUDA, and Triton builds, so a bundle from another stack unpacks
+    harmlessly but is never read: those fields must match. The product is in
+    the key because autotuning picks are timings taken on it. An MI300X's
+    picks load on an MI325X (both gfx942) but were not tuned there.
+    """
+
+    product: str
+    arch: str | None
+    torch: str
+    runtime: str  # "hip-7.2.53211" or "cuda-12.8"
+    triton: str
+
+    @property
+    def path(self) -> str:
+        return f"{self.product}/torch-{self.torch}-{self.runtime}-triton-{self.triton}"
+
+    @classmethod
+    def from_path(cls, product: str, rest: str) -> StackKey | None:
+        """The key of a bundle directory ``<product>/<rest>`` that has no
+        manifest; such a directory doesn't record its arch."""
+        m = _KEY_DIR.fullmatch(rest)
+        if m is None:
+            return None
+        return cls(product=product, arch=None, **m.groupdict())
+
+    def differences(self, other: StackKey) -> list[str]:
+        """What keeps ``other``'s entries from being read on this stack. A field
+        either side doesn't know (a manifest-less bundle's arch) is no reason."""
+        out = []
+        for name in ("arch", "torch", "runtime", "triton"):
+            here, there = getattr(self, name), getattr(other, name)
+            if here and there and here != there:
+                out.append(f"{name} {there} (here {here})")
+        return out
+
+
+def stack_key(device: int | None = None) -> StackKey:
+    """This process's key, from ``device`` (the current GPU by default)."""
+    import torch
+    import triton
+
+    from mstar.platforms import current_platform
+
+    if not torch.cuda.is_available():
+        raise RuntimeError("no CUDA or ROCm GPU is visible; a bundle is keyed by the GPU that compiled it")
+    name = torch.cuda.get_device_name(device)
+    product = re.sub(r"[^a-z0-9]+", "-", name.lower().replace("amd instinct", "")).strip("-")
+    runtime = f"hip-{torch.version.hip}" if torch.version.hip else f"cuda-{torch.version.cuda}"
+    return StackKey(product=product, arch=current_platform.arch(device), torch=torch.__version__,
+                    runtime=runtime, triton=triton.__version__)
+
+
+def cache_dirs() -> dict[str, Path]:
+    """Each kind's directory as a server started from this environment
+    resolves it: an exported variable, else M*'s default for Inductor and
+    Triton (set when :mod:`mstar.engine` is imported), else the library's own."""
+    # Importing torch._inductor writes torch's default into the environment, so M*'s goes in first.
+    importlib.import_module("mstar.engine")
+    from torch._inductor.runtime.cache_dir_utils import cache_dir
+
+    from mstar.platforms import current_platform
+
+    dirs = {
+        "inductor": Path(cache_dir()),
+        "triton": Path(os.environ.get("TRITON_CACHE_DIR") or Path.home() / ".triton" / "cache"),
+    }
+    if current_platform.is_rocm():
+        db = Path(os.environ.get("MIOPEN_USER_DB_PATH") or "~/.config/miopen").expanduser()
+        kernels = Path(os.environ.get("MIOPEN_CUSTOM_CACHE_DIR") or "~/.cache/miopen").expanduser()
+        dirs["miopen"] = db
+        if kernels.resolve() != db.resolve():
+            dirs["miopen-kernels"] = kernels
+    return dirs
+
+
+# --- reading and writing tars -------------------------------------------------------------------
+
+def _stamp(st: os.stat_result, link: bool) -> tuple[int, int]:
+    # what a tar member records of a file: a symlink's size is 0 there
+    return (0 if link else st.st_size, int(st.st_mtime))
+
+
+def _files(root: Path, skip: list[Path]) -> dict[str, tuple[int, int]]:
+    """Files and symlinks under ``root`` as tar names (``./a/b``), less the
+    marker and anything under ``skip``."""
+    skip_set = {p.resolve() for p in skip}
+    out: dict[str, tuple[int, int]] = {}
+    for top, dirnames, filenames in os.walk(root):
+        here = Path(top)
+        dirnames[:] = [d for d in dirnames if (here / d).resolve() not in skip_set]
+        for d in dirnames:
+            if (here / d).is_symlink():
+                rel = (here / d).relative_to(root).as_posix()
+                out[f"./{rel}"] = _stamp(os.lstat(here / d), True)
+        for f in filenames:
+            p = here / f
+            rel = p.relative_to(root).as_posix()
+            if rel == MARKER:
+                continue
+            st = os.lstat(p)
+            out[f"./{rel}"] = _stamp(st, p.is_symlink())
+    return out
+
+
+def _members(tar: Path) -> dict[str, tuple[int, int]]:
+    """A tar's file and symlink members by name; the last copy of a name wins,
+    as it does on extraction."""
+    out: dict[str, tuple[int, int]] = {}
+    with tarfile.open(tar, "r:") as tf:
+        while (m := tf.next()) is not None:
+            # tarfile keeps every header it reads, and a bundle's tar can hold a million
+            tf.members.clear()
+            if m.isfile() or m.issym():
+                name = m.name if m.name.startswith("./") else f"./{m.name}"
+                out[name] = (0 if m.issym() else m.size, int(m.mtime))
+    return out
+
+
+def _add(tar: Path, src: Path, names: list[str]) -> None:
+    """Append ``names`` (relative to ``src``) to ``tar``, or create it. A copy is
+    written and renamed into place, so a reader never sees a partial tar."""
+    tmp = tar.with_name(f"{tar.name}.{os.getpid()}")
+    with tempfile.NamedTemporaryFile("wb", suffix=".list") as listing:
+        listing.write(b"\0".join(n.encode() for n in names) + b"\0")
+        listing.flush()
+        try:
+            if tar.exists():
+                shutil.copyfile(tar, tmp)
+            subprocess.run(["tar", "-rf" if tar.exists() else "-cf", str(tmp), "-C", str(src),
+                            "--no-recursion", "--null", "-T", listing.name],
+                           check=True, capture_output=True, text=True)
+            os.replace(tmp, tar)
+        except subprocess.CalledProcessError as exc:
+            tmp.unlink(missing_ok=True)
+            raise RuntimeError(
+                f"tar failed adding {len(names)} files from {src} to {tar}: {exc.stderr.strip()}") from exc
+        except BaseException:
+            tmp.unlink(missing_ok=True)
+            raise
+
+
+_SQLITE = b"SQLite format 3\0"
+
+
+def _merge_db(old: Path, new: Path, out: Path) -> bool:
+    """Write ``new`` with the records of ``old`` it lacks to ``out``: a text
+    database (MIOpen's find and perf DBs, ``key=value`` lines) or SQLite (its
+    kernel DB). Where both have a record, ``new``'s wins. False when the result
+    holds nothing ``old`` doesn't; True for a file that is neither."""
+    head = new.read_bytes()[:16]
+    if head == _SQLITE:
+        shutil.copyfile(new, out)
+        con = sqlite3.connect(out)
+        try:
+            con.execute("ATTACH DATABASE ? AS old", (str(old),))
+            grew = False
+            tables = [r[0] for r in con.execute("SELECT name FROM old.sqlite_master WHERE type = 'table'")]
+            for table in tables:
+                # an INTEGER PRIMARY KEY is a row id, renumbered on insert; the
+                # tables' unique indexes say which rows are the same record
+                cols = [r[1] for r in con.execute(f"PRAGMA old.table_info(`{table}`)")
+                        if not (r[5] and r[2].upper() == "INTEGER")]
+                names = ", ".join(f"`{c}`" for c in cols)
+                con.execute(f"INSERT OR IGNORE INTO main.`{table}` ({names}) SELECT {names} FROM old.`{table}`")
+                n_old = con.execute(f"SELECT count(*) FROM old.`{table}`").fetchone()[0]
+                grew |= con.execute(f"SELECT count(*) FROM main.`{table}`").fetchone()[0] > n_old
+            con.commit()
+        finally:
+            con.close()
+        return grew
+    if new.name.endswith(".txt"):
+        records: dict[str, str] = {}
+        for line in old.read_text().splitlines() + new.read_text().splitlines():
+            key, sep, _ = line.partition("=")
+            if sep:
+                records[key] = line
+        merged = "".join(f"{line}\n" for line in records.values())
+        out.write_text(merged)
+        return merged != old.read_text()
+    shutil.copyfile(new, out)
+    return True
+
+
+def _merge_miopen(tar: Path, src: Path, names: list[str], have: dict, tmp: Path) -> list[str]:
+    """Stage the files of ``names`` (under ``src``) that differ from the
+    bundle's last copy under ``tmp/merged`` for appending to a MIOpen tar, each
+    database merged with that copy. MIOpen keeps one file per database, so a
+    cache warmed without the bundle's records would otherwise replace them.
+    Returns the names staged."""
+    if have:
+        subprocess.run(["tar", "-xf", str(tar), "-C", str(tmp)], check=True, capture_output=True)
+    keep = []
+    for name in names:
+        new, old, out = src / name, tmp / name, tmp / "merged" / name
+        out.parent.mkdir(parents=True, exist_ok=True)
+        if new.is_symlink() or name not in have or not old.is_file():
+            shutil.copy2(new, out, follow_symlinks=False)
+            keep.append(name)
+        elif new.read_bytes() != old.read_bytes() and _merge_db(old, new, out):
+            shutil.copystat(new, out)
+            keep.append(name)
+    return keep
+
+
+@contextlib.contextmanager
+def _locked(path: Path) -> Iterator[None]:
+    """One writer per bundle at a time, across hosts sharing the filesystem
+    when its locks work (they do on local disks and NFSv4)."""
+    with open(path, "a") as f:
+        fcntl.flock(f, fcntl.LOCK_EX)
+        try:
+            yield
+        finally:
+            fcntl.flock(f, fcntl.LOCK_UN)
+
+
+def _write_json(path: Path, data: dict) -> None:
+    tmp = path.with_name(f"{path.name}.{os.getpid()}")
+    tmp.write_text(json.dumps(data, indent=1) + "\n")
+    os.replace(tmp, path)
+
+
+def _now() -> str:
+    return datetime.datetime.now(datetime.timezone.utc).isoformat(timespec="seconds")
+
+
+# --- pack, unpack, find -------------------------------------------------------------------------
+
+@dataclass(frozen=True)
+class KindPack:
+    kind: str
+    source: str
+    added: int
+    members: int
+    tar_bytes: int
+
+
+def pack(root: Path, key: StackKey, dirs: dict[str, Path], kinds: list[str] | None = None,
+         note: str | None = None, log: Callable[[str], None] = print) -> tuple[Path, list[KindPack]]:
+    """Add the caches in ``dirs`` (``kinds`` of them, default all) to ``key``'s
+    bundle under ``root``: new files, and files whose size or mtime differ from
+    the bundle's last copy (Inductor and Triton entries never change; MIOpen
+    rewrites its databases in place, and a changed one is merged with the
+    bundle's copy). Run it with no server writing to the caches."""
+    bundle = root / key.path
+    bundle.mkdir(parents=True, exist_ok=True)
+    out = []
+    with _locked(bundle / LOCK):
+        manifest = read_manifest(bundle) or {"key": asdict(key), "created": _now(), "kinds": {}, "history": []}
+        for kind, src in dirs.items():
+            if kinds is not None and kind not in kinds:
+                continue
+            if not src.is_dir():
+                log(f"{kind}: {src} does not exist, left out")
+                continue
+            # a Triton dir inside Inductor's (M*'s default) goes in its own tar
+            inner = [d for k, d in dirs.items() if k != kind and d.resolve().is_relative_to(src.resolve())
+                     and d.resolve() != src.resolve()]
+            files = _files(src, inner)
+            tar = bundle / f"{kind}.tar"
+            have = _members(tar) if tar.exists() else {}
+            if kind.startswith("miopen"):
+                # a few small files, compared by content: two warm-ups can leave
+                # databases of the same size and mtime
+                with tempfile.TemporaryDirectory(dir=bundle) as tmp:
+                    new = _merge_miopen(tar, src, sorted(files), have, Path(tmp))
+                    if new:
+                        _add(tar, Path(tmp) / "merged", new)
+            else:
+                new = sorted(n for n, stamp in files.items() if have.get(n) != stamp)
+                if new:
+                    _add(tar, src, new)
+            members = len(have.keys() | set(new))
+            result = KindPack(kind, str(src), len(new), members, tar.stat().st_size if tar.exists() else 0)
+            out.append(result)
+            manifest["kinds"][kind] = {"members": result.members, "tar_bytes": result.tar_bytes}
+            log(f"{kind}: {result.added} files added from {src}, {result.members} in {tar.name} "
+                f"({result.tar_bytes / 2**30:.2f} GiB)")
+        manifest["updated"] = _now()
+        manifest["history"].append({
+            "time": manifest["updated"], "host": socket.gethostname(), "note": note,
+            "added": {r.kind: r.added for r in out}, "mstar": _mstar_commit(),
+        })
+        _write_json(bundle / MANIFEST, manifest)
+    return bundle, out
+
+
+def read_manifest(bundle: Path) -> dict | None:
+    try:
+        return json.loads((bundle / MANIFEST).read_text())
+    except FileNotFoundError:
+        return None
+
+
+def bundle_key(bundle: Path) -> StackKey | None:
+    """A bundle's key: its manifest's, else read off its path."""
+    manifest = read_manifest(bundle)
+    if manifest is not None:
+        return StackKey(**manifest["key"])
+    return StackKey.from_path(bundle.parent.name, bundle.name)
+
+
+def is_bundle(path: Path) -> bool:
+    return (path / MANIFEST).is_file() or any((path / f"{k}.tar").is_file() for k in KINDS)
+
+
+def find_bundles(root: Path) -> list[Path]:
+    """``root`` itself if it is a bundle, else the bundles in its
+    ``<product>/<stack>/`` directories."""
+    if is_bundle(root):
+        return [root]
+    return sorted(p for p in root.glob("*/*") if p.is_dir() and is_bundle(p))
+
+
+def select_bundle(src: Path, key: StackKey) -> Path:
+    """The bundle in ``src`` for ``key``: ``src`` itself when it is a bundle,
+    else ``src/<key path>``. Raises BundleMismatchError naming what ``src`` holds."""
+    if is_bundle(src):
+        return src
+    bundle = src / key.path
+    if is_bundle(bundle):
+        return bundle
+    have = [str(p.relative_to(src)) for p in find_bundles(src)]
+    raise BundleMismatchError(f"{src} has no bundle for {key.path}; it has: {', '.join(have) or 'none'}")
+
+
+@dataclass(frozen=True)
+class Unpacked:
+    bundle: Path
+    key: StackKey | None
+    kinds: dict[str, str]  # kind -> directory it went to
+
+
+def unpack(src: Path, key: StackKey, dirs: dict[str, Path], force: bool = False,
+           log: Callable[[str], None] = print) -> Unpacked:
+    """Extract the bundle for ``key`` from ``src`` into ``dirs``, over what is
+    there (the bundle's copy of a file wins), and leave a marker for the
+    startup line. A bundle for another stack is refused unless ``force``."""
+    bundle = select_bundle(src, key)
+    bkey = bundle_key(bundle)
+    diff = key.differences(bkey) if bkey is not None else []
+    if diff and not force:
+        raise BundleMismatchError(
+            f"{bundle} was built for {bkey.path}: {', '.join(diff)}; its entries would not be read")
+    done: dict[str, str] = {}
+    for kind in KINDS:
+        tar = bundle / f"{kind}.tar"
+        if not tar.is_file():
+            continue
+        # MIOpen's kernel cache shares the database dir unless it has one of its own here
+        dest = dirs.get(kind) or (dirs.get("miopen") if kind == "miopen-kernels" else None)
+        if dest is None:
+            log(f"{kind}: no {kind} cache on this platform, left out")
+            continue
+        dest.mkdir(parents=True, exist_ok=True)
+        p = subprocess.run(["tar", "-xf", str(tar), "-C", str(dest), "--no-same-owner"],
+                           capture_output=True, text=True, check=False)
+        if p.returncode != 0:
+            raise RuntimeError(f"tar failed extracting {tar} into {dest}: {p.stderr.strip()}")
+        done[kind] = str(dest)
+    dirs["inductor"].mkdir(parents=True, exist_ok=True)
+    _write_json(dirs["inductor"] / MARKER, {
+        "bundle": str(bundle), "key": asdict(bkey) if bkey else None, "unpacked": _now(), "kinds": done,
+    })
+    return Unpacked(bundle, bkey, done)
+
+
+def startup_line(key: StackKey, dirs: dict[str, Path]) -> tuple[bool, str]:
+    """What the server logs at startup: this stack's key and whether the last
+    bundle unpacked into its caches matches it. False when it doesn't."""
+    try:
+        info = json.loads((dirs["inductor"] / MARKER).read_text())
+    except FileNotFoundError:
+        return True, f"compile cache: stack {key.path}; no bundle unpacked into {dirs['inductor']}"
+    except (OSError, ValueError) as exc:
+        return False, f"compile cache: stack {key.path}; unreadable {MARKER} in {dirs['inductor']}: {exc}"
+    head = f"compile cache: stack {key.path}; bundle {info.get('bundle')} unpacked {info.get('unpacked')}"
+    if not info.get("key"):
+        return True, f"{head}, stack unknown"
+    bkey = StackKey(**info["key"])
+    diff = key.differences(bkey)
+    if diff:
+        return False, f"{head} was built for {bkey.path}: {', '.join(diff)}, so its entries are not used"
+    if bkey.product != key.product:
+        return True, f"{head} matches the stack, tuned on {bkey.product}"
+    return True, f"{head} matches"
+
+
+def log_startup(who: str, device: int | None = None) -> None:
+    """Log :func:`startup_line` for ``device``, as a warning when the bundle
+    doesn't match. Never fails a startup."""
+    try:
+        ok, line = startup_line(stack_key(device), cache_dirs())
+    except Exception as exc:
+        logger.warning("%s: compile cache: could not read this stack's key: %s", who, exc)
+        return
+    (logger.info if ok else logger.warning)("%s: %s", who, line)
+
+
+def _mstar_commit() -> str | None:
+    repo = Path(__file__).resolve().parents[2]
+    try:
+        p = subprocess.run(["git", "-C", str(repo), "rev-parse", "--short", "HEAD"],
+                           capture_output=True, text=True, timeout=10, check=False)
+    except (OSError, subprocess.TimeoutExpired):
+        return None
+    if p.returncode != 0:
+        return None
+    return p.stdout.strip() or None
diff --git a/mstar/worker/worker.py b/mstar/worker/worker.py
index 4d004e11c9..a5859a9479 100644
--- a/mstar/worker/worker.py
+++ b/mstar/worker/worker.py
@@ -72,7 +72,7 @@
 from mstar.profile.worker import WorkerProfileInfo
 from mstar.streaming.stream_buffer import StreamingEdge
 from mstar.streaming.topology import Connection, ProducerWalkCtx
-from mstar.utils import cpus, knobs
+from mstar.utils import compile_bundle, cpus, knobs
 from mstar.utils.client_errors import client_error_text
 from mstar.utils.containers import ParallelList, RecentSet
 from mstar.utils.device import DeviceCapability
@@ -569,6 +569,8 @@ def __init__(
         if self.device_capability.is_accelerator and self.device.index is not None:
             torch.accelerator.set_device_index(self.device)
         self._plan_step_cpus(worker_ids, local_worker_devices)
+        if self.device.type == "cuda":
+            compile_bundle.log_startup(f"Worker {worker_id}", self.device.index)
 
         # ``dist_init_method`` is normally provided by the conductor — it
         # picks a free TCP port at startup so multiple ``mstar`` runs on
diff --git a/test/modular/test_accelerator_graph_capture.py b/test/modular/test_accelerator_graph_capture.py
index 03c10357da..c1868e8e68 100644
--- a/test/modular/test_accelerator_graph_capture.py
+++ b/test/modular/test_accelerator_graph_capture.py
@@ -659,3 +659,44 @@ def test_geometric_batch_sizes(max_batch_size, expected):
 def test_geometric_batch_sizes_refuses_a_max_under_one():
     with pytest.raises(ValueError, match="at least 1"):
         geometric_batch_sizes(0)
+@pytest.mark.parametrize("dynamic, graphs", [(False, 6), (True, 2)])
+def test_compile_dynamic_compiles_one_graph_for_every_bucket(monkeypatch, dynamic, graphs):
+    """A ``compile_dynamic`` config's buckets share one compiled graph (size 1
+    gets its own), compiled under the platform's overrides for any count;
+    otherwise every bucket compiles, under its own count's overrides."""
+    from torch._dynamo.testing import CompileCounter
+
+    from mstar.engine import accelerator_graph_runner as agr
+    from mstar.engine.accelerator_graph_config import BatchedAcceleratorGraphConfig
+    from mstar.model.submodule_base import NodeInputs
+
+    torch._dynamo.reset()
+    counter = CompileCounter()
+    real_compile = torch.compile
+    # the counting backend takes neither a mode nor inductor options
+    monkeypatch.setattr(torch, "compile", lambda fn, **kw: real_compile(
+        fn, backend=counter, fullgraph=kw["fullgraph"], dynamic=kw["dynamic"]))
+    overrides = []
+    monkeypatch.setattr(agr, "with_inductor_overrides",
+                        lambda fn, num_tokens=None, rows=None, decode_rows=False: overrides.append((num_tokens, rows))
+                        or fn)
+
+    def encode(graph_walk, engine_inputs, feats):
+        return torch.nn.functional.gelu(feats @ feats.transpose(1, 2)).sum(-1)
+
+    runner = AcceleratorGraphRunner.__new__(AcceleratorGraphRunner)
+    runner._submodule = SimpleNamespace(forward_batched=encode)
+    runner._compiled_forwards = {}
+    config = BatchedAcceleratorGraphConfig(
+        "encode", single_request_inputs=NodeInputs(tensor_inputs={}, input_seq_len=1), compile_dynamic=dynamic,
+    )
+    sizes = [32, 16, 8, 4, 2, 1]
+    for bs in sizes:  # largest first, as the runner captures
+        static = {"feats": torch.randn(bs, 15, 8)}
+        if dynamic:
+            agr.mark_bucket_dims_dynamic(static, {(0, "feats"): 0}, config_idx=0, num_tokens=bs)
+        spec = SimpleNamespace(config=config, config_idx=0, num_tokens=bs, bs=bs)
+        runner._forward_for(spec)(graph_walk="encode", engine_inputs=None, **static)
+    torch._dynamo.reset()
+    assert counter.frame_count == graphs
+    assert overrides == ([(None, None)] * len(sizes) if dynamic else [(bs, bs) for bs in sizes])
diff --git a/test/modular/test_attn_prologue.py b/test/modular/test_attn_prologue.py
new file mode 100644
index 0000000000..982431d6aa
--- /dev/null
+++ b/test/modular/test_attn_prologue.py
@@ -0,0 +1,73 @@
+"""The packed attention prologue (``attn.prologue``): its reference does what a
+Qwen3.5 layer does unfused, and the rotary tables it reads agree with the
+step's. CPU only."""
+
+from __future__ import annotations
+
+import torch
+
+from mstar.engine.resources.attn import prologue
+from mstar.model.qwen3_5.components.rope import apply_partial_mrope, compute_3d_cos_sin, compute_inv_freq
+
+T, H, KH, D, HALF = 5, 4, 2, 32, 4
+
+
+def _args(gated=True):
+    shape = {"tokens": T, "heads": H, "kv_heads": KH, "head_dim": D, "half": HALF}
+    args = list(prologue._inputs(shape, torch.float32, "cpu", torch.Generator().manual_seed(0)))
+    if not gated:
+        args[0] = args[0][:, H * D:].contiguous()
+        args[-1] = False
+    return args
+
+
+def _unfused(qkv, qw, kw, cos, sin, pos, gated):
+    q_size = H * D * (2 if gated else 1)
+    q, k, v = qkv.split([q_size, KH * D, KH * D], -1)
+    gate = None
+    if gated:
+        q, gate = q.reshape(T, H, 2 * D).chunk(2, -1)
+    else:
+        q = q.reshape(T, H, D)
+
+    def norm(x, w):
+        return x * torch.rsqrt(x.pow(2).mean(-1, keepdim=True) + 1e-6) * (1 + w)
+
+    c, s = torch.cat([cos[pos], cos[pos]], -1), torch.cat([sin[pos], sin[pos]], -1)
+    q, k = apply_partial_mrope(norm(q, qw), norm(k.reshape(T, KH, D), kw), c, s)
+    return q, gate, k, v.reshape(T, KH, D)
+
+
+def test_the_reference_is_the_layers_unfused_math():
+    for gated in (True, False):
+        args = _args(gated)
+        q, gate, k, v = prologue.QKV_NORM_ROPE_CACHE(*args)
+        qkv, qw, kw, cos, sin, pos, key_cache, value_cache, slots = args[:9]
+        eq, egate, ek, ev = _unfused(qkv, qw, kw, cos, sin, pos, gated)
+        torch.testing.assert_close(q, eq, rtol=0, atol=0)
+        torch.testing.assert_close(k, ek, rtol=0, atol=0)
+        torch.testing.assert_close(v, ev, rtol=0, atol=0)
+        if gated:
+            torch.testing.assert_close(gate, egate, rtol=0, atol=0)
+        else:
+            assert gate.numel() == 0
+        block = key_cache.shape[1]
+        torch.testing.assert_close(key_cache[slots // block, slots % block], ek, rtol=0, atol=0)
+        torch.testing.assert_close(value_cache[slots // block, slots % block], ev, rtol=0, atol=0)
+        assert v.data_ptr() != qkv.data_ptr()
+
+
+def test_no_implementation_off_amd_gpus():
+    assert prologue.QKV_NORM_ROPE_CACHE.unavailable_reason() is not None
+
+
+def test_the_rope_tables_are_the_steps_tables():
+    """One position grid per token: the three MRoPE grids agree, so the 1D
+    table rounds to what ``build_cos_sin`` gives a decode step."""
+    inv_freq = compute_inv_freq(64, 1e7)
+    pos = torch.tensor([0, 1, 17, 4095])
+    cos, sin = compute_3d_cos_sin(pos[None].expand(3, -1), inv_freq, [11, 11, 10], torch.bfloat16)
+    freqs = inv_freq[None, :].float() * torch.arange(4096).float()[:, None]
+    table_cos, table_sin = freqs.cos().to(torch.bfloat16), freqs.sin().to(torch.bfloat16)
+    assert torch.equal(cos[:, :32], table_cos[pos]) and torch.equal(cos[:, 32:], table_cos[pos])
+    assert torch.equal(sin[:, :32], table_sin[pos])
diff --git a/test/modular/test_compile_bundle.py b/test/modular/test_compile_bundle.py
new file mode 100644
index 0000000000..29d387ad70
--- /dev/null
+++ b/test/modular/test_compile_bundle.py
@@ -0,0 +1,230 @@
+"""Compile-cache bundles (mstar/utils/compile_bundle.py, ``mstar cache``): packing
+appends only what is new, unpacking restores the caches and marks them, a
+bundle for another stack is refused, and a bundle root laid out like the
+golden copies the AMD benchmark scripts seed from (uncompressed per-kind tars
+of ``./`` members, no manifest) reads as a bundle."""
+
+from __future__ import annotations
+
+import json
+import os
+import subprocess
+import tarfile
+from pathlib import Path
+
+import pytest
+
+from mstar.cli import cache as cache_cli
+from mstar.cli.main import build_parser
+from mstar.utils import compile_bundle as cb
+
+KEY = cb.StackKey(product="mi325x-vf", arch="gfx942", torch="2.12.1+rocm7.2", runtime="hip-7.2.53211",
+                  triton="3.7.1")
+
+
+def _write(path: Path, data: str, mtime: int = 1_700_000_000) -> None:
+    path.parent.mkdir(parents=True, exist_ok=True)
+    path.write_text(data)
+    os.utime(path, (mtime, mtime))
+
+
+def _caches(root: Path) -> dict[str, Path]:
+    """M*'s default layout: Triton's dir inside Inductor's; MIOpen beside them."""
+    dirs = {"inductor": root / "inductor", "triton": root / "inductor" / "triton", "miopen": root / "miopen"}
+    _write(dirs["inductor"] / "fxgraph" / "ab" / "entry", "graph")
+    _write(dirs["inductor"] / "wq" / "kernel.py", "code")
+    _write(dirs["triton"] / "HASH" / "k.hsaco", "binary")
+    _write(dirs["miopen"] / "gfx942130.ukdb", "kernels")
+    _write(dirs["miopen"] / "gfx942130.HIP.ufdb.txt", "conv-a=Solver1:0.5\n")
+    return dirs
+
+
+def _names(tar: Path) -> list[str]:
+    with tarfile.open(tar) as tf:
+        return [m.name for m in tf.getmembers() if m.isfile()]
+
+
+def _tree(root: Path) -> dict[str, str]:
+    return {p.relative_to(root).as_posix(): p.read_text() for p in sorted(root.rglob("*")) if p.is_file()}
+
+
+def test_key_path_round_trips_and_names_what_differs():
+    assert KEY.path == "mi325x-vf/torch-2.12.1+rocm7.2-hip-7.2.53211-triton-3.7.1"
+    parsed = cb.StackKey.from_path(*KEY.path.split("/"))
+    assert parsed == cb.StackKey(**{**KEY.__dict__, "arch": None})
+    assert KEY.differences(parsed) == []  # a manifest-less bundle doesn't know its arch
+    other = cb.StackKey("mi300x", "gfx942", KEY.torch, KEY.runtime, "3.6.0")
+    assert KEY.differences(other) == ["triton 3.6.0 (here 3.7.1)"]  # the product alone is no reason
+    assert cb.StackKey.from_path("x", "not-a-stack") is None
+
+
+def test_pack_then_unpack_restores_the_caches(tmp_path):
+    dirs = _caches(tmp_path / "built")
+    bundle, results = cb.pack(tmp_path / "bundles", KEY, dirs, note="warm-up", log=lambda _: None)
+    assert bundle == tmp_path / "bundles" / KEY.path
+    assert {r.kind: r.added for r in results} == {"inductor": 2, "triton": 1, "miopen": 2}
+    # Triton's dir inside Inductor's goes in its own tar only
+    assert sorted(_names(bundle / "inductor.tar")) == ["./fxgraph/ab/entry", "./wq/kernel.py"]
+    manifest = json.loads((bundle / "manifest.json").read_text())
+    assert manifest["key"]["arch"] == "gfx942" and manifest["history"][0]["note"] == "warm-up"
+
+    fresh = {"inductor": tmp_path / "new" / "i", "triton": tmp_path / "new" / "t", "miopen": tmp_path / "new" / "m"}
+    done = cb.unpack(tmp_path / "bundles", KEY, fresh, log=lambda _: None)
+    assert done.bundle == bundle and set(done.kinds) == {"inductor", "triton", "miopen"}
+    assert _tree(fresh["triton"]) == _tree(dirs["triton"])
+    assert _tree(fresh["miopen"]) == _tree(dirs["miopen"])
+    restored = _tree(fresh["inductor"])
+    marker = json.loads(restored.pop(cb.MARKER))
+    assert marker["bundle"] == str(bundle) and marker["key"]["triton"] == "3.7.1"
+    assert restored == {k: v for k, v in _tree(dirs["inductor"]).items() if not k.startswith("triton/")}
+    assert os.stat(fresh["triton"] / "HASH" / "k.hsaco").st_mtime == 1_700_000_000
+
+
+def test_packing_inductor_alone_still_leaves_its_triton_dir_out(tmp_path):
+    dirs = _caches(tmp_path / "built")
+    bundle, results = cb.pack(tmp_path / "b", KEY, dirs, kinds=["inductor"], log=lambda _: None)
+    assert [r.kind for r in results] == ["inductor"]
+    assert sorted(_names(bundle / "inductor.tar")) == ["./fxgraph/ab/entry", "./wq/kernel.py"]
+
+
+def test_a_second_pack_appends_only_new_and_changed_files(tmp_path):
+    dirs = _caches(tmp_path / "built")
+    root = tmp_path / "bundles"
+    bundle, _ = cb.pack(root, KEY, dirs, log=lambda _: None)
+    _, again = cb.pack(root, KEY, dirs, log=lambda _: None)
+    assert all(r.added == 0 for r in again)
+    _write(dirs["inductor"] / "fxgraph" / "cd" / "entry", "another graph")
+    _write(dirs["miopen"] / "gfx942130.HIP.ufdb.txt", "conv-a=Solver2:0.4\n", mtime=1_700_000_100)
+    _, third = cb.pack(root, KEY, dirs, log=lambda _: None)
+    assert {r.kind: (r.added, r.members) for r in third} == {"inductor": (1, 3), "triton": (0, 1), "miopen": (1, 2)}
+    assert len(_names(bundle / "miopen.tar")) == 3  # appended, not rewritten
+    out = {"inductor": tmp_path / "o" / "i", "triton": tmp_path / "o" / "t", "miopen": tmp_path / "o" / "m"}
+    cb.unpack(bundle, KEY, out, log=lambda _: None)
+    assert (out["miopen"] / "gfx942130.HIP.ufdb.txt").read_text() == "conv-a=Solver2:0.4\n"  # this host's record
+    assert len(json.loads((bundle / "manifest.json").read_text())["history"]) == 3
+
+
+def _kernel_db(path: Path, rows: list[tuple[str, str]]) -> None:
+    """MIOpen's kernel DB schema: a row id and a unique (name, args)."""
+    import sqlite3
+
+    path.parent.mkdir(parents=True, exist_ok=True)
+    con = sqlite3.connect(path)
+    con.execute("CREATE TABLE IF NOT EXISTS `kern_db` (`id` INTEGER PRIMARY KEY ASC, `kernel_name` TEXT NOT NULL, "
+                "`kernel_args` TEXT NOT NULL, `kernel_blob` BLOB NOT NULL)")
+    con.execute("CREATE UNIQUE INDEX IF NOT EXISTS `idx_kern_db` ON kern_db(kernel_name, kernel_args)")
+    con.executemany("INSERT INTO kern_db (kernel_name, kernel_args, kernel_blob) VALUES (?, ?, ?)",
+                    [(n, a, b"bin") for n, a in rows])
+    con.commit()
+    con.close()
+
+
+def test_miopen_databases_from_separate_warm_ups_are_merged(tmp_path):
+    """Two hosts warmed different models from empty caches: MIOpen's find DB and
+    kernel DB are one file each, so packing the second must keep the first's
+    records rather than replace the file."""
+    import sqlite3
+
+    first, second = tmp_path / "a" / "miopen", tmp_path / "b" / "miopen"
+    _write(first / "x.ufdb.txt", "conv-a=S1:0.5\nconv-b=S1:0.7\n")
+    _kernel_db(first / "x.ukdb", [("k1", "-O3"), ("k2", "-O3")])
+    _write(second / "x.ufdb.txt", "conv-b=S2:0.6\nconv-c=S1:0.2\n", mtime=1_700_000_100)
+    _kernel_db(second / "x.ukdb", [("k2", "-O3"), ("k3", "-O3")])
+    cb.pack(tmp_path / "bundles", KEY, {"miopen": first}, log=lambda _: None)
+    _, res = cb.pack(tmp_path / "bundles", KEY, {"miopen": second}, log=lambda _: None)
+    assert res[0].added == 2
+    out = {"inductor": tmp_path / "o" / "i", "miopen": tmp_path / "o" / "m"}
+    cb.unpack(tmp_path / "bundles", KEY, out, log=lambda _: None)
+    assert (out["miopen"] / "x.ufdb.txt").read_text() == "conv-a=S1:0.5\nconv-b=S2:0.6\nconv-c=S1:0.2\n"
+    con = sqlite3.connect(out["miopen"] / "x.ukdb")
+    assert sorted(r[0] for r in con.execute("SELECT kernel_name FROM kern_db")) == ["k1", "k2", "k3"]
+    con.close()
+    # packing the merged result again adds nothing
+    _, res = cb.pack(tmp_path / "bundles", KEY, {"miopen": out["miopen"]}, log=lambda _: None)
+    assert res[0].added == 0
+
+
+def test_a_bundle_for_another_stack_is_refused_unless_forced(tmp_path):
+    dirs = _caches(tmp_path / "built")
+    other = cb.StackKey(KEY.product, KEY.arch, "2.11.0+rocm7.1", KEY.runtime, KEY.triton)
+    bundle, _ = cb.pack(tmp_path / "bundles", other, dirs, log=lambda _: None)
+    out = {"inductor": tmp_path / "o" / "i", "triton": tmp_path / "o" / "t"}
+    with pytest.raises(cb.BundleMismatchError, match="has no bundle for"):
+        cb.unpack(tmp_path / "bundles", KEY, out, log=lambda _: None)
+    with pytest.raises(cb.BundleMismatchError, match="torch 2.11.0"):
+        cb.unpack(bundle, KEY, out, log=lambda _: None)
+    cb.unpack(bundle, KEY, out, force=True, log=lambda _: None)
+    ok, line = cb.startup_line(KEY, out)
+    assert not ok and "torch 2.11.0+rocm7.1 (here 2.12.1+rocm7.2)" in line
+
+
+def test_a_golden_copy_layout_reads_as_a_bundle(tmp_path):
+    """$WORK/.cache/compile-golden/<product>/<stack>/{inductor,triton,miopen}.tar,
+    made with ``tar -cf X.tar -C dir .``: no manifest, a ``./`` directory member."""
+    built = _caches(tmp_path / "built")
+    golden = tmp_path / "golden" / KEY.path
+    golden.mkdir(parents=True)
+    for kind in ("inductor", "triton", "miopen"):
+        exclude = ["--exclude=./triton"] if kind == "inductor" else []
+        subprocess.run(["tar", "-cf", str(golden / f"{kind}.tar"), *exclude, "-C", str(built[kind]), "."],
+                       check=True)
+    assert cb.find_bundles(tmp_path / "golden") == [golden]
+    assert cb.bundle_key(golden) == cb.StackKey(**{**KEY.__dict__, "arch": None})
+    out = {"inductor": tmp_path / "o" / "i", "triton": tmp_path / "o" / "t", "miopen": tmp_path / "o" / "m"}
+    cb.unpack(tmp_path / "golden", KEY, out, log=lambda _: None)
+    assert _tree(out["miopen"]) == _tree(built["miopen"])
+    # packing into it appends to the same tars and adds a manifest
+    _write(built["triton"] / "NEW" / "k.hsaco", "binary 2")
+    _, res = cb.pack(tmp_path / "golden", KEY, built, log=lambda _: None)
+    assert {r.kind: r.added for r in res} == {"inductor": 0, "triton": 1, "miopen": 0}
+    assert (golden / "manifest.json").is_file()
+
+
+def test_startup_line_says_whether_the_unpacked_bundle_matches(tmp_path):
+    dirs = _caches(tmp_path / "built")
+    ok, line = cb.startup_line(KEY, dirs)
+    assert ok and "no bundle unpacked" in line and KEY.path in line
+    cb.pack(tmp_path / "b", KEY, dirs, log=lambda _: None)
+    cb.unpack(tmp_path / "b", KEY, dirs, log=lambda _: None)
+    ok, line = cb.startup_line(KEY, dirs)
+    assert ok and line.endswith("matches")
+    elsewhere = cb.StackKey("mi300x", KEY.arch, KEY.torch, KEY.runtime, KEY.triton)
+    ok, line = cb.startup_line(elsewhere, dirs)
+    assert ok and line.endswith("tuned on mi325x-vf")
+
+
+def test_miopen_kernels_without_a_dir_of_their_own_go_to_the_db_dir(tmp_path):
+    built = {"inductor": tmp_path / "i", "miopen": tmp_path / "db", "miopen-kernels": tmp_path / "kc"}
+    _write(built["miopen"] / "x.ufdb.txt", "db")
+    _write(built["miopen-kernels"] / "3.5.1" / "x.ukdb", "kernels")
+    bundle, _ = cb.pack(tmp_path / "b", KEY, built, log=lambda _: None)
+    assert _names(bundle / "miopen-kernels.tar") == ["./3.5.1/x.ukdb"]
+    shared = {"inductor": tmp_path / "o" / "i", "miopen": tmp_path / "o" / "m"}
+    cb.unpack(bundle, KEY, shared, log=lambda _: None)
+    assert (shared["miopen"] / "3.5.1" / "x.ukdb").read_text() == "kernels"
+
+
+def test_cli_wires_cache_subcommands(monkeypatch, tmp_path, capsys):
+    calls = {}
+    monkeypatch.setattr(cb, "stack_key", lambda device=None: KEY)
+    monkeypatch.setattr(cb, "cache_dirs", lambda: {"inductor": tmp_path / "i", "triton": tmp_path / "t"})
+
+    def pack(root, key, dirs, kinds=None, note=None):
+        calls["pack"] = (root, dirs, kinds, note)
+        return root, []
+
+    monkeypatch.setattr(cb, "pack", pack)
+    args = build_parser().parse_args(["cache", "pack", str(tmp_path / "r"), "--kinds", "triton", "--note", "n"])
+    args.func(args)
+    assert calls["pack"] == (tmp_path / "r", {"inductor": tmp_path / "i", "triton": tmp_path / "t"}, ["triton"], "n")
+
+    def refuse(*a, **k):
+        raise cb.BundleMismatchError("nothing for this stack")
+
+    monkeypatch.setattr(cb, "unpack", refuse)
+    args = build_parser().parse_args(["cache", "unpack", str(tmp_path)])
+    with pytest.raises(SystemExit) as exc:
+        args.func(args)
+    assert exc.value.code == cache_cli.EXIT_NO_MATCH
+    args = build_parser().parse_args(["serve", "qwen3_5_9b", "--compile-bundle", str(tmp_path)])
+    assert args.compile_bundle == str(tmp_path)
diff --git a/test/modular/test_conv_decode_route.py b/test/modular/test_conv_decode_route.py
new file mode 100644
index 0000000000..5f761cfc39
--- /dev/null
+++ b/test/modular/test_conv_decode_route.py
@@ -0,0 +1,39 @@
+"""Which conv kernel a decode step takes: AMD GPUs the decode kernel where it fits, every other platform today's."""
+
+from types import SimpleNamespace
+
+import pytest
+import torch
+
+from mstar.engine.resources.linear_attn import wrappers
+from mstar.engine.resources.linear_attn.wrappers import GDNDecodeWrapper
+from mstar.utils import causal_conv1d
+from mstar.utils.causal_conv1d.kernels import decode_kernel_fits
+
+
+def _wrapper():
+    w = GDNDecodeWrapper.__new__(GDNDecodeWrapper)
+    w._plan_state = wrappers.GDNDecodePlan(slots=torch.tensor([1, 2]), num_rows=2)
+    w._null_slot_id = 0
+    return w
+
+
+@pytest.mark.parametrize(("rocm", "expected"), [(True, "decode"), (False, "update")])
+def test_route(monkeypatch, rocm, expected):
+    called = []
+    # the module's reference, not an attribute of the shared proxy, which undoing would leave pinned
+    monkeypatch.setattr(wrappers, "current_platform", SimpleNamespace(is_rocm=lambda: rocm))
+    monkeypatch.setattr(causal_conv1d, "causal_conv1d_decode", lambda **kw: called.append("decode"))
+    monkeypatch.setattr(causal_conv1d, "causal_conv1d_update", lambda **kw: called.append("update"))
+    _wrapper().run_conv(torch.zeros(2, 8), torch.zeros(3, 8, 3), torch.zeros(8, 4))
+    assert called == [expected]
+
+
+def test_fits():
+    x, state, weight = torch.zeros(2, 8), torch.zeros(3, 8, 3), torch.zeros(8, 4)
+    assert decode_kernel_fits(x, state, weight)
+    # a strided column, another filter width, another state length, a second dtype
+    assert not decode_kernel_fits(torch.zeros(2, 16)[:, ::2], state, weight)
+    assert not decode_kernel_fits(x, torch.zeros(3, 8, 4), torch.zeros(8, 5))
+    assert not decode_kernel_fits(x, torch.zeros(3, 8, 4), weight)
+    assert not decode_kernel_fits(x.to(torch.bfloat16), state, weight)
diff --git a/test/modular/test_model_kv_layouts.py b/test/modular/test_model_kv_layouts.py
index a4d4347fbc..5baae78203 100644
--- a/test/modular/test_model_kv_layouts.py
+++ b/test/modular/test_model_kv_layouts.py
@@ -457,6 +457,17 @@ def test_higgs_lays_out_the_prefill_step_it_runs():
     assert sum(span.length for span in layout.spans) == text + audio
 
 
+@pytest.mark.parametrize("drafts", [0, 3])
+def test_qwen3_5_keys_its_prompts_only_without_speculation(drafts):
+    # a speculative node keeps no prefix, and the deployment check refuses a keyed label there
+    from mstar.model.base import keyed_labels
+    from mstar.model.qwen3_5.qwen3_5_model import Qwen3_5DenseModel
+
+    model = object.__new__(Qwen3_5DenseModel)
+    model.num_draft_tokens = drafts
+    assert bool(keyed_labels(model)) == (drafts == 0)
+
+
 def _omni_submodule():
     return SimpleNamespace(config=SimpleNamespace(im_end_token_id=0, talker=SimpleNamespace(codec_eos_token_id=0)))
 
diff --git a/test/modular/test_platform_dispatch.py b/test/modular/test_platform_dispatch.py
index d925e23f81..d90bfb0fd0 100644
--- a/test/modular/test_platform_dispatch.py
+++ b/test/modular/test_platform_dispatch.py
@@ -212,7 +212,7 @@ class _SmallGemmsOnAten(RocmPlatform):
     def __init__(self):
         self.queries: list[int | None] = []
 
-    def inductor_overrides(self, num_tokens=None, rows=None):
+    def inductor_overrides(self, num_tokens=None, rows=None, decode_rows=False):
         self.queries.append(num_tokens)
         if num_tokens is not None and num_tokens <= 32:
             return {"max_autotune_gemm_backends": "ATEN"}
@@ -233,6 +233,14 @@ def test_no_overrides_returns_the_callable_itself(on_platform, platform, num_tok
     assert with_inductor_overrides(_read_gemm_backends, num_tokens) is _read_gemm_backends
 
 
+def test_a_verify_bucket_compiles_as_decode_by_its_rows(on_platform):
+    on_platform(RocmPlatform())
+    # 64 rows of three tokens: past the packed limit, inside the decode one
+    assert RocmPlatform().inductor_overrides(192, 64) == {}
+    assert RocmPlatform().inductor_overrides(192, 64, True)["max_autotune_gemm_backends"] == "ATEN"
+    assert RocmPlatform().inductor_overrides(600, 200, True) == {}
+
+
 @pytest.mark.parametrize(
     ("num_tokens", "rows", "aten_only"),
     [
@@ -285,7 +293,7 @@ def fails():
 
 
 class _EverythingOnAten(_SmallGemmsOnAten):
-    def inductor_overrides(self, num_tokens=None, rows=None):
+    def inductor_overrides(self, num_tokens=None, rows=None, decode_rows=False):
         self.queries.append(num_tokens)
         return {"max_autotune_gemm_backends": "ATEN"}
 
@@ -323,7 +331,7 @@ def __init__(self):
 def _bucket(num_tokens: int, bs: int | None = None):
     return SimpleNamespace(
         config=SimpleNamespace(
-            compile=True, compile_mode=None, capture_forward_method="forward_batched",
+            compile=True, compile_mode=None, capture_forward_method="forward_batched", compile_dynamic=False,
         ),
         config_idx=0,
         num_tokens=num_tokens,
diff --git a/test/modular/test_qwen3_5_mixed.py b/test/modular/test_qwen3_5_mixed.py
index 1a75420177..58c02bdaea 100644
--- a/test/modular/test_qwen3_5_mixed.py
+++ b/test/modular/test_qwen3_5_mixed.py
@@ -14,7 +14,8 @@
 
 
 def test_text_prefill_and_decode_share_the_mixed_walk():
-    assert Qwen3_5DenseModel.get_combined_graph_walks(None) == {
+    model = Qwen3_5DenseModel.__new__(Qwen3_5DenseModel)
+    assert model.get_combined_graph_walks() == {
         "LLM": {LLM_MIXED: {"prefill_text", "decode"}},
     }
 
@@ -88,3 +89,17 @@ def test_text_prefill_chunks_under_a_budget_that_fits_the_captures():
     # decode rows count inside the budget, and a step over a bucket pads to the next one up
     assert LLMSubmodule.MAX_BATCH_TOKENS in LLMSubmodule.PREFILL_TOKEN_BUCKETS
 
+
+def test_a_speculative_node_chunks_its_prefill_and_names_no_mixed_walk():
+    from mstar.model.qwen3_5.qwen3_5_model import Qwen3_5DenseModel
+
+    model = Qwen3_5DenseModel.__new__(Qwen3_5DenseModel)
+    model.config = SimpleNamespace(max_position_embeddings=4096)
+    model.vision_config = None
+    model.num_draft_tokens = 2
+    policies = model.chunking_policy_of()
+
+    assert policies[("LLM", "prefill_text")].chunkable
+    # a verify row is k + 1 tokens, so there is no mixed walk to carry a budget for
+    assert model.get_combined_graph_walks() == {}
+    assert ("LLM", LLM_MIXED) not in policies
diff --git a/test/modular/test_rocm_skinny_gemm_pass.py b/test/modular/test_rocm_skinny_gemm_pass.py
index 846dea270a..a5108f71eb 100644
--- a/test/modular/test_rocm_skinny_gemm_pass.py
+++ b/test/modular/test_rocm_skinny_gemm_pass.py
@@ -140,6 +140,15 @@ def test_rocm_installs_the_passes_for_sixteen_rows_or_fewer(monkeypatch, num_tok
         assert overrides["max_autotune_gemm_backends"] == "ATEN"
 
 
+def test_a_verify_bucket_takes_the_passes_by_its_rows(monkeypatch):
+    pre, post = SkinnyLinearPrePass(), SkinnyGemmPass()
+    monkeypatch.setattr(rocm_gemm, "skinny_gemm_passes", lambda: (pre, post))
+    # eight rows of three tokens: its draft GEMMs run at 8 rows, its target's at 24
+    overrides = RocmPlatform().inductor_overrides(24, 8, True)
+    assert overrides["pre_grad_custom_pass"] is pre
+    assert "pre_grad_custom_pass" not in RocmPlatform().inductor_overrides(24, 8)
+
+
 def test_rocm_without_the_kernel_keeps_hipblaslt_only(monkeypatch):
     monkeypatch.setattr(rocm_gemm, "skinny_gemm_passes", lambda: None)
     assert RocmPlatform().inductor_overrides(1) == {"max_autotune_gemm_backends": "ATEN"}
diff --git a/test/modular/test_spec_decode_resource.py b/test/modular/test_spec_decode_resource.py
new file mode 100644
index 0000000000..265b3f55b9
--- /dev/null
+++ b/test/modular/test_spec_decode_resource.py
@@ -0,0 +1,734 @@
+"""Speculative decoding's bookkeeping, driven the way a serving step drives it,
+on CPU tensors.
+
+The real KV manager, AITER's plan staging, the recurrent pool, the GDN verify
+wrapper and the speculative resource run their admit/plan/commit; a toy target
+and a toy draft stand in for the layers. Both read their context back out of a
+fake paged cache through the device-side addressing the resource writes (KV
+write slots, the block table, per-row KV lengths), and a toy recurrent kernel
+starts each row from the slot ``num_accepted`` names. So a wrong position, page,
+length or start slot changes what the toy models see, and the test sees it:
+
+* greedy speculation emits exactly what plain greedy decoding does, for k = 1..4,
+  whether the host learns each step's acceptance before the next step is
+  planned or one step late (the async worker), with page crossings, padding
+  rows, a stop id and a token budget;
+* a draft that is right except at known positions is accepted exactly up to
+  the first of them, every step (the draft's own KV chain is addressed right);
+* at every admit the host's KV length covers what the step writes, and the
+  host gives back exactly what it over-counted (reserve, then trim);
+* each request holds k + 1 recurrent slots and resumes from the one its last
+  step accepted up to (slot re-pointing without a copy).
+"""
+
+from __future__ import annotations
+
+import importlib.util
+import pathlib
+import sys
+from types import SimpleNamespace
+
+import numpy as np
+import pytest
+import torch
+from torch import nn
+
+sys.path.insert(0, ".")
+
+from mstar.engine.resources.attn.aiter import (  # noqa: E402
+    _TILE_TOKENS,
+    _image_numel,
+    _PlanImage,
+    kernel_rows,
+    stage_plan,
+)
+from mstar.engine.resources.kv.config import KVReqConfig, KVStep, PagedKVConfig  # noqa: E402
+from mstar.engine.resources.kv.manager import KVManager  # noqa: E402
+from mstar.engine.resources.kv.plan import SINK_PAGE  # noqa: E402
+from mstar.engine.resources.linear_attn.gdn_kernels import GDNKernels  # noqa: E402
+from mstar.engine.resources.linear_attn.wrappers import GDNVerifyWrapper  # noqa: E402
+from mstar.engine.resources.recurrent.config import (  # noqa: E402
+    RecurrentBlockConfig,
+    RecurrentStateConfig,
+    RecurrentStep,
+)
+from mstar.engine.resources.recurrent.pool import SINK_SLOT, RecurrentStatePool  # noqa: E402
+from mstar.engine.resources.spec_decode.config import SpecDecodeStep, SpecPhase  # noqa: E402
+from mstar.engine.resources.spec_decode.resource import SpecDecodePlan, SpecDecodeResource  # noqa: E402
+from mstar.engine.resources.step import BucketKey, Segment, SlotLease, StepContext  # noqa: E402
+from mstar.model.qwen3_5.qwen3_5_model import KV_CACHE, Qwen3_5DenseModel  # noqa: E402
+
+_spec = importlib.util.spec_from_file_location(
+    "kv_prefix_decode_helpers", pathlib.Path(__file__).with_name("test_kv_prefix_decode.py"),
+)
+H = importlib.util.module_from_spec(_spec)
+_spec.loader.exec_module(H)
+
+CPU = torch.device("cpu")
+PAGE = 4
+VOCAB = 97
+EOS = VOCAB - 1
+NUM_PAGES = 512
+PAD_COLS = -(-_TILE_TOKENS // PAGE)
+
+
+@pytest.fixture(autouse=True)
+def _stub_transfer(monkeypatch):
+    monkeypatch.setattr(H.manager_mod, "KVTransferManager", H._StubTransfer)
+
+
+def next_token(seq) -> int:
+    """The toy target: a deterministic function of the whole prefix, never EOS
+    unless the test plants it."""
+    acc = len(seq)
+    for i, t in enumerate(seq):
+        acc = (acc * 31 + (int(t) + 7) * (i + 3)) % 1_000_003
+    return acc % (VOCAB - 1)
+
+
+def greedy(prompt: list[int], budget: int, stop_at: int | None = None) -> list[int]:
+    """Plain greedy decoding: the prompt's token, then ``budget`` more, cut
+    after the first token at position ``stop_at`` (made EOS)."""
+    seq, out = list(prompt), []
+    for _ in range(budget + 1):
+        t = EOS if stop_at is not None and len(seq) == stop_at else next_token(seq)
+        out.append(t)
+        seq.append(t)
+        if t == EOS:
+            break
+    return out
+
+
+class ToyState:
+    """A recurrent state per slot: a running hash of the tokens it has seen."""
+
+    @staticmethod
+    def step(state: float, token: int) -> float:
+        return (state * 1.000_37 + token + 1) % 10_007.0
+
+
+class _NoKernels(GDNKernels):
+    qk_l2norm = False
+
+    def chunked_prefill(self):
+        raise NotImplementedError
+
+    def captured_prefill(self, bs, num_tokens):
+        raise NotImplementedError
+
+    def decode(self, bs):
+        raise NotImplementedError
+
+    def verify(self):
+        return None
+
+
+class Harness:
+    """One node's resources and a fake two-layer token cache (layer 0 the
+    target's, layer 1 the draft head's)."""
+
+    def __init__(self, k: int, *, stop_at: int | None = None, wrong: frozenset[int] = frozenset()):
+        self.k, self.t = k, k + 1
+        self.stop_at = stop_at
+        # absolute token positions the toy draft guesses wrong
+        self.wrong = wrong
+        self.kv = KVManager(
+            cfg=PagedKVConfig(
+                num_layers=2, num_kv_heads=1, head_dim=8, max_seq_len=1024,
+                max_num_pages=NUM_PAGES, page_size=PAGE,
+            ),
+            name="kv", joint_comm_group=None, transfer_engine_info=None,
+            device=CPU, dtype=torch.float32,
+        )
+        self.kv._cg_max_seq_len = 256
+        self.spec = SpecDecodeResource(
+            device=CPU, num_draft_tokens=k, max_requests=8, page_size=PAGE,
+            kv_key="kv", attn_key="attn", stop_token_ids=(EOS,),
+        )
+        self.pool = RecurrentStatePool(CPU, RecurrentStateConfig(
+            num_layers=1,
+            blocks={
+                "state": RecurrentBlockConfig(shape=(1,), dtype=torch.float64),
+                "conv": RecurrentBlockConfig(shape=(2, 3), dtype=torch.float32, speculative_window=True),
+            },
+            max_slots=1 + 8 * (k + 1),
+            spec_tokens=k,
+        ))
+        self.gdn: dict = {}
+        self.cache = torch.full((NUM_PAGES, PAGE, 2), -1, dtype=torch.long)
+        self.slot = 0
+        self.max_seen_writes = 0
+
+    # one step's resources, as the runner sweeps them
+
+    def _ctx(self, rids, padded, walk, bucket=None):
+        self.slot ^= 1
+        lease = None if bucket is None else SlotLease(slot=self.slot, bucket=bucket)
+        ctx = StepContext(request_ids=tuple(rids), graph_walk=walk, slot=self.slot, capture=False,
+                          slot_lease=lease)
+        if padded is not None:
+            ctx.set_padded_rids(tuple(padded))
+        return ctx
+
+    def _attn(self, kv_out, lease):
+        rows = kernel_rows(kv_out.cpu_indptrs, PAGE)
+        num_rows = rows.qo_indptr.shape[0] - 1
+        width = NUM_PAGES // 4 + PAD_COLS if lease is not None else int(np.diff(rows.kv_indptr).max()) + PAD_COLS
+        image = _PlanImage(_image_numel(num_rows, width), CPU)
+        return stage_plan(image, rows, width, PAD_COLS, self.t, 1024)
+
+    def prefill(self, prompts: dict[int, list[int]], budgets: dict[int, int]):
+        rids = list(prompts)
+        for rid in rids:
+            self.kv.ingest_request(rid, KVReqConfig())
+            self.spec.ingest_request(rid)
+            self.pool.ingest_request(rid)
+        segs = tuple(Segment(rid, "main", len(prompts[rid])) for rid in rids)
+        ctx = self._ctx(rids, None, "prefill_text")
+        kv_step = KVStep(segments=segs, commit_adjust={(r, "main"): self.spec.seed_lookahead for r in rids})
+        spec_step = SpecDecodeStep(segments=segs, phase=SpecPhase.SEED, budgets=budgets)
+        pool_step = RecurrentStep(segments=segs)
+        assert self.kv.admit(kv_step, ctx).ok and self.pool.admit(pool_step, ctx).ok
+        assert self.spec.admit(spec_step, ctx).ok
+        ctx.plan_results["kv"] = self.kv.plan(kv_step, ctx)
+        ctx.plan_results["gdn_state"] = self.pool.plan(pool_step, ctx)
+        self.spec.plan(spec_step, ctx)
+        # the forward: target tokens and draft-layer tokens at prompt positions,
+        # each row's recurrent state over its prompt, then the seed
+        state = self.pool.block("state", 0)
+        first, drafts = [], []
+        for i, rid in enumerate(rids):
+            prompt = prompts[rid]
+            t0 = EOS if self.stop_at == len(prompt) else next_token(prompt)
+            view = ctx.plan_results["kv"]["main"].views[i]
+            for p, tok in enumerate(prompt):
+                page = view.page_idxs[p // PAGE]
+                self.cache[page, p % PAGE, 0] = tok
+                self.cache[page, p % PAGE, 1] = (prompt + [t0])[p + 1]
+            h = 0.0
+            for tok in prompt:
+                h = ToyState.step(h, tok)
+            state[self.pool._slots[rid]["main"].index, 0] = h
+            first.append(t0)
+            drafts.append(self._draft_guess(prompt + [t0], len(prompt) + 1))
+        self.spec.seed(torch.tensor(first), torch.tensor(drafts))
+        self.kv.commit(kv_step, ctx)
+        self.pool.commit(pool_step, ctx)
+        self.spec.commit(spec_step, ctx)
+        return first
+
+    def _draft_guess(self, seq: list[int], position: int) -> int:
+        """The toy draft's guess for the token at ``position``, from ``seq``
+        (the tokens before it); wrong where the test says."""
+        true = EOS if self.stop_at == position else next_token(seq[:position])
+        return (true + 1) % (VOCAB - 1) if position in self.wrong else true
+
+    def verify(self, rids: list[int], pad: int = 0):
+        """One verify step and its draft passes; returns (tokens, accepted, done)
+        rows for the real requests, and checks the reservation invariant."""
+        padded = rids + [-(i + 1) for i in range(pad)]
+        t = self.t
+        retract = {r: self.spec.pending_retract(r) for r in rids if self.spec.pending_retract(r)}
+        segs = tuple(Segment(r, "main", t) for r in padded)
+        bucket = BucketKey("decode", len(padded), len(padded) * t) if pad else None
+        ctx = self._ctx(rids, padded if pad else None, "decode", bucket)
+        lease = ctx.slot_lease
+        kv_step = KVStep(segments=segs, commit_adjust={(r, "main"): -g for r, g in retract.items()})
+        pool_step = RecurrentStep(segments=segs)
+        spec_step = SpecDecodeStep(segments=segs, phase=SpecPhase.VERIFY, retract=retract)
+        assert self.kv.admit(kv_step, ctx).ok and self.pool.admit(pool_step, ctx).ok
+        assert self.spec.admit(spec_step, ctx).ok
+        ctx.plan_results["kv"] = kv_out = self.kv.plan(kv_step, ctx)
+        ctx.plan_results["attn"] = {"main": self._attn(kv_out["main"], lease)}
+        ctx.plan_results["gdn_state"] = addressing = self.pool.plan(pool_step, ctx)
+        plan = self.spec.plan(spec_step, ctx)
+        wrapper = self.gdn.setdefault(lease is not None, GDNVerifyWrapper(
+            CPU, pad_slot_id=-1, tokens_per_row=t, bs=len(padded) if pad else None,
+            accelerator_graph=lease is not None, null_slot_id=SINK_SLOT, kernels=_NoKernels(),
+        ))
+        wrapper.plan(addressing["main"].host_slot_groups, plan.num_accepted)
+
+        ids, positions = self.spec.verify_inputs()
+        b = len(padded)
+        base = positions.view(b, t)[:, 0]
+        views = kv_out["main"].views
+        for i in range(len(rids)):
+            # the host's optimistic length covers every write of this step
+            assert views[i].length >= int(base[i]) + 2 * self.k, (views[i].length, int(base[i]))
+            assert len(views[i].page_idxs) * PAGE >= int(base[i]) + 2 * self.k
+        self._write(plan, ids, layer=0)
+        self._recurrent(wrapper, ids, b)
+        targets = self._target(plan, b, positions)
+        out = self.spec.accept_tokens(targets)
+        drafted = out.target.reshape(-1)
+        keep = self.spec.draft_slots()
+        for j in range(self.k):
+            where = self.spec.draft_address(j)
+            self._write_flat(where.slots, drafted, layer=1)
+            query = where.positions[keep] if j == 0 else where.positions
+            guess = self._draft(plan, b, query, where.kv_lens)
+            drafted = self.spec.record_draft(j, torch.nn.functional.one_hot(guess, VOCAB).float())
+        self.kv.commit(kv_step, ctx)
+        self.pool.commit(pool_step, ctx)
+        self.spec.commit(spec_step, ctx)
+        n = len(rids)
+        return out.tokens[:n], out.accepted[:n], out.done[:n]
+
+    def current_state(self, rid: int) -> float:
+        """The recurrent state the request's next step starts from."""
+        group = self.pool._slots[rid]["main"].group
+        accepted = int(self.spec._num_acc[self.spec._row_of[rid]])
+        return float(self.pool.block("state", 0)[group[accepted - 1], 0])
+
+    def _write(self, plan, tokens, layer):
+        pages, offsets = plan.write_pages, plan.write_offsets
+        real = pages != SINK_PAGE
+        self.cache[pages[real], offsets[real], layer] = tokens[real]
+
+    def _write_flat(self, slots, tokens, layer):
+        """A draft pass's KV write, one flat slot per token."""
+        pages, offsets = torch.div(slots, PAGE, rounding_mode="floor"), slots % PAGE
+        real = pages != SINK_PAGE
+        self.cache[pages[real], offsets[real], layer] = tokens[real]
+
+    def _read(self, plan, row: int, upto: int, layer: int) -> list[int]:
+        """Tokens of ``layer`` at positions ``0 .. upto - 1`` of a row, through
+        the block table the step's attention reads."""
+        table = plan.block_table[row]
+        return [int(self.cache[int(table[p // PAGE]), p % PAGE, layer]) for p in range(upto)]
+
+    def _target(self, plan, b, positions) -> torch.Tensor:
+        """Query i of row r sees keys 0 .. kv_len - (k + 1) + i, causal from
+        the bottom right as the kernel aligns it."""
+        t = self.t
+        lens = plan.kv_lens
+        out = torch.zeros(b * t, dtype=torch.long)
+        for r in range(b):
+            if not bool(plan.image.real[r]):
+                continue
+            for i in range(t):
+                seen = int(lens[r]) - t + i + 1
+                ctx_tokens = self._read(plan, r, seen, layer=0)
+                assert -1 not in ctx_tokens, f"row {r} attends an unwritten key"
+                out[r * t + i] = EOS if self.stop_at == seen else next_token(ctx_tokens)
+        return out
+
+    def _draft(self, plan, b, positions, kv_lens) -> torch.Tensor:
+        """A row's one query, at position q, reads the draft layer's keys 0 ..
+        q (the token after each) and guesses the token at q + 2."""
+        out = torch.zeros(b, dtype=torch.long)
+        for r in range(b):
+            if not bool(plan.image.real[r]):
+                continue
+            seen = int(kv_lens[r])
+            q = int(positions[r])
+            assert seen == q + 1, "a draft query must see its own position and none after"
+            seq = self._read(plan, r, 1, layer=0) + self._read(plan, r, seen, layer=1)
+            out[r] = self._draft_guess(seq, q + 2) if -1 not in seq else 0
+        return out
+
+    def _recurrent(self, wrapper: GDNVerifyWrapper, ids, b):
+        """The verify kernel's contract on the toy state: start from the
+        accepted slot, keep the state after every token."""
+        plan = wrapper._plan_state
+        state = self.pool.block("state", 0)
+        for r in range(b):
+            row = plan.slot_map[r]
+            if int(row[0]) < 0:
+                continue
+            h = float(state[int(row[int(plan.num_accepted[r]) - 1]), 0])
+            for i in range(self.t):
+                h = ToyState.step(h, int(ids[r * self.t + i]))
+                state[int(row[i]), 0] = h
+
+
+def run(h: Harness, prompts: dict[int, list[int]], budget: int, lag: int, pad: int = 0):
+    """Generate to completion; ``lag`` steps between a step and the host
+    learning what it accepted. Returns each request's emitted tokens and the
+    accepted count of every step."""
+    rids = list(prompts)
+    first = h.prefill(prompts, {r: budget for r in rids})
+    emitted = {r: [first[i]] for i, r in enumerate(rids)}
+    live = [r for i, r in enumerate(rids) if first[i] != EOS]
+    pending: list[tuple[list[int], list[int]]] = []
+    history: dict[int, list[int]] = {r: [] for r in rids}
+    for _ in range(4 * budget + 8):
+        if not live:
+            break
+        tokens, accepted, done = h.verify(live, pad=pad)
+        pending.append((list(live), accepted.tolist()))
+        if len(pending) > lag:
+            h.spec.learn(*pending.pop(0))
+        still = []
+        for i, rid in enumerate(live):
+            row = [int(x) for x in tokens[i] if int(x) >= 0]
+            emitted[rid] += row
+            history[rid].append(int(accepted[i]))
+            if not bool(done[i]):
+                still.append(rid)
+                # it has run every committed token but the last emitted one
+                expect = 0.0
+                for tok in prompts[rid] + emitted[rid][:-1]:
+                    expect = ToyState.step(expect, tok)
+                assert h.current_state(rid) == pytest.approx(expect), rid
+        live = still
+    assert not live, "generation never finished"
+    return emitted, history
+
+
+PROMPTS = {0: [3, 1, 4, 1, 5, 9, 2], 1: [2, 7, 1, 8, 2, 8, 1, 8, 2, 8, 4]}
+
+
+@pytest.mark.parametrize("k", [1, 2, 3, 4])
+@pytest.mark.parametrize("lag", [0, 1])
+def test_greedy_speculation_reproduces_greedy(k, lag):
+    budget = 23
+    h = Harness(k, wrong=frozenset(range(9, 200, 5)))
+    emitted, history = run(h, PROMPTS, budget, lag)
+    for rid, prompt in PROMPTS.items():
+        assert emitted[rid] == greedy(prompt, budget), rid
+        assert all(1 <= a <= k + 1 for a in history[rid])
+
+
+@pytest.mark.parametrize("k", [1, 3, 4])
+def test_padding_rows_leave_real_rows_alone(k):
+    budget = 17
+    h = Harness(k, wrong=frozenset(range(12, 200, 4)))
+    emitted, _ = run(h, PROMPTS, budget, lag=1, pad=2)
+    for rid, prompt in PROMPTS.items():
+        assert emitted[rid] == greedy(prompt, budget), rid
+
+
+@pytest.mark.parametrize("k", [2, 3, 4])
+def test_a_perfect_draft_is_accepted_whole(k):
+    h = Harness(k)
+    budget = 6 * (k + 1)
+    _, history = run(h, {0: PROMPTS[0]}, budget, lag=1)
+    # The prompt's step drafts one token (the rest of the first row repeat
+    # it), so the first verify takes at most two; every later step but the
+    # last (cut by the budget) takes all k + 1.
+    assert history[0][0] <= 2
+    assert history[0][1:-1] == [k + 1] * (len(history[0]) - 2)
+
+
+@pytest.mark.parametrize("k", [1, 2, 3, 4])
+def test_acceptance_stops_at_the_first_wrong_draft(k):
+    wrong = frozenset({11, 12, 17, 23, 24, 25, 31})
+    h = Harness(k, wrong=wrong)
+    prompt = PROMPTS[0]
+    _, history = run(h, {0: prompt}, 30, lag=1)
+    # the first verify's x_0 (the prompt's token) sits at len(prompt); its
+    # drafts past the first repeat it (see the perfect-draft test)
+    base = len(prompt) + history[0][0]
+    for accepted in history[0][1:-1]:
+        drafts = [base + i for i in range(1, k + 1)]
+        expect = 1
+        for pos in drafts:
+            if pos in wrong:
+                break
+            expect += 1
+        assert accepted == expect, (base, accepted, expect)
+        base += accepted
+
+
+@pytest.mark.parametrize("k", [1, 3])
+def test_a_stop_id_inside_a_step_ends_it_there(k):
+    prompt = PROMPTS[1]
+    stop_at = len(prompt) + 6
+    h = Harness(k, stop_at=stop_at)
+    emitted, _ = run(h, {1: prompt}, 40, lag=1)
+    assert emitted[1] == greedy(prompt, 40, stop_at=stop_at)
+    assert emitted[1][-1] == EOS
+
+
+@pytest.mark.parametrize("k", [2, 4])
+def test_host_gives_back_exactly_what_it_overcounted(k):
+    h = Harness(k, wrong=frozenset(range(10, 300, 3)))
+    prompt = PROMPTS[0]
+    rid = 0
+    h.prefill({rid: prompt}, {rid: 40})
+    stream = h.kv._streams[rid]["main"]
+    assert stream.stored_len == len(prompt) + k - 1
+    true_len = len(prompt)
+    for _ in range(6):
+        _, accepted, _ = h.verify([rid])
+        true_len += int(accepted[0])
+        # nothing learned yet: the step committed its whole row
+        assert stream.stored_len == true_len + (k - 1) + (k + 1 - int(accepted[0]))
+        h.spec.learn([rid], [int(accepted[0])])
+        assert h.spec.pending_retract(rid) == k + 1 - int(accepted[0])
+        # the next verify gives it back in its commit, then over-counts anew
+    _, accepted, _ = h.verify([rid])
+    true_len += int(accepted[0])
+    assert h.spec.pending_retract(rid) == 0
+    assert stream.stored_len == true_len + (k - 1) + (k + 1 - int(accepted[0]))
+
+
+def test_draft_passes_attend_one_query_a_row():
+    """Pass 0 writes every verify position and attends slot n alone; a later
+    pass runs one token a row, at the position after it. Both reach the
+    attention with each row's keys up to and including its query."""
+    k, rows, heads, kv_heads, dim = 2, 2, 2, 1, 8
+    t = k + 1
+    spec = SpecDecodeResource(device=CPU, num_draft_tokens=k, max_requests=rows, page_size=PAGE,
+                              kv_key="kv", attn_key="attn", stop_token_ids=(EOS,))
+    for r in range(rows):
+        spec.ingest_request(r)
+    image = spec._new_image(rows)
+    image.rows.copy_(torch.tensor([spec._row_of[r] for r in range(rows)], dtype=torch.int32))
+    image.real.fill_(1)
+    width = 16
+    table = (1 + torch.arange(rows * width, dtype=torch.int32)).view(rows, width)
+    from mstar.engine.resources.spec_decode.resource import SpecDecodePlan
+
+    spec._current = SpecDecodePlan(
+        phase=SpecPhase.VERIFY, num_rows=rows, image=image, block_table=table,
+        kv_lens=torch.zeros(rows, dtype=torch.int32), write_pages=torch.zeros(rows * t, dtype=torch.long),
+        write_offsets=torch.zeros(rows * t, dtype=torch.long),
+    )
+    regs = slice(1, rows + 1)
+    spec._seq_len[regs] = torch.tensor([10, 21], dtype=torch.int32)
+    spec._budget[regs] = 1 << 20
+    # row 0 accepts everything, row 1 refuses its first draft
+    spec._tokens[:, regs] = torch.tensor([[1, 2], [3, 4], [5, 6]])
+    ids, _ = spec.verify_inputs()
+    spec.accept_tokens(torch.tensor([3, 5, 9, 8, 8, 8]))
+
+    calls = []
+
+    class Attn:
+        def attend_rows(self, q, layer, kv_lens):
+            calls.append((q.clone(), kv_lens.tolist()))
+            return q
+
+    pages, slots_per_page = 32, PAGE
+    layer = torch.zeros(pages, 2, slots_per_page, kv_heads, dim)
+    spec._kv = type("KV", (), {"layer_view": lambda self, idx: layer})()
+    spec._attn = Attn()
+    cos, sin = torch.ones(64, 2), torch.zeros(64, 2)
+    width_qkv = (2 * heads + 2 * kv_heads) * dim
+    gen = torch.Generator().manual_seed(0)
+    q_w, k_w = torch.zeros(dim), torch.zeros(dim)
+    args = (q_w, k_w, cos, sin, heads, kv_heads, dim, 1e-6, True, 0)
+
+    full = torch.randn(rows * t, width_qkv, generator=gen)
+    out0, gate0 = spec.attend_draft(0, full, *args)
+    assert out0.shape == (rows, heads, dim) and gate0.shape == (rows, heads, dim)
+    # row 0's last accepted slot is 2, row 1's is 0
+    assert spec.draft_slots().tolist() == [2, 3]
+    assert calls[0][1] == [13, 22]
+    pos_pages = [int(table[0, p // PAGE]) for p in (10, 11, 12)]
+    written = layer[pos_pages, 0, [10 % PAGE, 11 % PAGE, 12 % PAGE]]
+    assert written.abs().sum(dim=(1, 2)).min() > 0, "pass 0 writes every verify position of a row"
+
+    one = torch.randn(rows, width_qkv, generator=gen)
+    out1, _ = spec.attend_draft(1, one, *args)
+    assert out1.shape == (rows, heads, dim)
+    assert calls[1][1] == [14, 23]
+    where = spec.draft_address(1)
+    assert where.positions.tolist() == [13, 22]
+    assert where.slots.tolist() == [int(table[r, p // PAGE]) * PAGE + p % PAGE for r, p in enumerate((13, 22))]
+
+
+def test_attend_rows_reads_one_query_a_row_off_a_verify_plan(monkeypatch):
+    from mstar.engine.resources.attn import aiter
+
+    manager = aiter.AiterPagedAttentionManager.__new__(aiter.AiterPagedAttentionManager)
+    manager._device = CPU
+    manager._row_starts, manager._retired_row_starts = None, []
+    manager._reserve_row_starts(3)
+    cu, lens = torch.tensor([0, 3, 6, 9], dtype=torch.int32), torch.zeros(3, dtype=torch.int32)
+    table = torch.zeros(3, 8, dtype=torch.int32)
+    manager._current_plans = {"main": aiter._Staged(cu, lens, table, 3, 1024, 9)}
+    seen = {}
+    monkeypatch.setattr(manager, "_unified", lambda q, plan, layer: seen.update(plan=plan) or q)
+    kv_lens = torch.tensor([7, 8, 9], dtype=torch.int32)
+    manager.attend_rows(torch.zeros(3, 2, 4), None, kv_lens)
+    plan = seen["plan"]
+    assert plan.cu_seqlens_q.tolist() == [0, 1, 2, 3]
+    assert plan.seqused_k is kv_lens and plan.max_seqlen_q == 1 and plan.num_tokens == 3
+    assert plan.block_table is table and plan.max_seqlen_k == 1024
+
+
+def _verify_plan(max_requests: int, rows: int, k: int = 2) -> SpecDecodeResource:
+    """A resource with a verify plan over ``rows`` padding rows."""
+    t = k + 1
+    spec = SpecDecodeResource(device=CPU, num_draft_tokens=k, max_requests=max_requests, page_size=16,
+                              kv_key="kv", attn_key="attn", stop_token_ids=(2,))
+    spec._current = SpecDecodePlan(
+        phase=SpecPhase.VERIFY, num_rows=rows, image=spec._new_image(rows),
+        block_table=torch.zeros(rows, 4, dtype=torch.int32), kv_lens=torch.zeros(rows, dtype=torch.int32),
+        write_pages=torch.zeros(rows * t, dtype=torch.long), write_offsets=torch.zeros(rows * t, dtype=torch.long),
+    )
+    return spec
+
+
+def test_a_bucket_wider_than_the_registers_argmaxes_every_row():
+    rows, t, vocab = 8, 3, 16
+    spec = _verify_plan(max_requests=4, rows=rows)
+    spec.verify_inputs()
+    logits = torch.zeros(rows * t, vocab)
+    logits[:, 3] = 1
+    accepted = spec.accept(logits)
+    assert accepted.target.shape == (rows, t) and accepted.target.eq(3).all()
+    assert spec.record_draft(0, torch.zeros(rows, vocab).index_fill(1, torch.tensor([5]), 1.0)).eq(5).all()
+
+
+def test_a_head_with_no_output_gate_drafts():
+    rows, t, heads, kv_heads, dim = 2, 3, 2, 1, 8
+    spec = _verify_plan(max_requests=4, rows=rows)
+    spec.verify_inputs()
+    spec.accept(torch.zeros(rows * t, 16))
+    layer = torch.zeros(32, 2, 16, kv_heads, dim)
+    spec._kv = SimpleNamespace(layer_view=lambda idx: layer)
+    spec._attn = SimpleNamespace(attend_rows=lambda q, layer, kv_lens: q)
+    qkv = torch.randn(rows * t, (heads + 2 * kv_heads) * dim)
+    weight = torch.zeros(dim)
+    out, _ = spec.attend_draft(0, qkv, weight, weight, torch.ones(64, 2), torch.zeros(64, 2),
+                               heads, kv_heads, dim, 1e-6, False, 0)
+    assert out.shape[0] == rows
+
+
+def test_speculation_refuses_a_context_past_the_rotary_table():
+    model = Qwen3_5DenseModel.__new__(Qwen3_5DenseModel)
+    model.config = SimpleNamespace(max_position_embeddings=1024)
+    model.num_draft_tokens = 2
+    fits, past = ({"resources": {KV_CACHE: {"max_seq_len": n}}} for n in (1024, 1025))
+    model.validate_config_yaml(fits, "x.yaml")
+    with pytest.raises(ValueError, match="max_seq_len=1025"):
+        model.validate_config_yaml(past, "x.yaml")
+    # plain decode builds its rotations from the position
+    model.num_draft_tokens = 0
+    model.validate_config_yaml(past, "x.yaml")
+
+
+def test_an_fp8_head_keeps_a_clear_argmax(monkeypatch):
+    from mstar.engine.resources.spec_decode import draft_head as module
+
+    class Platform:
+        name = "test"
+
+        def fp8_dtype(self):
+            return torch.float8_e4m3fnuz
+
+    monkeypatch.setattr(module, "current_platform", Platform())
+    gen = torch.Generator().manual_seed(0)
+    weight = torch.randn(256, 64, generator=gen, dtype=torch.bfloat16) * 0.02
+    head = module.Fp8DraftHead(weight)
+    assert head.weight.dtype == torch.float8_e4m3fnuz
+    # the engine moves the whole submodule to its dtype
+    bits = head.weight.clone()
+    head = nn.Sequential(head).to(dtype=torch.bfloat16)[0]
+    assert head.weight.dtype == torch.float8_e4m3fnuz and torch.equal(head.weight, bits)
+    # rows along a row of the head: its own dot product stands well clear of the rest
+    x = weight[[3, 200]].float() + 0.001 * torch.randn(2, 64, generator=gen)
+    logits = head(x.to(torch.bfloat16))
+    assert logits.shape == (2, 256) and logits.dtype == torch.bfloat16
+    assert logits.argmax(dim=-1).tolist() == [3, 200]
+
+
+def test_requests_hold_slot_groups_and_return_them():
+    k = 3
+    h = Harness(k)
+    free = h.pool.num_free_slots
+    h.prefill({0: PROMPTS[0], 1: PROMPTS[1]}, {0: 4, 1: 4})
+    groups = [h.pool._slots[r]["main"].group for r in (0, 1)]
+    assert all(len(g) == k + 1 for g in groups)
+    assert len(set(groups[0]) | set(groups[1])) == 2 * (k + 1)
+    assert h.pool.num_free_slots == free - 2 * (k + 1)
+    # the conv window holds k more columns; the state block is as declared
+    assert h.pool.block("conv", 0).shape[-1] == 3 + k
+    assert h.pool.block("state", 0).shape[-1] == 1
+    for r in (0, 1):
+        h.pool.remove_request(r)
+        h.spec.remove_request(r)
+    assert h.pool.num_free_slots == free
+    assert h.spec.pending_retract(0) == 0 and not h.spec._row_of
+
+
+def _filtered(logits: torch.Tensor, temperature: float, top_p: float, top_k: int) -> torch.Tensor:
+    """The sampler's distribution: top-k, then top-p over the renormalised rest."""
+    probs = (logits.double() / temperature).softmax(-1)
+    order = probs.argsort(descending=True)
+    keep = torch.zeros_like(probs, dtype=torch.bool)
+    if top_k > 0:
+        keep[order[:top_k]] = True
+    else:
+        keep[:] = True
+    probs = torch.where(keep, probs, 0.0)
+    probs = probs / probs.sum()
+    sorted_p = probs[order]
+    mass_before = sorted_p.cumsum(0) - sorted_p
+    cut = torch.zeros_like(keep)
+    cut[order] = mass_before < top_p
+    probs = torch.where(cut & keep, probs, 0.0)
+    return probs / probs.sum()
+
+
+def _reference_draw(logits, temperature, top_k, top_p, seed, offset):
+    """`draw_tokens` on CPU: each row's filtered distribution, drawn with a
+    generator seeded from its seed and offset (argmax at temperature 0)."""
+    out = torch.empty(logits.shape[0], dtype=torch.long)
+    for i in range(logits.shape[0]):
+        t = float(temperature[i])
+        if t == 0:
+            out[i] = logits[i].argmax()
+            continue
+        k = int(top_k[i])
+        probs = _filtered(logits[i], t, float(top_p[i]), k if k < logits.shape[1] else 0)
+        gen = torch.Generator().manual_seed(int(seed[i]) * 1_000_003 + int(offset[i]) * 7919 + i)
+        out[i] = torch.multinomial(probs.float(), 1, generator=gen)[0]
+    return out
+
+
+@pytest.mark.parametrize("params", [(1.0, 1.0, 0), (0.7, 0.9, 0), (1.3, 1.0, 5), (0.8, 0.6, 4)])
+def test_sampled_verification_follows_the_target_distribution(params, monkeypatch):
+    """k = 1 against fixed logits: the first emitted token is a draw from the
+    target's filtered distribution at position 0, whether or not the draft is
+    accepted, and the token after an accepted draft one from position 1's."""
+    from mstar.engine.resources.spec_decode import resource as spec_mod
+
+    monkeypatch.setattr(spec_mod.sampling, "draw_tokens", _reference_draw)
+    torch.manual_seed(0)
+    temperature, top_p, top_k = params
+    rows, vocab, draws = 400, 12, 30
+    spec = SpecDecodeResource(device=CPU, num_draft_tokens=1, max_requests=rows, page_size=PAGE,
+                              kv_key="kv", attn_key="attn", sampling=True)
+    for r in range(rows):
+        spec.ingest_request(r)
+    image = spec._new_image(rows)
+    image.rows.copy_(torch.tensor([spec._row_of[r] for r in range(rows)], dtype=torch.int32))
+    image.real.fill_(1)
+    table = torch.arange(1, rows * 4 + 1, dtype=torch.int32).view(rows, 4)
+    state = torch.zeros(rows * 2, dtype=torch.long)
+    from mstar.engine.resources.spec_decode.resource import SpecDecodePlan
+
+    spec._current = SpecDecodePlan(
+        phase=SpecPhase.VERIFY, num_rows=rows, image=image, block_table=table,
+        kv_lens=torch.zeros(rows, dtype=torch.int32), write_pages=state, write_offsets=state.clone(),
+    )
+    regs = slice(1, rows + 1)
+    spec._temperature[regs], spec._top_p[regs], spec._top_k[regs] = temperature, top_p, top_k
+    spec._seed[regs] = torch.arange(rows)
+    spec._budget[regs] = 1 << 20
+    logits = torch.randn(2, vocab) * 2
+    p0, p1 = (_filtered(logits[i], temperature, top_p, top_k) for i in range(2))
+    draft = int(p0.argmax())
+    first, after = torch.zeros(vocab), torch.zeros(vocab)
+    for _ in range(draws):
+        spec._tokens[:, regs] = torch.tensor([[3], [draft]])
+        spec._seq_len[regs] = 10
+        spec.verify_inputs()
+        out = spec.accept(logits.repeat(rows, 1), sample=True)
+        first += torch.bincount(out.tokens[:, 0], minlength=vocab)
+        accepted = out.accepted == 2
+        after += torch.bincount(out.tokens[accepted, 1], minlength=vocab)
+    n = rows * draws
+    tv_first = 0.5 * (first / n - p0).abs().sum()
+    tv_after = 0.5 * (after / after.sum() - p1).abs().sum()
+    assert tv_first < 0.02, (tv_first, first / n, p0)
+    assert tv_after < 0.03, (tv_after, after / after.sum(), p1)
+    # filtered-out tokens are never drawn
+    assert not first[p0 == 0].any() and not after[p1 == 0].any()
+    # the draft is accepted as often as the target would have drawn it
+    assert abs(float(accepted.float().mean()) - float(p0[draft])) < 0.1
diff --git a/test/modular/test_spec_decode_verify.py b/test/modular/test_spec_decode_verify.py
new file mode 100644
index 0000000000..f8166a7c6e
--- /dev/null
+++ b/test/modular/test_spec_decode_verify.py
@@ -0,0 +1,236 @@
+"""CPU tests for benchmark/spec_decode: the verify primitives and the speculative loop, on bigram table models.
+
+The fakes assert the loop's bookkeeping as it runs: every draft position the loop extends must pair a target
+position with the token that followed it, chain positions must be rolled back before the next extend, and the
+target must be fed exactly the token the draft was told would come next.
+"""
+from __future__ import annotations
+
+import itertools
+import math
+
+import pytest
+import torch
+
+from benchmark.spec_decode.spec_loop import run_speculative
+from benchmark.spec_decode.verify import SamplingParams, accept_draft, sampling_probs
+
+V = 5
+H = 3
+
+
+class TableTarget:
+    """Bigram model: next-token logits depend on the last token only; hidden states are a fixed embedding."""
+
+    def __init__(self, logits: torch.Tensor, hidden: torch.Tensor):
+        self.logits = logits
+        self.hidden = hidden
+        self.fed: list[int] = []
+        self.expect_next: int | None = None
+
+    def prefill(self, tokens):
+        self.fed = list(tokens)
+        return self.logits[tokens[-1]].clone(), self.hidden[tokens]
+
+    def step(self, token):
+        if self.expect_next is not None:
+            assert token == self.expect_next, "the target was fed a token other than the one the draft paired"
+            self.expect_next = None
+        self.fed.append(token)
+        return self.logits[token].clone(), self.hidden[token]
+
+
+class TableDraft:
+    """Bigram drafter keyed by the token paired with each position; checks the cache bookkeeping of the loop."""
+
+    def __init__(self, logits: torch.Tensor, target: TableTarget):
+        self.logits = logits
+        self.target = target
+        self.length = 0
+        self.chained = 0
+        self.extends = 0
+
+    def reset(self):
+        self.length = 0
+        self.chained = 0
+
+    def extend(self, hiddens, tokens):
+        assert self.chained == 0, "chain positions must be rolled back before extending"
+        n = tokens.shape[0]
+        assert hiddens.shape == (n, H)
+        fed = self.target.fed
+        # Position p pairs the target's hidden state at p with token p+1; the last token is still pending.
+        assert tokens[:-1].tolist() == fed[self.length + 1:], "tokens do not follow the positions they pair with"
+        assert self.length + n == len(fed), "extend must cover exactly the positions the target processed"
+        self.target.expect_next = int(tokens[-1])
+        self.length += n
+        self.extends += 1
+        return self.logits[int(tokens[-1])].clone()
+
+    def chain(self, token):
+        self.length += 1
+        self.chained += 1
+        return self.logits[token].clone()
+
+    def rollback(self, n):
+        assert n == self.chained, "rollback must drop exactly the chain positions"
+        self.length -= n
+        self.chained = 0
+
+
+def make_tables(seed: int, scale: float = 2.0, noise: float = 1.0):
+    g = torch.Generator().manual_seed(seed)
+    target = torch.randn(V, V, generator=g) * scale
+    draft = target + torch.randn(V, V, generator=g) * noise
+    hidden = torch.randn(V, H, generator=g)
+    return target, draft, hidden
+
+
+def greedy_reference(logits: torch.Tensor, prompt: list[int], max_new: int, stop: set[int]) -> list[int]:
+    out: list[int] = []
+    last = prompt[-1]
+    while len(out) < max_new:
+        last = int(logits[last].argmax())
+        out.append(last)
+        if last in stop:
+            break
+    return out
+
+
+@pytest.mark.parametrize("k", [1, 2, 3, 4])
+@pytest.mark.parametrize("seed", [0, 1, 2])
+def test_greedy_speculation_matches_plain_greedy(k, seed):
+    target_logits, draft_logits, hidden = make_tables(seed)
+    target = TableTarget(target_logits, hidden)
+    draft = TableDraft(draft_logits, target)
+    prompt = [1, 2, 3]
+    stop = {4}
+    generated, stats = run_speculative(target, draft, prompt, k, SamplingParams(), 40, frozenset(stop), None)
+    assert generated == greedy_reference(target_logits, prompt, 40, stop)
+    assert stats.tokens == len(generated)
+    # Each step emits accepted + 1 tokens, except a final step cut short by the stop token or the length cap.
+    assert 0 <= stats.steps + stats.accepted_total - (len(generated) - 1) <= 1
+    # Greedy decisions have accept_prob 0 or 1, so both estimators agree where a position was reached at all.
+    for rate, expected in zip(stats.position_rates(), stats.position_expected_rates(), strict=True):
+        assert rate == expected or (math.isnan(rate) and math.isnan(expected))
+    assert all(stats.evaluated[j] >= stats.accepted[j] for j in range(k))
+    assert all(stats.evaluated[j + 1] <= stats.accepted[j] for j in range(k - 1))
+    assert sum(stats.accepted) == stats.accepted_total
+
+
+def test_perfect_draft_accepts_everything():
+    target_logits, _, hidden = make_tables(0)
+    target = TableTarget(target_logits, hidden)
+    draft = TableDraft(target_logits.clone(), target)
+    generated, stats = run_speculative(target, draft, [0, 1], 3, SamplingParams(), 31, frozenset(), None)
+    assert generated == greedy_reference(target_logits, [0, 1], 31, set())
+    assert stats.accepted == stats.evaluated
+    assert stats.position_rates() == [1.0, 1.0, 1.0]
+    assert stats.tokens_per_step == pytest.approx(1 + stats.accepted_total / stats.steps)
+
+
+def test_hopeless_draft_accepts_nothing():
+    target_logits, _, hidden = make_tables(1)
+    # The draft's argmax is always one token off the target's.
+    draft_logits = torch.zeros(V, V)
+    for t in range(V):
+        draft_logits[t, (int(target_logits[t].argmax()) + 1) % V] = 1.0
+    target = TableTarget(target_logits, hidden)
+    draft = TableDraft(draft_logits, target)
+    generated, stats = run_speculative(target, draft, [2], 4, SamplingParams(), 20, frozenset(), None)
+    assert generated == greedy_reference(target_logits, [2], 20, set())
+    assert stats.accepted_total == 0
+    assert stats.steps == len(generated) - 1
+    assert stats.evaluated == [stats.steps, 0, 0, 0]
+
+
+def test_sampling_probs():
+    logits = torch.tensor([3.0, 1.0, 0.5, -2.0, 2.0])
+    greedy = sampling_probs(logits, SamplingParams())
+    assert greedy.tolist() == [1.0, 0.0, 0.0, 0.0, 0.0]
+    plain = sampling_probs(logits, SamplingParams(1.0, 1.0))
+    assert torch.allclose(plain, torch.softmax(logits, -1))
+    scaled = sampling_probs(logits, SamplingParams(0.7, 1.0))
+    assert torch.allclose(scaled, torch.softmax(logits / 0.7, -1))
+    # top_p keeps the smallest prefix reaching the mass, the crossing token included: 0.754 then 0.181.
+    nucleus = sampling_probs(logits, SamplingParams(0.7, 0.9))
+    assert (nucleus > 0).tolist() == [True, False, False, False, True]
+    assert nucleus.sum().item() == pytest.approx(1.0)
+    assert (nucleus[0] / nucleus[4]).item() == pytest.approx((scaled[0] / scaled[4]).item())
+    assert (sampling_probs(logits, SamplingParams(0.7, 0.5)) > 0).tolist() == [True, False, False, False, False]
+
+
+def test_accept_draft_is_deterministic_for_one_hot():
+    p = torch.tensor([0.0, 1.0, 0.0])
+    q_same = torch.tensor([0.0, 1.0, 0.0])
+    q_other = torch.tensor([0.0, 0.0, 1.0])
+    for seed in range(5):
+        g = torch.Generator().manual_seed(seed)
+        assert accept_draft(1, q_same, p, g).accepted
+        decision = accept_draft(2, q_other, p, g)
+        assert not decision.accepted and decision.replacement == 1 and decision.accept_prob == 0.0
+
+
+def test_accept_draft_preserves_target_distribution():
+    g = torch.Generator().manual_seed(0)
+    p = torch.tensor([0.5, 0.2, 0.2, 0.1])
+    q = torch.tensor([0.1, 0.4, 0.3, 0.2])
+    n = 20000
+    counts = torch.zeros(4)
+    accepted = 0
+    for _ in range(n):
+        d = int(torch.multinomial(q, 1, generator=g))
+        decision = accept_draft(d, q, p, g)
+        counts[d if decision.accepted else decision.replacement] += 1
+        accepted += decision.accepted
+    assert (counts / n - p).abs().max() < 0.015
+    assert abs(accepted / n - torch.minimum(p, q).sum().item()) < 0.015
+
+
+def exact_two_token_distribution(logits: torch.Tensor, first: int, params: SamplingParams) -> dict:
+    p1 = sampling_probs(logits[first], params)
+    return {
+        (x1, x2): (p1[x1] * sampling_probs(logits[x1], params)[x2]).item()
+        for x1, x2 in itertools.product(range(V), repeat=2)
+    }
+
+
+def total_variation(a: dict, b: dict) -> float:
+    return 0.5 * sum(abs(a.get(key, 0.0) - b.get(key, 0.0)) for key in set(a) | set(b))
+
+
+@pytest.mark.parametrize("k", [1, 2, 3])
+def test_speculative_sampling_preserves_target_distribution(k):
+    target_logits, draft_logits, hidden = make_tables(3, scale=1.5, noise=1.5)
+    params = SamplingParams(temperature=0.8, top_p=0.9)
+    prompt = [2]
+    exact = exact_two_token_distribution(target_logits, prompt[-1], params)
+    # The draft must disagree with the target enough that accepting its tokens blindly would be visible.
+    assert total_variation(exact, exact_two_token_distribution(draft_logits, prompt[-1], params)) > 0.15
+    runs = 3000
+    counts: dict = {}
+    for r in range(runs):
+        target = TableTarget(target_logits, hidden)
+        draft = TableDraft(draft_logits, target)
+        g = torch.Generator().manual_seed(r)
+        generated, _ = run_speculative(target, draft, prompt, k, params, 2, frozenset(), g)
+        assert len(generated) == 2
+        counts[tuple(generated)] = counts.get(tuple(generated), 0) + 1
+    empirical = {key: n / runs for key, n in counts.items()}
+    assert total_variation(empirical, exact) < 0.08
+
+
+@pytest.mark.parametrize("k", [1, 2, 4])
+def test_sampling_run_bookkeeping_with_stop_token(k):
+    target_logits, draft_logits, hidden = make_tables(5)
+    target = TableTarget(target_logits, hidden)
+    draft = TableDraft(draft_logits, target)
+    g = torch.Generator().manual_seed(1)
+    generated, stats = run_speculative(
+        target, draft, [0, 3], k, SamplingParams(0.7, 0.95), 25, frozenset({1}), g
+    )
+    assert stats.prompts == 1 and stats.tokens == len(generated) <= 25
+    if 1 in generated:
+        assert generated.index(1) == len(generated) - 1
+    assert 0 <= stats.steps + stats.accepted_total - (len(generated) - 1) <= 1
+    assert draft.extends >= 1
diff --git a/test/rocm/test_attn_prologue.py b/test/rocm/test_attn_prologue.py
new file mode 100644
index 0000000000..ad60b664c6
--- /dev/null
+++ b/test/rocm/test_attn_prologue.py
@@ -0,0 +1,101 @@
+"""``mstar::qkv_norm_rope_cache`` on AMD GPUs (AITER's fused kernel) against its
+torch reference at Qwen3.5's decode shapes: q, gate, k, v and the cache rows it
+writes, and a captured replay that writes the next step's slots."""
+
+import pytest
+import torch
+
+from mstar.engine.resources.attn import aiter as _aiter_attention  # noqa: F401  registers the op
+from mstar.engine.resources.attn import prologue
+from mstar.platforms import current_platform
+
+pytestmark = pytest.mark.skipif(not current_platform.is_rocm(), reason="needs an AMD GPU")
+
+DEV = "cuda"
+PAGE = 128
+
+
+def _case(tokens, heads, kv_heads, head_dim=256, half=32, pages=64, seed=0):
+    g = torch.Generator(device=DEV).manual_seed(seed)
+    qkv = torch.randn(tokens, (2 * heads + 2 * kv_heads) * head_dim, device=DEV, generator=g).to(torch.bfloat16)
+    qw = (torch.randn(head_dim, device=DEV, generator=g) * 0.1).to(torch.bfloat16)
+    kw = (torch.randn(head_dim, device=DEV, generator=g) * 0.1).to(torch.bfloat16)
+    inv = 1.0 / (1e7 ** (torch.arange(0, 2 * half, 2, device=DEV).float() / (2 * half)))
+    freqs = inv[None, :] * torch.arange(32768, device=DEV).float()[:, None]
+    cos, sin = freqs.cos().to(torch.bfloat16), freqs.sin().to(torch.bfloat16)
+    positions = torch.randint(0, 32768, (tokens,), device=DEV, generator=g)
+    slots = torch.randperm(pages * PAGE, device=DEV, generator=g)[:tokens]
+    cache = torch.zeros(pages, 2, PAGE, kv_heads, head_dim, device=DEV, dtype=torch.bfloat16)
+    return qkv, qw, kw, cos, sin, positions, cache, slots
+
+
+def _run(fn, qkv, qw, kw, cos, sin, positions, cache, slots, heads, kv_heads, head_dim=256):
+    out = fn(qkv, qw, kw, cos, sin, positions, cache[:, 0], cache[:, 1], slots, heads, kv_heads, head_dim, 1e-6, True)
+    return out, cache
+
+
+@pytest.mark.parametrize("tokens", [1, 4, 64, 128])
+@pytest.mark.parametrize("heads, kv_heads", [(16, 4), (8, 2)])
+def test_matches_the_reference(tokens, heads, kv_heads):
+    qkv, qw, kw, cos, sin, positions, cache, slots = _case(tokens, heads, kv_heads)
+    impl = prologue.QKV_NORM_ROPE_CACHE.implementation()
+    got, got_cache = _run(impl, qkv, qw, kw, cos, sin, positions, cache.clone(), slots, heads, kv_heads)
+    want, want_cache = _run(prologue._reference, qkv, qw, kw, cos, sin, positions, cache.clone(), slots,
+                            heads, kv_heads)
+    for name, a, b in zip(("q", "gate", "k", "v"), got, want, strict=True):
+        # the kernel rounds the normalised q once before rotating: a bf16 ulp or so
+        torch.testing.assert_close(a.float(), b.float(), rtol=1.6e-2, atol=1.6e-2, msg=name)
+    torch.testing.assert_close(got_cache.float(), want_cache.float(), rtol=1.6e-2, atol=1.6e-2)
+    untouched = torch.ones(cache.shape[0] * PAGE, dtype=torch.bool, device=DEV)
+    untouched[slots] = False
+    assert got_cache.permute(0, 2, 1, 3, 4).reshape(-1, 2, kv_heads, 256)[untouched].abs().max() == 0
+
+
+def test_a_captured_call_writes_the_slots_it_is_replayed_with():
+    heads, kv_heads, tokens = 16, 4, 8
+    qkv, qw, kw, cos, sin, positions, cache, slots = _case(tokens, heads, kv_heads)
+    impl = prologue.QKV_NORM_ROPE_CACHE.implementation()
+    stream = torch.cuda.Stream()
+    with torch.cuda.stream(stream):
+        _run(impl, qkv, qw, kw, cos, sin, positions, cache, slots, heads, kv_heads)
+    torch.cuda.synchronize()
+    graph = torch.cuda.CUDAGraph()
+    with torch.cuda.graph(graph, stream=stream):
+        _run(impl, qkv, qw, kw, cos, sin, positions, cache, slots, heads, kv_heads)
+    cache.zero_()
+    fresh = torch.arange(100, 100 + tokens, device=DEV)
+    slots.copy_(fresh)
+    graph.replay()
+    torch.cuda.synchronize()
+    want, want_cache = _run(prologue._reference, qkv, qw, kw, cos, sin, positions,
+                            torch.zeros_like(cache), fresh, heads, kv_heads)
+    torch.testing.assert_close(cache.float(), want_cache.float(), rtol=1.6e-2, atol=1.6e-2)
+
+
+@pytest.mark.parametrize("keys", [40, 300, 5000, 20000])
+@pytest.mark.parametrize("rows", [1, 4, 64])
+def test_the_decode_segment_pin_reads_every_key(rows, keys):
+    """A row past the pin is covered by wider segments, not cut: the same answer as the cache's whole context."""
+    from mstar.engine.resources.attn.aiter import DECODE_SEGMENT_PIN
+
+    if rows * keys > 400000:
+        pytest.skip("more KV than this test needs")
+    g = torch.Generator(device=DEV).manual_seed(rows + keys)
+    pages = rows * (keys // PAGE + 1) + 2
+    cache = torch.randn(pages, 2, PAGE, 4, 256, device=DEV, generator=g).to(torch.bfloat16)
+    q = torch.randn(rows, 8, 256, device=DEV, generator=g).to(torch.bfloat16)
+    lens = torch.full((rows,), keys, dtype=torch.int32, device=DEV)
+    table = torch.zeros(rows, pages + 2, dtype=torch.int32, device=DEV)
+    for r in range(rows):
+        n = keys // PAGE + 1
+        table[r, :n] = torch.arange(1 + r * n, 1 + (r + 1) * n, dtype=torch.int32, device=DEV)
+    qptr = torch.arange(rows + 1, dtype=torch.int32, device=DEV)
+    outs = []
+    for pin in (32768, DECODE_SEGMENT_PIN):
+        out = torch.empty_like(q)
+        torch.ops.mstar.aiter_unified_attention(
+            q, cache[:, 0], cache[:, 1], out, qptr, lens, table, 1, pin, 256 ** -0.5,
+        )
+        outs.append(out.float())
+    torch.cuda.synchronize()
+    torch.testing.assert_close(outs[1], outs[0], rtol=1e-2, atol=1e-2)
diff --git a/test/rocm/test_conv_decode.py b/test/rocm/test_conv_decode.py
new file mode 100644
index 0000000000..b66b8a5c8a
--- /dev/null
+++ b/test/rocm/test_conv_decode.py
@@ -0,0 +1,104 @@
+"""``causal_conv1d_decode`` against ``causal_conv1d_update``, its general kernel: output and state bit-equal at decode's
+shapes, null slots and strided rows included, eager and from a captured replay."""
+
+import pytest
+import torch
+
+from mstar.platforms import current_platform
+from mstar.utils.causal_conv1d import causal_conv1d_decode, causal_conv1d_update
+
+pytestmark = pytest.mark.skipif(not current_platform.is_rocm(), reason="needs an AMD GPU")
+
+DEV = "cuda"
+DT = torch.bfloat16
+
+
+def _case(rows, dim, seed=0, gap=0):
+    g = torch.Generator(device=DEV).manual_seed(seed)
+    wide = torch.randn(rows, dim + gap, device=DEV, generator=g).to(DT)
+    state = torch.randn(rows + 3, dim, 3, device=DEV, generator=g).to(DT)
+    weight = (torch.randn(dim, 4, device=DEV, generator=g) * 0.5).to(DT)
+    bias = (torch.randn(dim, device=DEV, generator=g) * 0.1).to(DT)
+    slots = torch.randperm(rows + 2, device=DEV, generator=g).to(torch.int32) + 1
+    slots = slots[:rows]
+    return wide, state, weight, bias, slots
+
+
+@pytest.mark.parametrize("rows", [1, 2, 3, 4, 8, 16, 32, 48, 64, 128])
+@pytest.mark.parametrize("dim", [1536, 3072, 6144, 12288])
+def test_bit_equal(rows, dim):
+    wide, state, weight, bias, slots = _case(rows, dim)
+    x1, s1 = wide.clone(), state.clone()
+    x2, s2 = wide.clone(), state.clone()
+    causal_conv1d_update(x1, s1, weight, bias=bias, activation="silu", conv_state_indices=slots, null_block_id=0)
+    causal_conv1d_decode(x2, s2, weight, bias=bias, activation="silu", conv_state_indices=slots, null_block_id=0)
+    torch.cuda.synchronize()
+    assert torch.equal(x1, x2) and torch.equal(s1, s2)
+
+
+@pytest.mark.parametrize("gap", [0, 2080])
+def test_null_slots_and_strided_rows(gap):
+    rows, dim = 16, 6144
+    wide, state, weight, bias, slots = _case(rows, dim, gap=gap)
+    slots[[3, 15]] = 0
+    x1, s1 = wide.clone(), state.clone()
+    x2, s2 = wide.clone(), state.clone()
+    common = dict(bias=bias, activation="silu", conv_state_indices=slots, null_block_id=0)
+    causal_conv1d_update(x1[:, :dim], s1, weight, **common)
+    out = causal_conv1d_decode(x2[:, :dim], s2, weight, **common)
+    torch.cuda.synchronize()
+    assert torch.equal(x1, x2) and torch.equal(s1, s2)
+    # in place, and a null row keeps its input and its state
+    assert out.data_ptr() == x2.data_ptr()
+    assert torch.equal(x2[3, :dim], wide[3, :dim])
+    assert torch.equal(s2[0], state[0])
+
+
+def test_no_bias_no_activation():
+    wide, state, weight, _, slots = _case(4, 3072)
+    x1, s1 = wide.clone(), state.clone()
+    x2, s2 = wide.clone(), state.clone()
+    causal_conv1d_update(x1, s1, weight, conv_state_indices=slots, null_block_id=0)
+    causal_conv1d_decode(x2, s2, weight, conv_state_indices=slots, null_block_id=0)
+    torch.cuda.synchronize()
+    assert torch.equal(x1, x2) and torch.equal(s1, s2)
+
+
+def test_other_operands_take_the_general_kernel():
+    wide, state, weight, bias, slots = _case(4, 3072)
+    state5 = torch.randn(7, 3072, 4, device=DEV).to(DT)
+    weight5 = (torch.randn(3072, 5, device=DEV) * 0.5).to(DT)
+    x1, s1 = wide.clone(), state5.clone()
+    x2, s2 = wide.clone(), state5.clone()
+    causal_conv1d_update(x1, s1, weight5, bias=bias, activation="silu", conv_state_indices=slots, null_block_id=0)
+    causal_conv1d_decode(x2, s2, weight5, bias=bias, activation="silu", conv_state_indices=slots, null_block_id=0)
+    torch.cuda.synchronize()
+    assert torch.equal(x1, x2) and torch.equal(s1, s2)
+
+
+def test_captured_replay_matches_eager_steps():
+    rows, dim = 64, 6144
+    source, state, weight, bias, slots = _case(rows, dim)
+    x1, s1 = source.clone(), state.clone()
+    x2, s2 = source.clone(), state.clone()
+    buf = source.clone()
+
+    def step():
+        buf.copy_(source)
+        common = dict(bias=bias, activation="silu", conv_state_indices=slots, null_block_id=0)
+        return causal_conv1d_decode(buf, s2, weight, **common)
+
+    step()
+    s2.copy_(state)
+    torch.cuda.synchronize()
+    graph = torch.cuda.CUDAGraph()
+    with torch.cuda.graph(graph):
+        step()
+    s2.copy_(state)
+    torch.cuda.synchronize()
+    for _ in range(3):
+        x1.copy_(source)
+        causal_conv1d_update(x1, s1, weight, bias=bias, activation="silu", conv_state_indices=slots, null_block_id=0)
+        graph.replay()
+        torch.cuda.synchronize()
+        assert torch.equal(x1, buf) and torch.equal(s1, s2)
diff --git a/test/rocm/test_spec_decode_rocm.py b/test/rocm/test_spec_decode_rocm.py
new file mode 100644
index 0000000000..74f1878ee6
--- /dev/null
+++ b/test/rocm/test_spec_decode_rocm.py
@@ -0,0 +1,420 @@
+"""A speculative verify step's device side on an AMD GPU.
+
+* GDN through the pool and the manager as the engine drives them: a verify row
+  of k + 1 tokens starts from the slot its last step accepted up to and leaves
+  the state after every token in its own slot, against the fp32 reference run
+  token by token, including a step that starts after a rejected tail.
+* The step's device work (the speculative resource's in-graph operations, the
+  GDN verify kernels and a whole draft pass's attention) issues no host sync:
+  run under ``torch.cuda.set_sync_debug_mode("error")``.
+* A captured step follows the device registers: replayed after the registers
+  move, its positions, KV slots and lengths move with them, with no re-plan.
+"""
+
+from __future__ import annotations
+
+from types import SimpleNamespace
+
+import pytest
+import torch
+
+from mstar.engine.resources.attn.aiter import AiterPagedAttentionManager, _Staged
+from mstar.engine.resources.base import EngineResourceInfo
+from mstar.engine.resources.kv.manager import KVPlanState
+from mstar.engine.resources.linear_attn import (
+    LinearAttnConfig,
+    LinearAttnManager,
+    LinearAttnSpec,
+    LinearAttnStep,
+    LinearAttnVariant,
+)
+from mstar.engine.resources.recurrent import (
+    DeltaNetGeometry,
+    RecurrentStateConfig,
+    RecurrentStatePool,
+    RecurrentStateSpec,
+    RecurrentStep,
+)
+from mstar.engine.resources.spec_decode.config import SpecPhase
+from mstar.engine.resources.spec_decode.draft_head import Fp8DraftHead
+from mstar.engine.resources.spec_decode.resource import SpecDecodePlan, SpecDecodeResource
+from mstar.engine.resources.step import Segment, StepContext
+from mstar.platforms import current_platform
+
+pytestmark = pytest.mark.skipif(
+    not (torch.cuda.is_available() and current_platform.is_rocm()),
+    reason="AITER's GDN kernels need an AMD GPU",
+)
+
+DEV = torch.device("cuda")
+POOL, BACKEND, SPEC = "gdn_state", "linear_attn", "spec"
+GEO = DeltaNetGeometry(num_k_heads=4, num_v_heads=8, head_k_dim=128, head_v_dim=128, conv_kernel_size=4)
+OUT_TOL = dict(atol=2e-2, rtol=2e-2)
+STATE_TOL = dict(atol=1e-2, rtol=1e-2)
+PAGE = 16
+
+
+@pytest.fixture(autouse=True)
+def exact_fp32_matmul():
+    precision = torch.get_float32_matmul_precision()
+    torch.set_float32_matmul_precision("highest")
+    yield
+    torch.set_float32_matmul_precision(precision)
+
+
+def l2norm(x):
+    x = x.float()
+    return x * torch.rsqrt((x * x).sum(-1, keepdim=True) + 1e-6)
+
+
+def reference(q, k, v, a, b, a_log, dt_bias, h0, scale):
+    """Token by token from ``h0``; the output and the state after each token."""
+    rep = v.shape[1] // q.shape[1]
+    g = -torch.exp(a_log.float()) * torch.nn.functional.softplus(a.float() + dt_bias.float())
+    beta = torch.sigmoid(b.float())
+    qn, kn = l2norm(q).repeat_interleave(rep, 1), l2norm(k).repeat_interleave(rep, 1)
+    h = h0.float().clone()
+    outs, states = [], []
+    for t in range(v.shape[0]):
+        h = h * torch.exp(g[t])[:, None, None]
+        delta = (v[t].float() - torch.einsum("hvk,hk->hv", h, kn[t])) * beta[t][:, None]
+        h = h + delta[:, :, None] * kn[t][:, None, :]
+        outs.append(torch.einsum("hvk,hk->hv", h, qn[t] * scale))
+        states.append(h.clone())
+    return torch.stack(outs), torch.stack(states)
+
+
+def build(k: int, rows: int):
+    spec = RecurrentStateSpec(POOL, {"llm"}, RecurrentStateConfig(
+        num_layers=1, blocks=GEO.to_blocks(), max_slots=1 + rows * (k + 1), spec_tokens=k,
+    ))
+    pool = RecurrentStatePool.build(spec, EngineResourceInfo(device=DEV))
+    manager = LinearAttnManager.build(
+        LinearAttnSpec(BACKEND, {"llm"}, LinearAttnConfig(
+            recurrent_state=POOL, variant=LinearAttnVariant.GDN, spec_decode=SPEC,
+        )),
+        EngineResourceInfo(device=DEV, dependencies={POOL: spec}),
+    )
+    # once per request, as the worker ingests: a second ingest starts the request over
+    for r in range(rows):
+        pool.ingest_request(r)
+    return pool, manager
+
+
+def inputs(tokens: int, seed: int):
+    gen = torch.Generator(device=DEV).manual_seed(seed)
+
+    def randn(*shape):
+        return torch.randn(*shape, generator=gen, device=DEV).to(torch.bfloat16)
+
+    return dict(q=randn(tokens, GEO.num_k_heads, GEO.head_k_dim), k=randn(tokens, GEO.num_k_heads, GEO.head_k_dim),
+                v=randn(tokens, GEO.num_v_heads, GEO.head_v_dim), a=randn(tokens, GEO.num_v_heads),
+                b=randn(tokens, GEO.num_v_heads))
+
+
+def step(pool, manager, rids, span, num_accepted=None):
+    segs = tuple(Segment(r, "main", span) for r in rids)
+    ctx = StepContext(request_ids=rids, graph_walk="decode" if num_accepted is not None else "prefill",
+                      slot=0, capture=False)
+    rstep = RecurrentStep(segments=segs)
+    assert pool.admit(rstep, ctx).ok
+    ctx.plan_results[POOL] = pool.plan(rstep, ctx)
+    if num_accepted is not None:
+        ctx.plan_results[SPEC] = SimpleNamespace(phase=SpecPhase.VERIFY, num_accepted=num_accepted)
+    manager.plan(LinearAttnStep(segments=segs), ctx)
+    pool.commit(rstep, ctx)
+
+
+def run(manager, pool, x, w):
+    return manager.run(x["q"], x["k"], x["v"], x["a"], x["b"], state_layer=pool.block("state", 0),
+                       a_log=w[0], dt_bias=w[1])
+
+
+@pytest.mark.parametrize("k", [1, 3, 4])
+def test_verify_rows_resume_from_the_accepted_slot(k):
+    rows, t, prompt = 3, k + 1, 9
+    pool, manager = build(k, rows)
+    rids = list(range(rows))
+    gen = torch.Generator(device=DEV).manual_seed(5)
+    w = (torch.randn(GEO.num_v_heads, generator=gen, device=DEV) * 0.5,
+         torch.randn(GEO.num_v_heads, generator=gen, device=DEV).to(torch.bfloat16).float())
+    scale = GEO.head_k_dim ** -0.5
+
+    # the prompts land in each request's first slot, as a prefill leaves them
+    xp = inputs(rows * prompt, seed=1)
+    step(pool, manager, rids, prompt)
+    run(manager, pool, xp, w)
+    state = pool.block("state", 0)
+    groups = [pool._slots[r]["main"].group for r in rids]
+    starts = [state[g[0]].clone() for g in groups]
+
+    accepted = torch.ones(rows, dtype=torch.int32, device=DEV)
+    for round_ in range(2):
+        x = inputs(rows * t, seed=10 + round_)
+        step(pool, manager, rids, t, num_accepted=accepted)
+        out = run(manager, pool, x, w)
+        torch.cuda.synchronize()
+        next_starts = []
+        for r, g in enumerate(groups):
+            sl = slice(r * t, (r + 1) * t)
+            ref_out, ref_states = reference(x["q"][sl], x["k"][sl], x["v"][sl], x["a"][sl], x["b"][sl], *w,
+                                            starts[r], scale)
+            torch.testing.assert_close(out[sl].float(), ref_out, **OUT_TOL)
+            for i in range(t):
+                torch.testing.assert_close(state[g[i]].float(), ref_states[i], **STATE_TOL)
+            # the next step accepted a row-dependent prefix of this one
+            n = (r + round_) % t
+            next_starts.append(ref_states[n])
+        accepted = torch.tensor([(r + round_) % t + 1 for r in range(rows)], dtype=torch.int32, device=DEV)
+        starts = next_starts
+
+
+def test_the_verify_kernel_replays_in_a_graph():
+    """FlyDSL's chain kernel from a captured graph: a real row as in eager, a
+    padding row (slot -1) skipped, its pool slots and output left alone."""
+    k, rows, t = 2, 2, 3
+    pool, manager = build(k, rows)
+    rids = list(range(rows))
+    w = (torch.zeros(GEO.num_v_heads, device=DEV), torch.zeros(GEO.num_v_heads, device=DEV))
+    step(pool, manager, rids, 5)
+    run(manager, pool, inputs(rows * 5, seed=1), w)
+    accepted = torch.ones(rows, dtype=torch.int32, device=DEV)
+    step(pool, manager, rids, t, num_accepted=accepted)
+    state = pool.block("state", 0)
+    x = inputs(rows * t, seed=2)
+    # the second row is a padding row from here on
+    wrapper = manager._current["main"]
+    wrapper._plan_state.slot_map[1] = -1
+    groups = [pool._slots[r]["main"].group for r in rids]
+    before = state.clone()
+    eager = run(manager, pool, x, w).clone()
+    after = state[groups[0]].clone()
+    torch.cuda.synchronize()
+
+    state.copy_(before)
+    graph = torch.cuda.CUDAGraph()
+    with torch.cuda.graph(graph):
+        out = run(manager, pool, x, w)
+    state.copy_(before)
+    graph.replay()
+    torch.cuda.synchronize()
+    torch.testing.assert_close(out, eager, rtol=0, atol=0)
+    torch.testing.assert_close(state[groups[0]], after, rtol=0, atol=0)
+    assert torch.equal(state[groups[1]], before[groups[1]])
+    assert torch.isfinite(out[t:].float()).all()
+
+
+def _attention(rows: int, width: int, table: torch.Tensor) -> AiterPagedAttentionManager:
+    """The real manager's kernel call over a plan with this block table."""
+    manager = AiterPagedAttentionManager.__new__(AiterPagedAttentionManager)
+    manager._device, manager._dtype, manager._scale = DEV, torch.bfloat16, 256 ** -0.5
+    manager._row_starts, manager._retired_row_starts = None, []
+    manager._reserve_row_starts(rows)
+    manager._current_plans = {"main": _Staged(
+        manager._row_starts[: rows + 1], torch.ones(rows, dtype=torch.int32, device=DEV), table, 1, 32768, rows,
+    )}
+    return manager
+
+
+def test_one_query_a_row_is_its_slot_of_a_verify_row():
+    """A draft pass attends one query a row off the verify's block table: what
+    the slot would read in a row of k + 1 queries, up to the kernel's rounding."""
+    t, rows, heads, kv_heads, dim, page = 3, 2, 16, 4, 256, 128
+    gen = torch.Generator(device=DEV).manual_seed(3)
+    base = torch.tensor([100, 201], dtype=torch.int32, device=DEV)
+    pages = 8
+    layer = torch.randn(pages * rows, 2, page, kv_heads, dim, device=DEV, generator=gen).to(torch.bfloat16)
+    table = torch.arange(pages * rows, dtype=torch.int32, device=DEV).view(rows, pages)
+    manager = _attention(rows, pages, table)
+    multi = _Staged(
+        torch.arange(rows + 1, dtype=torch.int32, device=DEV) * t, base + t, table, t, 1024, rows * t,
+    )
+    q = torch.randn(rows * t, heads, dim, device=DEV, generator=gen).to(torch.bfloat16)
+    full = manager._unified(q, multi, layer)
+    for n in range(t):
+        keep = torch.arange(rows, device=DEV) * t + n
+        one = manager.attend_rows(q[keep], layer, base + n + 1)
+        torch.testing.assert_close(one.float(), full[keep].float(), rtol=2e-2, atol=2e-2)
+
+
+def _spec(k: int, rows: int, width: int):
+    """The speculative resource with a verify plan over fake attention and KV
+    buffers: a block table of distinct pages per row."""
+    spec = SpecDecodeResource(device=DEV, num_draft_tokens=k, max_requests=rows, page_size=PAGE,
+                              kv_key="kv", attn_key="attn", stop_token_ids=(7,), sampling=True)
+    for r in range(rows):
+        spec.ingest_request(r)
+    t = k + 1
+    image, _ = spec._image_for(StepContext(request_ids=list(range(rows)), graph_walk="decode", slot=0,
+                                           capture=False), rows)
+    image.rows[:rows].copy_(torch.tensor([spec._row_of[r] for r in range(rows)], device=DEV))
+    image.real[:rows].fill_(1)
+    table = (1 + torch.arange(rows * width, device=DEV, dtype=torch.int32)).view(rows, width)
+    kv_lens = torch.zeros(rows, dtype=torch.int32, device=DEV)
+    state = KVPlanState(token_to_page=torch.zeros(rows * t, dtype=torch.long, device=DEV),
+                        token_to_cache=torch.zeros(rows * t, dtype=torch.long, device=DEV))
+    spec._current = SpecDecodePlan(
+        phase=SpecPhase.VERIFY, num_rows=rows, image=image, block_table=table, kv_lens=kv_lens,
+        write_pages=state.token_to_page, write_offsets=state.token_to_cache,
+    )
+    regs = slice(1, rows + 1)
+    spec._seq_len[regs] = torch.tensor([30 + 17 * r for r in range(rows)], dtype=torch.int32, device=DEV)
+    spec._tokens[:, regs] = torch.arange(rows * t, device=DEV).view(t, rows) % 5
+    spec._num_acc[regs] = 1
+    spec._budget[regs] = 1 << 20
+    spec._stop_on[regs] = 1
+    return spec
+
+
+def _verify_step(spec, logits_rows, vocab, k, sample=False):
+    """Every device operation of a verify step and its draft passes, as the
+    forward issues them, with stand-ins for the layers."""
+    ids, positions = spec.verify_inputs()
+    # a target that predicts each token's successor
+    logits = logits_rows.scatter(1, ((ids + 1) % vocab)[:, None], 1.0)
+    out = spec.accept(logits, sample=sample)
+    drafted = out.target.reshape(-1)[spec.draft_slots()]
+    rows = drafted.shape[0]
+    for j in range(k):
+        spec.draft_address(j)
+        drafted = spec.record_draft(j, logits_rows[:rows].scatter(1, ((drafted + 1) % vocab)[:, None], 1.0))
+    return out
+
+
+@pytest.mark.parametrize("k", [1, 3])
+def test_a_verify_step_issues_no_host_sync(k):
+    rows, vocab = 4, 64
+    pool, manager = build(k, rows)
+    spec = _spec(k, rows, width=8)
+    rids = list(range(rows))
+    t = k + 1
+    step(pool, manager, rids, 5)
+    xp = inputs(rows * 5, seed=2)
+    w = (torch.zeros(GEO.num_v_heads, device=DEV), torch.zeros(GEO.num_v_heads, device=DEV))
+    run(manager, pool, xp, w)
+    step(pool, manager, rids, t, num_accepted=spec._current.num_accepted)
+    noise = torch.zeros(rows * t, vocab, device=DEV)
+    x = inputs(rows * t, seed=3)
+    conv_x = torch.randn(rows * t, GEO.conv_dim, device=DEV).to(torch.bfloat16)
+    conv_w = torch.randn(GEO.conv_dim, GEO.conv_kernel_size, device=DEV).to(torch.bfloat16)
+    head_dim, heads, kv_heads = 256, 16, 4
+    layer = torch.zeros(64, 2, PAGE, kv_heads, head_dim, dtype=torch.bfloat16, device=DEV)
+    spec._kv = SimpleNamespace(layer_view=lambda idx: layer)
+    spec._attn = _attention(rows, width=8, table=spec._current.block_table)
+    packed = (rows * t, (2 * heads + 2 * kv_heads) * head_dim)
+    qkv = [torch.randn(*packed, device=DEV).to(torch.bfloat16),
+           torch.randn(rows, packed[1], device=DEV).to(torch.bfloat16)]
+    weight = torch.zeros(head_dim, dtype=torch.bfloat16, device=DEV)
+    freqs = torch.arange(1024, device=DEV).float()[:, None] * torch.linspace(0, 1, 32, device=DEV)[None, :]
+    cos, sin = freqs.cos().to(torch.bfloat16), freqs.sin().to(torch.bfloat16)
+    layout = (heads, kv_heads, head_dim, 1e-6, True, 0)
+
+    def step_with_attention():
+        out = _verify_step(spec, noise, vocab, k)
+        for j in range(k):
+            spec.attend_draft(j, qkv[min(j, 1)], weight, weight, cos, sin, *layout)
+        return out
+
+    # compile the kernels before the tripwire: a first launch syncs
+    step_with_attention()
+    torch.cuda.synchronize()
+    spec._seq_len[1:rows + 1] = torch.tensor([30 + 17 * r for r in range(rows)], dtype=torch.int32, device=DEV)
+    spec._num_acc[1:rows + 1] = 1
+    torch.cuda.synchronize()
+    torch.cuda.set_sync_debug_mode("error")
+    try:
+        step_with_attention()
+        manager.run_conv(conv_x, conv_layer=pool.block("conv", 0), weight=conv_w)
+        run(manager, pool, x, w)
+    finally:
+        torch.cuda.set_sync_debug_mode("default")
+    torch.cuda.synchronize()
+
+
+@pytest.mark.parametrize("rows", [1, 2, 4])
+def test_the_fp8_draft_head_issues_no_host_sync(rows):
+    head = Fp8DraftHead(torch.randn(2048, 1024, device=DEV).to(torch.bfloat16))
+    x = torch.randn(rows, 1024, device=DEV).to(torch.bfloat16)
+    # compile the kernels before the tripwire: a first launch syncs
+    head(x)
+    torch.cuda.synchronize()
+    torch.cuda.set_sync_debug_mode("error")
+    try:
+        head(x)
+    finally:
+        torch.cuda.set_sync_debug_mode("default")
+    torch.cuda.synchronize()
+
+
+def test_a_captured_step_follows_the_registers():
+    k, rows, vocab, width = 3, 2, 64, 24
+    spec = _spec(k, rows, width=width)
+    t = k + 1
+    noise = torch.zeros(rows * t, vocab, device=DEV)
+    plan = spec._current
+    side = torch.cuda.Stream()
+    side.wait_stream(torch.cuda.current_stream())
+    with torch.cuda.stream(side):
+        _verify_step(spec, noise, vocab, k)
+    torch.cuda.current_stream().wait_stream(side)
+    graph = torch.cuda.CUDAGraph()
+    with torch.cuda.graph(graph):
+        out = _verify_step(spec, noise, vocab, k)
+    for base in ([100, 7], [129, 263]):
+        spec._seq_len[1:rows + 1] = torch.tensor(base, dtype=torch.int32, device=DEV)
+        spec._num_acc[1:rows + 1] = 1
+        spec._tokens[:, 1:rows + 1] = torch.tensor([[5, 6, 7, 8], [1, 2, 9, 9]], device=DEV).T
+        graph.replay()
+        torch.cuda.synchronize()
+        # The toy target predicts each token's successor: row 0's drafts are
+        # all right (the stream grows by 4) but it emits through the stop id 7
+        # only; row 1's second draft is wrong.
+        assert out.accepted.tolist() == [4, 2]
+        assert out.tokens.tolist() == [[6, 7, -1, -1], [2, 3, -1, -1]]
+        assert out.done.tolist() == [True, False]
+        assert spec._seq_len[1:rows + 1].tolist() == [base[0] + 4, base[1] + 2]
+        # draft pass j addressed slot n at P + n + j through the table
+        for j in range(k):
+            where = spec.draft_address(j)
+            for r in range(rows):
+                n = int(out.accepted[r]) - 1
+                pos = base[r] + n + j
+                assert int(where.kv_lens[r]) == pos + 1
+                if j:
+                    assert int(where.positions[r]) == pos
+                    assert int(where.slots[r]) == int(plan.block_table[r, pos // PAGE]) * PAGE + pos % PAGE
+
+
+def test_a_captured_sampled_step_draws_fresh_tokens_each_replay():
+    """The sampled verification (FlashInfer's top-k/top-p sampler with device
+    seeds and offsets) captures, draws anew on every replay, and keeps greedy
+    rows on the argmax."""
+    k, rows, vocab = 2, 2, 64
+    spec = _spec(k, rows, width=24)
+    t = k + 1
+    regs = slice(1, rows + 1)
+    # row 0 samples flat logits; row 1 is greedy
+    spec._temperature[regs] = torch.tensor([1.0, 0.0], device=DEV)
+    # top-k and top-p cut by FlashInfer's sampler inside the graph
+    spec._top_p[regs], spec._top_k[regs] = 0.9, 40
+    spec._seed[regs] = torch.tensor([11, 12], device=DEV)
+    flat = torch.zeros(rows * t, vocab, device=DEV)
+    side = torch.cuda.Stream()
+    side.wait_stream(torch.cuda.current_stream())
+    with torch.cuda.stream(side):
+        _verify_step(spec, flat, vocab, k, sample=True)
+    torch.cuda.current_stream().wait_stream(side)
+    graph = torch.cuda.CUDAGraph()
+    with torch.cuda.graph(graph):
+        out = _verify_step(spec, flat, vocab, k, sample=True)
+    drawn = set()
+    for _ in range(20):
+        spec._seq_len[regs] = 50
+        spec._num_acc[regs] = 1
+        spec._tokens[:, regs] = torch.tensor([[5, 6, 7], [1, 2, 3]], device=DEV).T
+        graph.replay()
+        torch.cuda.synchronize()
+        drawn.add(int(out.target[0, 0]))
+        # the greedy row's target is each token's successor, so it accepts all three
+        assert out.target[1].tolist() == [2, 3, 4] and int(out.accepted[1]) == 3
+    assert len(drawn) > 5, drawn