diff --git a/.agents/parity-ledger.md b/.agents/parity-ledger.md index 258035eaf..514b186c7 100644 --- a/.agents/parity-ledger.md +++ b/.agents/parity-ledger.md @@ -838,3 +838,9 @@ Columns: | 2026-07-29 (**DeepSeek-V4-Flash W2c — rewire the forward to CONSUME the keep-quant tower (the last CODE brick before the single-Spark run; fixes the ~1 TiB f32-tower OOM W8-final proved)**; `CLAIM-DEEPSEEK-V4-W2C`; isolated worktree `/tmp/.../scratchpad/wt-w2c` (branch `claim-deepseek-v4-w2c`) off local `main` `328e6a50` confirmed via `git rev-parse HEAD`; CPU-only Release build + doctest gate; NO GPU/download; additive forward/loader edits + test + records; NOT pushed; closing commit: this checkpoint) | **What it does.** (1) `LoadDeepseekV4FromGguf` (`deepseek_v4_weights.cpp`) NO LONGER f32-expands the big MLA/MoE/lm_head weights into `weights.host` — only the small NON-GEMM tensors (norms, sinks, MHC/DSA mixing, ape, `tid2eid`/`exp_probs_b`, weights_proj, embed-gather) dequant; a load-time `VT_CHECK` asserts every big host slot is EMPTY (a regression that rebuilds the f32 tower fails LOUDLY); adds `DeepseekV4{Host,Gguf}ResidentBytes`. (2) `deepseek_v4.cpp` adds a `V4Backend::gguf` weight source + three keep-quant GEMM helpers — `Gemm` (`Y=X@W^T` off `OwnedTensor.View()` via `vt::MatmulBT`, which dispatches to the CPU `kMatmulBTQuant` CIQ GEMM for a block weight — no f32 expansion; f32-MatVec fallback bit-identical for the host path), `GemmRowSlice` (per-expert/per-group whole-block row slice), `GroupedOutputLoraGguf` (per-group wo_a slice + wo_b). `ForwardComposeImpl` routes the 512-wide MLA linears (wq_a/wq_b/wkv/wo_a/wo_b), compressor/indexer projections, router gate, shared + 256 routed experts, and lm_head to the COMPRESSED `weights.gguf` blocks; the small tensors stay f32 from the small host tower. New public `DeepseekV4ForwardGguf`; `DeepseekV4Model::Forward` gates on `has_gguf_weights` (safetensors/NVFP4 + tiny-synthetic host path byte-identical). | Mirrors how our other GGUF models run keep-quant in the forward without a full f32 tower: `src/vllm/model_executor/models/qwen3_5.cpp:786-838` (host `MatmulBT` off an `OwnedTensor.View()`), `qwen3_5_gguf_weights.cpp:57-101` (`OwnGgufQuantBlocks` row slice), `src/vt/ops.cpp:134-201` (`MatmulBT`→`MatmulBTQuant` block-quant dispatch), `src/vt/cpu/cpu_quant_gemm.cpp` (the CIQ GEMM). vLLM structural ground: `models/deepseek_v4/nvidia/model.py:1080-1148` (forward) + `nvidia/ops/o_proj.py:58-73` (grouped o-LoRA). Edited: `include/vllm/model_executor/models/deepseek_v4.h` + `src/.../deepseek_v4.cpp` + `src/.../deepseek_v4_weights.cpp` + `tests/vllm/models/test_deepseek_v4_gguf_load.cpp`. | **PASS (CPU keep-quant forward + memory gate) — `tests/vllm/models/test_deepseek_v4_gguf_load.cpp` 7 cases / 185 assertions GREEN** (CPU Release full-library `-Werror`-clean; +2 cases over the W8-final 6/168). (a) `DeepseekV4ForwardGguf` over the Q8_0 keep-quant tower RUNS finite + deterministic. (b) **near-tie (RED-first):** keep-quant (Q8_0 vec_dot) vs the dequant oracle (bf16 `MatmulBT`, the expand-all load) RelL2 = **0.0116** (< 0.05 band, ~4x margin; the only difference is the keep path also quantizes the activation to q8_0); a `kNoAttnSink` miswire DIVERGES (RelL2 0.122). (c) **MEMORY-BOUNDED:** at tiny shape host **23,980 B** < keep-quant **141,676 B** (weights NOT in host); a test asserts the big host slots empty + `host_bytes < the f32-expanded big-weight image`. **PROJECTION to the real config:** the 256 routed experts alone = 3·256·2048·4096·43·4 B ≈ **~1032 GiB** f32 (OOM-reboots the 119 GiB pool, the bug) vs the keep-quant `UD-IQ2_XXS` **~90.9 GiB** + small f32 host **< 3 GiB** = **93.9 GiB < 119 GiB → memory-FEASIBLE** on ONE GB10. **SACRED inertness PROVEN:** only `deepseek_v4.{h,cpp}` + `deepseek_v4_weights.cpp` + the test changed; the W3-W6 primitive host references (the correctness oracle) UNTOUCHED and still pass (dsa 13/38, mhc 14/125, moe 12/716, compressor 12/164, forward 6/26, scaffold 4/40); shared MLA/MoE + CUDA W7-device empty-diff; `test_cuda_deepseek_v4` compiles + skips on the CPU box. **NOT APPLICABLE for speed** (memory/correctness enabler on CPU at tiny shape; `benchmark_binding=false`). REUSES the landed `kMatmulBTQuant` — no new kernel-matrix row ⇒ no `check-agent-record.py` count bump. Honest 3-state: tiny keep-quant forward = DERIVED + BUILD-VERIFIED; the real 91 GB `UD-IQ2_XXS` run (download + GB10 generate + benchmark) stays the operational W8-run, now MEMORY-FEASIBLE. Model-matrix DeepSeek-V4 row STAYS `SPIKE`. Not pushed; FULL SHA reported. | | 2026-07-29 (**xgrammar structured-output backend W0 spike + W1 CPU brick**; `CLAIM-TOOLS-XGRAMMAR`; isolated worktree `/home/mudler/_git/vllm.cpp-xgrammar` (branch `tools-xgrammar`) off local `main` `6d54f242` confirmed via `git rev-parse HEAD`; CPU-only Release `-Werror` (`build-cpu`, `-DVLLM_CPP_CUDA=OFF`); no GPU/download (host-side masking); additive backend + converter + test + records; NOT pushed; closing commit: this checkpoint) | **What it does.** Opens HIGH-priority feature-gap #4 (the xgrammar structured-output backend, vLLM's DEFAULT `auto`). Adds `XgrammarStructuredOutputBackend` as a 2nd registerable `StructuredOutputBackend` behind the shared seam. **§9 DECISION (recorded): mirror xgrammar's algorithm PORTABLY, do NOT vendor the mlc-ai/xgrammar C++ library** — xgrammar IS "grammar → pushdown automaton → per-step token bitmask", already ours (the native pushdown-FSM + token-byte trie); the backend REUSES that matcher and ports only the xgrammar-faithful front-end where the two diverge: the JSON-schema→EBNF converter (`xgrammar_json_schema.{h,cpp}`) preserving property DECLARATION order (`nlohmann::ordered_json`, vs the native path's lexicographic sort) + the `any_whitespace` `ws` rule + the `basic_*` set VERBATIM — closing the key-order/whitespace/exotic-schema parity gap. GRAMMAR/REGEX/CHOICE/STRUCTURAL_TAG delegate to the native compile paths. Backend selection mirrored: `ResolveStructuredOutputBackend` (`auto`→`xgrammar`) + `MakeStructuredOutputBackendFactory` + `MakeXgrammarBackendFactory`. New: `include/vllm/v1/structured_output/{backend_xgrammar,xgrammar_json_schema}.h` + `src/vllm/v1/structured_output/{backend_xgrammar,xgrammar_json_schema}.cpp` + `tests/vllm/v1/structured_output/test_backend_xgrammar.cpp` + CMake wiring. | `vllm/v1/structured_output/backend_xgrammar.py:36,78,128` (`XgrammarBackend.compile_grammar`/`allocate_token_bitmask`); JSON→EBNF SEMANTICS from xgrammar `cpp/json_schema_converter.cc` @ mlc-ai/xgrammar `a32ac892676d2eedc0327416105b9b06edfb94b2` (the SHA cited across `backend_xgrammar.py`); selection `vllm/sampling_params.py:932-949,1024-1061`, `vllm/v1/structured_output/__init__.py:133-165`, `vllm/config/structured_outputs.py:13`. All @ pin `555967922` (vLLM 0.26.0.dev0). Test intent from `tests/entrypoints/llm/test_struct_output_generate.py:214`. | **PASS (correctness brick, host CPU; `benchmark_binding=false` — structured output is host-side per-step masking, no throughput owed; the only GPU work is the already-gated `apply_grammar_bitmask` scatter, unchanged).** `tests/vllm/v1/structured_output/test_backend_xgrammar.cpp` **6 cases / 39 assertions GREEN**: (1) the seam produces a per-step bitmask allowing EXACTLY the grammar-valid next tokens; (2) **key-order parity + RED-first** — after `{"` the xgrammar (declaration-order) grammar allows `n` (name) and forbids `a` (age), while the native (key-sorting) backend on the SAME schema admits `a` and forbids `n`, i.e. a sorting compiler's mask allows a token xgrammar forbids; (3) `disable_any_whitespace` collapses the `ws` rule; (4) json_object opens with `{`; (5) the converter EBNF carries `basic_*` + declaration order; (6) `auto`→xgrammar selection constrains JSON. **RED-first PROVEN operationally**: swapping the converter's `ordered_json`→`json` (sorted keys) rebuilds green but the gate goes 3/6 cases FAIL (5 assertions, first `CHECK(Allowed(*gx,kN))`); reverting → 6/6. No regression: `test_backend_native` 4878, `test_structured_output` 90, `test_json_schema_to_gbnf` 656 unchanged. Clean CPU `-Wall -Wextra -Werror` 0-warn (both new TUs recompiled from scratch). No new stable row (advanced existing `TOOLS-XGRAMMAR` `INVENTORIED`→`ACTIVE`) ⇒ no `check-agent-record.py` ENGINE count bump. Record checkers rc=0. Honest residuals (spec §Work breakdown): optional object properties + strict-compact separators (W2); the `has_xgrammar_unsupported_json_features` guard + `validate_xgrammar_grammar` feeding the `auto` fallback + `model_loader.cpp` production wiring (W2); xgrammar-specific regex/structural-tag parity (W3); GPU oracle parity (W4, DGX offline). Not pushed; FULL SHA reported. | | 2026-07-29 (**fp8 KV cache W0 spike + W1 CPU brick — HIGH-priority feature-gap #5, the standard memory/throughput lever that halves the KV footprint**; `CLAIM-KV-FP8`; isolated worktree `.claude/worktrees/kv-fp8-w1` (branch `kv-fp8-w1`) off local `main` `6d54f242` confirmed via `git rev-parse HEAD`; CPU-only Release `-Werror`; NO GPU/download; additive vt-op + CPU-kernel + config edits + one new test + records; NOT pushed; closing commit: this checkpoint) | **What it does.** Stores the paged K/V cache as fp8-e4m3 with a per-tensor dequant scale and dequants on the attention read, mirroring vLLM 1:1. (1) `include/vt/fp8_kv.h` (NEW) — the `Fp8KVCacheDataType` interpretation enum (mirror `csrc/attention/dtype_fp8.cuh:9-13`) + the fp8-e4m3 codec + `StoreKvFp8E4M3`/`LoadKvFp8E4M3` with the vLLM scale convention `FP8=Quantize(hp/scale)` / `Dequant(fp8)*scale` (`csrc/quantization/w8a8/fp8/nvidia/quant_utils.cuh:296-308`), bit-identical to the landed `vllm::F32ToF8E4M3`/`F8E4M3ToF32` (vt does not depend on vllm). (2) `vt::ReshapeAndCacheFp8` (NEW op `kReshapeAndCacheFp8`, wrapper + CPU kernel `src/vt/cpu/cpu_cache.cpp`) — the fp8 store, mirroring the fp8 branch of `reshape_and_cache_flash_kernel` (`csrc/libtorch_stable/cache_kernels.cu:241-252,314-401`). (3) additive default-inert `PagedAttentionArgs.{kv_cache_dtype,k_scale,v_scale}` + the CPU paged-attention read dequant (`src/vt/cpu/cpu_paged_attn.cpp`). (4) `vllm::v1::ParseCacheDType`/`IsQuantizedKvCache` (`include/vllm/v1/kv_cache_dtype.h`) — the `cache_dtype` config parse mirroring `CacheDType` (`vllm/config/cache.py:19-36,76`) + `is_quantized_kv_cache`. Storage is 1-byte fp8 as `DType::kI8` (the dtype.h "the byte never guesses its semantic type" rule), NOT a new `DType` enumerator. | vLLM 1:1: config `vllm/config/cache.py:19-36,76`; scale handling `vllm/model_executor/layers/quantization/kv_cache.py:42,108-191` (per-tensor k/v scale, default 1.0, checkpoint-loaded); store `cache_kernels.cu:241-252,314-401`; read/dequant `quant_utils.cuh:296-308`; interpretation enum `dtype_fp8.cuh:9-13`. Reuses the landed e4m3 math (no re-port). Tests ported from `vllm/tests/kernels/attention/test_cache.py::test_reshape_and_cache` (the fp8 branch, `:97-165`). | **PASS (CPU correctness brick) — `tests/vt/test_ops_fp8_kv_cache.cpp` 8 cases / 511 assertions GREEN** (CPU full-library `-Wall -Wextra -Werror` clean): store→read round-trips within the e4m3 band (upstream `atol=0.001, rtol=0.1`), fp8 tracks the bf16 KV baseline (NMSE < 1%), paged-attention over the fp8 cache matches the bf16-cache output within 5%, `ParseCacheDType` mirrors the CacheDType surface, e5m2 refused (later brick), auto (no-dequant) read of an fp8 cache refused. **RED-first PROVEN:** a wrong store direction (`hp*scale` instead of `hp/scale`) fails 3 cases / 480 assertions; a wrong read `v_scale` diverges > 0.05 from the baseline. **No sibling regression:** `test_ops_reshape_cache` 12/12, `test_ops_paged_attn` 14/14 (the additive default-inert fields keep every float caller byte-identical). **NOT APPLICABLE for speed** (CPU codec brick; `benchmark_binding=false`). **RESIDUALS (honest, named W2-W5 in the spec):** the CUDA fp8 store + fp8 paged-attention read (the GPU memory-halving path, DGX-blocked); the runner/spec integration (half-sized KV blocks + checkpoint `k_scale`/`v_scale` threading + `--kv-cache-dtype`/`--calculate-kv-scales`); fp8_e5m2 CPU compute + per-attention-head scales. No new counted-matrix ROW ⇒ no `check-agent-record.py` count bump; `KV-FP8` engine row INVENTORIED→ACTIVE, `QUANT-KV-FP8` INVENTORIED→PARTIAL. Not pushed; FULL SHA reported. | +| 2026-07-27 (**Local Qwen3.5-4B post-pull REVALIDATION on discrete Blackwell + a VOIDED prior series and the harness hole that produced it**; `CLAIM-LOCAL-4B-POSTPULL`; local workspace host only (RTX 5070 Ti sm_120 16 GiB, driver 595.71.05), no dgx, no worktree (single-agent task on branch `bench-lever-sampled-token-20260727` off `main` `7f620e74`); ALL 18 benchmark legs + the correctness tier under ONE `flock /tmp/gpu`; oracle = local `.venv-vllm` **vLLM 0.24.0**, which is BEHIND the project parity pin `555967922`/0.26.0.dev0 and is labelled as such; NOT pushed; closing commit: this checkpoint) | **What it does.** Discharges the revalidation `docs/BENCHMARKS.md` recorded as PENDING for the `c317237a` transplant (whose two commits, `a131de03` + `b6f1efc8`, are now upstream, so the measured tree is plain current `main`), and repairs three measurement defects found doing it. (1) **Harness idle gate**: `prepare_leg` gated GPU idleness on `nvidia-smi --query-compute-apps`, which enumerates CUDA contexts ONLY — a GRAPHICS consumer is invisible to it. All nine performance legs of the 2026-07-25 binding series ran at 11-13% utilization with 611 MiB extra resident VRAM and PASSED that gate; today's nine ran at 0%. Now also gated on `utilization.gpu` (`GPU_IDLE_UTIL_MAX`, default 2%); re-parsed it reads 0 for today and 12 for 2026-07-25. (2) **Summarizer**: read historical token legs under a `perf-` prefix the harness has never written (every historical comparison raised `FileNotFoundError`), and required a `vllm_production` key its own current output does not emit — so the tool could not read its own aggregates. Both fixed. (3) **NEW same-binary A/B harness** `tools/bench/run_qwen35_4b_ab.sh` (one lock across the series, arms interleaved with the order flipped on even reps so neither arm is systematically the warmer one, per-leg token ids captured so an arm that changed the OUTPUT is caught rather than celebrated, same idle gate). Also corrects the oracle version label: the local venv is 0.24.0, not the 0.25.0 the prior evidence claimed — that series' own `vllm-version.txt` already said 0.24.0. | Ours: `tools/bench/run_qwen35_4b_compare.sh` (`prepare_leg`), `tools/bench/summarize_qwen35_4b_compare.py`; the residual's upstream mirror is `vllm/v1/worker/gpu/states.py:64` (`last_sampled_tokens` is a GPU tensor unconditionally, no integrated/discrete branch) + `states.py:132` (`remove_request` frees an index; upstream NEVER condenses, which is why its device-resident state needs no permutation) + `input_batch.py:296-360,449-473` (the combine/scatter kernels we already ported), read from the installed release tree `${VLLM_SOURCE}` = `.venv-vllm/.../site-packages/vllm` (0.24.0, cited as installed-release source, NOT as the pin). | **How verified (idle box, one lock, 3 reps/arm).** **Correctness UNCHANGED bit for bit:** `test_qwen35_plain_weights --no-skip` 3/3 (1672 asserts), `test_ops_gdn` 66/66 (4242), `test_ops_paged_attn` 25/25 (454,474), `test_gdn_packed_decode_triton` 1/1 (10), `test_combine_tokens` 7/7 (14), `test_input_batch` 22/22 (163) — every count identical to the 2026-07-25 record. **Output token-IDENTICAL to that series: 128/128 requests per repetition for BOTH direct-ON and direct-OFF**, so 109 upstream commits moved no token on this workload; our arm is deterministic rep-to-rep and ON==OFF 128/128, while the oracle is NOT self-deterministic (102/128 on one pair). **Speed (means, ours ON / vLLM 0.24.0):** total tput 6600.66 / 6722.24 = **0.9819x FAIL** (was 0.9864x), output tput 729.88 / 743.33 = 0.9819x FAIL, req/s 5.700 / 5.807 = 0.9815x FAIL, mean TTFT 729.24 / 913.55 = **0.7982x PASS**, mean TPOT/ITL 38.22 / 33.53 = **1.1400x FAIL** (was 1.1341x). Spread across repetitions is 0.13% (ours: 6602.6/6603.9/6595.5) and 0.17% (vLLM: 6716.3/6722.9/6727.6), so these ratios are NOT noise. Memory PASSES: peak PSS 2.282 vs OFF 8.594 (-73.4%), stable PSS 0.761 vs 8.591 (-91.1%), peak VRAM 12850.0 vs vLLM 12933.3; direct-ON now also cuts mean TTFT 12.7% vs OFF, which the contended series could not resolve. **The 2026-07-25 ABSOLUTE numbers are VOID** (contended series, see above); its RATIOS, its same-binary component attributions (H32 AOT +4.5906%, decode graph +0.3873%, ratio-4 FA2 +1.6004%) and its profiling attribution SURVIVE, each having been internal to one uniformly-contended series. **Residual, unchanged and now specified:** TPOT fails because this GPU is DISCRETE — `is_integrated_gpu()` is false, so both ENG-ASYNC-SCHED W3 device call sites (`runner.cpp:821` combine, `runner.cpp:1763` scatter) take their host fallback and `sample_tokens_async` must synchronize the main stream (497 syncs, 20.975 s, 42.20 ms/call in the prior node-mode trace); the async scheduler IS engaged (`max_concurrent_batches=2`) so depth-2 overlaps nothing here. Scoped as ENG-ASYNC-SCHED W4 in [async-discrete-device-combine.md](specs/async-discrete-device-combine.md), including the second per-step barrier that must go with it (the CUDA embedding out-of-range flag's `cudaMalloc`+`cudaStreamSynchronize`+`cudaFree`, 531 calls). `benchmark_binding=true` for the 4B local row ONLY; no 4B result implies support or speed for the 27B/35B gates, which stay hardware-unavailable here. Evidence root `/tmp/qwen35-postpull-7f620e74`, index [qwen35-4b-postpull-20260727.md](../docs/bench-evidence/qwen35-4b-postpull-20260727.md). NOTE: `scripts/check-agent-record.py` reports 6 errors on PRISTINE `main` (DONE rows citing closing commits `164453a2`/`7a3f04b2` absent from this clone) — verified identical with and without this change, so pre-existing and not introduced here. Not pushed; FULL SHA reported. | +| 2026-07-27 (**ENG-ASYNC-SCHED W4 — discrete-CUDA device-resident sampled tokens IMPLEMENTED + correctness-GATED, opt-in DEFAULT OFF; and the 2026-07-25 attribution that motivated it DISPROVED**; `CLAIM-ASYNC-SCHED-W4`; local workspace host only (RTX 5070 Ti sm_120 16 GiB, discrete — `is_integrated_gpu()` FALSE), no dgx, single-agent on branch `bench-lever-sampled-token-20260727` off `main` `7f620e74`; every GPU run under one `flock /tmp/gpu`; NOT pushed; closing commit: this checkpoint) | **What it does.** Gives the discrete-CUDA async path the device-resident `last_sampled_tokens` upstream keeps on EVERY platform (`states.py:64`), so the sampled ids need not round-trip the host: runner-owned device buffers for `last_sampled`/`prefill_len`/`query_start_loc`/`seq_lens`/`input_ids`, the W3 combine/scatter kernels pointed at them, `ModelForwardInput::device_token_ids` so the Qwen3.5 forward embeds the patched DEVICE ids, and a NEW `vt::cuda::LaunchApplyLastSampledOps` replaying `InputBatch`'s structural row edits (seed/move/swap, recorded in a new `last_sampled_ops` log) onto the mirror in stream order — which is how a device-resident buffer survives our `condense()`, something upstream never needs because `states.py:132` frees slot INDICES into a pool and never permutes. W4e additionally removes the per-step embedding barrier: `EmbeddingKernelCuda` no longer does `cudaMalloc` + `cudaStreamSynchronize` + `cudaFree` for its out-of-range flag on every call, but uses a persistent RING of flag slots with a deferred, `cudaEventQuery`-gated check; the reporting contract changed deliberately (an out-of-range id is raised no later than the NEXT embedding, same message, same id; the gather always clamped, so the offending call never read out of bounds) and its CUDA test was rewritten to that contract. THREE design corrections forced during implementation, all recorded in the spec: the embed PATCHES a prefix rather than embedding the runner's buffer (the decode graph embeds a PADDED id vector); the per-step uploads copy from PAGEABLE host memory ON PURPOSE (a shared pinned staging buffer is a RACE — pinned copies are truly async, so the next upload can overwrite bytes an in-flight DMA has not read, and with depth-2 that window spans steps); and the forward reads the ids through an RAII scoped override consumed on first use instead of a parameter threaded through five entry points plus the decode-graph class. | `vllm/v1/worker/gpu/states.py:64` (`last_sampled_tokens` is a GPU tensor unconditionally — no integrated/discrete branch exists upstream), `states.py:132` (`remove_request` frees an index; never condenses), `vllm/v1/worker/gpu/input_batch.py:296-360,449-473` (the combine + post_update kernels we already ported), read from `${VLLM_SOURCE}` = the installed release tree (vLLM 0.24.0, cited as installed-release source, NOT the pin `555967922`). | **How verified (RED-first by accident, then by design).** The FIRST gate run FAILED loudly and productively: token identity mirror-ON vs OFF was **0/128**, output garbage from the first decode. Instrumenting the branch showed the mirror stayed all-zero because the scatter in `sample_tokens_async` NEVER EXECUTED — `vllm-bench` drives the SYNCHRONOUS `LLMEngine::step()` loop (`sample_tokens()`), not `AsyncLLM`'s depth-2 `step_with_batch_queue`. **That DISPROVES the 2026-07-25 attribution** which assigned 497 `cudaStreamSynchronize` calls (20.975 s, 42.20 ms/call) to `sample_tokens_async`'s discrete host path: that function is not on the benchmarked path at all, so those synchronizations are something else and MUST be re-attributed before another lever is chosen from that trace. Fixed by feeding the mirror from whichever sampler ran (the sync path uploads the ids it has already downloaded). FINAL gates, one lock, same binary: **token identity 128/128 in BOTH directions**; `test_qwen35_plain_weights --no-skip` 3/3 (1672 asserts); `test_input_batch` 25/25 (183 — +3 NEW W4 cases pinning the op log, including a replay-vs-host-array composition case that catches a MISORDERED op, which each op checked individually would not); `test_combine_tokens` 7/7; `test_ops_gdn` 66/66 (4242); `test_ops_paged_attn` 25/25 (454,474); clean `-Werror` full CUDA rebuild, 0 warnings. **SPEED: NEUTRAL and structurally unmeasurable on this harness** — the sync loop has no overlap for W4 to unlock, so it adds four small uploads + two kernels per step and removes nothing; paired runs gave 6612.31 vs 6600.68 and 6602.22 vs 6612.15 tok/s (equal and opposite = run noise). Hence **opt-in, DEFAULT OFF** (`VT_ASYNC_DEVICE_MIRROR=1`), the same disposition the W3 device kernels landed with before their DGX A/B; production is byte-identical. The binding measurement is a SERVING A/B over `AsyncLLM` and is **PENDING** (`tools/bench/run_serve_low.py`). `benchmark_binding=false`. **Pre-existing failure found, verified NOT introduced:** `test_cuda_ops` "CUDA matmul (cuBLASLt) matches CPU on odd sizes" fails 11 elements on the bf16/bf16 17x31x13 case on this discrete sm_120; reproduced on a pristine stash-clean build of the same commit (436 asserts, same single failure). The project develops on GB10/sm_121a, so this is an unattributed per-architecture numerics difference on consumer Blackwell. `ENG-ASYNC-SCHED` stays `DONE` for its W3 scope; W4 is `ACTIVE` (opt-in, serving A/B owed). Spec: [async-discrete-device-combine.md](specs/async-discrete-device-combine.md). Not pushed; FULL SHA reported. | +| 2026-07-27 (**HARDEN-DETECTOR-LANES — dynamic detector lanes + allocator blind-spot bypass, adopted from an external C++/Rust porting study**; `CLAIM-HARDEN-LANES`; local workspace host, CPU-tier infrastructure, no GPU claim; NOT pushed; closing commit: this checkpoint) | **What it does.** Closes the one gap an external hardening guide identified that this project genuinely had. The guide (`../float-harness/public-export/float-porting-analysis/`, distilled from defects in two unrelated tensor-engine ports) was assessed section by section; most of its rules this project already satisfies or exceeds — model metadata IS validated as a schema at load, reference-vs-optimized differential tests ARE the house parity tier, and the benchmark protocol already demands an idle box, one lock per series, >=3 repetitions and VOIDs contended numbers. The real gap was §8/§9: the test matrix was ONE default build, with no ASan, UBSan or TSan anywhere in CI. Adds `VLLM_CPP_SANITIZE` (`address` / `undefined` / `address,undefined` / `thread`) as a HOST lane that refuses to configure with the CUDA backend on (a host sanitizer runtime does not instrument nvcc device TUs and false-positives against the CUDA driver), with `-fno-sanitize-recover=all` so a finding ABORTS and ctest fails rather than burying the report in passing output, and a `sanitize-cpu` CI matrix running ASan+UBSan and TSan as separate jobs. Adds `VT_POOL_BYPASS=1`, which makes every `DevicePool` Get an exact-size driver allocation and every Put a real Free: the caching, size-class-rounding, never-returning pool is a compute-sanitizer BLIND SPOT twice over (a write past the last tensor row stays inside the block; a use-after-free reads memory that is still mapped and may already belong to another op). Explicitly a debugging lane, never a timing configuration — it reinstates the per-op device-sync storm the pool exists to remove. | External study (read-only, outside this repo): `HARDENING-GUIDE.md` §4, §8, §9; `LANGUAGE-ANALYSIS.md` §6-§7. No vLLM counterpart — this is house build/test infrastructure. | **How verified.** Configure-tier: the lane's guards were exercised directly (a lane value outside the allowlist and a lane with CUDA on both FATAL_ERROR by construction). `VT_POOL_BYPASS` and `VT_ASYNC_DEVICE_MIRROR` both pass `scripts/check-env-doc.py` (164 production env vars documented or classified) and are documented in `docs/ENVIRONMENT.md`. Default builds are BYTE-IDENTICAL: both switches default off, and the full CUDA `-Werror` rebuild plus the GPU correctness tier above ran with them unset. The CI lanes land `continue-on-error: true` because they have never run and their FIRST run is a survey, not a gate; removing that flag after triaging the first run is the recorded closing step of the row. `benchmark_binding=false` (infrastructure, no perf axis). Surveyed and found already prevented, recorded so nobody re-audits: no fixed-size stack scratch sized independently of input shape in the host layer, and 5 of 8 `std::atoi` env sites already guarded by an allowlist, a range check or a clamp (3 unvalidated but non-memory-safety sites remain, tracked). Assessment: [hardening-adoption-2026-07-27.md](specs/hardening-adoption-2026-07-27.md). Not pushed; FULL SHA reported. | +| 2026-07-27 (**W4 same-binary A/B + nsys RE-ATTRIBUTION — W4 is NEUTRAL on the sync loop, and the synchronizations the row was scoped from are the depth-1 engine's own per-step wait**; `CLAIM-ASYNC-SCHED-W4-AB`; local RTX 5070 Ti sm_120 discrete, one `flock /tmp/gpu` per series, idle box enforced by the new utilization gate; NOT pushed; closing commit: this checkpoint) | **What it does.** Records the binding same-binary A/B for W4 and replaces the profiling attribution the row was originally scoped from. A/B: 3 interleaved repetitions per arm, order flipped on even reps, 20 s cooldown before each leg, identical corpus. NEW nsys capture of the benchmarked path with `--cuda-graph-trace=node`. | Ours: `tools/bench/run_qwen35_4b_ab.sh`; the re-attributed claim supersedes the `sample_tokens_async` attribution in [qwen35-4b-main-repair-20260725.md](../docs/bench-evidence/qwen35-4b-main-repair-20260725.md). | **How verified.** **A/B (`/tmp/w4-ab-final`):** W4 ON vs OFF = **0.9996x** total throughput (6618.047 vs 6620.797), **0.9996x** output throughput, **1.0004x** mean TPOT (38.110 vs 38.093), **1.0002x** mean TTFT; per-arm repetitions ON 6610.4/6621.3/6622.4 and OFF 6622.3/6622.6/6617.5, so the 0.04% gap is an order of magnitude inside each arm's own spread → **NEUTRAL**. **Token identity 384/384 requests across all three pairs.** **RE-ATTRIBUTION (`/tmp/w4-attrib.nsys-rep`):** on the benchmarked path `cudaStreamSynchronize` is **112 calls, 1.134 s, 10.12 ms avg** over ~64 decode steps plus prefill and warm-up — about ONE PER ENGINE STEP; scaled to the binding workload (~512 steps at ~38 ms TPOT) that is ~500 calls at ~40 ms, i.e. the 2026-07-25 trace's 497 calls at 42.20 ms almost exactly. The COUNT and TIME were right, the ATTRIBUTION to `sample_tokens_async` was WRONG (that function is never called under `vllm-bench`, which drives the synchronous `LLMEngine::step()` loop — established by instrumenting the branch and observing it never executes). It is the depth-1 loop waiting for its own sampling. Corroborating: the per-call embedding `cudaMalloc`+`cudaStreamSynchronize`+`cudaFree` is ABSENT from the new trace (W4e working); whole-run `cudaMalloc` 818 calls / 29 ms and `cudaFree` 512 / 28 ms. **Consequence for the lever list:** there is no per-step synchronize left to delete on the synchronous path — the wait IS the step — so overlap requires running the ASYNC engine loop, which W4 now makes legal on a discrete GPU. Next step is a SERVING A/B over `AsyncLLM`, blocked by the HARNESS (`run_serve_low.py` needs a pinned SGLang container and accepts only 27B/35B model keys), not by the engine. `benchmark_binding=true` for the W4 A/B disposition (NEUTRAL, opt-in default OFF); serving A/B PENDING. Evidence: [w4-async-mirror-20260727.md](../docs/bench-evidence/w4-async-mirror-20260727.md). Not pushed; FULL SHA reported. | +| 2026-07-28 (**Oracle REBUILT FROM SOURCE at the parity pin; the published local ratio was PESSIMISTIC — 0.9819x -> 0.9970x with no code change**; `CLAIM-PINNED-ORACLE`; local workspace host (RTX 5070 Ti sm_120 discrete, driver 595.71.05), one `flock /tmp/gpu` across all 18 legs, all 9 performance legs recorded 0% GPU utilization before starting; NOT pushed; closing commit: this checkpoint) | **What it does.** Replaces a pip-installed PyPI-release oracle with one built from source AT the pin. The local oracle was vLLM **0.24.0**; the pin `555967922` is a vLLM **main** commit with no release tag and **no prebuilt wheel on any platform** (x86_64 verified here, aarch64 by the pin spec), so `pip install vllm` could never have reached it — the drift was structural, not neglect, and the records had additionally mislabelled the venv "0.25.0". Builds `vllm 0.23.1rc1.dev1511+g555967922.cu132` (editable, sm_120 only) into a SEPARATE `.venv-vllm-pin` with the pinned stack (torch 2.13.0, torchvision 0.28.0, triton 3.7.1, transformers 5.14.1, flashinfer 0.6.15.post1 +cubin, nvidia-cutlass-dsl 4.6.0, tilelang 0.1.9, quack-kernels, humming-kernels, tokenspeed-mla); the 0.24.0 venv is UNTOUCHED and retained as a rollback denominator. `${VLLM_SOURCE}` becomes a real checkout at the pin (`/home/rich/c/vllm-upstream`). Harness gains `VLLM_CUDA_HOME` because a venv carrying its own CUDA must not inherit the Nix 12.9 toolkit its extensions were not compiled against. | Pin `5559679229bc961848b121ccdeaa8fa5d79bec98` (2026-07-26); target stack per [pin-advance.md](specs/pin-advance.md) §1 + `requirements/cuda.txt` at that commit; FA2 arch rule `cuda_archs_loose_intersection(FA2_ARCHS "8.0+PTX" ...)` in the vendored vllm-flash-attn. | **How verified.** Oracle proven RUNNABLE before use as a denominator: `vllm._custom_ops` loads, platform resolves, and a real greedy generation on Qwen3.5-4B returns coherent text via the **FLASH_ATTN** backend (selected from `['FLASH_ATTN','FLASHINFER','TRITON_ATTN','FLEX_ATTENTION']`). **BINDING RESULT (3 reps/arm, interleaved):** total + output throughput **0.9970x** (6618.160 vs 6638.129 tok/s), req/s 0.9969x, mean TTFT **0.7731x PASS** (729.217 vs 943.198 ms), mean TPOT **1.1241x FAIL** (38.107 vs 33.900 ms); memory PASSES (peak PSS 2.488 vs vLLM 8.149, stable 0.759 vs 4.565). Spread 0.11% (pin: 6640.3/6633.4/6640.7) and 0.08% (ours: 6615.3/6618.7/6620.5), so the ratios are not noise. **THE FINDING: vLLM at the pin is 0.9875x the 0.24.0 release**, so the previous denominator was the FASTER vLLM and the published 0.9819x was UNDERSTATING us; the throughput gap is 0.3%, not 1.8%, with zero code change, and TPOT (+12.4%) is confirmed as the one real gap. **Control:** our own arm reproduces the 2026-07-27 series exactly (1.0027x, token-identical 128/128 per rep), which is what makes the delta attributable to the oracle. Token identity: ours ON vs OFF 128/128 per rep; vLLM-pin vs vLLM-0.24.0 92/92/95 (the oracle's own output moved between versions, expected); ours vs pin 89/89/98 (the established near-tie). **CUDA toolkit ceiling recorded (cost two full builds):** the usable version is set by what the DRIVER can JIT — vLLM's `8.0+PTX` FA2 is JIT-compiled at load, driver 595.71.05 tops out at 13.2 and rejects nvcc-13.3 PTX (`cudaErrorUnsupportedPtxVersion`, RUNTIME-only after a clean build); CUDA 13.0 is separately unusable (headers predate glibc 2.42 `rsqrt`; fixed from 13.1 via `_NV_RSQRT_SPECIFIER`); the whole toolkit must match (cccl hard-errors on a mixed compiler/header pair); install `--no-deps` or pip re-resolves the runtime down to 13.0 post-build; pip CUDA wheels need unversioned `.so` dev symlinks for `find_library`. **Honest non-fidelity:** the oracle logs a missing optional `triton_kernels` package, which is absent from vLLM's own `requirements/cuda.txt` so a stock install lacks it too (and it is NOT the repo directory shadowing — the harness invokes by path, `sys.path[0]` is `tools/bench`, verified). `benchmark_binding=true` for the local 4B row against the PIN denominator; no 4B result implies 27B/35B support or speed. Evidence: [qwen35-4b-pinned-oracle-20260728.md](../docs/bench-evidence/qwen35-4b-pinned-oracle-20260728.md). Not pushed; FULL SHA reported. | +| 2026-07-29 (**Post-rebase RE-VALIDATION of the 4B lever onto `main` `f3ecbe70d`, +139 upstream commits: nothing moved**; `CLAIM-POSTREBASE-REVALIDATE`; local working branch `bench-lever-sampled-token-20260727` rebased from `7f620e74`, build `build-nix-cuda-transplant-triton` clean 925/925 with 0 warnings; ALL GPU under one `flock /tmp/gpu`, all 9 performance legs at 0% utilization; NOT pushed; closing commit: this checkpoint) | **What it does.** Re-measures the Qwen3.5-4B matched comparison after the branch was rebased across 139 upstream commits, against the SAME source-built oracle at the parity pin (`0.23.1rc1.dev1511+g555967922`; pin re-checked and unchanged before the run). Result is a null: 0.9972x total throughput (was 0.9970x), TTFT 0.7701x PASS, TPOT 1.1247x FAIL — every axis inside noise. Records-and-evidence only; no engine change. | vLLM `555967922` (0.26.0.dev0) built from source at the pin, run through the same `run_qwen35_4b_compare.sh` harness as the 2026-07-28 series; no upstream code ported. | **VERIFIED, and the control is the point.** Our output is token-identical 128/128 per repetition to the pre-rebase series on BOTH the direct-ON and direct-OFF arms, which is the semantic check on three rebase conflict resolutions (two append-only record files, plus a real `ModelForwardInput` field collision with upstream's `mm`). All three arms drifted down ~0.13% (ours 0.9988x, OFF 0.9983x, pin 0.9986x) — the pin arm is an UNCHANGED binary, so that drift is ambient, not code. Gates re-run green: `test_qwen35_plain_weights --no-skip` 3/3·1672/1672, `test_input_batch` 25/25·183/183, `test_combine_tokens` 7/7·14/14. Evidence: [docs/bench-evidence/qwen35-4b-postrebase-20260729.md](../docs/bench-evidence/qwen35-4b-postrebase-20260729.md). | diff --git a/.agents/specs/async-discrete-device-combine.md b/.agents/specs/async-discrete-device-combine.md new file mode 100644 index 000000000..b49447991 --- /dev/null +++ b/.agents/specs/async-discrete-device-combine.md @@ -0,0 +1,174 @@ +# ENG-ASYNC-SCHED W4 — discrete-CUDA device-resident sampled tokens + +Row: `ENG-ASYNC-SCHED` W4 (engine matrix). Prerequisite W3 is landed and +default-ON. This spike scopes the remaining leaf: making the async-scheduling +overlap real on a **discrete** CUDA GPU, where W3 currently degrades to a +synchronizing host path. + +## The defect this closes + +W3 landed the device combine/scatter kernels +(`src/vt/cuda/cuda_combine_tokens.cu`) and wired them into the runner, but both +call sites are gated on `is_integrated_gpu()`: + +- `src/vllm/v1/worker/gpu/runner.cpp:821` — device combine, else host combine. +- `src/vllm/v1/worker/gpu/runner.cpp:1763` — device scatter, else host scatter + preceded by `vt::GetBackend(dev.type).Synchronize(queue_)`. + +The gate is correct as written: the kernels take the runner's **host** +`std::vector` buffers directly, which is only legal where the platform reports +pageable device access (GB10's UMA). A discrete GPU answers `false` and takes +the host branch, so: + +1. the host must know step N's sampled token before it can build step N+1's + `input_ids`, which forces a full-device `Synchronize` inside + `sample_tokens_async`; and +2. the depth-2 `step_with_batch_queue` loop (`src/vllm/v1/engine/core.cpp:115`) + therefore cannot overlap anything — the scheduler is async, the runner is not. + +Measured on the local RTX 5070 Ti (`is_integrated_gpu()` == false), 2026-07-25, +attribution-complete node-mode nsys: 497 `cudaStreamSynchronize` calls following +the 256-byte sampled-ID D2H, **20.975 s total, 42.20 ms/call**, against a binding +mean TPOT of 43.72 ms. The matched vLLM trace instead waits on sampled output via +`cudaEventSynchronize` (1,342 calls) AFTER launching the next batch. Evidence: +`docs/bench-evidence/qwen35-4b-main-repair-20260725.md`. + +## What upstream does (the mirror obligation) + +Read from `${VLLM_SOURCE}` (installed release tree, vLLM 0.24.0 — cited as +installed-release source, NOT as the parity pin): + +- `vllm/v1/worker/gpu/states.py:64` — + `self.last_sampled_tokens = torch.zeros(max_num_reqs, 1, dtype=torch.int64, + device=device)`. It is a **GPU tensor**, unconditionally, on every platform. + There is no integrated/discrete branch upstream at all. +- `vllm/v1/worker/gpu/input_batch.py:296-360` — the combine kernel loads + `last_sampled_tokens_ptr + req_state_idx` on device and stores into the device + `input_ids`. +- `vllm/v1/worker/gpu/input_batch.py:449-473` — the post_update kernel stores the + freshly sampled id back into `last_sampled_tokens_ptr + req_state_idx` on + device. +- `vllm/v1/worker/gpu/states.py:132` `remove_request` — **upstream never + condenses**. A finished request's slot goes onto `free_indices` and is reused; + the req_state index is stable for the request's lifetime, which is exactly why + the GPU tensor never has to be permuted, and why `idx_mapping` (batch row -> + req_state) exists. + +So the discrete path is not a new design; it is the upstream design, which our +W3 leaf implemented only for the UMA case. + +## The one real complication: our batch condenses + +`InputBatch::condense()` (`src/vllm/v1/worker/gpu/input_batch.cpp:554`) moves the +last live row into a freed slot (`last_sampled_tokens[empty] = +last_sampled_tokens[last]`, line 641), and `swap_states` (line 737) swaps two +rows. Upstream's free-index pool has neither. With a device-resident buffer the +host no longer holds the values it would need to perform those moves, and reading +them back would reintroduce the synchronize this row exists to delete. + +Resolution: keep the moves on the HOST as bookkeeping, but record them and replay +them ON DEVICE in stream order. `InputBatch` gains an ordered pending-op log +(seed on `add_request`, move on `condense`, swap on `swap_states`); the runner +drains it each step and applies it with one small kernel BEFORE the combine. The +ops are host-known (indices, and a host-known value for the seed), so no device +read is needed, and stream ordering makes the replay exact rather than racy. +Rejected alternative: put `last_sampled_tokens` in pinned mapped host memory so +the existing kernels work unchanged — the host `condense` read of a value the +device wrote is then an unsynchronized read, which is precisely the +removal/condensation hazard, and it would be latent rather than loud. + +## Work breakdown + +- **W4a** Runner-owned persistent device buffers sized to the batch bound: + `last_sampled` [max_num_reqs], `prefill_len` [max_num_reqs], `query_start_loc` + [max_num_reqs+1], `seq_lens` [max_num_reqs], `input_ids` + [max_num_batched_tokens], plus a pinned host staging buffer so the per-step H2D + is a real async copy rather than a pageable staging copy. +- **W4b** `InputBatch` pending-op log + the device replay kernel + (`LaunchApplyLastSampledOps`), with RED-first unit coverage of seed/move/swap + ordering. +- **W4c** `ModelForwardInput::device_token_ids` (default `nullptr`, so every + other model and the whole non-async path is byte-identical) honored by the + Qwen3.5 dense/MoE `EmbedInto`, so the forward embeds the DEVICE ids the combine + patched instead of a host vector. +- **W4d** Flip both `is_integrated_gpu()` gates to select + UMA-in-place / discrete-device-buffer rather than device / host, and delete the + `Synchronize` on the discrete branch. +- **W4e** Remove the other per-step barrier: `EmbeddingKernelCuda` + (`src/vt/cuda/cuda_ops.cu:638-669`) does a `cudaMalloc` + `cudaMemcpyAsync` + + **`cudaStreamSynchronize`** + `cudaFree` for its out-of-range flag on EVERY + call (531 calls, 12.6 ms total in the same trace). That is negligible while the + engine is serialized and becomes a hard barrier the moment it is not, so W4 + is not measurable until it is gone. Fix: a persistent per-device flag plus a + deferred check (read the previous call's flag through a completed event), which + keeps the loud failure at the cost of at most one step of latency. **W4e must + land before the W4 A/B, or the A/B measures the wrong thing.** + +## Gates + +- Correctness: `test_qwen35_plain_weights --no-skip` 3/3 unchanged; the CPU-tier + combine/scatter/input-batch suites unchanged; a new RED-first unit test for the + pending-op replay. +- Identity: token-for-token identity across `{default, VT_ASYNC_RUNNER=1, + VT_ASYNC_RUNNER=0}` on the 128-request benchmark corpus, including a run whose + requests finish at staggered lengths so `condense` actually fires. +- Speed: same-binary A/B under one `flock /tmp/gpu`, >=3 repetitions, against the + matched vLLM arm on the identical corpus. The target is the TPOT axis; TTFT + must be watched (upstream pays a TTFT premium for async scheduling). +- No 4B result implies anything about the 27B/35B release gates. + +## What was actually built (2026-07-27) + +Implemented as scoped, with three deviations worth recording because each was +forced by something the spike did not anticipate: + +1. **The forward reads the device ids through a scoped override, not a + parameter.** `ModelForwardInput::device_token_ids` reaches the two Qwen3.5 + registry forwards, which publish it as an RAII, thread-local + `detail::DeviceTokenIdsScope` that the embed consults. The alternative was a + defaulted pointer parameter on five entry points plus the decode-graph class. + The override is consumed on first use so a second, unrelated embed inside the + same forward (the multimodal helper embeds a prompt and then single tokens) + cannot pick up ids that were never meant for it. +2. **The embed PATCHES a prefix rather than embedding the runner's buffer + directly.** The decode-graph path does not embed `token_ids` as given — it + embeds a version padded up to the captured batch size, real rows first. So + the correct operation is "overwrite the first `count` rows", which is right + for the padded case and degenerates to "overwrite everything" on the eager + path. Embedding the runner's buffer directly would have silently mis-shaped + the graph path. +3. **The per-step uploads copy from PAGEABLE host memory on purpose.** The + spike said "pinned staging so the copies are real async copies". That is a + correctness trap here: pinned copies are truly asynchronous, so reusing one + staging buffer for the next upload can overwrite bytes an in-flight DMA has + not read, and with a depth-2 scheduler that window spans steps. For a + pageable source the driver stages the bytes before `cudaMemcpyAsync` returns, + so the source is immediately reusable. These arrays are a few kilobytes at + the front of a step; the staged copy is the better trade against per-upload + regions plus a per-step event. + +A fourth correction came from the shared-layer device-leakage ratchet: the +mirror's enable predicate first asked "is this device CUDA and not integrated", +which put a `kCUDA` token in the device-agnostic layer and failed the ratchet. +The right question was already available and is more precise — `vt::Backend:: +UnifiedMemory()`, i.e. "is device memory addressable from the host". A unified +device (GB10, and the CPU backend trivially) keeps the in-place path; a device +with separate memory needs the mirror. Same behaviour, no device-type test. + +Also landed alongside: `VT_ASYNC_DEVICE_MIRROR=0`, the same-binary rollback that +returns a discrete GPU to the pre-W4 host path WITHOUT disabling async scheduling +(which `VT_ASYNC_RUNNER=0` would also do). That separation is what makes an +honest A/B of W4 alone possible. + +W4e (the embedding barrier) landed as a persistent ring of flag slots with a +deferred, event-queried check. The reporting contract changed deliberately and +its test changed with it: an out-of-range id is now raised no later than the NEXT +embedding on the queue, carrying the same message and the same id, instead of on +the offending call. The gather itself is unchanged, so the offending call never +read out of bounds either way. + +## Status + +`ACTIVE` — implemented; gates and the A/B are the closing step. Nothing here may +be called DONE until the token-identity gate and the same-binary A/B are on the +record. diff --git a/.agents/specs/hardening-adoption-2026-07-27.md b/.agents/specs/hardening-adoption-2026-07-27.md new file mode 100644 index 000000000..e049fb0a5 --- /dev/null +++ b/.agents/specs/hardening-adoption-2026-07-27.md @@ -0,0 +1,130 @@ +# Hardening adoption review — external C++/Rust porting study, 2026-07-27 + +Row: `HARDEN-DETECTOR-LANES` (engine matrix, cross-cutting quality). +Scope: read an external, independently written hardening guide distilled from +defects found in two unrelated tensor-engine ports, decide item by item what +transfers to vllm.cpp, and land the transferable prevention. + +Source (read-only, outside this repository): +`../float-harness/public-export/float-porting-analysis/` — `HARDENING-GUIDE.md` +(the prevention rules), `LANGUAGE-ANALYSIS.md` §6 confirmed action items and §7 +auditor blind-spot guide, `REPORT.md` (the experiment the defects came from). +The study's subject engines are audio-driven portrait models, not inference +servers, so the transfer is by DEFECT CLASS, never by copying a fix. + +## Verdict per guide section + +| § | Guide rule | State here | Action | +|---|---|---|---| +| 1 | Validate once, preserve the proof in types | Partial. `vt::Tensor` carries dtype/shape/device and `VT_CHECK` validates at op entry, but token/slot/block counts are bare `int32_t`/`int64_t` throughout the runner. | NO CODE CHANGE. A units/newtype pass over the runner is a large refactor with a real regression surface and no observed defect behind it. Recorded as a standing preference for new interfaces, not a retrofit. | +| 2 | Model metadata is a schema, validated at load | ALREADY STRONGER. The safetensors/GGUF loaders validate the whole manifest (dtype, rank, extents, byte size) and the model gates are token-exact against the vLLM oracle. | None. | +| 3 | Encode pipeline state and policy | Partial by construction: the engine mirrors vLLM's staged pipeline, and policies (quantization scheme, attention backend, cache mode) are already closed enums resolved once. | None. | +| 4 | Confine unsafe memory; scoped arenas | `DevicePool` is a per-block cache, not a suballocating arena, so the guide's worst case (many logical tensors inside ONE allocation) does not exist. Two weaker forms of the same blind spot DO: size-class rounding hands back a block up to 6.25% larger than the tensor, and blocks are never returned to the driver. | **LANDED** — `VT_POOL_BYPASS=1` (`include/vllm/model_executor/models/device_pool.h`). | +| 5 | Harden parsers and FFI boundaries | ALREADY STRONGER on parsers (bounded readers, fuzz-covered GGUF/safetensors). The C ABI checks handles and translates failures. | None. | +| 6 | Separate semantic correctness from memory safety | ALREADY THE HOUSE RULE: every kernel keeps a CPU reference and the parity tiers are differential tests against it, plus the vLLM oracle above that. Fixed reduction order is a live concern, not a gap: the recorded 119/128 benchmark-token identity under Triton AOT is exactly the completion-order effect the guide warns about, and it is already attributed to reduction order with operator coverage supplying the equivalence proof. | None. | +| 7 | Make state ownership and concurrency visible | Partial. Profiling counters are per-object, and the aux-stream pool split (`AuxPool()`) is precisely the guide's "hand out disjoint capabilities" rule applied to CUDA streams. Untested: concurrent independent engines, callback re-entry. | Covered by the TSan lane below; a dedicated concurrency matrix is left as a follow-up row. | +| 8 | Make the build observable; use a build MATRIX | **THE REAL GAP.** Build observability is already strong (the bench harness records commit, dirty state, CMake cache, `ldd`, dataset/binary SHA-256, GPU/thermal state). But the test matrix was ONE default build: no ASan, no UBSan, no TSan, anywhere in CI. | **LANDED** — `VLLM_CPP_SANITIZE` + the `sanitize-cpu` CI matrix. | +| 9 | Build tests around a defect matrix | ALREADY STRONGER in the behavioural tier (parity taps per stage, ported upstream tests as the spec). Weakest in the DYNAMIC tier, which is the same gap as §8. | Same action as §8. | +| 10 | Treat performance and process defects as defects | ALREADY STRONGER than the guide asks: `.agents/benchmark-protocol.md` requires an idle box, one exclusion lock per series, >=3 repetitions, same-binary A/B, a recorded repro recipe, and it VOIDS contended or non-reproducing numbers. | None. | + +## What was landed + +### 1. Dynamic detector lanes (§8, §9 — the highest-priority gap) + +`VLLM_CPP_SANITIZE` selects a HOST instrumentation lane +(`address`, `undefined`, `address,undefined`, `thread`), and CI job +`sanitize-cpu` runs the full CPU suite under ASan+UBSan and under TSan as two +separate jobs. Design points, all from the guide: + +- ASan/UBSan and TSan are separate builds; their runtimes are mutually exclusive. +- The lane REFUSES to configure with the CUDA backend on: the host runtime does + not instrument nvcc device TUs and false-positives against the CUDA driver. + The CUDA tier's equivalent detector is `compute-sanitizer`. +- `-fno-sanitize-recover=all`, so a finding ABORTS and ctest fails the lane. The + guide's specific warning is that a print-and-continue sanitizer produces a + green run containing a real report. +- The lane keeps `-ffp-contract=off`, so a bit-identity assertion still means in + the lane what it means in production. +- `-Werror` is dropped IN THIS LANE ONLY, in `cmake/CompilerWarnings.cmake` + where the per-target warning flags are set (a global `-Wno-error` is overridden + by those PRIVATE target options and does nothing). Sanitizer instrumentation + changes inlining enough that GCC's range and initialization analyses fire + inside LIBSTDC++ on correct code: a one-element + `std::vector v = {x}` in `voxtral.cpp` draws "forming offset 4 is out + of the bounds [0, 4]" for a 4-byte read of a 4-byte array, and `` draws + 26 `-Wmaybe-uninitialized` reports from `std::function` internals. None is + project code, and the plain build — the one that DOES enforce `-Werror` — stays + clean. Letting a false positive in a system header stop the runtime detectors + would defeat the point. +- Landed `continue-on-error: true` because the lanes have never run: their FIRST + run is a survey, and a pre-existing finding must not block unrelated work. The + closing step of this row is triaging that first run and REMOVING the flag. + +**The lane paid for itself before it was even committed.** Its first real build +failed on `-Werror=unused-function`: `AsyncDeviceMirrorEnvDefault()`, added by the +W4 work in the same session, is referenced only inside `#ifdef VLLM_CPP_CUDA` and +is therefore dead on a CPU build. That would have broken the EXISTING +`build-test-cpu` CI job too — the sanitizer lane just happened to compile the +CPU configuration first. Fixed by guarding the definition with the same `#ifdef` +as its use. This is precisely the guide's §8 point: a single default build never +compiles the other `cfg` branches. + +**Verified end to end**, not merely configured: `address,undefined` and `thread` +both configure; the two guards fire (a lane value outside the allowlist, and a +lane with the CUDA backend on, are both `FATAL_ERROR`); the ASan+UBSan lane +builds `test_input_batch`, `test_combine_tokens` and `test_arena` clean and all +three PASS under `ASAN_OPTIONS=detect_leaks=1 UBSAN_OPTIONS=print_stacktrace=1` +(25/25 + 183 asserts, 7/7 + 14, 4/4 + 18), which includes the new W4 op-log +cases. The plain CPU build was re-verified clean and green in the same pass. + +### 2. Allocator-blind-spot bypass (§4) + +`VT_POOL_BYPASS=1` makes every `DevicePool` Get an exact-size driver allocation +and every Put a real Free. Default OFF and production-inert. It exists because a +caching, size-class-rounding, never-returning pool defeats the detector twice: +an overrun past the last tensor row stays inside the block, and a use-after-free +reads memory that is still mapped (and may already belong to another op). This +is the guide's single highest-priority blind spot ("the most likely place a real +bug is hiding undetected"), in the milder form our design permits. It is a +DEBUGGING lane, never a timing configuration — it reinstates the per-op +cudaMalloc/cudaFree device-sync storm the pool exists to remove. + +Intended use: `VT_POOL_BYPASS=1 compute-sanitizer --tool memcheck `. + +## Surveyed and found already prevented + +The guide's greps were run against this tree; two classes it flags do not +reproduce here, which is worth recording so a later reader does not re-audit: + +- **Fixed-size stack scratch sized independently of the input shape** (the + study's `srow[4096]` stack smash): no instances in the host layer. Scratch is + pooled device memory sized from the actual shape. +- **Unvalidated `atoi` on environment variables**: 8 sites, and 5 are already + guarded downstream by an allowlist (`VT_GDN_DECODE_NW` accepts only 1/2/4/8), + a range check (`VT_GDN_OCC_BLOCK` 32..1024), or a clamp + (`VLLM_CPP_CPU_THREADS` -> `std::clamp(n, 1, kMaxThreads)`). Three are + unvalidated and remain open, all low severity because none is a memory-safety + input: `VT_MOE_SHARED_AUX_THRESHOLD` (`qwen3_5.cpp`, garbage -> 0 silently + changes an overlap threshold), `VT_VULKAN_DEVICE` (`vulkan_context.cpp`, + garbage -> 0 silently forces device 0 instead of failing), and the LMCache + remote port (`remote_client.cpp`, garbage -> port 0). Tracked as follow-up. + +## Not adopted, with reasons + +- **Rust/Miri items** — no Rust in this project. +- **`panic = "abort"` / `catch_unwind` boundary policy** — Rust-specific; our C + ABI's exception policy is already explicit. +- **FIL-C lane** — its separate ABI and rebuild cost do not fit a CUDA project + whose hot path is device code the tool cannot see. ASan/UBSan plus + compute-sanitizer cover the same classes for us. +- **Units/newtype retrofit (§1)** — see the table; deliberate, not an oversight. + +## Follow-up rows + +1. Triage the first `sanitize-cpu` run and remove `continue-on-error` (closes + this row). +2. Validate the three residual `atoi` sites. +3. A concurrency matrix per §7 (independent concurrent engines, callback + re-entry, forced worker failure). +4. Run a gate test under `VT_POOL_BYPASS=1` + `compute-sanitizer` and record the + result; that is the evidence the bypass lane is worth keeping. diff --git a/.agents/state.md b/.agents/state.md index 4bc297d80..396292845 100644 --- a/.agents/state.md +++ b/.agents/state.md @@ -32362,3 +32362,298 @@ GPU memory-halving path, DGX-blocked); the runner/spec integration (half-sized K `k_scale`/`v_scale` threading + `--kv-cache-dtype`/`--calculate-kv-scales`); fp8_e5m2 CPU compute + per-attention-head scales. W1 is a CPU correctness brick — the real memory/throughput WIN is the GPU path + halved-block runner integration. Not pushed; FULL SHA reported to the caller. + +## 2026-07-27 — Qwen3.5-4B post-pull revalidation on discrete Blackwell, and a VOIDED prior series + +Workspace fast-forwarded 109 commits onto `upstream/main` `7f620e74`; the local +`c317237a` transplant's two commits are now upstream (`a131de03`, `b6f1efc8`), so +the measured tree is plain current `main`. Rebuilt +`build-nix-cuda-transplant-triton` clean, ran the correctness tier and the full +18-leg matched comparison under one `flock /tmp/gpu`. + +**Correctness: unchanged, bit for bit.** `test_qwen35_plain_weights --no-skip` +3/3 (1672 asserts), `test_ops_gdn` 66/66 (4242), `test_ops_paged_attn` 25/25 +(454,474), `test_gdn_packed_decode_triton` 1/1, `test_combine_tokens` 7/7, +`test_input_batch` 22/22 — every count identical to the 2026-07-25 record. And +the benchmark output is token-identical to that series: 128/128 requests per +repetition for BOTH direct-ON and direct-OFF. 109 upstream commits moved no token +on this workload. + +**Performance: no meaningful regression, and no real gain either.** Against the +oracle in the SAME series, total throughput 0.9864x -> 0.9819x and the TPOT +excess +13.41% -> +14.00%. Absolute numbers rose ~14% on BOTH arms, which is the +whole story below. + +**THE FINDING — the 2026-07-25 series was contended, so its absolute numbers are +VOID.** All nine of its performance legs recorded 11-13% GPU utilization and +611 MiB of extra resident VRAM BEFORE the leg started; all nine of today's +recorded 0%. `prepare_leg` gated idleness on `nvidia-smi --query-compute-apps`, +which enumerates CUDA contexts only, so a GRAPHICS consumer was invisible to it +and an entire binding series passed the gate while sharing the GPU. Both arms +gained ~14% on a genuinely idle box, which is exactly why the RATIO moved almost +not at all. Closed by adding an explicit `utilization.gpu` check +(`GPU_IDLE_UTIL_MAX`, default 2%); re-parsed, it reads 0 for today and 12 for +2026-07-25. The prior series' ratios, same-binary component attributions and +profiling attribution survive — each was internal to one uniformly-contended +series — but no absolute number from it may be published. + +Also corrected: the local oracle is **vLLM 0.24.0**, not the 0.25.0 the previous +evidence and BENCHMARKS entry claimed. That series' own recorded +`vllm-version.txt` already said 0.24.0. The project parity pin is now +`555967922` / 0.26.0.dev0, so this denominator is behind the pin and is labelled +as such. + +Two summarizer defects fixed on the way: it read historical token legs under a +`perf-` prefix the harness has never written, and it demanded a +`vllm_production` key its own current output does not contain. Either one made +summarizing a series against the previous one impossible. + +**Residual and next step.** TPOT is the only failing latency axis and its +mechanism is specific, not diffuse: this GPU is DISCRETE, `is_integrated_gpu()` +is false, so both ENG-ASYNC-SCHED W3 device call sites (`runner.cpp:821` +combine, `runner.cpp:1763` scatter) take their host fallback and +`sample_tokens_async` must synchronize the main stream before the host can read +the sampled ids. The async scheduler IS engaged (`max_concurrent_batches=2`), so +depth-2 currently overlaps nothing here. Upstream has no such branch: +`states.py:64` keeps `last_sampled_tokens` GPU-resident unconditionally and +`states.py:132` never condenses slots (free-index pool), which is why it needs no +host round-trip; our `InputBatch::condense` is the one real complication and the +spec resolves it by replaying the recorded row moves on device in stream order. +Scoped, not implemented: [specs/async-discrete-device-combine.md](specs/async-discrete-device-combine.md). +The spec also names the SECOND per-step barrier that must go with it — the CUDA +embedding out-of-range flag does `cudaMalloc` + `cudaStreamSynchronize` + +`cudaFree` on EVERY call, which is 23.7 us while the engine is serialized and a +hard barrier the moment it is not. + +Evidence root `/tmp/qwen35-postpull-7f620e74`; index +[docs/bench-evidence/qwen35-4b-postpull-20260727.md](../docs/bench-evidence/qwen35-4b-postpull-20260727.md). + +Pre-existing and NOT introduced here: `scripts/check-agent-record.py` reports 6 +errors on pristine `main` (DONE rows citing closing commits `164453a2` / +`7a3f04b2`, which do not exist in this clone). Verified identical with and +without this change. + +## 2026-07-27 (later) — ENG-ASYNC-SCHED W4 implemented and gated; the lever's own attribution was WRONG + +Built the discrete-CUDA device-resident sampled-token path the earlier entry +scoped, gated it, and in the process disproved the premise it was scoped from. + +**What landed.** `last_sampled_tokens` gains a device mirror on discrete CUDA +(what upstream keeps unconditionally, `states.py:64`); the combine and scatter +run on it; `InputBatch` records its structural row edits (seed/move/swap) and the +runner replays them on the device in stream order, which is how a device-resident +buffer survives our `condense()` (upstream never condenses — `states.py:132` uses +a free-index pool). The forward reads the patched device ids through +`ModelForwardInput::device_token_ids`. W4e removed the other per-step barrier: +the CUDA embedding's out-of-range flag no longer does cudaMalloc + stream-sync + +cudaFree per call, but uses a persistent ring of flag slots with a deferred, +event-queried check (reporting contract changed deliberately; its test changed +with it). + +**Three design corrections forced during implementation**, each recorded in the +spec: the embed PATCHES a prefix rather than embedding the runner's buffer +(the decode graph embeds a PADDED id vector, so "replace the whole buffer" would +have mis-shaped it); the uploads copy from PAGEABLE host memory on purpose (a +shared pinned staging buffer is a race — pinned copies are truly async, so the +next upload can overwrite bytes an in-flight DMA has not read, and with depth-2 +that window spans steps); and the forward reads the ids through an RAII scoped +override consumed on first use, rather than a parameter threaded through five +entry points plus the decode-graph class. + +**THE FINDING — the lever was aimed at code the benchmark never runs.** First +gate run: token identity mirror-ON vs OFF was **0/128**, output garbage from the +first decode. Instrumenting showed the device mirror stayed all-zero: the scatter +branch in `sample_tokens_async` NEVER EXECUTED. Cause: `vllm-bench` drives the +SYNCHRONOUS `LLMEngine::step()` loop, which calls `sample_tokens()`, not +`AsyncLLM`'s depth-2 `step_with_batch_queue` -> `sample_tokens_async`. So on the +benchmarked path there is no async sampler, no depth-2 overlap, and **no +`sample_tokens_async` synchronize to remove**. The 2026-07-25 attribution of 497 +`cudaStreamSynchronize` calls (20.975 s, 42.20 ms/call) to that function is +therefore WRONG for this workload; those synchronizations come from somewhere +else and must be re-attributed before another lever is chosen from that trace. +Fixed by feeding the mirror from whichever sampler ran (the sync path uploads the +ids it already downloaded); identity is now **128/128 in both directions**. + +**Disposition: opt-in, DEFAULT OFF** (`VT_ASYNC_DEVICE_MIRROR=1`), exactly as the +W3 device kernels landed before their DGX A/B. On the sync loop W4 is correct but +can only cost (four small uploads + two kernels per step, nothing removed); +paired runs measured 6612.31 vs 6600.68 and 6602.22 vs 6612.15 tok/s — equal and +opposite, run noise. The binding measurement is a SERVING A/B over `AsyncLLM`, +PENDING; `tools/bench/run_serve_low.py` is the harness for it. + +**Pre-existing failure found and verified, NOT introduced here:** `test_cuda_ops` +"CUDA matmul (cuBLASLt) matches CPU on odd sizes" fails 11 elements on the +bf16/bf16 17x31x13 case on this discrete sm_120. Reproduced on a pristine +stash-clean build of the same commit (436 assertions, same single failure). The +project develops on GB10/sm_121a, so this is an unattributed per-architecture +numerics difference on consumer Blackwell that nobody has recorded. + +Also adopted from the external hardening study: `VLLM_CPP_SANITIZE` host lanes +(ASan+UBSan, TSan) plus the `sanitize-cpu` CI matrix, and `VT_POOL_BYPASS=1` so +compute-sanitizer can see through the caching device-scratch pool. Item-by-item +transfer decisions: [specs/hardening-adoption-2026-07-27.md](specs/hardening-adoption-2026-07-27.md). + +## 2026-07-27 (close) — W4 A/B and the re-attribution that replaces the lever's premise + +**A/B (`/tmp/w4-ab-final`, 3 interleaved reps/arm, order flipped on even reps, +one lock, idle box):** W4 ON vs OFF is **0.9996x** total and output throughput, +**1.0004x** TPOT, **1.0002x** TTFT. Per-rep spread within each arm is larger than +the gap between them. **NEUTRAL.** Token identity across all three pairs: +**384/384 requests identical.** + +**Re-attribution (`/tmp/w4-attrib.nsys-rep`, `--cuda-graph-trace=node`).** The +2026-07-25 record blamed 497 `cudaStreamSynchronize` calls (20.975 s, 42.20 +ms/call) on `sample_tokens_async`. A fresh profile of the benchmarked path shows +112 `cudaStreamSynchronize` over ~64 decode steps plus prefill and warm-up at +10.12 ms each — about ONE PER ENGINE STEP. Scaled to the binding workload (~512 +steps at ~38 ms TPOT) that is ~500 calls at ~40 ms, which is the old trace's +numbers almost exactly. So the COUNT and the TIME were right and the ATTRIBUTION +was wrong: it is the depth-1 `LLMEngine::step()` loop waiting for its own +sampling (the sampler's greedy download, `sampler.cpp:191`), not an async-sampler +defect. Also visible: the per-call embedding `cudaMalloc`+sync+`cudaFree` is gone +from the trace, which is W4e working; `cudaMalloc` is down to 818 calls and +`cudaFree` to 512 for the whole run, at 29 ms and 28 ms total. + +**So the lever list changes.** There is no per-step synchronize to delete on the +synchronous path — the wait IS the step. Overlap requires running the ASYNC +engine loop, and W4 is what makes that loop legal on a discrete GPU. The next +step is therefore a SERVING A/B over `AsyncLLM` (`examples/server`), not another +runner-level kernel change. It is blocked by the harness, not by the engine: +`tools/bench/run_serve_low.py` needs a pinned SGLang container image and accepts +only the 27B/35B model keys, neither available here. A minimal serving driver for +the 4B on this host is the unblocking task. + +Evidence: [../docs/bench-evidence/w4-async-mirror-20260727.md](../docs/bench-evidence/w4-async-mirror-20260727.md). + +## 2026-07-27 (close) — OPEN LEAD from the same trace: cuBLASLt picks Ampere-class GEMM kernels on sm_120 + +Recorded as a lead, NOT a claim, from `/tmp/w4-attrib.nsys-rep`. The two largest +GPU kernels on the 4B workload are + +- `cutlass_80_tensorop_bf16_s16816gemm_relu_bf16_256x128_32x3_tn_align8`, + 1,664 instances, 2.099 s, **47.1%** of GPU kernel time; +- `cutlass_80_tensorop_bf16_s16816gemm_relu_bf16_128x256_32x3_tn_align8`, + 544 instances, 0.570 s, 12.8%. + +Both are `cutlass_80` — SM80 / Ampere-class tensorop kernels — selected by +cuBLASLt's heuristic and executing on a Blackwell sm_120 device. The house +parity-lever protocol already names this exact phenomenon (`nvjet_sm121` vs +`cutlass_80`) as something only an execution trace reveals. + +Two reasons this is a LEAD and not a lever yet. First, these instance counts and +the `_relu_` epilogue point at the PREFILL projections, and prefill is the axis +we already PASS (TTFT 0.798x vs the oracle); the failing axis is TPOT. Second, +nobody has checked what the vLLM oracle resolves to on the SAME workload on THIS +device — if vLLM also lands on `cutlass_80` there is no gap to close, and the +protocol is explicit that the oracle's trace, not ours alone, decides. The +honest next step is a matched pair of traces (ours and vLLM's) on the identical +corpus, diffed by kernel name, before anyone tries to force a tactic. + +## 2026-07-28 — oracle rebuilt AT the parity pin; the published ratio was PESSIMISTIC, not flattering + +The local oracle was two versions behind the pin and, on investigation, could +never have matched it by installation: the pin `555967922` is a vLLM **main** +commit with no release tag and **no prebuilt wheel on any platform** (x86_64 +verified here; the pin spec verified aarch64). `pip install vllm` reaches only +PyPI releases, so the 2026-07-09 provisioning got 0.24.0 and drifted from there. +The records compounded it by labelling that venv "0.25.0". + +**Built from source at the pin** into a SEPARATE venv (`.venv-vllm-pin`, +`vllm 0.23.1rc1.dev1511+g555967922.cu132`, sm_120 only) with the pinned stack +(torch 2.13.0, torchvision 0.28.0, triton 3.7.1, transformers 5.14.1, flashinfer +0.6.15.post1, cutlass-dsl 4.6.0, tilelang 0.1.9). The 0.24.0 venv is untouched +and retained as a rollback denominator. `${VLLM_SOURCE}` is now a real checkout +at the pin (`/home/rich/c/vllm-upstream`), so pin-era mirror claims no longer +have to be hedged. + +**THE RESULT — the direction is the surprise.** vLLM at the pin is **0.9875x** +the 0.24.0 release on this workload, i.e. the old denominator was the FASTER +vLLM. So the published 0.9819x was understating us. Against the true pin: +total/output throughput **0.9970x** (6618.160 vs 6638.129), req/s 0.9969x, TTFT +**0.7731x PASS**, TPOT **1.1241x FAIL**. The throughput gap is 0.3%, not 1.8%, +with no change to our code; TPOT (+12.4%) is confirmed as the one real gap. +Control: our own arm reproduces the previous series exactly (1.0027x, +token-identical 128/128 per repetition), which is what makes the delta +attributable to the oracle rather than to us. Spread 0.11% (pin) / 0.08% (ours). + +**The CUDA toolkit ceiling, recorded because it cost two full builds.** The +usable CUDA version is set by what the DRIVER can JIT, not by what is newest. +vLLM ships FA2 as `8.0+PTX`, so the driver JIT-compiles its PTX for Blackwell at +load; driver 595.71.05 tops out at 13.2 and rejects nvcc-13.3 PTX with +`cudaErrorUnsupportedPtxVersion` — visible only at RUNTIME, after a clean build. +CUDA 13.0 is separately unusable (headers predate glibc 2.42's `rsqrt`; fixed +from 13.1 via `_NV_RSQRT_SPECIFIER`). 13.2 is the only version clearing both, the +whole toolkit must match (cccl hard-errors on a mixed compiler/header pair), and +vLLM must be installed `--no-deps` or pip re-resolves the runtime down to 13.0 +after the build and reintroduces the mismatch. pip's CUDA wheels also ship no +unversioned `.so`, so `find_library` needs dev symlinks. + +`run_qwen35_4b_compare.sh` gains `VLLM_CUDA_HOME` for a venv carrying its own +toolkit; unset, behaviour is unchanged. + +**Correction to an earlier claim in this session:** I reported that the repo's +`triton_kernels/` directory shadows Triton's package in benchmark runs. It does +NOT. The harness invokes the metrics script BY PATH, so `sys.path[0]` is +`tools/bench` and the repo root is never on `sys.path` — verified directly. The +shadowing only occurred in an ad-hoc test that piped a script via stdin. The +`Failed to import Triton kernels` warning is a genuinely absent optional package +that is not in vLLM's own `requirements/cuda.txt`, so a stock install lacks it +too; the oracle is faithful. + +Evidence: [../docs/bench-evidence/qwen35-4b-pinned-oracle-20260728.md](../docs/bench-evidence/qwen35-4b-pinned-oracle-20260728.md). + +## 2026-07-29 — rebased onto 139 upstream commits and re-benchmarked: nothing moved + +Rebased `bench-lever-sampled-token-20260727` from `main` `7f620e74` onto +`f3ecbe70d`. Three conflicts. `.agents/parity-ledger.md` and `.agents/state.md` +were pure append collisions (upstream appended 76 ledger rows and ~4,580 log +lines at the tail, this branch appended its own) — resolved by keeping both with +upstream first, which is what the ledger's own header prescribes ("append-only … +Newest last") and what a rebase means, then verified additively: `diff` against +`main` shows ZERO lines lost from upstream. The third was real: +`ModelForwardInput` gained a field on both sides (upstream's `mm`, ours +`device_token_ids`). Both kept, ours LAST — it has to be the final member, which +is the same constraint that broke positional aggregate initializers when it was +first written mid-struct. + +**Verification of the merge, in the order that makes it binding.** Clean CUDA +rebuild 925/925, exit 0, no warnings. `test_input_batch` 25/25·183/183, +`test_combine_tokens` 7/7·14/14, `test_qwen35_plain_weights --no-skip` +3/3·1672/1672 — all identical to pre-rebase. Then the strongest statement +available: the 18-leg comparison reproduced our output **token-identical +128/128 in every repetition** against the pre-rebase series, for both the direct-ON +and direct-OFF arms. A 139-commit rebase that leaves generated tokens bit-stable +is a semantic check the model gate alone cannot give. + +**Benchmark: a null result, and the useful kind.** Pin re-checked first and +unchanged (`555967922`), so `.venv-vllm-pin` is still the right denominator. +0.9972x total throughput (was 0.9970x), TTFT 0.7701x PASS (was 0.7731x), TPOT +1.1247x FAIL (was 1.1241x). Ours 6610.270 vs pin 6628.651 tok/s. Per-rep spread +0.05% (ours) and 0.07% (pin). + +The control that makes "nothing moved" a measurement rather than an assumption: +all three arms drifted down ~0.13% (ours 0.9988x, direct-OFF 0.9983x, pin +0.9986x). The pin arm is the SAME binary on the SAME corpus and cannot have +changed, yet it drifted the same amount — so the drift is ambient thermal/clock +state, not code, and it is the size of the arms' own repetition spread. + +No movement was expected. The only upstream commit in the window naming this gap +(`2b00866a4`) concluded the decode gap is batch composition, not a decode-kernel +deficiency, and changed records rather than code; the rest of the window is +breadth on paths this dense 4B decode workload does not touch. TPOT stays owned +by ENG-ASYNC-SCHED and the W4 finding stands. + +**Harness note worth keeping.** The first attempt at this series was killed +partway by a 10-minute background-command ceiling, after the memory legs but +before any performance leg, leaving an output tree with no aggregate. Re-run +detached under `setsid` with an explicit wait for three consecutive idle GPU +reads first, because the killed run left the card draining at 3% and the harness +(correctly) refuses any leg above 2%. All nine performance legs then recorded 0%. + +Pre-existing and NOT introduced here: `check-agent-record` (6 errors), +`check-env-doc` (`VLLM_GEMMA4_MM_DEBUG`) and `check-device-leakage` (DSR +`kcuda=8` vs baseline 0) all fail identically on pristine `main`, verified in a +throwaway worktree. The DSR bucket counts are byte-identical with and without +this branch, so this branch adds zero device leakage. + +Evidence: [../docs/bench-evidence/qwen35-4b-postrebase-20260729.md](../docs/bench-evidence/qwen35-4b-postrebase-20260729.md). diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 1531e8246..7468df142 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -144,3 +144,41 @@ jobs: run: cmake --build build -j 2 - name: Test run: ctest --test-dir build --output-on-failure + sanitize-cpu: + # The DYNAMIC detector lanes (VLLM_CPP_SANITIZE, see CMakeLists.txt). The + # plain build-test-cpu job above proves the suite PASSES; it cannot see a + # read one past a tensor row, a use-after-free of an engine-owned container, + # a signed-overflow index computation, or an unsynchronized access between + # the serving threads — every one of which stays green under -O2 until it + # silently corrupts a token stream. These jobs run the SAME suite with + # ASan+UBSan and with TSan. + # + # Two jobs, not one: the runtimes are mutually exclusive. Both are CPU-tier + # (the lane refuses to configure with the CUDA backend on) and both are + # `continue-on-error` for their FIRST landing so a pre-existing finding + # cannot block unrelated work — the finding is triaged, then this flag is + # removed and the lane becomes binding. Removing it is tracked as the + # closing step of the hardening-adoption row. + strategy: + fail-fast: false + matrix: + lane: ["address,undefined", "thread"] + continue-on-error: true + runs-on: ubuntu-latest + steps: + - uses: actions/checkout@v4 + - name: Configure + run: | + cmake -S . -B build-sanitize \ + -DVLLM_CPP_BUILD_TESTS=ON \ + -DVLLM_CPP_CUDA=OFF \ + -DVLLM_CPP_SANITIZE='${{ matrix.lane }}' + - name: Build + run: cmake --build build-sanitize -j 2 + - name: Test + # ctest runs the suite serially: the sanitizer runtimes multiply peak RSS, + # and a parallel run OOM-kills the runner before it reports a finding. + env: + UBSAN_OPTIONS: print_stacktrace=1 + ASAN_OPTIONS: detect_leaks=1:strict_string_checks=1 + run: ctest --test-dir build-sanitize --output-on-failure diff --git a/.gitignore b/.gitignore index be2a7fc75..f4784e52f 100644 --- a/.gitignore +++ b/.gitignore @@ -7,10 +7,13 @@ compile_commands.json *.so .venv/ .venv-vllm/ +.venv-vllm-pin/ .hf-cache/ .torchinductor-cache/ .triton-cache/ +.triton-cache-pin/ .vllm-cache/ +.vllm-cache-pin/ .flashinfer-cache/ __pycache__/ diff --git a/CMakeLists.txt b/CMakeLists.txt index 9985c35f6..fe8f2d045 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -125,6 +125,54 @@ else() set(VLLM_CPP_CUDA OFF) endif() +# ─── Host instrumentation lanes ─────────────────────────────────────────────── +# One default build is shallow coverage: it reaches neither the dynamic +# memory/UB detectors nor the race detector, so a whole defect class (bad index +# arithmetic, use-after-free across the engine's owning containers, unsynchronized +# access from the serving threads) can only be caught by a failing parity test, +# if at all. This selects an instrumentation lane for the HOST C++ build so those +# detectors run as their own CI jobs alongside the plain build. +# +# ASan/UBSan and TSan are deliberately SEPARATE lanes: their runtimes are +# mutually exclusive and their environment requirements differ. The lane is +# HOST-ONLY — nvcc-compiled device TUs are not instrumented, and a host sanitizer +# runtime reports false positives against the CUDA driver — so anything other +# than OFF requires the resolved VLLM_CPP_CUDA to be OFF. The CUDA tier's +# equivalent detector is compute-sanitizer, run against the normal build. +# +# -fno-sanitize-recover=all makes a finding ABORT instead of print-and-continue, +# so ctest FAILS the lane rather than burying the report inside passing output. +# Frame pointers plus -g keep the traces attributable; the lane keeps the +# project's -ffp-contract=off numerics so a bit-identity test still means what it +# means in the production build. +set(VLLM_CPP_SANITIZE "OFF" CACHE STRING + "Host sanitizer lane: OFF | address | undefined | address,undefined | thread") +set_property(CACHE VLLM_CPP_SANITIZE PROPERTY STRINGS + OFF address undefined address,undefined thread) +if(NOT VLLM_CPP_SANITIZE STREQUAL "OFF") + set(_vllm_cpp_sanitize_allowed address undefined address,undefined thread) + if(NOT VLLM_CPP_SANITIZE IN_LIST _vllm_cpp_sanitize_allowed) + message(FATAL_ERROR + "VLLM_CPP_SANITIZE='${VLLM_CPP_SANITIZE}' is not one of: " + "OFF;${_vllm_cpp_sanitize_allowed}") + endif() + if(VLLM_CPP_CUDA) + message(FATAL_ERROR + "VLLM_CPP_SANITIZE requires VLLM_CPP_CUDA=OFF: the host sanitizer runtime " + "does not instrument nvcc device TUs and reports false positives against " + "the CUDA driver. Use compute-sanitizer for the CUDA tier instead.") + endif() + add_compile_options( + $<$:-fsanitize=${VLLM_CPP_SANITIZE}> + $<$:-fno-sanitize-recover=all> + $<$:-fno-omit-frame-pointer> + $<$:-g>) + # -Werror itself is dropped for this lane in cmake/CompilerWarnings.cmake, where + # the per-target warning flags are set; see the rationale there. + add_link_options(-fsanitize=${VLLM_CPP_SANITIZE}) + message(STATUS "Host sanitizer lane: ${VLLM_CPP_SANITIZE} (CUDA backend off)") +endif() + # Resolve the tri-state VLLM_CPP_METAL the same way (BACKEND-METAL-MLX W0). The # Metal backend is ObjC++ against the Metal/Foundation frameworks, so it needs # APPLE plus a working OBJCXX compiler; on any other host AUTO resolves OFF and diff --git a/cmake/CompilerWarnings.cmake b/cmake/CompilerWarnings.cmake index d641a61d2..b257e6730 100644 --- a/cmake/CompilerWarnings.cmake +++ b/cmake/CompilerWarnings.cmake @@ -1,6 +1,22 @@ function(vllm_cpp_set_warnings target) + # The instrumentation lanes (VLLM_CPP_SANITIZE) keep the warnings but drop + # -Werror. This has to happen HERE rather than as a global -Wno-error, because + # these are PRIVATE target options and therefore land after anything + # add_compile_options contributed — a global -Wno-error would simply be + # overridden. Sanitizer instrumentation changes inlining enough that GCC's + # range and initialization analyses fire inside LIBSTDC++ on correct code (a + # one-element `std::vector v = {x}` draws "forming offset 4 is out of + # the bounds [0, 4]" for a 4-byte read of a 4-byte array; draws 26 + # -Wmaybe-uninitialized reports from std::function internals). None is project + # code, and the plain build — which IS the one that enforces -Werror — is + # clean. Letting a false positive in a system header stop the RUNTIME detectors + # would defeat the point of the lane. + set(_vllm_cpp_werror -Werror) + if(NOT VLLM_CPP_SANITIZE STREQUAL "OFF") + set(_vllm_cpp_werror "") + endif() target_compile_options(${target} PRIVATE - $<$:-Wall -Wextra -Werror> + $<$:-Wall -Wextra ${_vllm_cpp_werror}> # OBJCXX (.mm — the Metal backend) is a SEPARATE COMPILE_LANGUAGE from CXX, # so the CXX genex above does not reach it. Without this line the Metal TUs # would be the only unwarned code in the tree (BACKEND-METAL-MLX W0). diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index 96d5131e7..b6a04c828 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -4941,43 +4941,191 @@ The local Qwen3.5-4B checkpoint corpus is exactly `/tmp/qwen35-4b-sharegpt-1024.json`, SHA-256 `9ea13603767c62c267e3f381fbccf42d0c9ca0c393655c37533eadca7aefca0c`. -### Local Qwen3.5-4B direct-load checkpoint, 2026-07-25 +### Local Qwen3.5-4B direct-load checkpoint, 2026-07-27 (revalidated on `main` @ `7f620e74`) -**GATING / speed-pending; current-main transplant revalidation PENDING.** The -measurements below bind the original repair commit `8964bde7` based on -`72f9fb13`, not its current-main transplant `c317237a` based on `c39d78a6`. -That path includes the exact H32 Triton-AOT recurrence, plain-BF16 decode graph -capture/replay and ratio-4 FA2 split-KV decode. On the corpus above at -concurrency 32, 128 output tokens/request and three repetitions, direct -ON/OFF/vLLM-0.25 means were: +**GATING / speed-pending.** This SUPERSEDES the 2026-07-25 entry and discharges +the revalidation it recorded as pending; the transplant's commits are now +upstream, so the measured tree is plain current `main`. Corpus above, +concurrency 32, 128 output tokens/request, three repetitions per arm, one +`flock /tmp/gpu` across all 18 legs, oracle **vLLM 0.24.0** (the local venv's +own recorded version; the previous entry's "0.25.0" label was wrong). -| Axis | Direct ON | Direct OFF | vLLM 0.25 stable confirmation | Disposition | +| Axis | Direct ON | Direct OFF | vLLM 0.24.0 | Disposition | |---|---:|---:|---:|---| -| Total throughput (tok/s) | 5769.99 | 5660.70 | 5849.80 | FAIL, 0.9864x | -| Output throughput (tok/s) | 638.03 | 625.94 | 646.85 | FAIL, 0.9864x | -| Requests/s | 4.9867 | 4.8900 | 5.0536 | FAIL, 0.9868x | -| Mean TTFT (ms) | 834.91 | 943.43 | 1047.13 | PASS | -| Mean TPOT/ITL (ms) | 43.72 | 43.84 | 38.55 | FAIL | -| Peak PSS (GiB) | 2.406 | 8.592 | 7.662 | PASS | -| Stable PSS (GiB) | 0.759 | 8.589 | 4.029 | PASS | -| Peak VRAM (MiB) | 12850.7 | 12843.3 | 12942.7 | PASS vs vLLM and within +8 MiB vs OFF | - -Direct ON and OFF are output-identical for 128/128 requests in every pair. -Same-binary component gains are H32 AOT **+4.5906%**, decode graph -**+0.3873%**, and ratio-4 FA2 **+1.6004%**. The final attribution-complete -node trace contains 453 graph launches and 200,972 graph-child kernels; local -FA2 averages 180.28 us/call versus vLLM's 178.40 us/call. CUDA API attribution -finds the residual: 497 sampled-ID D2H synchronizations consume 20.975 s -(42.20 ms/call) because the discrete path updates host request rows -immediately; vLLM defers the equivalent wait through events. Root, hashes, -exact commands and the stable-denominator rationale: +| Total throughput (tok/s) | 6600.66 | 6481.34 | 6722.24 | FAIL, 0.9819x | +| Output throughput (tok/s) | 729.88 | 716.69 | 743.33 | FAIL, 0.9819x | +| Requests/s | 5.700 | 5.600 | 5.807 | FAIL, 0.9815x | +| Mean TTFT (ms) | 729.24 | 835.01 | 913.55 | PASS, 20.2% lower | +| Mean TPOT/ITL (ms) | 38.22 | 38.21 | 33.53 | FAIL, 14.0% higher | +| Peak PSS (GiB) | 2.282 | 8.594 | 7.666 | PASS | +| Stable PSS (GiB) | 0.761 | 8.591 | 4.039 | PASS | +| Peak VRAM (MiB) | 12850.0 | 12846.0 | 12933.3 | PASS vs vLLM and within +4 MiB vs OFF | + +Every generated token is IDENTICAL to the 2026-07-25 run, 128/128 requests per +repetition in both arms, so 109 upstream commits moved no token here. Our arm is +deterministic rep-to-rep and ON==OFF 128/128; the oracle is not self-deterministic +(102/128 on one pair). Spread is 0.13% (ours) and 0.17% (vLLM) across +repetitions, so the ratios are not noise. + +**The 2026-07-25 absolute numbers are VOID: that series was contended.** All +nine of its performance legs ran with the GPU at 11-13% utilization and 611 MiB +of extra resident VRAM; today's nine ran at 0%. Both arms gained ~14% on the +idle box, which is exactly why the ratio barely moved (0.9864x -> 0.9819x). The +cause was a harness hole - `prepare_leg` gated idleness on +`--query-compute-apps`, which cannot see a GRAPHICS consumer - now closed by an +explicit `utilization.gpu` check (`GPU_IDLE_UTIL_MAX`, default 2%), carried by +the new same-binary A/B harness `tools/bench/run_qwen35_4b_ab.sh` too - where +running it promptly proved the gate works, refusing an arm at 14% because that +harness sampled idleness BEFORE its inter-leg cooldown and so read the previous +leg still draining; the cooldown now precedes the snapshot. The prior +entry's RATIOS, its same-binary component attributions (H32 AOT **+4.5906%**, +decode graph **+0.3873%**, ratio-4 FA2 **+1.6004%**) and its profiling +attribution survive, because each was internal to one uniformly-contended series. + +The residual is unchanged and now specified. TPOT fails because this GPU is +DISCRETE: `is_integrated_gpu()` is false, so both ENG-ASYNC-SCHED W3 device call +sites take their host fallback and `sample_tokens_async` must synchronize the +main stream (497 syncs, 20.975 s, 42.20 ms/call in the node-mode trace). The +async scheduler IS engaged (`max_concurrent_batches=2`), so depth-2 currently +buys nothing here. Upstream keeps `last_sampled_tokens` GPU-resident +unconditionally and never condenses request slots. Scoped as `ENG-ASYNC-SCHED` +W4 in [.agents/specs/async-discrete-device-combine.md](../.agents/specs/async-discrete-device-combine.md). +Root, hashes, per-leg contention table, exact commands: +[Qwen3.5-4B post-pull revalidation](bench-evidence/qwen35-4b-postpull-20260727.md). +Prior series (ratios and attributions only): [Qwen3.5-4B main repair evidence](bench-evidence/qwen35-4b-main-repair-20260725.md). -Revalidate `c317237a` with the exact `flock /tmp/gpu -tools/bench/run_qwen35_4b_compare.sh /tmp/qwen35-main-final-fa2-` recipe -in that evidence document before publishing these values as binding for the -development branch. This 4B diagnostic does not establish 27B/35B support. +### ENG-ASYNC-SCHED W4 discrete-CUDA device-resident sampled tokens (2026-07-27) - correctness GATED, speed NEUTRAL here, serving A/B PENDING + +**Implemented, opt-in (`VT_ASYNC_DEVICE_MIRROR=1`), DEFAULT OFF.** On a discrete +GPU the W3 device combine/scatter fell back to a host path that must synchronize +the main stream for the sampled ids; W4 gives that path a device-resident +`last_sampled_tokens` (what upstream does unconditionally, `states.py:64`), with +the condense/swap row edits recorded on the host and replayed on the device in +stream order. + +Correctness gate on the RTX 5070 Ti, one `flock /tmp/gpu`, same binary: +**token-identical, 128/128 requests, in both directions** on the 128-request +corpus; `test_qwen35_plain_weights --no-skip` 3/3 (1672), `test_input_batch` +25/25 (183, +3 new W4 cases), `test_combine_tokens` 7/7, `test_ops_gdn` 66/66 +(4242), `test_ops_paged_attn` 25/25 (454,474). + +Speed on this workload is NEUTRAL and cannot be otherwise, which is why the +default is OFF. `vllm-bench` drives the SYNCHRONOUS `LLMEngine::step()` loop, so +`sample_tokens_async` is never called and there is no depth-2 overlap for W4 to +unlock; it adds four small uploads and two kernels per step and removes nothing. +Two paired runs measured 6612.31 vs 6600.68 and 6602.22 vs 6612.15 tok/s - equal +and opposite, i.e. run noise. The win it was built for needs the ASYNC serving +loop (`AsyncLLM` -> `step_with_batch_queue`), so the binding measurement is a +SERVING A/B and is **PENDING**. + +**This also corrects a wrong attribution in the record.** The 2026-07-25 evidence +assigned the 497 `cudaStreamSynchronize` calls (20.975 s, 42.20 ms/call) to +`GPUModelRunner::sample_tokens_async`'s discrete host path. That function is not +called on the benchmarked path at all - verified directly, by instrumenting the +branch and observing it never executes under `vllm-bench`. Whatever those +synchronizations are, they are not the async sampler's. A fresh +attribution-complete profile re-attributes them: 112 `cudaStreamSynchronize` +calls over ~64 decode steps plus prefill and warm-up at 10.12 ms each - **about +one per engine step**, which scaled to the binding workload (~512 steps at ~38 ms) +is ~500 calls at ~40 ms, i.e. the 2026-07-25 numbers almost exactly. The count +and the time in that trace were right; the attribution was not. It is the depth-1 +engine loop waiting for its own sampling, not an async-sampler defect. The +per-call embedding barrier is absent from the new trace, which is W4e working. +The mirror's enable predicate asks the BACKEND whether memory is unified rather +than testing for CUDA, which is both more precise and what keeps the shared +layer's device-leakage ratchet at its floor. +Numbers, profile and reproduction: +[W4 evidence](bench-evidence/w4-async-mirror-20260727.md). Spec and deviations: +[.agents/specs/async-discrete-device-combine.md](../.agents/specs/async-discrete-device-combine.md). + +### Pinned-oracle provisioning (2026-07-28) - infrastructure, measurement PENDING + +The local oracle was **vLLM 0.24.0**, two versions behind the parity pin +(`555967922` / 0.26.0.dev0), because it was pip-installed from a PyPI RELEASE on +2026-07-09 while the pin tracks a vLLM **main commit** that has no release tag +and no prebuilt wheel on any platform. A pip install could never have matched it. + +An oracle at the actual pin now exists in a SEPARATE venv (`.venv-vllm-pin`, +built from source at `5559679229bc...`, `vllm-0.23.1rc1.dev1511+g555967922.cu132`) +with the pinned stack: torch 2.13.0, torchvision 0.28.0, triton 3.7.1, +transformers 5.14.1, flashinfer 0.6.15.post1, nvidia-cutlass-dsl 4.6.0. The +working 0.24.0 venv is untouched. + +`run_qwen35_4b_compare.sh` gains `VLLM_CUDA_HOME`. The harness previously derived +the vLLM arm's CUDA toolkit from the Nix CMake cache (12.9), which is correct for +a venv that borrows Nix's CUDA and WRONG for one carrying its own: it would put a +12.9 toolkit ahead of the 13.x the venv's extensions were compiled against. + +**Toolkit ceiling, recorded because it is not obvious:** the usable CUDA version +here is set by what the DRIVER can JIT, not by what is newest. vLLM ships +FlashAttention-2 as `8.0+PTX`, so the driver JIT-compiles its PTX for Blackwell at +load; driver 595.71.05 tops out at CUDA 13.2 and rejects nvcc-13.3 PTX with +`cudaErrorUnsupportedPtxVersion` — a failure that appears only at RUNTIME, after a +completely clean build. CUDA 13.0 is also unusable: its headers predate glibc +2.42's `rsqrt` declaration and collide with it. 13.2 is the only version in range +that clears both, and the whole toolkit must match (cccl refuses a mixed +compiler/header pair). Install vLLM with `--no-deps`, or pip re-resolves the CUDA +runtime downward after the build and reintroduces the mismatch. + +**RE-MEASURED (2026-07-28): the 0.9819x was PESSIMISTIC.** vLLM at the pin is +0.9875x the 0.24.0 release on this workload, so the old denominator was the +FASTER vLLM and was understating us. Against the true pin: + +| Axis | Direct ON | vLLM @ pin | ratio | Disposition | +|---|---:|---:|---:|---| +| Total throughput (tok/s) | 6618.160 | 6638.129 | 0.9970x | FAIL | +| Output throughput (tok/s) | 731.817 | 734.026 | 0.9970x | FAIL | +| Requests/s | 5.717 | 5.735 | 0.9969x | FAIL | +| Mean TTFT (ms) | 729.217 | 943.198 | 0.7731x | PASS | +| Mean TPOT (ms) | 38.107 | 33.900 | 1.1241x | FAIL | + +So the throughput gap is **0.3%, not 1.8%**, with no change to our code, and TPOT +(+12.4%) is confirmed as the one real gap. Our own arm reproduces the previous +series exactly (1.0027x, token-identical 128/128 per rep), which is the control +that makes the oracle delta attributable to the oracle. Spread is 0.11% (pin) and +0.08% (ours). Full provenance, the CUDA toolkit ceiling that constrained the +build, and token identity: +[pinned-oracle comparison](bench-evidence/qwen35-4b-pinned-oracle-20260728.md). + +**RE-VALIDATED (2026-07-29) after rebasing onto 139 upstream commits** (`main` +`7f620e74` -> `f3ecbe70d`): 0.9972x total throughput, TTFT 0.7701x PASS, TPOT +1.1247x FAIL — every axis inside noise of the row above, and our output +token-identical 128/128 per repetition to the previous series. The control is +that the vLLM arm, whose pinned binary did NOT change, drifted by the same ~0.13% +as ours (0.9986x vs 0.9988x), so the drift is ambient rather than a code effect. +No movement was the expected result: the only upstream commit in that window +naming this gap (`2b00866a4`) concluded the decode gap is batch composition and +changed records, not code. +[post-rebase re-validation](bench-evidence/qwen35-4b-postrebase-20260729.md). + +### OPEN LEAD - cuBLASLt resolves Ampere-class GEMM kernels on sm_120 (2026-07-27) - NOT MEASURED + +From the same profile: `cutlass_80_tensorop_bf16_s16816gemm_relu_bf16_256x128` +is 47.1% of GPU kernel time (1,664 instances, 2.099 s) and its 128x256 sibling a +further 12.8%. Both are SM80 / Ampere-class kernels chosen by cuBLASLt's +heuristic on a Blackwell device. Recorded as a LEAD, not a claim: the instance +counts and `_relu_` epilogue indicate the PREFILL projections, and prefill is the +axis we already pass (TTFT 0.798x), while the failing axis is TPOT. No matched +vLLM trace exists on this device yet, and if the oracle resolves the same kernels +there is no gap. The next step is a matched pair of traces on the identical +corpus, diffed by kernel name. + +### Hardening detector lanes (2026-07-27, `HARDEN-DETECTOR-LANES`) - infrastructure, NOT APPLICABLE + +No performance claim: `VLLM_CPP_SANITIZE` (ASan/UBSan/TSan host lanes, CPU-tier +only) and `VT_POOL_BYPASS=1` (exact-size, really-freed `DevicePool` allocations +so `compute-sanitizer` can see tensor boundaries and use-after-free) are both +default-OFF and production-inert. `VT_POOL_BYPASS` is explicitly NOT a timing +configuration - it reinstates the per-op `cudaMalloc`/`cudaFree` device-sync +storm the pool exists to remove. Rationale and the item-by-item transfer +decision: [.agents/specs/hardening-adoption-2026-07-27.md](../.agents/specs/hardening-adoption-2026-07-27.md). +The lanes are verified end to end (both configure, both guards fire, ASan+UBSan +builds and passes three CPU suites with leak detection on), and their first real +build already caught a `-Werror=unused-function` break in the W4 commit that +would also have failed the existing `build-test-cpu` job. + ## `M3c-1` CUDA-path neutrality on GB10 (2026-07-27) - NO MEASURABLE REGRESSION **Discharges the debt the `M3c-1` commit declared.** That change edited two files diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index dd4db9682..7a0887999 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -49,6 +49,7 @@ portable/reference path. In normal operation leave them unset. |---|---|---| | `VT_ASYNC_RUNNER` | on | Synchronous model runner (no async/overlap execution) | | `VT_ASYNC_SCHED` | on | Synchronous scheduling (no scheduler/execution overlap). The documented first-line workaround for a suspected scheduling bug | +| `VT_ASYNC_DEVICE_MIRROR` | off | `=1` engages the discrete-CUDA device-resident sampled-token mirror (ENG-ASYNC-SCHED W4), so the async serving loop's sampled ids stay on the device instead of round-tripping the host. DEFAULT OFF pending a serving A/B, exactly as the W3 device kernels landed: the synchronous `LLMEngine::step()` loop has no overlap for it to unlock, so it can only cost there. Token-identical to the default path either way. No effect on an integrated GPU (which keeps its in-place path) or on CPU | | `VLLM_CPP_CUDAGRAPH` | on (CUDA) | Eager launches instead of a captured CUDA graph | | `VLLM_CPP_DENSE_DECODE_GRAPH` | on (CUDA dense) | Non-graphed dense decode | | `VT_MM_DECODE_EAGER` | off (graph on) | Set to `1` to force the eager per-step multimodal (Qwen3.6-27B image/video) decode instead of routing it through the captured dense decode graph. Rollback / A-B knob; the graphed path is token-exact with the eager path | @@ -80,6 +81,7 @@ Read-only observability; none change output. | `VT_OP_PROVIDER_DISABLE` | (none) | Comma-separated provider names to disable, forcing fallback (diagnostic) | | `VT_GDN_VALIDATE` | off | Run the GDN validation/cross-check path (slower; for kernel debugging) | | `VT_FP4_AUTOTUNE_VERBOSE` | off | Log the NVFP4 GEMM autotuner's tactic selection | +| `VT_POOL_BYPASS` | off | `=1` makes every device-scratch pool allocation an exact-size driver `Alloc` and every release a real `Free`, so `compute-sanitizer` can see tensor boundaries and use-after-free that the caching, size-class-rounding pool hides. DEBUGGING ONLY: it reinstates the per-op `cudaMalloc`/`cudaFree` device-sync storm the pool exists to remove, so it is never a timing configuration | ## Kernel-internal knobs (deferred) diff --git a/docs/STATUS.md b/docs/STATUS.md index 3cd093b28..740ccf167 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -37,7 +37,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, decode at-parity, prefill speed-pending | Token-exact greedy; decode at or beyond vLLM, remaining gap is prefill TTFT | | Qwen3 / Qwen2 dense (BF16) | Correctness-complete, speed-pending | Near-tie-robust token-exact vs vLLM (Qwen3-0.6B, Qwen3-4B); c1 effective parity, c8 decode residual | -| Qwen3.5-4B plain BF16 direct loading on discrete CUDA | Correctness-complete, speed-pending | Direct ON and OFF are token-identical; direct loading cuts peak/stable host PSS by 72.0%/91.2%. The pre-transplant H32 AOT, decode-graph and ratio-4 FA2 result reached 0.9864x vLLM 0.25.0; the current-main development branch is pending revalidation | +| Qwen3.5-4B plain BF16 direct loading on discrete CUDA | Correctness-complete, speed-pending | Direct ON and OFF are token-identical, and token-identical to the previous series; direct loading cuts peak/stable host PSS by 73.4%/91.1% and mean TTFT by 12.7%. Against an oracle built at the actual parity pin: 0.9970x total throughput, TTFT passes (0.773x), TPOT 12.4% high; re-validated unchanged (0.9972x / 1.1247x) after rebasing onto 139 upstream commits. The failing axis is the discrete-GPU async-overlap gap, scoped as ENG-ASYNC-SCHED W4 | | 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 | | 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) | @@ -547,20 +547,79 @@ The remaining recent-dense families are the trivial tail only: **Yi** (`YiForCausalLM`, a Llama alias) and **InternLM3** (`InternLM3ForCausalLM`, InternLM2 plus a sliding window). +## Build and test lanes + +Measured against an oracle built from source at the ACTUAL parity pin, the gap +is 0.9970x total throughput, not the 0.9819x published against the older +pip-installed 0.24.0 release: that release is 1.25% faster than the pin, so +the old denominator was understating us. TPOT (+12.4%) is the one real gap. +Re-validated 2026-07-29 after rebasing onto 139 upstream commits: 0.9972x / +TPOT 1.1247x, every axis inside noise and token-identical to the prior series. + +One open lead is on record from the same profiling pass: cuBLASLt resolves +Ampere-class GEMM kernels on this Blackwell device. It is unmeasured and may be +a non-issue, since it lands on the axis we already pass. + +Device-specific references in the device-agnostic layer stay at their ratchet +floor: capability questions ("is memory unified?") rather than device-type +tests. + +Benchmark provenance is gated as well as measured: both the comparison and the +same-binary A/B harness refuse to start a leg on a GPU that is not idle, after +a cooldown rather than before it. + +Alongside the default build, `-DVLLM_CPP_SANITIZE=address,undefined` and +`-DVLLM_CPP_SANITIZE=thread` build the CPU tier under the dynamic detectors, and +CI runs both as separate jobs. Verified end to end: the ASan+UBSan lane builds +and passes `test_input_batch`, `test_combine_tokens` and `test_arena` with leak +detection on. The lanes keep the warnings but drop `-Werror`, because sanitizer +instrumentation makes GCC's range and initialization analyses fire inside +libstdc++ on correct code; the plain build is the one that enforces `-Werror`. The lane refuses to configure with the CUDA +backend on, because a host sanitizer runtime does not instrument nvcc device +translation units and reports false positives against the CUDA driver; the CUDA +tier's equivalent is `compute-sanitizer`, and `VT_POOL_BYPASS=1` makes the device +scratch pool hand out exact-size, really-freed allocations so that tool can see +tensor boundaries and use-after-free the caching pool otherwise hides. + +Known failing on discrete sm_120 (RTX 5070 Ti), pre-existing and not introduced +by the lanes: `test_cuda_ops` "CUDA matmul (cuBLASLt) matches CPU on odd sizes" +fails 11 of 442 elements on the bf16-in/bf16-out 17x31x13 case. Verified present +on a pristine build of the same commit. The project's development GPU is GB10 / +sm_121a, so this is an unattributed per-architecture numerics difference, not a +regression. + ## Performance detail -**Local Qwen3.5-4B plain BF16 direct loader, speed-pending:** the measured -pre-transplant repair on an RTX 5070 Ti reached 5769.99 total tok/s versus vLLM -0.25.0 at 5849.80 tok/s (0.9864x) on the identical 128-request workload. Direct -loading reduced peak PSS from 8.59 to 2.41 GiB and stable PSS from 8.59 to 0.76 -GiB. Mean TPOT/ITL was 43.72 ms versus vLLM's 38.55 ms. Profiling attributes the -residual to the discrete-CUDA sampled-token D2H path synchronizing the main -stream instead of retaining vLLM's event-overlapped device mapping. The repair -has since been transplanted onto current `main`; that development branch is -pending the same-series revalidation, so these numbers do not yet bind to it. +**Local Qwen3.5-4B plain BF16 direct loader, speed-pending:** revalidated on +current `main` (`7f620e74`) on an RTX 5070 Ti at 6600.66 total tok/s versus vLLM +0.24.0 at 6722.24 tok/s (0.9819x) on the identical 128-request workload. Direct +loading reduces peak PSS from 8.59 to 2.28 GiB, stable PSS from 8.59 to 0.76 +GiB, and mean TTFT from 835.0 to 729.2 ms. Mean TPOT/ITL is 38.22 ms versus +vLLM's 33.53 ms. Every token matches the previous series exactly (128/128 per +repetition, both arms), so the 109-commit upstream advance moved no output here. + +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. Closing it means running the async engine loop, which is +what the (opt-in) device-resident sampled-token mirror now makes legal on a +discrete GPU; the measurement that would bind it is a serving A/B and is +pending a harness this host can run. + +Two things to know about the numbers. First, the earlier 5769.99 tok/s figure is +VOID, not superseded by an improvement: that whole series ran against a GPU held +at 11-13% utilization by a graphics consumer the harness's idle check could not +see, and both arms gained ~14% once measured on a genuinely idle box. The check +now fails on non-idle utilization. Second, the residual is a specific missing +mechanism rather than a diffuse gap: this GPU is discrete, so the async-scheduling +device combine/scatter falls back to a host path that must synchronize the main +stream for the sampled ids, and the depth-2 scheduler therefore has nothing to +overlap. Upstream keeps that state GPU-resident unconditionally. Scoped as +ENG-ASYNC-SCHED W4. + This local 4B diagnostic does not establish 27B/35B support. Exact evidence and reproduction: -[Qwen3.5-4B main repair](bench-evidence/qwen35-4b-main-repair-20260725.md). +[Qwen3.5-4B post-pull revalidation](bench-evidence/qwen35-4b-postpull-20260727.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-pinned-oracle-20260728.md b/docs/bench-evidence/qwen35-4b-pinned-oracle-20260728.md new file mode 100644 index 000000000..dfe58b063 --- /dev/null +++ b/docs/bench-evidence/qwen35-4b-pinned-oracle-20260728.md @@ -0,0 +1,134 @@ +# Qwen3.5-4B against an oracle AT the parity pin, 2026-07-28 + +Immutable evidence index for the first local comparison whose denominator is the +project's actual parity pin rather than a pip-installed PyPI release. + +## Headline + +**The published 0.9819x was PESSIMISTIC, not flattering.** vLLM at the pin is +1.25% SLOWER than 0.24.0 on this workload, so measuring against the older release +was understating us. Against the true pin we are **0.9970x** on total throughput +— a 0.3% gap, not 1.8%. + +TPOT remains the failing axis and the real gap: **+12.41%**. + +## Why the oracle was wrong before + +The local oracle was pip-installed on 2026-07-09 and was vLLM **0.24.0**. The +parity pin is `555967922`, a vLLM **main** commit with no release tag and no +prebuilt wheel on ANY platform (x86_64 checked here, aarch64 checked by the pin +spec). `pip install vllm` reaches only PyPI releases, so the local oracle could +never have matched the pin by installation — it has to be built from source. The +records additionally mislabelled the venv as "0.25.0", which hid the drift. + +## Oracle provenance + +- Source: `/home/rich/c/vllm-upstream` at + `5559679229bc961848b121ccdeaa8fa5d79bec98` (2026-07-26), the pin. +- Built: `vllm 0.23.1rc1.dev1511+g555967922.cu132`, editable, sm_120 only + (`TORCH_CUDA_ARCH_LIST=12.0`), venv `.venv-vllm-pin`. The 0.24.0 venv is + untouched and still present. +- Stack, all as the pin specifies: torch 2.13.0+cu130, torchvision 0.28.0, + triton 3.7.1, transformers 5.14.1, flashinfer 0.6.15.post1 (+cubin), + nvidia-cutlass-dsl 4.6.0, tilelang 0.1.9, quack-kernels 0.6.1, + humming-kernels 0.1.10, tokenspeed-mla 0.1.8. +- Attention backend selected at runtime: **FLASH_ATTN**, out of + `['FLASH_ATTN', 'FLASHINFER', 'TRITON_ATTN', 'FLEX_ATTENTION']`. +- Binary SHA-256 (ours) + `31a22725f99926e3ad533baefca2eba225c2cd569069b86979bdecde95b366fe`; + dataset `9ea13603767c62c267e3f381fbccf42d0c9ca0c393655c37533eadca7aefca0c`. +- All 9 performance legs recorded **0% GPU utilization** before starting; one + `flock /tmp/gpu` across the whole 18-leg series. + +### The CUDA toolkit ceiling (recorded — it is not obvious and it cost two builds) + +The usable CUDA version is set by **what the DRIVER can JIT**, not by what is +newest: + +- vLLM ships FlashAttention-2 as `8.0+PTX` + (`cuda_archs_loose_intersection(FA2_ARCHS "8.0+PTX" ...)`), so the driver + JIT-compiles its PTX for Blackwell at load time. Driver 595.71.05 tops out at + CUDA 13.2 and rejects nvcc-13.3 PTX with `cudaErrorUnsupportedPtxVersion` — + which surfaces only at RUNTIME, after a completely clean build. +- CUDA **13.0** is separately unusable: its headers predate glibc 2.42's `rsqrt` + declaration and collide with it (`exception specification is incompatible`). + The fix (`_NV_RSQRT_SPECIFIER`) is present from 13.1 on. +- So **13.2** is the only version clearing both, and the whole toolkit must match + — cccl hard-errors on a mixed compiler/header pair + (`"CUDA compiler and CUDA toolkit headers are incompatible"`). +- Install vLLM with `--no-deps`, or pip re-resolves the CUDA runtime DOWN to + 13.0 after the build and reintroduces the mismatch. +- pip's CUDA wheels ship no unversioned `.so`, so `find_library` fails until dev + symlinks are added (`CUDA_nvrtc_LIBRARY ... NOTFOUND`). + +## Binding result + +Three performance repetitions per arm, interleaved ON / vLLM / OFF. + +| Axis | Direct ON | Direct OFF | vLLM @ pin | ON vs pin | Disposition | +|---|---:|---:|---:|---:|---| +| Total throughput (tok/s) | 6618.160 | 6504.9 | 6638.129 | 0.9970x | FAIL | +| Output throughput (tok/s) | 731.817 | — | 734.026 | 0.9970x | FAIL | +| Requests/s | 5.717 | — | 5.735 | 0.9969x | FAIL | +| Mean TTFT (ms) | 729.217 | — | 943.198 | 0.7731x | PASS, 22.7% lower | +| Mean TPOT (ms) | 38.107 | — | 33.900 | 1.1241x | FAIL, 12.4% higher | +| Peak PSS (GiB) | 2.488 | 8.595 | 8.149 | | PASS | +| Stable PSS (GiB) | 0.759 | 8.591 | 4.565 | | PASS | +| Peak VRAM (MiB) | 12850.7 | 12846.0 | 12826.7 | | ours +24 MiB | + +Repetition spread is tight enough that these ratios are not noise: pin +6640.3 / 6633.4 / 6640.7 (0.11%), ours 6615.3 / 6618.7 / 6620.5 (0.08%). + +## Oracle version delta, measured + +| | vLLM 0.24.0 (2026-07-27) | vLLM @ pin (2026-07-28) | +|---|---:|---:| +| Total throughput (tok/s) | 6722.241 | 6638.129 | +| Mean TTFT (ms) | 913.552 | 943.198 | +| Mean TPOT (ms) | 33.526 | 33.900 | + +The pin is **0.9875x** the 0.24.0 release on total throughput, and slightly worse +on both latency axes. Our own arm is unchanged between the two series +(1.0027x, i.e. reproduces itself), which is the control that makes the oracle +delta attributable to the oracle. + +Consequently our published ratio moves **0.9819x -> 0.9970x** with no change to +our code. + +## Token identity + +| Comparison | Result | +|---|---| +| Ours ON vs OFF, per rep | 128/128, 128/128, 128/128 | +| Ours ON vs the 2026-07-27 series | 128/128, 128/128, 128/128 (our arm reproduces exactly) | +| vLLM @ pin vs vLLM 0.24.0 | 92/128, 92/128, 95/128 | +| Ours vs vLLM @ pin, per rep | 89/128, 89/128, 98/128 (the established near-tie) | + +The oracle's own output moved between versions on ~28% of requests, which is +expected across two vLLM versions and consistent with the oracle not being +self-deterministic on this workload. Correctness remains grounded in the model +gate and the near-tie contract, not in forced benchmark token identity. + +## Known non-fidelity, recorded rather than hidden + +The oracle logs `Failed to import Triton kernels ... No module named +'triton_kernels.matmul_ogs'`. That package is not in vLLM's own +`requirements/cuda.txt`, so a stock install lacks it too — the oracle is faithful +to a stock vLLM, and the affected MoE paths are not on this dense 4B workload. +(This is NOT the repo's own `triton_kernels/` directory shadowing it: the harness +invokes the metrics script by path, so `sys.path[0]` is `tools/bench` and the +repo root is never on `sys.path`. Verified directly.) + +## Reproduction + +```sh +nix develop .#cuda --command bash -c 'REQUIRE_TRITON_AOT=1 \ + CPP_BENCH="$PWD/build-nix-cuda-transplant-triton/examples/vllm-bench" \ + CMAKE_CACHE="$PWD/build-nix-cuda-transplant-triton/CMakeCache.txt" \ + VLLM_PYTHON="$PWD/.venv-vllm-pin/bin/python" \ + VLLM_CUDA_HOME="$PWD/.venv-vllm-pin/lib/python3.12/site-packages/nvidia/cu13" \ + flock /tmp/gpu tools/bench/run_qwen35_4b_compare.sh /tmp/qwen35-pinned-oracle-' +``` + +Root `/tmp/qwen35-pinned-oracle-13729626`, aggregate `aggregate.json`. +No 4B result implies support or speed for the 27B/35B gate checkpoints. diff --git a/docs/bench-evidence/qwen35-4b-postpull-20260727.md b/docs/bench-evidence/qwen35-4b-postpull-20260727.md new file mode 100644 index 000000000..3badd9bf4 --- /dev/null +++ b/docs/bench-evidence/qwen35-4b-postpull-20260727.md @@ -0,0 +1,178 @@ +# Qwen3.5-4B post-pull revalidation, 2026-07-27 + +Immutable evidence index for the local discrete-Blackwell +`LOAD-SAFETENSORS-DIRECT-DENSE` checkpoint after fast-forwarding the workspace +109 commits onto `upstream/main` at `7f620e74`. This is the revalidation +`docs/BENCHMARKS.md` recorded as PENDING for the `c317237a` transplant; that +transplant's two commits are now upstream (`a131de03`, `b6f1efc8`), so the +measured tree is plain current `main`. + +The checkpoint remains **GATING / speed-pending**. Correctness and the +direct-loader memory goal pass. The oracle still wins total/output throughput +and TPOT/ITL, so this is not a parity claim. + +## Headline + +1. **No correctness regression.** Every generated token is IDENTICAL to the + 2026-07-25 run: 128/128 requests per repetition for both the direct-ON and + direct-OFF arms. 109 upstream commits moved no token on this workload. +2. **No meaningful performance regression.** Against the oracle measured in the + SAME series, total throughput went 0.9864x -> 0.9819x and the TPOT excess + went +13.41% -> +14.00%. +3. **The 2026-07-25 series was CONTENDED and its absolute numbers are VOID.** + All nine of its performance legs ran against a GPU already at 11-13% + utilization with 611 MiB of extra resident VRAM. Today's nine legs ran at 0%. + Both arms gained ~14% on the idle box, which is why the RATIO barely moved. + No absolute number from that series may be published; the ratios remain usable + because the contention hit both arms. + +## Workload and environment + +- GPU: NVIDIA GeForce RTX 5070 Ti, sm_120, 16 GiB, driver 595.71.05. + **Discrete**: `is_integrated_gpu()` is FALSE, which selects the host fallback + of the ENG-ASYNC-SCHED combine/scatter (see the residual, below). +- Model: cached `Qwen/Qwen3.5-4B` snapshot `851bf6e8...` under `.hf-cache`. +- Dataset: `/tmp/qwen35-4b-sharegpt-1024.json`, SHA-256 + `9ea13603767c62c267e3f381fbccf42d0c9ca0c393655c37533eadca7aefca0c`. +- Workload: 128 requests, 131,784 input tokens, 128 output tokens per request, + concurrency 32, greedy, prefix caching disabled. +- Project build: `RelWithDebInfo`, CUDA arch `120a`, CUTLASS, FlashAttention-2, + Triton AOT + AOT regeneration. Binary SHA-256 + `c95ecac92acf7f13c4c123de6917626a07fb00773f5fea22d68798522b48427a`; + CMake cache `f03148416721d92ff4592fe271bf877773bec2385b461e575fb6744e8596b527`. +- Oracle: local `.venv-vllm`, **vLLM 0.24.0** (`vllm-0.24.0.dist-info`, and the + series' own `vllm-version.txt`). NOTE: the 2026-07-25 evidence and the + BENCHMARKS entry derived from it called this oracle "vLLM 0.25.0"; that label + was wrong, as that series' own recorded `vllm-version.txt` also says 0.24.0. + The project parity pin is now `555967922` / vLLM 0.26.0.dev0, so this local + denominator is behind the pin and is labelled as such rather than treated as + a pin-era result. +- The whole 18-leg series held one `flock /tmp/gpu`. +- Root: `/tmp/qwen35-postpull-7f620e74`, aggregate `aggregate.json`. + +## Binding result + +Three performance repetitions per arm, interleaved ON / vLLM / OFF. + +| Axis | Direct ON | Direct OFF | vLLM 0.24.0 | ON vs vLLM | Disposition | +|---|---:|---:|---:|---:|---| +| Total throughput (tok/s) | 6600.663 | 6481.337 | 6722.241 | 0.9819x | FAIL | +| Output throughput (tok/s) | 729.883 | 716.687 | 743.327 | 0.9819x | FAIL | +| Requests/s | 5.700 | 5.600 | 5.807 | 0.9815x | FAIL | +| Mean TTFT (ms) | 729.240 | 835.007 | 913.552 | 0.7982x | PASS, 20.2% lower | +| Mean TPOT (ms) | 38.220 | 38.207 | 33.526 | 1.1400x | FAIL, 14.0% higher | +| Mean ITL (ms) | 38.220 | 38.207 | 33.526 | 1.1400x | FAIL | +| Peak PSS (GiB) | 2.282 | 8.594 | 7.666 | | PASS | +| Stable PSS (GiB) | 0.761 | 8.591 | 4.039 | | PASS | +| Peak VRAM (MiB) | 12850.0 | 12846.0 | 12933.3 | | PASS vs vLLM; ON is +4.0 MiB vs OFF | + +Run-to-run spread is very tight, so the ratios above are not noise: our ON arm +measured 6602.6 / 6603.9 / 6595.5 tok/s (0.13% spread) and the vLLM arm +6716.3 / 6722.9 / 6727.6 tok/s (0.17% spread). + +Direct loading still cuts peak host PSS by 73.4% and stable PSS by 91.1% +versus direct-OFF, and it now also LOWERS mean TTFT by 12.7% (835.0 -> 729.2 ms), +which the contended series could not resolve. + +## Token identity + +| Comparison | Result | +|---|---| +| Ours ON, rep-to-rep | 128/128, 128/128, 128/128 (deterministic) | +| Ours OFF, rep-to-rep | 128/128 (deterministic) | +| Ours ON vs OFF, per rep | 128/128, 128/128, 128/128 | +| Ours ON vs 2026-07-25 ON, per rep | **128/128, 128/128, 128/128** | +| Ours OFF vs 2026-07-25 OFF, per rep | **128/128, 128/128, 128/128** | +| vLLM, rep-to-rep | 128/128, 128/128, 102/128 (NOT deterministic) | +| Ours ON vs vLLM, per rep | 87/128, 87/128, 92/128 (the established near-tie) | + +Project-vLLM correctness is grounded in the real model gate and the ratified +near-tie contract, not in forced benchmark token identity, because the oracle is +not self-deterministic on this workload. + +## Contention finding, and the harness hole it came through + +Per-leg `nvidia-smi` snapshots, all nine performance legs of each series: + +| Series | VRAM used before | GPU utilization before | +|---|---|---| +| 2026-07-25 | 1150 MiB, every leg | 11-13%, every leg | +| 2026-07-27 | 539 MiB, every leg | 0%, every leg | + +`prepare_leg` gated idleness on `nvidia-smi --query-compute-apps`, which lists +CUDA contexts ONLY. A graphics consumer does not appear there, so a persistently +busy GPU passed the gate for an entire series. `.agents/benchmark-protocol.md` +voids contended runs, so this was a live measurement defect, not a footnote. + +Fixed in the same change: `prepare_leg` now also fails when `utilization.gpu` +exceeds `GPU_IDLE_UTIL_MAX` (default 2%). Re-parsed against both series, the gate +reads 0 for today (passes) and 12 for 2026-07-25 (would have refused to run). +The new same-binary A/B harness `tools/bench/run_qwen35_4b_ab.sh` carries the +same gate. + +Consequence for the record: the 2026-07-25 ABSOLUTE numbers (5769.99 tok/s and +the rest) are VOID and are superseded by the table above. Its RATIOS, its +same-binary component attributions (H32 AOT +4.5906%, decode graph +0.3873%, +ratio-4 FA2 +1.6004%) and its profiling attribution remain valid, because each +was an internal comparison within one contended-but-uniform series. + +## The residual is unchanged, and now has a spec + +TPOT is the failing axis, at the same relative size as before. The mechanism +identified by the 2026-07-25 node-mode trace still applies verbatim: on a +DISCRETE GPU, `is_integrated_gpu()` is false, so both ENG-ASYNC-SCHED W3 device +call sites (`runner.cpp:821` combine, `runner.cpp:1763` scatter) take their host +fallback, and `sample_tokens_async` must synchronize the main stream before the +host can read the sampled ids. The async scheduler is engaged (the engine logs +`Asynchronous scheduling is enabled (max_concurrent_batches=2)`) but the runner +cannot overlap, so depth-2 buys nothing here. + +Upstream has no such branch: `vllm/v1/worker/gpu/states.py:64` keeps +`last_sampled_tokens` as a GPU tensor unconditionally, and never condenses its +request slots (`states.py:132` uses a free-index pool), which is why it needs no +host round-trip. Scoped as `ENG-ASYNC-SCHED` W4 in +[.agents/specs/async-discrete-device-combine.md](../../.agents/specs/async-discrete-device-combine.md), +including the second per-step barrier that must go with it (the CUDA embedding +out-of-range flag's `cudaStreamSynchronize`). + +## Reproduction + +```sh +nix develop .#cuda --command cmake --build build-nix-cuda-transplant-triton -j6 + +nix develop .#cuda --command bash -c 'flock /tmp/gpu env HF_HOME="$PWD/.hf-cache" \ + LD_LIBRARY_PATH=/run/opengl-driver/lib:$LD_LIBRARY_PATH \ + build-nix-cuda-transplant-triton/tests/test_qwen35_plain_weights --no-skip' + +nix develop .#cuda --command bash -c 'REQUIRE_TRITON_AOT=1 \ + CPP_BENCH="$PWD/build-nix-cuda-transplant-triton/examples/vllm-bench" \ + CMAKE_CACHE="$PWD/build-nix-cuda-transplant-triton/CMakeCache.txt" \ + flock /tmp/gpu tools/bench/run_qwen35_4b_compare.sh /tmp/qwen35-postpull-' + +python3 -m tools.bench.summarize_qwen35_4b_compare \ + --root /tmp/qwen35-postpull- \ + --historical-root /tmp/qwen35-postpull-7f620e74 \ + --output /tmp/qwen35-postpull-/aggregate.json +``` + +The summarizer needed two repairs to run at all: it read the historical token +legs under a `perf-` prefix the harness has never written, and it required a +`vllm_production` key that its own current output does not contain. Both are +fixed in the same change, so a series can now be summarized against the previous +one. + +## Correctness suites, same build, same lock + +| Suite | Result | +|---|---| +| `test_qwen35_plain_weights --no-skip` | 3/3 cases, 1672/1672 assertions | +| `test_ops_gdn` | 66/66 cases, 4242/4242 assertions | +| `test_ops_paged_attn` | 25/25 cases, 454,474/454,474 assertions | +| `test_gdn_packed_decode_triton` | 1/1, 10/10 | +| `test_combine_tokens` | 7/7, 14/14 | +| `test_input_batch` | 22/22, 163/163 | + +Every count matches the 2026-07-25 record exactly. + +No 4B result implies support or speed for the 27B/35B gate checkpoints, which +remain hardware-unavailable on this host. diff --git a/docs/bench-evidence/qwen35-4b-postrebase-20260729.md b/docs/bench-evidence/qwen35-4b-postrebase-20260729.md new file mode 100644 index 000000000..7d42d2f2f --- /dev/null +++ b/docs/bench-evidence/qwen35-4b-postrebase-20260729.md @@ -0,0 +1,111 @@ +# Qwen3.5-4B re-validated after rebasing onto 139 upstream commits, 2026-07-29 + +Immutable evidence index for the re-measurement of the 4B lever after the branch +was rebased from `main` `7f620e74` onto `f3ecbe70d`. + +## Headline + +**Nothing moved.** 139 upstream commits later, every axis reproduces the +2026-07-28 pinned-oracle series inside run-to-run noise. The published ratio +stands: **0.9972x** total throughput (was 0.9970x), TTFT **PASS**, TPOT still the +one failing axis at **+12.5%**. + +This is a null result, and it is the useful kind: it says the rebase did not +regress the lever, and it says none of the upstream work in that window touched +this workload. + +## What was measured + +- Ours: `13906f189` (this branch rebased onto `main` `f3ecbe70d`), clean CUDA + rebuild 925/925, 0 warnings. +- Oracle: `.venv-vllm-pin`, vLLM `0.23.1rc1.dev1511+g555967922` — built from + source at the parity pin. Pin re-checked before the run and unchanged + (`AGENTS.md` still specifies `555967922`), so the same denominator remains + correct. +- Three performance repetitions per arm, interleaved, one `flock /tmp/gpu` across + the whole 18-leg series. +- All nine performance legs recorded **0% GPU utilization** and 539 MiB resident + (the X server) before starting. + +## Binding result + +| Axis | ours | vLLM @ pin | ratio | 2026-07-28 ratio | Disposition | +|---|---:|---:|---:|---:|---| +| Total throughput (tok/s) | 6610.270 | 6628.651 | 0.9972x | 0.9970x | FAIL | +| Output throughput (tok/s) | 730.947 | 732.978 | 0.9972x | 0.9970x | FAIL | +| Requests/s | 5.710 | 5.726 | 0.9971x | 0.9969x | FAIL | +| Mean TTFT (ms) | 729.963 | 947.850 | 0.7701x | 0.7731x | PASS | +| Mean TPOT (ms) | 38.153 | 33.923 | 1.1247x | 1.1241x | FAIL | +| Peak PSS (GiB) | 2.539 | 7.867 | | | PASS | + +Direct-load OFF averages 6493.390 tok/s and 8.592 GiB peak PSS, so direct loading +still cuts host peak PSS by 70.4%. + +Per-repetition total throughput: ours 6610.9 / 6611.6 / 6608.3 (spread 0.05%), +pin 6629.8 / 6625.7 / 6630.4 (0.07%), direct-OFF 6498.8 / 6496.4 / 6484.9 (0.21%). + +## The control that makes "nothing moved" a measurement rather than an assumption + +All three arms drifted down by the same ~0.13% against the previous series: + +| Arm | now | 2026-07-28 | drift | +|---|---:|---:|---:| +| ours (direct ON) | 6610.270 | 6618.160 | 0.9988x | +| ours (direct OFF) | 6493.390 | 6504.600 | 0.9983x | +| vLLM @ pin | 6628.651 | 6638.129 | 0.9986x | + +The vLLM arm is the same pinned binary running the same corpus — **its code did +not change between the two series**, yet it drifted by the same amount as ours. +A uniform drift across an arm that cannot have changed is ambient (thermal/clock +state of the box), not a code effect, and it is the same size as the arms' own +repetition spread. TPOT drifts identically small: ours 1.0012x, the pin 1.0007x. + +So the ~0.1% is the noise floor of this vehicle, and every ratio above sits well +inside it. + +## Token identity + +| Comparison | Result | +|---|---| +| Ours ON vs OFF, per rep | 128/128, 128/128, 128/128 | +| Ours ON vs the 2026-07-28 series | 128/128, 128/128, 128/128 | +| Ours OFF vs the 2026-07-28 series | 128/128, 128/128, 128/128 | +| Ours ON vs vLLM @ pin, per rep | 89/128, 89/128, 89/128 (the established near-tie) | +| vLLM @ pin vs its own 2026-07-28 run | 128/128, 128/128, 95/128 | + +Our output is bit-stable across a 139-commit rebase — the strongest available +statement that the merge resolutions were semantically correct, beyond the +1672/1672 model gate. The oracle's third repetition differing from its own prior +run is the already-recorded fact that vLLM is not self-deterministic on this +workload. + +## Why no movement was the expected outcome + +The one upstream commit in the window that names this gap — +`2b00866a4 explore(perf): decode TPOT/ITL lever` — is an exploration that +concluded the SGLang decode gap is **batch composition, not a decode-kernel +deficiency**. It changed records, not code. The rest of the window is breadth +(DeepSeek-V4-Flash, Gemma-4, Kimi K3, TP, LoRA, AWQ/GPTQ/MXFP4, fp8 KV, xgrammar) +on paths this dense 4B decode workload does not touch. + +TPOT therefore remains owned by ENG-ASYNC-SCHED, and the W4 finding stands: the +lever needs the ASYNC engine loop, which `vllm-bench` does not drive. + +## Reproduction + +```sh +nix develop .#cuda --command bash -c 'REQUIRE_TRITON_AOT=1 \ + CPP_BENCH="$PWD/build-nix-cuda-transplant-triton/examples/vllm-bench" \ + CMAKE_CACHE="$PWD/build-nix-cuda-transplant-triton/CMakeCache.txt" \ + VLLM_PYTHON="$PWD/.venv-vllm-pin/bin/python" \ + VLLM_CUDA_HOME="$PWD/.venv-vllm-pin/lib/python3.12/site-packages/nvidia/cu13" \ + flock /tmp/gpu tools/bench/run_qwen35_4b_compare.sh /tmp/qwen35-postrebase-' + +python3 tools/bench/summarize_qwen35_4b_compare.py \ + --root /tmp/qwen35-postrebase- \ + --historical-root /tmp/qwen35-pinned-oracle-13729626 \ + --output /tmp/qwen35-postrebase-/aggregate.json +``` + +Evidence root `/tmp/qwen35-postrebase-13906f189`, aggregate `aggregate.json`. +No 4B result implies support or speed for the 27B/35B gate checkpoints. diff --git a/docs/bench-evidence/w4-async-mirror-20260727.md b/docs/bench-evidence/w4-async-mirror-20260727.md new file mode 100644 index 000000000..63d6f1489 --- /dev/null +++ b/docs/bench-evidence/w4-async-mirror-20260727.md @@ -0,0 +1,116 @@ +# ENG-ASYNC-SCHED W4 evidence, 2026-07-27 + +Immutable evidence index for the discrete-CUDA device-resident sampled-token +mirror: its correctness gate, its same-binary A/B, and the profiling that +re-attributes the synchronizations the row was originally scoped from. + +Commit `87160099` on branch `bench-lever-sampled-token-20260727`, off `main` +`7f620e74`. GPU: RTX 5070 Ti, sm_120, 16 GiB, driver 595.71.05, DISCRETE +(`is_integrated_gpu()` FALSE). Every run under one `flock /tmp/gpu` on an idle +box (the harness now refuses a leg above 2% utilization). + +## Disposition + +**Correctness GATED. Speed NEUTRAL on the only runnable vehicle, and +structurally unmeasurable there. Ships OPT-IN, DEFAULT OFF +(`VT_ASYNC_DEVICE_MIRROR=1`). The binding serving A/B is PENDING.** + +## Correctness + +| Gate | Result | +|---|---| +| Token identity, W4 ON vs OFF, 3 paired repetitions | **384/384 requests identical** | +| Earlier single pair, W4 OFF vs the pre-W4 baseline | 128/128 identical | +| `test_qwen35_plain_weights --no-skip` | 3/3 cases, 1672/1672 assertions | +| `test_input_batch` | 25/25 cases, 183/183 (includes 3 new W4 cases) | +| `test_combine_tokens` | 7/7, 14/14 | +| `test_ops_gdn` | 66/66, 4242/4242 | +| `test_ops_paged_attn` | 25/25, 454,474/454,474 | +| Clean `-Werror` CUDA rebuild | 0 warnings | +| ASan+UBSan lane, `test_input_batch` (the W4 op log) | 25/25, 183/183, leaks on | + +The W4-OFF-vs-pre-W4-baseline row matters on its own: that binary already +contained W4e (the embedding barrier removal, which is NOT behind the opt-in), so +it establishes that W4e is output-identical to the pre-W4 engine. + +## Speed, same-binary A/B (`/tmp/w4-ab-final`) + +Three interleaved repetitions per arm, order flipped on even repetitions, 20 s +cooldown before each leg, identical corpus (128 requests, 128 output tokens, +concurrency 32, greedy). + +| Axis | W4 ON | W4 OFF | ON/OFF | +|---|---:|---:|---:| +| Total throughput (tok/s) | 6618.047 | 6620.797 | 0.9996 | +| Output throughput (tok/s) | 731.807 | 732.107 | 0.9996 | +| Mean TPOT (ms) | 38.110 | 38.093 | 1.0004 | +| Mean TTFT (ms) | 728.603 | 728.423 | 1.0002 | + +Per-repetition: ON 6610.4 / 6621.3 / 6622.4, OFF 6622.3 / 6622.6 / 6617.5. The +0.04% difference is an order of magnitude inside the arms' own spread. +**NEUTRAL**, and necessarily so — see below. + +## Why it is neutral here, and what that says about the original attribution + +`vllm-bench` drives the SYNCHRONOUS `LLMEngine::step()` loop, which calls +`GPUModelRunner::sample_tokens()`. Only `AsyncLLM`'s depth-2 +`step_with_batch_queue` calls `sample_tokens_async()`. This was established +directly rather than inferred: instrumenting the W4 scatter branch inside +`sample_tokens_async` showed it **never executes** under `vllm-bench` (the first +gate run failed at 0/128 token identity precisely because the mirror was fed only +there and so stayed all-zero). + +So on the benchmarked path there is no depth-2 overlap for W4 to unlock. It adds +four small uploads and two kernels per step and removes nothing, which is exactly +what the A/B measures. + +**This falsifies the attribution the row was scoped from.** The 2026-07-25 +evidence assigned 497 `cudaStreamSynchronize` calls (20.975 s, 42.20 ms/call) to +`sample_tokens_async`'s discrete host-bookkeeping path. That function is not on +the benchmarked path at all. + +A fresh attribution-complete profile (`/tmp/w4-attrib.nsys-rep`, +`nsys profile --cuda-graph-trace=node --trace=cuda`, 32 requests x 64 output +tokens, concurrency 32) says what those synchronizations really are: + +| CUDA API | Calls | Total | Avg | +|---|---:|---:|---:| +| `cudaStreamSynchronize` | 112 | 1.134 s | 10.12 ms | +| `cudaMemcpyAsync` | 1,228 | 0.398 s | 0.32 ms | +| `cudaMalloc` | 818 | 0.029 s | 0.035 ms | +| `cudaFree` | 512 | 0.028 s | 0.054 ms | + +112 synchronizations over ~64 decode steps plus prefill and warm-up is **about +one per engine step**, at ~10 ms each on this smaller batch. Scaled to the +binding workload (~512 steps at ~38 ms TPOT) that is ~500 calls at ~40 ms — the +2026-07-25 numbers almost exactly. The COUNT and the TIME in that trace were +right; only the attribution was wrong. It is the depth-1 engine loop's per-step +wait for its own sampling, not an async-sampler defect. + +The per-call embedding barrier is gone from this trace, which is W4e working: +before it, every step also paid a `cudaMalloc` + `cudaStreamSynchronize` + +`cudaFree` for the out-of-range flag. + +## The lever this actually leaves + +Overlap on this hardware needs the engine to run the ASYNC loop; W4 is what makes +that loop legal on a discrete GPU, and it is now built and gated. The measurement +that would bind it is a SERVING A/B over `AsyncLLM` (`examples/server`), which +`tools/bench/run_serve_low.py` cannot currently drive here — it requires a pinned +SGLang container image and accepts only the 27B/35B model keys, and neither is +available on this host. That harness gap, not W4, is what blocks the number. + +## Reproduction + +```sh +nix develop .#cuda --command cmake --build build-nix-cuda-transplant-triton -j6 + +nix develop .#cuda --command bash -c 'export \ + LD_LIBRARY_PATH=/run/opengl-driver/lib:$LD_LIBRARY_PATH HF_HOME=$PWD/.hf-cache; \ + CPP_BENCH="$PWD/build-nix-cuda-transplant-triton/examples/vllm-bench" \ + CMAKE_CACHE="$PWD/build-nix-cuda-transplant-triton/CMakeCache.txt" \ + flock /tmp/gpu tools/bench/run_qwen35_4b_ab.sh /tmp/w4-ab- \ + w4on VT_ASYNC_DEVICE_MIRROR=1 w4off ""' +``` + +No 4B result implies support or speed for the 27B/35B gate checkpoints. diff --git a/include/vllm/model_executor/models/device_pool.h b/include/vllm/model_executor/models/device_pool.h index caff1929b..511e814c9 100644 --- a/include/vllm/model_executor/models/device_pool.h +++ b/include/vllm/model_executor/models/device_pool.h @@ -47,6 +47,25 @@ namespace vllm { class DevicePool { public: void* Get(vt::Backend& b, size_t bytes) { + // BYPASS lane (VT_POOL_BYPASS=1) — the pool is a DETECTOR BLIND SPOT and + // this is how you see through it. Two ways it hides a real defect from + // compute-sanitizer: + // 1. size-class rounding hands back a block up to 6.25% LARGER than the + // logical tensor, so a write past the last row lands inside the same + // driver allocation and memcheck reports nothing; + // 2. blocks are never returned to the driver, so a use-after-free of a + // released DBuf reads memory that is still legally mapped — and, worse, + // may already have been handed to an unrelated op. + // Under bypass every Get is an EXACT-size driver allocation and every Put is + // a real Free, which restores both boundaries for the detector. It is a + // debugging lane only: it reinstates the per-op cudaMalloc/cudaFree sync + // storm this pool exists to remove, so it is never a timing configuration. + // The backend is remembered here so the no-backend Put overload (the + // cross-step shared_ptr deleter) can free through it. + if (Bypass()) { + backend_ = &b; + return b.Alloc(bytes); + } const size_t key = ClassOf(bytes); { std::lock_guard lk(mu_); @@ -66,6 +85,14 @@ class DevicePool { // logits / MTP hidden handed off via a shared_ptr deleter). Bytes are always // returned to the free list — the cross-step buffers are not cap-evicted. void Put(size_t bytes, void* p) { + // Bypass: free for real so a later use-after-free traps. `backend_` is set by + // the Get that produced `p`, so it is non-null whenever a Put can be reached; + // the null guard keeps the lane from leaking a block if that ever stops + // holding rather than dereferencing a null backend. + if (Bypass()) { + if (backend_ != nullptr) backend_->Free(p); + return; + } const size_t key = ClassOf(bytes); std::lock_guard lk(mu_); retained_ += key; @@ -78,6 +105,10 @@ class DevicePool { // When a discrete GPU sets a bound, scratch over the cap is freed to the driver // rather than pooled, so the reuse pool self-limits without a model edit. void Put(vt::Backend& b, size_t bytes, void* p, size_t cap) { + if (Bypass()) { + b.Free(p); + return; + } const size_t key = ClassOf(bytes); std::lock_guard lk(mu_); if (cap != 0 && retained_ + key > cap) { @@ -100,6 +131,18 @@ class DevicePool { } private: + // VT_POOL_BYPASS=1 turns every Get/Put into a raw driver Alloc/Free (see Get). + // Read once: it must not change between an allocation and its matching free, + // or a pooled block would be handed to Backend::Free (or a driver block leaked + // into the free list). + static bool Bypass() { + static const bool on = [] { + const char* e = std::getenv("VT_POOL_BYPASS"); + return e != nullptr && e[0] == '1'; + }(); + return on; + } + // Round `bytes` up so it keeps at most kClassBits leading significant bits. // Exact keying when VT_POOL_EXACT=1 (A/B). Small sizes (< 2^kClassBits) key // exactly — there are few of them and the waste would be proportionally large. @@ -118,6 +161,10 @@ class DevicePool { } std::mutex mu_; + // Backend the last Get allocated through, so the no-backend Put overload can + // free under bypass. One device per process (see ResolveDevicePoolPolicy), so + // this is stable; unused when bypass is off. + vt::Backend* backend_ = nullptr; std::unordered_map> free_; size_t retained_ = 0; // bytes (class-rounded) held in free_, for the soft cap std::atomic hits_{0}; diff --git a/include/vllm/model_executor/models/model_registry.h b/include/vllm/model_executor/models/model_registry.h index 5f4614cde..dc23dd8c1 100644 --- a/include/vllm/model_executor/models/model_registry.h +++ b/include/vllm/model_executor/models/model_registry.h @@ -217,6 +217,17 @@ struct ModelForwardInput { // tower/merge/MRoPE/DeepStack-conditioned decode. Default-nullopt keeps every // existing (text) call site byte-identical by construction. std::optional mm = std::nullopt; + // ENG-ASYNC-SCHED W4: when non-null, the [token_ids.size()] input ids already + // live in THIS device buffer and `token_ids` is stale for decode rows. The + // async runner's device combine spliced each decode row's sampled token here + // on the main queue, so the host copy never saw it — which is the whole point, + // since materializing it on the host is the synchronize W4 removes. + // + // A model that honors this embeds from the device pointer instead of uploading + // `token_ids`; a model that ignores it is simply never given one (the runner + // only sets it on the discrete-CUDA async path, which the Qwen3.5 gate vehicle + // owns). Null on every other path, so every other forward is byte-identical. + const int32_t* device_token_ids = nullptr; }; using ModelConfigHook = void (*)(const HfConfig& config); diff --git a/include/vllm/v1/worker/gpu/input_batch.h b/include/vllm/v1/worker/gpu/input_batch.h index a01ae9c6c..068e76d4b 100644 --- a/include/vllm/v1/worker/gpu/input_batch.h +++ b/include/vllm/v1/worker/gpu/input_batch.h @@ -235,6 +235,33 @@ class InputBatch { // writes it each step; add_request seeds it for a resumed/PD-disagg request // (0 < num_computed <= prefill_len) so its first decode reads the right id. std::vector last_sampled_tokens; + // ENG-ASYNC-SCHED W4: the ordered log of STRUCTURAL edits made to + // last_sampled_tokens since the runner last drained it — the seed at + // add_request, the row move at condense, the row swap at swap_states. + // + // On a DISCRETE GPU the authoritative copy of last_sampled_tokens is a device + // buffer (upstream keeps it device-resident on every platform, + // states.py:64), so the host no longer holds the VALUES those edits move + // around. It does know the INDICES, which is all a device replay needs. The + // runner drains this each step and applies it in stream order before the + // combine reads the buffer; see vt::cuda::LaunchApplyLastSampledOps. + // + // Upstream needs no equivalent because it never condenses: states.py:132 + // returns a finished request's slot to a free list and the slot index is + // stable for the request's lifetime. This log is the price of our condensed + // dense batch, not a deviation in what the state MEANS. + // + // Inert unless the async runner path is engaged on a device that mirrors the + // array: the ops are recorded unconditionally (a few ints per admitted or + // finished request, off the per-token path) and simply discarded otherwise. + struct LastSampledOp { + enum Kind : int32_t { kSeed = 0, kMove = 1, kSwap = 2 }; + int32_t kind = kSeed; + int32_t a = 0; // seed/move destination, or the first swapped slot + int32_t b = 0; // move source, or the second swapped slot + int32_t value = 0; // seed value; unused by move/swap + }; + std::vector last_sampled_ops; // prefill_len[slot]: the number of tokens KNOWN at admission (prompt + any // pre-existing output = num_tokens() at add_request), fixed for the request's // lifetime. combine gates on seq_len > prefill_len to tell a decode row (splice diff --git a/include/vllm/v1/worker/gpu/runner.h b/include/vllm/v1/worker/gpu/runner.h index 80ff19f17..284cf72d2 100644 --- a/include/vllm/v1/worker/gpu/runner.h +++ b/include/vllm/v1/worker/gpu/runner.h @@ -386,6 +386,8 @@ class GPUModelRunner final : public ModelRunnerBase { // cache is sized by this (one recurrent state per sequence), decoupled from // the attention num_blocks. See remap_gdn_state_slots. int max_num_reqs_ = 0; + // Flattened-token bound for one step; sizes the W4 device input_ids mirror. + int max_num_batched_tokens_ = 0; int64_t gdn_state_slots_ = 0; // Compact GDN state-slot allocator: request identity (req_id) -> slot in // [0, gdn_state_slots_); free list of unused slots. Keyed on the sequence, not @@ -420,6 +422,49 @@ class GPUModelRunner final : public ModelRunnerBase { // sample; freed in the dtor. Mirrors torch's caching device/pinned allocators. std::unique_ptr async_output_pool_; AsyncOutputPool& get_or_create_async_output_pool(); + + // ─── ENG-ASYNC-SCHED W4: discrete-CUDA device-resident async inputs ───────── + // On an INTEGRATED GPU the W3 combine/scatter kernels operate on the runner's + // host arrays in place, because pageable host memory is device-addressable + // there. A DISCRETE GPU cannot do that, and the host fallback has to + // synchronize the main stream to read the sampled ids — which is precisely + // what makes the depth-2 async scheduler overlap nothing on this hardware. + // + // These are the discrete equivalent: the small per-step inputs the combine + // reads, held in persistent device buffers, with last_sampled_tokens the + // AUTHORITATIVE copy (the scatter writes it on the main queue and the next + // step's combine reads it there, so no sampled id ever crosses to the host on + // the critical path). That is exactly what upstream does on every platform — + // vllm/v1/worker/gpu/states.py:64 makes last_sampled_tokens a GPU tensor + // unconditionally. + // + // Sized once from the batch bound; all zero-initialized. The per-step uploads + // copy from the caller's pageable host buffers on purpose (see stage_upload). + struct AsyncDeviceInputs { + int32_t* last_sampled = nullptr; // [max_num_reqs], AUTHORITATIVE + int32_t* prefill_len = nullptr; // [max_num_reqs] + int32_t* query_start_loc = nullptr; // [max_num_reqs + 1] + int32_t* seq_lens = nullptr; // [max_num_reqs] + int32_t* input_ids = nullptr; // [max_num_batched_tokens] + int32_t* ops = nullptr; // [4 * max_num_reqs] structural replay + int64_t input_ids_capacity = 0; // elements in `input_ids` + int32_t max_reqs = 0; + }; + std::unique_ptr async_device_inputs_; + // Allocate on first use, or return nullptr when this device does not need the + // mirror (integrated GPU, non-CUDA backend, or async not engaged). Caller + // treats nullptr as "take the pre-W4 path". + AsyncDeviceInputs* get_or_create_async_device_inputs(); + // True when this runner must mirror the async inputs onto the device: CUDA, + // async engaged, and the platform is NOT integrated. Memoized. + bool async_device_mirror() const; + mutable int async_device_mirror_cached_ = -1; // -1 unknown, 0 no, 1 yes + // Push the recorded InputBatch structural edits (seed/move/swap) to the device + // mirror in stream order, then clear the log. No-op without a mirror. + void replay_last_sampled_ops(AsyncDeviceInputs& dev); + // Upload `src` into `dst` through the pinned staging buffer on the main queue. + void stage_upload(AsyncDeviceInputs& dev, int32_t* dst, const int32_t* src, + int64_t count); // Assemble the [num_reqs, vocab] logits the sampler runs on (the three-case // device/host gather from the stashed forward result) and apply the grammar // bitmask, IN the exact order the sync path uses. Shared by sample_tokens and diff --git a/include/vt/cuda/combine_tokens.h b/include/vt/cuda/combine_tokens.h index 778d570af..562ba0552 100644 --- a/include/vt/cuda/combine_tokens.h +++ b/include/vt/cuda/combine_tokens.h @@ -58,6 +58,28 @@ void LaunchScatterLastSampled(Queue& queue, int32_t* last_sampled_tokens, const int64_t* sampled_ids, const int32_t* idx_mapping, int num_reqs); +// W4 (discrete CUDA): replay InputBatch's STRUCTURAL edits to last_sampled_tokens +// onto the device mirror, in stream order. +// +// Upstream never needs this. vllm/v1/worker/gpu/states.py:132 frees a request's +// slot index into a pool and reuses it, so a request's req_state index is stable +// for its lifetime and the GPU tensor is never permuted. Our InputBatch instead +// CONDENSES (moves the last live row into the freed slot) and swaps rows in the +// decode-first reorder. Once the values live on the device the host cannot +// perform those moves without reading them back — which is the synchronize this +// whole row exists to delete — so the host records what it did and the device +// replays it here. +// +// `ops` is a flat [4 * num_ops] int32 device array of (kind, a, b, value): +// kind 0 SEED: last_sampled[a] = value (add_request) +// kind 1 MOVE: last_sampled[a] = last_sampled[b] (condense) +// kind 2 SWAP: swap(last_sampled[a], last_sampled[b]) (swap_states) +// Applied STRICTLY IN ORDER by a single thread: the ops are not independent (a +// move can read a slot a previous move wrote), and there are at most a handful +// per step, so serial application is both correct and free. +void LaunchApplyLastSampledOps(Queue& queue, int32_t* last_sampled_tokens, + const int32_t* ops, int num_ops); + } // namespace vt::cuda #endif // VT_CUDA_COMBINE_TOKENS_H_ diff --git a/src/vllm/model_executor/models/qwen3_5.cpp b/src/vllm/model_executor/models/qwen3_5.cpp index 2a3cef038..72d64734c 100644 --- a/src/vllm/model_executor/models/qwen3_5.cpp +++ b/src/vllm/model_executor/models/qwen3_5.cpp @@ -336,6 +336,16 @@ int64_t detail::ValidateGdnStateCacheLayout( return state_slots; } +// ENG-ASYNC-SCHED W4 (see qwen3_5_internal.h for why this is a scoped override +// rather than a parameter on five entry points). Thread-local: one host thread +// drives a forward, and a serving process may drive independent engines from +// different threads, so a process-global would let one engine's device ids leak +// into another's embed. +detail::DeviceTokenIds& detail::DeviceTokenIdsOverride() { + thread_local DeviceTokenIds ids; + return ids; +} + vt::DType detail::ResolveMambaSsmCacheDType(const HfConfig& config, vt::DType conv_dtype) { const std::string& dtype = config.mamba_ssm_dtype; @@ -5712,6 +5722,35 @@ MoeBlockOutput RunMoeBlock(vt::Queue& queue, const MoeBlockWeights& weights, return r; } +// ENG-ASYNC-SCHED W4: overwrite the REAL prefix of a freshly uploaded input-id +// buffer with the device-resident ids the async runner's combine produced. +// +// Why patch a prefix instead of embedding straight from the runner's buffer: the +// decode-graph path does not embed `token_ids` as given — it embeds a version +// PADDED up to the captured batch size, whose first B rows are the real requests +// and whose tail is inert. So the correct operation is "replace the first +// ov.count rows", which is exactly right for the padded case AND degenerates to +// "replace everything" on the eager path where ov.count == T. The stale host +// upload that precedes it is a handful of int32s and its real rows are +// immediately overwritten. +// +// The override is CONSUMED here. A forward can reach a second, unrelated embed +// (the multimodal helper embeds a prompt and then single tokens); consuming on +// first use means those cannot be handed ids that were never meant for them. +// The first embed in a registry forward is always the step's own. +static void ApplyDeviceTokenIdsOverride(Dev d, DBuf& dids, int64_t T) { + const detail::DeviceTokenIds ov = detail::DeviceTokenIdsOverride(); + if (ov.ids == nullptr) return; + detail::DeviceTokenIdsOverride() = detail::DeviceTokenIds{}; + // A device buffer LONGER than the embed's input would run past the end. That + // can only mean the runner and the model disagree about this step's shape, so + // fail loudly rather than corrupt the embedding. + VT_CHECK(ov.count <= T, + "qwen3_5 embed: device input ids longer than the embed input"); + d.b.Copy(d.q, dids.ptr(), ov.ids, + static_cast(ov.count) * sizeof(int32_t)); +} + // Embed: hidden[T,H] bf16 = embed_tokens[token_ids] (device-resident table). // KEPT OUTSIDE THE CUDA-GRAPH (M2.5 Phase 2): the CUDA Embedding op allocates a // device bounds-check flag (cudaMalloc/cudaFree) and syncs the stream, all of @@ -5724,7 +5763,14 @@ static void EmbedInto(Dev d, DBuf& hidden, const std::vector& token_ids const int64_t H = config.hidden_size; const int64_t vocab = config.vocab_size; Tensor dtab = ResidentWeight(d, weights.embed_tokens, {vocab, H}); + // ENG-ASYNC-SCHED W4: when the async runner has already placed this step's + // input ids on the device (and spliced each decode row's sampled token into + // them there), embed straight from that buffer. `token_ids` is stale for + // decode rows in that case BY DESIGN — materializing it on the host is the + // synchronize W4 removes — so it must not be uploaded here. Its SIZE is still + // authoritative: the runner sized the device buffer from the same step. DBuf dids(d, DType::kI32, {T}, token_ids.data()); + ApplyDeviceTokenIdsOverride(d, dids, T); vt::Embedding(d.q, hidden.t(), dtab, dids.t()); } @@ -6517,6 +6563,7 @@ static void DenseEmbedInto(Dev d, DBuf& hidden, const int64_t vocab = config.vocab_size; Tensor dtab = ResidentWeight(d, weights.embed_tokens, {vocab, H}); DBuf dids(d, DType::kI32, {T}, token_ids.data()); + ApplyDeviceTokenIdsOverride(d, dids, T); vt::Embedding(d.q, hidden.t(), dtab, dids.t()); } diff --git a/src/vllm/model_executor/models/qwen3_5_dense.cpp b/src/vllm/model_executor/models/qwen3_5_dense.cpp index 13d94f1ab..d09314449 100644 --- a/src/vllm/model_executor/models/qwen3_5_dense.cpp +++ b/src/vllm/model_executor/models/qwen3_5_dense.cpp @@ -18,6 +18,7 @@ #include "vllm/model_executor/models/qwen3_5_common.h" // kQwen3_5Info, helpers #include "vllm/model_executor/models/qwen3_5_dense.h" #include "vllm/model_executor/models/qwen3_5_gguf_weights.h" +#include "qwen3_5_internal.h" // W4 DeviceTokenIdsScope #include "vllm/model_executor/models/qwen3_5_mtp.h" // SPEC-MTP I5d-pre draft #include "vllm/platforms/interface.h" // GetPlatform(device.type) memory-model seam @@ -109,6 +110,14 @@ ForwardLogits ForwardQwen3_5Dense(LoadedModel& model, auto& qwen = static_cast(model); const Qwen3_5DenseWeights& weights = qwen.weights(); + // ENG-ASYNC-SCHED W4: publish the async runner's device-resident input ids for + // the duration of THIS forward, so the embed at the top of every route below + // (eager, gathered, tap, multi-tap, decode-graph replay) reads them instead of + // uploading the host vector, which is stale for decode rows on that path. + // Null on every other path, and RAII-scoped so it cannot outlive the call. + const detail::DeviceTokenIdsScope device_ids_scope( + input.device_token_ids, static_cast(input.token_ids.size())); + // SPEC-MTP I5d-pre hidden-state tap. When the spec verify forward requests the // drafter's [T,H] post-final-norm hidden (I5d), route to the EXISTING // ForwardDeviceTap: byte-identical logits to ForwardDevice, plus the hidden diff --git a/src/vllm/model_executor/models/qwen3_5_internal.h b/src/vllm/model_executor/models/qwen3_5_internal.h index 4c4b9b809..18f41c764 100644 --- a/src/vllm/model_executor/models/qwen3_5_internal.h +++ b/src/vllm/model_executor/models/qwen3_5_internal.h @@ -150,4 +150,44 @@ void ValidateGdnDecodeGraphState( const v1::GDNAttentionMetadata& metadata, const std::vector& state_caches, int64_t real_batch); +// ─── ENG-ASYNC-SCHED W4: device-resident input ids for the embed ───────────── +// +// `ModelForwardInput::device_token_ids` says "the input ids for this step are +// ALREADY on the device; the host vector is stale for decode rows". Only the +// embed at the very top of the forward cares, and it sits under five layers of +// entry points (eager / gathered / tap / multi-tap / decode-graph replay), each +// of which takes `const std::vector& token_ids` and passes it down. +// +// Rather than add a defaulted pointer parameter to every one of those and to the +// decode-graph class, the two Qwen3.5 registry forwards establish this SCOPED +// override for the duration of one forward and the embed consults it. The +// trade-off is deliberate and bounded: it is thread-local (a forward runs on one +// host thread), strictly RAII so it cannot leak past the call that set it, and +// set ONLY from the registry entry points that receive the ModelForwardInput — +// so its lifetime is exactly the forward's, not process state. It is null on +// every path except the discrete-CUDA async runner. +// The COUNT travels with the pointer so the embed can prove the buffer is the +// one meant for it. A forward can reach a second, unrelated embed over different +// ids (the multimodal generate helper embeds a prompt and then single tokens); +// an override that matched on "non-null" alone would silently feed that embed +// the wrong row count. Length disagreement means "not mine" and falls back to +// the host upload, which is always correct. +struct DeviceTokenIds { + const int32_t* ids = nullptr; + int64_t count = 0; +}; + +DeviceTokenIds& DeviceTokenIdsOverride(); + +struct DeviceTokenIdsScope { + DeviceTokenIdsScope(const int32_t* ids, int64_t count) + : prev(DeviceTokenIdsOverride()) { + DeviceTokenIdsOverride() = DeviceTokenIds{ids, count}; + } + ~DeviceTokenIdsScope() { DeviceTokenIdsOverride() = prev; } + DeviceTokenIdsScope(const DeviceTokenIdsScope&) = delete; + DeviceTokenIdsScope& operator=(const DeviceTokenIdsScope&) = delete; + DeviceTokenIds prev; +}; + } // namespace vllm::detail diff --git a/src/vllm/model_executor/models/qwen3_5_moe.cpp b/src/vllm/model_executor/models/qwen3_5_moe.cpp index cbf183fec..51800b52b 100644 --- a/src/vllm/model_executor/models/qwen3_5_moe.cpp +++ b/src/vllm/model_executor/models/qwen3_5_moe.cpp @@ -16,6 +16,7 @@ #include "vllm/model_executor/models/qwen3_5.h" // ForwardLogits, Qwen3_5Model #include "vllm/model_executor/models/qwen3_5_common.h" // kQwen3_5Info, helpers #include "vllm/model_executor/models/qwen3_5_gguf_weights.h" +#include "qwen3_5_internal.h" // W4 DeviceTokenIdsScope #include "vllm/model_executor/models/qwen3_5_mtp.h" // SPEC-MTP I5d-pre draft #include "vllm/model_executor/models/qwen3_5_weights.h" #include "vllm/platforms/interface.h" // GetPlatform(device.type) memory-model seam @@ -91,6 +92,11 @@ ForwardLogits ForwardQwen3_5Moe(LoadedModel& model, auto& qwen = static_cast(model); const Qwen3_5MoeWeights& weights = qwen.weights(); + // ENG-ASYNC-SCHED W4 (see qwen3_5_dense.cpp): scope the async runner's + // device-resident input ids to this forward so the embed reads them. + const detail::DeviceTokenIdsScope device_ids_scope( + input.device_token_ids, static_cast(input.token_ids.size())); + // SPEC-MTP I5d-pre hidden-state tap (see qwen3_5_dense.cpp). Non-null routes to // the EXISTING ForwardDeviceTap (byte-identical logits + the [T,H] post-norm // hidden); null (every spec-off run) is byte-identical to the path below. diff --git a/src/vllm/v1/worker/gpu/input_batch.cpp b/src/vllm/v1/worker/gpu/input_batch.cpp index a231f0a37..a2eaff42d 100644 --- a/src/vllm/v1/worker/gpu/input_batch.cpp +++ b/src/vllm/v1/worker/gpu/input_batch.cpp @@ -208,6 +208,11 @@ int InputBatch::add_request(const CachedRequestState& request) { } else { last_sampled_tokens[static_cast(req_index)] = 0; } + // W4: record the seed so a device mirror starts from the same value. The value + // is host-known here, so this is the one op kind that carries data. + last_sampled_ops.push_back(LastSampledOp{ + LastSampledOp::kSeed, req_index, 0, + last_sampled_tokens[static_cast(req_index)]}); // Sampling metadata (pooling DEFERRED — T0 always has sampling_params). const SamplingParams& sp = request.sampling_params; @@ -652,6 +657,10 @@ void InputBatch::condense() { // to the dense req_state index combine reads). last_sampled_tokens[static_cast(empty_index)] = last_sampled_tokens[static_cast(last_req_index)]; + // W4: the moved VALUE may only exist on the device mirror, so record the + // move by index and let the device replay it in stream order. + last_sampled_ops.push_back( + LastSampledOp{LastSampledOp::kMove, empty_index, last_req_index, 0}); prefill_len[static_cast(empty_index)] = prefill_len[static_cast(last_req_index)]; block_table.move_row(last_req_index, empty_index); @@ -749,6 +758,8 @@ void InputBatch::swap_states(int i1, int i2) { // reorder, so combine's dense req_state index stays correct). std::swap(last_sampled_tokens[static_cast(i1)], last_sampled_tokens[static_cast(i2)]); + // W4: same reasoning as the condense move — indices, not values. + last_sampled_ops.push_back(LastSampledOp{LastSampledOp::kSwap, i1, i2, 0}); std::swap(prefill_len[static_cast(i1)], prefill_len[static_cast(i2)]); diff --git a/src/vllm/v1/worker/gpu/runner.cpp b/src/vllm/v1/worker/gpu/runner.cpp index 7435bea0e..db2fa24c3 100644 --- a/src/vllm/v1/worker/gpu/runner.cpp +++ b/src/vllm/v1/worker/gpu/runner.cpp @@ -310,6 +310,7 @@ GPUModelRunner::GPUModelRunner( group_block_sizes(kv_cache_config), group_block_sizes(kv_cache_config)) { max_num_reqs_ = max_num_reqs; + max_num_batched_tokens_ = max_num_batched_tokens; // SPEC-MTP I5e: the async input-combine splices the device-resident // last_sampled token over each decode row's input id with // num_new_sampled_tokens==1; it is NOT spec-aware and would overwrite the @@ -342,6 +343,7 @@ GPUModelRunner::GPUModelRunner( group_block_sizes(kv_cache_config), group_block_sizes(kv_cache_config)) { max_num_reqs_ = max_num_reqs; + max_num_batched_tokens_ = max_num_batched_tokens; // SPEC-MTP I5e: the async input-combine splices the device-resident // last_sampled token over each decode row's input id with // num_new_sampled_tokens==1; it is NOT spec-aware and would overwrite the @@ -864,6 +866,10 @@ std::optional GPUModelRunner::execute_model( // it is capture-safe. Default OFF: production keeps the byte-identical sync // host path (both give the same id, since sample_tokens writes the same token // to token_ids_cpu and last_sampled_tokens). + // Non-null only on the W4 discrete-CUDA path below: the device input-id buffer + // the combine patched, handed to the forward so it embeds the spliced ids + // instead of the (deliberately stale) host vector. + const int32_t* device_input_ids = nullptr; if (async_input_combine_ && num_reqs > 0) { #ifdef VLLM_CPP_CUDA // S7: the DEVICE combine splices input ids from device-ADDRESSABLE host @@ -888,6 +894,39 @@ std::optional GPUModelRunner::execute_model( input_batch_.last_sampled_tokens.data(), step.query_start_loc.data(), step.seq_lens.data(), input_batch_.prefill_len.data(), num_reqs, /*num_new_sampled_tokens=*/1); + } else if (AsyncDeviceInputs* dev = get_or_create_async_device_inputs(); + dev != nullptr) { + // W4 DISCRETE device combine. Same kernel, same semantics; the difference + // is WHERE the operands live. `last_sampled` is already on the device (the + // previous step's scatter wrote it there and nothing read it back), so the + // three host-known inputs plus the freshly built input_ids are uploaded + // through the pinned staging buffer, the recorded structural edits are + // replayed first so the mirror's row order matches this step's batch, and + // the combine patches the DEVICE input_ids the forward will embed. + // + // Ordering, all on the MAIN queue and therefore exact: replay -> uploads + // -> combine -> forward. The forward is handed `device_input_ids` below, + // so the host copy of step.input_token_ids is deliberately left stale for + // decode rows; nothing on this path reads it (the rejection-sampler path + // that does is spec-only, and spec forces the sync runner). + replay_last_sampled_ops(*dev); + const int64_t num_tokens = + static_cast(step.input_token_ids.size()); + VT_CHECK(num_tokens <= dev->input_ids_capacity, + "async device mirror: step tokens exceed max_num_batched_tokens"); + VT_CHECK(num_reqs <= dev->max_reqs, + "async device mirror: step requests exceed max_num_reqs"); + stage_upload(*dev, dev->input_ids, step.input_token_ids.data(), num_tokens); + stage_upload(*dev, dev->query_start_loc, step.query_start_loc.data(), + static_cast(num_reqs) + 1); + stage_upload(*dev, dev->seq_lens, step.seq_lens.data(), num_reqs); + stage_upload(*dev, dev->prefill_len, input_batch_.prefill_len.data(), + num_reqs); + vt::cuda::LaunchCombineSampledAndDraftTokens( + queue_, dev->input_ids, /*idx_mapping=*/nullptr, dev->last_sampled, + dev->query_start_loc, dev->seq_lens, dev->prefill_len, num_reqs, + /*num_new_sampled_tokens=*/1); + device_input_ids = dev->input_ids; } else #endif { @@ -900,6 +939,14 @@ std::optional GPUModelRunner::execute_model( } } + // W4: the structural-op log has exactly one consumer, the device mirror's + // replay above. On every other configuration (integrated GPU, CPU backend, the + // VT_ASYNC_DEVICE_MIRROR=0 rollback, async off) nothing drains it, so drop it + // here rather than let it grow for the life of a serving process. Deliberately + // NOT cleared when the mirror is on: a step with no requests replays nothing, + // and its ops must survive to the next step that does. + if (!async_device_mirror()) input_batch_.last_sampled_ops.clear(); + // Full-attention KV group metadata (M1.6 MakeCommonAttentionMetadata). int fa_cols = 0; const std::vector fa_bt = @@ -1048,6 +1095,11 @@ std::optional GPUModelRunner::execute_model( exec_state_.spec_aux.layer_ids = dflash_tap_layer_ids_; forward_input.aux_tap = &exec_state_.spec_aux; } + // W4: non-null only on the discrete-CUDA async path, where the combine above + // patched the DEVICE ids and `token_ids` is deliberately stale for decode rows. + // Set after construction because the field sits at the END of the struct, where + // it cannot shift the positional aggregate initializers other callers use. + forward_input.device_token_ids = device_input_ids; // KV-EXTERNAL-CACHE (LMCache): apply any external-prefix loads recorded by the // scheduler's connector for THIS step into the freshly-allocated KV blocks // BEFORE the forward reads them (load-before-compute, base.py:293). Inert when @@ -1383,6 +1435,30 @@ ModelRunnerOutput GPUModelRunner::sample_tokens( input_batch_.last_sampled_tokens[static_cast(i)] = toks.back(); } } + // ENG-ASYNC-SCHED W4: keep the DEVICE mirror coherent when the SYNCHRONOUS + // sampler ran. + // + // The mirror exists so the async sampler can leave the ids on the device. But + // whether the engine drives sample_tokens() or sample_tokens_async() is the + // ENGINE's choice, not the runner's: LLMEngine::step() (which vllm-bench and + // every synchronous embedding of the library use) calls this function, while + // only AsyncLLM's depth-2 step_with_batch_queue calls the async one. Feeding + // the mirror from just one of them leaves the device combine reading a buffer + // nobody wrote — every decode row splices a zero, and the output stream is + // garbage from the first generated token. So whichever sampler ran, the mirror + // is fed here from the values this loop just wrote. + // + // This is an upload rather than a device scatter on purpose: the synchronous + // path ALREADY has the ids on the host (it downloaded them to build `out`), so + // there is no round-trip left to remove, and a few hundred bytes of H2D is + // cheaper and simpler than a second kernel. It also makes the mirror exactly + // the host array on this path, which is what makes the ON/OFF token-identity + // gate meaningful. + if (AsyncDeviceInputs* dinp = get_or_create_async_device_inputs(); + dinp != nullptr && num_reqs > 0) { + stage_upload(*dinp, dinp->last_sampled, + input_batch_.last_sampled_tokens.data(), num_reqs); + } // SPEC-MTP I5d: propose drafts after a plain (no-draft, e.g. first) decode step // so the next step verifies them. Each generating row sampled exactly one token // (num_sampled=1, num_rejected=0); discarded prefill-chunk rows are skipped @@ -1731,6 +1807,19 @@ GPUModelRunner::~GPUModelRunner() { if (async_copy_queue_.id != 0) { vt::DestroyQueue(async_copy_queue_); } + // W4 device mirror. Freed here rather than leaked like the scratch pool: these + // are per-runner, and a serving process can construct more than one runner. + if (async_device_inputs_ != nullptr) { + vt::Backend& b = vt::GetBackend(queue_.device.type); + for (int32_t* p : {async_device_inputs_->last_sampled, + async_device_inputs_->prefill_len, + async_device_inputs_->query_start_loc, + async_device_inputs_->seq_lens, + async_device_inputs_->input_ids, + async_device_inputs_->ops}) { + if (p != nullptr) b.Free(p); + } + } } vt::Queue& GPUModelRunner::get_or_create_async_copy_queue() { @@ -1757,6 +1846,143 @@ AsyncOutputPool& GPUModelRunner::get_or_create_async_output_pool() { return *async_output_pool_; } +// ─── ENG-ASYNC-SCHED W4: discrete-CUDA device-resident async inputs ────────── + +// W4 opt-in, DEFAULT OFF — deliberately, and following the precedent W3 set for +// exactly this situation: the W3 device combine/scatter kernels also landed +// default OFF and were flipped on only after a measured A/B on the hardware they +// targeted. W4's target is the depth-2 overlap of the ASYNC serving loop +// (AsyncLLM -> step_with_batch_queue -> sample_tokens_async). The synchronous +// LLMEngine::step() loop that vllm-bench and every synchronous embedding drive +// has no overlap to unlock, so on that path W4 is correct but cannot pay: it +// adds four small uploads and two kernels per step and removes nothing. Until a +// SERVING A/B shows the win it was built for, production keeps the byte-identical +// host path and this is the switch that turns the mechanism on. +// +// VT_ASYNC_DEVICE_MIRROR=1 engages it. Distinct from VT_ASYNC_RUNNER, which would +// also turn off async scheduling itself; keeping them separate is what makes an +// honest A/B of W4 alone possible — same binary, same scheduler, one mechanism. +#ifdef VLLM_CPP_CUDA +// Guarded with its only use below: on a CPU build the mirror cannot exist, and +// an unused static function is a -Werror=unused-function break there. +static bool AsyncDeviceMirrorEnvDefault() { + const char* value = std::getenv("VT_ASYNC_DEVICE_MIRROR"); + return value != nullptr && value[0] == '1'; +} +#endif + +bool GPUModelRunner::async_device_mirror() const { + if (async_device_mirror_cached_ >= 0) return async_device_mirror_cached_ != 0; + bool on = false; +#ifdef VLLM_CPP_CUDA + // The question is not "is this CUDA" but "is device memory addressable from + // the host", which is exactly what the BACKEND already answers. A unified + // memory device (GB10, and the CPU backend trivially) keeps the W3 in-place + // path, because its host arrays ARE device-addressable and mirroring would + // only add copies; a device with separate memory needs the mirror, because the + // alternative is the host fallback's main-stream Synchronize. Asking the + // capability rather than the device type is also what keeps this file out of + // the shared-layer device-leakage ratchet. + on = async_input_combine_ && AsyncDeviceMirrorEnvDefault() && + !vt::GetBackend(queue_.device.type).UnifiedMemory(); +#endif + async_device_mirror_cached_ = on ? 1 : 0; + return on; +} + +GPUModelRunner::AsyncDeviceInputs* +GPUModelRunner::get_or_create_async_device_inputs() { + if (!async_device_mirror()) return nullptr; + if (async_device_inputs_ != nullptr) return async_device_inputs_.get(); + + const int reqs = max_num_reqs_ > 0 ? max_num_reqs_ : 1; + const int toks = max_num_batched_tokens_ > 0 ? max_num_batched_tokens_ : 1; + vt::Backend& b = vt::GetBackend(queue_.device.type); + auto dev = std::make_unique(); + dev->max_reqs = reqs; + dev->input_ids_capacity = toks; + + auto alloc_i32 = [&](int64_t count) { + void* p = b.Alloc(static_cast(count) * sizeof(int32_t)); + b.Memset(queue_, p, 0, static_cast(count) * sizeof(int32_t)); + return static_cast(p); + }; + dev->last_sampled = alloc_i32(reqs); + dev->prefill_len = alloc_i32(reqs); + dev->query_start_loc = alloc_i32(static_cast(reqs) + 1); + dev->seq_lens = alloc_i32(reqs); + dev->input_ids = alloc_i32(toks); + dev->ops = alloc_i32(4LL * reqs); + + // The mirror starts from whatever the host array already holds. In production + // that is all zeros (no request has been admitted yet), but seeding from the + // host makes the mirror correct even if a runner is switched on mid-flight, + // and it costs one copy for the process. + b.Copy(queue_, dev->last_sampled, input_batch_.last_sampled_tokens.data(), + static_cast(reqs) * sizeof(int32_t)); + // The seeds recorded so far are already reflected by that copy; dropping them + // here keeps the log from replaying them a second time. + input_batch_.last_sampled_ops.clear(); + + async_device_inputs_ = std::move(dev); + return async_device_inputs_.get(); +} + +void GPUModelRunner::stage_upload(AsyncDeviceInputs& dev, int32_t* dst, + const int32_t* src, int64_t count) { + (void)dev; + if (count <= 0) return; + // Copied straight from the caller's PAGEABLE host buffer, deliberately: for a + // pageable source the driver stages the bytes before cudaMemcpyAsync returns, + // so the caller may reuse or destroy that buffer immediately, while the copy + // itself stays ordered on the queue. + // + // A shared PINNED staging buffer would be the usual optimization and is WRONG + // here: pinned copies are truly asynchronous, so the next upload's memcpy into + // the shared buffer could overwrite bytes an in-flight DMA had not yet read, + // and with a depth-2 scheduler that window spans steps. Making that safe needs + // per-upload regions plus an event per step; these arrays are a few kilobytes + // on the front of a step, so the staged pageable copy is the better trade. + vt::GetBackend(queue_.device.type) + .Copy(queue_, dst, src, static_cast(count) * sizeof(int32_t)); +} + +void GPUModelRunner::replay_last_sampled_ops(AsyncDeviceInputs& dev) { +#ifdef VLLM_CPP_CUDA + std::vector& ops = input_batch_.last_sampled_ops; + if (ops.empty()) return; + // Flatten to (kind, a, b, value) quads. The log is bounded by the number of + // admissions/removals in one step, so it fits the [4 * max_num_reqs] buffer; + // if a step ever exceeded that, replaying a truncated prefix would silently + // corrupt the mirror, so drain it in whole chunks instead. + const int64_t cap_ops = 4LL * dev.max_reqs; + size_t done = 0; + std::vector flat; + while (done < ops.size()) { + const size_t chunk = + std::min(ops.size() - done, static_cast(dev.max_reqs)); + flat.clear(); + flat.reserve(chunk * 4); + for (size_t i = 0; i < chunk; ++i) { + const InputBatch::LastSampledOp& op = ops[done + i]; + flat.push_back(op.kind); + flat.push_back(op.a); + flat.push_back(op.b); + flat.push_back(op.value); + } + VT_CHECK(static_cast(flat.size()) <= cap_ops, + "async device mirror: structural-op chunk exceeds its buffer"); + stage_upload(dev, dev.ops, flat.data(), static_cast(flat.size())); + vt::cuda::LaunchApplyLastSampledOps(queue_, dev.last_sampled, dev.ops, + static_cast(chunk)); + done += chunk; + } + ops.clear(); +#else + (void)dev; +#endif +} + std::unique_ptr GPUModelRunner::sample_tokens_async( const std::optional& grammar_output) { // When async is NOT engaged (production default), degenerate to the byte- @@ -1842,6 +2068,33 @@ std::unique_ptr GPUModelRunner::sample_tokens_async( } input_batch_.num_tokens_no_spec[static_cast(i)] += 1; } + } else if (AsyncDeviceInputs* dinp = get_or_create_async_device_inputs(); + dinp != nullptr) { + // W4 DISCRETE device scatter. Identical in effect to the integrated branch, + // but the destination is the device mirror rather than a device-addressable + // host array. This is what DELETES the `Synchronize` below on a discrete GPU: + // the sampled ids stay on the device, main-stream-ordered with the next + // step's combine, and the host bookkeeping in this loop needs no device read. + // + // The host `last_sampled_tokens` is deliberately NOT updated here — it would + // require reading the ids back, which is the cost being removed. It stays + // valid as a structural array (its rows still move with condense/swap, and + // those moves are replayed onto the mirror), but its VALUES are stale on this + // path and nothing reads them; the sampled ids reach the engine through the + // async output's own copy, as upstream does. + vt::cuda::LaunchScatterLastSampled(queue_, dinp->last_sampled, + static_cast(dev_ids), + /*idx_mapping=*/nullptr, num_reqs); + for (int i = 0; i < num_reqs; ++i) { + const std::string& req_id = exec_state_.req_ids[static_cast(i)]; + skeleton.req_ids.push_back(req_id); + skeleton.req_id_to_index[req_id] = i; + if (i < static_cast(exec_state_.discard.size()) && + exec_state_.discard[static_cast(i)]) { + continue; + } + input_batch_.num_tokens_no_spec[static_cast(i)] += 1; + } } else #endif { diff --git a/src/vt/cuda/cuda_combine_tokens.cu b/src/vt/cuda/cuda_combine_tokens.cu index ed4e2a58c..d3789cc67 100644 --- a/src/vt/cuda/cuda_combine_tokens.cu +++ b/src/vt/cuda/cuda_combine_tokens.cu @@ -85,6 +85,29 @@ __global__ void ScatterLastSampledKernel(int32_t* last_sampled_tokens, last_sampled_tokens[req_state_idx] = static_cast(sampled_ids[i]); } +// W4 structural replay (see combine_tokens.h). ONE thread, strictly in order: +// the ops are dependent (a condense move can read a slot an earlier move wrote) +// and there are a handful per step at most. +__global__ void ApplyLastSampledOpsKernel(int32_t* last_sampled_tokens, + const int32_t* ops, int num_ops) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + for (int i = 0; i < num_ops; ++i) { + const int32_t kind = ops[4 * i + 0]; + const int32_t a = ops[4 * i + 1]; + const int32_t b = ops[4 * i + 2]; + const int32_t value = ops[4 * i + 3]; + if (kind == 0) { + last_sampled_tokens[a] = value; + } else if (kind == 1) { + last_sampled_tokens[a] = last_sampled_tokens[b]; + } else if (kind == 2) { + const int32_t tmp = last_sampled_tokens[a]; + last_sampled_tokens[a] = last_sampled_tokens[b]; + last_sampled_tokens[b] = tmp; + } + } +} + } // namespace void LaunchCombineSampledAndDraftTokens(Queue& queue, int32_t* input_ids, @@ -112,4 +135,12 @@ void LaunchScatterLastSampled(Queue& queue, int32_t* last_sampled_tokens, Check(cudaGetLastError(), "ScatterLastSampledKernel launch"); } +void LaunchApplyLastSampledOps(Queue& queue, int32_t* last_sampled_tokens, + const int32_t* ops, int num_ops) { + if (num_ops <= 0) return; + ApplyLastSampledOpsKernel<<<1, 1, 0, AsStream(queue)>>>(last_sampled_tokens, ops, + num_ops); + Check(cudaGetLastError(), "ApplyLastSampledOpsKernel launch"); +} + } // namespace vt::cuda diff --git a/src/vt/cuda/cuda_ops.cu b/src/vt/cuda/cuda_ops.cu index df4793541..548aebcc9 100644 --- a/src/vt/cuda/cuda_ops.cu +++ b/src/vt/cuda/cuda_ops.cu @@ -635,8 +635,79 @@ cudaError_t LaunchEmbeddingIn(cudaStream_t s, Tensor& out, const Tensor& table, return LaunchEmbedding(s, out, table, ids, err); } +// Out-of-range reporting WITHOUT a per-call barrier. +// +// The original shape was cudaMalloc + kernel + D2H + cudaStreamSynchronize + +// cudaFree on EVERY call. All three of those driver calls SYNCHRONIZE the device, +// and embedding runs once per engine step at the very front of the forward, so +// the sequence is a hard barrier between consecutive steps. It costs almost +// nothing while the engine is serialized anyway (measured 23.7 us/call), and it +// costs the ENTIRE overlap the moment the engine stops being serialized — which +// is exactly what ENG-ASYNC-SCHED is for. See +// .agents/specs/async-discrete-device-combine.md W4e. +// +// The replacement is a small RING of persistent slots. Each slot owns a device +// flag, a PINNED host mirror and an event; a call takes the next slot, resets and +// launches into it, and records the event — no sync, no allocation. The check is +// then DEFERRED: before reusing a slot, its event is consumed. If the event has +// already completed the check is free (cudaEventQuery); only a genuinely +// in-flight slot blocks, which bounds the ring's memory without ever making the +// common path wait. +// +// Semantics change in exactly one way, deliberately: a bad id is reported up to +// kSlots calls later than it was, instead of on the offending call. It is still +// LOUD (the same exception, the same message, on the same queue) and it is still +// memory-safe on the offending call itself, because the kernel clamps the gather +// — an out-of-range id has never produced an out-of-bounds read, only a wrong +// row. The one report the ring cannot deliver is an error in the final embedding +// of a process that then exits without another embedding; that residue is +// covered by the token-exactness gates, which compare the produced ids. +struct EmbeddingErrSlot { + EmbeddingErr* dev = nullptr; // device-side flag the kernel atomically sets + EmbeddingErr* host = nullptr; // pinned mirror the async D2H lands in + cudaEvent_t done = nullptr; // completion of that D2H + bool pending = false; // event recorded and not yet consumed + int64_t vocab = 0; // table rows of the call that armed the slot +}; + +// One embedding call per engine step per queue, so a handful of slots covers any +// realistic overlap depth; the ring only has to outlive the in-flight window. +constexpr int kEmbeddingErrSlots = 4; + +struct EmbeddingErrRing { + EmbeddingErrSlot slots[kEmbeddingErrSlots]; + int next = 0; +}; + +// One ring per process: the forward runs on a single device and a single host +// thread (the same assumption DevicePool and the resident-weight caches make). +EmbeddingErrRing& ErrRing() { + static EmbeddingErrRing ring; + return ring; +} + +// Consume a slot's outstanding result, blocking only if `force`. Returns the +// error to report, if any, WITHOUT throwing: the caller decides when to throw so +// a partially-armed slot is never left behind. +bool ConsumeEmbeddingErr(EmbeddingErrSlot& slot, bool force, EmbeddingErr* err_out, + int64_t* vocab_out) { + if (!slot.pending) return false; + if (!force) { + const cudaError_t q = cudaEventQuery(slot.done); + if (q == cudaErrorNotReady) return false; // still in flight: check it later + if (q != cudaSuccess) Check(q, "embedding flag event query"); + } else { + Check(cudaEventSynchronize(slot.done), "embedding flag event sync"); + } + slot.pending = false; + if (slot.host->status == 0) return false; + *err_out = *slot.host; + *vocab_out = slot.vocab; + return true; +} + void EmbeddingKernelCuda(Queue& q, Tensor& out, const Tensor& table, const Tensor& ids) { - // Validate dtypes before allocating the flag buffer so a throw cannot leak it. + // Validate dtypes before touching the ring so a throw cannot leave a slot armed. VT_CHECK(table.dtype == DType::kF32 || table.dtype == DType::kBF16, "cuda embedding: unsupported table dtype (f32/bf16 only)"); VT_CHECK(out.dtype == DType::kF32 || out.dtype == DType::kBF16, @@ -648,23 +719,65 @@ void EmbeddingKernelCuda(Queue& q, Tensor& out, const Tensor& table, const Tenso VT_CHECK(table.shape[0] > 0, "cuda embedding: empty table (vocab 0) with nonempty ids"); cudaStream_t s = AsStream(q); - EmbeddingErr* derr = nullptr; - Check(cudaMalloc(&derr, sizeof(EmbeddingErr)), "cudaMalloc embedding flag"); - EmbeddingErr herr{}; - cudaError_t st = cudaMemsetAsync(derr, 0, sizeof(EmbeddingErr), s); + EmbeddingErrRing& ring = ErrRing(); + EmbeddingErrSlot& slot = ring.slots[ring.next]; + ring.next = (ring.next + 1) % kEmbeddingErrSlots; + + if (slot.dev == nullptr) { + // First use of this slot: the ONLY allocation this path ever makes. Pinned + // host memory so the D2H is a real async copy rather than a staged one. + Check(cudaMalloc(&slot.dev, sizeof(EmbeddingErr)), "cudaMalloc embedding flag"); + Check(cudaHostAlloc(reinterpret_cast(&slot.host), sizeof(EmbeddingErr), + cudaHostAllocDefault), + "cudaHostAlloc embedding flag mirror"); + Check(cudaEventCreateWithFlags(&slot.done, cudaEventDisableTiming), + "cudaEventCreate embedding flag"); + slot.host->status = 0; + } + + // Reusing this slot means its previous result must be consumed first. Force the + // wait: the slot is about to be overwritten, so skipping the check here would + // DROP a report rather than defer it. With kSlots slots in the ring this only + // blocks when more than kSlots embeddings are genuinely in flight. + EmbeddingErr prev{}; + int64_t prev_vocab = 0; + const bool had_prev = ConsumeEmbeddingErr(slot, /*force=*/true, &prev, &prev_vocab); + + cudaError_t st = cudaMemsetAsync(slot.dev, 0, sizeof(EmbeddingErr), s); if (st == cudaSuccess) { - st = table.dtype == DType::kF32 ? LaunchEmbeddingIn(s, out, table, ids, derr) - : LaunchEmbeddingIn<__nv_bfloat16>(s, out, table, ids, derr); + st = table.dtype == DType::kF32 + ? LaunchEmbeddingIn(s, out, table, ids, slot.dev) + : LaunchEmbeddingIn<__nv_bfloat16>(s, out, table, ids, slot.dev); } if (st == cudaSuccess) { - st = cudaMemcpyAsync(&herr, derr, sizeof(EmbeddingErr), cudaMemcpyDeviceToHost, s); + st = cudaMemcpyAsync(slot.host, slot.dev, sizeof(EmbeddingErr), cudaMemcpyDeviceToHost, s); } - if (st == cudaSuccess) st = cudaStreamSynchronize(s); - cudaFree(derr); // best-effort; the primary error (if any) is reported below + if (st == cudaSuccess) st = cudaEventRecord(slot.done, s); + if (st == cudaSuccess) { + slot.pending = true; + slot.vocab = table.shape[0]; + } + + // A launch/copy failure is reported before a deferred out-of-range id: it is + // the more immediate fault, and it may be why the older flag never arrived. Check(st, "embedding"); - if (herr.status != 0) { - throw std::runtime_error("vt cuda: embedding: id " + std::to_string(herr.id) + - " out of range [0, " + std::to_string(table.shape[0]) + ")"); + if (had_prev) { + throw std::runtime_error("vt cuda: embedding: id " + std::to_string(prev.id) + + " out of range [0, " + std::to_string(prev_vocab) + ")"); + } + + // Opportunistically drain every OTHER slot whose copy has already landed, so a + // bad id surfaces at the next call rather than only when its slot comes round + // again. Free: cudaEventQuery on a completed event does not block. + for (int i = 0; i < kEmbeddingErrSlots; ++i) { + EmbeddingErrSlot& other = ring.slots[i]; + if (&other == &slot) continue; + EmbeddingErr err{}; + int64_t vocab = 0; + if (ConsumeEmbeddingErr(other, /*force=*/false, &err, &vocab)) { + throw std::runtime_error("vt cuda: embedding: id " + std::to_string(err.id) + + " out of range [0, " + std::to_string(vocab) + ")"); + } } } diff --git a/tests/vllm/v1/worker/test_input_batch.cpp b/tests/vllm/v1/worker/test_input_batch.cpp index 62e1df1f6..61bad7d09 100644 --- a/tests/vllm/v1/worker/test_input_batch.cpp +++ b/tests/vllm/v1/worker/test_input_batch.cpp @@ -267,6 +267,124 @@ TEST_CASE("swap_states swaps last_sampled_tokens + prefill_len") { CHECK(batch.prefill_len[0] == 1); } +// ─── ENG-ASYNC-SCHED W4: the structural-op log a device mirror replays ─────── +// +// On a discrete GPU last_sampled_tokens lives on the device, so the host cannot +// perform condense's row move or swap_states' row swap itself — it does not hold +// the values any more. It records what it did; the runner replays the record +// onto the device buffer in stream order. These cases pin the record, because a +// missing or misordered op silently feeds the NEXT step's combine the wrong +// request's token, which shows up as a corrupted output stream and nothing else. +TEST_CASE("W4: add_request records the last_sampled seed with its value") { + InputBatch batch = make_batch(); + using Op = InputBatch::LastSampledOp; + + CachedRequestState fresh = make_req("fresh", {10, 11, 12}, {}, {3}); + fresh.num_computed_tokens = 0; + const int i0 = batch.add_request(fresh); + + REQUIRE(batch.last_sampled_ops.size() == 1); + CHECK(batch.last_sampled_ops[0].kind == Op::kSeed); + CHECK(batch.last_sampled_ops[0].a == i0); + CHECK(batch.last_sampled_ops[0].value == 0); + + // A resumed request seeds a real token, and the op must carry that VALUE (it + // is the one op kind whose data the device cannot derive from indices). + CachedRequestState resumed = make_req("resumed", {20, 21, 22}, {23, 24}, {7}); + resumed.num_computed_tokens = 4; + const int i1 = batch.add_request(resumed); + + REQUIRE(batch.last_sampled_ops.size() == 2); + CHECK(batch.last_sampled_ops[1].kind == Op::kSeed); + CHECK(batch.last_sampled_ops[1].a == i1); + CHECK(batch.last_sampled_ops[1].value == 23); + CHECK(batch.last_sampled_ops[1].value == + batch.last_sampled_tokens[static_cast(i1)]); +} + +TEST_CASE("W4: condense records the row move, swap_states records the swap") { + InputBatch batch = make_batch(); + using Op = InputBatch::LastSampledOp; + + CachedRequestState r0 = make_req("r0", {100}, {}, {10}); + r0.num_computed_tokens = 1; + CachedRequestState r1 = make_req("r1", {200, 201}, {}, {20}); + r1.num_computed_tokens = 2; + CachedRequestState r2 = make_req("r2", {300, 301}, {302}, {30}); + r2.num_computed_tokens = 3; + batch.add_request(r0); + batch.add_request(r1); + batch.add_request(r2); + + // Drain the admission seeds the way the runner does, so what follows is only + // the structural edits under test. + batch.last_sampled_ops.clear(); + + batch.remove_request("r1"); + batch.condense(); // r2 (slot 2) slides into the freed slot 1 + + REQUIRE(batch.last_sampled_ops.size() == 1); + CHECK(batch.last_sampled_ops[0].kind == Op::kMove); + CHECK(batch.last_sampled_ops[0].a == 1); // destination + CHECK(batch.last_sampled_ops[0].b == 2); // source + // Replaying that move on a mirror must reproduce what the host array holds. + CHECK(batch.last_sampled_tokens[1] == 302); + + batch.last_sampled_ops.clear(); + batch.swap_states(0, 1); + + REQUIRE(batch.last_sampled_ops.size() == 1); + CHECK(batch.last_sampled_ops[0].kind == Op::kSwap); + CHECK(batch.last_sampled_ops[0].a == 0); + CHECK(batch.last_sampled_ops[0].b == 1); +} + +TEST_CASE("W4: replaying the op log reproduces the host array exactly") { + // The whole contract in one place: run a realistic admit / finish / reorder + // sequence, replay the recorded ops onto an INDEPENDENT array with the same + // semantics the device kernel implements, and require the two to agree. This + // is what catches an op that is recorded in the wrong ORDER — each op + // individually looks right, and only the composition disagrees. + InputBatch batch = make_batch(); + using Op = InputBatch::LastSampledOp; + std::vector mirror(batch.last_sampled_tokens.size(), 0); + + auto replay = [&] { + for (const Op& op : batch.last_sampled_ops) { + const size_t a = static_cast(op.a); + const size_t b = static_cast(op.b); + if (op.kind == Op::kSeed) { + mirror[a] = op.value; + } else if (op.kind == Op::kMove) { + mirror[a] = mirror[b]; + } else { + std::swap(mirror[a], mirror[b]); + } + } + batch.last_sampled_ops.clear(); + }; + + CachedRequestState a = make_req("a", {1, 2}, {3}, {10}); + a.num_computed_tokens = 3; // seeds 3 + CachedRequestState b = make_req("b", {5, 6}, {7}, {20}); + b.num_computed_tokens = 3; // seeds 7 + CachedRequestState c = make_req("c", {8, 9}, {11}, {30}); + c.num_computed_tokens = 3; // seeds 11 + batch.add_request(a); + batch.add_request(b); + batch.add_request(c); + replay(); + CHECK(mirror == batch.last_sampled_tokens); + + // Reorder, then finish the middle request and condense: the move must be + // replayed AFTER the swap, or the mirror picks up the pre-swap occupant. + batch.swap_states(0, 2); + batch.remove_request("b"); + batch.condense(); + replay(); + CHECK(mirror == batch.last_sampled_tokens); +} + TEST_CASE("condense is a no-op when only the last request was removed") { InputBatch batch = make_batch(); batch.add_request(make_req("a", {1}, {}, {0})); diff --git a/tests/vt/test_cuda_ops.cpp b/tests/vt/test_cuda_ops.cpp index eb2bf5ce1..0ea99e09e 100644 --- a/tests/vt/test_cuda_ops.cpp +++ b/tests/vt/test_cuda_ops.cpp @@ -409,7 +409,7 @@ TEST_CASE("CUDA embedding matches CPU (i32 and i64 ids)") { } } -TEST_CASE("CUDA embedding: out-of-range device id throws with the id") { +TEST_CASE("CUDA embedding: out-of-range device id is reported with the id") { if (!HasCuda()) { MESSAGE("no CUDA backend registered; skipping"); return; @@ -418,25 +418,55 @@ TEST_CASE("CUDA embedding: out-of-range device id throws with the id") { const int64_t v = 8, h = 4; const auto tf = RandomF32(static_cast(v * h), 42); + // CONTRACT (changed deliberately, see cuda_ops.cu EmbeddingErrRing): the CUDA + // embedding no longer synchronizes the stream to read its out-of-range flag, + // because that made a once-per-step op a hard barrier between engine steps. + // The report is now DEFERRED: an out-of-range id is raised no later than the + // NEXT Embedding on the queue, carrying the same message and the same id. The + // gather itself is unchanged — bad ids are still clamped in-kernel, so the + // offending call never reads out of bounds. + // + // The Synchronize below is what makes this deterministic rather than a race + // with the flag's device-to-host copy: after it the copy has certainly landed, + // so the next Embedding MUST raise. for (int32_t bad : {int32_t{8}, int32_t{-3}}) { CAPTURE(bad); - std::vector ids = {1, bad, 2}; + std::vector bad_ids = {1, bad, 2}; + const std::vector good_ids = {1, 0, 2}; QueueGuard gq(gpu); DeviceTensor dtab(gpu, gq.q, DType::kF32, {v, h}, tf.data()); - DeviceTensor dids(gpu, gq.q, DType::kI32, {3}, ids.data()); + DeviceTensor dbad(gpu, gq.q, DType::kI32, {3}, bad_ids.data()); + DeviceTensor dgood(gpu, gq.q, DType::kI32, {3}, good_ids.data()); DeviceTensor dout(gpu, gq.q, DType::kF32, {3, h}); + + // The offending call is allowed to raise here or to defer; either is the + // contract, so only the id is asserted when it does raise. bool threw = false; + std::string msg; try { - vt::Embedding(gq.q, dout.tensor(), dtab.tensor(), dids.tensor()); + vt::Embedding(gq.q, dout.tensor(), dtab.tensor(), dbad.tensor()); } catch (const std::runtime_error& e) { threw = true; - const std::string msg = e.what(); - CAPTURE(msg); - CHECK(msg.find("embedding") != std::string::npos); - CHECK(msg.find(std::to_string(bad)) != std::string::npos); + msg = e.what(); } + // The kernel clamps bad ids, so the stream stays healthy either way. + CHECK_NOTHROW(gpu.Synchronize(gq.q)); + + if (!threw) { + try { + vt::Embedding(gq.q, dout.tensor(), dtab.tensor(), dgood.tensor()); + } catch (const std::runtime_error& e) { + threw = true; + msg = e.what(); + } + } + CAPTURE(msg); CHECK(threw); - // The kernel clamps bad ids, so the stream stays healthy after the throw. + CHECK(msg.find("embedding") != std::string::npos); + CHECK(msg.find(std::to_string(bad)) != std::string::npos); + // Once reported, the error is CONSUMED: a subsequent clean call is clean, + // so a single bad id cannot poison every later step. + CHECK_NOTHROW(vt::Embedding(gq.q, dout.tensor(), dtab.tensor(), dgood.tensor())); CHECK_NOTHROW(gpu.Synchronize(gq.q)); } } diff --git a/tools/bench/run_qwen35_4b_ab.sh b/tools/bench/run_qwen35_4b_ab.sh new file mode 100755 index 000000000..9e1a24129 --- /dev/null +++ b/tools/bench/run_qwen35_4b_ab.sh @@ -0,0 +1,118 @@ +#!/usr/bin/env bash +# Same-binary A/B on the exact Qwen3.5-4B comparison workload. +# +# The comparison harness (run_qwen35_4b_compare.sh) answers "where are we versus +# vLLM". This answers "did THIS toggle move us", which needs a different shape: +# one binary, one lock across the WHOLE series, the two arms interleaved so a +# thermal or clock drift hits both arms equally, and the generated token ids +# captured per leg so an arm that changed the OUTPUT is caught rather than +# celebrated as a speedup. +# +# Invoke once, under the lock, from inside the CUDA dev shell: +# flock /tmp/gpu tools/bench/run_qwen35_4b_ab.sh /tmp/qwen35-ab- \ +# A_NAME 'ENV=VAL ...' B_NAME 'ENV=VAL ...' +# +# Both arms get the identical request corpus, sampling and concurrency; the ONLY +# difference is the per-arm environment assignments. +set -euo pipefail + +root=$(git rev-parse --show-toplevel) +cd "$root" +out=${1:?usage: run_qwen35_4b_ab.sh OUTPUT_DIR A_NAME A_ENV B_NAME B_ENV} +a_name=${2:?missing A arm name} +a_env=${3?missing A arm env} +b_name=${4:?missing B arm name} +b_env=${5?missing B arm env} +reps=${REPS:-3} + +model=${MODEL:-$root/.hf-cache/hub/models--Qwen--Qwen3.5-4B/snapshots/851bf6e806efd8d0a36b00ddf55e13ccb7b8cd0a} +dataset=${DATASET:-/tmp/qwen35-4b-sharegpt-1024.json} +cpp=${CPP_BENCH:-$root/build-nix-cuda-transplant-triton/examples/vllm-bench} +cmake_cache=${CMAKE_CACHE:-$(dirname "$(dirname "$cpp")")/CMakeCache.txt} + +test ! -e "$out" || { echo "refusing to overwrite $out" >&2; exit 2; } +test -x "$cpp" +test -d "$model" +test -f "$dataset" +test -f "$cmake_cache" +test -e /run/opengl-driver/lib/libcuda.so.1 +git -C "$root" diff --quiet +mkdir -p "$out" + +# Same rationale as the comparison harness: the live host driver must precede +# Nix's link-only libcuda stub or a Triton AOT object silently misses CUDA +# registration and the arm measures a CPU fallback. +ld_path=/run/opengl-driver/lib +if test -n "${LD_LIBRARY_PATH:-}"; then + ld_path="$ld_path:$LD_LIBRARY_PATH" +fi + +cpp_args=( + --model "$model" --dataset-path "$dataset" --num-prompts 128 + --output-len 128 --concurrency 32 --temperature 0 + --max-num-batched-tokens 2048 --num-blocks 1280 +) + +git -C "$root" rev-parse HEAD >"$out/commit.txt" +git -C "$root" status --porcelain=v1 >"$out/git-status.txt" +sha256sum "$dataset" "$cpp" "$cmake_cache" >"$out/sha256.txt" +grep -E '^(CMAKE_BUILD_TYPE|VLLM_CPP_CUDA_ARCHITECTURES|VLLM_CPP_FLASH_ATTN|VLLM_CPP_SANITIZE|VLLM_CPP_TRITON|VLLM_CPP_TRITON_REGEN|VLLM_CPP_TRITON_VENDORED_ARCH):' \ + "$cmake_cache" >"$out/cpp-build-config.txt" +printf 'A=%s env=[%s]\nB=%s env=[%s]\nreps=%s\n' \ + "$a_name" "$a_env" "$b_name" "$b_env" "$reps" >"$out/arms.txt" + +gpu_snapshot() { + nvidia-smi --query-gpu=name,driver_version,pstate,memory.used,utilization.gpu,temperature.gpu,power.draw \ + --format=csv,noheader >"$1.csv" + nvidia-smi --query-compute-apps=pid,process_name,used_memory \ + --format=csv,noheader >"$1.compute-apps.csv" +} + +run_leg() { + local name=$1 arm_env=$2 rep=$3 + local leg="$name-r$rep" + # Cooldown FIRST, then snapshot. Without it the second arm inherits the first + # arm's thermal and clock state and the A/B measures the order, not the toggle + # - and the idle check below would read the previous leg's GPU still draining + # and refuse to run at all, which is exactly what it did when this slept after + # sampling instead of before. + sleep "${COOLDOWN:-20}" + gpu_snapshot "$out/$leg.gpu-before" + if test -s "$out/$leg.gpu-before.compute-apps.csv"; then + echo "GPU is not compute-idle before $leg" >&2 + return 1 + fi + # Same hole the comparison harness had: --query-compute-apps lists CUDA + # contexts only, so a graphics consumer keeps the GPU busy invisibly. An A/B is + # if anything MORE sensitive to it than a comparison, because a drifting + # background load can look exactly like an arm effect. + local util + util=$(cut -d, -f5 "$out/$leg.gpu-before.csv" | tr -dc '0-9') + if test -n "$util" && test "$util" -gt "${GPU_IDLE_UTIL_MAX:-2}"; then + echo "GPU is not idle before $leg: utilization ${util}% exceeds" \ + "${GPU_IDLE_UTIL_MAX:-2}%" >&2 + return 1 + fi + printf '%s\n' "$arm_env" >"$out/$leg.env" + # shellcheck disable=SC2086 # arm_env is a deliberate list of NAME=VALUE words + env LD_LIBRARY_PATH="$ld_path" VT_RELEASE_HOST_WEIGHTS=1 \ + VT_DIRECT_DEVICE_LOAD=1 $arm_env \ + "$cpp" "${cpp_args[@]}" --output-token-ids "$out/$leg.tokens.json" \ + >"$out/$leg.log" 2>&1 + gpu_snapshot "$out/$leg.gpu-after" +} + +gpu_snapshot "$out/series-before" +for rep in $(seq 1 "$reps"); do + # Interleaved, and the order FLIPS on even repetitions, so neither arm is + # systematically the one that runs on a warmer GPU. + if test $((rep % 2)) -eq 1; then + run_leg "$a_name" "$a_env" "$rep" + run_leg "$b_name" "$b_env" "$rep" + else + run_leg "$b_name" "$b_env" "$rep" + run_leg "$a_name" "$a_env" "$rep" + fi +done +gpu_snapshot "$out/series-after" +echo "A/B series complete: $out" diff --git a/tools/bench/run_qwen35_4b_compare.sh b/tools/bench/run_qwen35_4b_compare.sh index 4775df3f8..9e88e89e3 100755 --- a/tools/bench/run_qwen35_4b_compare.sh +++ b/tools/bench/run_qwen35_4b_compare.sh @@ -42,11 +42,32 @@ fi ninja=$(sed -n 's/^CMAKE_MAKE_PROGRAM:[^=]*=//p' "$cmake_cache") nvcc=$(sed -n 's/^CMAKE_CUDA_COMPILER:[^=]*=//p' "$cmake_cache") host_cxx=$(sed -n 's/^CMAKE_CUDA_HOST_COMPILER:[^=]*=//p' "$cmake_cache") -cudart=$(sed -n 's/^CUDA_CUDART:[^=]*=//p' "$cmake_cache") -curand=$(sed -n 's/^CUDA_curand_LIBRARY:[^=]*=//p' "$cmake_cache") test -x "$ninja" test -x "$nvcc" test -x "$host_cxx" +libstdcpp=$($host_cxx -print-file-name=libstdc++.so.6) +test -f "$libstdcpp" + +# A venv may carry its OWN CUDA toolkit rather than borrowing the Nix one. The +# pinned oracle does: installing torch from the CUDA-13 wheels brings nvcc, +# cudart, nvrtc and the headers into site-packages/nvidia. For that venv the +# symlink farm below is not merely redundant, it is WRONG — it would put a 12.9 +# toolkit ahead of the 13.x the venv's extensions were compiled against, and +# FlashInfer's JIT would then compile against mismatched headers. VLLM_CUDA_HOME +# selects the venv's own toolkit and skips the farm entirely. +if test -n "${VLLM_CUDA_HOME:-}"; then + cuda_combined=$VLLM_CUDA_HOME + test -x "$cuda_combined/bin/nvcc" + test -d "$cuda_combined/include" + test -d "$cuda_combined/lib" + 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 + vllm_nix_ldflags="-L$cuda_combined/lib -L/run/opengl-driver/lib" +else +cudart=$(sed -n 's/^CUDA_CUDART:[^=]*=//p' "$cmake_cache") +curand=$(sed -n 's/^CUDA_curand_LIBRARY:[^=]*=//p' "$cmake_cache") test -f "$cudart" test -f "$curand" cuda_home=$(dirname "$(dirname "$nvcc")") @@ -79,10 +100,12 @@ ln -sfn /run/opengl-driver/lib/libcuda.so \ if test -d "$cuda_home/nvvm"; then ln -sfn "$cuda_home/nvvm" "$cuda_combined/nvvm" fi -libstdcpp=$($host_cxx -print-file-name=libstdc++.so.6) -test -f "$libstdcpp" vllm_path=$(dirname "$ninja"):$(dirname "$nvcc"):$(dirname "$host_cxx"):$PATH vllm_ld_library_path=$(dirname "$libstdcpp"):$(dirname "$cudart"):$(dirname "$curand"):/run/opengl-driver/lib +vllm_cpath="$cudart_home/include:$curand_include" +vllm_library_path="$cudart_home/lib:$(dirname "$curand")" +vllm_nix_ldflags="-L$cudart_home/lib -L$(dirname "$curand") -L/run/opengl-driver/lib" +fi cpp_args=( --model "$model" --dataset-path "$dataset" --num-prompts 128 @@ -101,9 +124,9 @@ vllm_env=( LD_LIBRARY_PATH="$vllm_ld_library_path" CUDA_HOME="$cuda_combined" CUDA_PATH="$cuda_combined" - CPATH="$cudart_home/include:$curand_include" - LIBRARY_PATH="$cudart_home/lib:$(dirname "$curand")" - NIX_LDFLAGS="-L$cudart_home/lib -L$(dirname "$curand") -L/run/opengl-driver/lib" + CPATH="$vllm_cpath" + LIBRARY_PATH="$vllm_library_path" + NIX_LDFLAGS="$vllm_nix_ldflags" HF_HOME="$root/.hf-cache" XDG_CACHE_HOME="$root/.vllm-cache" TORCHINDUCTOR_CACHE_DIR="$root/.torchinductor-cache" @@ -140,6 +163,24 @@ prepare_leg() { echo "GPU is not compute-idle before $name" >&2 return 1 fi + # The compute-apps check above is NOT sufficient, and a whole series has + # already been measured through the hole: `--query-compute-apps` lists CUDA + # contexts only, so a GRAPHICS consumer (compositor, browser, X client) is + # invisible to it. The 2026-07-25 series ran every one of its nine legs against + # a GPU sitting at 11-13% utilization with 611 MiB of extra VRAM resident, and + # passed this gate; re-measuring on a genuinely idle GPU moved BOTH arms by + # ~14%. So also require the device to be actually idle. The ratio survived that + # contention because it hit both arms, but no absolute number from a contended + # series may be published, per .agents/benchmark-protocol.md. + local util + util=$(cut -d, -f6 "$out/$name.gpu-before.csv" | tr -dc '0-9') + if test -n "$util" && test "$util" -gt "${GPU_IDLE_UTIL_MAX:-2}"; then + echo "GPU is not idle before $name: utilization ${util}% exceeds" \ + "${GPU_IDLE_UTIL_MAX:-2}% (a graphics consumer does not appear in" \ + "--query-compute-apps; stop it, or set GPU_IDLE_UTIL_MAX to accept it" \ + "and record the contention with the result)" >&2 + return 1 + fi awk '/^MemAvailable:/{print}' /proc/meminfo >"$out/$name.mem-before.txt" printf '%s\n' "$EPOCHREALTIME" >"$out/$name.started" } diff --git a/tools/bench/summarize_qwen35_4b_compare.py b/tools/bench/summarize_qwen35_4b_compare.py index 59a52bf69..0bc4cbff4 100755 --- a/tools/bench/summarize_qwen35_4b_compare.py +++ b/tools/bench/summarize_qwen35_4b_compare.py @@ -151,9 +151,13 @@ def token_comparisons( ] for arm in ("cpp-on", "cpp-off", "vllm") } + # Both roots are produced by run_qwen35_4b_compare.sh and use the SAME leg + # names. This side used to read a `perf-` prefix the harness has never + # written, so every historical comparison raised FileNotFoundError and the + # tool could not summarize any series against a previous one. historical = { arm: [ - load_tokens(historical_root / f"perf-{arm}-r{rep}.tokens.json") + load_tokens(historical_root / f"performance-{arm}-r{rep}.tokens.json") for rep in range(1, 4) ] for arm in ("cpp-on", "cpp-off", "vllm") @@ -206,14 +210,20 @@ def main() -> int: "vllm": memory(args.root, "vllm"), }, "tokens": token_comparisons(args.root, args.historical_root), + # The historical block accepts EITHER schema. An aggregate produced by an + # older revision of this tool carried the vLLM arm under + # `vllm_production` plus a separate `vllm_client`; the aggregate this + # tool writes today carries one `vllm` key and no client split. Insisting + # on the old names made the tool unable to read its OWN output, so a + # series could never be summarized against the previous one. "historical": { "commit": old_aggregate["commit"], "cpp_on": old_aggregate["cpp_on"], "cpp_off": old_aggregate["cpp_off"], - "vllm": old_aggregate["vllm_production"], + "vllm": old_aggregate.get("vllm_production", old_aggregate.get("vllm")), "vllm_client": { key: value - for key, value in old_aggregate["vllm_client"].items() + for key, value in old_aggregate.get("vllm_client", {}).items() if key != "per_request" }, "memory": old_aggregate["memory"],