From 1f3b71487ad9e9a6db8c47fd4cae8bcb968814f9 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sat, 8 Aug 2026 19:46:07 +0000 Subject: [PATCH] fix(vt,hooks): async readback is a Backend capability + unbreak the pre-push sandbox (#159) Two guards on main were RED and between them blocked every open external contributor PR (#127, #154, #155) and every push. Both premises were verified in the tree before changing anything. 1) check-device-leakage: src/vllm/v1/worker/gpu/runner.cpp named vt::DeviceType::kCUDA in the device-agnostic shared layer (DSR bucket 'kcuda' 1 > baseline 0). It came in with the QueueSupportsAsyncInputCombine rescope during the PR #140 fix round - ours, not the contributors'; richiejp reported it in #127's honest gaps. Fixed the way the guard's own message prescribes, mirroring the SupportsAuxStream precedent: ask the backend, not the device. New vt::Backend::SupportsAsyncSampledTokenReadback() (base false) answers whether the host may validly read the sampled token id back between steps; CPU overrides true (host and device memory are one allocation) and CUDA overrides true (the id is device-mirrored). The runner asks vt::TryGetBackend(queue.device.type), whose nullptr for a device absent from the build also subsumes the old #ifdef VLLM_CPP_CUDA guard. SEMANTICS UNCHANGED: CPU async-ON, CUDA async-ON, discrete non-CUDA (ROCm gfx1201) async-OFF - the "!"-token hazard stays closed. 2) .githooks/pre-push ran check-policy.py inside a PARTIAL export (README.md docs scripts .agents), but policy_contract.py:428 asserts AGENTS.md is a non-symlink regular file and resolves its Markdown links against that sandbox. AGENTS.md and its .env.example link were both missing, so the hook failed closed on content that is fine in the real tree - every push refused. EXPORT_PATHS is now a superset of what the CHECKERS read. Gates: check-device-leakage RED->GREEN (kcuda=0, DSR 32 == baseline 32); all four hook checkers OK in the reproduced sandbox; test_async_llm 8/8-347, test_engine_core 6/6-44, test_llm_engine 11/11-204 (CPU still resolves async-ON); clean -Werror CPU build; full 11-gate record battery green. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude (Opus 5) via Claude Code --- .agents/NOW.md | 2 +- .agents/state.md | 32 +++++++++++++++++++++++++++++- .githooks/pre-push | 8 +++++++- docs/BENCHMARKS.md | 1 + docs/STATUS.md | 5 +++-- include/vt/backend.h | 19 ++++++++++++++++++ src/vllm/v1/worker/gpu/runner.cpp | 33 +++++++++++-------------------- src/vt/cpu/cpu_backend.cpp | 6 ++++++ src/vt/cuda/cuda_backend.cu | 3 +++ 9 files changed, 83 insertions(+), 26 deletions(-) diff --git a/.agents/NOW.md b/.agents/NOW.md index bf7ca04de..b23e75c6a 100644 --- a/.agents/NOW.md +++ b/.agents/NOW.md @@ -28,7 +28,7 @@ Working head: `row/backend-rocm-w0` (#41). Prior: benchmark checkpoint | Supported-models list | **LANDED**: FEATURES arch table CI-bound (33 archs) | — | | `/v1/videos` OpenAI shape | **MERGED** (#71): Sora `model`/`size`/`seconds` + `GET /{id}/content` | `row/SERVE-VIDEOS-REFS` PR open: reference conditioning | | `BACKEND-ROCM` | **(b) fix in; #140 gfx1201 hipBLAS + Gemma-4 MoE landed (contributor, authorship-preserved); W0 green 4 archs** | compile + M2 ([spec](specs/rocm-unified-memory-b.md)) | -| TP spike #287 (PR #143) | **TP-W1 LANDED**: rank-group table + `LoadedModel` TP handle (6/6) | dispatch TP-W2 (linears + loader) | +| TP spike #287 (PR #143) | **TP-W1 LANDED**: rank-group table + TP handle (6/6); DSR leak FIXED (unblocks #127/#154/#155) | TP-W2 (linears + loader) | | Release | SPIKE; 30/30 | #129 | | Surface coverage (`ARCH-ONE-SURFACE`) | ROW 8 + #139 IN; **ROW 6 IN REVIEW (#137): embeddings LIVE — `LlamaModel` arch, PoolingRunner in the step, `vllm_embed` v15, `/v1/embeddings`, fold gate 4/4-231, 9 kills** | Merge #137; real-ckpt oracle cosine residual | diff --git a/.agents/state.md b/.agents/state.md index 0e50fef63..a7ae2d214 100644 --- a/.agents/state.md +++ b/.agents/state.md @@ -43289,4 +43289,34 @@ inertness and thread_local isolation. No regression: existing `test_tp_forward` 2/2·60 and `test_communicator` 8/8·50 unchanged (the LoadedModel field is additive/byte-neutral). Header-only, no CMake wiring beyond the test. Next: TP-W2 (row/input-dim shard + rank-0 bias + per-rank Hq/Hkv + QKV kv-replication + vocab -embed/lm_head + logits all-gather). \ No newline at end of file +embed/lm_head + logits all-gather). + +## 2026-08-08 — FIX: device-leakage regression on main (async readback becomes a Backend capability) — unblocks 3 contributor PRs + + +`check-device-leakage` was RED on main (`1a021b1b`, and back through `c05cee1d`): +`src/vllm/v1/worker/gpu/runner.cpp:107` named `vt::DeviceType::kCUDA` in the +device-agnostic shared layer (DSR bucket `kcuda` 1 > baseline 0). It came in with +the `QueueSupportsAsyncInputCombine` rescope during the PR #140 fix round — our +regression, and it was failing CI on all three open contributor PRs (#127, #154, +#155), exactly as richiejp reported in #127's "Honest gaps". Premise verified in +the tree before acting (not taken on report). + +FIX mirrors the `SupportsAuxStream` precedent the guard's own message prescribes: +move the question onto the backend. New +`vt::Backend::SupportsAsyncSampledTokenReadback()` (backend.h, base **false**), +overridden true by CPU (host and device memory are one allocation) and CUDA (the +sampled id is device-mirrored). The runner now asks +`vt::TryGetBackend(queue.device.type)` — nullptr (device not built) yields false, +which also subsumes the old `#ifdef VLLM_CPP_CUDA` guard. SEMANTICS UNCHANGED: +CPU async-ON, CUDA async-ON, discrete non-CUDA (ROCm gfx1201) async-OFF — the +"!"-token hazard that motivated the original rescope stays closed, and the ROCm +TODO now points at the override rather than at a device branch. + +Gates: `check-device-leakage` RED→GREEN (`kcuda=0`, DSR 32 == baseline 32); +`test_async_llm` 8/8·347, `test_engine_core` 6/6·44, `test_llm_engine` 11/11·204 +(CPU still resolves async-ON, the regression this could have caused); clean +-Werror CPU build. SEPARATE pre-existing main breakage recorded, not fixed here: +`check-doc-checkpoint` is RED on `1a021b1b` itself (that policy commit touched +70+ `.agents/` files without docs/BENCHMARKS.md + docs/FEATURES.md); it validates +HEAD, so it self-heals as compliant commits land. \ No newline at end of file diff --git a/.githooks/pre-push b/.githooks/pre-push index 2eecf8802..472bd15f8 100755 --- a/.githooks/pre-push +++ b/.githooks/pre-push @@ -25,7 +25,13 @@ ZERO='0000000000000000000000000000000000000000' # The checkers and the paths they read. Keep both lists in sync with # scripts/agent-preflight.sh's CHECKERS entries for the same two gates. CHECKERS=(check-policy.py check-prompt-contract.py check-readme-structure.py check-public-doc-tables.py) -EXPORT_PATHS=(README.md docs scripts .agents) +# AGENTS.md + .env.example are exported too. check-policy.py asserts AGENTS.md is +# a non-symlink regular file (policy_contract.py:428) and resolves its Markdown +# links against this sandbox, so omitting either made EVERY push fail here while +# passing in the real tree ("AGENTS.md must be a non-symlink regular file", then +# "broken Markdown link: .env.example"). Keep this list a SUPERSET of what the +# CHECKERS read, or the hook fails closed on content that is actually fine. +EXPORT_PATHS=(README.md AGENTS.md .env.example docs scripts .agents) shas=() while read -r _local_ref local_sha _remote_ref _remote_sha; do diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index 145c78a44..f6f568047 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -7,6 +7,7 @@ | **Binary release matrix (spiked)** | `ENG-RELEASE-BINARIES`: primary host-ABI fat-CUDA + adaptive-CPU static-core bundles; optional per-SM diagnostics; experimental literal-static musl CPU | **PENDING:** pins 10-SM fat CUDA, adaptive no-AVX2 CPU, W1-W13/W10-W12 policy, public pending states; 30 tests GREEN. No archive, staged smoke, runtime, correctness, or performance evidence | n/a | | **Developer agent entry point (implemented)** | `DOCS-AGENT-PROTOCOL-ENTRYPOINT`: public contribution guide + synchronized, mutation-gated pre-claim intake rule | Rebased documentation/protocol only; benchmark void | n/a | | **ARCH audit: ABI is text-only** | 4 capabilities (H3 video, Laguna, Kimi-Linear, DeepSeek-V4) reachable only from `examples/`, none registry-backed. No gate asks whether a CONSUMER can reach a capability. Documentation only | +| **DSR fix: async readback capability (2026-08-08)** | **No number owed**: behavior-neutral (CPU/CUDA async-ON, discrete non-CUDA async-OFF, unchanged); moves a `kCUDA` check onto `Backend`, unblocking red CI on #127/#154/#155 | | **`ROAD-V1-MEM` M1+M2 (2026-08-08)** | KV auto-sizing CPU brick: `--kv-cache-memory` sizes the pool from a byte budget via the group-aware `KVBytesPerBlock` divisor (ABI v16, CPU-gated). M3 profile run dgx-gated | | **Record repair 2026-08-07** | `main` was red on `check-agent-record` + `check-env-doc`, blocking every PR. Dangling `kda-chunk-aot/` link and two undocumented env vars. No behaviour change | | **vLLM** | Qwen3.6-27B NVFP4, GB10 | ahead 4.5% at c1, **tie** at c2 to c32 | identical | diff --git a/docs/STATUS.md b/docs/STATUS.md index 6ccd1857d..e7e0dd6ef 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -431,10 +431,11 @@ tp=1). The **real TP-2 GPU run + NCCL build-verify remain HW-blocked** (no parallelism mode is enumerated and ranked in [parallelism-modes](../.agents/specs/parallelism-modes.md), noting vLLM's "sequence parallel" is a TP compilation pass, not an axis). +**Shared layer names no device**: async readback is a `Backend` capability. + **Tensor parallelism is scoped end-to-end at the pin** (#287, [tensor-parallelism-spike](../.agents/specs/tensor-parallelism-spike.md)): -~40% of the TP surface already landed/reusable; the landed TP wiring -dead-ends at the layer boundary (no loader shards yet); **TP-W1 LANDED** (group table); W2..W4+W7 +~40% of the TP surface landed/reusable; **TP-W1 LANDED** (group table); W2..W4+W7 (the engine-level TP2-on-CPU token-exact gate) are CPU-completable NOW; only NCCL + gate-model perf wait on hardware. The DSpark speculator (DFlash-derived block drafter for our Qwen3 + Gemma4 families) is grounded diff --git a/include/vt/backend.h b/include/vt/backend.h index 21d884a82..cf4ff60ee 100644 --- a/include/vt/backend.h +++ b/include/vt/backend.h @@ -121,6 +121,25 @@ class Backend { // the DSR scan) so the model file stops naming a device at the aux-stream gate. virtual bool SupportsAuxStream() const { return false; } + // Whether the host may validly read the SAMPLED TOKEN ID back between steps, + // which is what the depth-2 async input-combine path requires + // (gpu/runner.cpp: QueueSupportsAsyncInputCombine). Base false — a DISCRETE + // non-CUDA GPU (e.g. ROCm gfx1201) is the hazard: the non-CUDA leg of + // sample_tokens_async Synchronizes and then host-dereferences `dev_ids`, a + // device Alloc that is garbage off-device (the "!"-token corruption on the lab + // R9700, 2026-08-07), so those queues MUST stay synchronous. Overridden true + // by CPU (host and device memory are the same allocation, so the read is + // always valid) and by CUDA (the sampled id is device-mirrored, + // async_device_mirror()). This is the capability the runner's + // `device == kCUDA` gate actually asked; it lives on Backend (src/vt, off the + // DSR scan) so the device-agnostic shared layer stops naming a device — the + // same move SupportsAuxStream made for the aux-stream gate. + // TODO(rocm): an INTEGRATED non-CUDA GPU reports UnifiedMemory()==true (see + // row/ROCM-UNIFIED-MEMORY-B), where the alias IS valid; such a backend may + // override this true once a HIP sampled-token mirror or a D2H copy of dev_ids + // lands. + virtual bool SupportsAsyncSampledTokenReadback() const { return false; } + // Optional graph/command capture (CUDA Graphs / Metal ICB / Vulkan CB). virtual bool SupportsGraphCapture() const { return false; } virtual void BeginCapture(Queue& q); diff --git a/src/vllm/v1/worker/gpu/runner.cpp b/src/vllm/v1/worker/gpu/runner.cpp index 1a2524ff6..4bf531bdb 100644 --- a/src/vllm/v1/worker/gpu/runner.cpp +++ b/src/vllm/v1/worker/gpu/runner.cpp @@ -86,28 +86,19 @@ static bool AsyncRunnerEnvDefault() { } // Async input-combine reads the sampled token id back on the host between -// steps. Where that read is valid the default-ON async path stays on: -// - kCPU: host and device memory are the same allocation, so the read is -// always valid. This path was correct, default-ON, and contract-tested -// before the ROCm work; it MUST stay true (else the CPU backend silently -// regresses to synchronous depth-1). -// - kCUDA: the sampled id is device-mirrored (async_device_mirror()). -// A DISCRETE non-CUDA GPU (e.g. ROCm gfx1201) is the real hazard: the non-CUDA -// leg of sample_tokens_async Synchronizes and then host-dereferences dev_ids, -// which is a device Alloc — valid on CPU/UMA, garbage off-device. That is the -// root cause of the "!" tokens on the lab R9700 (2026-08-07), not an embed -// race. Keep those queues synchronous until a HIP sampled-token mirror or a D2H -// copy of dev_ids lands. -// TODO(rocm): an INTEGRATED non-CUDA GPU reports UnifiedMemory()==true (see -// row/ROCM-UNIFIED-MEMORY-B), where the alias is valid and async would be safe; -// route it through the backend UnifiedMemory() seam once reachable here. +// steps. Whether that read is VALID is a backend CAPABILITY, not a device name: +// ask the backend (vt::Backend::SupportsAsyncSampledTokenReadback, backend.h), +// which answers true for CPU (host and device memory are one allocation) and +// CUDA (the sampled id is device-mirrored, async_device_mirror()), and false for +// a DISCRETE non-CUDA GPU (e.g. ROCm gfx1201) whose sample_tokens_async leg +// host-dereferences a device Alloc — the root cause of the "!" tokens on the lab +// R9700 (2026-08-07). An absent backend (device not built into this binary) +// yields nullptr and therefore false, which also subsumes the old +// #ifdef VLLM_CPP_CUDA guard. Keeping the question on the backend is what stops +// this device-agnostic shared layer from naming a device (check-device-leakage). static bool QueueSupportsAsyncInputCombine(const vt::Queue& queue) { - if (queue.device.type == vt::DeviceType::kCPU) return true; -#ifdef VLLM_CPP_CUDA - if (queue.device.type == vt::DeviceType::kCUDA) return true; -#endif - (void)queue; - return false; + const vt::Backend* backend = vt::TryGetBackend(queue.device.type); + return backend != nullptr && backend->SupportsAsyncSampledTokenReadback(); } // GDN step-geometry diagnostic (default OFF). When VT_GDN_DIAG_STEP_LOG=1, each diff --git a/src/vt/cpu/cpu_backend.cpp b/src/vt/cpu/cpu_backend.cpp index 96274a1d9..4a1591b2d 100644 --- a/src/vt/cpu/cpu_backend.cpp +++ b/src/vt/cpu/cpu_backend.cpp @@ -9,7 +9,13 @@ namespace vt::cpu { namespace { class CpuBackend final : public Backend { + // Host and device memory are the SAME allocation here, so reading the sampled + // token id back between steps is always valid: the depth-2 async input-combine + // path stays default-ON for CPU (it was correct and contract-tested before the + // ROCm work; regressing it would silently drop the CPU backend to depth-1). public: + bool SupportsAsyncSampledTokenReadback() const override { return true; } + void* Alloc(size_t bytes) override { VT_CHECK(bytes <= SIZE_MAX - 63, "cpu alloc size overflow"); void* p = std::aligned_alloc(64, ((bytes + 63) / 64) * 64); // 64B-aligned, padded size diff --git a/src/vt/cuda/cuda_backend.cu b/src/vt/cuda/cuda_backend.cu index aeb8fd342..06a0f7b78 100644 --- a/src/vt/cuda/cuda_backend.cu +++ b/src/vt/cuda/cuda_backend.cu @@ -188,6 +188,9 @@ class CudaBackend final : public Backend { bool SupportsGraphCapture() const override { return true; } // S7: CUDA has a secondary stream for the MoE shared-expert overlap fork. bool SupportsAuxStream() const override { return true; } + // The sampled token id is device-mirrored (async_device_mirror()), so the + // between-steps host readback the depth-2 async input-combine needs is valid. + bool SupportsAsyncSampledTokenReadback() const override { return true; } void BeginCapture(Queue& q) override { Check(cudaStreamBeginCapture(AsStream(q), cudaStreamCaptureModeThreadLocal), "cudaStreamBeginCapture");