From dd4c2034e9d63fbef546e093e1f6ac9da11442a0 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Thu, 6 Aug 2026 07:10:51 +0000 Subject: [PATCH] record(spec-decode): enumerate full vLLM speculative-method surface; 9 gaps rowed INVENTORIED Enumerate ALL speculative-decoding methods supported by pinned vLLM 555967922 (0.26.0.dev0) from SOURCE, cross-reference our support, and row every gap. - New master inventory `.agents/specs/spec-decode-inventory.md`: 13 vLLM SpeculativeMethod strings + acceptance/draft-sample/dynamic-k config axes, each x {upstream anchor, V1 status, our status}. Key findings: mlp_speculator has NO V1 proposer (dispatch ValueError) = upstream-deprecated; dflash/dspark are REAL vLLM methods (our DFlash mirrors vLLM, not ours-beyond). - 9 new engine-matrix rows (INVENTORIED, value-ordered): SPEC-NGRAM-GPU, SPEC-SUFFIX, SPEC-EAGLE, SPEC-MTP-FAMILY, SPEC-ACCEPT-VARIANTS, SPEC-DYNAMIC, SPEC-CUSTOM-CLASS, SPEC-EXTRACT-HIDDEN, SPEC-MLP-SPECULATOR. ENGINE_ROWS 131->140; engine summary Total + area rows updated. - roadmap ROAD-V1-D3 + gap paragraph enumerate the nine; FEATURES spec rows corrected (Medusa CPU-brick->spike drift fixed, n-gram ->DONE, EAGLE->EAGLE1/3, +breadth row); BENCHMARKS spec table +ngram +breadth rows. - STATUS method-surface paragraph added, PAID by collapsing 4 superseded DeepSeek-V4 decode-lever paragraphs (dup in BENCHMARKS, superseded by the 1.14x BEATS-ds4 binding) = -5772 chars; STATUS_RATCHET 287838->282066. - Records-only, no C++. check-agent-record / doc-checkpoint / public-doc-tables / now-current / state-order all green. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode] --- .agents/NOW.md | 17 +++--- .agents/engine-matrix.md | 13 +++- .agents/roadmap_v1.md | 26 +++++++- .agents/specs/spec-decode-inventory.md | 82 ++++++++++++++++++++++++++ .agents/state.md | 47 +++++++++++++++ docs/BENCHMARKS.md | 2 + docs/FEATURES.md | 7 ++- docs/STATUS.md | 10 +--- scripts/check-agent-record.py | 2 +- scripts/check-public-doc-tables.py | 2 +- 10 files changed, 184 insertions(+), 24 deletions(-) create mode 100644 .agents/specs/spec-decode-inventory.md diff --git a/.agents/NOW.md b/.agents/NOW.md index 7871368f..5c7db2e4 100644 --- a/.agents/NOW.md +++ b/.agents/NOW.md @@ -26,8 +26,10 @@ checkpoint on `upstream/main` at `59674cf1d`. | ROW-SERVE-ASYNC-DENSE-MIRROR | **LANDED+dgx-VERIFIED** (`f9c969ae`): #31 async mirror on classic dense Qwen3; gate RED→GREEN, SACRED 184/184 | Residual: sibling scope one-liner | | MXFP4 parity goal | binding <1.0x on c2-c8 (best c8 0.953); c1 crosses to parity+. Residual = grouped-Marlin +7-9% (only real lever) + host slice (VT_LOOP_TRACE: NOT engine-loop) | NEXT: grouped→dense-direct-A E1 marlin | -In-flight branches (gated default-OFF, not pushed): `laguna-fp4proj-prod` -(fp4 opt-in), laguna bf16/legacy/pipeline-gemv, `ds4-hc-expand-fuse`. +In-flight branches (default-OFF, not pushed): `laguna-fp4proj-prod` (fp4), +laguna bf16/legacy/pipeline-gemv, `ds4-hc-expand-fuse`. Records: +`row/SPEC-DECODE-INVENTORY`: 13 vLLM spec-decode methods enumerated from +source, 9 gaps `INVENTORIED` ([inventory](specs/spec-decode-inventory.md)). ## Current gate @@ -52,12 +54,11 @@ MEASUREMENT arbitrates; distrust aggregate bytes/time and CROSS-TOOL comparisons f32-out caller) once the Laguna fix proves the mechanism. 4. **Restore `local-ai-worker`** on dgx when the GPU campaign ends (`docker update --restart=always` + `docker start`). -5. **Protocol substrate — partly done.** Claim triage DONE; `docs/STATUS.md` - under a shrink-only ratchet; roadmap compacted; `AGENTS.md` tiered. REMAINING: - anchor backfill (98 rows `SPIKE`/`ACTIVE`, need code/test anchors; 6 model rows - need a DECISION, architecture unregistered); record-era rollover BLOCKED on - `check-agent-record.py` binding `DONE` rows to `parity-ledger.md` LINE anchors - (re-anchor by ROW ID first; `state.md`/`benchmark-record.md` can roll now). +5. **Protocol substrate — partly done.** `docs/STATUS.md` shrink-ratchet; + roadmap compacted; `AGENTS.md` tiered. REMAINING: anchor backfill (98 + `SPIKE`/`ACTIVE` rows need anchors; 6 model rows need a DECISION); record-era + rollover BLOCKED on `check-agent-record.py` `DONE`->`parity-ledger.md` LINE + anchors (re-anchor by ROW ID first). **Operator/helper protocol** ([spec](specs/operator-helper-protocol.md)): roles DECLARED then MATERIALIZED diff --git a/.agents/engine-matrix.md b/.agents/engine-matrix.md index d9e1e45e..9ee5313b 100644 --- a/.agents/engine-matrix.md +++ b/.agents/engine-matrix.md @@ -41,12 +41,12 @@ forensics: roadmap_v1.md and the parity ledger. | Parallelism | 6 | 0 | 0 | 0 | 1 | 0 | 0 | 0 | 5 | | Sampling and generation | 15 | 4 | 2 | 0 | 0 | 3 | 0 | 1 | 5 | | Structured output and tools | 7 | 0 | 3 | 0 | 0 | 2 | 0 | 0 | 2 | -| Speculative decoding | 12 | 0 | 0 | 1 | 0 | 4 | 0 | 4 | 2 | +| Speculative decoding | 21 | 0 | 0 | 1 | 0 | 4 | 0 | 4 | 11 | | Serving, API, CLI, library | 22 | 10 | 2 | 1 | 0 | 2 | 2 | 1 | 4 | | LoRA and adapters | 2 | 1 | 0 | 0 | 0 | 0 | 0 | 0 | 1 | | Long context and attention | 10 | 5 | 0 | 0 | 1 | 0 | 1 | 0 | 3 | | Loading, tokenizer, config | 9 | 3 | 3 | 0 | 0 | 0 | 1 | 1 | 1 | -| **Total** | **131** | **36** | **16** | **4** | **4** | **21** | **8** | **9** | **32** | +| **Total** | **140** | **36** | **16** | **4** | **4** | **21** | **8** | **9** | **41** | ## Engine core and scheduling @@ -179,6 +179,15 @@ claims it. |---|---|---|---|---|---|---|---|---| | `SPEC-DRAFT-MODEL` | Generic model-agnostic SEPARATE draft-model proposer (`method="draft_model"`): a full smaller standalone LM runs K autoregressive greedy steps to propose K draft tokens; the target verifies in one forward and the longest-accepted-prefix is emitted. Distinct from MTP/EAGLE/DFlash — `pass_hidden_states_to_model=False`, shares NEITHER embeddings NOR lm_head with the target. **W0 spike + W1 CPU brick LANDED 2026-07-29 (`CLAIM-SPEC-DRAFT-MEDUSA`, NOT pushed):** the greedy k-step autoregressive propose (`DraftModelProposeGreedy`/`Batch`) over a `DraftLogitsFn` next-token oracle, reusing the LANDED `SPEC-REJECTION` verify/accept UNCHANGED (only the proposer is net-new, mirror of the SPEC-NGRAM shape). Unit-gated RED-first: propose->verify->accept equivalence (accepted tokens == the target's own greedy run, every draft/target (dis)agreement pattern) + full-acceptance on a matching draft (num_sampled==k+1) + the RED witness that full acceptance DEPENDS on the autoregressive feed-back (5/6 fail with feed-back dropped). `ParseSpeculativeConfigJson` accepts `"draft_model"` (requires `model` + `num_speculative_tokens`). Additive + default-inert (no runner construction; engine byte-identical with no `SpeculativeConfig`). Clean CPU `-Werror`. **RESIDUAL (W3, DGX-offline):** the real draft-model forward behind the oracle (paged KV + CUDA-graph) + e2e greedy our-ON==vLLM-ON token-exact gate + throughput speed gate. | T2 | `vllm/v1/spec_decode/draft_model.py:19` (pass_hidden_states :29, no shared embed/lm_head :108-115); propose `vllm/v1/spec_decode/llm_base_proposer.py:502-767` (`_greedy_sample` :428-438, `set_inputs_first_pass` :838-851, K-1 feed-back :682-761); config `vllm/config/speculative.py:684,692-701,1195`; runner `vllm/v1/worker/gpu_model_runner.py:604-609`; e2e `tests/v1/e2e/spec_decode/test_spec_decode.py:500-561` @ `555967922` | `include/vllm/v1/spec_decode/draft_model_proposer.h`; `src/vllm/v1/spec_decode/draft_model_proposer.cpp`; `src/vllm/config/speculative.cpp` (`draft_model` accept); reuses `src/vllm/v1/spec_decode/rejection_sampler.{h,cpp}` — anchor `src/vllm/v1/spec_decode/draft_model_proposer.cpp:27` | `tests/vllm/v1/spec_decode/test_draft_model_proposer.cpp` (6/6, 41 assertions, CPU; RED-first 5/6 fail with feed-back dropped); ledger [parity-ledger.md](parity-ledger.md) 2026-07-29 — anchor `tests/vllm/v1/spec_decode/test_draft_model_proposer.cpp:143` | [specs/draft-model-medusa-spec.md](specs/draft-model-medusa-spec.md) | `ACTIVE` | `CLAIM-SPEC-DRAFT-MEDUSA` | | `SPEC-MEDUSA` | Medusa multi-head speculator (`method="medusa"`): the target carries N extra Medusa LM heads, each predicting ONE future position from the SAME target hidden state in a single (non-autoregressive) pass; `draft_tokens = stack([argmax(head_logits)])` -> `[batch, num_heads]`, `num_speculative_tokens == num_heads`. Verify/accept is the SAME `SPEC-REJECTION` loop (linear, not tree, at this pin). **W0 spike ONLY (`CLAIM-SPEC-DRAFT-MEDUSA`, 2026-07-29):** proposer scoped in [specs/draft-model-medusa-spec.md](specs/draft-model-medusa-spec.md); deferred to W2 because its multi-head target-tap propose needs the target model's Medusa heads (a model change) a pure host brick cannot meaningfully stand up. No code yet. | T2 | `vllm/v1/spec_decode/medusa.py:18` (propose :40-58: `model(hidden)`->per-head `compute_logits`->stacked argmax); config `vllm/config/speculative.py:822,888-889`; runner `vllm/v1/worker/gpu_model_runner.py:642-645` @ `555967922` | - (spike only, not implemented) | - (W2) | [specs/draft-model-medusa-spec.md](specs/draft-model-medusa-spec.md) | `SPIKE` | `CLAIM-SPEC-DRAFT-MEDUSA` | +| `SPEC-NGRAM-GPU` | On-device n-gram proposer (`method="ngram_gpu"`): the KMP suffix match + batch propose run on GPU (pinned idx/val buffers + a device token table) rather than the CPU `SPEC-NGRAM` path. Draft-FREE, gate-model usable — the cheapest breadth gap (value #1). | T2 | `vllm/v1/spec_decode/ngram_proposer_gpu.py:217`; dispatch `vllm/v1/worker/gpu_model_runner.py:610-626` | - (enumerated, not implemented) | - | [specs/spec-decode-inventory.md](specs/spec-decode-inventory.md) | `INVENTORIED` | - | +| `SPEC-SUFFIX` | Suffix Decoding proposer (`method="suffix"`, arXiv 2411.04975): a per-prompt suffix tree over prompt+output proposes drafts (tree-depth / spec-factor / min-token-prob knobs). Draft-FREE but depends on the EXTERNAL `arctic_inference` package (lazy import). Value #2. | T2 | `vllm/v1/spec_decode/suffix_decoding.py:9,26`; dispatch `vllm/v1/worker/gpu_model_runner.py:634-635` | - | - | [specs/spec-decode-inventory.md](specs/spec-decode-inventory.md) | `INVENTORIED` | - | +| `SPEC-EAGLE` | EAGLE(1) proposer (`method="eagle"`), distinct from `SPEC-EAGLE3`: a separate small draft consuming the target's LAST hidden state (single aux tap) that runs K AR draft steps; `EagleProposer` is shared with eagle3/mtp/dspark via `use_eagle()`. Separate draft weights; the eagle-draft `max_position_embeddings` clamp is already ported (`SpeculativeConfig::MaybeOverrideDraftMaxPositionEmbeddings`) ahead of the loader. Value #3. | T2 | `vllm/v1/spec_decode/eagle.py:10`; `use_eagle()` `vllm/config/speculative.py:1324-1328`; models `vllm/model_executor/models/{llama_eagle,mistral_eagle,cohere_eagle}.py` | - | - | [specs/spec-decode-inventory.md](specs/spec-decode-inventory.md) | `INVENTORIED` | - | +| `SPEC-MTP-FAMILY` | MTP method breadth beyond Qwen3.5 (`SPEC-MTP` DONE): the canonical `mtp` method dispatched per family by `draft_model_config.hf_config.model_type`. DeepSeek-V4 nextn is already ACTIVE-W1 (weight-blocked GGUF, model row `MODEL-SPEC-deepseek-v4-deep-seek-v4-mtp`); the ~18 other family heads (glm4_moe / ernie / nemotron_h / longcat_flash / bailing_hybrid / exaone / mimo / hy_v3 / minimax_m3 / inkling / openpangu / qwen3_next / gemma4 / step3p5) are model-matrix-INVENTORIED. gemma4_mtp / step3p5_mtp take dedicated proposers. | T2 | `vllm/config/speculative.py:37-59` (MTPModelTypes); dispatch `vllm/v1/worker/gpu_model_runner.py:627-637`; models `vllm/model_executor/models/*_mtp.py` | - | - | [specs/spec-decode-inventory.md](specs/spec-decode-inventory.md) | `INVENTORIED` | - | +| `SPEC-ACCEPT-VARIANTS` | Acceptance / draft-sampling variants beyond `SPEC-REJECTION`'s greedy `standard`: `rejection_sample_method` ∈ {standard, synthetic, block} and `draft_sample_method` ∈ {greedy, probabilistic} (typical / rejection acceptance for non-greedy sampling). Cross-cuts every proposer. | T2 | `vllm/config/speculative.py:77-78,216,283`; MRV2 `vllm/v1/worker/gpu/spec_decode/rejection_sampler.py:82-91` | - | - | [specs/spec-decode-inventory.md](specs/spec-decode-inventory.md) | `INVENTORIED` | - | +| `SPEC-DYNAMIC` | Dynamic speculation length (`num_speculative_tokens_per_batch_size`): a batch-size -> k schedule that adapts draft length to load (`uses_dynamic_speculative_decoding()`). A config axis layered on top of any drafter. | T2 | `vllm/config/speculative.py:1336-1337`; `vllm/v1/spec_decode/dynamic/utils.py:7,77` | - | - | [specs/spec-decode-inventory.md](specs/spec-decode-inventory.md) | `INVENTORIED` | - | +| `SPEC-CUSTOM-CLASS` | Pluggable custom proposer (`method="custom_class"`): the proposer class path is read from `speculative_config.model` and must expose a `propose` method — serves the extensibility-first priority (drop-in third-party speculators). | T2 | `vllm/v1/spec_decode/custom_class_proposer.py:12`; dispatch `vllm/v1/worker/gpu_model_runner.py:596-599` | - | - | [specs/spec-decode-inventory.md](specs/spec-decode-inventory.md) | `INVENTORIED` | - | +| `SPEC-EXTRACT-HIDDEN` | Hidden-state extraction proposer (`method="extract_hidden_states"`): an offline / data-collection path that runs the target with aux hidden-state outputs (for training EAGLE/MTP drafts) rather than accelerating decode; model row `MODEL-SPEC-extract-hidden-states-extract-hidden-states-model`. | T2 | `vllm/v1/spec_decode/extract_hidden_states.py:29`; dispatch `vllm/v1/worker/gpu_model_runner.py:646-650` | - | - | [specs/spec-decode-inventory.md](specs/spec-decode-inventory.md) | `INVENTORIED` | - | +| `SPEC-MLP-SPECULATOR` | MLP speculator (`method="mlp_speculator"`): a small MLP head predicting K tokens. UPSTREAM-DEPRECATED at the pin — the enum still resolves it but the V1 dispatch has NO branch, so it raises `ValueError` (`gpu_model_runner.py:651-654`); V0-only. Rowed for enumeration completeness, lowest priority. | T2 | `vllm/model_executor/models/mlp_speculator.py`; enum `vllm/config/speculative.py:69,890` | - | - | [specs/spec-decode-inventory.md](specs/spec-decode-inventory.md) | `INVENTORIED` | - | ## Serving surface, CLI, and library diff --git a/.agents/roadmap_v1.md b/.agents/roadmap_v1.md index 9519ae9c..c97b4312 100644 --- a/.agents/roadmap_v1.md +++ b/.agents/roadmap_v1.md @@ -72,7 +72,7 @@ models we already ship + benchmark. Full seam map + M0–M5 W-plan: | 10a | `ROAD-V1-D1-CUDA` | **CUDA-arch derive-and-ship + community testing matrix (user-directed 2026-07-27):** the arch-additivity framework already makes a new CUDA device a near-verbatim add, so DERIVE-AND-SHIP the fast-path kernels rather than scope-and-wait — port 1:1 from vLLM (Ampere FA2/Marlin/AllSpark/scaled-mm-C2x; Hopper FA3/C3x-Sm90/Machete/MoE; datacenter-Blackwell C3x-Sm100a-tcgen05/MLA/fp4), and the fp16/non-tensor-core bodies from llama.cpp for the arches vLLM DROPS (Pascal/Volta/Turing = MORE than vLLM), build-verify (compile + `cuobjdump` SASS), and SHIP LABELED with a 3-state signal: RUNTIME-VERIFIED / DERIVED+BUILD-VERIFIED (testing-welcome) / NOT-YET-BUILDABLE (a build is never a runtime claim). Runtime-gate boards: **AGX Orin (sm_87, reachable)** + **NVIDIA Thor (Blackwell, reachable, SM to confirm on-box)** — the first two non-GB10 runtime proofs, spanning Ampere + Blackwell. DeepGEMM (runtime JIT) is the one honest NOT-YET-BUILDABLE. | [backend matrix](backend-matrix.md), [Ampere](specs/cuda-arch-ampere-fastpath.md), [Hopper+DC](specs/cuda-arch-datacenter-fastpath.md), [breadth](specs/cuda-arch-breadth-fp16.md) | 3 committed spikes (2026-07-27); backend rows moved INVENTORIED/build-only -> SPIKE; **FIRST NON-GB10 RUNTIME PROOF: Thor sm_110 portable bf16 path RUNTIME-VERIFIED (token-exact Llama-3.2-1B, `CLAIM-CUDA-SM110-RUNTIME`)**; **FIRST DERIVE-AND-SHIP BRICK: WA-1 FA2 Ampere enablement DERIVED+BUILD-VERIFIED (fa2 cell widened to 8.x; sm_87/sm_80 FA2 cubins, sm_121a SACRED 16/16 unchanged, no Ampere board ran it)**; **Turing sm_75 W1 BUILD-VERIFIED 2026-07-28 (`CLAIM-CUDA-TURING-SM75`): the bf16-WMMA TU guard landed, single-arch `75` `-Werror` 0-warn + `cuobjdump` real `sm_75` cubin, sm_121a byte-identical — DERIVED+BUILD-VERIFIED (testing-welcome), no board ran it**; **DC1 datacenter-Blackwell sm_100a NVFP4 tcgen05 GEMM BUILD-VERIFIED 2026-07-28 (`CLAIM-CUDA-SM100-NVFP4`): faithful 1:1 port of vLLM `Fp4GemmSm100` (`ArchTag=Sm100`+`KernelScheduleAuto`→tcgen05) as a NEW dedicated `cutlass-nvfp4-sm100` FEATURE-TABLE cell (100a-only), single-arch `100a` `-Werror`-equiv 0-warn + `cuobjdump` real `sm_100a` cubin, RED (HEAD 100a fp4 DISABLED), sm_121a neutral — DERIVED+BUILD-VERIFIED (testing-welcome), no B200 board ran it; native `fp4-mma` stays sm_12x-only (mma.sync≠tcgen05)**; **DC2 Hopper sm_90a CUTLASS C3x FP8 scaled-mm BUILD-VERIFIED 2026-07-28 (`CLAIM-CUDA-SM90-C3X`): faithful 1:1 port of vLLM `cutlass_3x_gemm_sm90_fp8` (`ArchTag=Sm90`+`KernelTmaWarpSpecialized*FP8FastAccum`→wgmma/TMA) as a NEW dedicated `scaledmm-c3x-sm90` FEATURE-TABLE cell (90a-only), single-arch `90a` `-Werror` 0-warn + `cuobjdump` real `sm_90a` cubin (ptxas C7510 names `wgmma.mma_async`), RED (HEAD had no cell), sm_121a neutral (`cutlass-fp8` unchanged) — DERIVED+BUILD-VERIFIED (testing-welcome), no Hopper board ran it; sm90 int8/blockwise + sm100 C3x legs are residual bricks**; **DC3 datacenter-Blackwell sm_100a CUTLASS C3x FP8 scaled-mm tcgen05 GEMM BUILD-VERIFIED 2026-07-28 (`CLAIM-CUDA-SM100-C3X`): the intersection of DC1's tcgen05 arch + DC2's C3x fp8 kernel — a faithful 1:1 port of vLLM `cutlass_3x_gemm_sm100_fp8` (`ArchTag=Sm100`+`KernelScheduleAuto`→tcgen05, `sm100_fp8_config_{default,M256,M64}`, 2SM `<_2,_2,_1>` default) as a NEW dedicated `scaledmm-c3x-sm100` FEATURE-TABLE cell (100a-only), single-arch `100a` `-Werror all-warnings` 0-warn EXIT=0 + `cuobjdump` real `sm_100a` cubin (`Sm100TmaUmma`/`SM100_MMA_F8F6F4_2x1SM`/`TMEM` symbols, `LDTM`/`tmem` SASS), RED (HEAD had no cell), sm_121a neutral (`cutlass-fp8`/`cutlass-nvfp4`/`cutlass-nvfp4-sm100` unchanged) — DERIVED+BUILD-VERIFIED (testing-welcome), no B200 board ran it; sm100 int8/blockwise + MoE/MXFP4/MLA legs are residual bricks**; Volta/Pascal need CUDA<13 | ACTIVE | **SECOND NON-GB10 RUNTIME PROOF: Orin sm_87 portable bf16 SYNC path RUNTIME-VERIFIED 2026-07-28 (`CLAIM-CUDA-ORIN-SM87-RUNTIME`)** — built portable-only for sm_87 on a real Jetson AGX Orin (CUDA-13 container BLOCKED by the 12.6 driver → l4t-jetpack r36.4.0 nvcc 12.6 + g++-13; all fast paths EMPTY for [87]), and `test_llama_paged_engine` (Llama-3.2-1B bf16) is **13/16 strict token-exact vs the vLLM 0.25.0 oracle, 16/16 near-tie gate, 0 divergent** on the sm_87 GPU (exceeds Thor's 12/16), + test_cuda_backend/test_cuda_ops 461 assertions of real on-device kernels; ONE honest sm_87 bug found — the DEFAULT async runner crashes (`cudaFree illegal memory access`), SYNC path (`VT_ASYNC_RUNNER=0`) is the verified one (unblock item). Residual Orin W-plan: fix async runner -> GGUF leg -> FA2 (needs cutlass) -> Marlin-int4 -> vs llama.cpp; Thor fp8/fp4 fast paths once SSH-reachable | | 10b | `ROAD-V1-D1-GDN-AOT` | **Per-arch Triton-AOT GDN packed-decode cubins (user-directed 2026-07-28):** the vendored Triton-AOT GDN fast-path cubins — the MEASURED codegen-win decode (Triton REG:205/0-spill vs hand-CUDA REG:255+STACK:48 spills) plus the delta_h/chunk_o/kkt/tril/wu FLA set — existed for **`sm_121a` ONLY**, so GDN-hybrid models (Qwen3.6) ran GDN decode correct but on the SPILLING hand kernel on every other arch (all cross-family arch builds ship `-DVLLM_CPP_TRITON=OFF`). Regenerate the full GDN AOT set per-arch via the sanctioned `-DVLLM_CPP_TRITON_REGEN=ON` pipeline and vendor each `sm_XX/` tree so a `-DVLLM_CPP_TRITON=ON` build selects the non-spilling FLA path additively (build-time cubin selection = `sm_${VLLM_CPP_CUDA_ARCHITECTURES}`, no code change). | [kernel matrix](kernel-matrix.md), [backend matrix](backend-matrix.md) | [per-arch Triton AOT](specs/triton-aot-per-arch.md) | `DONE (DERIVED+BUILD-VERIFIED)` | **LANDED 2026-07-28 (`CLAIM-TRITON-AOT-PER-ARCH`, base `308c312a`):** full GDN AOT set regenerated + vendored for `sm_80/86/89/90a/100a` on dgx GB10 (Triton 3.6.0 / ptxas 12.8, cross-compiled — no target board needed; 57 artifacts + MANIFEST per arch matching the `sm_121a` fileset), `cuobjdump` proves real per-target SASS (`sm=80/86/89/90/100`, decode REG 209–217/0-spill, under the hand-CUDA spill floor); builder-path configure selects + integrity-verifies each tree; `check-triton-aot-drift.sh` rc=0 across all six trees; `sm_121a` byte-untouched (SACRED 27B/35B gate structurally unchanged). **HONEST: DERIVED+BUILD-VERIFIED — no non-`sm_121` board runs a GDN model here, so decode parity on these arches is build-verified, NOT runtime-measured.** Residual: runtime GDN-model verification on Ampere/Ada/Hopper/DC-Blackwell silicon when reachable | | 11 | `ROAD-V1-D2` | Tensor/multi-GPU and MoE sequence parallelism | [engine matrix](engine-matrix.md), [coverage view §3](feature-matrix.md#3-parallelism--scale-out) | TP spec written; v0.25.0 non-DP MoE sequence-parallel path is inventoried as `PAR-SEQUENCE-MOE` but unspiked | `READY` | acquire 2-GPU target and claim Phase 0 mock/ABI; its execution trace determines whether sequence parallel is part of the first performance slice | -| 12 | `ROAD-V1-D3` | Spec-decode breadth: ngram and EAGLE3 | [engine matrix](engine-matrix.md), [coverage view §8](feature-matrix.md#8-speculative-decoding), [spec](specs/spec-decode-breadth-d3.md) | [spec-decode-breadth-d3.md](specs/spec-decode-breadth-d3.md) | `DONE` | **DONE 2026-07-27 (`CLAIM-ROADMAP-D3`).** Both reachable methods closed reusing the LANDED MTP/DFlash verify/reject/`take_draft_token_ids` loop. **ngram (`SPEC-NGRAM` ACTIVE/gated):** draft-FREE suffix-ngram proposer (1:1 port of `ngram_proposer.py`) wired as a third `--speculative-config` method; 27B gate 5/5 STRICT our-ngram-ON == vLLM-ngram-ON + 180/180 drafts accepted; unit 19/19; spec-OFF byte-identical (SACRED 235/235 + MTP 9/9 + DFlash 27/27); host-side, no new kernel, `-Werror` clean. **EAGLE3 (`SPEC-EAGLE3` BLOCKED):** honest reachable-blocked (W0 RUN-verify) — no ungated oracle-runnable EAGLE3 draft arch/checkpoint for a Qwen3.6 gate model at pin `555967922` (registry has no `Eagle3Qwen3_5*`; z-lab published DFlash not EAGLE3); port scoped for when a checkpoint lands (Command-R HF-gate pattern). Named tail (DSpark `SPEC-DSPARK`, TLI `SPEC-TLI`) remains under `ROAD-V1-C3`. | +| 12 | `ROAD-V1-D3` | Spec-decode breadth: ngram and EAGLE3 | [engine matrix](engine-matrix.md), [coverage view §8](feature-matrix.md#8-speculative-decoding), [spec](specs/spec-decode-breadth-d3.md) | [spec-decode-breadth-d3.md](specs/spec-decode-breadth-d3.md) | `DONE` | **DONE 2026-07-27 (`CLAIM-ROADMAP-D3`).** Both reachable methods closed reusing the LANDED MTP/DFlash verify/reject/`take_draft_token_ids` loop. **ngram (`SPEC-NGRAM` ACTIVE/gated):** draft-FREE suffix-ngram proposer (1:1 port of `ngram_proposer.py`) wired as a third `--speculative-config` method; 27B gate 5/5 STRICT our-ngram-ON == vLLM-ngram-ON + 180/180 drafts accepted; unit 19/19; spec-OFF byte-identical (SACRED 235/235 + MTP 9/9 + DFlash 27/27); host-side, no new kernel, `-Werror` clean. **EAGLE3 (`SPEC-EAGLE3` BLOCKED):** honest reachable-blocked (W0 RUN-verify) — no ungated oracle-runnable EAGLE3 draft arch/checkpoint for a Qwen3.6 gate model at pin `555967922` (registry has no `Eagle3Qwen3_5*`; z-lab published DFlash not EAGLE3); port scoped for when a checkpoint lands (Command-R HF-gate pattern). Named tail (DSpark `SPEC-DSPARK`, TLI `SPEC-TLI`) remains under `ROAD-V1-C3`. **FULL SURFACE ENUMERATED 2026-08-09 (`row/SPEC-DECODE-INVENTORY`, [inventory](specs/spec-decode-inventory.md)):** all 13 vLLM `SpeculativeMethod` strings plus the acceptance / dynamic-k config axes cross-referenced from source; nine ABSENT methods newly rowed `INVENTORIED` in the engine matrix — see the D3-breadth gap paragraph below. | | 13 | `ROAD-V1-D4` | **KV persistent state to disk, and external KV-cache provider interoperability with LMCache** (user-directed 2026-07-22: "let's do the KV persistent state to disk support, and LMCache support too", under the standing same-featureset-as-vLLM-and-better bar) | [`KV-OFFLOAD`](engine-matrix.md), [`KV-EXTERNAL-CACHE`](engine-matrix.md), [`KV-CONNECTORS`](engine-matrix.md), [coverage view §2](feature-matrix.md#2-kv-cache--memory), [LMCache quickstart](https://docs.lmcache.ai/getting_started/quickstart.html) | spike ACCEPTED [kv-persistence-lmcache.md](specs/kv-persistence-lmcache.md) — 60 enumerated features across `vllm/v1/kv_offload/`, the `KVConnectorBase_V1` ABI and the LMCache integration, each with a DONE/PARTIAL/MISSING verdict read out of our source. **The two halves of the user's ask are NOT the same kind of work.** Disk persistence is a faithful MIRROR job and is tractable: vLLM's `fs` tier is ~101 lines of `open`/`write`/`readv` with nothing Python-specific in the byte path, one raw file per block, temp-file + atomic rename under `O_DIRECT`. LMCache is NOT: the vLLM-facing glue is vendored in-tree (~2396 lines) but every file of it imports the EXTERNAL PyPI package at module scope, and the storage engine, wire protocol, config schema and CUDA-IPC handoff all live outside the tree with no upstream test that runs without it — so it is scoped as an interop STUDY with a go/no-go, never a from-scratch client. **REOPENED 2026-07-23 ([LMCache client wire analysis](specs/lmcache-cpp-client-connector.md)), and the "no specified wire protocol" half of that verdict is REFUTED by reading the LMCache package: vLLM connects to a RUNNING LMCache over TWO fully-specified portable wires — the `lm://` remote-store (plain TCP + fixed `struct` header + raw KV bytes, no ZMQ/msgpack/pickle/CUDA-IPC) and the MP server (ZMQ + `msgspec.msgpack` + CUDA-IPC, the user's "zmq" recollection). A from-scratch C++ client is FEASIBLE with ZERO `lmcache` in-process; both wires sidestep the hash blocker because LMCache keys on its own blake3 token hash. Recommend the `lm://` mode first. Residual risk is that LMCache is an unpinned moving target — an interop feature with a version-sync cost, not a mechanical core port. LMCACHE-CLIENT W1 LANDED 2026-07-23 (`CLAIM-LMCACHE-CPP-CLIENT`, `KV-EXTERNAL-CACHE` `SPIKE`→`ACTIVE`): the pure-CPU `lm://` wire codec — fixed-`struct` `ClientMetaMessage`/`ServerMetaMessage` framing, the `CacheEngineKey` string, the blake3 rolling token hash (vendored BLAKE3 1.5.5), and the `KV_2LTD` `[2,L,T,D]` repack — is BYTE/BIT-EXACT vs fixtures from the real Python codec (`test_lmcache_codec` 6/6, 2074 assertions), blake3 verified byte-identical on x86-64 + aarch64, and INERT (no call site; the connector is client-W3). LMCACHE-CLIENT W2 LANDED 2026-07-23 — the go/no-go PASSED: a blocking POSIX-socket `LMCacheRemoteClient` (PUT/GET/EXIST/HEALTH/LIST + `KV_2LTD` repack + `VT_LMCACHE_*` config) round-trips a REAL `lmcache.v1.server` (`8570aad`, run headless from source in a throwaway venv — torch imported before lmcache to dodge a torch circular import, the compiled `c_ops` ext stubbed as unused by the lm:// CPU store) byte-identical (`test_lmcache_client` 36/36), with BIDIRECTIONAL interop proven against LMCache's OWN Python protocol codec; the always-on CI gate is a same-binary C++ mock-server round-trip (45/45, no Python). STILL `ACTIVE`, not DONE. Resume at client-W3 (wire as a `KVConnector` over the parent W5 seam, then key-agreement + DGX every-axis gates).** **Blocking correction found in OUR source:** `NONE_HASH` is seeded from `std::random_device` with no escape hatch, so every block hash differs across processes and a content-addressed disk tier would score 0% hits on restart — we are WORSE than vLLM here, which at least exposes `PYTHONHASHSEED`. This also FALSIFIES the caching spike's §B2 claim that we are deterministic by construction. **Two upstream weaknesses recorded as beyond-parity targets:** the `fs` tier's `config.json` is written and never read (its only identity check is a path digest that omits checkpoint content, weight quantization, rope config and `sliding_window` — a silent-wrong-output hazard we will not copy), and the disk tier has no capacity accounting and no eviction. Three matrix rows `INVENTORIED` -> `SPIKE`; `SharedStorageConnector` found RENAMED to `ExampleConnector` and `P2pNcclConnector` found DELETED at the pin, both stale in the prior record **W1-W3 IMPLEMENTED 2026-07-22, CPU-only.** W1 deterministic block hashes: `init_none_hash` now resolves explicit arg > `$VLLM_PREFIX_CACHING_HASH_SEED` > `$PYTHONHASHSEED` > a fixed built-in default, so hashes are identical across processes with ZERO configuration — the blocking correction is CLOSED, and we now BEAT upstream on this axis rather than trailing it (upstream is random-by-default and documents `PYTHONHASHSEED` as the operator's problem). Proven by comparing hash chains emitted by SEPARATELY LAUNCHED processes, with a negative control confirming the opt-in `=random` mode genuinely disagrees. W2 CPU primary tier: `CachePolicy` (LRU + ARC) with the `ref_cnt == -1` tri-state and the ATOMIC evict, `CPUOffloadingManager` incl. the `prepare_store -> nullopt` skip control path, pinned backing store, and a side-queue event-polled device/host transfer worker. W3 disk `fs` tier: one raw file per block, temp-file + atomic rename publish, self-healing unlink, dual-queue read/write pool. **BOTH recorded upstream weaknesses are now EXCEEDED rather than merely noted** — the identity block is a VERIFIED header read on EVERY open that REFUSES on mismatch across 27 fields (tested per field, with a positive control), and the tier carries a byte budget with policy eviction honoured across restarts. `O_DIRECT` deliberately NOT ported (a header+payload file breaks its alignment requirement); the GIL-releasing batch-lookup C extension is unconditionally unnecessary without a GIL. `KV-OFFLOAD` `SPIKE` -> `PARTIAL`. **W4 IMPLEMENTED 2026-07-23:** the TIERING MANAGER (ONE manager over CPU primary + disk secondary — disk→CPU promotion RETRY→flush→HIT, cascade demotion, reset drains the secondary first and never resets it so a persisted cache survives) and the CONNECTOR/SCHEDULER HALF (`OffloadingConnector`, the semantics of `KVConnectorBase_V1`'s scheduler hooks — nullopt third state, `block_hashes` striding, load-before-compute — wired OPT-IN + DEFAULT-OFF into the scheduler). First measured offload speedup: a restarted-prefix workload through the REAL scheduler saved 32/48 prefill tokens (2/3 blocks HIT from disk), promoted bytes byte-identical to the cold store; identity refusal holds through a promotion. Ported the SEMANTICS not the Python plugin ABI (compile-time wiring); the full abstract ABI is W5. **W5 IMPLEMENTED 2026-07-23, CPU-only:** the connector seam is now a first-class C++ ABI — the abstract `KVConnector` base carrying the full scheduler + worker method set of `KVConnectorBase_V1` (the scheduler methods load-bearing, the worker hooks defaulted no-ops for our synchronous runner, documented), a compile-time `KVConnectorFactory` + `REGISTER_KV_CONNECTOR` (the C++ analogue of vLLM's `importlib` module path), and a `KVTransferConfig` selection surface (default `kv_connector` empty == no connector == zero behaviour change, `kv_role` validation, `fail`-default load policy). The W4 disk connector was refactored ONTO this base behaviour-identically — the restart-hit e2e reproduces byte-for-byte and a config-selected owning connector shortcuts prefill by the identical 32/48. `KV-CONNECTORS` `SPIKE`→`ACTIVE`. This closes the seam so LMCache client-W3 is 'implement the abstract `KVConnector` with the landed W2 `lm://` client'. **LMCACHE-CLIENT W3 LANDED 2026-07-23 — the `lm://` client wired as a `KVConnector` over the W5 seam (`LMCacheConnector`, `REGISTER_KV_CONNECTOR("LMCacheConnector", …)`, default OFF), the FIRST time the whole chain engine -> connector -> W2 client -> a running lm:// server -> back runs.** Scheduler side computes rolling-blake3 chunk hashes and `Exist`-probes the remote store for the longest cached prefix (synchronous `(n, false)`, mirroring `lmcache_connector.py:230-259`); worker `StoreChunk`/`LoadChunk` drive the W2 client with foreign-block REFUSAL. **Gate ACHIEVED = the connector-level round-trip: STORE a prefix -> a fresh "restarted" connector LOOKS UP + shortcuts prefill through the REAL scheduler (32/48 tokens saved) -> LOAD byte-identical; foreign-key REFUSAL; default-off INERT** (`test_lmcache_connector` 5 cases / 50 assertions vs an in-process mock; store->load ALSO GREEN vs a REAL `lmcache.v1.server` 8570aad, 16 assertions, `VT_LMCACHE_LIVE_*`). **LMCACHE-CLIENT W4 LANDED 2026-07-23 — REAL peer KEY-AGREEMENT + a peer->us interop LOAD, both PROVEN — the interop-correctness milestone is COMPLETE; `KV-EXTERNAL-CACHE` stays `ACTIVE` for the DGX full-model output-invariance + throughput arm.** The actual `lm://` key derivation is `ChunkedTokenDatabase` (NOT the blake3 MP hasher): chunk_size 256, a rolling prefix-hash over `(prefix_int, tuple(tokens), extra=())` keyed by vLLM's OWN hash (portable `sha256_cbor`), folded to uint64 each step, `NONE_HASH=fold8(sha256_cbor(str(PYTHONHASHSEED)))`. Mirrored byte-exact (`chunked_token_database.{h,cpp}`, reusing `CborValue`+`sha256_cbor`) and wired as connector `key_mode=kVllmSha256Cbor` (chunk 256) alongside W3's kept-green blake3 path. Key-agreement GREEN: `test_lmcache_key_agreement` 4/85 == the REAL lmcache `ChunkedTokenDatabase.process_tokens()` BYTE-FOR-BYTE (fixtures dumped from the unmodified real driver + vLLM's pinned `sha256_cbor`/`init_none_hash`), sample `meta-llama/Llama-3.1-8B@1@0@33d6862800fff40c@bfloat16`. Peer->us LOAD GREEN over the wire: a REAL lmcache `ChunkedTokenDatabase` derives a key + PUTs KV to a REAL `lmcache.v1.server`, our C++ re-derives the SAME key and GETs the 512 B byte-identical (`run_key_interop.sh`). ASan+UBSan clean. Text-only (mm-hash extra_keys deferred). **LMCACHE-CLIENT W5 LANDED 2026-07-24 — the LAST open arm, connector-ON full-model OUTPUT-INVARIANCE + throughput in a REAL generation loop, is CLOSED (spec gates 4/6 met).** The worker side is now wired into the engine: `GPUModelRunner::execute_model` calls `ConnectorLoadExternalKv` before the forward (writes the external-prefix KV into the allocated GPU blocks, load-before-compute) and `ConnectorStorePromptKv` after (stores each newly-complete prompt block), and `LoadedEngine` builds the connector from an `EngineParams` `KVTransferConfig` and wires it to BOTH the scheduler and the runner. **OUTPUT-INVARIANCE PROVEN on a real OPT-125m loop vs a live `lmcache.v1.server`: connector-ON generated tokens are BIT-IDENTICAL to connector-OFF (cold full prefill) — first-divergence index -1 — on BOTH (a) a store->restart->load cycle in one process AND (b) a genuinely cold second process that only hits the server; prefill saved on the hit = 48 tokens (3×16-token blocks).** `tests/vllm/models/test_lmcache_output_invariance.cpp` PASSES both modes via `scripts/lmcache/run_output_invariance.sh`. THROUGHPUT reported HONESTLY: on a 125M model the wall-clock delta is noise-dominated (fixed TCP/copy overhead ~ tiny compute saved), so NO binding speedup is claimed — a real speed number is owed by an every-axis grid on a larger model + long shared-prefix corpus (docs/BENCHMARKS.md). No-regression: OPT SACRED UNCHANGED default-off (6/6, 96/96, 63/63); connector unit tests green (codec 6/6, client 3/3, connector 5/5, key-agreement 4/4, kv_offload_connector 11/11); ASan+UBSan clean on the connector path; CUDA `-Werror` 0 warnings. Additive + default-off inert (all worker/loader changes are behind a null-connector guard)| `PARTIAL` | **W7 the one genuine beyond-parity item (imperative named per-sequence save/restore), which now has the verified header it depends on; and a binding every-axis LMCache throughput grid on a larger model vs vLLM's `--kv-transfer-config`.** W5 (the abstract ABI) is DONE; LMCache client W1 (codec) + W2 (client) + W3 (connector round-trip) + W4 (peer key-agreement + interop load) + W5 (full-model output-invariance) are DONE. **The benchmark blocker is CLEARED:** the caching spike's W1 prefix-cache counters landed earlier, so the W4 offload arm proved its hits | | 13a | `ROAD-V1-D4-APC` | **Prompt / prefix caching to full vLLM parity, then beyond (user-directed 2026-07-22: "same featureset of vLLM and better")** — the headline user-facing caching feature, previously mentioned only once in this roadmap despite being a shipped, default-ON behaviour for dense models | [`KV-PREFIX-CACHE`](engine-matrix.md), [`KV-BLOCK-POOL`](engine-matrix.md), [`KV-HYBRID-COORD`](engine-matrix.md), [`KV-MAMBA-ALIGN`](engine-matrix.md), [`KV-EVENTS`](engine-matrix.md), [`KV-PREFIX-MATCH-UNIT`](engine-matrix.md), [`ENG-CASCADE-ATTN`](engine-matrix.md), [coverage view §2](feature-matrix.md#2-kv-cache--memory) | umbrella spike ACCEPTED [prefix-prompt-caching-parity.md](specs/prefix-prompt-caching-parity.md) — enumerates the complete pinned-vLLM caching surface (38 features) with a per-feature DONE/PARTIAL/MISSING verdict grounded in our source. **The ported core is deeper than the record claimed** (chain hashing, block pool, all three coordinators, the full hybrid fixed-point intersection, four single-type managers); the real gaps are narrower and different: block-hash extra keys are a no-op stub, there are NO prefix-cache statistics at any level, KV events are inert, `cache_salt` and 3 of 4 hash algos are absent, and `reset_prefix_cache` is implemented but unreachable. Three matrix rows corrected, two of them in our favour. `ENG-CASCADE-ATTN` DISPOSITIONED as not owed (default-off, absent from the MRV2 runner we port, unreachable on Blackwell). llama.cpp comparison completed: its "prompt cache" is session/slot state serialization, strictly weaker than APC on every reuse axis, and vLLM already covers disk persistence via the `kv_offload` fs tier — the ONE genuine capability neither vLLM nor we have is an imperative named per-sequence save/restore **W1 IMPLEMENTED 2026-07-22: prefix-cache statistics exist for the first time.** `BaseCacheStats`/`PrefixCacheStats`/`CachingMetrics` ported 1:1 from `vllm/v1/metrics/stats.py:35-142`, recorded in `get_computed_blocks`, flagged by `reset_prefix_cache`, taken-and-swapped per step and folded into a 1000-request sliding window exposed on `Scheduler`/`EngineCore`/`LLMEngine`. Per the standing parity-enabler rule `log_stats` is DEFAULTED ON (mirroring upstream's `disable_log_stats=False`), so no benchmark arm is void for want of a counter. `Request::num_preemptions` un-deferred to feed the mutually-exclusive `preempted_*` triple. **FIRST MEASURED HIT RATE: 0.75** (1920 of 2560 queried tokens over 16 requests sharing a 128-token prefix), with a caching-OFF 0.0 negative control — the first demonstration in this project that APC actually serves cached tokens. The hard blocker on [`BACKEND-GATE-CUDA-SGLANG-PREFIX`](backend-matrix.md) is CLOSED **W2 DONE 2026-07-27 (`CLAIM-ROADMAP-D4APC`, CPU-gated on dgx GB10, NOT pushed):** `generate_block_hash_extra_keys` ported 1:1 (`kv_cache_utils.py:451-591`) — mm hash + LoRA name + `cache_salt`, fixed order lora->mm->salt (prompt_embeds deferred: no prompt-embeds path); `cache_salt`/`lora_name` carried on `Request`/`EngineCoreRequest`, set before the first hash in `FromEngineCoreRequest` (fixes a latent ordering bug). RED-first no-false-share PROVEN: with the stub a differently-salted request false-hits the prior tenant's 48 cached tokens (`n1==48`), with extra keys `n1==0`. Ported extra-key/ordering cases + hash- and manager-level no-false-share (`test_kv_cache_utils.cpp` 29/29, `test_kv_cache_manager.cpp` 10/10). **This unblocks the MM + LoRA cache consumers.** **W3 DONE 2026-07-27 (`CLAIM-ROADMAP-D4APC-W3`, dgx GB10, NOT pushed) — the FIRST-EVER cache-ON model gate:** on `Qwen/Qwen3-4B` (dense, full-attention, APC-default-ON — the vehicle the prior "vehicle-blocked" note missed) a shared-prefix workload runs APC-ON and APC-OFF through the full paged engine, gating token-identity + hits + prefill drop. **NO engine code changed** (`git diff --stat` = tests+scripts+goldens) ⇒ pure GATE over the already-shipped default-ON path; binary byte-identical ⇒ SACRED unaffected. RESULT (`test_qwen3_apc_e2e` 2/2, 84/84 asserts): APC-ON hits **2240/2777 (rate 0.807)**, APC-OFF 0; APC-ON == APC-OFF EXACT on 5/6 (the 1 diff a vLLM-confirmed 0.125-nat near-tie, RCA'd = attention-kernel-path near-tie flip, not a cache bug); **== vLLM-APC-ON** teacher-forced (APC-OFF 6/6 max gap 0.0 nats = exact argmax, APC-ON 6/6 max gap 0.125 nats, 0 outside top-20); **TTFT drop 70.1→39.9 ms = 1.76×** on a cache hit. Existing 4B SACRED gate 16/16 GREEN (no regression). Oracle vLLM 0.25.0 (0.26 venv broken — editable source disk-reclaimed; 4B byte-stable across the pin). | `DONE` (headline) | **Row DONE for the default dense APC path (block hashing incl. extra_keys, pool, coordinators, stats, scheduling, cache-ON e2e all gated).** Named NON-BLOCKING tails tracked in their own rows / future items: W4 KV events (`KV-EVENTS` — event GENERATION + `msgpack` PAYLOAD DONE 2026-07-27 `CLAIM-ROADMAP-D4-KV-EVENTS`, `SPIKE`→`ACTIVE`, byte-exact vs `msgspec`; live ZMQ transport + engine batch wiring DEFERRED), W5 partial-block primitive (upstream dead-code), W6 Mamba-`align` hybrid cache-on (`KV-MAMBA-ALIGN`, SPIKE — feeds `BACKEND-GATE-CUDA-SGLANG-PREFIX`), W7 `reset_prefix_cache` dev-endpoint + `--prefix-caching-hash-algo` + `skip_reading_prefix_cache`, W8/W9 the beyond-vLLM named session save/restore. The every-axis cache-on grid vs vLLM/SGLang is a separate perf follow-on under `ROAD-V1-A`. No `/metrics` route yet (`SERVE-METRICS`), so the hit rate is read from the engine API. **`--prefix-match-unit` (0.26-new fine-grained matching unit) W0 spike + W1 resolver LANDED 2026-07-28 (`CLAIM-PREFIX-MATCH-UNIT`, `KV-PREFIX-MATCH-UNIT` PARTIAL): `resolve_kv_cache_block_sizes` computes `hash_block_size = prefix_match_unit if set else gcd(group_block_sizes)`, RED-first unit-gated; config/CLI/ABI field (W2) + scheduler threading of `hash_block_size != block_size` (W3, needs the `KV-BLOCK-POOL` align path) + benchmark (W4) deferred.** | | 14 | `ROAD-V1-D5` | LoRA, local KV/weight offload, expert streaming, wider model zoo | [engine matrix](engine-matrix.md), [model matrix](model-matrix.md) | corrected expert-streaming spike accepted (`ENG-EXPERT-STREAM` READY): bank-only safetensors→Marlin bank, fixed contiguous cache slots matching Marlin dense strides, logical→slot remap after explicit router D2H, chunked C [`SPEC-SUFFIX`](engine-matrix.md) (Arctic +suffix decoding, draft-free but external `arctic_inference`, S) -> +[`SPEC-EAGLE`](engine-matrix.md) (EAGLE1, distinct from the reachable-blocked +`SPEC-EAGLE3`; separate draft + one aux tap, M, checkpoint-gated like eagle3) -> +[`SPEC-MTP-FAMILY`](engine-matrix.md) (MTP breadth beyond Qwen3.5; DeepSeek-V4 +nextn already ACTIVE-W1 weight-blocked, ~18 other family heads inventoried, M) +-> [`SPEC-ACCEPT-VARIANTS`](engine-matrix.md) (synthetic/block rejection + +probabilistic draft sampling, cross-cuts every proposer, M) -> +[`SPEC-DYNAMIC`](engine-matrix.md) (per-batch-size k schedule, S) -> +[`SPEC-CUSTOM-CLASS`](engine-matrix.md) (pluggable proposer, extensibility, S) -> +[`SPEC-EXTRACT-HIDDEN`](engine-matrix.md) (offline hidden-state capture, not a +decode accelerator, S) -> [`SPEC-MLP-SPECULATOR`](engine-matrix.md) +(UPSTREAM-DEPRECATED — no V1 proposer, dispatch `ValueError`; enumeration +completeness only). `SPEC-DSPARK`/`SPEC-TLI` were already inventoried. Confirmed NON-gap: vLLM has REMOVED prompt adapters. Other MED/LOW gaps (DP/EP+EPLB [`PAR-DP`](engine-matrix.md)/[`PAR-EP-EPLB`](engine-matrix.md), KV offload [`KV-OFFLOAD`](engine-matrix.md), external KV connectors / PD diff --git a/.agents/specs/spec-decode-inventory.md b/.agents/specs/spec-decode-inventory.md new file mode 100644 index 00000000..49d3ff85 --- /dev/null +++ b/.agents/specs/spec-decode-inventory.md @@ -0,0 +1,82 @@ +# Speculative-decoding method inventory (vLLM `555967922` / 0.26.0.dev0) + +Enumerated from SOURCE, not memory, on 2026-08-09 (`row/SPEC-DECODE-INVENTORY`). +Authoritative surfaces at the parity pin: + +- `vllm/config/speculative.py:37-78` — the `SpeculativeMethod` Literal (every + accepted `method` string) plus `RejectionSampleMethod` / `DraftSampleMethod`. +- `vllm/v1/worker/gpu_model_runner.py:596-655` — the DEFAULT V1 runner's + method -> proposer dispatch (raises `ValueError` on an unhandled method). +- `vllm/v1/worker/gpu/spec_decode/__init__.py:8-40` — the OPT-IN "Model Runner + V2" (`VLLM_USE_V2_MODEL_RUNNER=1`) factory; a NARROWER subset (dflash, dspark, + gemma4_mtp, mtp, eagle/eagle3 only). +- `vllm/v1/spec_decode/` — the V1 proposer classes. +- `vllm/model_executor/models/` — the per-family speculator model classes. +- `vllm/transformers_utils/configs/speculators/algos.py:4-15` — the HF + `speculators`-library checkpoint-format integration. + +**V1/V0 note.** At the pin the engine is V1-only (no V0 model runner). "V1 +status" below therefore means: does the DEFAULT V1 runner instantiate a proposer +for the method. `mlp_speculator` is the one method the enum still accepts that +has NO V1 proposer -> it is effectively deprecated / unreachable at the pin. + +**Not "ours-beyond-vLLM".** `dflash` (block-diffusion) and `dspark` are REAL +upstream vLLM methods (`vllm/v1/spec_decode/dflash.py`, +`vllm/v1/worker/gpu/spec_decode/dspark/speculator.py`); our DFlash port mirrors +vLLM, it is not an addition vLLM lacks. + +## Method x status + +| Method (`method=`) | Upstream anchor | Draft-weight requirement | Upstream V1 status | Our status (row) | +|---|---|---|---|---| +| `ngram` | `vllm/v1/spec_decode/ngram_proposer.py:12` | none (suffix match over prompt+output) | V1 default runner (not MRV2) | **DONE** `SPEC-NGRAM` (27B 5/5 strict) | +| `ngram_gpu` | `vllm/v1/spec_decode/ngram_proposer_gpu.py:217` | none (on-device n-gram) | V1 default runner | **ABSENT** `SPEC-NGRAM-GPU` | +| `suffix` | `vllm/v1/spec_decode/suffix_decoding.py:9` (Arctic Inference, ext. `arctic_inference`, lazy import :26) | none (suffix tree over prompt+output) | V1 default runner | **ABSENT** `SPEC-SUFFIX` | +| `medusa` | `vllm/v1/spec_decode/medusa.py:18`; model `vllm/model_executor/models/medusa.py` | head-on-base (extra LM heads on target, single non-AR pass) | V1 default runner | **SPIKE** `SPEC-MEDUSA` (W0 only; model row `MODEL-SPEC-medusa-medusa`) | +| `mlp_speculator` | model `vllm/model_executor/models/mlp_speculator.py`; enum `speculative.py:69,890` | separate MLP head | **NONE** — no V1 proposer; dispatch `ValueError` (`gpu_model_runner.py:651-654`); V0-only/deprecated | **ABSENT** `SPEC-MLP-SPECULATOR` (upstream-deprecated) | +| `draft_model` | `vllm/v1/spec_decode/draft_model.py:19` | separate standalone LM (no shared emb/head) | V1 default runner | **ACTIVE** `SPEC-DRAFT-MODEL` (W0+W1 CPU brick; W3 GPU DGX residual) | +| `custom_class` | `vllm/v1/spec_decode/custom_class_proposer.py:12` (class path from `speculative_config.model`) | proposer-defined (pluggable) | V1 default runner | **ABSENT** `SPEC-CUSTOM-CLASS` | +| `eagle` | `vllm/v1/spec_decode/eagle.py:10` (`use_eagle()` `speculative.py:1324-1328`); models `llama_eagle.py`, `cohere_eagle.py`, `mistral_eagle.py`, ... | separate draft + target aux hidden-state tap | V1 default + MRV2 | **ABSENT** `SPEC-EAGLE` (EAGLE1 distinct from eagle3) | +| `eagle3` | same `EagleProposer`; models `llama_eagle3.py`, `qwen3_eagle3.py`, `deepseek_eagle3.py` | separate draft + multi-layer aux taps | V1 default + MRV2 | **BLOCKED** `SPEC-EAGLE3` (no ungated gate-model checkpoint at pin) | +| `extract_hidden_states` | `vllm/v1/spec_decode/extract_hidden_states.py:29` | none (offline hidden-state capture utility) | V1 default runner | **ABSENT** `SPEC-EXTRACT-HIDDEN` | +| `mtp` (canonical) | `EagleProposer` via `use_eagle()`; special `Gemma4Proposer`/`Step3p5MTPProposer` (`gpu_model_runner.py:627-630`); MRV2 `MTPSpeculator` | head-on-base per family (nextn layer on target) | V1 default + MRV2 | **DONE** `SPEC-MTP` (Qwen3.5/3.6 k=1); family breadth `SPEC-MTP-FAMILY` | +| `deepseek_mtp` ... `inkling_mtp` (20 family strings) | `MTPModelTypes` `speculative.py:37-59` -> deprecate-remap to `mtp` (`:686-690`); models `deepseek_mtp.py`, `glm4_moe_mtp.py`, `ernie_mtp.py`, `nemotron_h_mtp.py`, ... | head-on-base per family | V1 default (dispatched by `draft_model_config.hf_config.model_type`) | **ACTIVE** DeepSeek-V4 (`MODEL-SPEC-deepseek-v4-deep-seek-v4-mtp`, W1, weight-blocked); rest **INVENTORIED** (`SPEC-MTP-FAMILY`, model-matrix `MODEL-SPEC-*-mtp`) | +| `dflash` | `vllm/v1/spec_decode/dflash.py` (DFlashProposer); models `laguna_dflash.py`, `qwen3_dflash.py`; speculators algo `algos.py:93` | separate block-diffusion draft + aux | V1 default + MRV2 | **DONE** `SPEC-DFLASH` (ported from vLLM; +GGUF `SPEC-DFLASH-GGUF`) | +| `dspark` | MRV2 `vllm/v1/worker/gpu/spec_decode/dspark/speculator.py:37`; models `gemma4_dspark.py`, `qwen3_dspark.py`; speculators algo `algos.py:133` | separate SAR block draft + aux | V1 default (via `use_eagle()`) + MRV2 | **INVENTORIED** `SPEC-DSPARK` | + +## User-visible config axes (apply on top of a method) + +| Axis | Upstream anchor | Values | Our status (row) | +|---|---|---|---| +| Acceptance / rejection sampler | `speculative.py:77,216` (`rejection_sample_method`) | `standard` \| `synthetic` \| `block` | `standard` **ACTIVE** `SPEC-REJECTION`; `synthetic`/`block` **ABSENT** `SPEC-ACCEPT-VARIANTS` | +| Draft sampling | `speculative.py:78,283` (`draft_sample_method`) | `greedy` \| `probabilistic` | `greedy` only; `probabilistic` **ABSENT** `SPEC-ACCEPT-VARIANTS` | +| Dynamic k (per batch size) | `speculative.py:1336-1337`; `vllm/v1/spec_decode/dynamic/utils.py` (`num_speculative_tokens_per_batch_size`) | schedule list | **ABSENT** `SPEC-DYNAMIC` | +| Heterogeneous draft/target vocab | `vllm/v1/spec_decode/vocab_mapping.py` | shared-token / ID translation | **INVENTORIED** `SPEC-TLI` | +| GDN spec metadata + slot-snapshot rollback | `vllm/v1/attention/backends/gdn_attn.py` | — | **ACTIVE** `SPEC-GDN-SEGMENTS` | + +## HF `speculators` checkpoint-format integration + +`vllm/transformers_utils/configs/speculators/algos.py` registers +`SUPPORTED_SPECULATORS_TYPES = {eagle3, peagle, dflash, dspark}` (`:15,55,93,133`) +and rewrites a HF `speculators`-format config into a vLLM draft arch. Our side +loads native draft configs directly (no `speculators`-format adapter); tracked +under the EAGLE/DFlash/DSpark rows as a loader-format residual, not a separate +method. + +## Our lifecycle summary (engine matrix §Speculative decoding) + +- **DONE**: `SPEC-MTP`, `SPEC-MTP-GGUF`, `SPEC-DFLASH`, `SPEC-DFLASH-GGUF`. +- **ACTIVE**: `SPEC-REJECTION`, `SPEC-GDN-SEGMENTS`, `SPEC-NGRAM`, + `SPEC-DRAFT-MODEL`. +- **SPIKE**: `SPEC-MEDUSA`. **BLOCKED**: `SPEC-EAGLE3`. +- **INVENTORIED**: `SPEC-DSPARK`, `SPEC-TLI`, and the nine enumerated here: + `SPEC-NGRAM-GPU`, `SPEC-SUFFIX`, `SPEC-EAGLE`, `SPEC-MTP-FAMILY`, + `SPEC-ACCEPT-VARIANTS`, `SPEC-DYNAMIC`, `SPEC-CUSTOM-CLASS`, + `SPEC-EXTRACT-HIDDEN`, `SPEC-MLP-SPECULATOR`. + +Each INVENTORIED row is spike-first: a `.agents/specs/` spike gate precedes any +implementation claim, per the tabular-inventory directive. Value order for the +nine gaps (gate-model-usable and draft-free first): `SPEC-NGRAM-GPU` -> +`SPEC-SUFFIX` -> `SPEC-EAGLE` -> `SPEC-MTP-FAMILY` -> `SPEC-ACCEPT-VARIANTS` -> +`SPEC-DYNAMIC` -> `SPEC-CUSTOM-CLASS` -> `SPEC-EXTRACT-HIDDEN` -> +`SPEC-MLP-SPECULATOR` (upstream-deprecated, completeness only). diff --git a/.agents/state.md b/.agents/state.md index c4edbfc5..fa9639c0 100644 --- a/.agents/state.md +++ b/.agents/state.md @@ -37180,3 +37180,50 @@ also runs) or within the cross-tool attribution boundary. NEXT lever (only real grouped→dense-direct-A marlin decode (#50 NO-GO'd the full port; a targeted E1 direct-A path is the open candidate). Box left clean (both locks free, GPU idle, worker down, disk 21G, tmux gone). +## SPEC-DECODE-INVENTORY: full vLLM speculative-method surface enumerated from source; 9 ABSENT methods rowed INVENTORIED + + +Records-only branch `row/SPEC-DECODE-INVENTORY` off origin/main `4dd4e206`; pinned +vLLM verified `555967922` / 0.26.0.dev0. No C++ — inventory/records only. + +Enumerated ALL vLLM speculative methods from SOURCE (not memory): +`vllm/config/speculative.py:66-76` (`SpeculativeMethod` Literal), +`vllm/v1/worker/gpu_model_runner.py:596-655` (default V1 method->proposer +dispatch), the opt-in MRV2 factory `vllm/v1/worker/gpu/spec_decode/__init__.py:8-40` +(narrower subset), `vllm/v1/spec_decode/*` proposers, `model_executor/models/*` +speculator classes, and the HF `speculators` lib `algos.py:15,55,93,133`. Full +method x status table: `.agents/specs/spec-decode-inventory.md`. + +Findings: 13 method strings + config axes (`rejection_sample_method` +standard/synthetic/block, `draft_sample_method` greedy/probabilistic, dynamic +`num_speculative_tokens_per_batch_size`). `mlp_speculator` has NO V1 proposer +(dispatch raises `ValueError`) = upstream-deprecated. `dflash`/`dspark` are REAL +vLLM methods — our DFlash MIRRORS vLLM, it is NOT ours-beyond-vLLM. + +Our side (verified vs code): DONE `SPEC-MTP` (Qwen3.5/3.6 k=1) + `SPEC-DFLASH` ++ GGUF variants; ACTIVE `SPEC-REJECTION`/`SPEC-GDN-SEGMENTS`/`SPEC-NGRAM` +(27B 5/5 strict)/`SPEC-DRAFT-MODEL` (W0+W1 CPU brick, W3 GPU DGX residual); +SPIKE `SPEC-MEDUSA` (W0 only, NO code — the FEATURES "CPU brick" mark was drift, +corrected); BLOCKED `SPEC-EAGLE3` (no ungated gate-model checkpoint at pin); +DeepSeek-V4 MTP ACTIVE-W1 weight-blocked (`MODEL-SPEC-deepseek-v4-deep-seek-v4-mtp`). +Model-matrix already inventories 49 `MODEL-SPEC-*` speculator model classes. + +Added 9 engine-matrix rows (all `INVENTORIED`, value-ordered): `SPEC-NGRAM-GPU`, +`SPEC-SUFFIX`, `SPEC-EAGLE`, `SPEC-MTP-FAMILY`, `SPEC-ACCEPT-VARIANTS`, +`SPEC-DYNAMIC`, `SPEC-CUSTOM-CLASS`, `SPEC-EXTRACT-HIDDEN`, `SPEC-MLP-SPECULATOR`. +`ENGINE_ROWS` 131->140 (check-agent-record.py); engine summary Total + area rows +updated (INVENTORIED 32->41, Speculative area 12->21). INVENTORIED (not SPIKE) +chosen to match the `SPEC-DSPARK`/`SPEC-TLI` precedent and because SPIKE demands +an active coordination claim; each row is spike-first per the inventory directive. + +Roadmap `ROAD-V1-D3` portfolio + D3 track updated + gap paragraph enumerating the +nine. FEATURES.md spec rows corrected (Medusa->spike, n-gram->DONE, EAGLE->EAGLE1/3, ++breadth row). BENCHMARKS spec table +ngram +breadth rows. STATUS: method-surface +paragraph added, PAID by collapsing the 4 superseded DeepSeek-V4 decode-lever +paragraphs (duplicated in BENCHMARKS, superseded by the 1.14x BEATS-ds4 binding) += -5772 chars; `STATUS_RATCHET` chars 287838->282066. + +NEXT (value order): implementable draft-free gaps first — `SPEC-NGRAM-GPU`, +`SPEC-SUFFIX` (needs external `arctic_inference`) — then `SPEC-EAGLE` +(checkpoint-gated like eagle3). Gates green: check-agent-record / doc-checkpoint +/ public-doc-tables / now-current. diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index a2117bdb..ef0d4338 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -254,6 +254,8 @@ in the tree, default-OFF, for reproducibility; detail in the benchmark record. |---|---|---|---| | MTP | Qwen3.6-27B NVFP4 | token-identical to vLLM MTP, **~4% faster at c1**; on-par at c2-c8 | `DONE` | | DFlash | Qwen3.6-27B NVFP4 | **2.9x over spec-off** (10.16 → 29.32 tok/s), at/above vLLM DFlash-on (**1.003x**, non-overlapping bands) | `DONE` | +| n-gram | Qwen3.6-27B NVFP4 | draft-free (`SPEC-NGRAM`); 27B 5/5 STRICT our-ngram-ON == vLLM-ngram-ON, 180/180 drafts accepted (correctness only, no speed row yet) | `DONE` | +| Breadth (EAGLE1/3, suffix, ngram-gpu, dspark, dynamic-k, ...) | n/a | enumerated from vLLM source + `INVENTORIED` 2026-08-09 (`.agents/specs/spec-decode-inventory.md`), unmeasured | `INVENTORIED` | ## How we measure diff --git a/docs/FEATURES.md b/docs/FEATURES.md index 2f101524..23ec7446 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -109,10 +109,11 @@ the registered engine forward. |---|---|---|---| | MTP (multi-token prediction) | ✅ token-identical, ~4% faster at c1 | ✅ | ✅ | | Draft model | ◐ CPU brick | ✅ | ✅ | -| Medusa | ◐ CPU brick | ✅ | ✅ | -| EAGLE | ☐ | ✅ | ✅ | +| Medusa | ☐ spike only | ✅ | ✅ | +| EAGLE / EAGLE3 | ☐ | ✅ | ✅ | | DFlash block diffusion | ✅ 2.9x over spec-off, at/above vLLM DFlash-on | ✅ | ☐ | -| n-gram / prompt lookup | ☐ | ✅ | ✅ | +| n-gram / prompt lookup | ✅ 27B 5/5 strict vs vLLM | ✅ | ✅ | +| Other methods (ngram-gpu, suffix, dspark, custom-class, dynamic-k, mlp-speculator) | ☐ inventoried | ✅ | ◐ | ## Structured output and tool calling diff --git a/docs/STATUS.md b/docs/STATUS.md index b8b23ddd..4e97eee4 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -126,6 +126,8 @@ concurrency-1 A/B our-on 29.32 tok/s vs vLLM-on 29.24, non-overlapping bands, D0 through D14 on the vLLM 0.26.0.dev0 stack (which resolves vllm#40898), and it remains gated behind a spike while its user-facing serving surface is finalized. +**Method surface (enumerated from vLLM source 2026-08-09, `.agents/specs/spec-decode-inventory.md`).** Of the 13 vLLM `SpeculativeMethod` strings we ship MTP (k=1), DFlash and n-gram; draft_model is a CPU brick and Medusa a spike; EAGLE1/EAGLE3, ngram-gpu, suffix, dspark, custom_class, extract_hidden_states, dynamic-k and the synthetic/block acceptance variants are INVENTORIED; mlp_speculator is upstream-deprecated (no V1 proposer). + **DeepSeek-V4 native MTP** (`DeepSeekV4MTPModel`, ACTIVE — W1 self-spec wiring, 2026-07-30) has its nextn draft head wired to the same lossless spec-decode path. Unlike V3's fused `eh_proj`, the V4 nextn layer keeps separate `e_proj`/`h_proj`, @@ -149,13 +151,7 @@ loop (R2) + device draft (R4). fp16 dequant cache refuted net-slower on GB10; MH a measured tie; routed-MoE we already win. See `.agents/specs/deepseek-beat-ds4-sweep-2026-07-30.md`. -**DeepSeek-V4-Flash decode — Q8_K activation-quant ds4-geometry port LANDED: +5.4%, byte-exact (2026-08-03, `CLAIM-DSV4-Q8K-PREQ`, GB10 sm_121a, base `9154fbdb`).** Updates the "~13.1-13.2 ceiling" above: the Q8_0 GEMV front is closed, but the routed-expert **Q8_K activation quant was NOT**. `QuantizeQ8KKernel` mapped one thread to a whole 256-element super-block, so a decode grouped-MoE quant (`nsb`≈8-28) ran a SINGLE 128-thread block with ≤28 active threads on the >100-SM GB10. New `QuantizeQ8KPreqKernel` ports ds4's `q8_K_quantize_kernel` grid geometry (`grid=(nsb,m)`, one 256-thread block/super-block, shared-mem reduction), gated `VT_V4_PREQ_FUSED` (default ON), **byte-exact** (the amax reduction breaks ties by LOWEST index → reproduces the legacy first-occurrence scan; decode ids IDENTICAL `=1` vs `=0`, 40 tok). **`QuantizeQ8K` 4.04→0.20 ms/step (~20×); total GPU-busy 71.85→67.95 ms/step; decode 13.54→14.27 tok/s (+5.4%, 3 paired drop_caches reps) = 82.1%→86.5% of ds4 ~16.5** (bar corrected from the unreproduced 17.13 anchor — see the MHC-SINK4 BAR CORRECTION note below). This is NOT the `QuantizeQ8K`→GEMV inline fusion of `CLAIM-DSV4-DECODE-LEVER2` (that was −22%: per-block re-quant of the shared row under the cudagraph) — ds4 itself keeps Q8_K quant a SEPARATE stage; the lever was the quant kernel's launch geometry, not fusion. Q8_0 GEMV (53.5%) untouched. See docs/BENCHMARKS.md `CLAIM-DSV4-Q8K-PREQ`. - -**DeepSeek-V4-Flash decode — MHC-pre FP64→FP32 fold LANDED: 85.4%→87.5% of ds4 ~16.5, bit-exact (2026-08-03, `CLAIM-DSV4-MHC-FOLD`, GB10 sm_121a).** The MHC glue (`MhcPreFinish`/`MhcPreDots`/`MhcPost`) was 11% of decode GPU time (nsys, `CLAIM-DSV4-DECODE-NSYS`) and ran the mix dot-products + sqrsum/norm reductions in **FP64**, which GB10 throttles to ~1/32–1/64 of FP32. New `MhcPreDotsFloatKernel`/`MhcPreFinishFloatKernel` mirror ds4's `hc4_split_one` + `hc_split_weighted_sum_norm_fused_kernel` in **float** (the same algebra ds4 runs — its fused kernel is float throughout; the Sinkhorn was already float, so only the mix dot + the sqrsum/norm reductions move double→float), gated `VT_V4_MHC_FUSED` (**default ON** per parity-enablers-ship-as-defaults; `=0` restores the bit-faithful FP64 path for a same-binary A/B). **Bit-exact on GB10:** decode ids token-IDENTICAL `VT_V4_MHC_FUSED=1` vs `=0` on the real IQ2XXS UD checkpoint; the unit gate is a characterized near-tie (float reduction reorder, `RelL2 < 1e-3`, `test_cuda_deepseek_v4` Brick B). **MEASURED decode: 85.4% → 87.5% of ds4 ~16.5** (GB10, drop_caches paired reps; recomputed onto the corrected ~16.5 bar from the 17.13-anchored 82.3%→84.3%, measured tok/s unchanged). Read once process-wide so the captured decode graph bakes a consistent path. See docs/BENCHMARKS.md `CLAIM-DSV4-MHC-FOLD`. - -**DeepSeek-V4-Flash decode — norm+RoPE FP64→FP32 fold LANDED: 14.02→14.87 tok/s (+6.1%), byte-exact (2026-08-03, `CLAIM-DSV4-ROPE-FLOAT`, GB10 sm_121a, base `6576814b`).** The fused per-head/kv norm+RoPE kernel (`NormRopeRowsKernel`, the resident-decode default from Brick 7) was already ONE launch, but ran the RMS reduction + the RoPE `pow`/`cos`/`sin` in **FP64** AND recomputed each pair's theta via a DOUBLE `O(pairs²)` left-fold recurrence (`theta_extrap *= theta_scale`) — GB10 throttles FP64 transcendentals ~1/32–1/64 of FP32. New `NormRopeRowsFloatKernel` mirrors ds4's `head_rms_norm_rope_tail_kernel` + `dsv4_qkv_rms_norm_rows_kv_rope_kernel` (`~/w8run/ds4/ds4_cuda.cu:5873/5779`) in **float** with ds4's DIRECT `powf(base,-i/r)` per pair (`powf(base,-2p/r) == theta_scale^p`, identical algebra), gated `VT_V4_ROPE_FLOAT` (**default ON** per parity-enablers-ship-as-defaults; `=0` restores the bit-faithful FP64 path for a same-binary A/B). Read once process-wide so the captured decode graph bakes a consistent path; NO ds4 attn_factor/mscale (we have none, matching the double kernel). **Byte-exact on GB10:** decode ids token-IDENTICAL `VT_V4_ROPE_FLOAT=1` vs `=0` on the real IQ2XXS UD checkpoint (60 tok, coherent "…Paris."); the unit gate is a characterized near-tie (`RelL2 < 1e-3`, `test_cuda_deepseek_v4` Brick 7 20/20·67073). **MEASURED (nsys 2-length 40↔140 diff, `--cuda-graph-trace=node`): the fused kernel 4.576 → 0.459 ms/step (~10×, 6.4%→0.7% of GPU time), 129 launches/step (3 calls × 43 layers) unchanged; decode 14.02 → 14.87 tok/s (+6.1%, 3 tight reps, 200-tok single-load) → 85.0% → 90.1% of ds4 ~16.5** (exceeds the ~14.7 lever-3 target; bar corrected from the unreproduced 17.13 anchor). See docs/BENCHMARKS.md `CLAIM-DSV4-ROPE-FLOAT`. - -**DeepSeek-V4-Flash decode — MHC-pre finish occupancy + sqrsum-fold LANDED: +0.7%, byte-exact; ds4-parity MHC bounded by sequential single-block launches (2026-08-03, `CLAIM-DSV4-MHC-LEAN`, GB10 sm_121a, base `6576814b`).** Lever 2. Post-lever-3 nsys attribution: `MhcPreFinishFloatKernel` is the #1 non-GEMM decode kernel — **9.5% of GPU time, ~6.4 ms/step** — launched `<<<1,256>>>` (ONE block = one SM) **86× SEQUENTIALLY per step** (2 sub-blocks × 43 layers; the layer chain is data-dependent so they cannot overlap). Two ds4-aligned micro-levers, gated `VT_V4_MHC_LEAN` (default ON=1024; `=0`→256 baseline; `512` selectable): (1) widen the finish block 256→1024 (32 warps) to hide HBM latency on its lone SM; (2) fold Σresidual² into `MhcPreDotsFloatKernel` block 0 (it already streams the whole residual) so the finish skips its DUPLICATE residual pass — mirroring ds4's mix-stage rms fold. **Byte-exact:** decode ids token-IDENTICAL `VT_V4_MHC_LEAN=1024` vs `=0` (60 tok; the sqrsum fold is the same 256-thread reduction relocated; the wide-block norm reorder is argmax-robust). Unit gate `test_cuda_deepseek_v4` 20/20·67073. **MEASURED (GB10): `MhcPreFinish` 6.52→5.98 ms/step (−8.3%); decode 14.85→14.96 tok/s (+0.7%, 3 tight reps each) → 90.0%→90.7% of ds4 ~16.5** (bar corrected from the unreproduced 17.13 anchor). HONEST BOUND: the kernel barely moved despite 4× the warps (−5.8%) AND a removed residual pass (extra −2.5%) — it is NOT warp-latency- or memory-pass-bound but **floored by the 86 sequential single-block launches**. The task's aspirational "fuse the whole MhcPre stack into ONE kernel mirroring ds4's `hc_split_weighted_sum_norm_fused`" would trade 258→86 MHC launches/step BUT serialize the 24 mix-dots into one block (net-uncertain, numerics-delicate) and leave the 6 ms finish core single-block; ds4's own decode kernel is also one-block-per-row yet cheaper per call (an ncu-level per-call analysis + a layer-chain-spanning restructure is the owed residual, high-risk). See docs/BENCHMARKS.md `CLAIM-DSV4-MHC-LEAN`. +**DeepSeek-V4-Flash decode levers (2026-08-03, byte-exact, default-ON; SUPERSEDED by the 1.14x BINDING below).** The Q8_K-preq launch-geometry port (+5.4%), the MHC-pre and norm+RoPE FP64->FP32 folds and the MHC-lean occupancy widen climbed decode ~13.5 -> 14.96 tok/s toward ds4 ~16.5; full per-lever forensics in docs/BENCHMARKS.md (`CLAIM-DSV4-Q8K-PREQ` / `-MHC-FOLD` / `-ROPE-FLOAT` / `-MHC-LEAN`). **BINDING 2026-08-05: `VT_V4_RESIDENT_W` (default-ON) BEATS ds4 — 18.69 vs 16.33 tok/s (1.144x), byte-exact:** the dense Q8_0 MLA/shared/lm_head proj tower was read from GGUF-mmap over ATS; staging it `cudaMalloc`-device once (Q8_0 per-launch ~20% each, ids-IDENTICAL) lifts decode 16.23→18.69 (median-of-3, drop_caches, PEAK flat 86.68 GiB, net move). Mirrors Laguna `VT_LAGUNA_RESIDENT_BF16W`; our GEMV was ATS-bound, not at ds4 parity. **Phase-2 routed-expert residency (`VT_V4_RESIDENT_EXPERTS`) MEASURED NEGATIVE, HELD default-OFF (2026-08-05):** the ~70 GiB IQ2/Q2_K expert slabs staged device-resident (madvise move-semantics) are byte-exact but **−3.4% steady** (18.76 vs 19.43 tok/s, same-binary median-of-3) + a one-time capture cost — the grouped-MoE kernels are dequant/latency-bound (~19-24% of DRAM peak), so residency (a bandwidth lever) cannot help, and device pinning cuts pool headroom (103→30 GiB avail). Prior (superseded, see record): PARITY 16.28 vs 16.33 (0.997x); `VT_V4_MHC_SINK4` +4.6% byte-exact; HC-expand + f16-DSA held default-OFF. diff --git a/scripts/check-agent-record.py b/scripts/check-agent-record.py index a6b294d9..353c9b31 100644 --- a/scripts/check-agent-record.py +++ b/scripts/check-agent-record.py @@ -252,7 +252,7 @@ # the same change — the heads/`SequencePooler`/`DispatchPooler` composite — not a # new row.) Bumped for a real new row, never to make a failing state transition # pass. -ENGINE_ROWS = 131 +ENGINE_ROWS = 140 MATRIX_PATHS = [ENGINE_MATRIX, *(path for path, _ in MATRICES.values())] REQUIRED = [ diff --git a/scripts/check-public-doc-tables.py b/scripts/check-public-doc-tables.py index 515d920d..4bdc653a 100755 --- a/scripts/check-public-doc-tables.py +++ b/scripts/check-public-doc-tables.py @@ -321,7 +321,7 @@ def features_errors(text: str) -> list[str]: # Lowering these numbers as the page is compacted is the gate closing. STATUS = ROOT / "docs/STATUS.md" STATUS_RATCHET = { - "chars": 287_838, + "chars": 282_066, "h2_sections": 11, "long_paragraphs": 89, "oversized_cells": 47,