diff --git a/.agents/NOW.md b/.agents/NOW.md index 809223e9b..f662d5244 100644 --- a/.agents/NOW.md +++ b/.agents/NOW.md @@ -9,8 +9,7 @@ benchmark record. Budget: 100 lines. ## Live claims -Working head: `row/backend-rocm-w0` (#41). Prior: benchmark checkpoint -`bench/qwen35-upstream-rebenchmark-20260805` on `upstream/main` @ `59674cf1d`. +Work: exact-chunks on main `1ce0d662b`; sm_120 measured at `3d2581551`. | Claim / track | State | Next command or step | |---|---|---| @@ -20,7 +19,7 @@ Working head: `row/backend-rocm-w0` (#41). Prior: benchmark checkpoint | MiniMax-H3 lane | **bf16 shards STREAM both towers (DiT + encoder); Q4_K_M enc cond cos 0.9975, 3.5° med, DIFFUSE** | render A/B on saved embeds | | Kimi-Linear-48B | **ROW 7 fold LANDS (#122 §21): engine==CLI 128/128; golden 122/128; SACRED green; v13 tokens ABI** | ACTIVE: 19.0 tok/s vs vLLM ~21 (~0.90×) | | 35B fresh grid | **BOUND** @`1ea26427`: 0.93-1.03x, c16 0.93x. INTAKE + Option A both NEGATIVE | Lever left: prefill glue (#61) | -| Qwen3.5-4B revalidation | 0.9971x @`59674cf1` (#35); TTFT/PSS pass, TPOT/ITL open | `docs/bench-evidence/` | +| Qwen3.5-4B sm_120 | Exact chunks ON: rebased-main reprofile 3.072x kernel / +2.272% run; sealed-vLLM throughput 1.021x PASS. Latency/VRAM OPEN | Spike residual 1.609x conv gap | | RPi5 A76 CPU | **R5 asm GREEN; llama NOT MET**: 0.461x pf, 0.653x dec, RSS -24% | W6: BF16 GEMM | | MXFP4 parity | c1 1.020, c2-c8 0.962-0.969. **#82 CLOSED: ptxas-lineage REFUTED (A/B ties our+vLLM PTX all ptxas/JIT; +10us=engine context, not codegen)** | TERMINAL: at parity | | ROW-SERVE-ASYNC-DENSE-MIRROR | **LANDED+dgx-VERIFIED** (`f9c969ae`): async mirror on classic dense Qwen3; SACRED 184/184 | Residual: sibling scope one-liner | @@ -30,7 +29,7 @@ Working head: `row/backend-rocm-w0` (#41). Prior: benchmark checkpoint | `BACKEND-ROCM` | **(b) fix in; #140 gfx1201 hipBLAS + Gemma-4 MoE landed (contributor, authorship-preserved); W0 green 4 archs** | compile + M2 ([spec](specs/rocm-unified-memory-b.md)) | | TP spike #287 (PR #143) | **LANDED** ([spec](specs/tensor-parallelism-spike.md)); DSpark rider grounded | dispatch TP-W1 (CPU-able) | | Release | SPIKE; 30/30 | #129 | -| Surface coverage (`ARCH-ONE-SURFACE`) | ROW 8 + #139 IN; **ROW 6 IN REVIEW (#137): embeddings LIVE — `LlamaModel` arch, PoolingRunner in the step, `vllm_embed` v15, `/v1/embeddings`, fold gate 4/4-231, 9 kills** | Merge #137; real-ckpt oracle cosine residual | +| Surface coverage (`ARCH-ONE-SURFACE`) | ROW 8 + #139 IN; **ROW 6 LANDED (#137): embeddings LIVE — `LlamaModel` arch, PoolingRunner in the step, `vllm_embed` v15, `/v1/embeddings`, fold gate 4/4-231, 9 kills** | Real-checkpoint oracle cosine residual | In-flight (default-OFF, not pushed): `laguna-fp4proj-prod`, laguna bf16/legacy/pipeline-gemv, `ds4-hc-expand-fuse`. @@ -49,8 +48,8 @@ both gate models, reproduced 2–3x on an idle box. See [gates.md](gates.md) and 0. **`ROAD-V1-MEM`** KV auto-sizing spike LANDED (`specs/kv-sizing.md`, `READY`). 1. **Spike the Parakeet encoder row** (vLLM carries it inside `nano_nemotron_vl.py`; the transducer half is NOT in vLLM: separate call). -2. **Qwen3.5-4B serving follow-up:** bind the default-ON async-serving path - against the same oracle before attributing the remaining TPOT gap. +2. **Qwen3.5-4B sm_120:** rebased branch is GREEN and reprofiled. Spike the + residual 1.609x conv gap; latency/VRAM and gate models stay open. 2. **Merge the invocation-parity prevention** (CI guard + AGENTS.md checklist); CUDA build-verify the byte-exact `kGemvHeuristicAlgos` refactor on dgx. 3. **Same-tool re-verify deepseek_v4's bf16 resident tower** (the one other diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md index 7f5b9b04b..0bcd79d9e 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -15124,8 +15124,6 @@ llama.cpp's Vulkan" when it is really "our CPU tier vs llama.cpp's Vulkan". The comparison becomes meaningful when native coverage closes — the progress metric is `vt::GetReferenceTierHits()` reaching 0, and the ops that matter for this model are the RoPE table build, the sampler tail, and the remaining norm/glue set. ->>>>>>> 814230a0 (bench(vulkan): VK-E unblocked with identical weights; ours quoted as NO RATIO) - #### CORRECTION (2026-08-07, same session): the vllm.cpp Vulkan arm is GPU-BOUND, not CPU-bound The entry above attributes vllm.cpp's slowness to "our HOST FALLBACK wearing a @@ -15983,3 +15981,76 @@ loads, or caching the K/V slice across the query heads that share a KV head — Recorded because the hypothesis was specific and the refutation is reusable: this is the second kernel this session where a barrier-count argument looked compelling and measured flat (the subgroup GEMV was the first). + +## 2026-08-07 — sm_120 exact GDN causal-conv chunks: 3.069x kernel, +2.152% enclosing + +**Disposition:** ACCEPTED and reproduced on clean current-main transplant +`upstream/main` `f91a5917a`. Exact chunks default ON; latency, VRAM and +27B/35B gates stay open. + +The metadata builder enumerates exact `(sequence, 8-token chunk)` work, uploads +its two i32 descriptors once per step and shares them across GDN layers. The +register kernel consumes one descriptor per `grid.y`. +`VT_CONV_EXACT_CHUNKS=0` is the same-binary whole-sequence rollback; +`VT_CONV_REG=0` selects tiled/scalar. + +**Correctness.** RED compile evidence: +`/tmp/vllm-agent-runs/gdn-exact-red-escalated.json`. Focused host metadata and +flags, affected Qwen fixtures, full CUDA GDN and cached Qwen3.5-4B 3/3·1672 +were green. Three production ON/OFF pairs had identical token files for all +128 requests × 128 outputs. + +**Same-binary profile.** Manifest +`/tmp/vllm-agent-runs/qwen35-conv-exact-ab-profile.json`; rollback/default +traces `/tmp/qwen35-conv-exact-{off,on}.nsys-rep`. Rollback: 1728 calls, +720.047171 ms, 416.694 us mean. Exact: 1728 calls, 234.607112 ms, 135.768 us. +That is **3.069x**, saving 485.440 ms. Exact `grid.y=279/280/282` work counts +replace dominant `(64,28..32,1)` whole-sequence grids, confirming the proposed +mechanism. Pinned-vLLM same-tool total is 145.421 ms; residual **1.613x**. + +**Enclosing A/B.** Manifest +`/tmp/vllm-agent-runs/qwen35-conv-exact-local-ab.json`; evidence root +`/tmp/qwen35-conv-exact-local-ab-20260807`. Three alternating pairs under one +GPU lock and a 25 GiB user-systemd scope: + +| Axis | rollback | exact default | ratio | +|---|---:|---:|---:| +| total | 6641.800 tok/s | 6784.743 tok/s | 1.02152x | +| output | 734.433 tok/s | 750.237 tok/s | 1.02152x | +| TTFT | 1048.927 ms | 1018.040 ms | 0.97055x | +| TPOT / ITL | 35.420 ms | 34.740 ms | 0.98080x | +| E2E | 5547.687 ms | 5430.180 ms | 0.97882x | + +Against sealed vLLM, exact local is **1.021246x** throughput, +**1.085812x** TTFT and **1.024597x** TPOT. Exact peak VRAM +13044/13058/13058 MiB, mean 13053.3, versus old local 13054 and vLLM 12820. + +**VOID oracle attempts.** `qwen35-conv-exact-default-full-compare.json` exposed +the missing live-driver link path; the harness now adds `/run/opengl-driver/lib` +to `LIBRARY_PATH`. `qwen35-conv-exact-default-full-compare-rerun.json` reached +13/18 legs before a transient Torch bytecode read invalidated Triton AOT cache +keys. Neither attempt supersedes the sealed denominator. Full evidence: +`docs/bench-evidence/qwen35-4b-sm120-main-20260807.md`. + +**Clean-transplant reproduction (`f91a5917a`).** Contained CPU/CUDA rebuild; +focused CPU 6/6, full CUDA GDN 66/66·4300, cached 4B 3/3·1672. Same-binary +graph-node traces `/tmp/qwen35-conv-exact-transplant-{off,on}.nsys-rep` and +token files reproduce the mechanism with byte identity: rollback 1728 calls / +720.216507 ms / 416.792 us, exact 1728 / 234.379395 ms / 135.636 us = +**3.072866x**. Profiled enclosing totals are 6587.66→6727.35 tok/s +(**1.021205x**), TTFT 1058.73→1025.46 ms, TPOT 35.70→35.04 ms and E2E +5592.69→5475.91 ms. The transplanted result is therefore reproduced, not merely +carried from its old branch. Against the sealed vLLM conv trace, residual is +**1.611730x**. + +**Post-rebase reproduction (`3d2581551` on `upstream/main` `48a54141f`).** The +contained rebuild and all three gates remain green: focused 6/6, CUDA GDN +66/66·4300, cached 4B 3/3·1672. Fresh graph-node traces and token files under +`/tmp/qwen35-conv-exact-rebase-3d2581551-{off,on}.*` are byte-identical. +Rollback is 1728 calls / 718.704016 ms / 415.917 us; exact is 1728 / +233.954533 ms / 135.390 us = **3.07198x**. Profiled enclosing totals are +6589.65→6739.34 tok/s (**1.02272x**), TTFT 1057.63→1022.70 ms, TPOT +35.70→34.99 ms and E2E 5590.92→5466.20 ms. The result therefore survives the +27-commit main advance; against the sealed vLLM conv trace the residual is +**1.60881x**. Trace SHA-256: rollback `6a5dde18e...f97c47`, exact +`f47fb9cc...7aecf9`; both token files `83fcdc45...453545`. diff --git a/.agents/coordination.md b/.agents/coordination.md index 335539aef..172cd1702 100644 --- a/.agents/coordination.md +++ b/.agents/coordination.md @@ -1619,7 +1619,7 @@ items a-runner/b stay with the async/GDN `runner.cpp` owners. | `CLAIM-POOLING` | `ENG-POOLER-SEQ` (INVENTORIED-implicit→ACTIVE, W1→**W2**), `ENG-POOLING-RUNNER` (**NEW row, ACTIVE, W3**), `SERVE-POOLING-ENDPOINTS` (INVENTORIED→SPIKE) | Claude Code (opus-4-8) | isolated worktree `.claude/worktrees/claim-pooling-w2w3` (CPU build `build-cpu` `-DVLLM_CPP_CUDA=OFF -DVLLM_CPP_SERVER=ON` Release + CPU run; NO dgx/GPU — the pooling reductions + activations + runner are host arithmetic) | branch `claim-pooling-w2w3`, base `main` `edf68c91` (confirmed via `git rev-parse HEAD`) | Pooling task class HIGH-priority feature-gap #2. W0 spike + W1 CPU pooler OP (prior pass); **W2 pooler HEADS composite + `SequencePooler`/`DispatchPooler` + `PoolerConfig`/`PoolingParams` and W3 pooling RUNNER path (this pass).** Owns ONLY: NEW `include/vllm/model_executor/layers/pooler/{pooling_metadata,methods,activations,common,pooling_params,pooler_config,heads,poolers,dispatch_pooler}.h` + `src/vllm/model_executor/layers/pooler/{methods,activations,heads,poolers,dispatch_pooler}.cpp` + NEW `include/vllm/v1/worker/gpu/pool/pooling_runner.h` + `src/vllm/v1/worker/gpu/pool/pooling_runner.cpp`; NEW `tests/vllm/model_executor/layers/pooler/{test_pooler,test_pooler_heads}.cpp` + `tests/vllm/v1/worker/gpu/pool/test_pooling_runner.cpp`; `CMakeLists.txt` (6 source lines) + `tests/CMakeLists.txt` (3 tests); NEW `.agents/specs/pooling-task-class.md`; the `ENG-POOLER-SEQ` + NEW `ENG-POOLING-RUNNER` engine-matrix rows + `SERVE-POOLING-ENDPOINTS` note + engine Serving/Total rollup (Serving 21→22/ACTIVE 6→7, Total 130→131/ACTIVE 47→48) + `scripts/check-agent-record.py` ENGINE 130→131; the record surfaces (this row, `roadmap_v1.md` gap #2, `docs/STATUS.md`, `docs/BENCHMARKS.md`, `feature-matrix.md` MODEL-POOLING note, `parity-ledger.md`, `state.md`). **NON-COLLISION:** additive NEW files only — the sole edits to existing compiled headers are ADDITIVE (methods.h defaulted virtuals, pooling_metadata.h new fields); ZERO edits to any existing production forward/runner path; NO pooling MODEL row created (concrete embedding model + real-oracle cosine gate is the named W3-model residual), so README/Metal/model-matrix rows untouched. | `ACTIVE` | 2026-07-29 — **W2 + W3 LANDED + CPU-GATED (foreground, NOT pushed).** `test_pooler_heads` 27/27 (240 asserts, Embedding/Classifier heads + SequencePooler + DispatchPooler incl. mixed embed+classify batch + ctor validation) and `test_pooling_runner` 5/5 (14 asserts, runner path + STRUCTURAL cosine-parity gate vs double-precision LAST+normalize ref) — plus W1 `test_pooler` 17/17 unchanged. RED-first proven: disable matryoshka slice + logit_mean → 8 cases/50 asserts fail (heads); CLS-instead-of-LAST drops cosine <0.5 + disable normalize → 2 unit-L2 asserts fail (runner). Clean CPU `-Wall -Wextra -Werror` 0-warn full-library build. **HONEST RESIDUAL:** the cosine gate is STRUCTURAL (synthetic weights) — the real-model `vllm.LLM(task="embed").encode` oracle cosine gate needs a registered concrete embedding model forward (W3-model, no number fabricated). Residuals (spec §Work breakdown): concrete pooling MODEL + real-oracle cosine gate (W3-model), endpoints /v1/embeddings+score+rerank+classify (W4), tokwise AllPool/StepPool (W5). Prior 2026-07-28 — W0 spike + W1 pooler OP LANDED + CPU-GATED: `test_pooler` 17/17 (50 asserts) vs double-precision refs, RED-first proven. | | `CLAIM-DSV4-GGUF-LOADER` | `QUANT-GGUF-IQ2_XXS` (INVENTORIED→ACTIVE), `QUANT-GGUF-Q2_K` (INVENTORIED→ACTIVE); cross-refs `MODEL-TEXT-deepseek-v4-deepseek-v4-for-causal-lm` (stays `SPIKE`, owned by `CLAIM-DEEPSEEK-V4-IMPL`) | Claude Code (opus-4-8) | isolated worktree `.claude/worktrees/gguf-iquant-dsv4` (CPU-only `build-cpu` `-DVLLM_CPP_CUDA=OFF`; NO GPU, NO 90 GB download — the dequant unit gate uses known packed bytes; the GGUF header was HTTP-range-read, no download) | branch `feat/gguf-iquant-dsv4`, base `main` `4d1be010` (confirmed via `git rev-parse HEAD`) | GGUF IQ2_XXS + Q2_K dequant, the DeepSeek-V4-Flash single-Spark GGUF quant-path brick (W1). Owns ONLY the GGUF/quant PATH (NOT the forward — the forward TUs stay owned by `CLAIM-DEEPSEEK-V4-*`): the two `dequantize_row_*` decoders + grid/sign tables in `src/vt/cpu/cpu_quant_dequant.cpp`; the `kQ2_K`/`kIQ2_XXS` vt block dtype registration in `src/vt/dtype.{h,cpp}` + `src/vt/ops.cpp`; the id-16 reader trait in `gguf_reader.cpp` + ids-10/16 dispatch in `gguf_dequant.cpp`; `tests/vllm/test_gguf_dequant.cpp` + `tests/vt/test_ops_quant_traits.cpp`; the two `QUANT-GGUF-*` rows; NEW `.agents/specs/gguf-iquant-dsv4.md`; the V4 GGUF-loadable note on the model-matrix V4 row (row stays SPIKE); the record surfaces. **NON-COLLISION:** additive within the existing GGUF dequant switch + vt block table — does NOT touch any DeepSeek-V4 forward TU (`deepseek_v4.{cpp,h}`/`_dsa`/`_weights`), README, or Metal; the k-quant/NVFP4 decoders are byte-unchanged. | `ACTIVE` | 2026-07-29 — **W1 LANDED + CPU-GATED (foreground, NOT pushed).** IQ2_XXS (id 16, codebook `iq2xxs_grid`+signs+4-bit scale) + Q2_K (id 10, nibble sub-scale/min) ported 1:1 from llama.cpp `ggml-quants.c` `237ad9b96`; both DEQUANT-ONLY (no vec_dot ⇒ route to expand-bf16). `test_gguf_dequant` **15/15·480** (hand-derived literals: IQ2_XXS grid[1] byte0=0x2b→5.375, ksigns[1] flips j=0,7→±3.0, db 0.125/0.375; Q2_K 5.75/-0.25/2.5/0.25) + `test_ops_quant_traits` **9/9·5643** (dequant-only contract). All 7 changed TUs clean under full `-Werror`; the `voxtral.cpp` GCC-13 `-Werror=array-bounds` FP PROVEN pre-existing (fails at base with this diff's `dtype.h` reverted), neutralized only to link the test binaries. **W2 (V4-GGUF loader) DERIVED not landed:** HTTP-range-read the real `UD-IQ2_XXS` header — `general.architecture=deepseek4`, `general.file_type=19` (=IQ2_XXS), `split.tensors.count=1328`, full `deepseek4.*` config-KV schema; the tensor NAME manifest is beyond the CDN range cap + uncached ⇒ the V4 registry GGUF reject STAYS. Residuals: V4 forward (W3-W8, multi-Spark) + the V4-GGUF name map (W2, manifest-blocked) + a vec_dot perf leaf. | -| `CLAIM-EMBEDDINGS-ONE-SURFACE` | `ENG-POOLING-RUNNER` (live engine-step invocation), `SERVE-POOLING-ENDPOINTS` (SPIKE→ACTIVE, `/v1/embeddings`), `MODEL-EMBED-llama-llama-for-causal-lm` (INVENTORIED→ACTIVE; PARTIAL on merge — only the `LlamaModel` membership registered); cross-refs `ENG-POOLER-SEQ` (stays `CLAIM-POOLING`, ops untouched) | Claude Code (fable-5) helper, task #285 | isolated worktree `/home/mudler/_git/vllm.cpp-embeddings-one-surface` (CPU-only; lean per-target builds under disk pressure) | branch `row/EMBEDDINGS-ONE-SURFACE`, base `main` `b44ad337`, DRAFT PR #137 (the reservation) | ARCH-ONE-SURFACE fold ROW 6: embeddings/pooling through the ONE surface. Owns: NEW `src/vllm/model_executor/models/llama_embedding_registry.cpp` + `LoadLlamaModelEmbeddingWeights` (llama_weights.cpp) + `Qwen3DenseModel::ForwardHidden` (qwen3.{h,cpp} additive tail); the ADDITIVE task-gated pooling plumb (`LoadedModel::pooler()`, runner `pooling_runner_`+`pool_tokens`, `Request/EngineCoreRequest::pooling_params`, `ModelRunnerOutput::pooler_output`, scheduler pooling stop, `EngineCoreOutput/RequestOutput::pooling_output`, `LLMEngine::add_pooling_request/embed`, `ResolveAsyncEnabled(is_pooling_model)`); `vllm_embed`/`vllm_embedding_result_free` ABI v15 (vllm.h + vllm_c.cpp incl. the refuse-both-directions guards + the v13 `vllm_complete_tokens` missing-guard fix); `handle_embeddings` + `set_embedder` + task-conditional route (api_server.{h,cpp}) + server main pooling dispatch; NEW fixture `tests/vllm/models/fixtures/llama_embed_e2e` + `scripts/mm/llama_embed_fixture_gen.py` + `tests/vllm/models/test_llama_embedding_fold.cpp`; test/guard updates (test_capi v15 section + floor pin >= 15, test_dlopen symbols, c_header_compile.c, test_api_server embeddings section, test_model_registry/gguf arch pins, check-supported-models ARCH_TOKEN_RE); allowlist row removal + FEATURES/STATUS/BENCHMARKS rows + matrices + specs. **NON-COLLISION:** every engine hook is task-gated on `is_pooling_model`/`pooling_params` (nullopt/false = byte-identical text path); no SACRED path rewritten; no example added. | `ACTIVE` | 2026-08-08 — CPU-LANDED on the branch: fold gate `test_llama_embedding_fold` 4/4-231 (engine path == direct registry path + f64 LAST+normalize ref + chunked is_valid arm), `test_capi` 48/48-462, `test_dlopen` 30/30, server suite 50/50, registry 24/24-820, engine suites green (scheduler 423, llm_engine 204, engine_core 44, output_processor 77, qwen3_forward 1557, async_llm 342, llama_forward 509); 9 mutation kills (floor pin, refuse both directions, route gating both ways, engine-step invocation, scheduler stop, registry info pin, async-off wire). RESIDUAL: real embedding checkpoint + `LLM(task="embed")` oracle cosine. | +| `CLAIM-EMBEDDINGS-ONE-SURFACE` | `ENG-POOLING-RUNNER` (live engine-step invocation), `SERVE-POOLING-ENDPOINTS` (SPIKE→ACTIVE, `/v1/embeddings`), the `LlamaModel` embedding membership (INVENTORIED→PARTIAL on merge — only that membership registered); cross-refs `ENG-POOLER-SEQ` (stays `CLAIM-POOLING`, ops untouched) | Claude Code (fable-5) helper, task #285 | isolated worktree `/home/mudler/_git/vllm.cpp-embeddings-one-surface` (CPU-only; lean per-target builds under disk pressure) | branch `row/EMBEDDINGS-ONE-SURFACE`, base `main` `b44ad337`, PR #137 MERGED | ARCH-ONE-SURFACE fold ROW 6: embeddings/pooling through the ONE surface. Owns: NEW `src/vllm/model_executor/models/llama_embedding_registry.cpp` + `LoadLlamaModelEmbeddingWeights` (llama_weights.cpp) + `Qwen3DenseModel::ForwardHidden` (qwen3.{h,cpp} additive tail); the ADDITIVE task-gated pooling plumb (`LoadedModel::pooler()`, runner `pooling_runner_`+`pool_tokens`, `Request/EngineCoreRequest::pooling_params`, `ModelRunnerOutput::pooler_output`, scheduler pooling stop, `EngineCoreOutput/RequestOutput::pooling_output`, `LLMEngine::add_pooling_request/embed`, `ResolveAsyncEnabled(is_pooling_model)`); `vllm_embed`/`vllm_embedding_result_free` ABI v15 (vllm.h + vllm_c.cpp incl. the refuse-both-directions guards + the v13 `vllm_complete_tokens` missing-guard fix); `handle_embeddings` + `set_embedder` + task-conditional route (api_server.{h,cpp}) + server main pooling dispatch; NEW fixture `tests/vllm/models/fixtures/llama_embed_e2e` + `scripts/mm/llama_embed_fixture_gen.py` + `tests/vllm/models/test_llama_embedding_fold.cpp`; test/guard updates (test_capi v15 section + floor pin >= 15, test_dlopen symbols, c_header_compile.c, test_api_server embeddings section, test_model_registry/gguf arch pins, check-supported-models ARCH_TOKEN_RE); allowlist row removal + FEATURES/STATUS/BENCHMARKS rows + matrices + specs. **NON-COLLISION:** every engine hook is task-gated on `is_pooling_model`/`pooling_params` (nullopt/false = byte-identical text path); no SACRED path rewritten; no example added. | `DONE` | 2026-08-08 — MERGED in PR #137: fold gate `test_llama_embedding_fold` 4/4-231 (engine path == direct registry path + f64 LAST+normalize ref + chunked is_valid arm), `test_capi` 48/48-462, `test_dlopen` 30/30, server suite 50/50, registry 24/24-820, engine suites green (scheduler 423, llm_engine 204, engine_core 44, output_processor 77, qwen3_forward 1557, async_llm 342, llama_forward 509); 9 mutation kills. RESIDUAL moved to the PARTIAL model row: real embedding checkpoint + `LLM(task="embed")` oracle cosine. | diff --git a/.agents/engine-matrix.md b/.agents/engine-matrix.md index 82a97d9b3..5ae1d8a34 100644 --- a/.agents/engine-matrix.md +++ b/.agents/engine-matrix.md @@ -204,7 +204,7 @@ claims it. | `SERVE-HTTP-TRANSPORT` | Serving-socket transport parity: mirror vLLM's uvicorn/asyncio default `TCP_NODELAY` on every accepted SSE socket so per-token stream frames are not held by Nagle against the peer's delayed ACK. Implemented + CPU-tested; the non-binding localhost A/B sizing is COMPLETE and NEUTRAL within noise on c1/c2 ITL/TPOT/throughput (loopback ACKs are instant, so Nagle never coalesces ~100 ms-cadence token frames) — no gate-axis credit expected; the mirror stays for real-network parity. Future keep-alive / read-write-timeout / listening-socket option parity noted, not done | T0 | vLLM serves via uvicorn over asyncio `vllm/entrypoints/launcher.py:71,76`, `vllm/entrypoints/openai/api_server.py:591,630`; asyncio disables Nagle per accepted TCP stream socket `asyncio/base_events.py:192-197` (`_set_nodelay`) called from `asyncio/selector_events.py:950`; cpp-httplib default-off `third_party/httplib/httplib.h:142`, applied on accept only when set `third_party/httplib/httplib.h:12083` | `src/vllm/entrypoints/openai/api_server.cpp:69` (`set_tcp_nodelay(true)` in the ApiServer setup) | behavioral accepted-socket `getsockopt(TCP_NODELAY)` case `tests/vllm/entrypoints/openai/test_api_server.cpp:1076` (helper `:380`); RED accepted `TCP_NODELAY` 0 → GREEN 1, full `test_openai_api_server` **22/22 cases / 242 assertions**; non-binding sizing root `~/work/vllm.cpp-tcpnodelay-sizing/ff915e8…` (raw-set SHA `f5b52900…2128`) neutral within noise; closure [ledger](parity-ledger.md#L451) | [serve-tcp-nodelay.md](specs/serve-tcp-nodelay.md) | `DONE` | `ff915e8` | | `SERVE-C-ABI` | Stable LocalAI-style C FFI (**19** exported `VLLM_API` symbols at `VLLM_ABI_VERSION 10`; blocking and nonblocking request handles. Count corrected 2026-07-24 from a stale `17`, which predated ABI v4/v5 adding `tool_parser`/`reasoning_parser` and the chat entry points; `include/vllm.h` is the source of truth and README:231 already said 19). **ABI v9 2026-07-28 (`CLAIM-CAPI-ENGINE-CONFIG-V9`): the ABI carried strictly LESS engine config than `EngineParams` does** - `max_num_batched_tokens`, the scheduler `scheduling_policy` (`fcfs` / `priority` / `lpm`), and `kv_transfer_config` (the external KV connector / LMCache JSON) were reachable from the bundled server's flags and from NO embedder. All three added, inert at their defaults (zero-filled v8 growth == byte-identical pre-v9 engine); the connector NAME is validated against `KVConnectorFactory` at load, mirroring the server's startup check. `tokenizer_config_path` stopped being a declared-since-v1 no-op and now selects the chat template's source file. Malformed `speculative_config`/`kv_transfer_config` documents now report `VLLM_ERR_INVALID_ARGUMENT` (the contract vllm.h documented since v6) instead of `VLLM_ERR_MODEL_LOAD`, via a catch scoped to the parse block so a real `FromModelDir` failure still reports MODEL_LOAD. Driver: the LocalAI vllm-cpp backend could not expose LMCache or the prefill budget in a model config) | T0 | Original project ABI; pinned vLLM has no C ABI | `include/vllm.h:143,181,207`; `src/capi/vllm_c.cpp:229,264,327,391` | `tests/capi/test_capi.cpp:320,428,505,574,606,640`; `tests/capi/test_dlopen.cpp:77,86`; `tests/capi/c_header_compile.c:1` | [c-api-library.md](specs/c-api-library.md) | `ANCHOR-BACKFILL` | `CLAIM-SERVE-C-ABI-SPIKE` | | `SERVE-CPP-API` | Rich `LLM` and `AsyncLLM` C++ API | T1 | `vllm/entrypoints/llm.py:66,422`; `vllm/v1/engine/async_llm.py:70` | - | - | `planned: specs/cpp-api.md` | `INVENTORIED` | - | -| `SERVE-CLI-BENCH` | Serve and latency/throughput/serve benchmark modes | T0 | `vllm/entrypoints/cli/serve.py:44`; `vllm/entrypoints/cli/benchmark/main.py:29` | separate binaries + explicit scheduler-capacity flags `examples/server/main.cpp:63,96,116,170`; `examples/bench/main.cpp:40,109`; `examples/bench/bench_core.h:96,468` | server help contract `examples/CMakeLists.txt:34`; benchmark `tests/examples/test_bench.cpp:15,48` | `planned: specs/cli-serve-bench.md` | `PARTIAL` | - | +| `SERVE-CLI-BENCH` | Serve and latency/throughput/serve benchmark modes | T0 | `vllm/entrypoints/cli/serve.py:44`; `vllm/entrypoints/cli/benchmark/main.py:29`; production queue `vllm/v1/engine/core.py:200-231,622-669` | separate binaries + explicit scheduler-capacity flags `examples/server/main.cpp:63,96,116,170`; production `AsyncLLM` benchmark frontend + auditable scheduler depth `examples/bench/bench_core.h:426,495,595`; `examples/bench/main.cpp:51` | server help contract `examples/CMakeLists.txt:34`; production-frontend and metric assertions `tests/examples/test_bench.cpp:18,29-32,61,81,97` | [CLI/serve/benchmark spike](specs/cli-serve-bench.md) | `PARTIAL` | - | | `SERVE-GATE-ONLINE` | Same-corpus online correctness, TTFT/TPOT/ITL, throughput and peak-memory gate vs vLLM v0.25.0 | T0 | `vllm/benchmarks/serve.py:1,581-615`; [v0.25 audit](sync/2026-07-12-702f481.md); `tests/benchmarks/test_serve_cli.py:1` | Schema-v5 harness plus [trace controller](../include/vt/cuda/cuda_profiler_control.h#L13), [production component driver](../scripts/dgx-gdn-packed-component.sh), and fail-closed [component finalizer](../tools/bench/gdn_packed_component.py) | **BINDING `9ecd9d0`: 114/124** (async default ON; mem 4/4, c1 20/20, c2 20/20, c16 19/20, c4 & c32 18/20, c8 15/20; `benchmark_binding` refers here, superseding `3f256ab` 55/124 and `246a23c` 49/124, both retained immutable). Two-grid totality with `f0fb727` (111/124) is 115/124 effective parity vs vLLM 0.25.0 (27B). Async CLOSED the c16/c32 ITL tails (ours now BEATS vLLM: c16 p99 1.055, c32 p90 1.034/p99 1.078) and leaves a stable c8 `p99_itl` ~0.86 residual, ROOT-CAUSED (2026-07-18, `CLAIM-C8-P99-TAIL-1`, [spec](specs/c8-p99-itl-tail-2026-07-18.md)) as IRREDUCIBLE-AS-MIRRORED: our deterministic synchronous forward keeps co-admitted c8 requests in byte-identical lockstep where vLLM's async-future jitter de-phases them; the c16/c32 INVERSION proves this is the trailing edge of the per-step determinism that wins c16/c32 + throughput, not a capability gap (scheduler + async placeholder byte-identical, `tests/vllm/v1/test_scheduler_wave.cpp:265`, [tail spec](specs/tail-stall-analysis-2026-07-16.md)). Full grid + per-binding forensics: roadmap_v1.md + parity ledger; no packed speed credit | [online serving gate](specs/cuda-online-serving-gate.md); [merged GDN projections](specs/gdn-merged-input-projections.md); [packed decode](specs/gdn-packed-decode.md) | `ANCHOR-BACKFILL` | CLAIM-SERVE-GATE-1 | | `SERVE-E2E-NIGHTLY` | Server conformance and real-model nightly suites for all release gates | T0 | `tests/entrypoints/openai/`; `tests/v1/e2e/`; `.buildkite/test-pipeline.yaml` | current unit/conformance tests only; no scheduled DGX suite | `tests/vllm/entrypoints/openai/test_conformance.cpp:1`; `tests/parity/test_qwen36_paged_engine.cpp:78`; `tests/parity/test_qwen27_paged_engine.cpp:110` | `planned: specs/server-e2e-nightly.md` | `INVENTORIED` | - | | `ENG-RELEASE-BINARIES` | Downloadable host-ABI-specific `vllm-server` bundles: adaptive CPU and fat CUDA primary artifacts, optional per-SM diagnostics, and literal-static feasibility boundary | T0 | vLLM release lanes `.buildkite/release-pipeline.yaml:1-18,34-170` @ `555967922`; release-image dependency boundary `docker/Dockerfile.cpu:262-290` | server target only `examples/CMakeLists.txt:54-64`; CPU per-TU/runtime-dispatch baseline `CMakeLists.txt:870-890`, `src/vt/cpu/cpu_matmul_elem.cpp:553-612`, `src/vt/cpu/cpu_quant_dot_arm.cpp:39-77`; cross-family CUDA fat/per-source-gencode and multi-SM AOT gaps remain; no install/archive/publish implementation | help smoke only `examples/CMakeLists.txt:59-63`; issue `#117`; user-reviewed fat-CUDA/adaptive-CPU matrix and gates in [release-binary-matrix.md](specs/release-binary-matrix.md) | [release-binary-matrix.md](specs/release-binary-matrix.md) | `SPIKE` | `CLAIM-ENG-RELEASE-BINARIES-SPIKE` | diff --git a/.agents/kernel-matrix.md b/.agents/kernel-matrix.md index 6a215b7c9..0a0cf92ec 100644 --- a/.agents/kernel-matrix.md +++ b/.agents/kernel-matrix.md @@ -164,6 +164,14 @@ host/sched. Detail: state `KERNEL-FA2-GQA-SWAP-FLIP`. | `KERNEL-DEPTHWISE-CONV1D` | **NON-CAUSAL depthwise `nn.Conv1d(C,C,K,groups=C)`**, the conformer convolution module's temporal mixer: centre-padded, stateless, activation-free, strided/dilatable. A deliberate SIBLING of `vt::CausalConv1dFwd` (Mamba/GDN: causal, persistent `conv_state`, folded SiLU), which is NOT modified — widening the causal op would have put a branch in a hot decode kernel and risked its byte-exactness | `modeling_parakeet.py:116` (padding :136, ctor :138-146, applied :180); vLLM-native sibling `conformer_encoder.py:229` | `vt::DepthwiseConv1d` / `OpId::kDepthwiseConv1d` (include/vt/ops.h:2101-2121, wrapper src/vt/ops.cpp:2262-2296); CPU kernel src/vt/cpu/cpu_conv1d_depthwise.cpp:63-95 | tests/vt/test_ops_conv1d_depthwise.cpp:1-338 — 5 cases / 1184 assertions, same byte-identity bar, both accepted weight layouts, plus a left-only-padding cross-check that the causal op reproduces element for element | [specs/parakeet-conformer-encoder.md](specs/parakeet-conformer-encoder.md) (row P2) | `ACTIVE` (CPU tier LANDED + byte-identity gated; the CUDA provider is the remaining work) | `CLAIM-PARAKEET-KERNELS-P1P3` | | `KERNEL-ATTN-RELPOS` | **Transformer-XL relative-position ENCODER self-attention** — no KV cache, no paging, no RoPE, non-causal. Every other attention path in `vt` is a decoder path (RoPE + paged/flash KV, `ops.h` kAttention/kPagedAttention/kMla*), so none could express it. The upstream `_rel_shift` pad/reshape/slice is carried as the CLOSED FORM `raw(i, T-1-i+j)`, so no `[T,2T-1]` scratch is materialised; the unit test's reference does the literal reshape, which is what proves it. The two upstreams differ only in where the scale lands, exposed as `AttentionRelPosArgs::scale_after_sum` rather than chosen | `modeling_parakeet.py:259` (forward :302-347, `_rel_shift` :349-355, `eager_attention_forward` :225-255); vLLM-native sibling `conformer_encoder.py:170` (:188-217, :179-186) | `vt::AttentionRelPos` / `OpId::kAttentionRelPos` (include/vt/ops.h:2123-2165, wrapper src/vt/ops.cpp:2298-2345); CPU kernel src/vt/cpu/cpu_attn_relpos.cpp:77-158 | tests/vt/test_ops_attn_relpos.cpp:1-450 — 7 cases / 368 assertions, byte-identity vs a LITERAL-`_rel_shift` reference, GQA ratios, padding mask, both scale placements, thread counts | [specs/parakeet-conformer-encoder.md](specs/parakeet-conformer-encoder.md) (row P3) | `ACTIVE` (CPU tier LANDED + byte-identity gated; the CUDA provider is the remaining work) | `CLAIM-PARAKEET-KERNELS-P1P3` | +**2026-08-07 `KERNEL-SSM-MAMBA` checkpoint.** Exact `(sequence, 8-token +chunk)` prefill-conv metadata/dispatch is default ON on sm_120 after byte-exact +gates. Rebased-main same-binary causal-conv is **718.704→233.955 ms (3.072x)** +and the profiled enclosing workload improves **2.272%**; pinned vLLM remains +145.421 ms (**1.609x open**). +Lifecycle stays `INVENTORIED` because generic Mamba coverage and the 27B/35B +release gates are unchanged. [Spec and evidence](specs/sm120-qwen35-conv-chunking-2026-08-07.md). + ## Count invariants - This table has exactly 35 practical kernel-family rows. diff --git a/.agents/model-matrix.md b/.agents/model-matrix.md index cc9d8d732..a827a38af 100644 --- a/.agents/model-matrix.md +++ b/.agents/model-matrix.md @@ -44,8 +44,8 @@ Rollup by lifecycle state (must equal the detailed per-state row counts): | State | Rows | |---|---| | INVENTORIED | 314 | -| PARTIAL | 19 | -| ACTIVE | 10 | +| PARTIAL | 20 | +| ACTIVE | 9 | | SPIKE | 6 | | BLOCKED | 5 | | DONE | 3 | @@ -293,7 +293,7 @@ Transformers compatibility is capability-driven and excluded from finite counts. | `MODEL-EMBED-bert-with-rope-gte-new-model` | `GteNewModel` | `registry.py:221`; `vllm/model_executor/models/bert_with_rope.py::GteNewModel` | embedding / text | encoder attention; pooler; FusedMoE/grouped GEMM | ☐ required | `INVENTORIED` | none | unassigned | | `MODEL-EMBED-jina-jina-embeddings-v5-model` | `JinaEmbeddingsV5Model` | `registry.py:222`; `vllm/model_executor/models/jina.py::JinaEmbeddingsV5Model` | embedding / text | encoder attention; pooler | ☐ required | `INVENTORIED` | none | unassigned | | `MODEL-EMBED-llama-llama-bidirectional-model` | `LlamaBidirectionalModel` | `registry.py:223`; `vllm/model_executor/models/llama.py::LlamaBidirectionalModel` | embedding / text | encoder attention; pooler; sliding-window attention | ☐ required | `INVENTORIED` | none | unassigned | -| `MODEL-EMBED-llama-llama-for-causal-lm` | `LlamaModel`, `CwmForCausalLM`, `InternLM3ForCausalLM`, `IQuestCoderForCausalLM`, `LlamaForCausalLM`, `LLaMAForCausalLM`, `TeleChat3ForCausalLM`, `MistralModel` | `registry.py:224-231`; `vllm/model_executor/models/llama.py::LlamaForCausalLM` | embedding / text | encoder attention; pooler; sliding-window attention | [embeddings-one-surface](specs/embeddings-one-surface.md) | `ACTIVE` | ARCH-ONE-SURFACE ROW 6 (2026-08-08, in flight on `row/EMBEDDINGS-ONE-SURFACE` PR #137): the `LlamaModel` membership is REGISTERED + LIVE (`as_embedding_model` mirror, adapters.py:230 — `is_pooling_model=true`, bare-prefix loader, pooling forward `Qwen3DenseModel::ForwardHidden`, engine-step `pool_tokens`, `vllm_embed` ABI v15 + `/v1/embeddings`); registration `src/vllm/model_executor/models/llama_embedding_registry.cpp:132`; fold gate `tests/vllm/models/test_llama_embedding_fold.cpp:206` 4/4-231 on the committed synthetic fixture. RESIDUALS: the other 7 memberships (incl. `MistralModel`) unregistered; REAL-checkpoint (e5-mistral class) + `LLM(task="embed")` oracle cosine gate not run (synthetic-fixture arm only; no cosine-vs-oracle number fabricated) | `CLAIM-EMBEDDINGS-ONE-SURFACE` | +| `MODEL-EMBED-llama-llama-for-causal-lm` | `LlamaModel`, `CwmForCausalLM`, `InternLM3ForCausalLM`, `IQuestCoderForCausalLM`, `LlamaForCausalLM`, `LLaMAForCausalLM`, `TeleChat3ForCausalLM`, `MistralModel` | `registry.py:224-231`; `vllm/model_executor/models/llama.py::LlamaForCausalLM` | embedding / text | encoder attention; pooler; sliding-window attention | [embeddings-one-surface](specs/embeddings-one-surface.md) | `PARTIAL` | ARCH-ONE-SURFACE ROW 6 LANDED in PR #137 (2026-08-08): the `LlamaModel` membership is REGISTERED + LIVE (`as_embedding_model` mirror, adapters.py:230 — `is_pooling_model=true`, bare-prefix loader, pooling forward `Qwen3DenseModel::ForwardHidden`, engine-step `pool_tokens`, `vllm_embed` ABI v15 + `/v1/embeddings`); registration `src/vllm/model_executor/models/llama_embedding_registry.cpp:132`; fold gate `tests/vllm/models/test_llama_embedding_fold.cpp:206` 4/4-231 on the committed synthetic fixture. RESIDUALS: the other 7 memberships (incl. `MistralModel`) unregistered; REAL-checkpoint (e5-mistral class) + `LLM(task="embed")` oracle cosine gate not run (synthetic-fixture arm only; no cosine-vs-oracle number fabricated) | `CLAIM-EMBEDDINGS-ONE-SURFACE` | | `MODEL-EMBED-modernbert-modern-bert-model` | `ModernBertModel` | `registry.py:232`; `vllm/model_executor/models/modernbert.py::ModernBertModel` | embedding / text | encoder attention; pooler; sliding-window attention | ☐ required | `INVENTORIED` | none | unassigned | | `MODEL-EMBED-bert-with-rope-nomic-bert-model` | `NomicBertModel` | `registry.py:233`; `vllm/model_executor/models/bert_with_rope.py::NomicBertModel` | embedding / text | encoder attention; pooler; FusedMoE/grouped GEMM | ☐ required | `INVENTORIED` | none | unassigned | | `MODEL-EMBED-phi3-phi3-for-causal-lm` | `Phi3ForCausalLM` | `registry.py:234`; `vllm/model_executor/models/phi3.py::Phi3ForCausalLM` | embedding / text | encoder attention; pooler | ☐ required | `INVENTORIED` | none | unassigned | diff --git a/.agents/parity-ledger.md b/.agents/parity-ledger.md index 282396c4e..767e70923 100644 --- a/.agents/parity-ledger.md +++ b/.agents/parity-ledger.md @@ -928,11 +928,7 @@ Columns: | 2026-08-08 (`row/ARCH-ONE-SURFACE`; H3 device-dispatch mutation-gate follow-up; lifecycle unchanged) | **Pins the single-queue and device-provenance contract through the real `MiniMaxH3VideoEngine::Load` path.** Adds a read-only engine `device()` query and a counting CUDA-backend fixture whose queue reports a distinctive CPU:7 device. | No vLLM-Omni behavior changes; this is test observability for the vllm.cpp backend seam above. | Reviewer mutation reproduced first: replacing `im.device = stream_queue.device` with a second `CreateQueue().device` remained GREEN at 5 cases / 135 assertions. RED-first test then reported `2 == 1`; restoring reuse is GREEN at 6/137. Independent mutation deleting queue-device provenance fails the CPU:7 assertion. DSR remains 32 and all 25 leakage-checker mutations pass. No GPU, download, generation, or performance claim. | | 2026-08-08 (`row/ARCH-ONE-SURFACE-DEVICE-LEAKAGE-V2`; `ARCH-ONE-SURFACE` ROW 8 follow-up; PR #139; CPU-only; lifecycle unchanged) | **Removes PR #136's seven shared CUDA literals without changing ABI-v14 device selection.** Wire 0/1/2 and public `auto`/`cpu`/`cuda` stay exact; internal slot 2 becomes a named-platform tag, `FindPlatformByName` resolves the canonical name, and shared loading propagates the registered `DeviceType`. Explicit CPU, absent-CUDA-before-I/O and CAPI slot 2 are preserved; H3 is untouched. | vLLM `config/device.py:13,61-66` @ `555967922`: supported device name assigned verbatim and never silently substituted. The registry lookup is the in-tree Platform/Backend portability seam; no upstream behavior delta. | RED: inherited DSR 39 (`kcuda=7`) and focused compile failures for the new enum/API/signature. Mutation pin: resolver input `kXPU` must return `kXPU`. GREEN: DSR 32 (`kcuda=0`) with baseline/allowlist unchanged; checker suite 25/25; CPU Release `-Werror` `test_platform` 11/11·85, `test_loaded_engine_dense` 9/9·65, `test_capi` 45/45·428. No CUDA runtime or performance claim; GPU A/B remains residual. | | 2026-08-06 (`row/BACKEND-CPU`; `BACKEND-CPU` R1; PR #65; lifecycle remains `PARTIAL`) | Adds `vllm-cpu-kernel-bench`, a developer-only vt-op benchmark substrate: deterministic quant-GEMM fixtures, calibrated batched timing, cache-pressure profiles, affinity, JSON, checksums, system metadata, and grouped generic/Cortex-A76 `perf_event_open` counters with explicit multiplex/unsupported status. Production dispatch and numerics are unchanged. | No vLLM behavior counterpart; vLLM remains the x86 semantic oracle and llama.cpp `237ad9b96` remains the Pi performance floor. The quant fixture invokes the existing `vt::MatmulBTQuant` contract unchanged. | **CPU-GATED, `benchmark_binding=false`.** GCC 15.2 `-Wall -Wextra -Werror` build; clang-format clean; `test_cpu_kernel_bench_cli` deterministic JSON schema/checksum + invalid-input + structured-counter cases; direct 1/4-thread x86 runs and real generic PMU counts. X86 timings are non-binding tool validation. Pi PMU execution, model correctness, throughput and memory all remain `PENDING`. | -<<<<<<< ours | 2026-08-08 (`row/SPIKE-TENSOR-PARALLELISM`; `CLAIM-TP-SPIKE-287`; task #287; records-only, lifecycle unchanged) | **Tensor-parallelism end-to-end scope spike at the CURRENT pin `555967922`** — `.agents/specs/tensor-parallelism-spike.md`: S1 at-pin inventory of everything TP touches (GroupCoordinator/linear-layer sharding/vocab+logits/weight-loader shard_id/head-split+KV-per-rank/MoE-TP-vs-EP/multiproc-executor/sampler) with per-item our-seam verdicts (~40% landed/reusable); the W1/W2 landed-vs-claimed audit (the tp handle dead-ends at the layer boundary; no production loader passes `tp`; the NCCL TU has never compiled); S2 decisions (thread-per-rank recorded deviation, additive TP>1 Forward branch with TP=1 byte-identical, `tensor_parallel_size` on `vllm_model_params` at the next ABI bump, per-weight-class shard map); S3 TP2-on-CPU token-exact gate design + the honest bar (upstream's own TP=2 test compares vs HF, not vs TP1 — `tests/basic_correctness/test_basic_correctness.py:204`; `use_all_gather()` defaults True at the pin ⇒ every rank samples full-vocab identically); S4 TP-W0..W7 ranked plan, TP-W1..W4+W7 CPU-completable. RIDER (USER 2026-08-08): DSpark speculator re-grounded at the pin (`specs/dspark-speculator-note.md`; `SPEC-DSPARK` stays INVENTORIED). Verified against upstream by reading the pinned tree directly (SHA-checked); no numerics claim, `benchmark_binding=false`. | -======= -| 2026-08-06 (`KERNEL-CPU-A76-Q8-DOT` R4-R5; `CLAIM-KERNEL-CPU-A76-Q8-DOT`; physical RPi5 Cortex-A76; closing commit: this checkpoint) | Adds an exact-order ACLE SDOT control and an original AAPCS64 two-block Q8_0×Q8_0 leaf behind Linux DotProd/MIDR dispatch. `auto` selects assembly only on Cortex-A76+DotProd; x86, non-DotProd and other Arm CPUs retain portable dispatch. Explicit `portable`/`sdot`/`a76-asm` same-binary controls remain. ARM64 builds/tests locally through buildx/QEMU; Pi is execution-only. | The integer structure is informed by llama.cpp `237ad9b96` `ggml/src/ggml-cpu/arch/arm/quants.c:1076-1160`, while this port deliberately retains the local portable function's per-block f32 reduction order. vLLM `555967922` supplies Qwen3.5 semantics, not a corresponding CPU microkernel. Local anchors: `src/vt/cpu/cpu_quant_dot_{sdot.cpp,a76.S}`, Q8 dispatch in `cpu_quant_dot.cpp`, direct tests in `tests/vt/test_ops_quant_dot.cpp`, and [immutable evidence](../docs/bench-evidence/rpi5-a76-q8-dot-20260806.md). | **PASS for the compiler-gap/component gate; row `GATING`, `benchmark_binding=true`.** Final binaries `9eb57cf...`/`a94dad30...`; QEMU focused suite 20/20, 150258 assertions; physical-Pi checksums exact. Assembly vs compiler SDOT wall/cycles/instructions: M1/T1 +3.66%/+3.17%/+10.10%, M128/T1 +5.08%/+4.61%/+10.24%, M128/T4 +3.69%/+3.69%/+9.74%. M1/T4 is an explicit −2.43% wall/−4.32% cycles residual despite 8.77% fewer instructions. All 64 Qwen tokens equal the x86 golden in all nine runs; median assembly vs SDOT TTFT −1.55%, TPOT −0.05% neutral, E2E −0.13%. Disassembly proves GCC's framed dependent one-block loop versus the stack-free independent two-block schedule. Same-file Pi llama.cpp, peak memory and concurrency remain `PENDING`; no competitor-floor binding is claimed. | -| 2026-08-06 (`KERNEL-CPU-A76-Q8-DOT` R6 competitor checkpoint; physical RPi5 Cortex-A76; lifecycle remains `GATING`; closing commit: this checkpoint) | Measures the separate four-core A76 same-file llama.cpp floor after the assembly leaf became default. No production code changes. The vllm.cpp nominal p16 request measures 17 input tokens, so the binding competitor uses pp17/tg64/pp17+tg64. A same-text CLI arm verifies 64-token greedy output equality. | Official llama.cpp tag b9892 `ee445f93d` reconstructed under QEMU because historical recorded fork object `237ad9b96` is unavailable; exact recorded anchors match (`quants.c:400`, `arch/arm/quants.c:1076`, `repack.cpp:2725`, `qwen35.cpp`). Local evidence: [Pi competitor record](../docs/bench-evidence/rpi5-a76-llamacpp-20260806.md). | **CORRECTNESS PASS, PERFORMANCE NOT MET, `benchmark_binding=true`.** Three clean unthrottled vllm.cpp reps: prefill 12.81 tok/s, decode 2.55 tok/s, output-equivalent E2E 2.46 tok/s, E2E 26,018.39 ms. llama.cpp three-sample p17/tg64/combined: 27.77 / 3.91 / 3.77 tok/s, E2E 16,998.49 ms. vllm.cpp ratios 0.461x prefill / 0.653x decode+E2E; peak RSS wins 2.841 vs 3.747 GiB (24.2% less). Same-text normalized output SHA `a5a630d7...` equal; all vllm performance tokens retain golden SHA `0ec98e...`. Intrusive 50 ms forked sampler run VOID; accepted timing has no sampler, RSS sampled separately at 1 Hz. Next lever: fresh both-engine profile, then BF16 GEMM; M1/T4 and concurrency remain. | ->>>>>>> theirs | 2026-08-06 (`KERNEL-CPU-A76-Q8-DOT` R4-R5; `CLAIM-KERNEL-CPU-A76-Q8-DOT`; physical RPi5 Cortex-A76; closing commit: this checkpoint) | Adds an exact-order ACLE SDOT control and an original AAPCS64 two-block Q8_0×Q8_0 leaf behind Linux DotProd/MIDR dispatch. `auto` selects assembly only on Cortex-A76+DotProd; x86, non-DotProd and other Arm CPUs retain portable dispatch. Explicit `portable`/`sdot`/`a76-asm` same-binary controls remain. ARM64 builds/tests locally through buildx/QEMU; Pi is execution-only. | The integer structure is informed by llama.cpp `237ad9b96` `ggml/src/ggml-cpu/arch/arm/quants.c:1076-1160`, while this port deliberately retains the local portable function's per-block f32 reduction order. vLLM `555967922` supplies Qwen3.5 semantics, not a corresponding CPU microkernel. Local anchors: `src/vt/cpu/cpu_quant_dot_{sdot.cpp,a76.S}`, Q8 dispatch in `cpu_quant_dot.cpp`, direct tests in `tests/vt/test_ops_quant_dot.cpp`, and [immutable evidence](../docs/bench-evidence/rpi5-a76-q8-dot-20260806.md). | **PASS for the compiler-gap/component gate; row `GATING`, `benchmark_binding=true`.** Final binaries `9eb57cf...`/`a94dad30...`; QEMU focused suite 20/20, 150258 assertions; physical-Pi checksums exact. Assembly vs compiler SDOT wall/cycles/instructions: M1/T1 +3.66%/+3.17%/+10.10%, M128/T1 +5.08%/+4.61%/+10.24%, M128/T4 +3.69%/+3.69%/+9.74%. M1/T4 is an explicit −2.43% wall/−4.32% cycles residual despite 8.77% fewer instructions. All 64 Qwen tokens equal the x86 golden in all nine runs; median assembly vs SDOT TTFT −1.55%, TPOT −0.05% neutral, E2E −0.13%. Disassembly proves GCC's framed dependent one-block loop versus the stack-free independent two-block schedule. Same-file Pi llama.cpp, peak memory and concurrency remain `PENDING`; no competitor-floor binding is claimed. | | 2026-08-06 (`KERNEL-CPU-A76-Q8-DOT` R6 competitor checkpoint; physical RPi5 Cortex-A76; lifecycle remains `GATING`; closing commit: this checkpoint) | Measures the separate four-core A76 same-file llama.cpp floor after the assembly leaf became default. No production code changes. The vllm.cpp nominal p16 request measures 17 input tokens, so the binding competitor uses pp17/tg64/pp17+tg64. A same-text CLI arm verifies 64-token greedy output equality. | Official llama.cpp tag b9892 `ee445f93d` reconstructed under QEMU because historical recorded fork object `237ad9b96` is unavailable; exact recorded anchors match (`quants.c:400`, `arch/arm/quants.c:1076`, `repack.cpp:2725`, `qwen35.cpp`). Local evidence: [Pi competitor record](../docs/bench-evidence/rpi5-a76-llamacpp-20260806.md). | **CORRECTNESS PASS, PERFORMANCE NOT MET, `benchmark_binding=true`.** Three clean unthrottled vllm.cpp reps: prefill 12.81 tok/s, decode 2.55 tok/s, output-equivalent E2E 2.46 tok/s, E2E 26,018.39 ms. llama.cpp three-sample p17/tg64/combined: 27.77 / 3.91 / 3.77 tok/s, E2E 16,998.49 ms. vllm.cpp ratios 0.461x prefill / 0.653x decode+E2E; peak RSS wins 2.841 vs 3.747 GiB (24.2% less). Same-text normalized output SHA `a5a630d7...` equal; all vllm performance tokens retain golden SHA `0ec98e...`. Intrusive 50 ms forked sampler run VOID; accepted timing has no sampler, RSS sampled separately at 1 Hz. Next lever: fresh both-engine profile, then BF16 GEMM; M1/T4 and concurrency remain. | +| 2026-08-07 (`SERVE-CLI-BENCH` + `KERNEL-SSM-MAMBA`; clean sm_120 exact-chunk transplant; local-4B binding only) | Makes the benchmark use production `AsyncLLM`, then ports vLLM's exact `(sequence, BLOCK_M=8 token chunk)` descriptors into shared GDN step metadata and maps one CUDA register-kernel program to each descriptor. `VT_CONV_EXACT_CHUNKS` defaults ON with a same-binary `=0` rollback; `VT_CONV_REG=0` retains tiled/scalar. | Pinned vLLM engine-core queued dispatch `vllm/v1/engine/core.py:200-231,622-669` and causal-conv descriptor mapping `causal_conv1d.py:15-28,71-79,123-124`; local anchors `examples/bench/bench_core.h`, `gdn_attn.{h,cpp}`, `qwen3_5.cpp`, `ops.{h,cpp}`, `cuda_gdn.cu`; [spike/result](specs/sm120-qwen35-conv-chunking-2026-08-07.md). | **ACCEPTED + REBASED-MAIN REPRODUCED.** Contained rebuild; CPU 6/6, CUDA GDN 66/66·4300, cached 4B 3/3·1672; exact/rollback token files identical. On `3d2581551` over `upstream/main` `48a54141f`, same-binary `nsys` reproduces conv **718.704→233.955 ms = 3.072x** and profiled total **6589.65→6739.34 tok/s = +2.272%**; vLLM 145.421 ms leaves **1.609x**. Binding three-pair A/B remains total/output **+2.152%**, TTFT **-2.945%**, TPOT **-1.920%**, E2E **-2.118%**, VRAM unchanged. Sealed-vLLM throughput **1.021246x PASS**; latency/VRAM OPEN. No gate-model extrapolation. [Evidence](../docs/bench-evidence/qwen35-4b-sm120-main-20260807.md). | diff --git a/.agents/roadmap_v1.md b/.agents/roadmap_v1.md index e70e21c79..159f38eca 100644 --- a/.agents/roadmap_v1.md +++ b/.agents/roadmap_v1.md @@ -61,7 +61,7 @@ models we already ship + benchmark. Full seam map + M0–M5 W-plan: | MEM | `ROAD-V1-MEM` | **Memory budgeting: auto-size to the declared workload by default, optional total-footprint cap, pre-flight error instead of an OOM (user-directed 2026-08-06, [#83](https://github.com/mudler/vllm.cpp/issues/83))** — the user-facing wart that every operator hits before they hit any perf question: vLLM makes you compute your own VRAM budget, express it as a PERCENT, and nail it or OOM | [coverage view §2](feature-matrix.md#2-kv-cache--memory), [porting inventory](porting-inventory.md) | — **M1+M2 LANDED 2026-08-08** ([`specs/kv-sizing.md`](specs/kv-sizing.md)): the absolute `--kv-cache-memory` knob sizes the pool via a group-aware `KVBytesPerBlock` divisor, `--num-blocks` is the override, both mirrored on the C ABI at v16; `ResolveNumBlocks` precedence `num_blocks > bytes > 256`, CPU-gated (`KVBytesPerBlock` 5/5 + `test_capi` v16). M3 (the `gpu_memory_utilization` profile run) stays dgx-gated. (M0 design grounded in vLLM `config/cache.py` + `gpu_worker.py:497-599`; GB10 unified-pool caveat carried) | `M1+M2 DONE` | **Source-verified 2026-08-06 (records-only, NO code).** WE ARE CURRENTLY BEHIND vLLM ON THIS AXIS, not ahead: there is NO memory profiling at all and the KV pool is a RAW BLOCK COUNT the user types by hand — `EngineParams::num_blocks = 256` (`include/vllm/entrypoints/model_loader.h:58`, beside `block_size = 32` `:57` / `max_model_len` `:59` / `max_num_seqs = 8` `:60`), exposed verbatim as `--num-blocks N` (`examples/server/main.cpp:100,203-204,370`), carried on the C ABI as `vllm_model_params.num_blocks` at the same 256 default (`src/capi/vllm_c.cpp:429,485`), landing as `BlockPool(num_gpu_blocks, ...)` which asserts `> 0` and otherwise TRUSTS it (`include/vllm/v1/core/block_pool.h:96,223`; `src/vllm/v1/core/block_pool.cpp:51`). So a user must convert "40 GB free, 32k context, concurrency 8" into a block count themselves — strictly worse ergonomics than a percentage. UPSTREAM HAS THREE KNOBS, all `config/cache.py`, all T0, all rowed at [porting-inventory.md](porting-inventory.md) §T0: `gpu_memory_utilization` (fraction of TOTAL, default 0.9), `kv_cache_memory_bytes` (absolute KV pool) and `num_gpu_blocks_override` (exact block pin), sized off a profile run as `total x utilization - non-torch - peak activation`. **Mirroring that is necessary but does NOT solve the complaint**, for three source-grounded reasons: (1) the fraction is of TOTAL not FREE, so on any shared card the right fraction is a function of what someone else already holds — exactly the arithmetic the user is being asked to do; (2) WEIGHTS LOAD BEFORE THE KNOB ENGAGES (utilization sizes the KV pool AFTER the model is resident), so an oversized model OOMs during load and never reaches the check — which is the failure operators actually hit; (3) 0.9 is taken whether or not it is needed (a 4B model on an 80 GB card reserves 72 GB it will never touch and blocks everything else on the device). THE DESIGN (user-ratified 2026-08-06) is three modes: **Mode 1 default = SIZE TO THE DECLARED WORKLOAD** — per-allocation-class accounting BEFORE allocating anything (weights from checkpoint metadata / safetensors header / GGUF manifest, known before reading a byte of tensor data; CUDA context measured at creation; peak activation from a profile run at `max_num_batched_tokens`; KV for `max_model_len x max_num_seqs` at the resolved `block_size`/KV dtype; CUDA-graph capture-set footprint) → allocate exactly that and LEAVE THE REMAINDER OF THE DEVICE FREE. This is the SURPASS over vLLM, which takes its 90% regardless of whether the workload needs 8 GiB or 80. **Mode 2 = a cap on the TOTAL ENGINE FOOTPRINT** (weights + activations + KV + graph pools + context), NOT on the KV pool alone — load-bearing, because a KV-only cap cannot prevent the weight-load OOM; three spellings of the same cap: `--memory-limit 40GiB` (absolute, the primary form), `--gpu-memory-utilization 0.85` (vLLM's exact flag name and fraction semantics so existing vLLM launch lines port unchanged, per [[mirror-vllm-always-no-asking]]) and `--num-gpu-blocks-override N` (upstream's reproducibility escape hatch — this is where today's `--num-blocks` GOES, demoted from primary knob to explicit override), with precedence spelled out and TESTED, not left to argument order. **Mode 3 = REFUSE BEFORE ALLOCATING** with the full per-class breakdown and remedies COMPUTED FROM THE ACTUAL BUDGET (`--max-model-len N` / `--max-num-seqs M` / `--kv-cache-dtype fp8` / smaller quant, each with the GiB it recovers) — "you are 43.9 GiB over and here are the three levers that close it" is the difference between a usable error and a stack trace. **UNIFIED-MEMORY HAZARD (not hypothetical):** on GB10 the ~119 GiB pool is UNIFIED, a fraction-of-total setting reserves HOST RAM too, and `gpu_memory_utilization=0.85` has HARD-REBOOTED our DGX three separate times ([[gb10-unified-memory-oom-reboots-box]]) — hence absolute bytes is the PRIMARY form with the percentage kept only for vLLM compatibility, and hence the accounting must know whether the pool is unified, which makes free/total + an is-unified predicate a PLATFORM-SEAM question belonging behind `ROAD-V1-C1`'s abstraction (note `Platform::needs_weight_staging()` was deliberately NOT `is_unified_memory()` because the latter FLIPS GB10 — the distinction matters here) rather than a CUDA-specific branch, since discrete and unified devices need different safety margins. CORRECTNESS: pool size changes preemption/scheduling TIMING but not emitted tokens, so the SACRED token-exact gates are unaffected — and M2's gate makes that explicit by re-running them with no block flag at all. **Next gate = M0 the `specs/kv-sizing.md` spike (accounting model + precedence rules + upstream `file:line`); then M1 a `MemoryBudget` computing required bytes per class WITHOUT allocating plus the Platform free/total + is-unified seam, unit-gated predicted-vs-actual weight bytes; M2 auto-sizing as the default with `--num-blocks` demoted to `--num-gpu-blocks-override`, gated by every existing model gate running with NO block flag and staying token-exact; M3 the three caps + precedence through the server flags and the C ABI, gated by our KV pool matching vLLM's own at a matched `--gpu-memory-utilization`; M4 the pre-flight refusal, gated by a deliberately over-subscribed config exiting cleanly on GB10 (non-zero exit, no OOM, NO BOX REBOOT) covering the unified-pool path specifically; M5 (optional) a runtime guard failing the REQUEST rather than the engine.** Docs (README, [STATUS](../docs/STATUS.md)) update in the SAME change as whichever milestone shifts externally-visible behaviour, per [[keep-readme-current]] | | 1 | `ROAD-V1-C1` | **Extensibility-first (USER PRIORITY 2026-07-18):** drop-in kernel ABI + the MISSING Platform seam + model self-registration — make new GPUs/archs/models ADDITIVE (plan: [extensibility-platform-seam-2026-07-18.md](specs/extensibility-platform-seam-2026-07-18.md)) | [`BACKEND-ABI-VT`](backend-matrix.md), [kernel matrix](kernel-matrix.md) | exhaustive kernel/dependency inventory and [raw-pointer adapter ABI](specs/dropin-kernel-abi.md) accepted; additive W0 implemented and CPU 94/94. `CLAIM-BACKEND-ABI-W0-GPU-1` repaired the GCC13/doctest blocker without runtime changes; exact sm_121a all-target build, focused CUDA/ABI sanitizer, and both gate-model tests pass at `1141b79`. Cross-arch/trace/A-B and scalar-forwarder/backend-shim debts remain explicit | `PARTIAL` | **★ NEW ORDER-1 HEAD (user-directed 2026-07-19): the PORTABLE AUTOMATIC OP-FUSION FRAMEWORK (`KERNEL-FUSION-FRAMEWORK`, spike [portable-fusion-framework.md](specs/portable-fusion-framework.md), `SPIKE`).** The extensibility cornerstone: fusions DECLARED ONCE (backend-agnostic `constexpr FusedRecipe` catalog above `vt::`, transcribing vLLM's finite pattern-pass set `passes/fusion/*` @ `pass_manager.py:138-200`, mirroring the `CustomOp` `forward_native`/`forward_cuda` seam `custom_op.py:103`) and REALIZED PER-BACKEND through the existing `vt::` op table (Tier-0 composite = the CPU oracle inherited free by any backend; Tier-1 interpreter = one kernel port per backend lights up every recipe). Makes a new vLLM fusion PR a ONE-declaration port, a new GPU an additive catalog realization, a new model an additive pattern declaration — the PR-#4 remedy, composed with the Platform/attn-registry/model-registry seams below. The TDR Phase-0 skeleton is already LANDED (`fused_recipe.h`/`recipes.h` one recipe + `OpId::kFusedChain` Tier-0/1 on CPU+CUDA + byte-exact `test_ops_fused_chain.cpp`). **W0 ADOPTED 2026-07-19 (`CLAIM-FUSION-FRAMEWORK-W0`):** the seam is now used in production at ONE real site — the 35B `RunLayerPaged` post-attention layernorm routes its plain add+residual+gemma-RMSNorm through `vt::FusedChain(kFusedAddRmsNorm)` (`VT_FUSED_CHAIN_ADOPT` default-ON / `=0` rollback), behaviour-preserving + byte-identical to the prior hand-call (DGX: clean CUDA `-Werror` 0 warn, byte-exact composite==interp==golden incl. H=2048, 35B 315/315 + 27B 235/235 token-exact BOTH arms, memcheck 0 errors) — proving the declare-once/realize-per-backend seam end-to-end; the current 3-opcode/4-role POD sufficed byte-identically for the plain add+rmsnorm class, so W0 needed NO generalization. **W1 GENERALIZED the POD 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W1`, `1115648`):** full activation/norm/quant/rope opcode set + indexed operand table; all 5 quant-fused chains declared byte-exact; Tier-0 composite = ONE device-agnostic walker (kills CPU/CUDA oracle drift); infrastructure only, no call site changed (DGX: `-Werror` 0-warn, byte-exact CPU 196 + CUDA 361, memcheck 0, 27B 235/235 + 35B 315/315 both arms). **W2 MIGRATED the hand-fusions 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W2`):** the framework now OWNS the fusion dispatch — each recipe binds (new backend-agnostic `FusedRecipe.fast_op`) to its EXISTING single-launch fused kernel, so `FusedChain(recipe)` dispatches to the SAME fast kernel the model called directly pre-migration (byte-identical + perf-neutral by construction; composite is the graceful fallback + oracle). SIX call sites migrated behind `VT_FUSED_CHAIN_ADOPT` (`kSiluMulFp4Quant`/`kSigmoidGateFp4Quant`/`kRmsNormGatedQuantFp8`×2/`kRmsNormQuantFp8`/`kAttnQkNormRopeGate`×2). DGX: `-Werror` 0-warn, byte-exact CPU 228 + CUDA 420, memcheck 0, 27B 235/235 + 35B 315/315 BOTH arms. **W3 MECHANICAL-SYNC PROOF LANDED 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W3`):** ported a NEW, previously-unported vLLM fusion pass — `SiluMulFp8StaticQuantPattern` (`act_quant_fusion.py:81` → `_C.silu_and_mul_quant`, the static-per-tensor-FP8 sibling of `kSiluMulFp4Quant`) — as ONE `constexpr FusedRecipe kSiluMulQuantFp8` + its byte-exact test, touching EXACTLY 2 shared files (`recipes.h` + `test_ops_fused_chain.cpp`), NO kernel/dispatch/composite-walker/model-site edit and NO new primitive (composite = existing `vt::MoeSiluMul` + `vt::QuantFp8Static`; `fast_op=kNoFastOp`). The PR-#4 additivity test made concrete: a whole new fusion pattern = one declaration. DGX: `-Werror` 0-warn, byte-exact CUDA 432, memcheck 0, no token regression (recipe declared-only) 27B 235/235 + 35B 315/315. **W4 BACKEND-ADDITIVITY PROOF LANDED 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W4`) — the W-series proof milestone is DONE:** made the additivity claim EXECUTABLE — new test `test_fused_chain_additivity.cpp` treats the EXISTING CPU backend AS the 'second backend' relative to CUDA (no mock `DeviceType` — that would edit the core enum + every switch, ironically non-additive) and, in ONE generic loop over the WHOLE catalog (all 7 recipes), asserts each runs byte-exact on the CPU backend via the Tier-0 composite — 4 CPU-full end-to-end + 3 fp8-prefix (byte-exact prefix + the FULL composite asserted to THROW on CPU, documenting the CUDA-only static-fp8 backend-negotiated tail, §3b/§6). Additivity evidence: catalog `recipes.h` grew 1→6→7 while the composite walker stayed ONE per-OPCODE function (12 `FOp::` cases) + the CPU/CUDA `kFusedChain` registration ONE line each + `cpu_ops.cpp` never `#include`s `recipes.h` — W3's whole new recipe `kSiluMulQuantFp8` is in ZERO backend TUs, inherited free. CPU `-Werror` 0-warn, `test_fused_chain_additivity` 17/17 + `test_ops_fused_chain` 228/228; engine byte-identical (no `src/`/`include/` change) ⇒ 27B 235/235 + 35B 315/315 structurally unchanged; memcheck N/A. Honest deferred (named, non-blocking the ORDER-1 milestone): Tier-1 perf interpreter for the quant chains (composite-only today), a REAL Metal/Vulkan catalog realization (M4 HW-blocked), and per-recipe fast single-launch kernels. Honest payoff: perf ceiling ~3.5%/step compute-bound on 35B (NOT a perf lever — tasks #61/#62; W0 is perf-neutral by construction); primary value = extensibility + mechanical upstream-sync + CPU/CUDA oracle-drift elimination. Incremental W0 adopt-one **(DONE)** → W1 generalize POD **(DONE)** → W2 migrate hand-fusions **(DONE)** → W3 mechanical-sync proof **(DONE)** → W4 backend-additivity proof **(DONE)** → Wn honest re-measure (optional, off the extensibility critical path). **W-SERIES ORDER-1 PROOF MILESTONE DONE 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W4`).** **PRIOR extensibility items (all LANDED, the seams this composes with):** **#1 extensibility item — extract the Platform seam — LANDED 2026-07-18** (`BACKEND-PLATFORM` `ACTIVE`, `CLAIM-BACKEND-PLATFORM-1`): `include/vllm/platforms/interface.h` + `src/vllm/platforms/{platform,cpu,cuda}.cpp` mirror `vllm/platforms/interface.py:134-229` 1:1; `CurrentPlatform()` self-registered per `DeviceType`; the 7 memory-model/residency `device.type == kCUDA` sites (of PR #4's ~37) now route through it → new-GPU memory model is ONE additive `platforms/.cpp`. Behavior-preserving (clean CPU build + `test_platform` + full CPU CTest green; DGX 235/235 + 315/315 pending). **Item 2 residency-as-Platform-capability LANDED 2026-07-19** (`CLAIM-BACKEND-PLATFORM-2`): the host-free / load-stream / DevicePool-cap decisions in `qwen3_5.cpp` now READ `GetPlatform(.device.type).residency_policy()` (per-device) via the pure `ShouldReleaseHostWeights`/`ShouldInterleaveLoadStream` helpers + `device_pool_cap_bytes`, not an inline `device.type`/env gate; `CudaPlatform.release_host_weights_after_upload` flipped false→true (now CONSUMED ⇒ reproduces today's GB10 host-free-after-Marlin + ~4 GiB load peak EXACTLY); `MarlinMoeEnabled()` stays the orthogonal kernel-path gate. **A new (discrete) GPU sets `residency_policy()` values ⇒ ZERO model edit.** Behavior-preserving (clean CPU build + `test_platform` consumption cases 7/43 + full CPU CTest + tools 164/164 green; **DGX-CONFIRMED @ `62fc0e0`: clean CUDA `-Werror`, 27B 235/235 + 35B 315/315 token-exact, 35B VmHWM ≈ 4.0 GiB load-stream win preserved, memcheck 0 errors**). Then item 3 drop-in ABI family migration. **Item 4 attn-backend registry LANDED 2026-07-19** (`CLAIM-ATTN-REGISTRY-1`, `BACKEND-ATTN-REGISTRY`): NEW `include/vllm/v1/attention/registry.{h,cpp}` `(DeviceType,name)` registry + `SelectAttentionBackendName` selector (mirror `registry.py` self-registration + `cuda.py:361-470` `get_attn_backend_cls`/`_get_backend_priorities`); `Platform::get_attn_backend_priority()` filled (was the item-1 STUB) → capability-ordered name lists on `CudaPlatform` (major-10 vs else) + `CpuPlatform`; FLASH_ATTN/GDN self-register. **Adding a backend's attention = 1 self-registering TU + 1 priority slot, ZERO selector/model/runner edit.** Behavior-preserving — the walk returns FLASH_ATTN on CUDA+CPU (the same FA2 attention runs); clean CPU build + `test_attn_backend_registry` (8/25) + full CPU CTest, tools 164/164, checkers green; **DGX-CONFIRMED @ `2c732e7`: 27B 235/235 + 35B 315/315 token-exact (FA2 sm_121a), memcheck 0/315**. **Item 5 model self-registration LANDED 2026-07-19** (`CLAIM-MODEL-SELFREG-1`, `MODEL-FACTORY-registry`): the fixed `kRegistrations` array → `REGISTER_VLLM_MODEL(...)` static-`Registrar` idiom (`model_registry.h:167-189`) + Qwen dense/MoE arch entry points split into per-variant TUs (`qwen3_5_dense.cpp`/`qwen3_5_moe.cpp`) over shared `qwen3_5_common.{h,cpp}`, so **adding a model = 1 new TU + 1 REGISTER line, ZERO shared-array edit**; behavior-preserving (clean CPU build + `test_model_registry` extension + full CPU CTest, tools 164/164, checkers green; DGX 27B/35B token-exact pending). Deep `qwen3_5.cpp` machinery factoring deferred. Metal/MLX bring-up proves the seams (needs M4). **★ THE ARCH HALF OF THIS ITEM IS NOW PROVEN BY MEASUREMENT, NOT ARGUED (2026-07-22, `CLAIM-CUDA-SM120-BRINGUP`, [spec §W8](specs/cuda-arch-additivity.md)):** consumer-Blackwell `sm_120a` was brought up as a BUILD-supported target and required **ZERO kernel, model, runner, sampler or feature-table edits** — the additive seams (per-arch FEATURE TABLE, capability-keyed tactic registry keyed on `major == 12`, Platform auto-probe, `pageable && integrated` residency classification) already covered it, so the entire diff is build configuration, a configure-tier test and records. That is the PR-#4 additivity test passed on a real second architecture. It is deliberately NOT a runtime-support claim: no sm_120 board exists here. **★ THE MODEL/QUANT HALF NOW ADVANCES TOO (2026-07-23, `CLAIM-BACKEND-SEAM-S4-1`): the `model_executor/layers/` `LinearMethod`/`QuantizationConfig` seam the [accelerator-seam audit](specs/accelerator-seam-audit.md) §9 named ABSENT now EXISTS in part.** `S4` landed byte-identical: the dense model's projections route through a `method.Apply()` chosen ONCE from the checkpoint (retiring the per-call `IsNvfp4()` tensor-name probe), and 18 shared-layer `device==kCUDA` availability gates became `vt::OpRegistered` op-table queries — the policy(scheme)/implementation(kernel) split the audit's binding rule requires. **DSR 86 → 67**; all six SACRED gates byte-identical on dgx (27B/35B/Coder/dense/OPT/DeepSeek-V2); the fragile 27B-W4A4/fp8-recipe gates are correctly deferred to `S6` behind `S5`'s reference tier. **★ `S6` ASSESSED 2026-07-23 (`CLAIM-BACKEND-SEAM-S6-1`) → NO-OP / BLOCKED, DSR stays 67 (§11):** the deferred fp4/fp8 gates convert ZERO sites byte-identically — every one bottoms out at a **dual-registered** (CPU+CUDA) bespoke op (none CUDA-only, unlike S4's convertible gates), so `OpRegistered(op,dev)` is TRUE on `kCPU` ⇒ the class-A swap is bit-changing on the CPU reference/emulation path (two numerics per device); S5's reference tier does not change this (those CPU kernels are present natively, never a miss). No `src/`/`include/`/test byte changed, no baseline moved. The genuine byte-identical unlock is re-scoped to **`S3`** (Platform capability fields mirroring `supports_fp8`/`cutlass_fp4_supported` — the audit's own class-D fix) and **`S7`** (layer extraction); the plan's `~37` S6 target assumed the class-A `OpRegistered` swap was byte-identical, which holds only for CUDA-only ops (all taken by S4). **★ `S3` LANDED 2026-07-23 (`CLAIM-BACKEND-SEAM-S3-1`) — the byte-identical unlock S6 re-scoped:** mirrors vLLM's `Platform` capability surface (`supports_fp8`/`cutlass_fp4_supported`/`opaque_attention_op`/`is_integrated_gpu`/`support_static_graph_mode`/`is_device_capability_family`, base false in `interface.h`; CudaPlatform answers GB10 values in `cuda.cpp`/`platform.cpp`) and converts **12** deferred `qwen3_5.cpp` gates onto it (7 fp4-act `cutlass_fp4_supported`, 3 fp8-fused `supports_fp8`, 2 decode-graph `support_static_graph_mode`) — byte-identical because a capability answers the base false off CUDA, exactly what `device==kCUDA` did (where S6's `OpRegistered` was TRUE on `kCPU`), and it DECOUPLES (a future accelerator answers for itself). **DSR 67 → 55** (`kcuda` 25→13), baseline lowered same commit, ratchet + 24-case suite green; all seven SACRED gates byte-identical on dgx (27B 235/235 · 35B 315/315 · Coder · dense-32B · OPT · DeepSeek-V2 · Llama), `test_platform` CUDA-leg proves each predicate == former `device==kCUDA` on GB10, memcheck 0 errors, clean CUDA+CPU `-Werror`. Residency/stream/FA2-dtype/merged-layout sites LEFT for `S7`. **★ `S7` LANDED 2026-07-23 (`CLAIM-BACKEND-SEAM-S7-1`) — the seam campaign's TERMINAL runtime-decoupling state, the closest this extensibility work comes to a finish line:** ALL 23 remaining runtime `kCUDA`/`is_cuda()` sites in the shared model layer hoisted onto capabilities — new `Platform::needs_weight_staging()` (the CUDA device-resident staging policy, NOT `is_unified_memory()` which would FLIP GB10; covers residency/merged-GDN/packed-decode/direct-load), `Platform::supports_fa2_attention()` (FA2 dtype), `Backend::SupportsAuxStream()` (MoE aux-stream), reusing S3's `supports_fp8`/`cutlass_fp4_supported`/`support_static_graph_mode`/`is_integrated_gpu` (runner combine/scatter) and `vt::OpRegistered(kMoeGroupedGemmBf16)` (a CUDA-only op). Each returns the former `device==kCUDA` value on GB10 → byte-identical. **DSR 55 → 32 — the IRREDUCIBLE build-gate floor:** the shared model layer holds ZERO runtime device tests; the 32 residual are all `#ifdef VT_*` compile-time gates for kernels that only build on one GPU family (a kernel that only compiles on one arch is legitimately irreducible), so the audit's `<10` is NOT reachable and this is the honest answer to "how additive can the shared layer get" — every runtime device coupling is gone. baseline lowered same commit, ratchet + 24-case suite green; all seven SACRED gates byte-identical on dgx (27B 235/235 · 35B 315/315 · Coder 6/6 · dense-32B 16/16 · OPT 6/6 · DeepSeek-V2 8/8 · Llama 16/16), new `test_platform`/`test_backend`/`test_cuda_backend` cases green, memcheck 0, clean CUDA+CPU `-Werror`. The `layers/`-library physical relocation (shrinking `qwen3_5.cpp` toward `qwen3_next.py`'s 802-line shape) is a follow-on refactor; the device coupling it was to remove is already gone | | 2 | `ROAD-V1-C2` | Model families: Llama/Qwen3/Mistral, MoE, Qwen3-Next | [model matrix](model-matrix.md) | current pin has 353 static IDs; v0.25.0 adds three sync-target rows (MOSS-Transcribe-Diarize, Laguna DFlash, Bailing hybrid MTP), yielding 356 after pin advance. **FIRST ADDITIVE-MODEL BRING-UP W0-W4 LANDED 2026-07-20 — CORRECTNESS COMPLETE (0.6B + 4B gates PASS 16/16; SPEED pending)** ([first-additive-model-qwen3-dense.md](specs/first-additive-model-qwen3-dense.md), `MODEL-TEXT-qwen3-qwen3-for-causal-lm` `ACTIVE`(correctness COMPLETE, speed pending), runner generalization `ENG-RUNNER-MODELSHAPE`, `CLAIM-MODEL-QWEN3-DENSE`) **MLA CAMPAIGN SPIKED 2026-07-21** ([mla-deepseek-campaign](specs/mla-deepseek-campaign.md), `CLAIM-MLA-DEEPSEEK`): five rows `INVENTORIED` -> `SPIKE` (DeepSeek V2 / V3+V3.2 / v1-MHA, Kimi-Linear, MiniMax-M2). **KIMI-LINEAR-48B W0 DEDICATED SPIKE 2026-08-05** ([kimi-linear.md](specs/kimi-linear.md), `CLAIM-KIMI-LINEAR-W0`): full dedicated W0 spike for `MODEL-TEXT-kimi-linear-*` (stays `SPIKE` — actively claimed) — the ONE Kimi text model that FITS one GB10 (91.5 GiB, 0.77x pool) with a real e2e SACRED gate; HEAVY reuse (MLA + sigmoid/`noaux_tc` MoE + GDN family + KDA host refs landed), NET-NEW = the KDA device kernel + NoPE-MLA branch + hybrid schedule/loader; W1 implementation can start. Answers the Tier-3 "MLA = new attention, new campaign" item in [breadth-sweep-plan](specs/breadth-sweep-plan.md) §B.3. Key determinations: GB10/sm_121 selects **`TRITON_MLA`** for dense MLA decode and **`FLASH_ATTN`** for MLA prefill (`vllm/platforms/cuda.py:129-133`, `mla/prefill/selector.py:74-77`), so the sm90/sm100-only MLA kernels are out of reach and out of scope; the cross-cutting cost is the **compressed-latent KV cache** (one 576-wide latent per token, `num_kv_heads=1`, no separate V), which our allocator and `vt::ReshapeAndCache`/`vt::PagedAttention` cannot express; and **only DeepSeek-V2-Lite (~29.3 GiB bf16) fits GB10** — V3/V3.2, Kimi-K2.5, MiniMax-M2/M3 are HW-BLOCKED e2e, Kimi-Linear-48B is HW-MARGINAL. W0-W10 plan recorded; nothing implemented. **GLM + DSA + LATEST-DEEPSEEK SPIKED 2026-07-21** ([glm-dsa-latest-deepseek](specs/glm-dsa-latest-deepseek.md), `CLAIM-GLM-DSA-LATEST-DEEPSEEK`): seven rows `INVENTORIED` -> `SPIKE` (ChatGLM, Glm, Glm4, Glm4Moe, Glm4MoeLite, GlmMoeDsa, DeepSeek-V4). Answers the user's "also glm, and deepseek latest versions" priority. Headline: **`Glm4MoeLiteForCausalLM` / `zai-org/GLM-4.7-Flash` (31.2B, 58.2 GiB bf16) FITS GB10 and is a SECOND MLA gate vehicle that closes BOTH coverage gaps the MLA campaign named as unit-gated-only** (it has `q_lora_rank=768` and `noaux_tc`/`e_score_correction_bias`, which DeepSeek-V2-Lite lacks). **DSA is DOUBLY BLOCKED on GB10:** the sparse XOR filter eliminates `TRITON_MLA` for sparse models leaving `FLASHINFER_MLA_SPARSE_SM120` as the sole candidate, and that path is non-functional on flashinfer 0.6.12 (XQA backend is dense-only, discards `sparse_mla_top_k`); GLM-5 is 1404 GiB and V3.2 642 GiB regardless. DeepSeek-V4 is a NEW architecture (Sinkhorn-normalized Manifold Hyper-Connections, CSA/HCA compressor with recurrent state, hash-routed MoE) and HW-BLOCKED at 148.7 GiB — but its TOKENIZER risk is LOW (standard HF fast BPE; only the chat template needs porting, with upstream golden fixtures). Glm4/Glm need two primitives we have NONE of: partial rotary factor and sandwich norms. Nothing implemented. **NVFP4A16 (W4A16)** on the already-done dense `Qwen3ForCausalLM` (`RedHatAI/Qwen3-32B-NVFP4A16`, 64L) — the QUANT-SCHEME additivity experiment, serving user priorities #2 (models) and #4 (quants) at once. KERNEL LAYER FULLY ADDITIVE (ZERO new kernel code: vLLM FORCES Marlin for `use_a16`, OBSERVED `Using MarlinNvFp4LinearKernel`, and that is the GEMM we already vendored for the 35B). **CORRECTNESS CLOSED 2026-07-21 (W4b):** the strict gate's 4/6 was diagnosed by the ratified TEACHER-FORCING isolation — all 29 divergent positions gap <= 0.0625 nats with 28/29 EXACTLY 0.0, one root flip an EXACT bf16 tie at which vLLM's teacher-forced argmax is OURS and vLLM contradicts its own greedy. **NOT a W4A16 defect: it is the pre-existing dense-forward bf16 near-tie drift, recorded against `MODEL-TEXT-qwen3-qwen3-for-causal-lm`.** Gate closes **6/6** under the ratified near-tie-robust bar with the nats evidence committed. SPEED still pending ⇒ `ACTIVE`, not `DONE`. **GEMMA FAMILY SPIKED 2026-07-24** ([sweep-gemma](specs/sweep-gemma.md), `CLAIM-SWEEP-GEMMA`): four rows `INVENTORIED` → `SPIKE` (Gemma 1/2/3/4). Answers the user's "and then we do gemma" ("gemma 4") next-target. **The newest registered Gemma is Gemma 4** (real, public, but ALL checkpoints multimodal-wrapped `Gemma4*ForConditionalGeneration`, ≥12B, 0.25.0 oracle-support unverified, needs a PLE/YOCO/MoE/k_eq_v stack) — it leads the characterization but is gate-BLOCKED as a first vehicle. **The recent-first gate vehicle that FITS + is oracle-certain is Gemma 3** (`Gemma3ForCausalLM` on `google/gemma-3-1b-it`). Headline: Gemma reduces MOSTLY to landed infra — gemma-RMSNorm, sandwich norms (glm4 `b568d20`), SentencePiece (names "Gemma"), sliding-window (FA-2 + SlidingWindow/ChunkedLocalAttention specs), the `kAttnQkNormRopeGate` QK-norm+rope recipe, tied embeddings are ALL REUSE; the one genuinely-new compute kernel is GeGLU (`gelu_pytorch_tanh`+mul; we have only SiLU), plus the final logit soft-cap + qpas/embed-scale scalars + dual-rope routing. Per-version delta: Gemma-2 has an attn logit soft-cap, Gemma-3 removed it for QK-norm. **GEMMA-3 W0-W2 LANDED 2026-07-24 — CORRECTNESS COMPLETE, the FIRST Gemma family** (`MODEL-TEXT-gemma3-gemma3-for-causal-lm` `ACTIVE`, speed pending): `Gemma3ForCausalLM` on `google/gemma-3-1b-it`. W1 = two additive default-inert vt ops `kGeluAndMul` (GeGLU `gelu_pytorch_tanh`) + `kMulScalar` (bf16 embed-scale), CUDA+CPU, unit 12/12. W2 = `gemma3.{h,cpp}`/`gemma3_weights.cpp`/`gemma3_registry.cpp` reusing the GLM-4 sandwich-norm layout + `dense_attn_block.h` + FA-only KV: GemmaRMSNorm `(1+w)`, per-head Gemma q/k norm, dual per-layer RoPE theta, qpas scale, per-layer sliding window, GeGLU MLP, `sqrt(hidden)` embed-scale, tied lm_head. **SACRED gate STRICT token-exact 48/48** greedy vs vLLM 0.25.0 (K=5 ALL-DETERMINISTIC → STRICT; BOS-verified; tokenizer-free like Mistral's `LOAD-SENTENCEPIECE` path). Loader 340 tensors, registry 23/23, clean `-Werror` 0 warn. **GEMMA-2 + GEMMA-1 W3-W6 LANDED 2026-07-24 — CORRECTNESS COMPLETE** (`MODEL-TEXT-gemma2-gemma2-for-causal-lm` + `MODEL-TEXT-gemma-gemma-for-causal-lm` `ACTIVE`, speed pending): W3 = the logit soft-cap primitives (`vt::SoftCap` final cap + `PagedAttentionArgs.logits_soft_cap` attention cap threaded into the native/CPU/FA-2 attention, default-inert). W4 `Gemma2ForCausalLM` (gemma-2-2b-it) = the inverse of Gemma-3 (BOTH soft-caps, NO QK-norm, single rope) — **near-tie-band SACRED 48/48** (44/48 strict + 4/48 at 0.0-nat vLLM-own ties, 0 forward-divergent; soft-cap PROVEN applied by a cap-on≠cap-off A/B). W5 `GemmaForCausalLM` (gemma-2b) = the original Gemma (two fused norms, head_dim scale) — **STRICT 48/48**. W6 = Gemma-4 honesty pass (HW/DEP-BLOCKED, not registered). Regressions byte-identical (Gemma-3 48/48, Qwen3-dense 184/184, OPT 63/63, Llama 92/92, Mistral 92/92) + DeepSeek-V2 asserts-on 223/223; compute-sanitizer 0; clean `-Werror` 0 warn. Gemma-4 stays `BLOCKED`. | `PARTIAL` | **ACTIVE: the first additive-model bring-up = Qwen3 dense (`Qwen3ForCausalLM`) on `Qwen3-0.6B` BF16 — W0+W1 landed 2026-07-20.** W0 (config+registry stub: new TU `qwen3_dense.cpp`+`qwen3.h`, one `REGISTER_VLLM_MODEL`, full-attention-only KV spec, forward stub) + W1 (the RUNNER GENERALIZATION `ENG-RUNNER-MODELSHAPE`) are DONE and gated: dgx CUDA `-Werror` 0-warn, **27B 235/235 + 35B 315/315 token-exact UNCHANGED** (behaviour-preserving), new CPU runner tests RED(SIGSEGV)→GREEN, registry resolves `Qwen3ForCausalLM`, ASan/UBSan clean. The runner is now MODEL-SHAPE-AGNOSTIC (a full-attention-only KV config allocates+steps without the hybrid GDN path) → every future dense/non-hybrid arch adds new-files-only. Qwen3-0.6B is the only standard-dense arch with a checkpoint + runnable vLLM 0.25.0 oracle on dgx TODAY (no Llama/Mistral checkpoint present → Llama-first needs a download, sequenced as W-next for genuine cross-family additivity). **W2 loader + W3 forward LANDED 2026-07-20:** dense forward `qwen3.cpp` (`Qwen3DenseModel::Forward/ForwardDevice`) composed from vt:: ops + the fusion catalog (2 new byte-exact recipes: `kFusedAddRmsNormStd`, `kAttnQkNormRope`); bf16 attention numerics mirror vLLM. The first pure-dense bf16 model forced out + FIXED 2 genuine latent bugs: tokenizer `kQwen2Classic` (classic Qwen2/Qwen3 pre-tokenizer was hard-rejected) and `cuda_paged_attn.cu` WMMA prefill mistokenizing at head_dim≠256 (now gated to the validated d=256). **W4 CORRECTNESS COMPLETE 2026-07-20 — near-tie-robust gate PASSES on 0.6B AND a bigger 4B.** The 2026-07-20 razor's "vLLM greedy non-deterministic" premise was a BATCHING artifact: per-prompt (batch=1, the gate regime) vLLM 0.25.0 greedy is DETERMINISTIC (0.6B 0-multi/K=10, 4B 0-multi/K=5). Forward correctness is PROVEN by teacher-forcing vLLM on OUR exact prefix (`scripts/qwen3-neartie-gap.py`): at all-but-2 positions vLLM's own argmax given our prefix IS our token (gap 0.0000, bit-identical logprobs — our forward matches vLLM's prefill logits); residual flips are bf16 near-ties (0.6B ≤0.125 nats, 4B ≤0.25) where vLLM's own prefill argmax disagrees with its decode. Gate = our token within 0.5 nats of vLLM's teacher-forced argmax (strict where equal): **Qwen3-0.6B 16/16** (strict 12 + near-tie 4) and the **bigger-model complete-correctness proof Qwen3-4B (36L, GQA 32/8, hidden 2560, same forward code) 16/16** (strict 10 + near-tie 6). Regression 27B 235/235 + 35B 315/315 UNCHANGED, CUDA `-Werror` 0-warn, memcheck 0. Correctness-complete. **SPEED — d128 FA2 PREFILL + DECODE DEFAULT-ON 2026-07-20 (`Qwen3-4B` vs vLLM 0.25.0 production/graphed, in1024/out128) — big gap-close, still below vLLM, `MODEL-TEXT-qwen3-qwen3-for-causal-lm` stays `ACTIVE` NOT `DONE`:** implemented the dominant prefill lever (a d128 FlashAttention-2 varlen prefill — generalized the vendored FA2 launcher to head_dim 128, `VT_FA2_PREFILL_QWEN3` default-ON) and flipped the FA2 varlen d128 decode default ON (near-tie gate re-passes 16/16 on 0.6B + 4B). Total tput now 0.90× (c1)/0.62× (c8) (up from 0.80×/0.48×), c1 decode at parity (TPOT 1.04×, ITL P99 0.98× win); prefill A/B = +7%/+41% total, −55%/−48% TTFT. STILL failing TTFT median 5.85×/10.2× + total <1×: the full prefill STEP (not the attention kernel, now vLLM's FA2 family) is ~6× vLLM = non-attention glue (GEMM/MLP fusion) + host-side launch overhead (un-graphed prefill); plus c8 decode batch efficiency (TPOT 1.38×). Dominant residual lever = portable prefill-step fusion + graphed prefill (secondary = c8 split-KV decode occupancy). **RoPE cos/sin cache flipped DEFAULT-ON 2026-07-20** (`VT_QWEN3_ROPE_CACHE`): the opt-in blocker (an alleged FA2-split-KV-combine run-to-run nondeterminism) was GROUNDED + DISPROVEN — the paged engine is byte-deterministic run-to-run and the combine never launches on the gate (`num_splits==1`); goldens regenerated on the canonical `$HOME/cutlass-4.5.0` build (the flashinfer cutlass copy tips the 27B tok6 razor to 234/235; cutlass-4.5.0 = 235/235), gate 16/16 both, 27B 235/235 + 35B 315/315 unchanged. RoPE-ON closes total tput 0.90×→0.97× (c1) / 0.62×→0.82× (c8) and c1 TTFT ratio 5.85×→2.27×, still `ACTIVE`. **SPEED RE-BOUND 2026-07-21 (same-session, matching-recipe) — TTFT residual RESOLVED, cutlass claim CORRECTED:** the "TTFT 2.27×/5.85×" + "c8 ITL 4.3×" were BAD-DENOMINATOR/num-prompts artifacts — a fresh same-session vLLM capture gives c1 TTFT ~152 ms & c8 ITL P99 ~130 ms, and OURS WINS TTFT at both concurrencies (c1 0.90×, c8 0.38×). **c1 = effective every-axis parity** (tput 0.98× / TPOT 1.01× / TTFT+ITL wins); **c8 residual = decode** (tput 0.93× / TPOT 1.10× / ITL P99 1.12×), which nsys shows is 93% GPU-busy/compute-bound (small-M=8 `cutlass_80_wmma` projections). The **qkv-merge** (new GQA `QkvSplit` op mirroring vLLM `QKVParallelLinear`) was implemented + MEASURED NEUTRAL (doesn't cut decode FLOPs) ⇒ default-OFF. **CUTLASS CLAIM CORRECTED: 27B `test_qwen27_paged_engine` = 235/235 on the FLASHINFER cutlass build** (the "flashinfer ⇒ 234/235" was a build artifact). Stays `ACTIVE`; named residual = c8 decode-GEMM efficiency (a decode-fusion sub-campaign). **SWEEP MODEL #1 — Qwen3-Coder-30B-A3B (`Qwen3MoeForCausalLM`) W0+W1 LANDED 2026-07-21** ([sweep-qwen3-coder-30b.md](specs/sweep-qwen3-coder-30b.md), `MODEL-TEXT-qwen3-moe-qwen3-moe-for-causal-lm` `ACTIVE`, `CLAIM-MODEL-QWEN3-CODER`): the first full-attention BF16 MoE, composed from the done dense attention + the done 35B MoE experts (ZERO runner change). W0 = registry stub (`qwen3_moe_registry.cpp` + `qwen3_moe.h`, one `REGISTER_VLLM_MODEL`, full-attn-only KV, `is_dense_model=false`, W2/W3 throwing stubs). W1 = three behaviour-preserving refactors making the two done pieces reusable: (#1) dense `AttnBlock` + glue extracted to `dense_attn_block.h` (Qwen3-dense byte-identical), (#2) bf16 `MoeBlock` exposed cross-TU via `RunMoeBlock` (`qwen3_5_moe_block.h`; 35B untouched), (#3) no-shared-expert guard (inert for the 35B). Gated: dgx CUDA `-Werror` 0-warn; Qwen3-dense 0.6B+4B 16/16 + 27B 235/235 + 35B 315/315 UNCHANGED; registry resolves `Qwen3MoeForCausalLM`; memcheck 0. W2 bf16 loader → W3 forward → W4 near-tie token-exact → W5 fast bf16 grouped-MoE GEMM remain. Then Llama dense (download), Mistral, MoE families **SWEEP MODEL — GLM-4-9B-0414 (`Glm4ForCausalLM`) G2 LANDED 2026-07-24 — CORRECTNESS COMPLETE** ([glm-dsa-latest-deepseek](specs/glm-dsa-latest-deepseek.md), `MODEL-TEXT-glm4-glm4-for-causal-lm` now `READY` per the [live-state audit](specs/live-state-audit-2026-08-06.md), `CLAIM-GLM-DSA-LATEST-DEEPSEEK` amended) **GLM-4.7-Flash (`Glm4MoeLiteForCausalLM`, 31.2B MLA+MoE) G1 LANDED 2026-07-24 — SACRED gate 8/8, CORRECTNESS COMPLETE** (`MODEL-TEXT-glm4-moe-lite-glm4-moe-lite-for-causal-lm` `ACTIVE`, speed pending): reuses the DeepSeek-V2 MLA stack; first e2e coverage of the q_lora branch + noaux_tc sigmoid router, closing the MLA campaign's two C2 coverage gaps: the FIRST GLM-family model. SACRED gate 16/16 vs vLLM 0.25.0 (STRICT 13/16 + near-tie 3/16, max gap 0 nats; vLLM K=5 self-deterministic ⇒ STRICT bar), speed PENDING. The two "new primitives" the spike named reduced to EXISTING infra: partial + interleaved `RopeFromCache` (`is_neox_style=false`, the DeepSeek-V2 decoupled-rope path) over `rotary_dim=64`; standalone `vt::RmsNorm` sandwich norms. Biased qkv, no QK-norm, GQA 32/2, untied lm_head. New files + one REGISTER, reusing the shared dense glue. **SWEEP MODEL — Llama-3.2 (`LlamaForCausalLM`) W0-W4 LANDED 2026-07-23 — CORRECTNESS COMPLETE** ([sweep-llama-3.2](specs/sweep-llama-3.2.md), `MODEL-TEXT-llama-llama-for-causal-lm` `ACTIVE`, `CLAIM-MODEL-LLAMA-3.2`): the roadmap's explicit "Llama-first" increment and the first mainstream non-Qwen/non-OPT dense arch. `LlamaForCausalLM` (Llama-3.2-1B) = the Qwen3-dense forward with exactly two ADDITIVE deltas — NO qk-norm (shared `AttnBlock` skips it when q_norm/k_norm empty) + llama3 rope-scaling (4 default-0 `RopeArgs` fields + a `Llama3ScaleFreq` kernel helper, no-op elsewhere; formula verified 2e-7 rel vs vLLM) — reusing the shared dense forward VERBATIM (`LlamaModel == Qwen3DenseModel`). 3 new Llama files, ZERO edit to runner/scheduler/platforms/attn-registry/`hf_config`/any qwen3-opt model. vLLM 0.25.0 greedy MEASURED DETERMINISTIC (K=6, 0 multi-valued cells) ⇒ STRICT token-exact bar, PASS **16/16 (12 strict + 4 near-tie band, max gap 0.0000 nats, 0 divergent)** — at all 13 divergent positions vLLM's own teacher-forced argmax given our prefix IS our token. A correctness-fatal tokenizer bug (Llama's `Sequence` post_processor wrapping `TemplateProcessing` ⇒ BOS 128000 never prepended, silently 1/16) was isolated by a CUDA prefill-argmax diagnostic (forward proven 4/4 correct given vLLM's exact tokens) and fixed byte-preservingly (Qwen/OPT/DeepSeek unaffected — all ByteLevel/top-level-TemplateProcessing). Regressions 27B 235/235 · 35B 315/315 · Coder 6/6 · Qwen3-dense 16/16 · OPT 6/6 · DeepSeek-V2 8/8 UNCHANGED; `-Werror` 0-warn; memcheck 0; DSR 67. SPEED pending (head_dim 64 → generic paged path; Llama-3.2-3B head_dim-128 is the FA2-toggle W-next). **MLA CAMPAIGN W0+W1 LANDED 2026-07-21** (`CLAIM-MLA-DEEPSEEK`; rows STAY `SPIKE` — W0/W1 make no model supported). **W0 grounded every fact the spike flagged as an unverified source read; ALL CONFIRMED, none contradicted:** DeepSeek-V2-Lite fetched to dgx (30 GB, 4 shards) and loading in the vLLM 0.25.0 oracle; the oracle's own DEBUG startup on sm_121 prints `Using TRITON_MLA attention backend out of potential backends: ['TRITON_MLA']` and `Using FLASH_ATTN MLA prefill backend` — so the dense-MLA decode + MLA-prefill targets are OBSERVED, not inferred, and the sm90/sm100-only MLA kernel class stays out of scope; the real `config.json` confirms every §5.1 number (`kv_lora_rank=512`, `qk_nope=128`, `qk_rope=64` -> the **576-wide latent**, `v_head_dim=128`, `q_lora_rank=null`, `n_group=topk_group=1`, softmax/greedy, 64+2 experts, 27 layers) plus `is_neox_style=False` and the mscale2 scale correction; and BOTH recorded coverage gaps (no `fused_qkv_a_proj` branch, no `e_score_correction_bias`) are confirmed real, so those pieces stay unit-gated only. **W1 = the behaviour-preserving spec-driven KV allocation, ZERO MLA math:** the attention cache is now sized `num_blocks * spec->page_size_bytes()` and viewed from the spec's own `block_size`/`num_kv_heads`/`head_size`/`dtype` instead of the hardcoded `2 * block * Hkv * Dh` with shape reconstructed from the HF config (`runner.cpp`), plus `MLAAttentionSpec` with upstream's factor-1 single-tensor page formula (`kv_cache_interface.py:397-398`) registered against the ORDINARY `FullAttentionManager` (`single_type_kv_cache_manager.py:1539`) — the spike's key finding, which is why block manager/prefix caching/eviction need no change. Gated: dgx clean CUDA `-Werror` 0 warnings/0 errors; **27B 235/235 + 35B 315/315 + Qwen3-Coder 6/6 + Qwen3-dense 16/16 ALL UNCHANGED** (behaviour-preserving proven, not assumed); `test_runner` 15/15, `test_kv_cache_interface` 21/21 (4 new MLA-spec cases), `test_llm_engine` 5/5; the new path is proven EXERCISED (not merely compiled) by `fa_page_size_bytes()` + a `page_size_padded` case no HF-config formula can produce. **W2 + W3 LANDED 2026-07-21** (base `a05437f`; rows STAY `SPIKE` — still no MLA attention math, no MLA model, no forward). **W2 = the MLA branch of `_get_backend_priorities` the pre-W2 comment deferred, ported as DATA:** the whole of `cuda.py:84-176` (BOTH branches — MLA sm_100 including the `:96-115` adaptive sparse tail, MLA sm_12x, MLA `else`, and the two pre-existing non-MLA arms) is now a TABLE in the new header `include/vllm/platforms/cuda_attn_priority.h`, one row per upstream arch arm keyed on `(use_mla, major)`, so a future arch is a ROW rather than a code path; putting it in a header (not the CUDA-only TU) also let the CPU test tier assert the REAL table and DELETED the hand-copied `FakeCudaPlatform` duplicate. On sm_121 a `use_mla=true` request now RESOLVES to `TRITON_MLA`, matching the W0 oracle observation. **The sparse/DSA seam is left OPEN and unit-proven:** GB10's row keeps both upstream entries and the sparse one loses to a real FILTER — `AttentionBackend::is_mla()`/`is_sparse()` checked against the request (`backend.py:307-360 validate_configuration`) — so a future DSA backend is selected purely by declaring `is_sparse() == true`, with ZERO edit to the table or the selector. `TritonMLABackend` lands the NAME plus upstream's 3-D `get_kv_cache_shape` (no K/V axis; `num_kv_heads != 1` REFUSED), `get_impl_cls()` deliberately still `nullptr`. MLA prefill priority ported too (GB10 -> `[FLASH_ATTN]` alone). **W3 = the two new `vt::` ops, both CPU-reference-gated.** `vt::ConcatAndCacheMla` mirrors `csrc/libtorch_stable/cache_kernels.cu:401-442` — and per the whole-chain rule this was VERIFIED, not assumed, to be vLLM's OWN csrc kernel (`_custom_ops.py:2532` -> `torch.ops._C_cache_ops`), with no flashinfer/cutlass variant in the dense-bf16 path; it concatenates the latent + rope part into ONE 576-wide entry, the write `ReshapeAndCache`'s K/V-pair signature cannot express, stride-driven so a per-layer cache slice and the two column halves of `kv_a_proj_with_mqa` both work copy-free. The **grouped-topk (`noaux_tc`) router** extension — flagged in `coordination.md` as SHARED with `CLAIM-GLM-DSA-LATEST-DEEPSEEK` and "must not be implemented twice" — is landed HERE and that claim now consumes it: additive `MoeRouterTopKArgs` fields + an optional `e_score_correction_bias`, with `num_expert_group == 0` still dispatching the ORIGINAL kernel so the 27B/35B/Coder/dense routers are byte-identical BY CONSTRUCTION. **Stated plainly: the `noaux_tc` correctness evidence is UNIT-ONLY.** V2-Lite has `n_group=topk_group=1` and no bias, so the e2e vehicle exercises none of it; the gate is `tests/vt/test_ops_moe_router_grouped.cpp` at DeepSeek-V3's REAL dimensions (256 experts, n_group=8, topk_group=4, sigmoid, scaling 2.5, WITH the bias) against an INDEPENDENT sort-based transcription of the upstream formula. **W4 LANDED 2026-07-22** (base `ed2c342`; rows STAY `SPIKE` — W4 adds a kernel and fills a `nullptr`, it makes no model supported). **`vt::MlaDecodeAttention` — the MQA decode over the compressed latent (QK 576 / V 512, `num_kv_heads=1`), a structure port of the two-stage split-KV pair W0 OBSERVED EXECUTING:** `MlaDecodeStage1` <- `_fwd_grouped_kernel_stage1` (`triton_decode_attention.py:278-458`, the `IS_MLA` branch whose `v = tl.trans(k)` at `:424-431` is the whole MLA trick — V is the leading 512 columns of the SAME latent row already loaded as K, so one shared-memory tile serves as both), `MlaDecodeStage2` <- `_fwd_kernel_stage2` (`:575-639`), `ComputeNumKvSplits` <- `_compute_num_kv_splits` (`triton_mla.py:40-47`), split workspace <- `_reserve_attn_logits_workspace` (`:57-78`) realized as the house grow-only per-stream scratch. **Honest reuse verdict:** our FA-2 split+combine machinery fit at the ALGORITHM level (the split schedule, the LSE merge algebra, the fixed-ascending no-atomicAdd determinism rule) and NOT at the code level — the vendored FA-2 launcher takes separate 4-D k/v caches and is instantiated for symmetric head_dim {128,256}, which cannot express a 3-D single-buffer cache with QK 576 / V 512; that is recorded in the TU header rather than forced. **Evidence is unit-level and deliberately strong** (there is no e2e model until W7): [`tests/vt/test_ops_mla_attn.cpp`](../tests/vt/test_ops_mla_attn.cpp), a port of `tests/kernels/attention/test_mla_decode_cpu.py` whose `ref_mla` becomes an INDEPENDENT TWO-PASS oracle (a different algorithm from the streaming online-softmax both impls use) plus its NaN-padding out-of-bounds detector, run at the REAL V2-Lite geometry (576/512/64, block 16, mscale^2 scale) over ragged / multi-block / single-block / every split boundary (`num_kv_splits` ∈ {1..512} incl. splits > seq_len) / 128-head V3 / non-BLOCK_H head counts / a 288-256 block-32 geometry / bf16 + f32, with run-to-run BIT-exactness. dgx sm_121: 11/11 cases, 2,303,193 assertions; `compute-sanitizer` memcheck **0 errors**, racecheck **0 hazards**, synccheck **0 errors**; clean CUDA build **0 warnings / 0 errors**; **27B 235/235 + 35B 315/315 + Coder 6/6 + Qwen3-dense 16/16 + OPT 6/6 ALL UNCHANGED**. `TritonMLABackend::get_impl_cls()` is no longer `nullptr` — it returns a real `TritonMLAImpl` whose `forward_mqa` is the 1:1 counterpart of `triton_mla.py:189-260`; PREFILL is W5 and `forward()` refuses a prefill-shaped batch BY NAME rather than producing wrong numbers. NO speed number — decode perf is W9. **W5 LANDED 2026-07-22** (base `5395203`; rows STAY `SPIKE`). **MLA PREFILL + the workspace-bounded CHUNKED-CONTEXT loop.** Three new ops — `vt::MlaPrefillAttention` (<- `mla/prefill/flash_attn.py:153-248`, the ONLY MLA prefill backend reachable on sm_121 and the one W0 OBSERVED the oracle logging), `vt::GatherMlaCache` (<- `cache_kernels.cu:992-1064`) and `vt::MergeAttnStates` (<- `merge_attn_states.cu:18-192`, both `-inf` edge cases verbatim) — plus the loop itself in the new `mla_chunked_context.h` (<- `mla_attention.py:1422-1451,1667-1745,2094-2199,2344-2425`), which is what keeps a long-context prefill inside a bounded workspace instead of materializing a 3 GB up-projected context. **The vendored FA-2 launcher WAS generalized, and W4's prediction that it would be tractable held for a reason worth recording: upstream does not ask FA-2 for asymmetric head dims either.** `requires_v_padding` is TRUE on GB10, so upstream ZERO-PADS V from 128 to 192 and slices the output back — the kernel stays a plain SYMMETRIC head_dim-192 instantiation. The whole change is two new explicit instantiations of an UNCHANGED generic template, one new launcher entry for the contiguous-varlen mode, and the pad/slice pair; the paged launcher every non-MLA prefill calls is textually untouched (211 insertions / **0 deletions**), and 27B 235/235 + 35B 315/315 + Coder 138/138 + Qwen3-dense 664/664 + OPT 36/36 are all UNCHANGED. Evidence is UNIT-ONLY (there is still no model): 4/4 cases / **2,377,052 assertions** and 5/5 / **306,037 assertions** on dgx sm_121 at the real QK 192 / V 128 geometry, against an INDEPENDENT double-precision two-pass oracle and — for the loop — a SINGLE-SHOT whole-sequence oracle that never chunks, over exact / +1 / -1 chunk boundaries, zero-context and zero-key-in-chunk requests, ADVERSARIAL reverse-interleaved block tables, NaN-poisoned outputs and run-to-run bit-exactness; memcheck **0**, racecheck **0 hazards**, synccheck **0**. A genuine upstream FA-2 quirk was found on the way and worked around rather than papered over: its EMPTY-K early exit ignores the unpadded-LSE flag, which a zero-key chunk request would turn into an out-of-bounds LSE write. **W6 LANDED 2026-07-22** (base `2846467`; rows STAY `SPIKE` — W6 adds an attention LAYER, not a model). **The MLA attention BLOCK + LOAD-TIME WEIGHT ABSORPTION — the piece that finally COMPOSES W3's cache write, W4's MQA decode and W5's MHA prefill into one layer:** the projections with BOTH `q_lora_rank` branches (`fused_qkv_a_proj` -> `q_a_layernorm` -> `q_b_proj`, or the direct `q_proj`), the two RMSNorms (the rope part deliberately NOT normed), the DECOUPLED RoPE (`is_neox_style=False`, only the trailing 64-dim slice rotates) with its YaRN cos/sin cache and the SEPARATE mscale^2 softmax-scale correction, the `kv_b_proj -> W_UK/W_UV` split, the prefill-MHA / decode-MQA dispatch with decode tokens packed FIRST, and the `kv_b_proj` up-projection callback W5 left open. **The spike's most useful prediction held: absorption needed NO new attention kernel** — it is a LOAD-TIME weight transform plus TWO batched GEMMs, so the entire new-kernel surface is two general primitives, `vt::BatchedMatmul` (<- `torch.bmm` at `mla_attention.py:789,1034`; on CUDA torch resolves that to cuBLAS `gemmStridedBatchedEx`, and ours is the cuBLASLt strided-batched form of the same GEMM) and `vt::ConcatMlaNopeRope` (<- `concat_mla_q`, generalized so one op also serves `_concat_k_nope_k_pe`). **The absorbed-vs-unabsorbed equivalence — the heart of W6 — is PROVEN NUMERICALLY, three independent ways, rather than argued:** an INDEPENDENT double-precision block oracle computes the attention BOTH ways and agrees to < 1e-11 (the identity itself); our absorbed decode reproduces the UNABSORBED oracle to < 2e-4 in f32; and — the strongest — the SAME batch is driven once through our ABSORBED MQA decode kernel (QK 576 / V 512, one KV head, K/V never materialized) and once through our UNABSORBED materialized-MHA prefill path (QK 192 / V 128 plus the chunked-context loop), agreeing to < 3e-4 (CPU f32) / < 4e-2 (CUDA bf16) with nothing but the weights shared between them. Evidence on dgx sm_121: `test_mla_attention_block.cpp` 10/10 cases / 2,372,644 assertions and `test_ops_mla_absorb.cpp` 9/9 / 1,644,807 (CUDA cases proven to EXECUTE; NaN-poisoned outputs; run-to-run BIT-exact), porting `tests/kernels/test_concat_mla_q.py` in both arms. memcheck 0, racecheck 0 hazards, synccheck 0 (the last needing `--num-cuda-barriers 65536` — the default table overflows on a binary driving this many kernel families and the tool then reports a bogus launch failure, a trap worth knowing). Clean CUDA build 0 warn / 0 err; **27B 235/235 + 35B 315/315 + Coder 138/138 + Qwen3-dense 664/664 + OPT 36/36 ALL UNCHANGED**. **Coverage stated plainly: the `q_lora` query branch has NO e2e coverage and cannot get any on GB10** — DeepSeek-V2-Lite has `q_lora_rank=null`, so it is unit-gated at DeepSeek-V3's real dimensions only; GLM-4.7-Flash (`q_lora_rank=768`, 58.2 GiB, fits) is what would close it. **W7 LANDED 2026-07-22** (base `ce43c51`; the row STILL stays `SPIKE`). **THE DEEPSEEK-V2 MODEL — registry + config parse + loader + forward: the first MLA model in the tree, and the first one that runs a real MLA checkpoint end to end.** Four new files plus ONE shared-code edit (a two-line additive condition in `runner.cpp` recognising a `kMlaAttention` KV group as the model's attention group — upstream maps MLA onto the ordinary `FullAttentionManager`, so block tables/prefix caching/eviction are untouched). **LOADER GATE PASSED on the real 4-shard DeepSeek-V2-Lite: 5291/5291 checkpoint tensors accounted for, ZERO unmapped and ZERO leftover** (4/4 cases / 37,331 assertions), every shape asserted including the LOAD-TIME `kv_b_proj -> W_UK_T [16,128,512]` / `W_UV [16,512,128]` absorption split — the same transform, at the same lifecycle point, as upstream's `process_weights_after_loading`. **V2-Lite takes the DIRECT `q_proj` query branch** (`q_lora_rank: null`), asserted with the fused branch EMPTY on every layer. **FORWARD GATE PASSED and obviously right, not merely finite: the real checkpoint prefill of `The capital of France is` -> argmax ` Paris`** (top-5 ` Paris`/` the`/` a`/` one`/` also`, run-to-run bit-exact) — the direct analogue of the Qwen3-Coder W3 sanity case. **BATCH-ORDERING GATE:** the ordering invariant W6 measured 0.86 relative error from is now VALIDATED, not assumed — `BuildMlaBatchSplit` throws (naming the request and citing the upstream line) if a decode follows a prefill or a with-context prefill follows a context-free one. **SHARED EXPERTS — new for this family and UNGATED unlike Qwen3.6's sigmoid-gated one — gated two ways:** a MoE layer with every routed expert zeroed is BIT-IDENTICAL to a dense layer holding the same MLP, and turning the shared expert off CHANGES the logits. **The CUDA path is EXERCISED, not merely compiled:** a case at the real MLA head geometry drives the CUDA MLA kernels and the CUDA-only grouped bf16 MoE GEMM, bit-exact on device and within 0.0061 worst relative logit error of the CPU reference path. 11/11 forward cases; memcheck/racecheck/synccheck all **0**; clean CUDA build **0 warn / 0 err**; **regression set UNCHANGED**. **Only `DeepseekV2ForCausalLM` is REGISTERED** — `DeepseekForCausalLM` (plain MHA), V3 (fp8/671B) and V3.2 (DSA indexer) are REFUSED BY NAME in the config parse rather than falsely claimed. A pre-existing tree-wide hazard was found on the way and recorded: the shared `DevicePool` is a process-wide singleton keyed only on a byte size class, so a single process driving BOTH a CPU and a CUDA forward hands the second backend the first's recycled pointers. **NEXT: W8 — the SACRED token-exact gate on DeepSeek-V2-Lite** (wire the paged engine to produce the MLA batch order the model already validates, capture oracle goldens, run the STRICT form W0 determined). A loading, forwarding model is NOT a supported model, so no model row moves until that gate passes. **W8 LANDED 2026-07-22 — THE SACRED CORRECTNESS GATE PASSES 8/8, and `MODEL-TEXT-deepseek-v2-deepseek-v2-for-causal-lm` moves `SPIKE` -> `ACTIVE` (correctness COMPLETE, speed PENDING). NOT `DONE` — that additionally requires vLLM-speed parity on every axis, which is W9 and has NO number yet; the other four campaign rows stay `SPIKE`.** An 8-prompt battery is driven through the FULL paged `LLMEngine` and compared to the pinned vLLM 0.25.0 oracle: **8/8 PASS — STRICT token-exact 5/8, near-tie band 3/8, 92/128 tokens strictly exact, max teacher-forced gap 0.25 nats, 0 forward-divergent** (223 assertions). **The bar was ARRIVED AT by measurement, not chosen:** vLLM is DETERMINISTIC on this model at batch=1 (W0's K=5 8/8, re-confirmed by W8's own capture at T=16 with 0 multi-valued cells), so the STRICT form ran FIRST and came out 5/8; the ratified TEACHER-FORCING diagnostic then showed **36 divergent positions with 35 at gap EXACTLY 0.0000 nats** — vLLM's own argmax GIVEN OUR PREFIX is our token, so they are the downstream tail of one earlier flip — **exactly ONE root flip with any gap at all (prompt[3] tok 9, 0.2500 nats, inside the ratified 0.5-nat band and equal to the landed Qwen3-dense 4B gate's worst)**, and **ZERO tokens outside vLLM's top-20**, with the per-position nats COMMITTED as goldens and anything beyond the band still FAILING. **W8's first job — the scheduler/runner wiring — turned out to need NO new code, for a non-accidental reason:** `runner.cpp:671` already reorders with `decode_threshold = 1`, exactly MLA's `reorder_batch_threshold` (`mla_attention.py:1420`), and its `decode -> short_extend -> long_extend -> pure_prefill` ordering satisfies BOTH MLA invariants (decodes form a batch prefix; with-context prefills lead the prefill tail). W8 PROVES that end to end rather than duplicating it, with new DIAGNOSTIC `MlaBatchSplitStats` counters and a non-vacuity bar: the battery is admitted CONCURRENTLY with staggered arrival, producing **7 genuinely MIXED decode+prefill steps at up to 8 concurrent requests** with `BuildMlaBatchSplit` (which throws naming the request) never firing, plus a prefix-cache-driven **with-context prefill**, and a phase-0 check that the engine really allocated the compact MLA cache (`fa_page_size_bytes = 36864`, no factor 2). **THE REAL BLOCKER WAS THE TOKENIZER, NOT THE MODEL:** the first run REFUSED to load (`unsupported normalizer "Sequence"`), and behind it sat a whole NEW pre-tokenizer family — DeepSeek's is a HF `Sequence` PIPELINE of SEVEN stages (five `Split(Isolated)` over ENUMERATED codepoint ranges, then `Digits(individual_digits=true)`, then `ByteLevel(use_regex=false)`), whose stage ORDER is load-bearing because stage 2's punctuation class spans 0x3A-0x7E and CONTAINS A-Z/a-z. Landed as `SplitPattern::kDeepSeek` with the five patterns compared VERBATIM at load, and MEASURED token-for-token against the REAL HF `tokenizers` library over a stage-stress corpus (**6/6 cases / 2461 assertions**). **The TOKENIZATION goldens earned their keep by REFUTING a fix that was already written:** `tokenizer_config.json` declares `add_bos_token: true`, which reads as exactly the OPT missing-BOS bug — but vLLM's resolved tokenizer (`TokenizersBackend`) adds NO BOS, our loader already matched bit-for-bit, and the "fix" would have BROKEN a passing gate; it was reverted and the measured behaviour PINNED by a guard case ([[ground-premises-before-dispatching]]). Regression set UNCHANGED (27B 235/235, 35B 315/315, Coder 6/6, Qwen3-dense 16/16, OPT 6/6, plus every tokenizer test — W8 touches SHARED tokenizer code, so that was proved, not assumed); clean CUDA rebuild 0 warn/0 err; local CPU suite 151/151; memcheck/racecheck/synccheck 0. Batch invariance is REPORTED (6/8) and deliberately NOT a bar, because the ORACLE itself changed on 3/8 of this battery under batched generation (W0). One W9 input recorded: the oracle must run `moe_backend='triton'` — vLLM's auto-selected FlashInfer CUTLASS unquantized MoE REBOOTED dgx three times on GB10's unified memory. **W9 SPEED CLOSE LANDED 2026-07-22 — the track has its FIRST binding speed number, and it is an ATTRIBUTED MISS: `MODEL-TEXT-deepseek-v2-deepseek-v2-for-causal-lm` STAYS `ACTIVE` (correctness COMPLETE, speed SHORT), NOT `DONE`** ([grid](../docs/BENCHMARKS.md), [spike §W9](specs/mla-deepseek-campaign.md)). Denominator SETTLED with evidence — CUTLASS MoE has now rebooted dgx **five times** (two more at W9, the second on a pristine box with a 0 GiB page cache and every mitigation applied, both deaths at the identical post-`torch.compile` phase), so `--moe-backend triton` IS vLLM's best STABLE GRAPHED configuration here and is the bar; the substitution does not flatter us, we lose to it. `nsys` (both sides, `--cuda-graph-trace=node`) overrode the plan: the lever was not the planned MLA fusion recipes but `MlaDecodeStage1` sitting at **44.7% of all GPU time and ~180x off its own memory-bound floor** on a **2-CTA grid at batch 1**; applying upstream's own occupancy target made it **18.3x faster** (837 -> 45.8 us) for **+69.5%/+53.3%/+32.0%/+19.5%** end-to-end at c1/c2/c4/c8, while the planned decode-graph sibling is worth only ~+2% (this decode is GPU-bound). Grid vs vLLM: output throughput **0.87/0.95/0.86/0.88** (was 0.50 at c1), TTFT **1.06/1.14/0.96/0.88** (we WIN at c4/c8), TPOT **1.11/0.97/1.16/1.17**. SACRED gate **8/8 UNCHANGED** with both levers default-ON; a real latent CUDA-graph use-after-free in the MLA metadata upload was found and fixed (its whole class now guarded); regression set UNCHANGED; clean rebuild 0 warn/0 err; sanitizers 0. **NEXT LEVER, NAMED: route the batch-1 dense projections off cuBLAS `gemvx` (31.8% of our GPU time) onto a tensor-core GEMM — vLLM splits the same work `gemvx` 12.7% + `nvjet_sm121_tst_mma_*` 6.6%.** **W10 BLOCKED-ROW HONESTY PASS LANDED 2026-07-22 — the campaign's W-plan is COMPLETE; records only (no code, no build, no GPU work, nothing downloaded, no number claimed).** Rows set to their final honest state: `MODEL-TEXT-deepseek-v2-deepseek-v3-for-causal-lm` (V3 + V3.2, and Kimi-K2/K2.5's text backbone by config composition) and `MODEL-TEXT-minimax-m2-mini-max-m2-for-causal-lm` move `SPIKE` -> `BLOCKED`, joined cross-claim by `MODEL-TEXT-deepseek-v2-glm-moe-dsa-for-causal-lm` (GLM-5) under `CLAIM-GLM-DSA-LATEST-DEEPSEEK`; each is HW-BLOCKED on 119 GiB (~642 GiB fp8 / ~428 GiB / 1404 GiB) and the two DSA models are additionally DEP-BLOCKED — for a SPARSE model the XOR filter eliminates `TRITON_MLA`, leaving `FLASHINFER_MLA_SPARSE_SM120` alone, whose sm12x dispatch goes to flashinfer's DENSE-ONLY XQA backend that discards `sparse_mla_top_k` (upstream's own test monkeypatches the probe and asserts nothing numerical). `MODEL-TEXT-deepseek-v2-deepseek-for-causal-lm` stays `SPIKE` with the record repaired to say it is plain MHA and needs NO MLA; Kimi-Linear stays `SPIKE` (MLA half unlocked, KDA a separate kernel campaign, HW-MARGINAL). Each blocked row states what CAN still be gated (config resolution, weight-map on a slice, unit parity at real dimensions) versus what CANNOT (anything e2e). **Two PERMANENT coverage gaps now stated in the rows:** the `noaux_tc` grouped router and the `q_lora` query branch have NO e2e coverage and are unit-gated only, because V2-Lite is `n_group=topk_group=1`/softmax with no `e_score_correction_bias` and `q_lora_rank=null`. **NAMED NEXT VEHICLE: GLM-4.7-Flash** (`MODEL-TEXT-glm4-moe-lite-glm4-moe-lite-for-causal-lm`, 31.2B / 58.2 GiB, FITS GB10) — the only reachable checkpoint that closes BOTH gaps. **BLOCK NOT CLOSEABLE, nothing archived:** the DeepSeek-V2 row is `ACTIVE`, not `DONE`, so the plan/spec stay LIVE; the single open item is the named `gemvx` -> tensor-core dispatch lever. **MISTRAL FIFTH FAMILY W0-W3 LANDED 2026-07-23** ([sweep-mistral](specs/sweep-mistral.md), `MODEL-TEXT-mistral-mistral-for-causal-lm` `ACTIVE`, `CLAIM-MODEL-MISTRAL`): the closest-to-Llama dense arch (vLLM `mistral.py` = "Mistral adaptation of the LLaMA architecture") — plain rope θ1e6 (no rope_scaling) + qk-norm-optional + untied lm_head + null sliding_window, all PRE-EXISTING ⇒ NO new primitive, 3 new files + additive CMake/registry-test rows only, ZERO shared-code edit. **MODEL forward gate 30/30 greedy tokens vs vLLM 0.25.0** (tokenizer-free: fed vLLM's exact prompt tokens through our CUDA prefill; 29 STRICT token-exact + 1 near-tie, 0 forward-divergent; vLLM greedy det 4/5 K=3). W2 loader real-weights 1541 assertions. **REAL FINDING:** Mistral's SentencePiece/Metaspace tokenizer is unsupported by our ByteLevel-BPE tokenizer → the FULL paged-engine SACRED gate is BLOCKED, the pre-inventoried `LOAD-SENTENCEPIECE` row (SentencePiece tokenizer family). `-Werror` 0-warn, DSR 32, regressions UNCHANGED (Llama paged 16/16, Qwen3-dense forward 1031, registry 299; MoE/GDN gates unaffected by construction). SPEED + full paged gate both PENDING (row `ACTIVE`, not `DONE`). **OLMo-2 SPIKED 2026-07-24** ([sweep-olmo2](specs/sweep-olmo2.md), `CLAIM-SWEEP-OLMO2`): one row `INVENTORIED` → `SPIKE` (`MODEL-TEXT-olmo2-olmo2-for-causal-lm`, covering `Olmo2ForCausalLM` + its `Olmo3ForCausalLM` alias). Answers the breadth-sweep §B.3 Tier-2 rank-8 "GLM4 / Olmo2-3" item (GLM-4 + Gemma landed; OLMo-2 next). **HEADLINE: OLMo-2 is the cleanest dense bring-up yet — ZERO new compute kernels.** The two distinctive facts both reduce to WIRING over landed ops: (1) the **pure post-norm (`norm_after`) placement** is a strict SUBSET of the GLM-4/Gemma sandwich (keeps ONLY the standalone-output-norm op `glm4.cpp:174-178` — the exact primitive flagged — DROPS the pre-norms, plain residual add); (2) the **QK-norm is FULL-WIDTH not per-head** → reuses `vt::RmsNorm` at a `[T,q_size]`/`[T,kv_size]` shape but CANNOT use the fused per-head `kAttnQkNormRopeGate`. Everything else REUSES (plain RMSNorm, SiLU SwiGLU, NeoX rope, GQA paged glue, Gemma-3 sliding-window for Olmo-3, tied embeddings, packed loader, ByteLevel BPE). Gate vehicle `allenai/OLMo-2-0425-1B` (1.485B, ~2.77 GiB, fits GB10 tight ~30 GiB free); Olmo-3 rides the same row (0.25.0 oracle-support UNVERIFIED). OLMo-1 (non-parametric LayerNorm), OLMoE/FlexOlmo (MoE), OlmoHybrid (SSM) stay `INVENTORIED`. Nothing implemented. | -| 2a | `ROAD-V1-C2-LOCAL-BF16` | Rebase the local discrete-Blackwell Qwen3.5 plain-BF16 diagnostic onto the current additive model/loader seams | [`MODEL-MM-qwen3-5-qwen3-5-for-conditional-generation`](model-matrix.md), [`LOAD-SAFETENSORS-DIRECT-DENSE`](engine-matrix.md) | Rebase conflicts are resolved by transplanting the relevant H32 AOT repair onto current `main`; CPU/CUDA gates and direct-OFF/ON token equivalence are green. H32 AOT (+4.59%), plain-BF16 decode graphs (+0.39%) and ratio-4 FA2 (+1.60%) are implemented and graph-node trace-proven. Final matched 4B ON/OFF/vLLM total is **5769.99/5660.70/5849.80 tok/s**, ON=OFF 128/128 outputs in all pairs, peak PSS **2.406/8.592/7.662 GiB**. Direct ON is **0.9864x** stable vLLM total/output throughput; TPOT/ITL remains **43.72 vs 38.55 ms** | `GATING` | Port vLLM's device-resident sampled-token mapping to discrete CUDA with request-compaction correctness, removing the measured immediate main-stream wait, then rerun the exact series. Sanitizer availability and external 27B/35B follow-ups remain open; no 4B-to-gate-model support extrapolation. Evidence: [2026-07-25 checkpoint](../docs/bench-evidence/qwen35-4b-main-repair-20260725.md) | +| 2a | `ROAD-V1-C2-LOCAL-BF16` | Close local discrete-Blackwell Qwen3.5 plain-BF16 production parity | [`MODEL-MM-qwen3-5-qwen3-5-for-conditional-generation`](model-matrix.md), [`LOAD-SAFETENSORS-DIRECT-DENSE`](engine-matrix.md), [`SERVE-CLI-BENCH`](engine-matrix.md), [`KERNEL-SSM-MAMBA`](kernel-matrix.md) | Exact `(sequence, 8-token chunk)` GDN conv dispatch is byte-identical and default ON. Rebased-main same-binary reprofile: **718.704→233.955 ms (3.072x)** and whole run **+2.272%**; sealed vLLM remains 145.421 ms (**1.609x residual**). The binding three-pair A/B improved total/output **2.152%**, TTFT **2.945%**, TPOT/ITL **1.920%**, no VRAM regression. Against sealed vLLM: throughput **1.021246x PASS**; TTFT, TPOT and VRAM OPEN | `GATING` | Spike/profile the residual causal-conv gap. Latency, VRAM and 27B/35B correctness remain gates; no 4B-to-gate-model extrapolation. Evidence: [production baseline and exact-chunk outcome](../docs/bench-evidence/qwen35-4b-sm120-main-20260807.md), [conv spike/result](specs/sm120-qwen35-conv-chunking-2026-08-07.md) | | 3 | `ROAD-V1-C3` | MTP k=1 + GDN speculative path, then DFlash, DSpark and heterogeneous-vocabulary TLI | [engine matrix](engine-matrix.md), [coverage view §8](feature-matrix.md#8-speculative-decoding) | MTP and DFlash specs exist. **M-mtp-0 CLOSED 2026-07-24** - the standalone MTP draft head is oracle-parity-proven on BOTH gate checkpoints (op-level parity vs a dumped k=1 vLLM oracle, not a token-generation SACRED gate). **I2 SCHEDULER-HALF LANDED 2026-07-24** ([mtp-spec-decode §2.7](specs/mtp-spec-decode.md)): host-side spec-decode scheduler/engine plumbing + the FROZEN spec-metadata ABI that I3 (rejection sampler) and I5 (verify/propose runner) build against - `SpeculativeConfig`, `DraftTokenIds`, `Request::spec_token_ids`/`NumTokensWithSpec`, the first population of `scheduled_spec_decode_tokens`, `Scheduler::update_draft_token_ids`, the `take_draft_token_ids` seam, `EngineCore::post_step`, `InputBatch::num_accepted_tokens`/`update_req_spec_token_ids`; DEFAULT-OFF and INERT (no `SpeculativeConfig` => `num_lookahead_tokens == 0` => byte-identical engine). `SPEC-MTP` **STAYS `GATING`** because M-mtp-1..4 (greedy rejection, GDN spec slots, k>1, CUDA graphs) are still open, so spec decode remains user-invisible. DSpark is user-promoted scope with DeepSeek-V4/Qwen3 draft models, reduced-vocabulary handling and full-CUDA-graph behavior inventoried under `SPEC-DSPARK`; tokenizer-agnostic target<->draft mapping is separately inventoried as `SPEC-TLI`. Their dedicated spikes are not written **I3 GREEDY REJECTION SAMPLER LANDED 2026-07-24** (`SPEC-REJECTION` `READY` -> `ACTIVE`): per-request logits expansion to `1 + k_i` rows plus the greedy accept rule (accept a draft iff it equals the target argmax; on the first mismatch emit the target argmax and stop; bonus token when all k accept), CUDA==CPU bit-exact at vocab 248320. **I4 GDN-HALF LANDED 2026-07-24** (`SPEC-GDN-SEGMENTS` -> `ACTIVE`): the GDN spec metadata split + decode->prefill reclassification, the `T>1`/`IS_SPEC` recurrence with per-timestep snapshots, the conv sliding window honouring `num_accepted`, k+1 slot allocation - bit-exact rollback. **I5a GDN LAYER ROUTING + runner spec-metadata upload LANDED 2026-07-24** (`CLAIM-SPEC-MTP-I5A`): `GdnBlockPaged` now routes a pure-spec batch through `vt::GdnSpecDecode`/`vt::CausalConv1dSpecUpdate` and the runner uploads I4's six spec device tensors - first of the scoped M-mtp-1 sub-increments (I5a GDN wiring -> I5b prepare_prefill -> I5c MTP paged propose -> I5d config+runner-loop+27B token gate, spec §5), DEFAULT-OFF INERT, bit-exact vs the I4 ops, no e2e loop yet. **I5b `prepare_prefill_inputs` LANDED 2026-07-24** (`CLAIM-SPEC-MTP-I5B`): the drafter prefill input-prep - shift each request's `input_ids` left one within its query span, splice the just-sampled next token, `query_len -= num_rejected`, emit last-token index / query_start_loc / seq_lens into the `SpecPrefillInputs` struct; a HOST routine (no new CUDA kernel; mirrors our DEVICE-NEUTRAL `prepare_inputs`/`combine_sampled_and_draft_tokens` family), unit-gated 7 cases / 27 assertions RED-first, DEFAULT-OFF INERT, additive by construction. **I5c MTP PAGED PROPOSE + DRAFT KV LAYER LANDED 2026-07-24** (`CLAIM-SPEC-MTP-I5C`): `Qwen3_5MTPModel::ForwardPaged` runs the head + one full_attention decoder layer over the head's OWN paged draft KV layer (ReshapeAndCache + PagedAttention over the target's block table / slot mapping); `MakeQwen3_5KVCacheSpec(num_spec>0)` adds that draft KV layer (`fa_draft` FullAttentionSpec group, index num_hidden_layers); `ForwardDeviceTap` exposes the `[T,H]` post-final-norm hidden tap (INERT); and `MtpProposePrefill` is the callable k=1 propose (I5b shift-splice -> one paged forward -> argmax at last_token_indices, early-exit). CORE PROOF: the paged forward reproduces I1's standalone head logits/argmax on BOTH gate checkpoints; a two-step drive proves the draft-KV write/read (RED control diverges). DEFAULT-OFF INERT (no spec config -> draft KV layer not allocated, tap nullptr, target forward byte-identical); NOT wired into the runner step loop. **I5d-pre REGISTRY/FORWARD-SEAM ENABLING REFACTOR LANDED 2026-07-25** (`CLAIM-SPEC-MTP-I5D-PRE`): a scoping pass found the model seam is fully TYPE-ERASED, so the runner cannot reach the concrete target weights / hidden-state tap / loaded MTP weights the I5d loop needs. Four ADDITIVE, inert-when-spec-off access paths + one latent-bug fix - the `hidden_tap` out-field on the type-erased `ModelForwardInput` (routes to the existing `ForwardDeviceTap`), a `LoadedModel::BuildMtpDraft` virtual (typed path to the draft, null for non-MTP), MTP weight loading + shard retention in `FromModelDir` behind `EngineParams::speculative_config`, and the `GPUModelRunner` ctor widened with optional draft/draft-KV/`SpeculativeConfig`; PLUS the latent `initialize_kv_cache` fix (select the FIRST non-eagle full-attn group as the target so a third `fa_draft` group can't displace it; byte-identical at num_spec==0). DEFAULT-OFF INERT, unit-gated RED-first, spec-off SACRED gates byte-identical. **I5d CONFIG RUNTIME + VERIFY/PROPOSE RUNNER LOOP LANDED as a spec-off-byte-identical PARTIAL 2026-07-25** (`CLAIM-SPEC-MTP-I5D`): `--speculative-config` JSON parse -> `EngineParams` -> `LoadedEngine` resolution (widened KV `MakeQwen3_5KVCacheSpec(num_spec>0)`, `BuildMtpDraft`, forced sync scheduling, `MakeScheduler(spec)`, `EngineCore(check_for_draft=true)`) + the full runner loop (draft splice, hidden-tap capture, GDN builder spec-overload feed, k+1 GDN state-slot remap + widened conv cache + draft-KV alloc, `MtpProposePrefill`, `take_draft_token_ids`, acceptance telemetry). CUDA `-Werror` 0 warnings, cutlass-ON; spec-OFF SACRED byte-identical (27B 235/235, 35B 315/315, Coder 138/138 + spec unit tests ALL PASS). **The three-way 27B token gate is NOT yet passing**: the spec-ON engine RUNS the loop end to end and MEASURES the blocker (`test_qwen27_spec_decode`) - it throws on the FIRST prefill step at `gdn_state_gather: working/cache row shapes must match` (`src/vt/ops.cpp:1773`) because I4's spec conv rollback needs the conv row widened to `(K-1)+num_spec` while the non-spec GDN conv ops assume `(K-1)`. `SPEC-MTP` STAYS `GATING`. **I5e LANDED 2026-07-25 (`CLAIM-SPEC-MTP-I5E`): the non-spec GDN conv ops made widened-cache-aware (mirror vLLM `state_len=KERNEL_WIDTH-1` + physical `stride_conv_state_tok`, leading `(K-1)` sub-window, byte-identical at `num_spec==0`) AND the async input-combine forced off under spec (it overwrote the verify batch's draft position with the committed token -> 0 acceptance, RCA'd on the real 27B). **THE THREE-WAY 27B SINGLE-REQUEST GREEDY GATE PASSES**: our-spec-ON == vLLM `--speculative-config mtp` greedy == our-spec-OFF token-for-token, **acceptance 16/16 drafts accepted** (~16 target steps saved); spec-OFF SACRED byte-identical (27B 235/235, 35B 315/315, Coder 138/138), compute-sanitizer 0 on the spec step. `SPEC-MTP` LEAVES `GATING` (single-request greedy correctness PROVEN); NOT `DONE` - the MIXED `GdnBlockPaged` concurrency split/merge + the throughput A/B vs vLLM same-config are I6. **I6 LANDED 2026-07-25 (`CLAIM-SPEC-MTP-I6`, `benchmark_binding=true`): the §5 c1 THROUGHPUT GATE — ours spec-ON AT/ABOVE vLLM spec-ON on EVERY measured axis at c1** (TPOT 66.2/62.95 vs 69.1/65.3 ms prose/code, ours ~1.04x faster; output tput +4.6%/+3.9%; ITL/TTFT lower; acceptance ours 0.85/0.92 vs vLLM 0.838, within noise; spec helps both ~1.5-1.6x TPOT; ours ~4% faster spec-OFF too), via an additive example-only `--speculative-config` bench flag (NO engine code touched). STAYS `ACTIVE`: the c>1 mixed-batch path is still refused + owes a c>1 A/B, and no server-facing spec flag yet. | `ACTIVE` | M-mtp-0, I2 scheduler-half, I3 rejection sampler, I4 GDN spec slots, I5a GDN layer routing, I5b prepare_prefill, I5c MTP paged propose + draft KV, I5d-pre the registry/forward enabling seam, and I5d config runtime + verify/propose runner loop (spec-off byte-identical) are landed; next (before `SPEC-MTP` leaves `GATING`) is closing the measured I5d gate blocker - make the non-spec GDN conv ops widened-cache-aware (mirror vLLM `causal_conv1d` `state_len=width-1+(seqlen-1)`) + the MIXED `GdnBlockPaged` split/merge - then the passing M-mtp-1 27B k=1 greedy three-way token gate + acceptance, then M-mtp-2 35B, then DFlash, the DSpark spike/gates and TLI. **DFlash D0-redo + D1 LANDED 2026-07-26 (`CLAIM-DFLASH-D0D1`, [dflash-spec-decode §0](specs/dflash-spec-decode.md)): `SPEC-DFLASH` UNBLOCKED + `ACTIVE` on the advanced pin `555967922`/vLLM 0.26.0.dev0.** The prior 0.25.0 ORACLE-BLOCKED verdict is SUPERSEDED — under `VLLM_USE_V2_MODEL_RUNNER=1` (vllm#40898 resolved) the mixed-SWA/full z-lab 27B draft CONSTRUCTS and the drafter is ALIVE (acceptance 2.21/8.80/4.75/4.57 > 1, `num_spec=16`, flashinfer-native fp8-KV; goldens committed). Gate FORM measured STRICT MODE-MATCHED (vLLM-ON run-deterministic K>=3 but != vLLM-OFF at k=16 near-ties — NOT the MTP three-way identity). **D1 `DF-AUX-TAPS` DONE:** the single hidden tap is generalized to the multi-tap `[T,H×taps]` (`ForwardDeviceMultiTap` capturing `(hidden+res)` at `target_layer_ids`), config-gated byte-identical off; unit gate 598 assertions (RED-first), CUDA 697/697 + sanitizer 0, 27B MTP e2e 9/9 + 27B SACRED 235/235 byte-identical (inertness). **D2-D5 LANDED 2026-07-26 (`CLAIM-DFLASH-D2`/`D3`/`D4D5`/`D5`):** the drafter model + the project's first non-causal in-block attention (D2, GPU parity vs the real vLLM draft), context-KV precompute + `prepare_dflash_inputs` (D3, GPU numeric-parity 61/61), the non-autoregressive whole-block propose brick + `dflash` config-select (D4), and the RUNNER-LOOP INTEGRATION + 27B e2e (D5): the full verify/propose loop is wired (separate z-lab draft load + target-shared bf16 embed/lm_head, aux-tap capture, per-request combined-feature context accumulation honoring num_rejected, `propose_drafts_dflash`) and RUNS end to end - `test_qwen27_dflash_spec_decode` 2/4 STRICT token-exact vs the vLLM-DFlash-ON golden + acceptance ~ vLLM on ALL 4 (19/39/29/25 vs 17/39/30/25). The 2 divergences are SINGLE bf16 near-tie flips (ratified near-tie ROOT = the D3 inline context-KV recompute envelope), NOT a wiring bug; inertness SACRED 235/235 + MTP 9/9 byte-identical; CUDA `-Werror` clean, no new kernel. NOT a clean strict-4/4 pass - STRICT 4/4 token-identity + the speed A/B = D6 (persistent paged draft-KV bit-matching vLLM's fused projections + the uniform-1+k FULL CG). Capture tool + goldens: `scripts/spec/d{0,2,3}_dflash_*.py`, `tests/parity/goldens/dflash_27b{,_draft,_kvprep}/`. **D6-D9 SPEED CAMPAIGN 2026-07-27:** D6 c1 A/B + STRICT-4/4 bf16-irreducibility RCA; D7 device-resident within-step forward (bit-identical); D8 acceptance RCA + FINAL golden A/B (ours 0.69× vLLM). **D9 (`CLAIM-DFLASH-D9`) PERSISTENT PAGED DRAFT-KV LANDED (bit-identical, +22.7%): `AppendContextKVHost` + `ForwardBlockLogitsWithPrecomputedKV` replace the O(context²) per-step recompute with an append-only per-request store; ours-ON 20.99→25.75 tok/s = 0.917× vLLM-ON (28.09, was 0.69×); e2e 27/27 SAME tokens, SACRED 235/235 + MTP 9/9 byte-identical, CUDA `-Werror` clean, no new kernel. D8's "bf16 acceptance ceiling" REFUTED — same-trajectory per-step acceptance == vLLM (ratio 1.00) and ours realized acceptance (3.68/step) > vLLM (3.31); the SOLE residual (~8%) is the FULL uniform-(1+k) CUDA graph (eager-vs-graphed), a closeable increment. SPEC-DFLASH stays `ACTIVE` (speed not yet ≥ vLLM).** **SPEC-MTP → `DONE` 2026-07-26 (`CLAIM-SPEC-MTP-DONE`, records-only, closing commit I7 `72f9fb1`):** the user ratified the c>1 near-tie+SPEED criterion, closing both I6-owed items (mixed-batch concurrency + server/CLI/C-ABI `--speculative-config`); MTP k=1 is COMPLETE + gated. **M-mtp-2 CLOSED 2026-07-26 (`CLAIM-SPEC-MTP-M-MTP-2`): the 35B `Qwen3_5MoeMTP` full e2e three-way token gate PASSES** — our spec-ON == our spec-OFF == vLLM 0.25.0 `--speculative-config mtp` greedy == vLLM spec-OFF, 16/16 vs the `greedy_ids` anchor (STRICT, c1), acceptance 16/16 both sides; c1 spec-ON 1.19x TPOT / +16.3% output-tput vs spec-OFF (acceptance 0.908) — MoE speedup transfers; spec-OFF byte-identical (test+docs-only). MTP is now `DONE` on BOTH gate models (`MODEL-SPEC-qwen3-5-mtp-qwen3-5-moe-mtp` `GATING`→`DONE`). **DFlash D11+D12 2026-07-27 — the FULL uniform-(1+k) CUDA graph is being built in three parts:** D11 (`CLAIM-DFLASH-D11`) landed Part A (the device-store primitive, CPU-gated); **D12 (`CLAIM-DFLASH-D12`) landed A-wire (the D11 device store is now the PRODUCTION path; GPU-gated e2e 27/27 all-exact acceptance 19/39/29/25 + SACRED 235/235 + MTP 9/9 byte-identical) + Part B (`vt::DFlashPagedBlockAttention`, the capture-safe paged kernel; `test_ops_dflash_paged_block_attn` 795648/795648 CPU==CUDA + cross-check vs materialized `DFlashBlockAttention` + compute-sanitizer 0; NO function-local host `cu_seqlens` upload = capture-UAF fixed).** Speed UNCHANGED 0.917× (A-wire eager + Part B not yet wired). The SOLE remaining piece is Part C (static-shape capture + device mask-scatter + `BeginCapture`/replay + the ≥vLLM c1 A/B); if ours-ON-graphed ≥ vLLM-ON → SPEC-DFLASH DONE → C3 complete. C3 stays `ACTIVE` (DFlash Part C + DSpark/TLI remain) | | 4 | `ROAD-V1-C4` | Quantization: llama.cpp breadth/speed, NVFP4/FP8/MX, MLX native | [quantization matrix](quantization-matrix.md) | coverage spike merged; `QUANT-GGUF-CPU-THREADPOOL` W1-W3 implemented and correctness-gated, still `GATING` (its reproduction now exists — same-binary 1-vs-20-thread A/B is prefill 12.47x / decode 8.05x / RSS 1.000x, so **decode misses the >=10x bar**). **GGUF COMPUTE-IN-QUANT IS NOW LIVE AND DEFAULT-ON (2026-07-22, `CLAIM-QUANT-GGUF-CIQ-G4-1`):** [compute-in-quant GEMM](specs/gguf-compute-in-quant-gemm.md) **G1-G4** — block dtypes + traits, the Q8_0/Q8_K activation quantizers, the six generic `vec_dot`, `kMatmulBTQuant`, and now the ROUTING (`vt::MatmulBT` dispatches a block-dtype weight to it) — plus [keep-quant loader](specs/gguf-keep-quant-loader.md) **L1-L4**, whose master switch defaults ON wherever that op is registered for the running device (CPU today; a CUDA build still expands). Six encodings (Q4_0, Q8_0, Q3_K, Q4_K, Q5_K, Q6_K) now carry `C` = `Y`. **Correctness held exactly**: the 35B GGUF gate is 16/16 token-exact vs the same-file llama.cpp oracle with the quant path on, and the bench model's output tokens are byte-identical across the pre-G4, post-G4 and `VT_CPU_REF=1` arms — no golden regenerated. **Binding CPU A/B** (idle dgx aarch64, one flock, same binary, 3 reps): decode **3.45x**, prefill **4.16x**, peak RSS **1.16x less**; vs llama.cpp we went from 11.7x / 34.1x / 2.66x behind to **3.38x / 8.20x / 2.29x**. The projected 9-17x did **not** hold, for a measured reason: 60 % of that file's weight bytes are `f16`, which no block encoding covers. Keep-quant loader **L4** is therefore MEASURED-and-NOT-MET on RSS; other leaf specs open. **THAT #1 LEVER IS NOW LANDED (2026-07-22, `CLAIM-KERNEL-CPU-ELEM-GEMM-1`, new row [`KERNEL-GEMM-CPU-ELEM`](kernel-matrix.md)):** [the elementwise CPU GEMM](specs/cpu-elementwise-gemm.md) **E1-E4** — per-dtype specialization out of the K loop, 16 independent accumulators instead of one, AArch64 NEON + x86-64 SSE2/F16C tiers behind a runtime probe, and M-blocking — all **BYTE-IDENTICAL** to the historical kernel (`memcmp` gate, exhaustive 65,536-pattern widening check, same token md5), so nothing was regenerated. Binding same-binary A/B: prefill **3.41x**, decode **3.11x**; op-level bf16 18-24 -> 69-351 GFLOP/s. **vs llama.cpp: decode AT PARITY (1.03x), prefill 2.34x behind, RSS 2.29x worse. THEN loader L5 LANDED (2026-07-23, `CLAIM-QUANT-GGUF-KEEPQ-L5-1`):** [keep-quant loader](specs/gguf-keep-quant-loader.md) **L5** — mmap in-place residency (borrow kept blocks out of the read-only mapping, refcounted), tied-head sharing (one bf16 vocab matrix for embed+lm_head), and a read-once page release (port of llama.cpp `unmap_fragment`) — took **peak RSS 6.401 -> 3.884 GiB (2.29x -> 1.39x llama.cpp)** with decode UNCHANGED and output tokens byte-identical (md5 `d235db12f2cd304007530286a1755c95`). The remaining ~1.09 GiB over llama.cpp is the f16 expansion (no block encoding covers f16). | `PARTIAL` | **THE OWED FRESH PROFILE IS DONE (2026-07-23) and it re-ranks the plan.** A `vt::GetOp` hook (100% of wall time, reverted before binding) on the CURRENT binary: prefill is **no longer GEMM-bound** — kMatmulBTQuant 37%, **kGdnPrefill 25%**, kMatmul 12%, kMatmulBT 10%, **kPagedAttention 10%**; the two non-GEMM kernels (GDN linear-attention recurrence + paged attention) run **SINGLE-THREADED** on the CPU and are now the top prefill levers. Decode is memory-bound matmul at parity, no kernel work owed. **RE-RANK: G5/G6/G7 all only speed the already-fast quant GEMM and rank BELOW the two serial non-GEMM kernels; the new #1 CPU lever is threading kGdnPrefill + kPagedAttention.** **THAT #1 LEVER IS NOW LANDED (2026-07-23, `CLAIM-CPU-THREAD-GDN-PAGED-1`, [two-kernel threading](specs/cpu-thread-gdn-paged-2026-07-23.md)):** kGdnPrefill chunks over the (sequence, value-head) axis and kPagedAttention over query-token rows, both via the existing `ParallelForRows`, both **byte-identical** (qwen35 output-token md5 `d235db12f2cd304007530286a1755c95` unchanged at threads 1/4/20 + `VT_CPU_REF=1`, determinism battery extended, CPU ctest 158/158). **Binding dgx aarch64 (idle): prefill 1.382× same-binary (73.0→100.9 t/s), 2.43×→1.76× behind llama.cpp pp128; decode at parity; op-scaling 1→20 GdnPrefill 7.08× / PagedAttention 8.96×; fresh profile shows the two kernels 35%→8.6% of prefill and re-ranks the NEW bottleneck to the GEMMs (kMatmulBTQuant 50% + kMatmul 16% + kMatmulBT 14% = 80%) ⇒ next CPU lever is the SIMD/repack GEMM tiers (G5/G6/G7).** **THE FIRST SUCH TIER IS NOW LANDED (2026-07-23, `CLAIM-QUANT-GGUF-CIQ-G6-1`, [compute-in-quant GEMM](specs/gguf-compute-in-quant-gemm.md) G6):** the Arm **i8mm mmla `nrc==2`** `vec_dot` tier for q8_0/q4_0/q4_K/q6_K (q3_K/q5_K have no upstream mmla → stay portable), 2x2-tiled into `kMatmulBTQuant` at even M,N (decode M=1 → portable, unchanged); runtime `HWCAP2_I8MM` probe + `VT_CPU_QUANT_MMLA` defeat + per-file `+i8mm`. **BYTE-IDENTICAL** where the math allows (q8_0/q4_0 bit-exact to the scalar tier, q4_K/q6_K within NMSE ≤ 5e-4), bit-identical across threads 1/2/4/20, e2e token md5 `d235db12f2cd304007530286a1755c95` byte-identical (mmla on/off/`VT_CPU_REF=1`), 35B GGUF gate 16/16 vs llama.cpp on both files. **Op-level portable→i8mm: q4_K 7–8.4×, q6_K 3.8–4.5×, q8_0 ~1.2×**; e2e prefill same-binary **1.084×** on the q8_0-dominant bench file (1.56×→**1.44× behind** llama.cpp pp128, Amdahl-bounded — the big k-quant win lands on the APEX 35B files). Fresh bottleneck: the elementwise f16/f32 GEMM (~30%, unchanged) is now co-dominant on this mixed file. CUDA `-Werror` 0-warn, regression set UNCHANGED. `docs/BENCHMARKS.md` ACCEPTED. RSS deficit closed to 1.39x by L5; the last RSS lever is an f16 keep-as-is compute path, not this loader. **THEN the GDN split-projection orientation LANDED (2026-07-23, `CLAIM-CPU-GDN-ORIENT-1`, [GDN projection orientation](specs/cpu-gdn-proj-orientation-2026-07-23.md)):** a fresh op-dispatch profile of the current binary (warm prefill, `vt::GetOp` hook + per-GEMM shape histogram, reverted before binding) found the four GDN input projections (`in_proj_qkv/z/b/a`, 72 GEMMs, **17.9%** of prefill: `kMatmulBTQuant 50.7% / kMatmul 17.9% / kMatmulBT 14.9%`) were the LAST weight family `LoadGdnGguf` still transposed into [K,N] (nk=false → the N-striding `kMatmul`, no M-blocking) after G4's `expand_nk` gave every other expanded weight the file's own [N,K] order. New `GgufLoadPolicy::gdn_expand_nk` + `MakeGdnProj` keep them [N,K] nk=true → the M-blocked `kMatmulBT`; **BYTE-IDENTICAL** (same sequential f32 K-reduction, only the weight offset differs — token md5 `d235db12f2cd304007530286a1755c95` unchanged across on/`VT_GGUF_GDN_NK=0`/`VT_CPU_REF=1` and threads 1/4/20), `test_qwen36_gguf_engine` 2/2·28/28·16/16 on APEX. **Binding same-binary prefill 1.090× / decode 1.09× (44.1→40.4 ms TPOT = 1.01× llama tg32, at parity), 1.44×→1.32× behind llama.cpp pp128, RSS unchanged.** Fresh post-change profile: `kMatmul` **17.9%→0% (72→0 calls, ELIMINATED)**, absorbed into `kMatmulBT` (14.9%→27.7%); **next CPU prefill lever = the quant GEMM (kMatmulBTQuant, now 55%): G7 repack-at-load.** **G7 LANDED 2026-07-23 (`CLAIM-QUANT-GGUF-CIQ-G7-1`, [compute-in-quant GEMM](specs/gguf-compute-in-quant-gemm.md) G7):** q8_0 repacked once at load into the i8mm `block_q8_0x4` interleave (ported llama.cpp `repack.cpp` `q8_0_4x8`), `kMatmulBTQuant` dispatches a pre-shuffled gemm/gemv with no per-block register shuffles. **BIT-IDENTICAL** (byte-permute weight + non-fused `vmlaq_f32`, 305-assertion memcmp across decode/prefill/bf16-out/threads, token md5 `d235db12f2cd304007530286a1755c95` unchanged on/`VT_CPU_QUANT_REPACK=0`/`VT_CPU_REF=1`; a `ResidentWeight`/`MakeTensor` flag-drop that produced all-zero tokens was caught by the E2E gate and fixed). Op-level q8_0 **3.7–5.9×** (518→2401 / 583→3456 / 514→1902 GFLOP/s); **E2E prefill 1.92× same-binary (1096→572 ms), 223.8 t/s vs llama.cpp pp128 177.3 = 1.26× — AT/BEYOND PARITY** (was ~1.5× behind), decode at parity, RSS 3.884 GiB unchanged. Fresh profile: q8_0 GEMM 55%→~21%; **the CPU prefill-lever search is CLOSED — the sole remaining gap to llama.cpp is peak RSS (1.39×), not prefill.** CUDA-inert (gated off any non-CPU-quant device), CUDA `-Werror` 0-warn, regression set UNCHANGED. **THE RSS GAP IS NOW CORRECTLY ATTRIBUTED (2026-07-23, `CLAIM-QUANT-GGUF-KEEPF16-L6-1`, [keep-quant loader](specs/gguf-keep-quant-loader.md) L6): it is NOT the f16 expansion.** L6 implemented keep-f16 residency (keep the file's 56 F16 weights + tied head resident as F16 and compute on them, mirroring llama.cpp `ggml_vec_dot_f16`) and MEASURED it **RSS-NEUTRAL** (3.884 → 3.832 GiB, −52 MB) and prefill-regressive (TTFT 577 → ~1000 ms, from 1.25× ahead of llama.cpp to 0.72× behind) — because L5's page-release had ALREADY dropped the f16 file pages, so keep-f16 only swaps an anonymous bf16 buffer for equal-size file-backed f16 pages. smaps attribution proves our weight residency is at llama.cpp parity (file-backed 2.63 ≈ 2.68 GiB); **the residual ~1.08 GiB is the engine's ANONYMOUS activation/KV workspace, not weights — the real, separate CPU RSS lever.** keep-f16 ships DEFAULT OFF (`VT_GGUF_KEEP_F16=1` opt-in), tokens byte-identical (md5 `d235db1…`), `test_gguf_keep_quant` 35/35 (x86+aarch64), regressions UNCHANGED (27B 235/235, 35B 315/315, Coder 138, dense 184, OPT 63, DeepSeek 223, Llama 92, GGUF engine 28/28). **NEXT CPU RSS lever: profile + shrink the engine's activation/KV working set, NOT the weight loader** | | 5 | `ROAD-V1-C5` | Sliding window, local attention, YaRN/long context | [engine matrix](engine-matrix.md), [coverage view §§2,11](feature-matrix.md#2-kv-cache--memory), [joint spike](specs/sliding-local-yarn-long-context.md) | **CUDA GPU CLOSURE 2026-07-27 (`CLAIM-ROADMAP-C5`, dgx GB10 sm_121a, clean build of `489f7771`, oracle vLLM 0.26.0.dev0):** the shared scaled-RoPE + local-mask CUDA path COMPILES `-Werror`-clean and RUNS on GB10; the C5 feature-positive correctness gates that were the stated `GATING` blocker now PASS — SWA (Gemma-2/Gemma-3 48/48), LongRoPE (Phi-4-mini 16/16, RED-first), llama3-rope (Llama-3.2-1B 16/16), dynamic-NTK (InternLM2 16/16); both RoPE 0.26-oracle recaptures BIT-IDENTICAL to goldens (zero drift). Leaves `ATTN-SLIDING-WINDOW`/`ATTN-ROPE-{LLAMA3,LONGROPE,DYNAMIC-NTK}`/`ATTN-YARN` → `ACTIVE` | `PARTIAL` | (RI) **Honest residual (vehicle-blocked, not skipped):** YaRN model e2e (no cached Nomic/gpt-oss consumer) + chunked-local model e2e (no Llama4 row) are REACHABLE-BLOCKED — operator/formula stay GPU/G3-gated; long-context positive-mask (prompt > W) SWA model e2e + the KV memory-optimization G8; and the roadmap-wide every-axis SPEED tail (all C5 leaves correctness-complete, speed-pending, mirroring their model consumers). Not row-DONE until speed + the blocked vehicles close | diff --git a/.agents/specs/cli-serve-bench.md b/.agents/specs/cli-serve-bench.md new file mode 100644 index 000000000..c02e0aea2 --- /dev/null +++ b/.agents/specs/cli-serve-bench.md @@ -0,0 +1,127 @@ +# CLI serving/benchmark parity — structured spike + +Stable row: `SERVE-CLI-BENCH`. This spike covers only the benchmark half of +the row: making the local closed-loop client exercise the same production +engine mode as the pinned vLLM denominator. The server CLI breadth remains the +existing `PARTIAL` residual. + +## Problem and measured discriminator + +`vllm-bench` constructs `LoadedEngine`, whose default configuration correctly +resolves asynchronous scheduling ON and `max_concurrent_batches=2`, but then +calls `loaded->engine()` and drives `LLMEngine::step()` +(`examples/bench/bench_core.h:485-512`). That path calls +`EngineCore::step()` directly (`src/vllm/v1/engine/llm_engine.cpp:141-143`) and +therefore never selects `EngineCore::step_with_batch_queue()`. + +The pinned vLLM engine makes the dispatch in its engine core: it sizes the +batch queue from `max_concurrent_batches` and selects +`step_with_batch_queue` whenever the queue exists +(`${VLLM_SOURCE}/vllm/v1/engine/core.py:200-231`). The queued path +schedules a new batch before consuming the oldest result +(`${VLLM_SOURCE}/vllm/v1/engine/core.py:622-669`). Our production +equivalent is already implemented behind `LoadedEngine::async_engine()`: +`AsyncLLM` constructs `EngineCoreProc` with the resolved queue depth +(`src/vllm/entrypoints/model_loader.cpp:800-811`), and `EngineCoreProc` selects +`step_with_batch_queue` at depth greater than one +(`src/vllm/v1/engine/core_proc.cpp:22-37`). + +The 2026-08-07 fresh RTX 5070 Ti run at `7ef5f1001` exposed the consequence. +The exact c32 workload completed the same 148,168 tokens in about 22.2 seconds +on both engines, but same-tool graph-node traces showed different sampling +waves: ours ran 529 sample steps, overwhelmingly batch 32; vLLM ran 650 steps +with a mean sampled batch near 25.3. Total steady GPU-kernel time was almost +equal (about 21.7 versus 22.0 seconds), while our client-observed mean TPOT was +38.133 ms versus 33.916 ms. These numbers are a structural discriminator, not +a binding parity result, because the frontends were mismatched. + +## Upstream and dependency chain + +| Layer | Pinned vLLM / dependency behavior | Our anchor | +|---|---|---| +| CLI/client | `LLM.llm_engine.add_request` + `engine.step`, DELTA output (`tools/bench/vllm_closed_loop_metrics.py:59-102`) | closed-loop admission and metrics (`examples/bench/bench_core.h:416-589`) | +| engine dispatch | queue size and `step_fn` selection (`vllm/v1/engine/core.py:200-231`) | `LoadedEngine::async_engine` passes resolved depth (`src/vllm/entrypoints/model_loader.cpp:800-811`) | +| batch queue | schedule-before-oldest-result (`vllm/v1/engine/core.py:622-669`) | `EngineCore::step_with_batch_queue` (`src/vllm/v1/engine/core.cpp:115-185`) | +| scheduler | `AsyncScheduler` placeholder accounting (`vllm/v1/core/sched/async_scheduler.py`) | `src/vllm/v1/core/sched/async_scheduler.cpp` | +| runner/sample | non-blocking sampled-token copy and device input update (`vllm/v1/worker/gpu/async_utils.py`, `gpu_model_runner.py`) | `GPUModelRunner::sample_tokens_async` and device mirror (`src/vllm/v1/worker/gpu/runner.cpp`) | +| CUDA graph | graph replay is selected inside the runner; tracing must expose child nodes | existing `--cuda-graph-trace=node` paired `nsys` recipe | + +No new kernel is proposed. CUTLASS, cuBLASLt, FlashAttention and GDN calls +must remain byte-identical; their aggregate and per-template times are the +post-change structural control. + +## Dispatch and implementation + +1. `RunBench` obtains `loaded->async_engine()` for every non-legacy benchmark + run. This keeps async-scheduling OFF meaningful: the same frontend then + drives a depth-1 `EngineCoreProc`, matching vLLM's `async_scheduling=False` + control rather than silently switching frontend APIs. +2. Admission stays deterministic and sequential. Keep a request-id-ordered map + of `AsyncRequest` collectors, scan them with `get_output_nowait()`, process + every ready DELTA, and refill immediately after terminal outputs. Yield only + when no collector is ready. Do not introduce one client thread per request; + racing initial submissions would change batch composition and invalidate the + same-workload claim. +3. Report the selected frontend, resolved async flag and batch-queue depth in + `BenchResult`/CLI output so a benchmark artifact proves which path ran. +4. Retain the existing synchronous `LLMEngine` as a library API and test seam; + only the comparison harness changes. + +## Files and tests + +- Modify `examples/bench/bench_core.h` and `examples/bench/main.cpp`. +- Extend `tests/examples/test_bench.cpp` to assert the async frontend was + exercised, all requests/tokens remain complete, and serialized output order + remains submission order. +- Reuse the production depth-2 correctness gates + `tests/parity/test_qwen3_dense_async_serving.cpp` and + `tests/vllm/v1/test_engine_core_proc.cpp`; no upstream test module directly + covers this original C++ benchmark, so there is no omitted upstream test to + port. +- Re-run `tests/vllm/models/test_qwen35_plain_weights.cpp` on the cached 4B + model before timing. + +RED-first mutation: replace `loaded->async_engine()` with `loaded->engine()` or +force the reported frontend false; the new benchmark contract test must fail. + +## Gates and hardware + +CPU gate: focused benchmark, async engine/core, scheduler and output-processor +tests, followed by the staged protocol gates. CUDA correctness: cached +Qwen3.5-4B plain-weight test, direct ON/OFF output IDs 128/128 per repetition, +and no new sanitizer finding where sanitizer support exists. + +Performance gate on the local RTX 5070 Ti (`sm_120`): hold `${GPU_LOCK}` across +the entire interleaved direct-ON / pinned-vLLM / direct-OFF series; run under the +validated user-systemd scope (`MemoryHigh=22G`, `MemoryMax=25G`, swap disabled); +three memory and three timed repetitions per arm; exact cached model, ShareGPT +digest, 128 requests, 128 output tokens, c32, 2048 batched-token cap, greedy. +Trace both engines with the same `nsys` and `--cuda-graph-trace=node`. + +Acceptance is token correctness plus no regression on any axis. The corrected +frontend materially changed the queue composition, so its production run +supersedes the diagnostic TPOT target. Paired same-tool traces selected batched +GDN prefill causal-conv total GPU time; the implemented exact-chunk lever and +its enclosing result are in +[the sm_120 chunking spike](sm120-qwen35-conv-chunking-2026-08-07.md). +Aggregate throughput, TTFT, TPOT/ITL, peak/stable PSS and VRAM remain guards. + +## Dependencies, risks, rollback, work breakdown + +Dependencies are already landed: `SERVE-ASYNC-LLM`, `ENG-CORE-BUSY-LOOP` and +`ENG-ASYNC-SCHED`. The cached model and pinned oracle are available locally. + +Risks: collector DELTA coalescing could distort ITL if the client fails to drain +promptly; the gate therefore checks output chunking and TPOT independently. +Busy polling could contend with the engine thread; the loop yields on no work +and CPU utilization is recorded. Async production correctness is protected by +the existing token-exact tests. Rollback is the single harness dispatch change; +the synchronous engine remains untouched. + +| Work item | Lifecycle / result | +|---|---| +| B0 fresh-main build, correctness, 18-leg diagnostic and paired trace | complete; non-binding because frontend mismatch discovered | +| B1 async-front-end benchmark repair + RED-first contract test | complete: `RunBench` uses `AsyncLLM`; report records frontend/flag/depth | +| B2 contained rebuild and focused correctness | complete: RED compile failure; GREEN build; benchmark 4/4·35, async 8/8·323, core 10/10·95, cached 4B 3/3·1672; real-model async smoke | +| B3 exact interleaved rerun + same-tool traces | complete: production c32 baseline bound; causal-conv 720.954/145.421 ms selected | +| B4 choose/implement first measured lever | complete and current-main revalidated: exact chunks 3.073x kernel / +2.121% profiled enclosing; three-pair binding A/B +2.152% | diff --git a/.agents/specs/sm120-qwen35-conv-chunking-2026-08-07.md b/.agents/specs/sm120-qwen35-conv-chunking-2026-08-07.md new file mode 100644 index 000000000..426246d94 --- /dev/null +++ b/.agents/specs/sm120-qwen35-conv-chunking-2026-08-07.md @@ -0,0 +1,152 @@ +# sm_120 Qwen3.5 batched prefill-conv chunking — structured spike + +**Rows:** `KERNEL-SSM-MAMBA`, feeding `ROAD-V1-C2-LOCAL-BF16`. +**Hardware/workload:** RTX 5070 Ti (`sm_120`), Qwen3.5-4B plain BF16, +128 ShareGPT requests, 128 output tokens, concurrency 32, +`max_num_batched_tokens=2048`, greedy. **Lifecycle:** implemented, locally +gated and default ON; 27B/35B release gates remain hardware-unavailable. + +## Measured selection + +The corrected production-frontend comparison at `3f35356e0` completed three +memory and three timed repetitions per direct-ON, pinned-vLLM and direct-OFF +arm under one `/tmp/gpu` lock. Direct ON and OFF are loader controls; both +reported `AsyncLLM`, async scheduling enabled and maximum concurrent batches 2. + +Direct ON versus pinned vLLM means are: + +| Axis | ours | vLLM | ours / vLLM | +|---|---:|---:|---:| +| total throughput | 6633.750 tok/s | 6643.593 tok/s | 0.998518x | +| output throughput | 733.540 tok/s | 734.630 tok/s | 0.998516x | +| mean TTFT | 1051.333 ms | 937.584 ms | 1.121322x | +| mean TPOT / ITL | 35.457 ms | 33.906 ms | 1.045721x | +| peak VRAM | 13054 MiB | 12820 MiB | 1.018253x | + +The same-tool graph-node traces isolate request processing to 22.370 seconds +locally and 22.269 seconds in vLLM. Both are saturated (98.54% and 99.34% GPU +busy by interval union), so a host-only polling change is not the first lever. +Sampling covers essentially identical work: 16,442 local versus 16,443 vLLM +rows. The largest clear non-GEMM kernel-family gap is batched GDN prefill +`causal_conv1d`: + +| Trace metric | ours `CausalConv1dFwdRegKernel` | vLLM `_causal_conv1d_fwd_kernel` | +|---|---:|---:| +| launches | 1,728 | 1,893 | +| total GPU time | **720.954 ms** | **145.421 ms** | +| mean launch | 417.219 us | 76.821 us | + +The selected primary micro-metric is therefore **steady-interval total GPU time +for batched GDN prefill `causal_conv1d` on the exact c32 workload**. Baseline +gap: **4.958x**. This metric is preferred over isolated TPOT because the old +synchronous frontend demonstrated that queueing can improve TTFT while making +TPOT worse without improving device work. + +## Whole-chain cause + +Pinned vLLM's Triton kernel takes `batch_ptr` and `token_chunk_offset_ptr` and +maps every program to exactly one `(sequence, BLOCK_M token chunk)` before +processing the feature tile +(`${VLLM_SOURCE}/vllm/model_executor/layers/mamba/ops/causal_conv1d.py:15-28,71-79,123-124`). +The scheduler/backend metadata constructs those exact chunk descriptors. + +Our register kernel has the same register-window arithmetic, but +`kConvRegChunkMaxSeqs=4`; `LaunchConvFwdReg` sets `gridZ=1` whenever the batch +contains more than four sequences, so each channel block serially walks an +entire request (`src/vt/cuda/cuda_gdn.cu:702-705,720-766,793-819`). The trace +shows that production c32 path directly: the dominant local shapes are +`grid=(64,28..32,1)` at roughly 410-466 us, while vLLM uses exact token-chunk +programs with feature-grid 32 at roughly 82-95 us for the large waves. + +The missing data is already an explicit recorded deviation: +`GDNAttentionMetadata` omits vLLM's `batch_ptr` and +`token_chunk_offset_ptr` because the original sequential C++ kernel had no +consumer (`include/vllm/v1/attention/backends/gdn_attn.h:33-43`). That +assumption is now refuted for performance on the production batched path. + +## Port, not reinvention + +Implement vLLM's exact chunk list through the existing GDN metadata and device +input path, then make `CausalConv1dFwdRegKernel` consume one descriptor per +program. Do not launch a rectangular `num_sequences * ceil(total_tokens/M)` +grid: it over-launches by the batch size and is not the upstream algorithm. +Preserve the current tap-order float accumulation, state read/write semantics, +BF16 I/O, and `VT_CONV_REG=0` rollback. Tune `BLOCK_M`/feature tile only after +the exact upstream work partition is measured; the first discriminator is the +metadata/dispatch port, not an arbitrary sm_120-only kernel. + +No GEMM claim is made from this trace. The GEMM templates differ between the +engines, so any later GEMM lever must separately satisfy the four-axis +invocation-parity gate (C dtype, compute/scale type, entry point/algo policy and +resolved same-tool template). + +## Tests and acceptance + +RED-first tests must cover: + +1. host metadata for unequal sequence lengths produces every `(sequence, + token-chunk)` exactly once, no missing or padded work; +2. CUDA chunked versus current register path is byte-identical for output and + final convolution state over BF16/F32 inputs, initial/fresh state, unequal + lengths, `T < K-1`, and batches above four sequences; +3. a mutant restoring `gridZ=1` above four sequences fails the structural + launch/metadata test; +4. cached Qwen3.5-4B correctness remains 3/3 cases and 1672/1672 assertions; + direct ON/OFF remains 128/128 identical per repetition. + +Performance is measured first as a same-binary, same-workload graph-node trace +under the 25 GiB user-systemd cap and one `${GPU_LOCK}`. The micro-metric must +improve outside run noise and move toward vLLM's 145.421 ms; a default flip also +requires the project's kernel-efficiency bar and byte-exactness. Then repeat +the full 18-leg comparison. No end-to-end axis may regress: total/output +throughput must reach or exceed vLLM, all latency axes must move toward or beat +vLLM, and peak VRAM must not exceed the 13054 MiB baseline or vLLM floor. The +27B and 35B gate-model correctness suites remain required before a shared +default change; this 4B run cannot extrapolate support to them. + +## Evidence and rollback + +- benchmark aggregate: `/tmp/qwen35-async-3f35356e0/aggregate.json`, SHA-256 + `d006d6ffd6d014fc1861a30c126f603117a1ff80b339172d63fab75f2ae07f1d` +- local trace: `/tmp/qwen35-async-3f35356e0-ours.nsys-rep`, SHA-256 + `182426e961ccaa53896e7cb7e9c4d2ba1cc35ba695ea3a9c3c727195326c8dab` +- vLLM trace: `/tmp/qwen35-async-3f35356e0-vllm.nsys-rep`, SHA-256 + `9d543390741629fc3873ad857706595cf880ee2fcce63af3e4be17e1a343355a` +- full record: [Qwen3.5-4B sm_120 evidence](../../docs/bench-evidence/qwen35-4b-sm120-main-20260807.md) + +Rollback keeps the existing `VT_CONV_REG=0` tiled/scalar path and adds a +same-binary switch for the new exact-chunk dispatch until the full gate closes. + +## Implementation and measured disposition + +The implementation follows the spike exactly: `ComputeCausalConv1dMetadata` +enumerates every `ceil(sequence_length / 8)` work item, the step metadata owns +and uploads the two i32 descriptor arrays once, and every GDN layer reuses them. +The register kernel consumes one descriptor per `grid.y` program. The +same-binary control is `VT_CONV_EXACT_CHUNKS`; it defaults ON and `=0` selects +the prior whole-sequence mapping. `VT_CONV_REG=0` remains the independent +tiled/scalar rollback. + +RED-first metadata coverage, affected model fixtures, full CUDA GDN tests and +cached Qwen3.5-4B correctness are green; default and rollback output are byte +identical in three production pairs. Same-binary `nsys` reduces causal-conv +GPU time from 720.047 to 234.607 ms (**3.069x**), leaving a **1.613x** gap to +the sealed pinned-vLLM 145.421 ms trace. Three alternating enclosing pairs +improve total/output throughput **2.152%**, TTFT **2.945%**, TPOT/ITL **1.920%** +and E2E latency **2.118%**, with no VRAM regression. The local default now +measures **1.021246x** the sealed vLLM throughput; TTFT (**1.085812x**) and +TPOT/ITL (**1.024597x**) remain slower, and VRAM remains above the vLLM floor. + +The attempted fresh 18-leg oracle reruns are VOID because the oracle's +FlashInfer/Torch/Triton JIT environment became unstable after 13/18 legs. The +accepted evidence is the same-binary trace plus three-pair local A/B, compared +to the sealed same-hardware/same-workload vLLM denominator. See the linked full +record for manifests and exact artifacts. No 4B result is extrapolated to the +unavailable release-gate models. + +The clean transplant onto `upstream/main` `f91a5917a` is revalidated: focused +CPU 6/6, CUDA GDN 66/66·4300, cached 4B 3/3·1672, and byte-identical same-binary +token files. Its fresh graph-node profile reproduces causal-conv +720.216507→234.379395 ms (**3.072866x**) and the profiled enclosing run +6587.66→6727.35 tok/s (**+2.1205%**). The old-branch gain therefore survives +the clean transplant; the sealed vLLM conv residual is **1.611730x**. diff --git a/.agents/state.md b/.agents/state.md index c54fa655a..f865e09f0 100644 --- a/.agents/state.md +++ b/.agents/state.md @@ -42374,6 +42374,38 @@ leaves (Kimi runner fold #279, Parakeet ASR #280). Reviewer findings 1-8 applied tests, ratchet equality pin, reachable-row-removal design note, meta-gap note). No CUDA build; no perf number owed; STATUS inside its char ratchet. +## 2026-08-07T21:41 — exact GDN chunks transplanted onto current main + + + +`ROAD-V1-C2-LOCAL-BF16` / `SERVE-CLI-BENCH` / `KERNEL-SSM-MAMBA`, clean +branch `row/KERNEL-SSM-MAMBA-EXACT-CHUNKS` from `upstream/main` `f91a5917a`. + +- **Clean transplant.** The unrelated profile-aware agent-efficiency commit and + five intermediate measurement-only commits are absent. This branch carries + the production-`AsyncLLM` benchmark correction used by the measurement, its + committed spike, the exact-chunk kernel change, tests, final evidence and the + current-main keyed-record reconciliation. +- **Upstream partition ported.** `GDNAttentionMetadata` constructs exact + `(sequence, 8-token chunk)` descriptors once per step, uploads once and + shares them across GDN layers. `VT_CONV_EXACT_CHUNKS=0` restores the legacy + mapping; `VT_CONV_REG=0` remains tiled/scalar. Default is ON. +- **Measured source result.** RED-first host/CUDA/model gates and cached 4B + 3/3·1672 were green on the source tree. Three production rollback/default + pairs were byte-identical. Same-binary `nsys`: conv **720.047→234.607 ms = + 3.069x**; vLLM 145.421 ms leaves **1.613x**. Enclosing total/output +2.152%, + TTFT -2.945%, TPOT -1.920%, E2E -2.118%; VRAM unchanged. Against sealed + vLLM, throughput **1.021246x PASS**; latency and VRAM remain open. +- **Fresh-oracle caveat.** Two 18-leg attempts were VOID on FlashInfer/Torch/ + Triton JIT infrastructure. The linker search is repaired; the sealed same-box + denominator remains binding. No 4B result extrapolates to 27B/35B. +- **Transplant gate GREEN.** Contained CPU/CUDA rebuild passes; focused CPU + 6/6, full CUDA GDN 66/66·4300, cached Qwen3.5-4B 3/3·1672. Current-main + same-binary graph-node reprofile is token-identical and reproduces conv + **720.217→234.379 ms (3.073x)**. Profiled total throughput + **6587.66→6727.35 tok/s (+2.121%)**, TTFT -3.142%, TPOT -1.849%, E2E -2.088%. + Next: spike the residual tile/register-pressure hypothesis. + ## 2026-08-07 — ARCH-ONE-SURFACE ROW 1: Parakeet ASR folded onto the ONE surface (PR #121) @@ -42982,8 +43014,23 @@ removing only `; #129: SPIKE∅`; the compact clause consumes the existing 279150-character ratchet exactly. `ENG-RELEASE-BINARIES` remains `SPIKE`: there is no archive, runtime, correctness or performance evidence. +## 2026-08-08 — sm_120 exact chunks rebased and revalidated + + +The measured code tree `3d2581551` was one commit above `upstream/main` +`48a54141f`; it was subsequently rebased code-identically onto `b38f78a77` +(the intervening release-binary merge changes only records/checkers/workflow). +Keyed records took each new main wholesale before the exact branch edits were +reapplied. The contained 881-target CUDA build is clean. Focused gates pass +6/6, full CUDA GDN passes 66/66 cases and 4300 assertions, and cached +Qwen3.5-4B passes 3/3 cases and 1672 assertions. -<<<<<<< ours +Fresh same-binary graph-node traces preserve byte-identical token files. +Rollback/exact causal-conv totals are 718.704016/233.954533 ms, **3.07198x**; +the enclosing run is 6589.65→6739.34 tok/s (**+2.272%**) with TTFT, TPOT and +E2E all improving. The result survives the main advance. The sealed vLLM +causal-conv denominator remains 145.421 ms, leaving **1.60881x** open; latency, +VRAM and both hardware-unavailable release-model gates remain open. ## 2026-08-08 — Tensor-parallelism end-to-end spike lands at the current pin (task #287) @@ -43024,37 +43071,6 @@ scoping stays a future spike. Rows moved: none in lifecycle (`BACKEND-DISTRIBUTED-TP`/`PAR-TP` stay `READY`, now pointing at the new spec; `SPEC-DSPARK` stays `INVENTORIED`). Next: dispatch TP-W1 (GroupCoordinator-analog) — CPU-completable, no HW wait. -======= - -## 2026-08-08T21:00 - richiejp stack LANDS: #65 + #79 merged with fix map (prep/land-65, prep/land-79) - - -The external RPi5/Cortex-A76 stack lands as two prepared branches off -`b38f78a7`: `prep/land-65` (PR #65, PMU kernel-bench harness + lane records) -and `prep/land-79` (PR #79, the Q8_0 SDOT/AAPCS64 tier), keyed records merged -by the resolution law (main wholesale + in-lane rows; state pure-append in -anchor order). The mutation review's fix map applied: F65-1 (ONE-SURFACE -allowlist entry for `examples/cpu_kernel_bench`, ratchet 8->9 with a dated -operator exception), F79-5 (KERNEL count 50->51 on top of main, their dated -comment transplanted), F79-6 (the two RPi5 BENCHMARKS.md paragraphs are keyed -table rows pointing at `docs/bench-evidence/rpi5-*.md`; the STATUS.md CPU cell -condensed inside the char/cell ratchets), and F79-1: the new -`test_ops_quant_dot` A76 byte-equality case referenced -`QuantTraits(kQ8_0).vec_dot` — the SELECTED kernel, i.e. the assembly tier on -a real A76 — so its CHECKs were a self-comparison. A true-portable seam -`vt::cpu::QuantQ8PortableVecDot()` (quants.c:400 order) is exported next to -the existing sdot/asm getters, the reference arm retargeted, and a new -seam-pinning case compares the portable kernel byte-equal against an -independent in-test transcription of the exact order on EVERY platform -(contraction-robust: mul and fma candidates). Mutation kill verified on x86: -reversing the portable block accumulation order REDs exactly that case -(2 assertions, blocks=3/64), green on revert; 23/23 cases, 150231 assertions. - -RESIDUALS, recorded not fixed (per the review map, operator-held): F79-3 and -F79-4 remain open on the landed tree; the review's merge-and-fix map is the -binding description. Pi concurrency, BF16 GEMM/speed closure (W6) stay open -as the lane's own next steps. ->>>>>>> theirs ## 2026-08-08T21:00 - richiejp stack LANDS: #65 + #79 merged with fix map (prep/land-65, prep/land-79) @@ -43085,7 +43101,6 @@ F79-4 remain open on the landed tree; the review's merge-and-fix map is the binding description. Pi concurrency, BF16 GEMM/speed closure (W6) stay open as the lane's own next steps. - ## 2026-08-08 — ROCm approach-(b): unified memory true by construction on integrated APUs diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index 4a621e9fc..f5d298809 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -8,7 +8,7 @@ | **Developer agent entry point (implemented)** | `DOCS-AGENT-PROTOCOL-ENTRYPOINT`: public contribution guide + synchronized, mutation-gated pre-claim intake rule | Rebased documentation/protocol only; benchmark void | n/a | | **ARCH audit: ABI is text-only** | 4 capabilities (H3 video, Laguna, Kimi-Linear, DeepSeek-V4) reachable only from `examples/`, none registry-backed. No gate asks whether a CONSUMER can reach a capability. Documentation only | | **`ROAD-V1-MEM` M1+M2 (2026-08-08)** | KV auto-sizing CPU brick: `--kv-cache-memory` sizes the pool from a byte budget via the group-aware `KVBytesPerBlock` divisor (ABI v16, CPU-gated). M3 profile run dgx-gated | -| **Record repair 2026-08-07** | `main` was red on `check-agent-record` + `check-env-doc`, blocking every PR. Dangling `kda-chunk-aot/` link and two undocumented env vars. No behaviour change | +| **Record/checker repair 2026-08-07–08** | Restored red record/env gates; made release AST semantic pins Python 3.12/3.13-stable; recorded merged Gemma-4 MoE as known merged-GEMM drift and closed the stale embeddings claim. No runtime/performance change | | **vLLM** | Qwen3.6-27B NVFP4, GB10 | ahead 4.5% at c1, **tie** at c2 to c32 | identical | | **vLLM** | Qwen3.6-35B-A3B NVFP4, GB10 | 0.93x to 1.03x: ahead at c4, worst c16 0.93x | identical | | **vLLM** | DeepSeek-V2-Lite (MLA), GB10 | 0.86x to 0.95x throughput, TTFT wins at c4/c8 | identical | @@ -32,7 +32,16 @@ The binding comparison. vLLM runs its **production graphed config**, never | Qwen3.6-27B | NVFP4 | 0.25.0 | **115/124** | Effective parity-or-better, two-grid totality | | Qwen3.6-35B-A3B | NVFP4 `modelopt_mixed` | 0.25.0 | 2/18 | 3-rep grid 2026-08-05 @`1ea26427`: 0.93-1.03x (c4 wins), c16 0.93x. Both c16 levers A/B'd NEG: drain event -1.9%, mirror 0.999x. ★ probe found a prod async batch-1 greedy DEGENERATION bug the mirror fixes | | DeepSeek-V2-Lite | bf16 MLA | 0.25.0 | 4/25 | Attributed miss, row stays `ACTIVE` | -| Qwen3.5-4B | bf16 direct-load | 0.26.0.dev0 | 3/9 | 0.9971x throughput after the upstream update; TTFT and host PSS win. TPOT/ITL 1.1244x and VRAM remain open ([evidence](bench-evidence/qwen35-4b-upstream-20260805.md)) | +| Qwen3.5-4B | bf16 direct-load | 0.26.0.dev0 | throughput + host PSS | Exact chunks ON: total **1.021x PASS**; TTFT **1.086x**, TPOT **1.025x**, VRAM **1.018x OPEN**; local A/B **+2.152%** ([evidence](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | + +### GDN prefill causal-convolution by GPU + +| GPU | Workload and basis | vllm.cpp | vLLM | Ratio | Status | +|---|---|---:|---:|---:|---| +| RTX 5070 Ti (`sm_120`) | Qwen3.5-4B BF16, c32, steady-interval total | 233.955 ms | 145.421 ms | **1.609x slower** | Rebased-main exact chunks ON; rollback 718.704 ms, so local is **3.072x faster** ([evidence](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | +| GB10 (`sm_121a`) | Qwen3.6-27B NVFP4, historical normalized prefill | 0.43 us/token/layer | 0.18 us/token/layer | **2.39x slower** | Directional only: unequal token clusters, older pin ([ledger](../.agents/parity-ledger.md)) | +| GB10 (`sm_121a`) | Qwen3.6-35B NVFP4, later local kernel A/B | 321.148 us c1; 960.313 us c6 | - | `PENDING` | Register vs tiled improved 4.7%/7.3%; no paired vLLM denominator ([record](../.agents/specs/gdn-prefill-conv-reg-2026-07-18.md)) | +| Jetson Thor (`sm_110`), AGX Orin (`sm_87`) | No matched GDN workload | - | - | `PENDING` | Runtime correctness only; no causal-conv speed trace | ### Qwen3.6-27B by concurrency diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index 66ffaf0cd..679e763d9 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -84,6 +84,7 @@ portable/reference path. In normal operation leave them unset. | `VT_GPU_SAMPLE` | on (CUDA) | Host-side sampling instead of on-GPU sampling | | `VT_GDN_PACKED_DECODE` | on (CUDA GDN) | Unpacked GDN decode path | | `VT_CONV_REG` | on (CUDA GDN) | The non-register-tiled short causal convolution | +| `VT_CONV_EXACT_CHUNKS` | on (CUDA GDN prefill) | Use `=0` for the legacy sequence-serial causal-conv mapping; default mirrors vLLM's exact `(sequence, 8-token chunk)` descriptor and is byte-identical | | `VT_FA2_PREFILL` | on (CUDA) | The portable prefill attention instead of the vendored FA2 | | `VT_FA2_DECODE` | on (CUDA) | The portable decode attention instead of the vendored FA2 | | `VT_FA2_DECODE_4B` | on (CUDA, Qwen3.5-4B) | The portable paged decode attention instead of the ratio-4 vendored FA2 path; the 27B and 35B selectors are unchanged | diff --git a/docs/FEATURES.md b/docs/FEATURES.md index 0877b8317..c9f507de5 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -94,7 +94,7 @@ speed-pending, which [BENCHMARKS.md](BENCHMARKS.md) tracks. | Architecture | Tested checkpoint(s) | Correctness gate | Speed vs reference | |---|---|---|---| -| `Qwen3_5ForConditionalGeneration` | Qwen3.6-27B (NVFP4, GDN hybrid) | strict 235/235 text, image+video 32/32 vs vLLM 0.25.0 | gate model: at or above vLLM. CUDA/CPU only; the off-CUDA host-pointer bug (#125) is fixed but unrun | +| `Qwen3_5ForConditionalGeneration` | Qwen3.6-27B NVFP4; Qwen3.5-4B BF16 | 27B strict 235/235 text + 32/32 image/video; 4B cached 3/3 | 27B at/above vLLM; 4B throughput 1.021x, latency/VRAM pending. CUDA/CPU only; the off-CUDA host-pointer bug (#125) is fixed but unrun | | `Qwen3_5MoeForConditionalGeneration` | Qwen3.6-35B-A3B (NVFP4, GDN MoE) | strict 315/315 text vs vLLM 0.25.0 | gate model: 0.93x to 1.03x grid | | `Qwen3ForCausalLM` | Qwen3 dense 0.6B/1.7B/4B/32B, NVFP4A16 | near-tie strict 16/16 vs vLLM 0.25.0 | c1 every-axis parity, c8 decode residual | | `Qwen3MoeForCausalLM` | Qwen3-Coder-30B-A3B | strict 6/6 vs vLLM 0.25.0 | 11/16 grid cells at or above graphed vLLM | @@ -122,7 +122,7 @@ speed-pending, which [BENCHMARKS.md](BENCHMARKS.md) tracks. | `LagunaForCausalLM` | poolside/Laguna-S-2.1-NVFP4, GGUF-Q4_K, Laguna-XS | byte-exact near-tie (distributional vs vLLM) | vLLM parity+ 1.03x, default on, via the `laguna-gen` CLI; the registered engine forward VT_CHECKs non-bf16 (`ARCH-ONE-SURFACE` fold) | | `KimiLinearForCausalLM` | Kimi-Linear-48B-A3B (KDA + NoPE-MLA + MoE) | **Folded onto the shared paged runner (ROW 7 §21, #122): engine==CLI 128/128 byte-identical; vs golden 122/128 (the intrinsic near-tie profile); FA2 paged MLA default-ON; SACRED post-fold green** | Served via `vllm_engine_load` + `vllm_complete_tokens` (ABI v13); server 19.0 tok/s wall vs vLLM ~21 (~0.90×), speed residual open | | `KimiK3ForConditionalGeneration` | Kimi-K3 (2.8T MoE) | scaffold: registry+config+enumeration gated, forward refuses | HW-infeasible (~1.56 TB); no run | -| `LlamaModel` | committed tiny synthetic embedding fixture (engine path == direct pooler path, identical vectors; f64 LAST+normalize reference); real checkpoint (e5-mistral class) is a NAMED residual | pooling/embed only, text paths refuse by task; `vllm_embed` + `/v1/embeddings` | n/a (CPU correctness-grade embeddings) | +| `LlamaModel` | landed tiny synthetic embedding fixture (engine path == direct pooler path, identical vectors; f64 LAST+normalize reference); real checkpoint (e5-mistral class) is a NAMED residual | pooling/embed only, text paths refuse by task; `vllm_embed` + `/v1/embeddings` | n/a (CPU correctness-grade embeddings) | | `ParakeetForCTC`, `ParakeetForRNNT`, `ParakeetForTDT` | nvidia/parakeet-ctc-0.6b/-1.1b, -rnnt-0.6b, -tdt-0.6b-v3 (transcribed, ids exact vs HF `generate()`, P4/P6 2026-08-07; not retained) + committed synthetic fold fixture | ASR transcription-only (`SupportsTranscription` mirror; text paths refuse by task); fold gate byte-identical to the pre-refactor pipeline | n/a (CPU correctness-grade ASR via `vllm_transcribe` + `/v1/audio/transcriptions`) | | `CohereForCausalLM` | Command-R / Cohere (and Cohere2) | scaffold: W0 tiny-random oracle run-verified; real-checkpoint gate blocked | no run | diff --git a/docs/STATUS.md b/docs/STATUS.md index 2e6cfc435..1ceb946d5 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -39,6 +39,10 @@ Startup-latency axis (2026-08-07): `MEASURED / provisional`. Cold launch to firs Protocol (2026-08-07): PR disposition — verified-good PRs MERGE in-session, superseded CLOSE with a reason; prompt pair tracked, 25 gate rows exact-pinned. +Protocol repair (2026-08-08): release AST pins pass 30 tests on Python +3.12/3.13; Gemma-4 MoE is known drift pending the shared merged-GeGLU fold; +embeddings #137 is landed/partial, not an active claim. No runtime change. + Supported-model registry guard (2026-08-06): the public per-architecture list in [FEATURES](FEATURES.md) is CI-bound to the C++ registry by `scripts/check-supported-models.py` (+ mutation test), so the 30 @@ -67,7 +71,7 @@ token-for-token correctness against the pinned oracle. | Qwen3.6-27B (NVFP4) text generation | Correctness-complete, at/above vLLM speed | Token-exact greedy on GB10; beats vLLM 0.25.0 total throughput at every concurrency (1.007-1.045x), effective parity 115/124 axes | | Qwen3.6-35B-A3B (NVFP4, GDN MoE) | Correctness-complete; 3-rep grid 0.93-1.03x. Async batch-1 token-0 degeneration FIXED: `VT_ASYNC_DEVICE_MIRROR` default ON | Token-exact SYNC+ASYNC (RED→GREEN); c16 0.93x; `VT_ASYNC_EXECUTOR` Option A (H2D out of capture) GREEN+RED but A/B NEUTRAL → OFF; c16 residual is prefill glue | | Qwen3 / Qwen2 dense (BF16) | Correctness-complete, speed-pending. Async-serving P0 FIXED (`ROW-SERVE-ASYNC-DENSE-MIRROR`): classic-dense `Qwen3ForCausalLM` now honors the async device token-ids mirror; CPU-only -Werror test-guard fixes x2 | Near-tie-robust token-exact vs vLLM (Qwen3-0.6B, Qwen3-4B); c1 effective parity, c8 decode residual. **Async device-mirror (`ROW-SERVE-ASYNC-DENSE-MIRROR`, `f9c969ae`): the #31 fix ported to the classic dense family, dgx-VERIFIED.** The shared dense `EmbedInto` (qwen3.cpp) raced the async combine's device input-ids write against a stale host upload → token-0 degeneration on the depth-2 AsyncLLM serving path (quant-independent). `EmbedInto` now consumes the device override published by `ForwardQwen3ForCausalLM`'s `DeviceTokenIdsScope` (27B-dense template); gate `test_qwen3_dense_async_serving` RED on `VT_ASYNC_DEVICE_MIRROR=0`, GREEN default, byte-identical mirror-off. dgx GB10: async gate RED→GREEN 0.6B+4B, SACRED 0.6B+4B 184/184 unchanged (byte-neutral sync path), memcheck 0 errors; Yi30/Qwen3-8B-MXFP4 default-config e2e coherent + 3/4 token-exact (p2 = oracle-ratified near-tie, gap 0.0000), closing the QUANT-CT-MXFP4 async-default residual. RESIDUAL: sibling InternLM2/Mistral/Llama scope one-liner; W4 bench RAN; FA2 GQA-swap default-ON, c2-c8 <1.0x. `FLASH-PTXAS` #82: codegen at PARITY (no ptxas lever); gap=engine context. **D1 (2026-07-31, `CLAIM-D1-BF16-MERGED-QKV`): the bf16 merged-QKV path (`Qwen3QkvMergeEnabled`/`VT_QWEN3_QKV_MERGE`) is now default-ON** — one `vt::MatmulBT` over the merged `[qdim+2kdim,H]` owner + a contiguous `vt::QkvSplit` (OLMo-2 exemplar), replacing three per-shard GEMMs. Bit-exact GEMM math (A/B unit `test_ops_qkv_merge` byte-identical, RED-first); the wider-N cuBLASLt K-reduction flips the 0.6B genuine bf16 near-tie so the SACRED 0.6B golden was regenerated (all tokens within the near-tie band, max 0.125 nats), while Qwen3-4B is byte-neutral (0 diffs, stays STRICT). Re-gated 0.6B 16/16 + 4B 16/16; consistency/launch-count fold (measured NEUTRAL on 4B decode), no new throughput owed | -| Qwen3.5-4B plain BF16 direct loading on discrete CUDA | Correctness-complete, speed-pending | Revalidated after merging current upstream: local throughput is unchanged at 0.99997x its prior run; against the freshly measured pinned oracle it is 0.9971x. TTFT 0.7719x and host PSS 0.3127x pass; TPOT/ITL 1.1244x and VRAM 1.0014x remain open. Direct ON/OFF outputs remain 128/128 identical | +| Qwen3.5-4B plain BF16 direct loading on discrete CUDA | Correctness-complete; throughput passes, latency/VRAM open | Exact GDN chunks default ON and byte-identical to rollback. Local A/B: total/output +2.152%, TTFT -2.945%, TPOT/ITL -1.920%; sealed-vLLM comparison 1.021x throughput, 1.086x TTFT, 1.025x TPOT, +233 MiB VRAM ([evidence](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | | Qwen3-Coder-30B-A3B MoE (BF16) | Correctness-complete, speed-pending | Near-tie-robust token-exact 6/6; 11 of 16 binding grid cells at or above vLLM. **D1 (2026-07-31): inherits the default-ON bf16 merged-QKV via the shared dense `AttnBlock` — byte-neutral (0 token diffs, golden UNCHANGED); re-gated 6/6** | | Llama-3.x dense (BF16) | Correctness-complete, speed-pending | Near-tie-robust token-exact 16/16 (Llama-3.2-1B); llama3 RoPE scaling | | Mistral dense (BF16) | Correctness-complete, speed-pending | Paged-engine token-exact 16/16 (Mistral-7B-v0.3) | @@ -1283,28 +1287,24 @@ regression. ## Performance detail -**Local Qwen3.5-4B plain BF16 direct loader, speed-pending:** on the current -`upstream/main` tree at `59674cf1d`, the uncontended three-repetition comparison -on an RTX 5070 Ti measured 6611.207 total tok/s versus the pinned vLLM oracle at -6630.481 tok/s (0.9971x). Mean TTFT is 730.403 versus 946.214 ms -(PASS), peak/stable host PSS is 2.531/0.739 versus 8.093/4.422 GiB (PASS), while -mean TPOT/ITL is 38.143 versus 33.924 ms and peak VRAM is 12850 versus 12832 MiB -(OPEN). Local throughput is 0.99997x its previous run, a null result. - -The failing TPOT axis has a corrected diagnosis as of the W4 work below: the -per-step synchronization it was blamed on is the synchronous engine loop -waiting for its own sampling, about one per step, not a removable defect in -the async sampler. The device-resident sampled-token mirror is now default ON; -this synchronous harness does not claim the remaining serving comparison. - -The earlier 5769.99 tok/s series remains VOID because a graphics consumer held -the GPU at 11-13% utilization. Both arms gained about 14% once measured on an -idle box; the harness now rejects that contention. The current 0.9971x result -passed the same idle gate on all 18 legs. +**Local Qwen3.5-4B plain BF16 direct loader, throughput passed; latency and +VRAM pending:** the benchmark now uses production `AsyncLLM`. Exact +`(sequence, 8-token chunk)` causal-conv metadata is built once per step, shared +across GDN layers and default ON. `VT_CONV_EXACT_CHUNKS=0` is a byte-identical +same-binary rollback. Rebased-main graph-node `nsys` confirms the mechanism: +causal-conv falls from **718.704 to 233.955 ms (3.072x)**, leaving **1.609x** +to vLLM. The profiled whole run improves **2.272%** with identical token files. + +The enclosing A/B improves total/output **2.152%**, TTFT **2.945%**, TPOT/ITL +**1.920%** and E2E latency **2.118%** without a local VRAM regression. Against +the sealed same-workload vLLM baseline, throughput is **1.021246x PASS**; TTFT +is **1.085812x OPEN**, TPOT/ITL **1.024597x OPEN**, and mean peak VRAM +13053.3/12820 MiB OPEN. Fresh 18-leg oracle attempts were VOID JIT-environment +runs and do not replace the sealed denominator. This local 4B diagnostic does not establish 27B/35B support. Exact evidence and reproduction: -[Qwen3.5-4B post-upstream revalidation](bench-evidence/qwen35-4b-upstream-20260805.md). +[Qwen3.5-4B exact-chunk outcome](bench-evidence/qwen35-4b-sm120-main-20260807.md). There is no front-page race clip yet; when one is produced it will follow the LocalAI house style (side-by-side, identical output, honest measured ratios). diff --git a/docs/bench-evidence/qwen35-4b-sm120-main-20260807.md b/docs/bench-evidence/qwen35-4b-sm120-main-20260807.md new file mode 100644 index 000000000..7865d9552 --- /dev/null +++ b/docs/bench-evidence/qwen35-4b-sm120-main-20260807.md @@ -0,0 +1,296 @@ +# Qwen3.5-4B on RTX 5070 Ti — fresh-main production baseline (2026-08-07) + +This run refreshed `upstream/main` to `68b394bc2`, applied only the +profile-aware agent-efficiency toolkit (`7ef5f1001`), rebuilt under a bounded +user-systemd cgroup, re-ran cached-model correctness, and executed the exact +three-repetition comparison. The first run found a benchmark frontend mismatch. +That run remains diagnostic evidence below; the corrected production-frontend +series at `3f35356e0` supersedes it and selects the first sm_120 optimization +metric. + +## Reproduction identity + +- Local source: `7ef5f10011dc34bc798883512bbea57238167658`, based directly on + `upstream/main` `68b394bc2`. +- Model: `Qwen/Qwen3.5-4B`, cached snapshot `851bf6e...`. +- Corpus: `/tmp/qwen35-4b-sharegpt-1024.json`, SHA-256 + `9ea13603767c62c267e3f381fbccf42d0c9ca0c393655c37533eadca7aefca0c`. +- Workload: 128 requests, 128 output tokens, concurrency 32, + `max_num_batched_tokens=2048`, greedy, direct ON / pinned vLLM / direct OFF. +- GPU exclusion: one `/tmp/gpu` lock across all 18 legs; no competing work. +- Containment: user-systemd scope with `MemoryHigh=22G`, `MemoryMax=25G`, + `MemorySwapMax=0`. The clean fresh-main build completed in 362 seconds and + the host remained responsive. +- Correctness before timing: `test_qwen35_plain_weights` 3/3 cases, + 1672/1672 assertions. +- Manifest: `/tmp/vllm-agent-runs/qwen35-baseline-7ef5f1001.json`; aggregate: + `/tmp/qwen35-main-68b394bc2-7ef5f1001/aggregate.json`. + +## Diagnostic result + +Means of three timed repetitions: + +| Axis | local direct ON | pinned vLLM | ratio / verdict | +|---|---:|---:|---:| +| total throughput | 6625.323 tok/s | 6643.621 tok/s | 0.997246x | +| output throughput | 732.607 tok/s | 734.633 tok/s | 0.997242x | +| mean TTFT | 719.887 ms | 936.576 ms | 0.76864x, local faster | +| mean TPOT / ITL | 38.133 ms | 33.916 ms | 1.12435x, local slower | +| peak PSS | 2.048 GiB | 7.855 GiB | 0.2607x | +| stable PSS | 0.769 GiB | 4.495 GiB | 0.1712x | +| peak VRAM | 12852 MiB | 12844 MiB | 1.00062x, +8 MiB | + +Direct ON versus direct OFF remained 128/128 token-identical in every +repetition. Direct ON total throughput was 1.00214x the 2026-08-05 historical +run; the source refresh itself was a null. + +## Same-tool structural trace + +Both arms used Nsight Systems 2025.1.3.140 with +`--trace=cuda,nvtx --cuda-graph-trace=node --sample=none --cpuctxsw=none` on the +same workload: + +- ours: `/tmp/qwen35-main-68b394bc2-7ef5f1001-ours.nsys-rep` +- vLLM: `/tmp/qwen35-main-68b394bc2-7ef5f1001-vllm.nsys-rep` + +The steady benchmark intervals contain about 21.7 seconds of local GPU kernel +time and 22.0 seconds of vLLM kernel time. Both are GPU-saturated and total +device work is effectively tied. Sampling-grid reconstruction nevertheless +differs: local has 529 sample steps (497 at batch 32; mean about 31.1), while +vLLM has 650 (mean about 25.3, with substantial batch 30/6/32 waves). That +composition is consistent with the near-equal throughput plus local TTFT win +and local TPOT loss. + +The cause is now source-proven: `vllm-bench` logs async scheduling enabled but +uses `LoadedEngine::engine()` and synchronous `LLMEngine::step()`. The pinned +vLLM engine selects `step_with_batch_queue` at `max_concurrent_batches=2`. +Therefore this trace compares different frontend/queue modes. The correction +and gates are specified in +[the `SERVE-CLI-BENCH` spike](../../.agents/specs/cli-serve-bench.md). The old +0.9971x public ratio and this 0.99725x refresh remain useful stability controls, +but neither is accepted as production async parity until the harness is rerun. + +## Superseding production-frontend comparison + +`RunBench` now uses `LoadedEngine::async_engine()`, and every local leg records +the frontend, resolved async flag and maximum concurrent batches. The corrected +series ran at `3f35356e0ac1c11689efefdb57f9cfc5af35275e`, with the same model, +corpus, workload and pinned vLLM. One `/tmp/gpu` lock covered all 18 legs. The +whole run was contained by a user-systemd scope with `MemoryHigh=22G`, +`MemoryMax=25G` and swap disabled; peak cgroup memory was 17.75 GiB and the host +remained responsive. + +Both `VT_DIRECT_DEVICE_LOAD=1` and `=0` report `AsyncLLM`, async scheduling ON +and maximum concurrent batches 2. They are direct-loader residency controls, +not async ON/OFF controls. The earlier “depth-1 OFF” wording is withdrawn. + +Means of three timed repetitions: + +| Axis | local direct ON | pinned vLLM | ratio / verdict | +|---|---:|---:|---:| +| total throughput | 6633.750 tok/s | 6643.593 tok/s | 0.998518x, open | +| output throughput | 733.540 tok/s | 734.630 tok/s | 0.998516x, open | +| mean TTFT | 1051.333 ms | 937.584 ms | 1.121322x, open | +| mean TPOT / ITL | 35.457 ms | 33.906 ms | 1.045721x, open | +| peak PSS | 2.328 GiB | 7.932 GiB | 0.29350x, pass | +| stable PSS | 0.782 GiB | 4.505 GiB | 0.17361x, pass | +| peak VRAM | 13054 MiB | 12820 MiB | 1.018253x, +234 MiB, open | + +Direct ON repetitions were 6630.42 / 6638.71 / 6632.12 tok/s; vLLM was +stable at 6643.593 tok/s mean. Direct OFF measured 6509.900 tok/s, 35.507 ms +mean TPOT and 8.612 GiB peak PSS. Every same-arm repetition was 128/128 +token-identical; direct ON versus OFF was also 128/128 in all three pairs. +The local and vLLM streams were 87/128 request-identical in the unprofiled +series. The cached correctness gate remains the binding correctness evidence: +3/3 cases and 1672/1672 assertions. No cross-engine token-exact claim is made +from this throughput corpus. + +The frontend correction changes the latency trade: compared with the invalid +synchronous diagnostic, total throughput rises only 0.127%, TPOT improves from +38.133 to 35.457 ms, and mean TTFT rises from 719.887 to 1051.333 ms. That is +why isolated TPOT is not the optimization objective. + +## Corrected paired trace and selected metric + +The same Nsight Systems build and graph-node trace settings captured both +production frontends under one GPU lock: + +- ours: `/tmp/qwen35-async-3f35356e0-ours.nsys-rep`, SHA-256 + `182426e961ccaa53896e7cb7e9c4d2ba1cc35ba695ea3a9c3c727195326c8dab` +- vLLM: `/tmp/qwen35-async-3f35356e0-vllm.nsys-rep`, SHA-256 + `9d543390741629fc3873ad857706595cf880ee2fcce63af3e4be17e1a343355a` + +Request-processing windows are 22.370 seconds local and 22.269 seconds vLLM. +Kernel-interval union gives 98.54% and 99.34% GPU busy, with 0.327 and 0.146 +seconds idle. Local uses 536 sampling steps and vLLM 650, but their sampled +row totals are effectively identical (16,442 / 16,443), so family totals are +comparable despite different queue waves. + +The first optimization metric is **total steady-interval GPU time in batched +GDN prefill `causal_conv1d`**: + +| | calls | total GPU time | mean call | +|---|---:|---:|---:| +| local `CausalConv1dFwdRegKernel` | 1728 | **720.954 ms** | 417.219 us | +| vLLM `_causal_conv1d_fwd_kernel` | 1893 | **145.421 ms** | 76.821 us | + +That is a 4.958x local gap and about 576 ms of visible headroom, large enough to +close the 0.15% end-to-end throughput miss if the gain survives in situ. It is +also source-explained: local chunking is disabled above four sequences and +dominant c32 launches serialize whole sequences as `grid=(64,28..32,1)`; +vLLM consumes exact `(sequence, token-chunk)` scheduler metadata and retains +token parallelism at batch 32. The full port and RED-first gate plan are in the +[sm_120 convolution-chunking spike](../../.agents/specs/sm120-qwen35-conv-chunking-2026-08-07.md). + +End-to-end acceptance remains multi-axis: token correctness, total/output +throughput, all latency quantiles, PSS and VRAM. A faster micro-kernel cannot +claim success by moving cost into TTFT or memory. + +## Corrected artifact identity + +- aggregate: `/tmp/qwen35-async-3f35356e0/aggregate.json`, SHA-256 + `d006d6ffd6d014fc1861a30c126f603117a1ff80b339172d63fab75f2ae07f1d` +- 18-leg wrapper manifest: + `/tmp/vllm-agent-runs/qwen35-async-baseline-3f35356e0.json`, SHA-256 + `619a3cb2ec95800a19371d0d929d5c9bd4fe23b2ffe2076fe40f4e8533f2ec63` +- paired-trace wrapper manifest: + `/tmp/vllm-agent-runs/qwen35-async-nsys-pair-3f35356e0.json`, SHA-256 + `53b10bdba75ab5bab3cee6a4109551547f95eb2ed6d567846d45219c5a735a07` + +## Exact chunk implementation outcome + +The upstream work partition is now implemented. `GDNAttentionMetadata` builds +the exact `(sequence, 8-token chunk)` descriptor list from the query sequence +lengths once per step, uploads it once, and shares it across all GDN layers. +`CausalConv1dFwdRegKernel` maps `grid.y` directly through those descriptors. +`VT_CONV_EXACT_CHUNKS=0` restores the legacy whole-sequence dispatch in the +same binary; `VT_CONV_REG=0` remains the tiled/scalar rollback. Exact chunks +are default ON after the gates below. + +The implementation was RED-first: the metadata test initially failed to +compile on the absent `ComputeCausalConv1dMetadata` seam. The focused host +metadata and flag tests, affected Qwen fixtures, full CUDA GDN suite, and cached +Qwen3.5-4B gate are green. The real-model gate is 3/3 cases and 1672/1672 +assertions. Three production ON/OFF benchmark pairs are byte-identical for all +128 requests and 128 generated tokens. + +### Same-binary causal-conv result + +Nsight Systems graph-node traces of the identical production workload give: + +| Arm | Calls | Total GPU time | Mean call | +|---|---:|---:|---:| +| legacy whole-sequence (`VT_CONV_EXACT_CHUNKS=0`) | 1,728 | 720.047 ms | 416.694 us | +| exact upstream chunks (default) | 1,728 | 234.607 ms | 135.768 us | +| pinned vLLM trace | 1,893 | 145.421 ms | 76.821 us | + +Exact dispatch is **3.069x faster** than the same local binary's rollback and +removes 485.440 ms of causal-conv work. The dominant grid changes from +`(64,28..32,1)` whole-sequence launches to exact `grid.y=279/280/282` chunk +counts, directly confirming the proposed mechanism. It explains about 84.5% +of the original local/vLLM causal-conv gap. The residual is **1.613x**; the +trace suggests feature-tile and register-pressure differences are the next +hypothesis, not yet an accepted claim. + +Artifacts: + +- manifest: `/tmp/vllm-agent-runs/qwen35-conv-exact-ab-profile.json`; +- rollback trace: `/tmp/qwen35-conv-exact-off.nsys-rep`; +- default trace: `/tmp/qwen35-conv-exact-on.nsys-rep`. + +### Enclosing production benchmark + +Three alternating rollback/default pairs were run under one GPU lock and the +same 25 GiB user-systemd cap, with file cache dropped and the GPU idle before +each leg: + +| Axis | rollback mean | exact default mean | default / rollback | +|---|---:|---:|---:| +| total throughput | 6641.800 tok/s | 6784.743 tok/s | **1.02152x** | +| output throughput | 734.433 tok/s | 750.237 tok/s | **1.02152x** | +| mean TTFT | 1048.927 ms | 1018.040 ms | **0.97055x** | +| mean TPOT / ITL | 35.420 ms | 34.740 ms | **0.98080x** | +| mean E2E latency | 5547.687 ms | 5430.180 ms | **0.97882x** | + +The default improves every measured local axis. Against the sealed pinned-vLLM +denominator above, total/output/input throughput is now **1.021246x** faster; +TTFT remains **1.085812x** slower and TPOT/ITL **1.024597x** slower. Peak local +VRAM across the completed memory legs is 13044/13058/13058 MiB (13053.3 MiB +mean), effectively unchanged from the 13054 MiB baseline and still above +vLLM's sealed 12820 MiB mean. Host PSS remains substantially better locally. + +The enclosing A/B manifest is +`/tmp/vllm-agent-runs/qwen35-conv-exact-local-ab.json`; results are under +`/tmp/qwen35-conv-exact-local-ab-20260807`. + +### Voided fresh-oracle attempts + +Two attempts to repeat the entire 18-leg local/vLLM/local series are retained +as **VOID infrastructure evidence**, not a new denominator. The first exposed +that FlashInfer's CUDA JIT linker also needs the venv CUDA directory in +`LIBRARY_PATH`; the harness now supplies it. The repaired run completed 13/18 +legs before a transient Torch bytecode read failure invalidated Torch/Triton +AOT cache keys and induced unrelated Triton compiler failures. No oracle model +or inference semantics were changed. Manifests: + +- `/tmp/vllm-agent-runs/qwen35-conv-exact-default-full-compare.json`; +- `/tmp/vllm-agent-runs/qwen35-conv-exact-default-full-compare-rerun.json`; +- repaired-oracle preflight: + `/tmp/vllm-agent-runs/qwen35-conv-exact-vllm-preflight.json`. + +Therefore the accepted change is grounded in the same-binary local profile and +three-pair local A/B; its cross-engine ratios reuse the sealed, same-hardware, +same-workload pinned-vLLM baseline. This 4B result does not establish the +hardware-unavailable 27B/35B release gates. + +## Clean current-main transplant reproduction + +The final code and required production-benchmark correction were transplanted +without the unrelated efficiency-tooling or intermediate measurement commits +onto `upstream/main` `f91a5917a`. A contained CPU/CUDA rebuild passes: focused +CPU 6/6, full CUDA GDN 66/66 cases and 4300 assertions, cached Qwen3.5-4B 3/3 +and 1672 assertions. + +A fresh same-binary graph-node profile reproduces the result, and rollback and +default token files compare byte-for-byte: + +| Arm | Calls | Total GPU time | Mean call | +|---|---:|---:|---:| +| legacy whole-sequence | 1,728 | 720.216507 ms | 416.792 us | +| exact chunks | 1,728 | 234.379395 ms | 135.636 us | + +That is **3.072866x** at the kernel. The profiled whole run improves from +6587.66 to 6727.35 tok/s (**2.1205%**), TTFT 1058.73→1025.46 ms, TPOT +35.70→35.04 ms and E2E 5592.69→5475.91 ms. Against the sealed vLLM conv trace, +the residual is **1.611730x**. Artifacts: + +- `/tmp/qwen35-conv-exact-transplant-off.nsys-rep`; +- `/tmp/qwen35-conv-exact-transplant-on.nsys-rep`; +- `/tmp/qwen35-conv-exact-transplant-{off,on}.tokens.json`. + +## Post-rebase reproduction on `upstream/main` `48a54141f` + +The branch was rebased onto `48a54141f` and rebuilt as `3d2581551` under the +same 22/25 GiB user-systemd limits. The focused suite is 6/6, full CUDA GDN is +66/66 cases and 4300/4300 assertions, and cached Qwen3.5-4B is 3/3 cases and +1672/1672 assertions. + +A fresh same-binary graph-node profile again produced byte-identical token +files and measured: + +| Arm | Calls | Total GPU time | Mean call | +|---|---:|---:|---:| +| legacy whole-sequence | 1,728 | 718.704016 ms | 415.917 us | +| exact chunks | 1,728 | 233.954533 ms | 135.390 us | + +That is **3.07198x** at the kernel, saving 484.749 ms. The profiled enclosing +run improves 6589.65→6739.34 tok/s (**2.272%**), TTFT 1057.63→1022.70 ms, +TPOT 35.70→34.99 ms and E2E 5590.92→5466.20 ms. Against the sealed vLLM +causal-conv trace, the residual is **1.60881x**. Artifact SHA-256 values: + +- rollback trace `6a5dde18eae3b7a8157def69cf8167dd4317e4fad3ec0d14b7d02e6976f97c47`; +- exact trace `f47fb9cc50d21324056659aea7f94eb1f88828f93dce109bec23cbc17f7aecf9`; +- both token files `83fcdc45f79ddb06a634c7d7d95eba3384543b3cd781a45a8db1fc4e2a453545`. + +The files are `/tmp/qwen35-conv-exact-rebase-3d2581551-{off,on}.nsys-rep` +and `/tmp/qwen35-conv-exact-rebase-3d2581551-{off,on}.tokens.json`. diff --git a/examples/bench/bench_core.h b/examples/bench/bench_core.h index 912737d35..274e3a3a5 100644 --- a/examples/bench/bench_core.h +++ b/examples/bench/bench_core.h @@ -44,6 +44,7 @@ #include #include #include +#include #include #include @@ -56,6 +57,7 @@ #include "vllm/tokenizer/bpe.h" #include "vllm/tokenizer/tokenizer.h" #include "vllm/transformers_utils/hf_config.h" +#include "vllm/v1/engine/async_llm.h" #include "vllm/v1/engine/llm_engine.h" #include "vt/dtype.h" #include "vt/tensor.h" @@ -113,6 +115,12 @@ struct RequestRecord { // ── Aggregated result (field names mirror serve.py BenchmarkMetrics). ────────── struct BenchResult { + // Auditable engine selection. The comparison harness uses the production + // AsyncLLM frontend even when async scheduling resolves OFF (then the core's + // queue depth is one), matching vLLM's frontend across the ON/OFF control. + bool async_frontend = false; + bool async_scheduling_enabled = false; + int max_concurrent_batches = 1; int completed = 0; double duration_s = 0.0; int64_t total_input = 0; @@ -428,9 +436,11 @@ inline SamplingParams MakeSampling(const BenchConfig& cfg, int req_index) { // ─────────────────────────────── The harness ────────────────────────────────── // Creates the engine (synthetic if cfg.model_path is empty, else loaded from the -// dir/.gguf), builds N prompts, then drives the V1 engine step() loop admitting -// up to C requests at a time until all N finish — timing everything with -// steady_clock. Returns the aggregated metrics. +// dir/.gguf), builds N prompts, then drives the production AsyncLLM frontend, +// admitting up to C requests at a time until all N finish — timing everything +// with steady_clock. Async scheduling OFF still uses AsyncLLM with a depth-1 +// core queue, so the ON/OFF comparison changes only scheduler/runner behavior. +// Returns the aggregated metrics. inline BenchResult RunBench(const BenchConfig& cfg) { using Clock = std::chrono::steady_clock; @@ -495,9 +505,10 @@ inline BenchResult RunBench(const BenchConfig& cfg) { } } - vllm::v1::LLMEngine& engine = loaded->engine(); + vllm::v1::AsyncLLM& engine = loaded->async_engine(); std::map records; + std::map active; const Clock::time_point t0 = Clock::now(); auto now_s = [&]() { return std::chrono::duration(Clock::now() - t0).count(); @@ -513,8 +524,9 @@ inline BenchResult RunBench(const BenchConfig& cfg) { RequestRecord rec; rec.arrival_s = now_s(); records[rid] = rec; - engine.add_request(rid, prompts[static_cast(next)], - MakeSampling(cfg, next)); + active.emplace( + rid, engine.add_request(rid, prompts[static_cast(next)], + MakeSampling(cfg, next))); ++next; ++in_flight; } @@ -522,7 +534,15 @@ inline BenchResult RunBench(const BenchConfig& cfg) { admit(); while (done < cfg.num_prompts) { - for (RequestOutput& out : engine.step()) { + bool observed_output = false; + for (auto it = active.begin(); it != active.end();) { + std::optional ready = engine.get_output_nowait(it->second); + if (!ready.has_value()) { + ++it; + continue; + } + observed_output = true; + RequestOutput& out = *ready; RequestRecord& rec = records[out.request_id]; if (rec.prompt_tokens == 0 && !out.prompt_token_ids.empty()) { rec.prompt_tokens = static_cast(out.prompt_token_ids.size()); @@ -546,9 +566,18 @@ inline BenchResult RunBench(const BenchConfig& cfg) { rec.completion_s = now_s(); --in_flight; ++done; + it = active.erase(it); + } else { + ++it; } } admit(); // keep C in flight as requests finish. + if (!observed_output) { + // The output-handler thread will publish the next per-request DELTA. + // Yield rather than block on an arbitrary request: blocking on one + // collector can delay ready outputs for other requests and distort ITL. + std::this_thread::yield(); + } } const double dur_s = now_s(); @@ -576,6 +605,9 @@ inline BenchResult RunBench(const BenchConfig& cfg) { } BenchResult res; + res.async_frontend = true; + res.async_scheduling_enabled = loaded->async_scheduling_enabled(); + res.max_concurrent_batches = loaded->max_concurrent_batches(); res.output_token_ids.resize(static_cast(cfg.num_prompts)); res.completed = done; res.duration_s = dur_s; @@ -651,6 +683,10 @@ inline void PrintReport(const BenchConfig& cfg, const BenchResult& r, }; std::fprintf(out, "\n============= vllm.cpp Benchmark Result =============\n"); + std::fprintf(out, "%-42s %-12s\n", "Engine frontend:", + r.async_frontend ? "AsyncLLM" : "LLMEngine"); + line_i("Async scheduling enabled:", r.async_scheduling_enabled ? 1 : 0); + line_i("Maximum concurrent batches:", r.max_concurrent_batches); line_i("Successful requests:", r.completed); line_i("Maximum request concurrency:", cfg.concurrency); line_f("Benchmark duration (s):", r.duration_s); diff --git a/examples/bench/main.cpp b/examples/bench/main.cpp index 93bc4d481..08fb0bb87 100644 --- a/examples/bench/main.cpp +++ b/examples/bench/main.cpp @@ -1,7 +1,7 @@ // vllm-bench — the M2.1 throughput/latency benchmark harness for vllm.cpp. // // This is the tool that MEASURES gate #1 (throughput parity vs vLLM, -// .agents/gates.md): it drives the V1 LLMEngine step() loop at a fixed +// .agents/gates.md): it drives the production V1 AsyncLLM frontend at a fixed // concurrency and reports request/output/total token throughput plus // TTFT/TPOT/ITL/E2EL, with a prefill-vs-decode split. The metrics + report // format mirror `vllm bench serve` / `vllm bench throughput` @@ -48,7 +48,7 @@ void Usage(const char* argv0, std::FILE* out) { " [--output-token-ids ]\n" " [--speculative-config ]\n" "\n" - "Throughput/latency benchmark over the vllm.cpp V1 LLMEngine, mirroring\n" + "Throughput/latency benchmark over the vllm.cpp V1 AsyncLLM, mirroring\n" "`vllm bench serve` metrics. With no --model, a synthetic CPU engine runs\n" "(numbers meaningless; smoke only). Defaults: --num-prompts 8 " "--input-len 16\n" diff --git a/include/vllm/v1/attention/backends/gdn_attn.h b/include/vllm/v1/attention/backends/gdn_attn.h index ebea24f4c..220ef0d38 100644 --- a/include/vllm/v1/attention/backends/gdn_attn.h +++ b/include/vllm/v1/attention/backends/gdn_attn.h @@ -32,15 +32,12 @@ // // ─── DEFERRED upstream fields (T0 gate models never exercise them) ────────── // * Triton-kernel launch metadata: chunk_indices / chunk_offsets (FLA chunk -// kernel) and nums_dict / batch_ptr / token_chunk_offset_ptr (Triton -// causal_conv1d). OMITTED: our sequential C++ vt::GdnPrefill / -// vt::GdnSpecDecode and -// vt::CausalConv1dFwd consume `prefill_query_start_loc` + a -// `has_initial_state` mask DIRECTLY (ops.h — GdnPrefill takes -// query_start_loc; CausalConv1dFwd takes query_start_loc + has_initial_state), -// so the chunk/conv-tile precompute has no C++ consumer. Upstream itself -// tolerates them being None — the FLA ops recompute on the fly -// (gdn-semantics.md §8). +// kernel) and nums_dict. The causal-conv batch_ptr / +// token_chunk_offset_ptr pair IS ported below and consumed by the optional +// exact-chunk CUDA dispatch; our sequential C++ vt::GdnPrefill / +// vt::GdnSpecDecode still consume query_start_loc directly. Upstream itself +// tolerates the remaining FLA fields being None — those ops recompute on +// the fly (gdn-semantics.md §8). // * the FLA/cutedsl prefill-backend selection. #ifndef VLLM_V1_ATTENTION_BACKENDS_GDN_ATTN_H_ #define VLLM_V1_ATTENTION_BACKENDS_GDN_ATTN_H_ @@ -77,6 +74,19 @@ inline constexpr int32_t kNullStateSlot = -1; std::tuple SplitDecodesAndPrefills( const CommonAttentionMetadata& m, int decode_threshold = 1); +// Exact causal-conv work descriptor from upstream +// utils.py::compute_causal_conv1d_metadata @ the parity pin. One program owns +// one (sequence, 8-token chunk); unequal sequence lengths therefore launch no +// rectangular padding work. Built from the already-host-resident cu_seqlens so +// it introduces no device-to-host synchronization. +inline constexpr int32_t kCausalConv1dBlockM = 8; +struct CausalConv1dMetadata { + std::vector batch_ptr; + std::vector token_chunk_offset_ptr; +}; +CausalConv1dMetadata ComputeCausalConv1dMetadata( + const std::vector& query_start_loc_cpu); + // The GDN prefill/decode/spec segmentation for one batched step. // (Upstream @dataclass GDNAttentionMetadata, gdn_attn.py:42-79.) Field names // mirror upstream 1:1. `std::optional` mirrors upstream's `torch.Tensor | None` @@ -131,6 +141,13 @@ struct GDNAttentionMetadata : AttentionMetadata { std::optional> prefill_state_indices; std::optional> prefill_has_initial_state; + // Flattened exact causal-conv program map over the WHOLE non-spec stream. + // Mixed batches include their leading decode rows because causal_conv1d_fn + // consumes that whole stream before recurrence segmentation. None when there + // is no prefill and the single-token update kernel is selected instead. + std::optional> batch_ptr; + std::optional> token_chunk_offset_ptr; + // ── Spec-decode segmentation (SPEC-MTP I4; gdn_attn.py:189-326) ──────────── // All nullopt / 0 unless the step actually carries drafts. // diff --git a/include/vt/ops.h b/include/vt/ops.h index e6dee5eb3..07ad6ca43 100644 --- a/include/vt/ops.h +++ b/include/vt/ops.h @@ -449,6 +449,13 @@ struct CausalConv1dArgs { // Upstream `activation` is "silu"/"swish" (→ silu) or None (→ identity); // Qwen GDN always uses silu (gdn-semantics.md §2). bool silu_activation = true; + // Optional exact upstream prefill work descriptor, both i32 [num_programs] + // on the queue device. Entry p owns sequence batch_ptr[p] and its + // token_chunk_offset_ptr[p]-th 8-token chunk. CUDA consumes these when the + // default CUDA path is selected; VT_CONV_EXACT_CHUNKS=0 restores the legacy mapping and + // CPU keeps its scalar reference. + const Tensor* batch_ptr = nullptr; + const Tensor* token_chunk_offset_ptr = nullptr; }; struct L2NormArgs { diff --git a/scripts/check-gate-commands.py b/scripts/check-gate-commands.py index 64e25e6c1..23a562bf9 100755 --- a/scripts/check-gate-commands.py +++ b/scripts/check-gate-commands.py @@ -244,10 +244,11 @@ def audit() -> list[dict]: "KERNEL-GEMM-CPU-ELEM", "KV-CHUNKED-LOCAL-SPEC", "KV-SLIDING-LOCAL-SPECS", - # ARCH-ONE-SURFACE ROW 6 (2026-08-08): embeddings-one-surface.md carries a - # runnable Gates section (preflight + the fold/capi/server suites) for the - # two rows it activates. - "MODEL-EMBED-llama-llama-for-causal-lm", + # ARCH-ONE-SURFACE ROW 6 (2026-08-08): SERVE remains gated, while the + # merged MODEL row legitimately moved ACTIVE -> PARTIAL because only one + # of eight upstream embedding memberships is live. Re-pin removes that + # model row from this lifecycle-scoped runnable population; its completed + # fold commands remain preserved in embeddings-one-surface.md. "SERVE-POOLING-ENDPOINTS", "KV-SLIDING-WINDOW-SPEC", "LOAD-SAFETENSORS-DIRECT-DENSE", diff --git a/scripts/check-release-binary-contract.py b/scripts/check-release-binary-contract.py index ce424ca17..9aae47587 100644 --- a/scripts/check-release-binary-contract.py +++ b/scripts/check-release-binary-contract.py @@ -566,23 +566,23 @@ } TEST_INVENTORY_BODY_DIGESTS = { - "PRIMARY_CUDA_SMS": "5dc05132f5f24f0b2add406cb5682031b05d1940d2189728543b53a375d15125", - "EXACT_MACHINE_FIELDS": "8fa2ec5fca092a1092052a1426c37b66b8b8be49290e0e6e414e555dcbdad8dc", - "EXPECTED_DEPS": "4f8d345df7b467312869355f15153e04b77a0fd7005ea60ad86db40cc55ae6a2", - "HUMAN_WORK_IDS": "363c9494815086a53c3112792afe3c1651256d06a519e6fbeed43b78e865d922", - "RECORD_ANCHORS": "0e15f59b43bd4e70535055ccbcb41c500f89006fe9df4c2e58b0cfd0a00205f1", - "LIFECYCLE_RECORD_MUTATIONS": "79de0584503bb8c1cc1463169d8741455714c1d252136a89294e532878ec996a", - "PUBLIC_PENDING_MUTATIONS": "507e58eb51f04ab03db0cea484513da5cb0a7fdd2d6f6f332004e09106682fd3", - "W10_W12_HUMAN_MUTATIONS": "1fcb8914c91dc01052f7f26a45cbbaebe01c211bc0e692f70df3d62b13a963bf", - "PRIMARY_ARTIFACT_PROSE_MUTATIONS": "1fcb8914c91dc01052f7f26a45cbbaebe01c211bc0e692f70df3d62b13a963bf", - "GUARD_MAP_KEYS": "06691dc7239166ac458c9999d545938f7c8afb5021cb347d0fa6bdd0dac2a082", - "INVENTORY_CONSUMER_METHODS": "5754b33de9ca699665d4f612f8089371dbc7d9ce583422ffeb48a33499575eab", - "CONSUMER_FLOW_MUTATIONS": "a2d05b5bea24a09c6c5313f4e65d259a9b7e984f6450aa74309f790b3e81de8a", - "UNKNOWN_MACHINE_FIELD_MUTATIONS": "1d9242dabe625e43b909709ed0a40915d55fd6610c06bab3e696dc803745e0b8", - "HUMAN_WORK_DEPS": "800c69c64c995030fff12ccfbe2bb0002193c17cbe28aeca52c584cf15c2b675", - "BACKEND_POLICY_PROSE_MUTATIONS": "df80167a6da67c499ba1fdebda15e8c5250b46eccebb329973d7e63f7c6d0763", - "PREFLIGHT_WIRING_MUTATIONS": "635f1f49d2e04eb662e61d97f8e4128569d98b8af327fd33d3093451f3f24f7e", - "CI_WIRING_MUTATIONS": "ce2cb4b0c5c71431861c9af57b15468205ee335387120b48b150d7244a008885", + "PRIMARY_CUDA_SMS": "43e348a6fefad920d5ac461ef34868d20c05af64d3c9f032c62af469a358dee9", + "EXACT_MACHINE_FIELDS": "f7389b004be2b5665456e893abfa8ebb1b404c84711e86aba9673fcf8775c971", + "EXPECTED_DEPS": "b7d4608bab17632a8a02e7da6f7b8f656415c9ed08dc4ee18f7268709ec91512", + "HUMAN_WORK_IDS": "49195d0f7cd3f40d48c9f1282e4b9ead9571ca3a546c449c427801cac8fc8bdd", + "RECORD_ANCHORS": "5d354a9ed8590deccdc62890e403af66a21c416cb45e9788aacc4dce14364500", + "LIFECYCLE_RECORD_MUTATIONS": "ab35f4e72fe180cd3ef4675939d5ff8c709aff4c0474412d47ee78988c61199d", + "PUBLIC_PENDING_MUTATIONS": "69a3fc11686ccea3856b61f473796499466dc234e4aca952fce79bef2157714d", + "W10_W12_HUMAN_MUTATIONS": "17cb0586bf5ea235ba668bd0a4ae90345a33e125f211909e7f97267ec9e59dc8", + "PRIMARY_ARTIFACT_PROSE_MUTATIONS": "17cb0586bf5ea235ba668bd0a4ae90345a33e125f211909e7f97267ec9e59dc8", + "GUARD_MAP_KEYS": "701e4821bee926c2e074dbf2b97ff4a93bebb610cc6bed76e06063cab8974758", + "INVENTORY_CONSUMER_METHODS": "916894a32d88026a883cc1f316d949eb116ee1fced36d635b585d7bf3372b01d", + "CONSUMER_FLOW_MUTATIONS": "6f69f9e361d38c325fbc455c31ec7211578131368624312e024448afdfc01e83", + "UNKNOWN_MACHINE_FIELD_MUTATIONS": "8d10128593c67c64cde8cbe7e58faa9a9ade0d4bde391d24474fd04d6392ed6b", + "HUMAN_WORK_DEPS": "54a501b903eb3c97023084393666f9f63d289ab9a78e22f389c32bfc1711573b", + "BACKEND_POLICY_PROSE_MUTATIONS": "c5fea18a668932c4768cb9feb4746fd444b3df7e7ec15df1a588141898d28f2d", + "PREFLIGHT_WIRING_MUTATIONS": "d442c6d188efd624bffc9e94a7750d6a527c7b693affde5cbc33304f9e95272e", + "CI_WIRING_MUTATIONS": "7e20ed4d041fee98f96bb751e4435ad266d75ec8fbc9c5e1197a5d21940a6424", } EXACT_MACHINE_FIELDS = { @@ -1137,9 +1137,28 @@ def _inventory_loop( return matches[0] if len(matches) == 1 else None +_VERSION_ONLY_AST_FIELDS = frozenset({"type_params"}) + + +def _canonical_ast(value: object) -> object: + """Return a Python-version-stable representation of an AST value.""" + if isinstance(value, ast.AST): + return ( + type(value).__name__, + tuple( + (name, _canonical_ast(field_value)) + for name, field_value in ast.iter_fields(value) + if name not in _VERSION_ONLY_AST_FIELDS + ), + ) + if isinstance(value, list): + return tuple(_canonical_ast(item) for item in value) + return value + + def _consumer_body_digest(loop: ast.For) -> str: body = ast.Module(body=loop.body, type_ignores=[]) - serialized = ast.dump(body, include_attributes=False).encode("utf-8") + serialized = repr(_canonical_ast(body)).encode("utf-8") return hashlib.sha256(serialized).hexdigest() diff --git a/scripts/merged-gemm-consistency-allowlist.txt b/scripts/merged-gemm-consistency-allowlist.txt index bbfc62ceb..339b0d072 100644 --- a/scripts/merged-gemm-consistency-allowlist.txt +++ b/scripts/merged-gemm-consistency-allowlist.txt @@ -30,6 +30,7 @@ minicpm # SwiGLU dense MLP -> UnquantizedMlpGateUpMethod; pending FOLD-M minicpm3 # SwiGLU dense MLP (MLA arch) -> UnquantizedMlpGateUpMethod; pending FOLD-MIGRATE phi3 # SwiGLU dense MLP -> UnquantizedMlpGateUpMethod (also on the glue allowlist); pending FOLD-MIGRATE gemma4_vision # GeGLU vision-tower MLP -> UnquantizedMlpGateUpGeluMethod + a clamp-epilogue hook (Tier C2); pending FOLD-MIGRATE +gemma4_moe # known-drift pending fold: the ROCm/Gemma-4 MoE path has a merged [2I,H] GeGLU operand but launches gate/up separately for host-backed and device-resident expert weights; fold both through a shared merged-GEMM MoE descriptor without changing its residency fallback laguna # known-drift pending fold: Laguna NVFP4 resident/graph decode hand-rolls the shared-expert SwiGLU epilogue; fold onto layers::MlpGateUp seam is part of the decode/runtime framework-routing port (AGENTS.md 3rd seam) diff --git a/src/vllm/model_executor/models/qwen3_5.cpp b/src/vllm/model_executor/models/qwen3_5.cpp index 132c53c8f..d306ab7d4 100644 --- a/src/vllm/model_executor/models/qwen3_5.cpp +++ b/src/vllm/model_executor/models/qwen3_5.cpp @@ -308,6 +308,15 @@ void detail::ValidateGdnAttentionMetadata( } VT_CHECK(prefill_qsl.back() == np_tok, "qwen3_5: prefill query offsets must span prefill tokens"); + VT_CHECK(metadata.batch_ptr.has_value() && + metadata.token_chunk_offset_ptr.has_value(), + "qwen3_5: missing exact causal-conv chunk metadata"); + const v1::CausalConv1dMetadata expected_conv = + v1::ComputeCausalConv1dMetadata(full_qsl); + VT_CHECK(*metadata.batch_ptr == expected_conv.batch_ptr && + *metadata.token_chunk_offset_ptr == + expected_conv.token_chunk_offset_ptr, + "qwen3_5: causal-conv chunk metadata does not exactly cover query offsets"); } bool detail::CanUseGdnDecodeGraphSize(int64_t real_batch, @@ -3198,6 +3207,9 @@ struct StepDevInputs { DBuf gdn_prefill_qsl; // i32 [num_prefills+1] DBuf gdn_prefill_has_initial; // i8 [num_prefills] bool has_gdn_prefill_meta = false; + DBuf gdn_conv_batch_ptr; // i32 [num exact conv programs] + DBuf gdn_conv_token_chunk_offsets; // i32 [num exact conv programs] + bool has_gdn_conv_chunks = false; bool indexed_gdn_state_io = false; // ── Spec-decode device tensors (SPEC-MTP I5a). Uploaded ONCE per step (shared // by every GDN layer's spec branch), only when the step carries drafts @@ -3260,6 +3272,9 @@ StepDevInputs BuildStepDevInputs(Dev d, const std::vector& positions, DBuf(d, DType::kI32, {1}), // prefill qsl stub DBuf(d, DType::kI8, {1}), // prefill has-initial stub false, + DBuf(d, DType::kI32, {1}), // exact conv batch-ptr stub + DBuf(d, DType::kI32, {1}), // exact conv chunk-offset stub + false, indexed_state_io, DBuf(d, DType::kI32, {1}), // spec state-idx stub DBuf(d, DType::kI32, {1}), // spec qsl stub @@ -3315,6 +3330,17 @@ StepDevInputs BuildStepDevInputs(Dev d, const std::vector& positions, gm.prefill_has_initial_state->data()); s.has_gdn_prefill_meta = true; } + if (gm.num_prefills > 0 && gm.batch_ptr.has_value() && + gm.token_chunk_offset_ptr.has_value()) { + s.gdn_conv_batch_ptr = DBuf( + d, DType::kI32, {static_cast(gm.batch_ptr->size())}, + gm.batch_ptr->data()); + s.gdn_conv_token_chunk_offsets = DBuf( + d, DType::kI32, + {static_cast(gm.token_chunk_offset_ptr->size())}, + gm.token_chunk_offset_ptr->data()); + s.has_gdn_conv_chunks = true; + } // ── Spec-decode tensor upload (SPEC-MTP I5a). The six device tensors the GDN // spec branch of GdnBlockPaged reads (mirror qwen_gdn_linear_attn.py: // 1344-1476). Uploaded once per step from I4's builder output; NONE of this @@ -3491,11 +3517,17 @@ DBuf GdnBlockPagedMixedSpec(Dev d, const GdnLayerWeights& w, const HfConfig& cfg } DBuf dconv_ns(d, convdt, {nns_tok, conv_dim}); { + VT_CHECK(sdi.has_gdn_conv_chunks, + "gdn paged mixed spec: exact causal-conv chunks must be uploaded"); + Tensor conv_batch_ptr = sdi.gdn_conv_batch_ptr.t(); + Tensor conv_chunk_offsets = sdi.gdn_conv_token_chunk_offsets.t(); + vt::CausalConv1dArgs conv_args{true, &conv_batch_ptr, + &conv_chunk_offsets}; DBuf dcs(d, DType::kF32, {np, conv_dim, Kw - 1}); vt::GdnStateGather(d.q, dcs.t(), state.conv_state, sdi.gdn_state_idx.t()); vt::CausalConv1dFwd(d.q, dconv_ns.t(), mixed_ns.t(), dcw, nullptr, dcs.t(), sdi.gdn_non_spec_qsl.t(), sdi.gdn_has_initial.t(), - vt::CausalConv1dArgs{true}); + conv_args); Tensor conv_cache = state.conv_state; vt::GdnStateScatter(d.q, conv_cache, dcs.t(), sdi.gdn_state_idx.t()); } @@ -3749,6 +3781,12 @@ DBuf GdnBlockPaged(Dev d, const GdnLayerWeights& w, const HfConfig& cfg, // cache on CUDA → upcast; f32 cache on CPU → direct), run the f32 // CausalConv1dFwd, then downcast + scatter back to the cache. const std::vector cs_shape = {nreq, conv_dim, Kw - 1}; + VT_CHECK(sdi.has_gdn_conv_chunks, + "gdn paged: exact causal-conv chunks must be uploaded"); + Tensor conv_batch_ptr = sdi.gdn_conv_batch_ptr.t(); + Tensor conv_chunk_offsets = sdi.gdn_conv_token_chunk_offsets.t(); + vt::CausalConv1dArgs conv_args{true, &conv_batch_ptr, + &conv_chunk_offsets}; if (indexed_state_io) { VT_CHECK(sdi.has_gdn_idx && sdi.has_gdn_prefill_meta, "indexed GDN conv requires persistent non-spec metadata"); @@ -3758,7 +3796,7 @@ DBuf GdnBlockPaged(Dev d, const GdnLayerWeights& w, const HfConfig& cfg, vt::CausalConv1dFwd(d.q, dconv.t(), mixed, dcw, nullptr, dcs.t(), sdi.gdn_non_spec_qsl.t(), sdi.gdn_has_initial.t(), - vt::CausalConv1dArgs{true}); + conv_args); Tensor conv_cache = state.conv_state; vt::GdnStateScatter(d.q, conv_cache, dcs.t(), sdi.gdn_state_idx.t()); @@ -3772,7 +3810,7 @@ DBuf GdnBlockPaged(Dev d, const GdnLayerWeights& w, const HfConfig& cfg, DBuf dhis(d, DType::kI32, {nreq}, his.data()); vt::CausalConv1dFwd(d.q, dconv.t(), mixed, dcw, nullptr, dcs.t(), dqsl.t(), dhis.t(), - vt::CausalConv1dArgs{true}); + conv_args); ScatterStateF32(d, state.conv_state, dcs, sidx, conv_row_elems); } } else { @@ -5781,6 +5819,9 @@ StepDevInputs BuildFullAttnStepDevInputs(Dev d, DBuf(d, DType::kI32, {1}), // gdn_prefill_qsl stub DBuf(d, DType::kI8, {1}), // gdn_prefill_has_initial stub false, // has_gdn_prefill_meta + DBuf(d, DType::kI32, {1}), // gdn_conv_batch_ptr stub + DBuf(d, DType::kI32, {1}), // gdn_conv_token_chunk_offsets stub + false, // has_gdn_conv_chunks false, // indexed_gdn_state_io (no GDN layers) DBuf(d, DType::kI32, {1}), // gdn_spec_state_idx stub DBuf(d, DType::kI32, {1}), // gdn_spec_qsl stub @@ -7135,6 +7176,10 @@ static std::vector VLGenerateCoreGdn( g.prefill_query_start_loc = std::vector{0, static_cast(qlen)}; g.prefill_state_indices = std::vector{0}; g.prefill_has_initial_state = std::vector{0}; + const v1::CausalConv1dMetadata conv = + v1::ComputeCausalConv1dMetadata(*g.non_spec_query_start_loc); + g.batch_ptr = conv.batch_ptr; + g.token_chunk_offset_ptr = conv.token_chunk_offset_ptr; return g; }; auto gdn_decode_meta = [&]() { @@ -7565,6 +7610,8 @@ void BuildPaddedDecode(int64_t S, const std::vector& tok, gm_out.prefill_query_start_loc.reset(); gm_out.prefill_state_indices.reset(); gm_out.prefill_has_initial_state.reset(); + gm_out.batch_ptr.reset(); + gm_out.token_chunk_offset_ptr.reset(); (void)B; } @@ -7792,6 +7839,9 @@ struct Qwen3_5DecodeGraph::Impl { CopyInPlace(gdn_meta.prefill_query_start_loc, gm.prefill_query_start_loc); CopyInPlace(gdn_meta.prefill_state_indices, gm.prefill_state_indices); CopyInPlace(gdn_meta.prefill_has_initial_state, gm.prefill_has_initial_state); + CopyInPlace(gdn_meta.batch_ptr, gm.batch_ptr); + CopyInPlace(gdn_meta.token_chunk_offset_ptr, + gm.token_chunk_offset_ptr); gdn_meta.num_prefills = gm.num_prefills; gdn_meta.num_prefill_tokens = gm.num_prefill_tokens; gdn_meta.num_decodes = gm.num_decodes; @@ -8110,6 +8160,9 @@ struct Qwen3_5DenseDecodeGraph::Impl { CopyInPlace(gdn_meta.prefill_query_start_loc, gm.prefill_query_start_loc); CopyInPlace(gdn_meta.prefill_state_indices, gm.prefill_state_indices); CopyInPlace(gdn_meta.prefill_has_initial_state, gm.prefill_has_initial_state); + CopyInPlace(gdn_meta.batch_ptr, gm.batch_ptr); + CopyInPlace(gdn_meta.token_chunk_offset_ptr, + gm.token_chunk_offset_ptr); gdn_meta.num_prefills = gm.num_prefills; gdn_meta.num_prefill_tokens = gm.num_prefill_tokens; gdn_meta.num_decodes = gm.num_decodes; diff --git a/src/vllm/v1/attention/backends/gdn_attn.cpp b/src/vllm/v1/attention/backends/gdn_attn.cpp index 74ce793ad..8e84addbe 100644 --- a/src/vllm/v1/attention/backends/gdn_attn.cpp +++ b/src/vllm/v1/attention/backends/gdn_attn.cpp @@ -53,6 +53,29 @@ std::tuple SplitDecodesAndPrefills( return {num_decodes, num_prefills, num_decode_tokens, num_prefill_tokens}; } +CausalConv1dMetadata ComputeCausalConv1dMetadata( + const std::vector& qsl) { + if (qsl.empty() || qsl.front() != 0) { + throw std::invalid_argument( + "causal-conv metadata: query_start_loc must start at zero"); + } + CausalConv1dMetadata out; + for (size_t s = 0; s + 1 < qsl.size(); ++s) { + const int32_t len = qsl[s + 1] - qsl[s]; + if (len < 0) { + throw std::invalid_argument( + "causal-conv metadata: query_start_loc must be monotonic"); + } + const int32_t chunks = + (len + kCausalConv1dBlockM - 1) / kCausalConv1dBlockM; + for (int32_t chunk = 0; chunk < chunks; ++chunk) { + out.batch_ptr.push_back(static_cast(s)); + out.token_chunk_offset_ptr.push_back(chunk); + } + } + return out; +} + GDNAttentionMetadata GDNAttentionMetadataBuilder::build( int common_prefix_len, const CommonAttentionMetadata& m, bool fast_build) { // The non-spec entry point IS the spec build with both spec arguments null: @@ -331,6 +354,15 @@ GDNAttentionMetadata GDNAttentionMetadataBuilder::build( meta.prefill_state_indices = non_spec_state_indices; meta.prefill_has_initial_state = has_initial_state; } + + // The conv forward consumes the WHOLE non-spec stream, including leading + // decode rows in a mixed step. Mirror upstream's exact flattened program + // descriptor once here; the runner uploads it once and all GDN layers reuse + // it. This replaces the old n<=4 rectangular-grid approximation. + const CausalConv1dMetadata conv = + ComputeCausalConv1dMetadata(*non_spec_query_start_loc); + meta.batch_ptr = conv.batch_ptr; + meta.token_chunk_offset_ptr = conv.token_chunk_offset_ptr; } // else: has_initial_state / prefill_* stay nullopt (gdn_attn.py:405). diff --git a/src/vt/cuda/cuda_gdn.cu b/src/vt/cuda/cuda_gdn.cu index 64cde855e..9d4706323 100644 --- a/src/vt/cuda/cuda_gdn.cu +++ b/src/vt/cuda/cuda_gdn.cu @@ -701,6 +701,7 @@ void LaunchConvFwdTiled(cudaStream_t s, Tensor& out, const Tensor& x, const Tens // directly — never a value the window mutated). See gdn_prefill_conv.h. constexpr int kConvRegN = 128; // channels per block (blockDim.x; coalesced x/out) constexpr int kConvRegM = 32; // token chunk per block (grid.z parallelism) +constexpr int kConvExactM = 8; // upstream compute_causal_conv1d_metadata BLOCK_M constexpr int kConvRegMaxW = 8; // max supported width (k-1); Qwen GDN k=4 -> 3 constexpr int64_t kConvRegChunkMaxSeqs = 4; // chunk the token axis only for few seqs @@ -709,21 +710,28 @@ __global__ void CausalConv1dFwdRegKernel(Tout* out, const Tin* x, const Tin* w, const Tin* bias, float* conv_state, const int32_t* qsl, const THas* his, int64_t c_dim, int64_t x_row_stride, int64_t k, bool silu, - int chunked) { + int chunked, const int32_t* batch_ptr, + const int32_t* token_chunk_offset_ptr, int exact) { const int64_t width = k - 1; - const int64_t s = blockIdx.y; // sequence + const int64_t program = blockIdx.y; + const int64_t s = exact ? batch_ptr[program] : program; // sequence const int64_t c = static_cast(blockIdx.x) * kConvRegN + threadIdx.x; // channel const bool active = c < c_dim; const int64_t begin = qsl[s]; const int64_t t_len = qsl[s + 1] - begin; // Token range this block owns. chunked: [chunk*M, chunk*M+M); else whole sequence. - const int64_t token_offset = chunked ? static_cast(blockIdx.z) * kConvRegM : 0; + const int64_t chunk_m = exact ? kConvExactM : kConvRegM; + const int64_t token_offset = exact + ? static_cast(token_chunk_offset_ptr[program]) * + kConvExactM + : (chunked ? static_cast(blockIdx.z) * kConvRegM + : 0); // Skip chunks past the sequence end — except chunk 0, which still runs the state // write-back (also the only path when t_len == 0). if (token_offset > 0 && token_offset >= t_len) return; const int64_t token_end = - (chunked && token_offset + kConvRegM < t_len) ? token_offset + kConvRegM : t_len; + ((chunked || exact) && token_offset + chunk_m < t_len) ? token_offset + chunk_m : t_len; const bool init = his[s] != 0; float* srow = active ? conv_state + (s * c_dim + c) * width : nullptr; @@ -790,13 +798,16 @@ bool ConvRegEnabled() { return ConvRegFlagIsOn(std::getenv("VT_CONV_REG")); } -// Register-window launcher (VT_CONV_REG=1). grid = (channel-tiles, sequences, chunks). -// grid.z chunks the token axis only for few (kConvRegChunkMaxSeqs) sequences, where -// one block per (channel-tile, seq) would under-occupy on a long prefill; gridZ must -// cover the LONGEST sequence — without a host-side max we bound it by -// cdiv(total_tokens, M) (exact for n==1, over-provisions by <= n for n>1, and the -// early-return blocks are ~free). For many sequences grid.z==1 and each block streams -// its whole sequence (the channel-tile x seq grid already occupies). +bool ConvExactChunksEnabled() { + return ConvExactChunksFlagIsOn(std::getenv("VT_CONV_EXACT_CHUNKS")); +} + +// Register-window launcher (VT_CONV_REG=1). The default exact descriptor maps +// grid.y to a flattened list of (sequence, 8-token chunk) programs, mirroring +// upstream and launching neither rectangular padding nor sequence-serial work. +// VT_CONV_EXACT_CHUNKS=0 restores the legacy grid=(channel tiles, sequences, +// chunks): it chunks grid.z only for <=4 sequences, and serially streams each +// whole sequence for larger batches. template void LaunchConvFwdReg(cudaStream_t s, Tensor& out, const Tensor& x, const Tensor& w, const Tensor* bias, Tensor& conv_state, const Tensor& qsl, @@ -808,26 +819,37 @@ void LaunchConvFwdReg(cudaStream_t s, Tensor& out, const Tensor& x, const Tensor const int64_t chan_tiles = (c + kConvRegN - 1) / kConvRegN; int64_t gridZ = 1; int chunked = 0; - if (n <= kConvRegChunkMaxSeqs) { + const bool exact = ConvExactChunksEnabled() && args.batch_ptr != nullptr; + int64_t gridY = n; + if (exact) { + gridY = args.batch_ptr->shape[0]; + VT_CHECK(gridY <= kMaxGridY, + "cuda causal_conv1d_fwd(reg): too many exact chunk programs"); + } else if (n <= kConvRegChunkMaxSeqs) { const int64_t z = (total_tokens + kConvRegM - 1) / kConvRegM; if (z >= 1 && z <= kMaxGridY) { gridZ = z; chunked = 1; } } - const dim3 grid(static_cast(chan_tiles), static_cast(n), + const dim3 grid(static_cast(chan_tiles), static_cast(gridY), static_cast(gridZ)); const dim3 block(kConvRegN); + const int32_t* batch_ptr = exact ? args.batch_ptr->Ptr() : nullptr; + const int32_t* token_chunk_offset_ptr = + exact ? args.token_chunk_offset_ptr->Ptr() : nullptr; if (his.dtype == DType::kI8) { CausalConv1dFwdRegKernel<<>>( out.Ptr(), x.Ptr(), w.Ptr(), bias != nullptr ? bias->Ptr() : nullptr, conv_state.Ptr(), - qsl.Ptr(), his.Ptr(), c, x_rs, k, args.silu_activation, chunked); + qsl.Ptr(), his.Ptr(), c, x_rs, k, args.silu_activation, chunked, + batch_ptr, token_chunk_offset_ptr, exact ? 1 : 0); } else { CausalConv1dFwdRegKernel<<>>( out.Ptr(), x.Ptr(), w.Ptr(), bias != nullptr ? bias->Ptr() : nullptr, conv_state.Ptr(), - qsl.Ptr(), his.Ptr(), c, x_rs, k, args.silu_activation, chunked); + qsl.Ptr(), his.Ptr(), c, x_rs, k, args.silu_activation, chunked, + batch_ptr, token_chunk_offset_ptr, exact ? 1 : 0); } Check(cudaGetLastError(), "causal_conv1d_fwd(reg) launch"); } diff --git a/src/vt/cuda/gdn_prefill_conv.h b/src/vt/cuda/gdn_prefill_conv.h index da84b1f9d..c87277998 100644 --- a/src/vt/cuda/gdn_prefill_conv.h +++ b/src/vt/cuda/gdn_prefill_conv.h @@ -71,6 +71,14 @@ inline bool ConvRegFlagIsOn(const char* env_value) { return env_value == nullptr || env_value[0] != '0'; } +// Exact upstream (sequence, token-chunk) dispatch: DEFAULT ON, with `0` as the +// same-binary rollback to the legacy sequence-serial / rectangular-grid mapping. +// On sm_120 Qwen3.5-4B c32 the paired graph-node trace measured the causal-conv +// family at 720.047 -> 234.607 ms (3.07x), with byte-identical output tokens. +inline bool ConvExactChunksFlagIsOn(const char* env_value) { + return env_value == nullptr || env_value[0] != '0'; +} + // Pure predicate for the VT_GDN_POSTCONV_SPLIT contract: DEFAULT OFF (OPT-IN). The // split post-conv kernel (GdnPostConvSplitKernel) is BIT-IDENTICAL (0-ulp) to the // shipped GdnPostConvKernel by construction, but the DGX nsys A/B measured it diff --git a/src/vt/ops.cpp b/src/vt/ops.cpp index 988739715..127303a7e 100644 --- a/src/vt/ops.cpp +++ b/src/vt/ops.cpp @@ -1676,6 +1676,16 @@ void CausalConv1dFwd(Queue& q, Tensor& out, const Tensor& x, const Tensor& weigh const int64_t n = conv_state.shape[0]; CheckI32Meta(q, query_start_loc, n + 1, "causal_conv1d_fwd", "query_start_loc"); CheckBoolMeta(q, has_initial_state, n, "causal_conv1d_fwd", "has_initial_state"); + VT_CHECK((args.batch_ptr == nullptr) == (args.token_chunk_offset_ptr == nullptr), + "causal_conv1d_fwd: batch_ptr and token_chunk_offset_ptr must be supplied together"); + if (args.batch_ptr != nullptr) { + const int64_t programs = args.batch_ptr->shape[0]; + CheckI32Meta(q, *args.batch_ptr, programs, "causal_conv1d_fwd", "batch_ptr"); + CheckI32Meta(q, *args.token_chunk_offset_ptr, programs, "causal_conv1d_fwd", + "token_chunk_offset_ptr"); + VT_CHECK(programs > 0, + "causal_conv1d_fwd: exact chunk descriptor must not be empty"); + } reinterpret_cast(GetOp(OpId::kCausalConv1dFwd, q.device.type))( q, out, x, weight, bias, conv_state, query_start_loc, has_initial_state, args); } diff --git a/tests/examples/test_bench.cpp b/tests/examples/test_bench.cpp index bc14f34a9..72cbbbc65 100644 --- a/tests/examples/test_bench.cpp +++ b/tests/examples/test_bench.cpp @@ -1,5 +1,5 @@ // Smoke test for the M2.1 benchmark harness (examples/bench/bench_core.h): drive -// the SYNTHETIC CPU engine through the full admission + step() measurement loop +// the SYNTHETIC CPU engine through the production AsyncLLM measurement loop // and assert it produces sane metrics. The NUMBERS are meaningless (toy weights) // — this asserts the HARNESS: all N requests finish, throughput > 0, TTFT > 0, // and the token accounting is coherent. The real parity numbers come from a GB10 @@ -26,6 +26,12 @@ TEST_CASE("bench: synthetic engine completes all requests with sane metrics") { const BenchResult r = RunBench(cfg); + // The comparison harness must exercise the production AsyncLLM frontend. + // Before the SERVE-CLI-BENCH B1 repair it called synchronous + // LoadedEngine::engine() even when async scheduling resolved enabled. + CHECK(r.async_frontend); + CHECK(r.max_concurrent_batches >= 1); + // All N requests finished through the engine loop. CHECK(r.completed == cfg.num_prompts); // Wall time advanced and throughput is positive. diff --git a/tests/scripts/test_check_release_binary_contract.py b/tests/scripts/test_check_release_binary_contract.py index 8336d82ef..5cdde14de 100644 --- a/tests/scripts/test_check_release_binary_contract.py +++ b/tests/scripts/test_check_release_binary_contract.py @@ -1024,6 +1024,20 @@ def test_each_semantic_inventory_consumer_is_pinned(self) -> None: self.assertIn(inventory, result.stdout + result.stderr) def test_each_semantic_inventory_consumer_body_is_pinned(self) -> None: + loop = ast.parse( + "for item in ITEMS:\n" + " def identity(value):\n" + " return value\n" + ).body[0] + self.assertIsInstance(loop, ast.For) + baseline = checker._consumer_body_digest(loop) + nested = next( + node for node in ast.walk(loop) if isinstance(node, ast.FunctionDef) + ) + nested._fields = (*nested._fields, "type_params") + nested.type_params = [] + self.assertEqual(checker._consumer_body_digest(loop), baseline) + for inventory, method in INVENTORY_CONSUMER_METHODS.items(): for mutation in CONSUMER_FLOW_MUTATIONS: with ( diff --git a/tests/vllm/models/test_qwen27_paged_forward.cpp b/tests/vllm/models/test_qwen27_paged_forward.cpp index 2641803f1..82654bd57 100644 --- a/tests/vllm/models/test_qwen27_paged_forward.cpp +++ b/tests/vllm/models/test_qwen27_paged_forward.cpp @@ -294,6 +294,10 @@ GDNAttentionMetadata PrefillGdnMeta(int64_t T, int32_t sidx) { g.prefill_query_start_loc = std::vector{0, static_cast(T)}; g.prefill_state_indices = std::vector{sidx}; g.prefill_has_initial_state = std::vector{0}; + const auto conv = + vllm::v1::ComputeCausalConv1dMetadata(*g.non_spec_query_start_loc); + g.batch_ptr = conv.batch_ptr; + g.token_chunk_offset_ptr = conv.token_chunk_offset_ptr; return g; } @@ -351,6 +355,10 @@ GDNAttentionMetadata ChunkGdnMeta(int64_t qlen, int32_t sidx, bool has_initial) g.prefill_query_start_loc = std::vector{0, static_cast(qlen)}; g.prefill_state_indices = std::vector{sidx}; g.prefill_has_initial_state = std::vector{hi}; + const auto conv = + vllm::v1::ComputeCausalConv1dMetadata(*g.non_spec_query_start_loc); + g.batch_ptr = conv.batch_ptr; + g.token_chunk_offset_ptr = conv.token_chunk_offset_ptr; return g; } @@ -603,6 +611,10 @@ TEST_CASE("qwen27 GDN metadata validates complete prefill suffixes before I/O") gm.prefill_state_indices = std::vector{1, 2}; gm.prefill_query_start_loc = std::vector{0, 2, 5}; gm.prefill_has_initial_state = std::vector{0, 1}; + const auto conv = + vllm::v1::ComputeCausalConv1dMetadata(*gm.non_spec_query_start_loc); + gm.batch_ptr = conv.batch_ptr; + gm.token_chunk_offset_ptr = conv.token_chunk_offset_ptr; CHECK_NOTHROW(vllm::detail::ValidateGdnAttentionMetadata( gm, /*state_slots=*/3, /*allow_inert_padding=*/false)); @@ -1164,6 +1176,10 @@ TEST_CASE("qwen27 dense paged: indexed GDN mixed turnover matches row-copy fallb gm.prefill_query_start_loc = std::vector{0, 2}; gm.prefill_state_indices = std::vector{1}; gm.prefill_has_initial_state = std::vector{0}; + const auto conv = + vllm::v1::ComputeCausalConv1dMetadata(*gm.non_spec_query_start_loc); + gm.batch_ptr = conv.batch_ptr; + gm.token_chunk_offset_ptr = conv.token_chunk_offset_ptr; const std::vector ids = {4, 11, 0}; const std::vector pos = {3, 0, 1}; @@ -1246,6 +1262,10 @@ TEST_CASE("qwen27 dense paged: GDN state zeroing protects a fresh req in a mixed gm.prefill_query_start_loc = gm.non_spec_query_start_loc; gm.prefill_state_indices = std::vector{0, 1}; gm.prefill_has_initial_state = std::vector{0, 0}; + const auto conv = + vllm::v1::ComputeCausalConv1dMetadata(*gm.non_spec_query_start_loc); + gm.batch_ptr = conv.batch_ptr; + gm.token_chunk_offset_ptr = conv.token_chunk_offset_ptr; const std::vector batch = Qwen3_5DenseModel::Forward( ids, pos, am, gm, pool.attn_kv, pool.gdn_state, w, c, q); diff --git a/tests/vllm/models/test_qwen35_paged_forward.cpp b/tests/vllm/models/test_qwen35_paged_forward.cpp index 4b8c2c2e8..d7d5dce4b 100644 --- a/tests/vllm/models/test_qwen35_paged_forward.cpp +++ b/tests/vllm/models/test_qwen35_paged_forward.cpp @@ -266,6 +266,10 @@ GDNAttentionMetadata PrefillGdnMeta(int64_t T, int32_t sidx) { g.prefill_query_start_loc = std::vector{0, static_cast(T)}; g.prefill_state_indices = std::vector{sidx}; g.prefill_has_initial_state = std::vector{0}; + const auto conv = + vllm::v1::ComputeCausalConv1dMetadata(*g.non_spec_query_start_loc); + g.batch_ptr = conv.batch_ptr; + g.token_chunk_offset_ptr = conv.token_chunk_offset_ptr; return g; } @@ -441,6 +445,10 @@ TEST_CASE("qwen35 paged: GDN state zeroing protects a fresh req in a mixed batch gm.prefill_query_start_loc = gm.non_spec_query_start_loc; gm.prefill_state_indices = std::vector{0, 1}; gm.prefill_has_initial_state = std::vector{0, 0}; + const auto conv = + vllm::v1::ComputeCausalConv1dMetadata(*gm.non_spec_query_start_loc); + gm.batch_ptr = conv.batch_ptr; + gm.token_chunk_offset_ptr = conv.token_chunk_offset_ptr; const std::vector batch = Qwen3_5Model::Forward( ids, pos, am, gm, pool.attn_kv, pool.gdn_state, w, c, q); diff --git a/tests/vllm/models/test_qwen3_5_gdn_spec_routing.cpp b/tests/vllm/models/test_qwen3_5_gdn_spec_routing.cpp index 1a204d2de..028fd3226 100644 --- a/tests/vllm/models/test_qwen3_5_gdn_spec_routing.cpp +++ b/tests/vllm/models/test_qwen3_5_gdn_spec_routing.cpp @@ -259,6 +259,10 @@ GDNAttentionMetadata MixedMeta(int Tp) { g.prefill_state_indices = std::vector{2}; g.prefill_query_start_loc = std::vector{0, Tp}; g.prefill_has_initial_state = std::vector{0}; + const auto conv = + vllm::v1::ComputeCausalConv1dMetadata(*g.non_spec_query_start_loc); + g.batch_ptr = conv.batch_ptr; + g.token_chunk_offset_ptr = conv.token_chunk_offset_ptr; return g; } @@ -274,6 +278,10 @@ GDNAttentionMetadata PrefillMeta(int Tp, int slot) { g.prefill_state_indices = std::vector{slot}; g.prefill_query_start_loc = std::vector{0, Tp}; g.prefill_has_initial_state = std::vector{0}; + const auto conv = + vllm::v1::ComputeCausalConv1dMetadata(*g.non_spec_query_start_loc); + g.batch_ptr = conv.batch_ptr; + g.token_chunk_offset_ptr = conv.token_chunk_offset_ptr; return g; } diff --git a/tests/vllm/v1/attention/test_gdn_metadata_builder.cpp b/tests/vllm/v1/attention/test_gdn_metadata_builder.cpp index 440f062b6..36c932763 100644 --- a/tests/vllm/v1/attention/test_gdn_metadata_builder.cpp +++ b/tests/vllm/v1/attention/test_gdn_metadata_builder.cpp @@ -22,6 +22,7 @@ #include "vllm/v1/worker/gpu/prepare_inputs.h" using vllm::v1::CommonAttentionMetadata; +using vllm::v1::ComputeCausalConv1dMetadata; using vllm::v1::GDNAttentionBackend; using vllm::v1::GDNAttentionMetadata; using vllm::v1::GDNAttentionMetadataBuilder; @@ -116,6 +117,16 @@ TEST_CASE("GDN build: mixed decode + prefill (decode-first)") { CHECK(*meta.prefill_has_initial_state == std::vector{0}); } +TEST_CASE("GDN metadata: causal-conv programs enumerate each sequence chunk once") { + // Unequal lengths distinguish the exact flattened descriptor from a rectangular + // sequence x chunk grid. BLOCK_M=8 yields 1+2+3 programs, with no padded work. + const auto chunks = ComputeCausalConv1dMetadata( + std::vector{0, 1, 10, 27}); + CHECK(chunks.batch_ptr == std::vector{0, 1, 1, 2, 2, 2}); + CHECK(chunks.token_chunk_offset_ptr == + std::vector{0, 0, 1, 0, 1, 2}); +} + // Decode-only batch: all query_len==1. num_prefills==0 ⇒ has_initial_state and // all prefill_* fields are None (gdn_attn.py:405) — the decode kernel reads the // state via state_indices and needs no has_initial_state mask. @@ -504,6 +515,8 @@ void CheckMetaEqual(const GDNAttentionMetadata& a, const GDNAttentionMetadata& b CHECK(a.prefill_query_start_loc == b.prefill_query_start_loc); CHECK(a.prefill_state_indices == b.prefill_state_indices); CHECK(a.prefill_has_initial_state == b.prefill_has_initial_state); + CHECK(a.batch_ptr == b.batch_ptr); + CHECK(a.token_chunk_offset_ptr == b.token_chunk_offset_ptr); CHECK(a.spec_query_start_loc == b.spec_query_start_loc); CHECK(a.spec_state_indices_tensor == b.spec_state_indices_tensor); CHECK(a.spec_state_indices_num_cols == b.spec_state_indices_num_cols); diff --git a/tests/vt/test_gdn_prefill_conv.cpp b/tests/vt/test_gdn_prefill_conv.cpp index 2a6423743..eb904fdd3 100644 --- a/tests/vt/test_gdn_prefill_conv.cpp +++ b/tests/vt/test_gdn_prefill_conv.cpp @@ -18,6 +18,7 @@ #include "vt/cuda/gdn_prefill_conv.h" using vt::cuda::ConvRegFlagIsOn; +using vt::cuda::ConvExactChunksFlagIsOn; using vt::cuda::GdnPostConvFastFlagIsOn; using vt::cuda::GdnPostConvSplitFlagIsOn; @@ -40,6 +41,15 @@ TEST_CASE("VT_CONV_REG defaults ON; only a '0'-leading value rolls back") { CHECK_FALSE(ConvRegFlagIsOn("00")); } +TEST_CASE("VT_CONV_EXACT_CHUNKS defaults ON; only a '0'-leading value rolls back") { + CHECK(ConvExactChunksFlagIsOn(nullptr)); + CHECK(ConvExactChunksFlagIsOn("")); + CHECK_FALSE(ConvExactChunksFlagIsOn("0")); + CHECK_FALSE(ConvExactChunksFlagIsOn("0abc")); + CHECK(ConvExactChunksFlagIsOn("1")); + CHECK(ConvExactChunksFlagIsOn("on")); +} + TEST_CASE("VT_GDN_POSTCONV_SPLIT defaults OFF (opt-in); a non-'0' value enables it") { // Default (unset) is OFF: GdnPostConvSplitKernel is BIT-IDENTICAL (0-ulp) to the // shipped GdnPostConvKernel by construction (byte-for-byte q/k L2-norm branch; same diff --git a/tests/vt/test_ops_gdn.cpp b/tests/vt/test_ops_gdn.cpp index 538450165..f952b20d5 100644 --- a/tests/vt/test_ops_gdn.cpp +++ b/tests/vt/test_ops_gdn.cpp @@ -2416,25 +2416,56 @@ void RunConvFwdRegByteExactCase(const std::vector& qsl, const std::vect i8_mask ? static_cast(his_i8.data()) : static_cast(his.data())); - auto run = [&](bool reg, std::vector& out_bytes, std::vector& st_bytes) { + std::vector batch_ptr; + std::vector chunk_offsets; + for (int32_t s = 0; s < n; ++s) { + const int32_t len = qsl[static_cast(s + 1)] - + qsl[static_cast(s)]; + const int32_t chunks = (len + 7) / 8; // upstream BLOCK_M=8 + for (int32_t chunk = 0; chunk < chunks; ++chunk) { + batch_ptr.push_back(s); + chunk_offsets.push_back(chunk); + } + } + DeviceTensor dbatch(gpu, gq.q, DType::kI32, + {static_cast(batch_ptr.size())}, + batch_ptr.data()); + DeviceTensor doffsets(gpu, gq.q, DType::kI32, + {static_cast(chunk_offsets.size())}, + chunk_offsets.data()); + + auto run = [&](bool reg, bool exact, std::vector& out_bytes, + std::vector& st_bytes) { ::setenv("VT_CONV_REG", reg ? "1" : "0", 1); + ::setenv("VT_CONV_EXACT_CHUNKS", exact ? "1" : "0", 1); DeviceTensor dst(gpu, gq.q, DType::kF32, {n, c, k - 1}, stb.data()); // fresh state per arm DeviceTensor dout(gpu, gq.q, cb.out, {t, c}); gpu.Memset(gq.q, dout.tensor().data, 0x5a, static_cast(t * c) * vt::SizeOf(cb.out)); + Tensor batch_tensor = dbatch.tensor(); + Tensor offsets_tensor = doffsets.tensor(); + CausalConv1dArgs run_args{silu}; + if (exact) { + run_args.batch_ptr = &batch_tensor; + run_args.token_chunk_offset_ptr = &offsets_tensor; + } vt::CausalConv1dFwd(gq.q, dout.tensor(), dx.tensor(), dw.tensor(), with_bias ? &db.tensor() : nullptr, dst.tensor(), dqsl.tensor(), - dhis.tensor(), args); + dhis.tensor(), exact ? run_args : args); out_bytes.resize(static_cast(t * c) * vt::SizeOf(cb.out)); st_bytes.resize(stb.size()); dout.Download(gq.q, out_bytes.data()); dst.Download(gq.q, st_bytes.data()); }; - std::vector out_tiled, st_tiled, out_reg, st_reg; - run(/*reg=*/false, out_tiled, st_tiled); - run(/*reg=*/true, out_reg, st_reg); + std::vector out_tiled, st_tiled, out_reg, st_reg, out_exact, st_exact; + run(/*reg=*/false, /*exact=*/false, out_tiled, st_tiled); + run(/*reg=*/true, /*exact=*/false, out_reg, st_reg); + run(/*reg=*/true, /*exact=*/true, out_exact, st_exact); ::unsetenv("VT_CONV_REG"); + ::unsetenv("VT_CONV_EXACT_CHUNKS"); CHECK(out_reg == out_tiled); // out activation byte-identical CHECK(st_reg == st_tiled); // rolled conv_state byte-identical + CHECK(out_exact == out_tiled); // exact descriptor changes only work assignment + CHECK(st_exact == st_tiled); } TEST_CASE("CUDA causal_conv1d_fwd register kernel (VT_CONV_REG) matches tiled 0-ulp") { diff --git a/tools/bench/run_qwen35_4b_compare.sh b/tools/bench/run_qwen35_4b_compare.sh index 9e88e89e3..f61bfe1d2 100755 --- a/tools/bench/run_qwen35_4b_compare.sh +++ b/tools/bench/run_qwen35_4b_compare.sh @@ -63,7 +63,11 @@ if test -n "${VLLM_CUDA_HOME:-}"; then vllm_path=$cuda_combined/bin:$(dirname "$ninja"):$(dirname "$host_cxx"):$PATH vllm_ld_library_path=$(dirname "$libstdcpp"):$cuda_combined/lib:/run/opengl-driver/lib vllm_cpath=$cuda_combined/include - vllm_library_path=$cuda_combined/lib + # FlashInfer JIT links the CUDA driver with plain `c++ -lcuda`. The venv + # toolkit carries cudart but the live driver belongs to the host, so expose + # both roots to the compiler's library search (LD_LIBRARY_PATH alone is only + # a runtime lookup and does not satisfy this link step). + vllm_library_path=$cuda_combined/lib:/run/opengl-driver/lib vllm_nix_ldflags="-L$cuda_combined/lib -L/run/opengl-driver/lib" else cudart=$(sed -n 's/^CUDA_CUDART:[^=]*=//p' "$cmake_cache")