Speculative Decoding (D5-R) — Plan

Status: D5-R CLOSED (2026-09-12, docs 83–88): all stages landed or closed by measurement. 14B d=2 = 1.42×/1.68× (95%/88% of llama's absolute speed, same-window); d=8 retired (0.63× e2e); graph-capture prize verified already banked (R3-B); open leads localized (doc 89): row marginal = q4_K 2.4 + q6_K 1.0 + norm/quant 0.7 ms/row; fix menu re-priced by measurement (doc 90): R-rows-per-block CLOSED (~0 at d=2 — chain gives act rows L2-hot; only small-M tensor-core mma remains as a lever of size), chain hygiene de-priced (~0.1 ms/row); mma audit (doc 91) + K-split (doc 92): BT GEMM is tensor-core; K-split shipped and auto-ksplit enabled on the default path by decision (doc 92 §3b: default C_T(9) 86.4→73.0; multi-MMVQ keeps nt 2..8, residual is a tile-geometry redesign; d=8 needs code-p≈0.755, measured 0.731); doc 93/94: greedy identity fixed (spec output now byte-identical to sequential: batched verify attention made bitwise position-invariant + the spec penalty window capped to repeat_last_n) and the d=8 door crossed with the q4_k_m draft (code 76.5% ≥ 0.755: d=8 = 46.6 tok/s vs d=2 43.6, 96% of llama; prose stays d=2 — deep acceptance collapses); doc 95: adaptive draft depth shipped (--spec-draft-adaptive) — per-round d from beta-smoothed per-depth acceptance + min-of-4 cost windows, 10% switch hysteresis, optimism for unseen depths; identity boundary pinned: verify nt ≤ 8 bitwise (multi-MMVQ), nt=9 (BT-MMQ lm_head) tolerance-class → adaptive capped at d=7; all four gates pass and adaptive beats the best static on both code cells (46.5/48.4 vs 43.9/46.6 tok/s — the true optimum dmean ≈ 3.5 was never in the static sweep, which a CLI bug had silently pinned to d=8); doc 96 Phase 0: the nt=9 verify path profiled — BT-MMQ = 84% of GPU time, smem scoreboard stalls 40–59% inside the BT kernels (stop-gate ≥30% → tile/staging work may proceed, re-priced separately; the bitwise-preserving lever is cp.async double-buffered staging); doc 97: conversation/server integration LANDED — --cnv --spec-draft and serve --spec-draft are live with byte-identical identity gates (all frontends now share one spec implementation; a pre-existing plain-server cross-request KV-contamination bug was found and fixed in passing). doc 98: the 96-Phase-1 cp.async staging lever measured NULL → reverted. doc 99: the fast-verify P0/P1 pass corrected the stall attribution (L1TEX latency at a register-file-capped 16 warps/SM — not smem, not staging, not scheduling), landed one bitwise micro-fix (B-plane XOR swizzle), measured three further scheduling levers NULL, and re-priced the tolerance-class knob down to single-digit percent → not built. doc 100: the last priced lever (aligned qs plane) measured timing-NULL under interleaved A/B, and exposed a ±7–12% clock-ramp band that invalidates sequential benchmark comparisons on dgxspark. The performance line stays closed — now with the mechanism mapped and the measurement method corrected (doc 101: time-budget warmups; the headline table re-measured tight — adaptive 35.9 prose / 44.8 code tok/s, 1.42×/1.77× sequential). The draft-scale sweep (doc 102) closes the last speculative-quality axis: bigger drafts lose (acceptance is bounded by the target's predictability), so the 0.5B Q4_K_M draft stays; cross-family drafts are identity-safe (4/4 byte-identical vs sequential) and the sweep flushed a latent qwen3-loader namespaced-registration bug. Multi-MMVQ verify + adaptive d (doc 95) is the final production configuration; doc 103 extends the same verify machinery to q8_0/q4_0-weighted targets (the q8_0 verify batch had been running the f32-activation kernel — 7B Q8_0 spec e2e jumps 26.8 → 68.0 tok/s, +154%); doc 104's p32 planes lift it again to 79.0 tok/s = 2.47× sequential. Speculative decoding reopened by decision after the doc 81 §4.3 errata invalidated the original closure's external pillar and doc 82 restored the batching invariant. The previous plan (closed 2026-09-10, "no loop plumbing will be built") is retired; its errors are recorded in the Appendix.

The reference study is LLAMA-CPP-SPECULATIVE-ANALYSIS.md (llama.cpp draft-simple, source-verified: speculator framework §3, draft model setup §4, drafting loop §5, verification §6, KV rollback §7, cost model §9, minfer port notes §11). Records: docs 80, 81, 82.

1. Why reopened — the corrected evidence

Every number below is measured (docs 80–82 + the doc 81 §4.3 corrected batteries):

• Doc 82 fixed the dispatch hole: verify amortization at nt=3 went 0.52× → 2.14× (14B q4_k_m, C_T(3)=56.6 ms vs C_T(1)=39.0); the batched verify is now a weights-once round, not a per-row restream.
• llama.cpp's own speculative decoding demonstrably wins on GB10: 1.53–1.60× (7B) / 1.86–2.43× (14B) with the generic Qwen2.5-0.5B q4_0 draft (the original "1.00× external anchor" was a measurement artifact — see Appendix).
• The engine-independent terms are known: per-position greedy acceptance p ≈ 0.74 (14B pair) / 0.69 (7B), draft cost C_D ≈ 2.9 ms/token — identical for both engines, as is C_T(1) (minfer 39.0 vs llama 41.8 ms on the 14B).
• Predicted minfer economics (measured components, doc 81 §4.3 addendum): d=2 → 1.34× on the 14B from plumbing alone today; d=8 → 1.09× until the verify marginal shrinks. The entire gap to llama is the batched-verify marginal (minfer 3.9/7.8 vs llama ~1.0–1.5 ms per extra token).
• The original go/no-go bar was mis-derived (doc 82 §5): it took the per-token marginal as ε. With the corrected cost model, break-even acceptance at d=2 with an efficient verify is p* ≈ 0.18–0.35 — the measured p clears it comfortably.

2. Scope

In scope (D5-R): draft-simple, greedy first

One draft GGUF (Qwen2.5-0.5B q4_0) drafting greedily for a target (Qwen2.5-14B q4_k_m primary — the 7B is marginal at d=2 — verified in one nt = d+1 batched forward through the existing forward_graph_cached primitive). d=2 is the gate configuration; d=8 the end state.

Deferred (not in D5-R)

draft-mtp / draft-eagle3 / draft-dflash (a trained draft is the obvious follow-up lever — llama.cpp natively supports all three spec types — but each needs target-side feature seams or new architectures); n-gram stretch; temp>0 target-authoritative verification; server multi-slot.

3. Mechanism — one round

With draft length d and accept count a (greedy: keep while the target's argmax equals the drafted token):

draft phase     draft model: d sequential nt==1 forwards (second GraphCache)
verification    target: ONE nt=d+1 forward, n_out=d+1 (rows: next_token +
                the d drafted candidates, positions contiguous)
accept/cut      greedy argmax per verify row; accept while equal to the
                drafted token; always emit ≥ 1 token (the deepest accepted
                row's logits = the bonus)
KV              NO rollback kernels: the target wrote d+1 KV rows in place;
                rejected slots are simply overwritten next round
                ("positions are data" — the graph design makes rollback a
                position-bookkeeping operation)
commit          append accepted tokens; next round seeds from the bonus

Greedy equivalence is the primary correctness gate, re-scoped by measurement (doc 83 §3.4): every emitted token is the target sampler's own decision on its verify row, but the verify graph (nt=d+1, Prefill kernels) differs numerically from the nt=1 decode graph by ~0.01–0.05 logits, so sub-margin tokens may flap — both streams are valid greedy chains of the same model (AGENTS rule-9 class). Exact identity is gated on (a) the accept-rule unit tests, (b) the d=0 fallback, (c) near-tie attribution — not on bit-equality across kernel assignments.

4. Design — stage ① minimal loop

Deliberately smaller than the retired plan's machinery:

• src/spec.rs — one concrete SpecEngine (no Speculator trait chain; the chain seat opens in a later stage if a second speculator type lands): holds the draft model (second load_model + its own GraphCache), runs d draft forwards + 1 verify forward per round, returns 1..=d+1 tokens plus the last row's logits.
• Vocab compatibility gate at load (kept from the D5 analysis §4.2): same vocab type, BOS/EOS, size delta ≤ 128, token-text equality.
• CLI: --spec-draft <model> + --spec-draft-n <d> (default 2), round stats (rounds / drafted / accepted / tokens-per-round) to stderr. The spec = off path stays byte-identical to today.
• Sampling: greedy argmax over the returned n_out rows (stage ① only; temp>0 chains deferred). Logits readback is nt×n_vocab×4 B ≈ 1.8 MB/round at d=2 — acceptable for the loop; on-GPU argmax is a stage ④ candidate.

5. Stages & gates

Each stage records per cuda_optimization_steps/STYLE.md (English, six sections, real numbers) — docs 83+.

6. Risks

7. Acceptance gates (campaign rules apply unchanged)

Greedy token-identity vs the non-spec path (the spec path is an optimization, never a behavior change, at temp=0); full suite green; interleaved same-window A/B medians; acceptance-rate and tokens/round instrumentation on every round; every stage documented per STYLE.md.

Appendix — what the D5 closure got wrong (2026-09-10 → 09-12)

Kept brief; details in docs 80/81 (§4.3) / 82 (§5):

1. The go/no-go bar was mis-derived. The 2.5×-amortization gate treated the per-token marginal (attention + q8 quantize + dp4a compute) as ε; the correct model is C_T(nt) ≈ weights + nt·token-work, so break-even acceptance at d=2 is far lower than the p* = 0.73 the old model printed.
2. The external anchor was a measurement artifact. The llama-cli batteries passed -md without --spec-type draft-simple (default none) — the draft was silently never loaded, so "llama.cpp also lands at 1.00×" was base-vs-base. Corrected: 1.53–2.43×.
3. The instrument's summary field was wrong (amortization computed as C_T(1)/C_T(nt), missing the nt factor; fixed 2026-09-12; doc tables were computed from raw medians and unaffected).

Net: the closure rested on a real 0.52× measurement but a miscalibrated bar and a void anchor. Doc 82 fixed the underlying dispatch hole (0.52× → 2.14×), the corrected batteries showed the strategy wins on GB10, and the campaign reopens as D5-R.

Decisions governing this document

This page is the current contract; the decisions behind it are frozen in the ADR corpus:

• ADR-0017 — Speculative decoding refuses the features its identity contract cannot carry