diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md index d3f285af6..0728bc454 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -22678,3 +22678,90 @@ consecutive reps** — a ~2x spread that averages to a plausible-looking and entirely fictional number. The figures above use 8 prompts with a discarded warmup, where both sides hold to ~±0.3 ms. A two-request rate harness is not a measurement of this axis; it is a coin flip with a mean. + +## ENG-CUDAGRAPH-BREAK W1 — the measurement that was NOT taken, and the one that was (2026-08-18, #1192) + +**No throughput number was taken, and none is owed.** This is a coverage and +correctness row. A speed claim from the break-point seam is admissible only after +naming a path that is BOTH currently eager AND currently host-bound, and stating +how the host-bound part is measured. Our prefill is neither: GB10 measured idle +between launches at 3.8% with GPU-busy above 96%, and the 27B prefill gap at +92.5% non-GEMM glue GPU work. Decode is already captured and already banked its +launch-overhead win. The refutation is dated and hardware-specific rather than +permanent; the burden is on a later claimant to name the path. + +**What W1 did measure is a CAPABILITY, not a rate.** W0 deliberately left open +whether CUDA permits `cudaStreamEndCapture` followed by `cudaStreamBeginCapture` +on the same stream mid-forward with eager work between them, on our stream +configuration. SGLang does exactly this on a production path at the pinned +revision, which is strong evidence and was not our measurement. + +Measured on `orin:gpu0` through an `rc` lease, driver `12060`, under +`cudaStreamCaptureModeThreadLocal` — the mode `src/vt/cuda/cuda_backend.cu:204` +uses. The probe runs the seam's own shape rather than a toy: segment, +host-dependent eager break on the same stream, RE-BEGIN, segment, bare zero-work +re-begin, segment. Then three replays with fresh inputs. Result: every re-begin +legal, 0 mismatches on all three replays. The bare re-begin +(`breakable_cuda_graph.py:370-374`) is legal too. + +**The first two probe runs REFUSED, and both refusals were the probe's.** This is +the entry worth reading before anyone re-runs this lever. Binding the CUDA driver +API through `dlsym` on the BARE symbol name gets `libcuda`'s LEGACY v1 entry +points, which are not capture-aware: `cuMemcpyDtoDAsync` (v1) returned +`CUDA_ERROR_STREAM_CAPTURE_UNSUPPORTED` inside a capture that is in fact legal, +and `cuStreamBeginCapture` (v1) takes NO capture-mode argument, so the probe +believed it was exercising the thread-local mode while exercising the global one. +Preferring `_v3` blindly then bound `cuCtxCreate_v3`, which takes two extra +parameters, and context creation failed `CUDA_ERROR_INVALID_DEVICE`. Both +readings presented as a verdict about the DESIGN. Bound by exact versioned name, +the criterion holds. An instrument that can fail toward a code verdict has to +assert its own precondition first. + +**A second run on GB10 (`dgx:gpu0`) compiled and did not execute**, for a reason +that is also not CUDA's: `nvcc` produced the binary under `/workspace`, which on +that host is a CIFS mount storing `file_mode=0664`, so the run exited 126 +`Permission denied`. `BUILD_STATUS=0` on `NVIDIA GB10` is recorded; the execution +belongs on container-local disk. + +**The correctness number that IS recorded** is not a rate either: the Qwen3 dense +forward run with a capture scope open produces logits BIT-IDENTICAL to the same +forward with no scope — 500 values compared, 0 differing +(`tests/vllm/models/test_qwen3_break_point.cpp`). That is what makes the stage +reversible, and it is the polarity AGENTS.md requires when a greedy path exists. + +**THE PROBE IS IN THE TREE, and here is how to run it.** A measurement that +produced two false refusals, each presenting as a verdict about the design, is +the last one anybody should have to reconstruct from prose. The first record of +it named neither an artifact nor a recipe; this one does. + +| Artifact | Path | sha256 | Lines | +|---|---|---|---| +| driver-API probe (`dlsym`, no toolkit, no headers) | `scripts/probe_cudagraph_rebegin.c` | `dbf95d69d396e7f3a41c754ccedda59ef076a1dd30960d2c738a1c26fe238297` | 246 | +| runtime-API probe (needs `nvcc`) | `scripts/probe_cudagraph_rebegin.cu` | `ce73250331af5bc3eb8b560176e60a07f77476dcff60b81a0b0d7caea6f20bed` | 167 | + +Both exercise the SEAM's shape rather than a toy — segment, host-dependent eager +break on the same stream, RE-BEGIN, segment, bare zero-work re-begin, segment, +then three replays with fresh inputs — and both print a single +`VERDICT: REBEGIN_HOLDS` or a named refusal. Recipe, on a leased device: + +```sh +# driver-API build: no CUDA toolkit needed, links libdl and binds libcuda at run time +cc -O1 -o probe scripts/probe_cudagraph_rebegin.c -ldl && ./probe; echo "PROBE_EXIT=$?" + +# runtime-API build +nvcc -O1 -arch=native -o probe scripts/probe_cudagraph_rebegin.cu || \ + nvcc -O1 -o probe scripts/probe_cudagraph_rebegin.cu +echo "BUILD_STATUS=$?"; ./probe; echo "PROBE_EXIT=$?" +``` + +Build it on CONTAINER-LOCAL disk. The `dgx:gpu0` run compiled under `/workspace`, +which on that host is a CIFS mount storing `file_mode=0664`, and exited 126 +`Permission denied` — `BUILD_STATUS=0` with nothing executed, which is a build +result wearing a measurement's clothes. + +**What is NOT recoverable, stated plainly rather than reconstructed.** The +`orin:gpu0` run's own driver script and its raw stdout were not retained. The +verdict above therefore rests on the recorded reading, and the only way to +re-derive it is to re-run the committed source under the recipe above on a leased +GPU. Nothing here was re-run for this record: the repair pass that added it had +no GPU lease. diff --git a/.agents/claims/CLAIM-ENG-CUDAGRAPH-BREAK-W1.md b/.agents/claims/CLAIM-ENG-CUDAGRAPH-BREAK-W1.md new file mode 100644 index 000000000..ba3345122 --- /dev/null +++ b/.agents/claims/CLAIM-ENG-CUDAGRAPH-BREAK-W1.md @@ -0,0 +1,5 @@ +# CLAIM-ENG-CUDAGRAPH-BREAK-W1 + +| Claim | Row IDs | Agent | Worktree / remote dir | Branch | Owned scope | State | Last update | +|---|---|---|---|---|---|---|---| +| `CLAIM-ENG-CUDAGRAPH-BREAK-W1` | `ENG-CUDAGRAPH-BREAK` (`ACTIVE`) | Claude Code (opus-5), fresh implementer | isolated worktree; `rc`-leased GPU for the exit criterion only (`orin:gpu0`), no registry, no publication | `row/ENG-CUDAGRAPH-BREAK-W1`, issue [#1192](https://github.com/mudler/vllm.cpp/issues/1192), parent [#1163](https://github.com/mudler/vllm.cpp/issues/1163) | Owns ONLY W1 of [`eng-cudagraph-break.md`](../specs/eng-cudagraph-break.md): NEW `include/vt/breakable_graph.h`, `src/vt/breakable_graph.cpp`, `tests/vt/breakable_graph` suite and its recording backend helper, `tests/vllm/models/test_qwen3_break_point.cpp`, the ONE break point at the dense attention entry in `src/vllm/model_executor/models/qwen3.cpp`, and the two CMake registrations. EXCLUDES every driver migration (W2 through W5), the `pure_decode` predicate at `src/vllm/v1/worker/gpu/runner.cpp:1341` (W6), the auxiliary-stream auto-join (D10, W4/W5), and graph-executable dedup, which is [#1162](https://github.com/mudler/vllm.cpp/issues/1162) on a parallel branch and whose files this claim never touches | `ACTIVE` | 2026-08-18 — exit criterion CONFIRMED on `orin:gpu0` (driver 12060, 3 replays, 0 mismatches); seam + 24-case ported suite (24 cases / 163 assertions, re-derived by `ninja test_breakable_graph && ./build/tests/test_breakable_graph`; the recorded 14-case figure never re-derived) + G2 reachability gate green (2/2 cases, 516 assertions); W2 through W6 unclaimed | diff --git a/.agents/engine-matrix.md b/.agents/engine-matrix.md index 9b8cd9400..81074c7dd 100644 --- a/.agents/engine-matrix.md +++ b/.agents/engine-matrix.md @@ -38,7 +38,7 @@ forensics: roadmap_v1.md and the parity ledger. | Area | Rows | `ANCHOR-BACKFILL` | `PARTIAL` | `SPIKE` | `READY` | `ACTIVE` | `GATING` | `DONE` | `INVENTORIED` | |---|---:|---:|---:|---:|---:|---:|---:|---:|---:| -| Engine and scheduling | 30 | 6 | 3 | 1 | 4 | 6 | 2 | 2 | 6 | +| Engine and scheduling | 30 | 6 | 3 | 1 | 3 | 7 | 2 | 2 | 6 | | KV cache and memory | 24 | 7 | 3 | 2 | 3 | 3 | 2 | 0 | 4 | | Parallelism | 6 | 0 | 0 | 0 | 1 | 0 | 0 | 0 | 5 | | Sampling and generation | 15 | 4 | 2 | 0 | 0 | 4 | 0 | 1 | 4 | @@ -48,7 +48,7 @@ forensics: roadmap_v1.md and the parity ledger. | LoRA and adapters | 2 | 0 | 0 | 0 | 0 | 1 | 0 | 0 | 1 | | Long context and attention | 10 | 5 | 0 | 0 | 1 | 0 | 1 | 0 | 3 | | Loading, tokenizer, config | 10 | 3 | 3 | 0 | 0 | 1 | 1 | 1 | 1 | -| **Total** | **162** | **35** | **17** | **5** | **13** | **33** | **8** | **11** | **39** | +| **Total** | **162** | **35** | **17** | **5** | **12** | **34** | **8** | **11** | **39** | ## Engine core and scheduling @@ -61,7 +61,7 @@ forensics: roadmap_v1.md and the parity ledger. | `ENG-PREEMPT-RECOMPUTE` | FCFS tail preemption with recompute | T0 | `vllm/v1/core/sched/scheduler.py:1142`; `tests/v1/core/test_scheduler.py:930` | `src/vllm/v1/core/sched/scheduler.cpp:102,157`; `src/vllm/v1/core/sched/request_queue.cpp:36` | `tests/vllm/v1/test_scheduler.cpp:247,295`; `tests/vllm/v1/test_request_queue.cpp:91` | `planned: specs/preemption.md` | `ANCHOR-BACKFILL` | - | | `ENG-CUDAGRAPH` | Decode graph capture/replay modes (host-cluster cleanup: capture-size set derived from `max_num_seqs` mirroring vLLM `_set_cudagraph_sizes`; 2026-07-18 graph-baked-scratch use-after-free fix — the 35B c2+ online-serving IMA blocker) | T0 | `vllm/config/compilation.py:53,1319,683-684,1438-1444`; `vllm/config/vllm.py:1667-1770`; `vllm/v1/worker/gpu/cudagraph_utils.py:116`; `tests/compile/test_config.py:122,229` | `src/vt/cuda/cuda_backend.cu:76,97,105`; `include/vllm/model_executor/models/decode_graph_sizes.h`; `src/vllm/model_executor/models/qwen3_5.cpp:3754,3952`; `src/vllm/v1/worker/gpu/runner.cpp:577,597`; graph-safe scratch (retire-on-grow so graph-baked scratch pointers stay valid) `src/vt/cuda/graph_safe_scratch.h`, `src/vt/cuda/cuda_moe_marlin.cu:75`, `src/vt/cuda/cuda_matmul_nvfp4.cu:766`, `src/vt/cuda/cuda_matmul_nvfp4_cutlass.cu:105`, `src/vt/cuda/cuda_matmul_fp8_cutlass.cu:95` | `tests/vt/test_cuda_backend.cpp:98`; `tests/vllm/models/test_decode_graph_sizes.cpp`; `tests/vt/test_graph_safe_scratch.cpp`; explicit 35B gate `tests/parity/test_qwen36_paged_engine.cpp:140` | [blocktable-host-cluster-cleanup.md](specs/blocktable-host-cluster-cleanup.md); [decode-graph-scratch-uaf-2026-07-18.md](specs/decode-graph-scratch-uaf-2026-07-18.md) | `PARTIAL` | **PREFILL capture REFUTED as a lever (2026-08-17, [#1161](https://github.com/mudler/vllm.cpp/issues/1161)).** vLLM's v1 default already captures prefill piecewise (`vllm/config/compilation.py:60-63,615,630` @ `555967922`) and it is in our denominator; SGLang reached the same coverage without `torch.compile` via BCG (`SGLANG-BCG` in [sglang-matrix.md](sglang-matrix.md)). Neither helps us: GB10 2026-07-09 measured prefill GPU-idle-between-launches at **3.8%** with GPU-busy >96% on both arms, and the 27B prefill gap at **92.5% non-GEMM glue GPU work** with the dominant GEMM at +0.17% and attention AHEAD. There are no launch bubbles in our prefill to collapse. Row stays `PARTIAL`; the real residuals are exec dedup ([#1162](https://github.com/mudler/vllm.cpp/issues/1162)) and the break-point seam ([#1163](https://github.com/mudler/vllm.cpp/issues/1163)). Spec [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | | `ENG-CUDAGRAPH-DEDUP` | Graph-executable dedup: hash each captured graph's topology and re-point ONE `cudaGraphExec` with `cudaGraphExecUpdate` on a signature hit, instead of instantiating one exec per padded bucket per model. A memory and capture-time change, NOT a throughput change — a deduped replay launches the same nodes | T2 | vLLM has no analogue (its execs come from `torch.compile`); SGLang `python/sglang/srt/model_executor/runner_backend/cuda_graph_dedup_mixin.py:219-242,358` @ `f63458b5be` | UNPORTED — `cudaGraphExecUpdate` appears nowhere in `src/` or `include/`; today `src/vt/cuda/cuda_backend.cu:222-232` instantiates a fresh exec per capture, over the 7 (`max_num_seqs=32`) or 11 (64) buckets of `include/vllm/model_executor/models/decode_graph_sizes.h:32-41`, times NINE drivers (count corrected 2026-08-18, [#1179](https://github.com/mudler/vllm.cpp/issues/1179); `9bc4d7f44` recorded eight, missing the DFlash draft graph `src/vllm/model_executor/models/qwen3_dflash.cpp:771,870,1038,1091,1095,1106`) | owed: same-binary A/B proving a deduped replay is byte-identical to a non-deduped one; a capture-count/exec-count log line so the ratio is observable | [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `INVENTORIED` | [#1162](https://github.com/mudler/vllm.cpp/issues/1162) | -| `ENG-CUDAGRAPH-BREAK` | One shared `vt` capture seam that accepts BREAK POINTS, so a forward containing a host-dependent op is still graphed instead of falling out entirely — and so the NINE hand-rolled drivers become one (count corrected 2026-08-18, [#1179](https://github.com/mudler/vllm.cpp/issues/1179); `9bc4d7f44` recorded eight). **Coverage AND CORRECTNESS row, not a throughput row** | T1 | mirror vLLM `CUDAGraphMode.PIECEWISE` splitting at `splitting_ops` (`vllm/config/compilation.py:60-63,517,615,630` @ `555967922`); construction from SGLang BCG `python/sglang/srt/model_executor/runner_backend_utils/breakable_cuda_graph/breakable_cuda_graph.py:204-243,246-274,309-333,335-367` @ `f63458b5be` (decorator + runtime stream capture, no compiler); its unit suite `test/registered/cuda_graph/breakable/test_breakable_cuda_graph.py:30,172,230` (305 lines, 11 unit cases) is mapped case for case in the spec's `## Tests to port` | all-or-nothing today: `src/vllm/v1/worker/gpu/runner.cpp:1338-1341` routes only `pure_decode`; drivers `qwen3_5.h:275`, `qwen3_5_dense.h:391`, `qwen3_moe.h:117`, `qwen3.h:243`, `deepseek_v2.h:324`, `voxtral.h:126`, plus `deepseek_v4.cpp`, `laguna.cpp` — and the spike found the NINTH already written, `src/vllm/model_executor/models/qwen3_dflash.cpp:771,1091`. The re-derivation is measured, not asserted: `StepDevInputs` (`src/vllm/model_executor/models/qwen3_5.cpp:3894`, the persistent DEVICE input path) exists in ONE driver and `grep -c` returns 0 in `qwen3_moe.cpp`, `qwen3.cpp`, `deepseek_v2.cpp` and `voxtral.cpp`, which is why `src/vllm/model_executor/models/qwen3.cpp:961-986` DECLINES the graph outright when the async device-token mirror is live. **That decline is why this is also a CORRECTNESS row** ([#1179](https://github.com/mudler/vllm.cpp/issues/1179)): a SHIPPED model has already lost its decode graph to the duplication, on the driver's own measurement (`depth-1, graph ON PASS 78/78`; `depth-2, graph OFF PASS 82/82`; `depth-2, graph ON FAIL, slots 1-3 degenerate`), and the fix its comment names is the sibling's `StepDevInputs`. The row still makes NO throughput claim: the prefill refutation on the `ENG-CUDAGRAPH` row (3.8% host idle, >96% GPU-busy, 92.5% glue) stands unchanged | owed: reachability mutation (delete the production call site, rerun the focused gate); bit-exactness vs eager on every migrated model over MORE than one replay; the host-lifetime contract of [decode-graph-scratch-uaf-2026-07-18.md](specs/decode-graph-scratch-uaf-2026-07-18.md) enforced AT the seam; the 11 ported SGLang unit cases; the break-function OUTPUT writeback (`replay_fn`/`_copy_output` `breakable_cuda_graph.py:231-235,172-201`, spec D9) without which a replay leaves the next segment reading capture-time data; the auxiliary-stream auto-join before every segment close (`:353-361`, spec D10), live at `src/vllm/model_executor/models/qwen3_5.cpp:6254-6255,6384` and `src/vllm/model_executor/models/laguna.cpp:2572-2576,2612` | spec [eng-cudagraph-break.md](specs/eng-cudagraph-break.md) (W0 spike DONE 2026-08-18: the existing `vt` capture vocabulary `include/vt/backend.h:208-222` expresses a SEGMENTED capture with NO new virtual, because `EndCaptureGraph` stores nothing (`src/vt/cuda/cuda_backend.cu:225-232`); a break point is expressible with one `thread_local` capture pointer plus a free function, no compiler and no decorator); analysis [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `READY` | [#1163](https://github.com/mudler/vllm.cpp/issues/1163), [#1020](https://github.com/mudler/vllm.cpp/issues/1020) | +| `ENG-CUDAGRAPH-BREAK` | One shared `vt` capture seam that accepts BREAK POINTS, so a forward containing a host-dependent op is still graphed instead of falling out entirely — and so the NINE hand-rolled drivers become one (count corrected 2026-08-18, [#1179](https://github.com/mudler/vllm.cpp/issues/1179); `9bc4d7f44` recorded eight). **Coverage AND CORRECTNESS row, not a throughput row** | T1 | mirror vLLM `CUDAGraphMode.PIECEWISE` splitting at `splitting_ops` (`vllm/config/compilation.py:60-63,517,615,630` @ `555967922`); construction from SGLang BCG `python/sglang/srt/model_executor/runner_backend_utils/breakable_cuda_graph/breakable_cuda_graph.py:204-243,246-274,309-333,335-367` @ `f63458b5be` (decorator + runtime stream capture, no compiler); its unit suite `test/registered/cuda_graph/breakable/test_breakable_cuda_graph.py:30,172,230` (305 lines, 11 unit cases) is mapped case for case in the spec's `## Tests to port` | all-or-nothing today: `src/vllm/v1/worker/gpu/runner.cpp:1338-1341` routes only `pure_decode`; drivers `qwen3_5.h:275`, `qwen3_5_dense.h:391`, `qwen3_moe.h:117`, `qwen3.h:243`, `deepseek_v2.h:324`, `voxtral.h:126`, plus `deepseek_v4.cpp`, `laguna.cpp` — and the spike found the NINTH already written, `src/vllm/model_executor/models/qwen3_dflash.cpp:771,1091`. The re-derivation is measured, not asserted: `StepDevInputs` (`src/vllm/model_executor/models/qwen3_5.cpp:3894`, the persistent DEVICE input path) exists in ONE driver and `grep -c` returns 0 in `qwen3_moe.cpp`, `qwen3.cpp`, `deepseek_v2.cpp` and `voxtral.cpp`, which is why `src/vllm/model_executor/models/qwen3.cpp:961-986` DECLINES the graph outright when the async device-token mirror is live. **That decline is why this is also a CORRECTNESS row** ([#1179](https://github.com/mudler/vllm.cpp/issues/1179)): a SHIPPED model has already lost its decode graph to the duplication, on the driver's own measurement (`depth-1, graph ON PASS 78/78`; `depth-2, graph OFF PASS 82/82`; `depth-2, graph ON FAIL, slots 1-3 degenerate`), and the fix its comment names is the sibling's `StepDevInputs`. The row still makes NO throughput claim: the prefill refutation on the `ENG-CUDAGRAPH` row (3.8% host idle, >96% GPU-busy, 92.5% glue) stands unchanged | owed: bit-exactness vs eager on every migrated model over MORE than one replay, on a real GPU (W2); the host-lifetime contract of `decode-graph-scratch-uaf-2026-07-18.md` enforced AT the seam — D1's INPUT half, making the intermediates a segment reads unavailable to the `DevicePool` free list (W2); the auxiliary-stream auto-join before every segment close (`:353-361`, spec D10), live at `src/vllm/model_executor/models/qwen3_5.cpp:6254-6255,6384` and `src/vllm/model_executor/models/laguna.cpp:2572-2576,2612` (W4, W5). **Delivered by W1** ([#1192](https://github.com/mudler/vllm.cpp/issues/1192)): the reachability mutation (performed; deleting the call site reds `tests/vllm/models/test_qwen3_break_point.cpp` and leaves the unit suite green); the ported SGLang unit cases with their arithmetic chains and post-replay assertions; and the break-function OUTPUT writeback (`replay_fn`/`_copy_output` `breakable_cuda_graph.py:231-235,172-201`, spec D9), whose destination is a `vt::BreakSlot` the seam owns rather than a caller reference it cannot outlive | spec [eng-cudagraph-break.md](specs/eng-cudagraph-break.md) (W0 spike DONE 2026-08-18: the existing `vt` capture vocabulary `include/vt/backend.h:208-222` expresses a SEGMENTED capture with NO new virtual, because `EndCaptureGraph` stores nothing (`src/vt/cuda/cuda_backend.cu:225-232`); a break point is expressible with one `thread_local` capture pointer plus a free function, no compiler and no decorator); **W1 DONE 2026-08-18 ([#1192](https://github.com/mudler/vllm.cpp/issues/1192)): the seam LANDS** — `vt::BreakableGraph`, `vt::GraphCaptureScope` and `vt::GraphBreak` (`include/vt/breakable_graph.h`, `src/vt/breakable_graph.cpp`), the SGLang unit suite ported case for case (`tests/vt/test_breakable_graph.cpp`, 24 cases / 163 assertions, re-derived 2026-08-18 by `ninja test_breakable_graph && ./build/tests/test_breakable_graph`; the recorded 14/81 never re-derived at any head of this branch), and ONE break point registered at the DENSE ATTENTION ENTRY of `Qwen3ForCausalLM` (`src/vllm/model_executor/models/qwen3.cpp`, inside `RunLayer`). **The exit criterion W0 deliberately left open is ANSWERED on a leased GPU:** `cudaStreamEndCapture` then `cudaStreamBeginCapture` on the SAME stream mid-forward with EAGER work between is LEGAL under `cudaStreamCaptureModeThreadLocal` (`src/vt/cuda/cuda_backend.cu:204-206`) — `orin:gpu0` via an `rc` lease, driver 12060, 3 replays with fresh inputs, 0 mismatches, bare zero-work re-begin legal too. G2 reachability is `tests/vllm/models/test_qwen3_break_point.cpp`, which drives the production `Qwen3DenseModel::Forward` with a scope open and counts `num_hidden_layers + 1` segments (mutation: delete the call site ⇒ 1 segment ⇒ RED), and holds G4 in the same case at 500 logits / 0 differing bit for bit. STAGED SLICE, named: the scope and the container are not yet ENTERED from a production step — no driver opens a scope until W2 migrates `Qwen3DenseDecodeGraph` — and the spec's `## Owed` lists it with W2 as owner, alongside the D10 auxiliary-stream auto-join (W4/W5), G5's ROCm/Tenstorrent arms (W3) and G1 on a real GPU (W2). **The capture-failure drain is NOT among them: it landed HERE**, as behaviour (`std::uncaught_exceptions()` compared against the depth recorded at scope entry, so a break function or ordinary model code throwing mid-capture destroys the partial container instead of handing back a forward that reports `captured() == true`) and as three gated arms (tests 13a, 13b, 13c). The spec's `## Owed` strikes the item through and reads DELIVERED in W1; this cell said the opposite until 2026-08-18 because `cba969857` re-derived field 6 alone. Still NO throughput claim; analysis [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `ACTIVE` | `CLAIM-ENG-CUDAGRAPH-BREAK-W1`; [#1163](https://github.com/mudler/vllm.cpp/issues/1163), [#1192](https://github.com/mudler/vllm.cpp/issues/1192), [#1020](https://github.com/mudler/vllm.cpp/issues/1020) | | `ENG-CUDAGRAPH-DIFFUSION` | Capture the LTX-2.5 denoise loop (fixed shapes, many identical iterations — the ideal graph target). **BLOCKED, and the blocker is ours:** the render does almost no device compute to capture | T2 | SGLang enabled BCG on this shape AFTER our pin — LTX-2 H200 two-stage 10.75s->6.90s (`d4be483efb`), SANA 1024px -26% (`6c7498113f`), SANA denoise 0.73->0.457s (`56ef810cad`). Dated events, NOT pinned evidence; their win is mostly PyTorch host tax we do not pay | NO capture at all: `grep` for capture across `src/vllm/model_executor/models/ltx2*.cpp` returns nothing | blocked by [#1024](https://github.com/mudler/vllm.cpp/issues/1024) (GPU util **exactly 0 in 321 of 347 samples**, 1.00 core of 20 held for 17+ min after staging), [#1007](https://github.com/mudler/vllm.cpp/issues/1007) (VAE decode has no device arm), [#1087](https://github.com/mudler/vllm.cpp/issues/1087) (**57-66% of wall** is ONE resolution-CONSTANT serial host phase), [#1010](https://github.com/mudler/vllm.cpp/issues/1010) (no phase-boundary log). Decision point is a MEASUREMENT of GPU-busy vs wall once device-resident, not an implementation | [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `INVENTORIED` | [#1164](https://github.com/mudler/vllm.cpp/issues/1164) | | `ENG-BATCH-INVARIANT` | Opt-in deterministic execution across scheduler batch sizes (`VLLM_BATCH_INVARIANT=1`): batch-invariant matmul/norm/attention/collectives plus persistent-scheduler NVFP4; production default remains off | T1 | default/env `vllm/envs.py:89,576-578`; initialization `vllm/v1/worker/gpu_worker.py:1262`; NVFP4 dispatch `csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_sm120_kernels.cu:212-220`; suite fixture `tests/v1/determinism/conftest.py:9-12`; operator/e2e `tests/v1/determinism/test_nvfp4_batch_invariant_scaled_mm.py`, `tests/v1/determinism/test_nvfp4_batch_invariant.py` @ `702f481` | - | [W3-C3R executed contract](specs/nvfp4-persistent-plan-cache.md#w3-c3r-batch-shape-localization-and-gate-correction-2026-07-13): production-default ours and vLLM both change outputs across batch shapes; no local opt-in implementation is claimed | `planned: specs/batch-invariant-execution.md` | `INVENTORIED` | - | | `ENG-ASYNC-SCHED` | Async/overlap scheduling (AsyncScheduler placeholders + depth-2 batch-queue step + async D2H on a copy stream); vLLM's DEFAULT at the pin — mirror obligation per B3. **Host-side machinery + runner device-input half + sampler-OUTPUT half LANDED + CPU-gated (2026-07-16):** `AsyncScheduler` placeholder accounting, `step_with_batch_queue` depth-2, `ResolveAsyncScheduling` default-ON-when-compatible + `MaxConcurrentBatches`, `VT_ASYNC_SCHED` rollback; the runner device-input path `combine_sampled_and_draft_tokens`; PLUS the sampler-OUTPUT half — `vt::Backend` event/pinned primitives (`AllocPinned`/events, CUDA cudaHostAlloc+cudaEvent, CPU sync-degeneration), `AsyncGPUModelRunnerOutput` (device sampled-id snapshot → non-blocking D2H on a copy queue + event; `get_output()` waits only that event; MAIN queue never blocked), `Sampler::forward(sampled_ids_out)` device-resident greedy, `GPUModelRunner::sample_tokens_async` + `runner_supports_async`, and the `Executor`+`step_with_batch_queue` seam resolving `get_output()` at CONSUME time. All behind `VT_ASYNC_RUNNER`/`set_async_input_combine`, default OFF. Sync path byte-identical (placeholder sites INERT while count 0; combine off; `sample_tokens_async` degenerates to sync when async off; `sampled_ids_out=nullptr`). **ENABLE-FLIP LANDED + CPU-gated (2026-07-16):** (1) `LoadedEngine` now reorders `runner_` before the scheduler and builds an `AsyncScheduler` + `max_concurrent_batches=2` when `ResolveAsyncScheduling(runner_.runner_supports_async())` resolves ON (else the byte-identical synchronous `Scheduler` + depth-1); the resolved mcb threads into `AsyncLLM`→`EngineCoreProc` (`step_with_batch_queue`) and the "Asynchronous scheduling is enabled/disabled" log mirrors vLLM for A/B audit; (2) the device combine/scatter kernel (`_combine_sampled_and_draft_tokens_kernel` + last_sampled scatter) is ported to CUDA (`src/vt/cuda/cuda_combine_tokens.cu`), main-stream-ordered on the CUDA async path so it DELETES `sample_tokens_async`'s pre-scatter `Synchronize`; the CPU backend keeps the host loop. `VT_ASYNC_RUNNER=1` engages full W3; `VT_ASYNC_SCHED=0` is the same-binary rollback. Production default (no env) stays synchronous byte-identical. **FULL W3 DGX proof RAN twice** — `f086b64` (5/5 gates PASS; c16 TPOT −5.4 ms WIN, tput neutral, TTFT +36 % = Little's-law repayment) and the 2026-07-16 re-proof on the THROUGHPUT-lever fix (persistent pooled sampled-id/pinned buffers + `Sampler` greedy scratch removing ALL per-step `cudaMalloc`/`cudaFree`/`cudaHostAlloc`/event-create from the sampled-id path, incl. the overlap-killing `cudaFree` inside `get_output`; mirrors `gpu_model_runner.py:873-878` + `async_utils.py:12-70`): token-exactness **6/6 PASS**, interleaved c16 **tput −0.32 % (gate ≥+1.5 % FAILS), TPOT −4.95 ms retained, TTFT +34.8 %** — the allocator lever is REFUTED as the tput unlock (≤0.1 % of a ~165 ms c16 step). **DEFAULT FLIPPED ON 2026-07-17** (`VT_ASYNC_RUNNER` default ON via the pure `AsyncRunnerFlagIsOn` predicate, mirroring `vllm/config/vllm.py:992-1044`): the discriminator (`6ea7856`) proved vLLM's own async pays the identical +26–31 % TTFT / −0.7 to −0.9 % tput / −2.6 to −4.3 ms TPOT envelope and W3-ON nets positive (both binding ITL-tail anomalies flip to PASS), so the "needs a throughput lever" ship-gate is RETIRED — W3 is a parity/mirror obligation with a tails+TPOT win. The flip is TOKEN-NEUTRAL (async-ON ≡ async-OFF bit-identical on DGX). `VT_ASYNC_RUNNER=0` = runner-level rollback, `VT_ASYNC_SCHED=0` = scheduler-level rollback. TTFT means rise into vLLM's async envelope BY DESIGN — the next binding grid runs async by default and its TTFT must NOT be misread as a regression. **ROBUSTNESS FIX 2026-07-20 (`discard_request_mask`):** the runner was missing vLLM's `discard_request_mask`, so `GPUModelRunner` emitted a sampled token for prefill-CHUNK requests too; under async this drained a `num_output_placeholders` never reserved (the `is_prefill_chunk` path adds none) → the `async_scheduler.cpp` `num_output_placeholders >= 0` assertion aborted on c8 + short-output (chunked prefill + preemption). FIX mirrors vLLM: `execute_model` computes `exec_state_.discard[i] = seq_len < num_tokens` (`gpu_model_runner.py:2048`); `sample_tokens` clears those rows to empty (`outputs.py:303`), the async path passes `invalid_req_indices` to `AsyncGPUModelRunnerOutput::get_output` (`gpu_model_runner.py:3625` + `outputs.py:303`). Scheduler UNCHANGED (assertion kept — it was correct once the runner honors `scheduler.py:1888-1890`). Sync/non-chunked decode byte-identical (mask all-zero); DGX 27B 235/235 + 35B 315/315, `vllm-bench` c8+short-output+chunked+kv-pressure no longer crashes, memcheck 0. Ledger [parity-ledger.md](parity-ledger.md) 2026-07-20 row | T1 | `vllm/v1/core/sched/async_scheduler.py:12`; `vllm/config/vllm.py:490,990,1038`; `vllm/v1/engine/core.py:519`; `vllm/v1/worker/gpu/input_batch.py:304-406`; `vllm/v1/worker/gpu/async_utils.py:12-70`; `vllm/v1/worker/gpu/gpu_model_runner.py:242-332`; `vllm/v1/outputs.py:298-307` | `src/vllm/v1/core/sched/async_scheduler.cpp:10,45`; placeholder plumbing `src/vllm/v1/core/sched/scheduler.cpp:148,164,605`; `src/vllm/v1/engine/core.cpp:91` (`step_with_batch_queue`, async-output seam); `src/vllm/v1/engine/core_proc.cpp:32,46`; config `include/vllm/config/scheduler.h:117,165,188`, `src/vllm/config/scheduler.cpp:12`; `include/vllm/v1/request.h:187`; runner input leaf `src/vllm/v1/worker/gpu/prepare_inputs.cpp`, `src/vllm/v1/worker/gpu/input_batch.cpp`; runner output leaf `include/vt/backend.h`+`src/vt/backend.cpp`+`src/vt/cuda/cuda_backend.cu` (event/pinned), `include/vllm/v1/worker/gpu/async_output.{h,cpp}` (`AsyncGPUModelRunnerOutput`), `src/vllm/v1/sample/sampler.cpp` (`sampled_ids_out`), `src/vllm/v1/worker/gpu/runner.cpp` (`sample_tokens_async`/`runner_supports_async`), `src/vllm/v1/executor/executor.cpp`+`include/vllm/v1/worker/gpu/model_runner_base.h` (async seam); enable-flip `include/vllm/entrypoints/model_loader.h`+`src/vllm/entrypoints/model_loader.cpp` (`runner_` before scheduler, `ResolveAsyncEnabled`/`MakeScheduler`, `AsyncScheduler`+mcb=2, log), `include/vllm/v1/engine/async_llm.h`+`src/vllm/v1/engine/async_llm.cpp` (mcb param → `EngineCoreProc`); device kernel `include/vt/cuda/combine_tokens.h`+`src/vt/cuda/cuda_combine_tokens.cu`, wired `src/vllm/v1/worker/gpu/runner.cpp` (CUDA combine/scatter branch removes the pre-sync) | `tests/vllm/v1/test_async_scheduler.cpp:1` (6 cases, 54 asserts; RED vs base Scheduler 2/6 fail); depth-2 engine cycle `tests/vllm/v1/test_engine_core_proc.cpp:479` (mcb=2, async-output seam); config resolution `tests/vllm/test_scheduler_config.cpp:75`; enable-flip construction matrix `tests/vllm/entrypoints/test_loaded_engine_dense.cpp` (runner×VT_ASYNC_SCHED → scheduler type + mcb; RED = un-flipped engine, 3/3 ON-arm asserts fail); runner input leaf `test_combine_tokens.cpp` (RED = stale → 5/7 fail), `test_input_batch.cpp`, `test_runner.cpp` (async-ON≡sync); output leaf `tests/vt/test_backend.cpp` (event/pinned contract), `tests/vllm/v1/worker/test_async_output.cpp` (materialize/flush/snapshot; RED = +1 splice), `test_runner.cpp` (`sample_tokens_async` decode ≡ sync); full CPU ctest 111/111, tools 164/164. Prior diagnostic `3812d8` six-leg control: total **1.002153×**, TTFT **0.862159×**, no GPU-time reduction (neutral for speed). **DEFAULT-FLIP (2026-07-17):** new pure CPU flag test [test_async_runner_flag.cpp](../tests/vllm/v1/worker/test_async_runner_flag.cpp) (11 asserts, default-ON/'0'-off); construction matrix [test_loaded_engine_dense.cpp](../tests/vllm/entrypoints/test_loaded_engine_dense.cpp) INVERTED (default → AsyncScheduler+mcb=2; RED verified 5 asserts fail vs un-flipped). CPU clean `-Werror` rebuild, full serial ctest **116/116**, tools **164/164**. **DGX re-confirmation** (evidence `dgx:~/work/vllm.cpp-async-flip`, CUTLASS+FA2 hard-verified, one flock): shipping default (async ON + RMSNorm-fast OFF) → **27B 235/235 + 35B 315/315** with the "Asynchronous scheduling is enabled (mcb=2)" log, and both rollback arms (`VT_ASYNC_RUNNER=0`, `VT_ASYNC_SCHED=0`) 235/235 + 315/315 log "disabled"; async arms BIT-IDENTICAL (token-neutral). Closing record [parity-ledger.md#L502](parity-ledger.md#L502) | [async-serving.md](specs/async-serving.md) | `DONE` | `6ea7856` | diff --git a/.agents/issue-index.md b/.agents/issue-index.md index 16bd934f9..16c0a8390 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -366,6 +366,7 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1179](https://github.com/mudler/vllm.cpp/issues/1179) | `ENG-CUDAGRAPH-BREAK` | The hand-rolled decode-graph driver count recorded in `9bc4d7f44` is **eight** and is actually **nine**, and the row it feeds was framed as coverage-only when it is also correctness. The ninth is the DFlash draft graph, file-local with no header declaration, at `src/vllm/model_executor/models/qwen3_dflash.cpp:771,870,1038,1091,1095,1106` — its own `int g_state = 0` three-state machine (`:771`), its own `VT_DFLASH_GRAPH` kill switch (`:870`) instead of the `VLLM_CPP_CUDAGRAPH` the six batched drivers read, its own invalidate-on-block-width-change (`:1038-1047`) and its own `try { EndCaptureGraph(); } catch (...) {}` drain (`:1106`). The eight-count is stated in four places, all corrected here: [`sglang-breakable-cuda-graph.md`](specs/sglang-breakable-cuda-graph.md) §4 and `## Owed`, [`.agents/engine-matrix.md`](engine-matrix.md) rows `ENG-CUDAGRAPH-BREAK` and `ENG-CUDAGRAPH-DEDUP` ("times eight drivers", which sizes #1162's signature table), and [`.agents/roadmap_v1.md`](roadmap_v1.md) track `C12`. The reframing is the substantive half: `ENG-CUDAGRAPH-BREAK` was recorded as a COVERAGE row, and the duplication has already cost a SHIPPED model its decode graph. `src/vllm/model_executor/models/qwen3.cpp:961-986` declines the decode graph outright whenever the asynchronous device-token mirror is live, on its own measured battery — `depth-1, graph ON PASS 78/78`; `depth-2, graph OFF PASS 82/82`; `depth-2, graph ON FAIL, slots 1-3 degenerate` — because `Step()` replays against the HOST `input.token_ids` and the combine has patched the DEVICE ids. The comment names the real fix as reading the identifiers at replay time from a stable device buffer, and that fix exists, in exactly one sibling driver, as `StepDevInputs` (`src/vllm/model_executor/models/qwen3_5.cpp:3894`): `grep -c StepDevInputs` returns 41 lines there and 0 in each of `qwen3_moe.cpp`, `qwen3.cpp`, `deepseek_v2.cpp` and `voxtral.cpp`. One capability, written once, unavailable to four models, with a live mitigation standing in its place. This does NOT weaken the framing rule that `ENG-CUDAGRAPH` established: the row still makes no throughput claim, and the prefill refutation (GB10 3.8% host-idle between launches, GPU-busy >96%, 27B prefill gap 92.5% non-GEMM glue) stands unchanged. Coverage AND correctness, never speed. Fixed in flow with the [`eng-cudagraph-break.md`](specs/eng-cudagraph-break.md) review repair ([#1163](https://github.com/mudler/vllm.cpp/issues/1163)) | record | | [#1181](https://github.com/mudler/vllm.cpp/issues/1181) | `FIX-READ-F32-SCALAR-GUARD` | `ReadF32Scalar` (`src/vllm/model_executor/models/qwen3_5_weights.cpp:312-318` @ `ab6e65216`) bounds its input with `t.data != nullptr && t.nbytes >= sizeof(float)`, a LOWER bound, and then `memcpy`s four bytes into a `float`. Two silent wrong-value paths follow and neither fails: an ARRAY is reduced to element 0, so a block-wise FP8 scale grid of shape `[ceil(N/128), ceil(K/128)]` passes and stands in for the whole weight (measured under [#1166](https://github.com/mudler/vllm.cpp/issues/1166) on `Qwen/Qwen3.8-27B-FP8` @ `017b9c7af6b5689d5dd426a76e0bc077eb5ca20a`, `q_proj.weight_scale_inv` is `[96, 40]`), and ANY dtype is reinterpreted, since that same tensor is `BF16` and its four bytes are two bf16 values read as one float. Both return a finite plausible float, so the output is fluent, plausible and wrong, which is what a token gate cannot see. Upstream makes both facts structural rather than optional: a per-tensor scale is a distinct parameter TYPE that asserts `loaded_weight.shape[0] == 1` (`vllm/model_executor/parameter.py:260-272,304-309` @ `555967922`, plus the `_assert_and_load` shape assert at `:93-96`), the slot is allocated `torch.float32` so a narrow on-disk dtype is VALUE-converted rather than reinterpreted (`utils/fp8_utils.py:1276`), and the declared strategy TENSOR/CHANNEL/BLOCK picks the parameter type before a byte is read (`compressed_tensors/schemes/compressed_tensors_w8a8_fp8.py:63,128`). The AUDIT corrects the issue's own framing twice. The 27 grep hits across five files are 5 definitions, 20 call sites and 2 comment references, and both counts are short: `ReadCtF32Scalar` (`include/vllm/model_executor/models/dense_weight_loaders.h:376`) is a SIXTH copy of the same defect under another name, reached from a SIXTH model file (`src/vllm/model_executor/models/qwen3_weights.cpp:100,126-128` through `LoadCtNvfp4W4A16`). Of the six, three check nothing, `LnReadF32Scalar`/`ShReadF32Scalar` check dtype but not count, and only `nemotron_h_weights.cpp:557-573` is correct, which makes it the model the shared guard generalizes. No call site legitimately passes a multi-element or non-F32 tensor, and every existing fixture emits rank-0 or `{1}` `F32`, so nothing in the tree needed the leniency. It is NOT merely latent: `dense_weight_loaders.h:73-74` and `docs/BENCHMARKS.md:52` both record `unsloth/Qwen3.6-27B-NVFP4` @ `ccdaab7e` as FP8 W8A8 throughout with BF16 PER-OUTPUT-CHANNEL scales, and `LoadAttnDense` branches on the weight dtype alone (`qwen3_5_dense_weights.cpp:478-480`), so those projections enter the per-tensor arm and hit both defects at once under the tensor name the loader actually asked for, with no misspelling to stop them. Fixed in flow by one `dense_loaders::ReadF32Scalar(get, name)` that refuses `numel != 1` naming the shape, refuses a non-`F32` dtype naming the dtype, and requires exactly four readable bytes, with the other five copies deleted onto it and `nemotron_h`'s `Loader`-based twin kept as the one tracked exception. A narrow dtype is refused rather than converted, because a one-element BF16 scale has never been read correctly here and the BF16 layout that IS shipped is per-channel, which the count check refuses first. Per-channel FP8, block-wise FP8 and any explicit narrow-dtype conversion stay owed. Spec [`read-f32-scalar-guard.md`](specs/read-f32-scalar-guard.md) | bug | | [#1185](https://github.com/mudler/vllm.cpp/issues/1185) | `ENV-ORACLE-WHEEL-IN-LEASE` | The pinned vLLM oracle BUILDS, installs, imports and sees the GPU inside an `rc` lease on `dgx:gpu0`, measured 2026-08-18, which falsifies the `nvcc` clause four records carried. [`lease-runtime-staging.md`](specs/lease-runtime-staging.md) said the oracle "needs `nvcc`, which the worker still lacks", and `.agents/environment.md`, [`mtp-k-gt-1.md`](specs/mtp-k-gt-1.md) and [`gpu-lease-methodology.md`](specs/gpu-lease-methodology.md) each derived a blocker from it. The build job (`buildvllm.sh`, staged sha256 prefix `15e140d41f44e7c2`) asserted the checkout against the pin BEFORE compiling, printing `PIN CONFIRMED` at `5559679229bc961848b121ccdeaa8fa5d79bec98` and aborting otherwise, took `nvcc` from the toolkit row `MODEL-NEMOTRON-H-ABI-A3-E2E` staged (`NVCC_RC=0`, CUDA `release 13.3, V13.3.73`) and produced `WHEEL_RC=0`, `PERSIST_RC=0` and a 434 MiB `vllm-0.1.dev1+g555967922.cu133-cp312-cp312-linux_aarch64.whl`, sha256 `7c58b339741a288fbb313f4f5196c9c92a9e3b3c3ebe2ea970b0ff50bb9bcba4`. The identity job (`oracleenv.sh`, prefix `6119f5223f5d818c`) asserted from `cd /`, outside any source tree: `vllm.__version__ = 0.1.dev1+g555967922`, `IDENTITY_RC=0`, `cuda True NVIDIA GB10`, `CUDA_RC=0`. SCOPE, and it carries the same weight as the result: RUNNING A MODEL IS UNTESTED. Only build, install, import and `torch.cuda.is_available()` are measured, and [`mtp-k-gt-1.md`](specs/mtp-k-gt-1.md) records that the last time an oracle reached this far it consumed the host in the step AFTER `torch.compile` and REBOOTED the box, at `gpu_memory_utilization` 0.75 and again at 0.30, so the fraction is not the lever. The version string is an OPEN discrepancy: `.agents/upstream-sync.md` records `vllm_runtime_version = 0.23.1rc1.dev1511+g555967922`, the commit segment matches and satisfies the pin's binding `+g` rule, and the prefix differs because a shallow fetch stops `setuptools_scm` counting commits since the last tag, so a full-string gate needs a deeper fetch or a recorded pretend-version. The venv is NOT staged, because that job was killed at a 90-minute ceiling mid-copy and its partial tree was removed, so only the WHEEL is durable. Four staging walls, all artifacts of the NAS rather than of CUDA: `cp -a` preserves `file_mode=0664` so `nvcc` exited 126. CIFS `nounix` stores no symlink so `include` and `lib64` vanished and CMake reported `Could NOT find CUDA (missing: CUDA_INCLUDE_DIRS CUDA_CUDART_LIBRARY) (found version "13.3")`, naming the version and denying the toolkit in one line. 32 library links `libfoo.so` and `libfoo.so.MAJOR` had to be rebuilt because only the `libfoo.so.X.Y.Z` real files survived. And `markupsafe` existed as a dist-info with NO package files from a `pip --target` killed at a 35-minute ceiling, so Marlin codegen died on `ModuleNotFoundError`. The `rc` worker container is REUSED between jobs, so a repair inside a staging branch is skipped on the next run (`nvcc already in place`) and an environment repair must be unconditional and assert its postcondition. CONSEQUENCE for the rows #1129 blocked, [#1003](https://github.com/mudler/vllm.cpp/issues/1003), [#915](https://github.com/mudler/vllm.cpp/issues/915), [#821](https://github.com/mudler/vllm.cpp/issues/821) and [#81](https://github.com/mudler/vllm.cpp/issues/81): UNBLOCKED FOR THE BUILD STEP and STILL BLOCKED FOR A MODEL RUN. None can take a measurement until a model run is demonstrated. Job details, walls and non-claims in [`oracle-wheel-in-lease.md`](specs/oracle-wheel-in-lease.md) | verification | +| [#1192](https://github.com/mudler/vllm.cpp/issues/1192) | `ENG-CUDAGRAPH-BREAK` | W1 of the break-point capture seam: `vt::BreakableGraph`, `vt::GraphCaptureScope` and `vt::GraphBreak` land in `include/vt/breakable_graph.h` + `src/vt/breakable_graph.cpp`, with the SGLang unit suite ported case for case and ONE break point registered on ONE model. **The stage's exit criterion was answered first, on a leased GPU, because W0 deliberately did not answer it**: `cudaStreamEndCapture` followed by `cudaStreamBeginCapture` on the SAME stream mid-forward with EAGER work between them is LEGAL under `cudaStreamCaptureModeThreadLocal`, our backend's mode (`src/vt/cuda/cuda_backend.cu:204-206`) — measured on `orin:gpu0` through an `rc` lease, driver 12060, three replays with fresh inputs and 0 mismatches, plus the bare zero-work re-begin (`breakable_cuda_graph.py:370-374`). The first probe REFUSED and that refusal was the probe's, not CUDA's: `dlsym` on the bare symbol name binds the LEGACY v1 entry points, which are not capture-aware, so `cuMemcpyDtoDAsync` returned `CUDA_ERROR_STREAM_CAPTURE_UNSUPPORTED` and `cuStreamBeginCapture` (v1) silently took no capture-mode argument — a broken instrument failing toward a verdict about the code. Bound by exact versioned name, the criterion holds. The break point is the DENSE ATTENTION ENTRY of `Qwen3ForCausalLM` (`src/vllm/model_executor/models/qwen3.cpp`, inside `RunLayer`), which is vLLM's boundary and not one this row invents (`splitting_ops` defaulted to the attention family, `vllm/config/compilation.py:517,764-772,1145` @ `5559679229`) registered in SGLang's form because we have no compiler to match operation names against (`layers/radix_attention.py:256` @ `f63458b5be`). The G2 reachability gate is `tests/vllm/models/test_qwen3_break_point.cpp`: it drives the production `Qwen3DenseModel::Forward` with a scope open and counts `num_hidden_layers + 1` segments, so deleting the one call site takes it to 1 and turns the gate RED — a unit test that constructs the container by hand could not have asked that. It also holds G4 in the same case: the logits are compared BIT FOR BIT against the unscoped forward, 500 values, 0 differing. **STAGED SLICE, named rather than implied:** `GraphCaptureScope` and `BreakableGraph` are not yet ENTERED from a production step, because no driver opens a scope until W2 migrates `Qwen3DenseDecodeGraph`; the break point itself is on the production path on every forward and takes the pass-through arm. Listed under `## Owed` in [`eng-cudagraph-break.md`](specs/eng-cudagraph-break.md). The capture-failure drain is NOT owed: it landed here, as behaviour and as three gated arms. The destructor `catch` alone covered a throwing `EndCaptureGraph` and nothing else, so a break function or ordinary model code throwing mid-capture left a container reporting `captured() == true` over a PARTIAL forward; comparing `std::uncaught_exceptions()` against the depth recorded at scope entry is what tells those apart, and tests 13a, 13b and 13c gate the three arms. Still owed by later stages and NOT claimed here: the auxiliary-stream auto-join before every segment close (spec D10, live at `qwen3_5.cpp:6254-6255,6384` and `laguna.cpp:2572-2576,2612`, which W4 and W5 migrate), the non-capturing-backend arm on ROCm and Tenstorrent (G5), G1 bit-exactness on a real GPU over more than one replay, and one residual the fresh review surfaced: an exception CAUGHT INSIDE the capture scope leaves the segment closed and the rest of the forward uncaptured while `captured()` stays true, which the drain cannot see because nothing is unwinding at scope exit (spec `## Owed`, W2). No throughput claim is made anywhere in this change: the prefill refutation stands (GB10 3.8% host-idle between launches, GPU-busy above 96%, 27B prefill gap 92.5% non-GEMM glue) | feature | | [#1189](https://github.com/mudler/vllm.cpp/issues/1189) | `VT-QUANT-FP8-GROUP` | Block-wise (128x128) FP8 so `Qwen/Qwen3.8-27B-FP8` runs instead of being refused. `weight_block_size` appears nowhere in `src/` or `include/`, so `469f38395` refuses the arm by name (#1166). Six independently landable milestones; M1 lands here. M1 is `vt::QuantFp8Group`, the dynamic per-token per-group activation quant, CPU and CUDA. The numerics mirror the kernel that ACTUALLY EXECUTES on a CUDA-alike platform with a contiguous input, which is the C++ custom op at `csrc/libtorch_stable/quantization/w8a8/fp8/per_token_group_quant.cu:42-96` and NOT the Triton kernel at `fp8_utils.py:95-150`: `fp8_utils.py:635-650` calls the former and returns before the latter. The two arms differ, and the difference is measurable rather than cosmetic. The CUDA kernel divides twice, `local_absmax / max_8bit` at `:68` and `static_cast(src) / y_s` at `:85`; the Triton kernel multiplies by `(1.0 / fp8_max)` at `fp8_utils.py:145` under a comment that names the 1-ULP gap. Upstream's own test tolerates the gap with `rtol=0.15` (`test_block_fp8.py:112-114`), so a value comparison cannot tell the two apart and only a byte comparison against a spelled-out reference can. `eps` is the reduction's INITIAL value (`:47`), not a post-clamp, which is what keeps an all-zero group from dividing by zero. Scope refused here and owed to later milestones: the block-scaled GEMM (M2), `Fp8BlockWeight` and the loader (M3), `Fp8BlockLinearMethod` and the Qwen3.5 wiring (M4), the mainloop-scaled CUTLASS kernel and the column-major/TMA-aligned scale layouts (M5), merged `gate_up`/QKV (M6). M1 lands UNREACHED: no production entry point dispatches `vt::QuantFp8Group` at its merge commit, M4 owns the wiring, and `.agents/specs/vt-quant-fp8-group.md` lists it under `## Owed`. The CUDA arm compiles and its on-hardware leg is owed too, because the row took no GPU lease by design: the CPU arm is the gateable one | feature | | [#1190](https://github.com/mudler/vllm.cpp/issues/1190) | `ENV-AGNOSTIC-CAMPAIGN` | One operator's hosts, share paths and addresses are written into 227 tracked files where a `.env` placeholder belongs, so a second developer who follows the protocol documents is told to reach a box on one home network. Re-derived at `fd64c76ee` with `git grep -cIE 'dgx\.casa\|nas_share\|192\.168\.\|thor:gpu0'`: `dgx.casa` 203 files, `nas_share` 32, `192.168.` 24, `thor:gpu0` 11. The mechanism was already there and almost unused: `.env.example` is tracked, `.env` and `.agents/developer-preferences.md` are ignored, `scripts/agent-onboard.py --env-set` already records one answered value and refuses an undeclared key, and `${VLLM_SOURCE}` and `${GPU_LOCK}` already resolve in 59 and 50 files while `${DEVICE_ARCH}`, `${DEVICE_TOOLKIT_ROOT}` and `${DEVICE_COMPILER}` resolve in none. This SCOPING row lands the rule and the mechanism, not the sweep. The rule is a substitution test: replace the literal with a second developer's value and if the sentence stays true it is CONFIGURATION and becomes `${KEY}`, and if it becomes false it is PROVENANCE and stays literal. That reverses the issue's own ranking, because all 34 hits in the densest guide file `.agents/environment.md` are provenance or named-profile definition and none is configuration, so density does not predict the defect and a blind `sed` would falsify records. Landed: three new keys `GATE_CHECKOUT`, `SHARED_STORAGE_ROOT` and `GATE_DEVICE` derived from the literals that recur and map to no existing key; the create-on-first-use route in `scripts/agent-start.py`, which printed `environment: missing` as a status label and then listed next actions that never mentioned it, so the fallback in practice was a host name copied from a document; the matching obligation in `AGENTS.md`; and one worked example. The worked example `scripts/dgx-bringup.sh` found a live defect rather than a cosmetic one: it defaulted `CUTLASS_DIR` to `$HOME/cutlass_probe` while `.agents/environment.md:389` records `$HOME/cutlass-4.5.0` as mandatory on the same box, and a configure that misses CUTLASS silently drops the sm120a NVFP4 GEMM and FlashAttention-2, which that file measures as moving the SACRED `test_qwen27_paged_engine` from 235/235 to 234/235 with the source untouched, so a stale hard-coded default is a false green. Its new test also caught that `set -a; . ./.env; set +a` over an `.env.example` copy blanks a value the caller exported, so the process-environment contract is now executable. Waves `ENV-AGNOSTIC-W1-TOOLING` through `ENV-AGNOSTIC-W5-LEDGERS` own the sweep, partition all 227 files, and open their own issues. Spec [`env-agnostic.md`](specs/env-agnostic.md) | bug | | [#1193](https://github.com/mudler/vllm.cpp/issues/1193) | `SPEC-DSPARK-QWEN3-ROUTING` | A Qwen3 DSpark draft declaring `architectures=["DSparkDraftModel"]` with `model_type` `qwen3` has no route. The pin forces every DSpark draft that is not `Qwen3DSparkModel` or `Gemma4DSparkModel` onto `model_type` `deepseek_v4` (`vllm/config/speculative.py:934-944` @ `555967922`), and vLLM PR 52197 (merged 2026-08-17 at `7075ddac`) replaced that with a leading branch normalizing the pair to `Qwen3DSparkModel`. We diverge from BOTH: the forced rewrite was never ported, so nothing in `src/vllm/entrypoints/model_loader.cpp` reads a draft config's `architectures` key at all, and `SpeculativeConfig::IsDsparkDraft` (`include/vllm/config/speculative.h:120-136`) has no production caller — every reference outside its header is in `tests/vllm/config/test_speculative_dspark.cpp:132-140`, and `ResolveSpecConfig` branches on `cli.method` alone. The checkpoint is real and gateable here: `RadixArk/Qwen3.8-27B-DSpark` at revision `85ef153be924f17ce4bf62726954eeaa4a73e854` carries exactly that config shape in one 2718576122-byte shard, drafting five layers for a 64-layer Qwen3.8-27B target | bug | diff --git a/.agents/specs/eng-cudagraph-break.md b/.agents/specs/eng-cudagraph-break.md index e884dc60f..27da7e8ea 100644 --- a/.agents/specs/eng-cudagraph-break.md +++ b/.agents/specs/eng-cudagraph-break.md @@ -283,10 +283,16 @@ Three types, all under `vt`, all backend-agnostic, none of them a new virtual on 1. **`vt::BreakableGraph`** — the container. Holds `std::vector segments` (each an opaque handle from `Backend::EndCaptureGraph`) and `std::vector> break_fns` (each a seam-built replay - closure, see item 3, never a caller's raw `fn`). Invariant, asserted: - `segments.size() == break_fns.size() + 1`. `Replay(Queue&)` walks - `ReplayGraph(segments[i])` then `break_fns[i]()` for each `i`, mirroring - `breakable_cuda_graph.py:255-263`. The destructor calls `DestroyGraph` on every + closure, see item 3, never a caller's raw `fn`). Invariant: + `segments.size() == break_fns.size() + 1`. It is ASSERTED, in the two places + where violating it would otherwise be silent: `GraphCaptureScope` refuses to + open on a container that already holds a capture (re-entry appended to it and + produced `break_count() == segment_count()`), and `Replay` refuses counts that + disagree (such a graph replays with its LAST BREAK DROPPED, whose only symptom + is a forward that skips one host-dependent operation). Documenting it was not + enough: the first draft said "asserted" in three records and asserted it + nowhere. `Replay(Queue&)` walks `ReplayGraph(segments[i])` then + `break_fns[i]()` for each `i`, mirroring `breakable_cuda_graph.py:255-263`. The destructor calls `DestroyGraph` on every segment. This is the direct analogue of `BreakableCUDAGraph` at `:246`. 2. **`vt::GraphCaptureScope`** — a resource-acquisition-is-initialization (RAII) @@ -333,13 +339,30 @@ Three types, all under `vt`, all backend-agnostic, none of them a new virtual on the two, and the seam must make the choice explicit rather than leave it to the caller: - - `GraphBreak(fn, out)` takes the destination the following segment reads, and - the seam copies `fn`'s result into `out` on every replay. This is the direct - port of `replay_fn` plus `_copy_output`. + - `GraphBreak(fn, slot)` takes the destination the following segment reads, and + the seam copies `fn`'s result into it on every replay. This is the direct + port of `replay_fn` plus `_copy_output`. **The destination is a + `vt::BreakSlot`, never a bare `Out&`,** and that is a correctness + requirement rather than a style choice. The destination must outlive the + `BreakableGraph`, because the following segment bakes its address; upstream + gets that from holding `captured_output` BY VALUE with its storage pinned by + the segment graphs' mempool (`:156-169,225-227`). A C++ seam taking a + reference gets the opposite, because the natural call site declares the + destination as a local of the function containing the break — which is + exactly what the first W1 site did, registering a stack slot that died on + the next `return` while the pooled block it named went back on the free + list. `BreakSlot` moves the value into a cell the replay closure owns, so + the rule holds by construction and the obvious call site cannot break it. + A destination type with no `CopyOutput` overload is a COMPILE error naming + the type, not a silent downgrade into the fallback below. - `GraphBreak(fn)` with no destination requires `fn` to write IN PLACE into a persistent buffer that the model owns and that no replay reallocates. The no-destination form is then only legal for a break function with no return - value; the bare marker below is its degenerate case. + value; the bare marker below is its degenerate case. Upstream's own + non-copyable fallback (`:201`, "return `src`, leave the destination alone") + stays reachable, but it has to be ASKED FOR by passing `vt::NoWriteback{}`, + because a fallback that is the DEFAULT makes a renamed, moved or misspelled + `CopyOutput` overload compile clean and be wrong only on replay N. `vt::BreakableGraph` therefore holds `std::vector>` whose elements are the SEAM's closures, never the caller's `fn` — and D9 records the @@ -483,12 +506,23 @@ migrated model. Excluded as redundant to a stronger gate, not as unaffordable. **The one unavoidable harness adaptation.** Every upstream class's `setUpClass` raises `unittest.SkipTest` without CUDA (`:34-36`, `:177-178`, `:235-236`), so upstream runs all eleven cases on a real device. T1 through T5, and T7 through T11, -run instead against a test backend that RECORDS the call sequence, so they gate on -this box and in continuous integration with no GPU. T6, the multi-replay case, has -both arms: the recording arm asserts the replay ORDER, and the GPU arm under G1 -asserts the VALUES. Nothing else about the upstream cases is adapted: the same break -counts, the same segment counts, the same arithmetic chains, the same in-place versus -assign split, and the same non-copyable fallback. +run instead against a test backend that RECORDS the call sequence and SIMULATES the +graph: `Record(fn)` files one captured operation against the open segment and does +NOT execute it, exactly as a real stream capture files a kernel without running it, +and `ReplayGraph` runs the filed work in order. That is what keeps upstream's +POST-REPLAY assertions literal — `x.fill_(5); graph.replay(); y == 6` is the same +assertion here — rather than degrading them into capture-time value checks, which +assert something upstream never asserted. The same break counts, the same segment +counts, the same arithmetic chains (`x=10 -> 11 -> 22 -> 25`, `x=5 -> 6 -> 7 -> 8 -> +16`, `3.0` then `33.0`, `x=10 -> 11 -> 13`), the same in-place versus assign split +and the same non-copyable fallback. + +**Two things the harness does NOT model, named rather than claimed away.** Real +device memory, asynchrony, and any error a real runtime would raise are absent; a +case that needs those is a GPU case and is owed rather than approximated. And +upstream's `eager_on_graph` always returns a fresh tensor, so every upstream break +is the DESTINATION form — the in-place form has no upstream case and is exercised +beside T2 and by the bare marker. The existing test surface the migration must keep green: @@ -514,6 +548,14 @@ numbered on from the ported set: already-instantiated segment released. This ports the recovery that `qwen3_5.cpp:9913` and `qwen3_dflash.cpp:1106` each hand-rolled. Upstream has no equivalent because Python's `finally` in `__exit__` (`:323-332`) covers it. + **The BEHAVIOUR was owed too, not only the test**, and an earlier draft of + `## Owed` said otherwise. The destructor's `catch` guarded a throwing + `EndCaptureGraph` and nothing else: an exception from a break function, or + from ordinary model code between two break points, left a PARTIAL capture + reporting `captured() == true`, which is replayable as half a forward. + Comparing `std::uncaught_exceptions()` against the depth recorded at scope + entry is what separates the two, and all three arms are gated (13a break + throws, 13b ordinary code throws mid-segment, 13c `EndCaptureGraph` throws). 14. **Non-capturing backend.** On a backend where `SupportsGraphCapture()` is false, the scope is inert and the forward runs eager. Vulkan (`vulkan_backend.cpp:16`) and Metal (`metal_backend.mm:13`) are the live cases. @@ -626,13 +668,92 @@ leased through `rc`. SGLang's `_end_current_segment` and `_begin_new_segment` pinned revision, which is strong evidence but is not our measurement. **This is W1's exit criterion and it needs a GPU lease.** It is named rather than skipped. -**W1, the seam plus its unit gate, and one break point on one model.** Land -`vt::BreakableGraph`, `vt::GraphCaptureScope` and `vt::GraphBreak` with tests 1 -through 5 of `## Tests to port`. Register exactly one break point, on one model, -reached from `ModelRegistry::Forward`. Confirm the CUDA re-begin behavior on a -leased GPU first; if it does not hold, W1 stops and reports rather than working -around it. G2's mutation applies from this stage on, because W1 is the first stage -that can be dead. +**W1, the seam plus its unit gate, and one break point on one model. DONE +2026-08-18, [#1192](https://github.com/mudler/vllm.cpp/issues/1192).** Landed +`vt::BreakableGraph`, `vt::GraphCaptureScope`, `vt::BreakSlot` and +`vt::GraphBreak` in `include/vt/breakable_graph.h` and +`src/vt/breakable_graph.cpp`, with tests 1 through 5 of `## Tests to port` plus +the `TestCopyOutput` set (T7 through T10), the bare marker (T11), the +replay-order case (test 12), the capture-failure drain (test 13, all three +arms), the non-capturing backend (test 14), the aliasing refusal with its two controls, +and the ownership case that proves every segment is released through +`Backend::DestroyGraph`. `tests/vt/test_breakable_graph.cpp`, 24 cases, 163 +assertions, exit 0, re-derived by `ninja test_breakable_graph && +./build/tests/test_breakable_graph`. + +**A fresh review returned FAIL on the first head and the repairs are part of +W1.** What the review found, and what closed it, because each one is a class of +defect a later stage can repeat: + +- The production break point's destination was a function-local + `std::optional`, so both scoped layers registered the SAME stack slot in + a frame that was gone before the scope closed, and the pooled block it named + went back on the `DevicePool` free list. Inert at W1 because no driver opens a + scope, and wrong the moment W2 does. Closed structurally: the destination is a + `vt::BreakSlot` the seam owns and there is no reference-taking overload left. + + **That closed the LIFETIME half of the rule, and this spec claimed the whole + rule was "unexpressible to violate". It was not, and a scoped re-review + measured the other half still open.** One slot reused for TWO break points in + one capture compiled: `PinForCapture` returns the cell it already made, so + both replay closures bind to the SAME address, `&*slot` is byte identical + after break 0 and after break 1, and after `Replay` the slot holds only break + 1's value — break 0's writeback overwritten, and any segment that baked break + 0's destination reading break 1's data. The production site is correct today + because its slot is a `RunLayer` local, but W2 through W5 add nine more + callers and the prose was telling them the shape was unwritable. Closed by a + REFUSAL rather than by the type, and the prose now names which half is which: + `GraphCaptureScope::AppendBreak` takes the destination as a REQUIRED + parameter, so no form can register without stating where it lands, and it + throws when a second break in the same capture names a cell already + registered. Two things the repair had to get right, both measured rather than + reasoned. The identity is the CELL and not the SLOT: the production slot is a + per-call local, every layer's slot lands at the same STACK address, and + slot-address identity therefore refused the correct program on layer 2 and + reddened the G2 gate. And the non-copyable fallback registers no destination, + because it pins no cell and its replay writes back nowhere. Red-first: the + case failed on the unrepaired head with the aliased capture ACCEPTED, + `break_count() == 2` into one slot. Two mutations, both detected — deleting + the refusal reds only the new case (23 of 24 still pass and G2 stays green, + which is exactly why nothing caught this before), and making it over-fire reds + two unit cases AND the G2 reachability gate, so the controls are not + vacuous. +- Interleaved replay was claimed, spec'd, and NOT gated. Replacing the + interleaved loop with "replay all segments, then run all breaks" left the suite + green, because break markers went into a vector the backend log knew nothing + about and the two sequences were asserted independently. Closed by putting the + markers into the backend's OWN log and asserting one trace. +- The writeback branch degraded SILENTLY: renaming `CopyOutput` dropped the site + into the non-copyable fallback, which is the D9 wrong-numerics path, with both + suites green. Closed by a `static_assert` in the destination form plus the + explicit `vt::NoWriteback` opt-in, and by moving the model's overload into a + header so the G2 gate can `static_assert` that the writeback branch is the one + selected. +- The `VLLM_CPP_CUDAGRAPH=0` kill switch was ungated (the existing case + substituted the OTHER conjunct), nesting was accepted and traced an illegal + capture sequence, the invariant three records called "asserted" was asserted + nowhere, and `Reset()` left the replay count stale. Each is now gated, the + switch through a child process because it is read once per process. +- Upstream fidelity was overstated. T3 had lost its arithmetic chain and T1 its + post-replay value, because a recording backend cannot re-execute a segment. The + backend now SIMULATES the graph, so every upstream chain and every post-replay + assertion ports literally, and the two remaining deviations are named above. + +**The exit criterion was answered FIRST, and it holds.** `cudaStreamEndCapture` +followed by `cudaStreamBeginCapture` on the SAME stream mid-forward with EAGER +work between them is legal under `cudaStreamCaptureModeThreadLocal`, which is the +mode `src/vt/cuda/cuda_backend.cu:204-206` uses. Measured on `orin:gpu0` through +an `rc` lease, driver `12060`: segment, host-dependent break, re-begin, segment, +bare zero-work re-begin, segment; then three replays with fresh inputs and +0 mismatches on every one. The BARE re-begin — end then immediately begin with no +work between, the degenerate case of `breakable_cuda_graph.py:370-374` — is legal +too. See `## Outcome` for what the first, refused probe actually measured. + +**The break point is the dense attention entry of `Qwen3ForCausalLM`** +(`src/vllm/model_executor/models/qwen3.cpp`, inside `RunLayer`), in its +destination-carrying form because `AttnBlock` returns a fresh pooled buffer on +every call. G2's mutation applies from this stage on and was performed; see +`## Outcome`. **W2, migrate `Qwen3DenseDecodeGraph` first.** It goes first for three reasons, in order. It is the SMALLEST batched driver by machinery, holding no @@ -817,13 +938,108 @@ point registered inside an unjoined fork window without this rule fails at ## Now -`READY`. The spec is committed, the design is grounded in both oracles, the -inventory is enumerated with anchors, the upstream unit suite is mapped case for -case in `## Tests to port`, and the work is decomposed into six landable stages. No -production code exists for this row. W1 needs an `rc` GPU lease to confirm its exit -criterion — that CUDA permits `cudaStreamEndCapture` followed by -`cudaStreamBeginCapture` mid-forward on our stream configuration — before -implementation starts. +`ACTIVE`. W0 (spike) and W1 (the seam, its ported unit gate, and one registered +break point) have landed; W2 through W6 remain, and `## Work breakdown` states +each. Owner: `.agents/claims/CLAIM-ENG-CUDAGRAPH-BREAK-W1.md`. + +W1's exit criterion — that CUDA permits `cudaStreamEndCapture` followed by +`cudaStreamBeginCapture` mid-forward on our stream configuration — is +CONFIRMED on a leased GPU and is no longer an open question for any later stage. + +## Owed + +Each item names the stage that owns it. Nothing here is claimed by W1. + +- **The seam is not yet ENTERED from a production step.** `GraphCaptureScope` and + `BreakableGraph` are constructed by the gates, not by a driver, because no + driver opens a scope until **W2** migrates `Qwen3DenseDecodeGraph`. The break + point itself IS on the production path — every `Qwen3ForCausalLM` forward + executes it and takes the pass-through arm — and + `tests/vllm/models/test_qwen3_break_point.cpp` is what holds that. This is the + staged slice AGENTS.md allows, named here rather than left silent. Owner: W2, + [#1163](https://github.com/mudler/vllm.cpp/issues/1163) tracks the handoff — + the parent, because #1192 CLOSES with W1 and a tracker that dies at merge + tracks nothing. +- **The auxiliary-stream auto-join before every segment close** (D10, the port of + `breakable_cuda_graph.py:353-361`). W1 registers its break point on a model + that forks no auxiliary queue, so the rule is not exercised and untested + machinery was not landed for it. A break point placed inside an unjoined fork + window today fails LOUDLY at `EndCaptureGraph`, which is the one failure mode + in this spec that is not silent. Owners: **W4** (`qwen3_5.cpp:6254-6255,6384`) + and **W5** (`laguna.cpp:2572-2576,2612`). +- ~~**The capture-failure drain as a GATED case** (test 13).~~ DELIVERED in W1, + and the record it replaces was wrong twice over. The destructor did NOT already + behave: its `catch` guarded a throwing `EndCaptureGraph` alone, and an + exception from a break function or from ordinary model code left a partial + capture reporting `captured() == true`. The behaviour and all three gated arms + landed together; see `## Tests to port` test 13. +- **The non-capturing arm on the OTHER capture backends** (G5). The inert path is + gated here through a recording backend reporting `SupportsGraphCapture()` + false; ROCm (`rocm_backend.hip:248`) and Tenstorrent + (`tenstorrent_backend.cpp:75-81`) are not exercised, and D6 records that + segmenting a ttnn mesh trace is UNVERIFIED on that runtime. Owner: **W3**. +- **G1, bit-exactness against eager on a real GPU over MORE than one replay.** + W1's bit-exactness is against the model's own eager forward on CPU, with the + scope open, 500 logits and 0 differing — which proves the seam changed no + numerics, and does NOT prove a replayed segmented capture matches. Owner: + **W2**, on the leased GPU its gate models already need. +- **An exception CAUGHT INSIDE the capture scope leaves a partial capture the + drain cannot see.** The `uncaught_exceptions()` comparison in + `~GraphCaptureScope` detects an exception that is PROPAGATING at scope exit. A + break function that throws and is caught between two break points inside the + scope leaves `segment_open_ == false`, so the remainder of the forward is + never captured, and nothing is unwinding at scope exit, so the drain does not + fire and `captured()` stays true over a forward that is missing its tail. The + header and this spec state the guarantee only for the propagating case, so + this is a RESIDUAL rather than a false claim — but a driver that wraps a break + in a `try` gets silently wrong numerics on replay, and W2 is the first stage + that opens a scope from a driver. Owner: **W2**. +- **The reuse hazard the seam must close** (D1): making the intermediates a + segment reads unavailable to the `DevicePool` free list for the life of the + `BreakableGraph`. W1 states the lifetime rules at the `GraphBreak` declaration + and enforces the OUTPUT half through the writeback contract; the INPUT half is + a pool change with its own argument. Owner: **W2**. + +## Outcome + +**What the first probe measured, and what it did not.** The exit-criterion probe +refused twice before it held, and neither refusal was CUDA's. Bound through +`dlsym` on the BARE symbol name, `libcuda` hands back the LEGACY v1 entry points, +which are not capture-aware: `cuMemcpyDtoDAsync` (v1) returned +`CUDA_ERROR_STREAM_CAPTURE_UNSUPPORTED` inside a legal capture, and +`cuStreamBeginCapture` (v1) takes NO capture-mode argument, so the probe believed +it was testing the thread-local mode while testing the global one. Preferring +`_v3` blindly then bound `cuCtxCreate_v3`, which takes two extra parameters, and +context creation failed with `CUDA_ERROR_INVALID_DEVICE`. Both readings presented +as a verdict about the design. Bound by EXACT versioned name the probe passes, +and the rule that survives is the general one: an instrument that can fail toward +a code verdict must assert its own precondition first. + +**Why the destination form, and not the in-place one, at the qwen3 site.** +`AttnBlock` returns a fresh `DBuf` from the device pool on every call, so the +in-place form's contract — write into a persistent buffer no replay reallocates — +is one the site cannot meet. Registering it in the in-place form would have +compiled, run, and been wrong on replay in exactly the way D9 describes. The +destination is `std::optional` with a `CopyOutput` overload in +`vllm::dense_attn`, found by argument-dependent lookup. + +**Why the writeback is a customization point rather than a fixed signature.** +Upstream's `_copy_output` (`:172-201`) reaches an object's fields through +`__dict__`; C++ has no reflection. Overloading on the output type reproduces the +three branches upstream has — in-place for a tensor, field-wise for a keyed set, +in-place-or-assign for a struct — and reproduces the fourth, the non-copyable +fallback (`:201`), as a type with no overload, for which the seam fabricates no +copy. T10 asserts the ABSENCE of that writeback deliberately: it is the warning a +future author needs, and asserting only its presence would have left the case +mute. + +**Why the counters exist.** G3 is not decoration. Without a segment count there +is no way to tell a two-segment capture from a fully eager step, and the second +G2 case uses `break_points_reached` to prove the model REACHES the seam even on +the pass-through arm, which is the only observable that distinguishes a +registered break point from a deleted one while no driver opens a scope. + +--- Revised 2026-08-18 after a fresh review returned `FAIL`. What changed: the upstream test suite is ported rather than declared absent; `## Port map` §3 states the output @@ -832,3 +1048,40 @@ stream fork-join rule is stated and gated (D10); four enumerations in `## Our baseline` and one `grep` count were re-derived and corrected against printed instruments; D6 now defers to W1's exit criterion instead of asserting it; and the row is recorded as coverage AND correctness per #1179. + +Revised 2026-08-18 again, when W1 landed (#1192): `## Work breakdown` W1 records +the confirmed exit criterion and what shipped, `## Now` moves the row to `ACTIVE`, +and `## Owed` and `## Outcome` are new. The framing is unchanged and is not +negotiable — coverage and correctness, never speed. No throughput gate is +declared and none was measured. + +Revised 2026-08-18 a third time, after a fresh review of the W1 head returned +`FAIL`. What changed: the destination-carrying break point takes a +`vt::BreakSlot` the seam owns instead of a caller reference it cannot outlive +(the rule the header stated and the only site using it broke); the interleaved +replay, the `VLLM_CPP_CUDAGRAPH=0` kill switch, the capture-failure drain, the +`segment == break + 1` invariant, the nesting refusal and the writeback-branch +selection are each GATED rather than described; the recording backend simulates +the graph so upstream's arithmetic chains and post-replay assertions port +literally; and the exit-criterion probe is committed with its recipe and hashes. +The framing is unchanged and is not negotiable — coverage and correctness, never +speed. No throughput gate is declared and none was measured. + +**The probe is COMMITTED, because the measurement that produced two false +refusals is the one that most needs a re-runnable form.** Both sources are in the +tree: `scripts/probe_cudagraph_rebegin.c` (the `dlsym` driver-API build, which +needs no CUDA toolkit and no headers, and is the one the exit criterion was +answered with) and `scripts/probe_cudagraph_rebegin.cu` (the runtime-API build, +which needs `nvcc`). `.agents/benchmark-record.md` carries their sha256 sums and +the exact recipe. **What was NOT retained is the leased run's own driver script +and its raw stdout**, so the verdict recorded above is reproducible only by +re-running the committed source under the recorded recipe on a leased GPU — which +is what a later claimant has to do rather than cite this paragraph. + +**The finding worth keeping at the top level.** The seam's most-documented rule — +that the destination must outlive the graph — was broken by the ONLY site that +used it, in the same change that wrote the paragraph explaining it. A rule stated +at a declaration and re-derived by every caller is not a mechanism. Where a +contract can be encoded in a type, encode it: W2 through W5 add nine more callers +and none of them should have to read that paragraph to be correct. + diff --git a/CMakeLists.txt b/CMakeLists.txt index 992447a8c..f425fb5f1 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1215,6 +1215,7 @@ add_library(vllm STATIC src/vllm/parser/parser_manager.cpp src/vt/dtype.cpp src/vt/backend.cpp + src/vt/breakable_graph.cpp src/vt/tensor.cpp src/vt/arena.cpp src/vt/op_provider.cpp diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index 065d8ee98..bb01a3dfb 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -13,6 +13,7 @@ | **NemotronH paged forward** (`MODEL-NEMOTRON-H-ABI-A2P`, [#810](https://github.com/mudler/vllm.cpp/issues/810)) | **No number on any axis, by the unit's own rule** ([spec](../.agents/specs/nemotron-h-a2p-paged-forward.md) §5) | **A3 gate PENDING: `nvcc` + checkpoint visibility in the rc container.** Two earlier causes here (contention, then "cannot build") were measured FALSE, see [benchmark-record](../.agents/benchmark-record.md) | CPU gate 12/12, 9/9 mutations RED. A3 driver `examples/nemotron_h_gen` exists; guards armed on a real engine (pass 0, divergence 1, short 4) | | **LoRA runtime W2** (`LORA-RUNTIME`, #278) | **No number owed:** correctness-only; a grid PENDS the W7 model gate | | **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 | +| **CUDA-graph break seam W1** (`ENG-CUDAGRAPH-BREAK`, [#1192](https://github.com/mudler/vllm.cpp/issues/1192)) | **No number owed, and none taken:** coverage and correctness row, no throughput gate declared | Capability, not a rate: mid-forward capture re-begin holds on a leased GPU; scoped forward matches eager, 500 logits, 0 differing. Probe committed, recipe and sha256 in the [record](../.agents/benchmark-record.md) | | **DSR fix: server TU profiler guards (2026-08-09)** | **No number owed:** comments only. #189 moved the server body into the shared layer with its 5 `VT_BENCH_PROFILE_CONTROL` guards, taking DSR 32 -> 37; they are `DSR-ALLOW`'d per site, baseline unchanged at 32 | | **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 | diff --git a/docs/STATUS.md b/docs/STATUS.md index 050469b2d..fb60c528c 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -181,6 +181,61 @@ token-for-token correctness against the pinned oracle. | SGLang parity (competitor floor + oracle) | Oracle STOOD UP + first floor MEASURED (cache-neutral, 27B, c8/c16): throughput/TTFT WIN, TPOT/ITL open GAP | SGLang v0.5.15 (`f63458b`) whole runtime surface inventoried (44 rows: 23 FUSED into our vLLM-derived engine, 8 SGLANG-DISTINCT opt-ins, 5 inventoried, 8 out-of-scope). The `v0.5.15-cu130` arm64 image (`@sha256:d0a667e`) PULLED and RAN the 27B-NVFP4 gate model on GB10 sm_121a with no from-source build. First reproduced SGLang-vs-ours cache-neutral comparison (27B, 3 reps, idle box, one flock, engines sequential): **ours beats the SGLang floor on total/output throughput + req/s (2.21×@c16, 1.44×@c8) and TTFT (6–12× lower), but SGLang wins per-token latency (TPOT/ITL 1.18–1.49× below ours) — a reproduced OPEN GAP.** SGLang is a competitor, not the mirror source (vLLM stays behavior truth). Residuals: 35B, c1/c2/c4 low-conc sweep, shared-prefix cache-ON arm, token-exact cross-check. Numbers + repro: `docs/BENCHMARKS.md`, `.agents/sglang-matrix.md`, `.agents/specs/sglang-parity-oracle.md`. **UPDATE 2026-07-28 (`CLAIM-DECODE-LATENCY-EXPLORE`, measurement only, no source changed): the TPOT/ITL gap is CONFIRMED batch-composition, NOT a decode-kernel deficiency** — on our engine ITL(decode-batch=1)=101.75 ms is already ≤ SGLang's op-point 104–105 ms and rises monotonically with batch (→158.5 ms @ B16); nsys shows every hot decode kernel sub-linear in batch (per-token cost ↓~10×); SGLang's effective decode concurrency is ~4 (not 16) due to its 33 s admission queue, so its low ITL is simply the ITL of a small batch. Our throughput win IS the ITL cost — same lever; knob `max_num_seqs`/`max_num_batched_tokens` already exists, latency-oriented point `max_num_seqs≈8` = ITL −21% at 1.38× SGLang throughput; default stays throughput-oriented. Full data: `.agents/specs/decode-latency-lever.md` **UPDATE 2026-08-16 ([#979](https://github.com/mudler/vllm.cpp/issues/979), records only, nothing measured): two entries were stale and are corrected.** `BACKEND-GATE-CUDA-SGLANG` moves `BLOCKED` to `PARTIAL`: its recorded blocker `SERVE-ASYNC-LLM` is discharged, so the reason "HTTP TTFT/ITL cannot be measured honestly yet" is retracted. The production server streams incrementally over `AsyncLLM` (`serving_completion.h:9`, `api_server.cpp:971-981`), the benchmark harness ENFORCES it rather than assuming it (`tools/bench/run_serve_low.py:296-310` refuses `first_chunk_s >= total_s`), and the run above demonstrates it: c16 mean TTFT 2980 ms against mean ITL 154.4 ms over 128 tokens means first byte preceded completion by about twenty seconds, which a buffered server cannot do. That run is NOT voided by [#931](https://github.com/mudler/vllm.cpp/issues/931), because the keepalive frame needs 15 s of silence and the worst observed p99 TTFT was 7220 ms. The SGLang oracle also moves to `gateable = yes`: `.agents/oracles/sglang.md` still said "no SGLang run has been recorded on this project's hardware" two and a half weeks after this measurement, and `docs/BENCHMARKS.md` still said "Never ran". This discharges the SGLang third of [#647](https://github.com/mudler/vllm.cpp/issues/647). A new `BACKEND-GATE-CUDA-LLAMACPP` row is filed `INVENTORIED` for the llama.cpp-on-current-CUDA arm, which had no owner. Campaign spec: `.agents/specs/bench-qwen38-27b-four-way.md` | +**CUDA graph capture takes BREAK POINTS** as of `ENG-CUDAGRAPH-BREAK` W1 +(`ACTIVE`, 2026-08-18, [spec](../.agents/specs/eng-cudagraph-break.md), #1192, +parent #1163). A forward can be captured as a SEQUENCE of segments split at +break points, so a forward containing a host-dependent operation is still +graphed except at that operation instead of falling out to eager for the whole +step. + +`vt::BreakableGraph`, `vt::GraphCaptureScope` and `vt::GraphBreak` +(`include/vt/breakable_graph.h`) replace the nine hand-rolled per-model drivers +as those migrate. **Coverage and correctness, not speed**: no throughput gate is +declared, because our prefill has no launch bubbles to collapse (3.8% host idle, +GPU-busy above 96%) and decode already banked its launch-overhead win. + +The boundary is vLLM's — its v1 default splits at `splitting_ops`, defaulted to +the attention family — and only the registration form is SGLang's, because vLLM +gets its split from Dynamo and FX and we have no compiler. W1 registers ONE +break point, at the dense attention entry of `Qwen3ForCausalLM`. + +The stage's exit criterion was measured rather than assumed: ending a capture +and BEGINNING A NEW ONE on the same stream mid-forward, with eager work between, +is legal under the thread-local capture mode our CUDA backend uses (`orin:gpu0` +via an `rc` lease, driver 12060, three replays with fresh inputs, zero +mismatches). + +Gates: `tests/vt/test_breakable_graph.cpp` ports SGLang's unit suite case for +case, with its arithmetic chains and its post-replay assertions intact (24 +cases, 163 assertions); `tests/vllm/models/test_qwen3_break_point.cpp` drives +the production forward with a scope open, counts one segment per layer plus one, +and compares the logits BIT FOR BIT against the unscoped forward (500 values, 0 +differing). + +The seam ENFORCES what it used to document. A destination-carrying break point +takes a `vt::BreakSlot` whose storage the seam owns, because the following +segment bakes that address and a caller's local dies first; a destination with no +`CopyOutput` overload is a compile error naming the type rather than a silent +drop into the no-writeback path; a second capture scope on one thread and a +re-entered container are refused; and a forward that throws OUT OF the scope +mid-capture leaves no partial graph reporting itself as captured. Interleaved +replay, the `VLLM_CPP_CUDAGRAPH=0` kill switch and each refusal above are gated, +every one proven by a mutation that reds them. + +Two break points in one capture writing through ONE destination are refused as +well, at registration. `BreakSlot` closed the lifetime half of that rule and +left the aliasing half writable: both replay closures bound to the same address, +so the earlier writeback was overwritten. One shape still escapes, and the spec +names it and assigns it to W2 — an exception CAUGHT INSIDE the scope leaves the +rest of the forward uncaptured while `captured()` stays true, because nothing is +unwinding at scope exit for the drain to see. + +**Not yet entered from a production step.** No driver opens a capture scope +until W2 migrates `Qwen3DenseDecodeGraph`; the break point itself runs on every +forward and takes the pass-through arm. The spec's `## Owed` names that with its +owner, along with the auxiliary-stream auto-join, the ROCm and Tenstorrent arms, +and GPU bit-exactness over more than one replay. + ## Speculative decoding **MTP draft DEPTH is configurable** as of `SPEC-MTP-K-GT-1` (`ACTIVE`, diff --git a/docs/USAGE.md b/docs/USAGE.md index 6aa5b1e04..1be6fccc3 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -4086,6 +4086,42 @@ on a GB10, because host and device share one pool. `cudaMemGetInfo` answers honestly, and its `total` is EXACTLY `/proc/meminfo MemTotal` (125442340 kB) times 1024. Do not size this from `nvidia-smi`. +## Turning CUDA graph capture off, including the break seam + +`VLLM_CPP_CUDAGRAPH=0` disables CUDA graph capture. It always did for the six +batched decode drivers that each read it, and as of `ENG-CUDAGRAPH-BREAK` W1 +(#1192) it is also the switch the shared break-point seam reads, once per +process, into a function-local static — so a process is in exactly one lane for +its whole life and nothing can toggle it mid-run. + +With capture off, or on a backend that reports no capture support (Vulkan, +Metal, and the CPU backend), a `vt::GraphCaptureScope` is INERT: it captures +nothing, every `vt::GraphBreak` inside it calls its function and returns, and +the forward runs eager exactly as before. That path is byte-identical to the +non-capturing forward and makes zero backend calls, which is what makes each +migration stage reversible. + +Nothing about this is new configuration to learn: there is no new flag, no new +config key and no new command. The seam is a library surface +(`include/vt/breakable_graph.h`), and W1 registers one break point at the dense +attention entry of `Qwen3ForCausalLM`. No production step opens a capture scope +yet — that arrives when the decode drivers migrate onto the seam — so today the +switch changes nothing about the break point beyond what it already changed +about the decode graphs. + +Building it needs no option. `src/vt/breakable_graph.cpp` is part of the core +`vllm` library on every platform, because the seam is backend-agnostic and asks +nothing new of any backend. + +The switch is GATED, and it is gated in a child process, because it is read once +per process into a function-local static and no test in a running process can +toggle it. `tests/vt/test_breakable_graph.cpp` re-executes itself with +`VLLM_CPP_CUDAGRAPH=0` and requires the inert behaviour on a backend that CAN +capture — the arm that proves the switch itself is what turns capture off, rather +than the backend's own lack of support. Asserting the backend arm instead +substitutes a different condition, and dropping the switch from the seam left the +whole suite green. + ## SSE keepalives on long prefill Async chat/completion streams can emit SSE **comment** frames (`:\n\n`) while diff --git a/include/vllm/model_executor/models/dense_attn_graph_break.h b/include/vllm/model_executor/models/dense_attn_graph_break.h new file mode 100644 index 000000000..e98315780 --- /dev/null +++ b/include/vllm/model_executor/models/dense_attn_graph_break.h @@ -0,0 +1,40 @@ +// The `vt::CopyOutput` customization point for a break point whose destination +// is a POOLED DEVICE BUFFER — the port of `_copy_output` +// (`breakable_cuda_graph.py:172-201` @ SGLang pin `f63458b5be`). +// +// Row ENG-CUDAGRAPH-BREAK W1, issue #1192, parent #1163. +// +// It lives in a header rather than beside its one call site for two reasons. +// Argument-dependent lookup has to find it wherever a dense model registers an +// attention break point, and W2 through W5 migrate eight more drivers that share +// this exact destination type. And a customization point that only one +// translation unit can see is one nobody can ASSERT is visible: the G2 gate +// static-asserts `vt::detail::HasCopyOutput>` against this +// declaration, which is what turns "the writeback branch is selected" from a +// claim into a compile-time fact. +#pragma once + +#include + +#include "vllm/model_executor/models/dense_device_glue.h" // DBuf +#include "vt/backend.h" +#include "vt/breakable_graph.h" + +namespace vllm { +namespace dense_attn { + +// Why it cannot be a rebind. On replay N the attention call returns a FRESH +// `DBuf` from the device pool, whose address is not the one the FOLLOWING +// segment baked at capture time. Rebinding the destination would leave that +// segment reading capture-time data forever while the break wrote elsewhere — +// wrong numerics, not a fault, and invisible to `compute-sanitizer` (spec D9). +// So the bytes are copied INTO the destination and the destination's address +// never moves. +inline void CopyOutput(vt::Backend& b, vt::Queue& q, std::optional& dst, + const std::optional& src) { + if (!dst.has_value() || !src.has_value()) return; + vt::CopyOutput(b, q, dst->t(), src->t()); +} + +} // namespace dense_attn +} // namespace vllm diff --git a/include/vt/breakable_graph.h b/include/vt/breakable_graph.h new file mode 100644 index 000000000..0484c7b12 --- /dev/null +++ b/include/vt/breakable_graph.h @@ -0,0 +1,447 @@ +// PORT of SGLang's breakable CUDA graph (BCG) at pin `f63458b5be`, +// `python/sglang/srt/model_executor/runner_backend_utils/breakable_cuda_graph/breakable_cuda_graph.py` +// (374 lines; every bare `:N` below is a line in that file). +// +// Row ENG-CUDAGRAPH-BREAK, spec `.agents/specs/eng-cudagraph-break.md`, +// issue #1192 (W1), parent #1163. +// +// WHAT THIS IS. A forward is captured as a SEQUENCE of graph segments split at +// break points, so a forward containing a host-dependent operation is still +// graphed except at that operation, instead of running eager for the whole step. +// +// WHERE THE BOUNDARY COMES FROM. vLLM, the primary oracle, at pin `5559679229`. +// `CUDAGraphMode.PIECEWISE` (`vllm/config/compilation.py:60-63`) keeps "the +// cudagraph incompatible ops (i.e. some attention ops) outside the cudagraph" +// (`:608-635`), and the split points are `splitting_ops` (`:517`), defaulted at +// `:1145` to the attention family listed at `:764-772`. That boundary transfers. +// Its MECHANISM does not: vLLM gets the split from Dynamo and FX, and we have no +// compiler. SGLang reaches the same coverage with no compiler, so SGLang is the +// secondary oracle for the CONSTRUCTION only, per `.agents/oracles/sglang.md`. +// +// WHAT THIS IS NOT. This is a COVERAGE and CORRECTNESS seam, not a speed lever, +// and no throughput claim is attached to it anywhere. GB10 measured prefill idle +// between launches at 3.8% with GPU-busy above 96%, and the 27B prefill gap at +// 92.5% non-GEMM glue GPU work; decode is already captured and already banked +// its launch-overhead win. A speed claim from this seam is admissible only after +// naming a path that is BOTH currently eager AND currently host-bound. +#pragma once + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "vt/backend.h" +#include "vt/dtype.h" // VT_CHECK +#include "vt/tensor.h" + +namespace vt { + +// --------------------------------------------------------------------------- +// The writeback contract. Port of `_copy_output` (`:172-201`). +// --------------------------------------------------------------------------- +// +// THIS IS LOAD-BEARING, NOT DECORATOR SUGAR (spec `## Risks/decisions` D9). On +// replay N the eager break operation returns a FRESH allocation, whose address +// is NOT the one the following segment baked at capture time. A container that +// replayed the caller's raw function and discarded its result would leave +// segment `i+1` reading capture-time data forever while the break function wrote +// somewhere else. The failure is WRONG NUMERICS, not a fault, which is the same +// detection class a clean `compute-sanitizer` run cannot see. +// +// Overload `CopyOutput` for a break function's output type; it is found by +// argument-dependent lookup. The tensor overload copies IN PLACE and preserves +// the destination's address, mirroring upstream's `assertIs(result, dst)` +// (`:190`). A type with NO overload cannot use the destination form at all: the +// static assertion in `GraphBreak` below refuses it by name, and the documented +// non-copyable fallback (`:201`) has to be asked for with `vt::NoWriteback{}`. +void CopyOutput(Backend& b, Queue& q, Tensor& dst, const Tensor& src); +void CopyOutput(Backend& b, Queue& q, std::vector& dst, + const std::vector& src); +void CopyOutput(Backend& b, Queue& q, std::map& dst, + const std::map& src); + +namespace detail { +template +struct HasCopyOutput : std::false_type {}; +template +struct HasCopyOutput< + T, std::void_t(), std::declval(), + std::declval(), std::declval()))>> + : std::true_type {}; + +// Bumps GraphBreakStats::break_points_reached. Declared here because the +// GraphBreak templates are header-defined and the counter is not. +void CountBreakPoint(); +} // namespace detail + +// The EXPLICIT opt-in to upstream's non-copyable fallback (`:201`), where +// `_copy_output` returns `src` and the destination is left untouched. Passing it +// says "this break has no device output and I know the seam writes nothing +// back". Without it a destination type with no `CopyOutput` overload is a +// COMPILE error rather than a silent downgrade to the wrong-numerics path, which +// is the whole reason the tag exists: `HasCopyOutput` as a plain `if constexpr` +// default made a misspelled, moved or shadowed overload compile clean and be +// wrong only on replay N. +struct NoWriteback {}; + +// --------------------------------------------------------------------------- +// vt::BreakSlot — the destination of a destination-carrying break point. +// Port of `captured_output = _weak_ref_if_tensor(output)` (`:156-169`, `:225`). +// --------------------------------------------------------------------------- +// +// WHY THIS TYPE EXISTS AND A REFERENCE DOES NOT. The destination must outlive +// the `BreakableGraph`, because the following segment bakes its address at +// capture time and every replay writes through that address. Upstream gets that +// for free: `captured_output` is held BY VALUE in the replay closure and its +// storage is pinned by the segment graphs' mempool. A C++ seam taking a bare +// `Out&` gets the opposite — the natural call site declares the destination as a +// local of the function containing the break, which dies on the very next +// `return`, and the first such site did exactly that. A comment at the +// declaration is not a mechanism; every caller re-derives it and one of them is +// wrong. +// +// So the seam takes a slot whose storage it can OWN. On the capturing path the +// value is moved into a heap cell before the eager call produces it, and the +// replay closure holds a `shared_ptr` to that cell, so the destination lives +// exactly as long as the break function that writes into it — by construction, +// not by convention. +// +// On the NON-capturing path (every production step until a driver opens a +// scope) the cell is never allocated and the value lives inline in the slot, so +// a pass-through forward allocates nothing extra and is byte-identical to the +// plain local it replaces. +// +// Read the value through `*slot` or `slot->`, never through a pointer cached +// before the break: pinning moves the value, exactly as the destination form's +// `out = fn()` rebinds it (`:222-227`). +template +class BreakSlot { + public: + BreakSlot() = default; + explicit BreakSlot(T v) : inline_(std::move(v)) {} + + T& operator*() { return cell_ ? *cell_ : inline_; } + const T& operator*() const { return cell_ ? *cell_ : inline_; } + T* operator->() { return cell_ ? cell_.get() : &inline_; } + const T* operator->() const { return cell_ ? cell_.get() : &inline_; } + + // True once the slot's storage is owned by the seam rather than by this + // object, which is the state the writeback contract requires. + bool pinned() const { return static_cast(cell_); } + + // Called by GraphBreak ONLY, and only on the capturing path, BEFORE the eager + // result is produced. Moving an as-yet-unwritten value is free, and from this + // point the address is stable for the life of the closure that holds the cell. + const std::shared_ptr& PinForCapture() { + if (!cell_) cell_ = std::make_shared(std::move(inline_)); + return cell_; + } + + private: + T inline_{}; + std::shared_ptr cell_; +}; + +// --------------------------------------------------------------------------- +// G3 observability (spec `## Gates` G3). +// --------------------------------------------------------------------------- +// Without a count there is no way to tell a two-segment capture from a fully +// eager step, and "the graph ran" is exactly the claim a broken instrument +// fabricates. Each existing driver carries a private `replay_count()` and no +// segment count exists at all. +struct GraphBreakStats { + int64_t break_points_reached = 0; // GraphBreak call sites executed, capturing or not + int64_t segments_captured = 0; + int64_t breaks_registered = 0; + int64_t replays = 0; // BreakableGraph::Replay calls +}; +GraphBreakStats GetGraphBreakStats(); +void ResetGraphBreakStats(); + +// The one kill switch, read ONCE per process into a function-local static. +// Today six drivers each read `VLLM_CPP_CUDAGRAPH` for themselves and the three +// single-shape drivers invented their own switch instead, so there is no one +// switch that turns capture off; this is that switch for the seam. +bool GraphCaptureEnabled(); + +class GraphCaptureScope; + +// --------------------------------------------------------------------------- +// vt::BreakableGraph — the segment container. Port of `BreakableCUDAGraph` +// (`:246-274`); fields `:251-252`, interleaved replay `:255-264`. +// --------------------------------------------------------------------------- +// +// Holds one opaque handle per segment (each from `Backend::EndCaptureGraph`) and +// one seam-built replay closure per break. The invariant is +// `segment_count() == break_count() + 1` for any capture that completed, and it +// is ASSERTED in two places rather than documented: `GraphCaptureScope` refuses +// to open on a container that already holds a capture, and `Replay` refuses a +// container whose counts disagree. A graph with `break_count() == segment_count()` +// would otherwise replay with its last break silently dropped. +// +// A segment handle is treated as OPAQUE and every acquisition and release goes +// through `Backend::EndCaptureGraph` and `Backend::DestroyGraph`, so +// ENG-CUDAGRAPH-DEDUP (#1162) can interpose at the backend without editing this +// container (spec `## Risks/decisions` D4). +class BreakableGraph { + public: + BreakableGraph() = default; + ~BreakableGraph(); + BreakableGraph(const BreakableGraph&) = delete; + BreakableGraph& operator=(const BreakableGraph&) = delete; + + // segment 0, break 0, segment 1, break 1, ..., segment N. Interleaved, never + // batched: the break between two segments runs BETWEEN their replays. + void Replay(Queue& q); + + size_t segment_count() const { return segments_.size(); } + size_t break_count() const { return break_fns_.size(); } + bool captured() const { return !segments_.empty(); } + int64_t replay_count() const { return replays_; } + + // Releases every segment through Backend::DestroyGraph, clears the breaks, and + // returns the container to its as-constructed state — replay count included, + // because a stale count on a reset container is a number that describes a + // graph that no longer exists. + void Reset(); + + private: + friend class GraphCaptureScope; + Backend* backend_ = nullptr; + std::vector segments_; + std::vector> break_fns_; + int64_t replays_ = 0; +}; + +// --------------------------------------------------------------------------- +// vt::GraphCaptureScope — the RAII capture scope. Port of +// `BreakableCUDAGraphCapture` (`:277-367`); `__enter__` `:309-320`, `__exit__` +// `:322-333`, segment open and close `:335-350` and `:352-367`. +// --------------------------------------------------------------------------- +// +// The scope pointer is THREAD-LOCAL, not global, for the same reason upstream +// chose a `ContextVar` (`:63`) and our CUDA backend chose +// `cudaStreamCaptureModeThreadLocal` (`src/vt/cuda/cuda_backend.cu:204-206`): +// capture is a property of one thread's stream, and a process-wide flag would +// make an unrelated thread's forward observe a capture it is not part of. +// +// INERT when the backend cannot capture (`SupportsGraphCapture()` false, which +// is Vulkan `vulkan_backend.cpp:16` and Metal `metal_backend.mm:13`) or when +// `VLLM_CPP_CUDAGRAPH=0`. An inert scope captures nothing, every `GraphBreak` +// inside it is a pass-through, and the forward runs eager exactly as today. +// +// THREE THINGS IT REFUSES, all `VT_CHECK`. The first two are refused at +// construction, before any backend call: +// +// * A NESTED scope on the same thread. `BeginCapture` on a stream that is +// already capturing is `CUDA_ERROR_ILLEGAL_STATE` on a real backend, and the +// sequence a nested pair traces is not a legal capture on any of them. +// * A container that already holds a capture. Re-entering appends to it, and +// the result has `break_count() == segment_count()`, whose last break +// `Replay` would silently drop. Call `Reset()` first. +// * TWO BREAK POINTS SHARING ONE DESTINATION, refused at registration. Both +// replay closures would write through the same address, so the earlier +// writeback is overwritten and any segment that baked the earlier +// destination reads the later break's data. See lifetime rule 1 below; +// this one throws from `AppendBreak`, mid-capture, and the drain below is +// what turns it into no capture at all. +// +// AND ONE THING IT DRAINS. If the scope exits by EXCEPTION — a break function +// that throws, or ordinary model code between two break points — the stream is +// taken out of capture and the half-built container is `Reset()`. A partial +// capture must never be reported as `captured()`, because a partial capture is +// replayable, and replaying half a forward is silently wrong numerics. +// +// THE DRAIN'S LIMIT, stated because it is not obvious. It compares +// `std::uncaught_exceptions()` against the depth at entry, so it sees an +// exception that is PROPAGATING at scope exit. An exception CAUGHT INSIDE the +// scope leaves `segment_open_ == false` with the rest of the forward +// uncaptured, and nothing is unwinding at scope exit, so `captured()` stays +// true over a forward missing its tail. Do not `try` around a break point +// inside a capture scope. Owed to W2 in the spec's `## Owed`. +class GraphCaptureScope { + public: + GraphCaptureScope(Backend& b, Queue& q, BreakableGraph& out); + ~GraphCaptureScope(); + GraphCaptureScope(const GraphCaptureScope&) = delete; + GraphCaptureScope& operator=(const GraphCaptureScope&) = delete; + + bool active() const { return active_; } + static GraphCaptureScope* Current(); + + Backend& backend() const { return *b_; } + Queue& queue() const { return *q_; } + + // Called by GraphBreak. Closes the segment that captured up to the break + // point, and opens a fresh one for the remainder of the forward. + void EndSegment(); + void BeginSegment(); + // `destination` is the CELL this break's replay closure writes through, or + // `nullptr` for a form whose replay writes back nowhere. It is a REQUIRED + // parameter, not an optional one, so a break cannot be registered without + // stating where it lands: two breaks naming one destination are refused here + // (see the aliasing paragraph at lifetime rule 1). + // + // The CELL and not the SLOT, because a slot's ADDRESS is not its identity. + // The production site declares its slot as a local of `RunLayer`, so every + // layer's slot occupies the SAME stack address and slot-address identity + // refuses the correct program on layer 2. The cell is heap storage the + // replay closures keep alive for the whole capture, so two live cells cannot + // share an address — and the cell is what actually aliases. + void AppendBreak(std::function fn, const void* destination); + + private: + Backend* b_; + Queue* q_; + BreakableGraph* g_; + bool active_ = false; + bool segment_open_ = false; + int uncaught_on_entry_ = 0; + // One entry per registered break that names a destination, for the life of + // this capture. A capture has as many breaks as the model has split points, + // so a linear scan here is cheaper than the map that would replace it. + std::vector destinations_; +}; + +// --------------------------------------------------------------------------- +// vt::GraphBreak — the break point. Port of `eager_on_graph` (`:204-243`), +// whose wrapper body is `:209-241`. +// --------------------------------------------------------------------------- +// +// The SITE IS THE REGISTRATION. vLLM registers `splitting_ops` by operation NAME +// because it has an FX graph to match names against; we have neither, so our +// equivalent is one line at the break site, exactly as SGLang's is +// (`layers/radix_attention.py:256`, `layers/radix_linear_attention.py:159`). +// +// OUTSIDE a capture scope every form calls the function and returns, so a +// non-capturing forward is byte-identical to today and makes ZERO backend calls. +// +// THREE LIFETIME RULES. Rule 1 used to be a comment; it is now the type system, +// because the one site that had to obey it did not. Rules 2 and 3 remain rules a +// new model author is least likely to know, and both produce wrong numbers +// rather than a fault: +// +// 1. The DESTINATION must outlive the `BreakableGraph`, no replay may +// reallocate it, and no two break points may share it. ENFORCED in two +// different ways, because the rule has two halves and only one of them is +// a lifetime: +// * LIFETIME, by the TYPE. The destination form takes a +// `vt::BreakSlot` and the seam owns its storage for the life of the +// replay closure. There is no overload taking a bare reference, so +// handing the break a destination that dies first is not writable. +// * ALIASING, by a REFUSAL. One slot reused for two break points in one +// capture is still writable — `PinForCapture` hands back the cell it +// already made — and it binds BOTH replay closures to the SAME +// address, so the earlier writeback is overwritten and any segment +// that baked the earlier destination reads the later break's data. +// `AppendBreak` throws on the second registration naming a cell +// already registered in this capture. The constraint is per CAPTURE, +// not per slot for life: a later capture bakes its own addresses and +// may reuse the slot. The NON-COPYABLE fallback is outside this +// refusal and says so at its own declaration. +// 2. A closure must capture DEVICE POINTERS or `Tensor` views that are stable +// across replays, never a reference to a host temporary that dies with the +// capturing call. +// 3. Capture bakes the HOST SOURCE ADDRESS of an upload as well as the device +// destination (spec D2, `.agents/specs/decode-graph-scratch-uaf-2026-07-18.md`). +// A break function that reads a host vector which is reallocated between +// capture and replay reads freed host memory on replay. A clean +// `compute-sanitizer` run is NOT evidence against any of these: the +// recorded incident reproduced 5 of 5 times in normal operation and 0 of 2 +// under memcheck. + +// The BARE marker. Splits the segment and runs nothing — the form a model uses +// when the host-dependent work already ran outside the forward. Port of +// `break_graph` (`:370-374`), an empty body under the decorator. +void GraphBreak(); + +// The IN-PLACE form. `fn` writes into a PERSISTENT buffer the model owns and +// that no replay reallocates, and returns nothing. The writeback contract is +// vacuous for it, because there is no fresh allocation to write back. +template ()())>>> +void GraphBreak(Fn&& fn) { + GraphCaptureScope* s = GraphCaptureScope::Current(); + if (s == nullptr || !s->active()) { + fn(); // pass-through: byte-identical to today, zero backend calls + detail::CountBreakPoint(); + return; + } + s->EndSegment(); + fn(); // run once eagerly, so the outputs hold real data (`:222-223`) + s->AppendBreak([f = std::forward(fn)]() mutable { f(); }, nullptr); + s->BeginSegment(); + detail::CountBreakPoint(); +} + +// The DESTINATION form. `fn` returns a fresh result and the seam copies it into +// the slot on every replay. Direct port of `replay_fn` plus `_copy_output` +// (`:231-235`, `:172-201`), with `BreakSlot` playing the part +// `_weak_ref_if_tensor` (`:156-169`) plays upstream: the capture-time +// destination that outlives the eager call. +// +// The static assertion is the one MEDIUM the fallback branch used to hide. When +// selecting the fallback was a silent `if constexpr` default, renaming or moving +// a `CopyOutput` overload left both suites green while the break dropped into +// the branch whose own comment says such a break must use the in-place form — +// the D9 wrong-numerics path, reached by a typo. +template +void GraphBreak(Fn&& fn, BreakSlot& out) { + static_assert(detail::HasCopyOutput::value, + "vt::GraphBreak destination form: no CopyOutput(Backend&, Queue&, Out&, " + "const Out&) is visible for this destination type. Declare one (found by " + "argument-dependent lookup) so every replay writes back into the address " + "the following segment baked, or use the in-place form GraphBreak(fn), or " + "ask for the non-copyable fallback explicitly with vt::NoWriteback{}."); + GraphCaptureScope* s = GraphCaptureScope::Current(); + if (s == nullptr || !s->active()) { + *out = fn(); // pass-through: a move, never a copy, and the cell stays unpinned + detail::CountBreakPoint(); + return; + } + s->EndSegment(); + // Pin BEFORE the eager call, so the address the following segment bakes is the + // one the replay closure owns. `:222-227`: the capture-time destination IS the + // first eager result. + std::shared_ptr cell = out.PinForCapture(); + *cell = fn(); + Backend* b = &s->backend(); + Queue* q = &s->queue(); + s->AppendBreak( + [b, q, cell, f = std::forward(fn)]() mutable { CopyOutput(*b, *q, *cell, f()); }, + cell.get()); + s->BeginSegment(); + detail::CountBreakPoint(); +} + +// The DESTINATION form's non-copyable fallback (`:201`), asked for by name. +// Upstream returns `src` and leaves the destination untouched; so does this. The +// slot then holds the CAPTURE-TIME value forever, which is correct only when +// nothing downstream of the break reads it on replay. +// +// It registers NO destination, and that is not an oversight: this form pins no +// cell and its replay closure writes back nowhere, so the aliasing refusal has +// nothing to compare. Two of these sharing one slot still overwrite each other +// at CAPTURE time, which is a narrower hazard than the writeback one and is not +// gated here. +template +void GraphBreak(Fn&& fn, BreakSlot& out, NoWriteback) { + GraphCaptureScope* s = GraphCaptureScope::Current(); + if (s == nullptr || !s->active()) { + *out = fn(); + detail::CountBreakPoint(); + return; + } + s->EndSegment(); + *out = fn(); + s->AppendBreak([f = std::forward(fn)]() mutable { (void)f(); }, nullptr); + s->BeginSegment(); + detail::CountBreakPoint(); +} + +} // namespace vt diff --git a/scripts/probe_cudagraph_rebegin.c b/scripts/probe_cudagraph_rebegin.c new file mode 100644 index 000000000..d8c76a078 --- /dev/null +++ b/scripts/probe_cudagraph_rebegin.c @@ -0,0 +1,246 @@ +/* ENG-CUDAGRAPH-BREAK W1 exit criterion. + * + * Question: does CUDA permit cuStreamEndCapture followed by cuStreamBeginCapture + * on the SAME stream mid-forward, with EAGER work between them, under the + * thread-local capture mode our CUDA backend uses (cuda_backend.cu:204-206)? + * + * Self-contained: binds the CUDA driver API through dlopen("libcuda.so"), so it + * needs no CUDA toolkit and no headers. Uses only capturable memcpy/memset + * stream operations, so it needs no nvcc-compiled kernel. + * + * The break function is genuinely host-dependent: it copies the mid buffer to + * the host, computes v*2+3 on the CPU, and copies it back. That is the exact + * class of operation a break point exists for. + */ +#include +#include +#include +#include + +typedef int CUresult; +typedef int CUdevice; +typedef unsigned long long CUdeviceptr; +typedef void *CUcontext; +typedef void *CUstream; +typedef void *CUgraph; +typedef void *CUgraphExec; + +static void *H; + +static CUresult (*p_cuInit)(unsigned); +static CUresult (*p_cuDriverGetVersion)(int *); +static CUresult (*p_cuDeviceGet)(CUdevice *, int); +static CUresult (*p_cuDeviceGetName)(char *, int, CUdevice); +static CUresult (*p_cuCtxCreate)(CUcontext *, unsigned, CUdevice); +static CUresult (*p_cuStreamCreate)(CUstream *, unsigned); +static CUresult (*p_cuMemAlloc)(CUdeviceptr *, size_t); +static CUresult (*p_cuMemcpyHtoD)(CUdeviceptr, const void *, size_t); +static CUresult (*p_cuMemcpyDtoH)(void *, CUdeviceptr, size_t); +static CUresult (*p_cuMemcpyDtoDAsync)(CUdeviceptr, CUdeviceptr, size_t, CUstream); +static CUresult (*p_cuMemcpyDtoHAsync)(void *, CUdeviceptr, size_t, CUstream); +static CUresult (*p_cuMemcpyHtoDAsync)(CUdeviceptr, const void *, size_t, CUstream); +static CUresult (*p_cuMemsetD32Async)(CUdeviceptr, unsigned, size_t, CUstream); +static CUresult (*p_cuStreamBeginCapture)(CUstream, int); +static CUresult (*p_cuStreamEndCapture)(CUstream, CUgraph *); +static CUresult (*p_cuStreamIsCapturing)(CUstream, int *); +static CUresult (*p_cuGraphInstantiate)(CUgraphExec *, CUgraph, unsigned long long); +static CUresult (*p_cuGraphLaunch)(CUgraphExec, CUstream); +static CUresult (*p_cuGraphDestroy)(CUgraph); +static CUresult (*p_cuGraphExecDestroy)(CUgraphExec); +static CUresult (*p_cuStreamSynchronize)(CUstream); +static CUresult (*p_cuGetErrorName)(CUresult, const char **); + + +/* Bind by EXACT symbol name. libcuda exports several versions of the same + entry point and they are NOT interchangeable: the legacy v1 + cuStreamBeginCapture takes no capture mode, the v1 memcpys are not + capture-aware, and cuCtxCreate_v3 takes two extra parameters. Guessing a + suffix binds a different function with the same name, which fails in a way + that reads like a CUDA verdict rather than a probe bug. */ +static void *sym(const char *n) { + void *f = dlsym(H, n); + if (!f) { printf("FAIL dlsym %s\n", n); } + return f; +} + +static const char *errname(CUresult r) { + const char *s = "?"; + if (p_cuGetErrorName) p_cuGetErrorName(r, &s); + return s ? s : "?"; +} + +static int fails = 0; +#define CK(expr) \ + do { \ + CUresult _r = (expr); \ + if (_r != 0) { \ + printf("FAIL %s -> %d (%s)\n", #expr, _r, errname(_r)); \ + fails++; \ + } \ + } while (0) + +#define N 256 +#define BYTES (N * (int)sizeof(int)) + +/* capture mode 1 == CU_STREAM_CAPTURE_MODE_THREAD_LOCAL, what cuda_backend.cu uses */ +#define MODE_THREAD_LOCAL 1 + +static CUstream S; +static CUdeviceptr dA, dB, dC; +static int hostbuf[N]; + +/* The break function: host-dependent, runs eagerly on the same stream. */ +static void break_fn(void) { + int i; + CK(p_cuMemcpyDtoHAsync(hostbuf, dB, BYTES, S)); + CK(p_cuStreamSynchronize(S)); + for (i = 0; i < N; i++) hostbuf[i] = hostbuf[i] * 2 + 3; + CK(p_cuMemcpyHtoDAsync(dB, hostbuf, BYTES, S)); + CK(p_cuStreamSynchronize(S)); +} + +int main(void) { + int drv = 0, i, rep, cap = -1; + char name[256]; + CUdevice dev; + CUcontext ctx; + CUgraph g0 = 0, g1 = 0, g2 = 0; + CUgraphExec e0 = 0, e1 = 0, e2 = 0; + int in[N], out[N]; + + H = dlopen("libcuda.so.1", RTLD_NOW); + if (!H) H = dlopen("libcuda.so", RTLD_NOW); + if (!H) { printf("FAIL dlopen libcuda: %s\n", dlerror()); return 2; } + + p_cuInit = sym("cuInit"); + p_cuDriverGetVersion = sym("cuDriverGetVersion"); + p_cuDeviceGet = sym("cuDeviceGet"); + p_cuDeviceGetName = sym("cuDeviceGetName"); + p_cuCtxCreate = sym("cuCtxCreate_v2"); + p_cuStreamCreate = sym("cuStreamCreate"); + p_cuMemAlloc = sym("cuMemAlloc_v2"); + p_cuMemcpyHtoD = sym("cuMemcpyHtoD_v2"); + p_cuMemcpyDtoH = sym("cuMemcpyDtoH_v2"); + p_cuMemcpyDtoDAsync = sym("cuMemcpyDtoDAsync_v2"); + p_cuMemcpyDtoHAsync = sym("cuMemcpyDtoHAsync_v2"); + p_cuMemcpyHtoDAsync = sym("cuMemcpyHtoDAsync_v2"); + p_cuMemsetD32Async = sym("cuMemsetD32Async"); + /* v2 is the one that takes the capture MODE; v1 takes only the stream. */ + p_cuStreamBeginCapture = sym("cuStreamBeginCapture_v2"); + p_cuStreamEndCapture = sym("cuStreamEndCapture"); + p_cuStreamIsCapturing = sym("cuStreamIsCapturing"); + p_cuGraphInstantiate = sym("cuGraphInstantiateWithFlags"); + p_cuGraphLaunch = sym("cuGraphLaunch"); + p_cuGraphDestroy = sym("cuGraphDestroy"); + p_cuGraphExecDestroy = sym("cuGraphExecDestroy"); + p_cuStreamSynchronize = sym("cuStreamSynchronize"); + p_cuGetErrorName = sym("cuGetErrorName"); + + if (!p_cuInit || !p_cuStreamBeginCapture || !p_cuStreamEndCapture || + !p_cuGraphInstantiate || !p_cuGraphLaunch || !p_cuMemcpyDtoDAsync) { + printf("FAIL symbol bind: init=%p begin=%p end=%p inst=%p launch=%p d2d=%p\n", + (void *)p_cuInit, (void *)p_cuStreamBeginCapture, + (void *)p_cuStreamEndCapture, (void *)p_cuGraphInstantiate, + (void *)p_cuGraphLaunch, (void *)p_cuMemcpyDtoDAsync); + return 2; + } + + CK(p_cuInit(0)); + if (fails) { printf("VERDICT: NO_GPU (cuInit failed)\n"); return 3; } + CK(p_cuDriverGetVersion(&drv)); + CK(p_cuDeviceGet(&dev, 0)); + name[0] = 0; + CK(p_cuDeviceGetName(name, sizeof name, dev)); + printf("driver_version=%d device=%s\n", drv, name); + CK(p_cuCtxCreate(&ctx, 0, dev)); + CK(p_cuStreamCreate(&S, 1 /* NON_BLOCKING */)); + CK(p_cuMemAlloc(&dA, BYTES)); + CK(p_cuMemAlloc(&dB, BYTES)); + CK(p_cuMemAlloc(&dC, BYTES)); + if (fails) { printf("VERDICT: SETUP_FAILED\n"); return 3; } + + for (i = 0; i < N; i++) in[i] = i + 1; + CK(p_cuMemcpyHtoD(dA, in, BYTES)); + + /* ---- CAPTURE PHASE: segment 0, eager break, segment 1 on the SAME stream ---- */ + printf("\n[capture] segment 0: begin capture (thread-local mode)\n"); + CK(p_cuStreamBeginCapture(S, MODE_THREAD_LOCAL)); + CK(p_cuStreamIsCapturing(S, &cap)); + printf("[capture] is_capturing during segment 0 = %d (expect 1/2 = active)\n", cap); + CK(p_cuMemcpyDtoDAsync(dB, dA, BYTES, S)); + CK(p_cuStreamEndCapture(S, &g0)); + printf("[capture] segment 0 ended, graph=%p\n", (void *)g0); + CK(p_cuStreamIsCapturing(S, &cap)); + printf("[capture] is_capturing after end 0 = %d (expect 0 = none)\n", cap); + CK(p_cuGraphInstantiate(&e0, g0, 0)); + + printf("[capture] BREAK: eager host-dependent work on the SAME stream\n"); + break_fn(); + if (fails) { printf("VERDICT: EAGER_BETWEEN_SEGMENTS_FAILED\n"); return 3; } + + printf("[capture] segment 1: RE-BEGIN capture on the same stream <-- exit criterion\n"); + { + CUresult r = p_cuStreamBeginCapture(S, MODE_THREAD_LOCAL); + if (r != 0) { + printf("FAIL re-begin -> %d (%s)\n", r, errname(r)); + printf("VERDICT: REBEGIN_REFUSED\n"); + return 3; + } + } + CK(p_cuStreamIsCapturing(S, &cap)); + printf("[capture] is_capturing during segment 1 = %d (expect active)\n", cap); + CK(p_cuMemcpyDtoDAsync(dC, dB, BYTES, S)); + CK(p_cuStreamEndCapture(S, &g1)); + printf("[capture] segment 1 ended, graph=%p\n", (void *)g1); + CK(p_cuGraphInstantiate(&e1, g1, 0)); + if (fails) { printf("VERDICT: SEGMENT1_FAILED\n"); return 3; } + + /* A third segment split by a BARE break (zero work between end and re-begin), + the degenerate case of breakable_cuda_graph.py:370-374. */ + printf("[capture] bare break: end then immediately re-begin with NO work between\n"); + { + CUresult r = p_cuStreamBeginCapture(S, MODE_THREAD_LOCAL); + if (r != 0) { + printf("FAIL bare re-begin -> %d (%s)\n", r, errname(r)); + printf("VERDICT: BARE_REBEGIN_REFUSED\n"); + return 3; + } + } + CK(p_cuMemsetD32Async(dA, 0u, N, S)); /* harmless, keeps the graph non-empty */ + CK(p_cuStreamEndCapture(S, &g2)); + CK(p_cuGraphInstantiate(&e2, g2, 0)); + if (fails) { printf("VERDICT: BARE_SEGMENT_FAILED\n"); return 3; } + + /* ---- REPLAY PHASE: three replays with DIFFERENT inputs ---- */ + printf("\n[replay] three replays, fresh input each time\n"); + for (rep = 0; rep < 3; rep++) { + int base = 100 * (rep + 1), bad = 0; + for (i = 0; i < N; i++) in[i] = base + i; + CK(p_cuMemcpyHtoD(dA, in, BYTES)); + + CK(p_cuGraphLaunch(e0, S)); /* segment 0: B <- A */ + CK(p_cuStreamSynchronize(S)); + break_fn(); /* break: B <- B*2+3 on the host */ + CK(p_cuGraphLaunch(e1, S)); /* segment 1: C <- B */ + CK(p_cuStreamSynchronize(S)); + + CK(p_cuMemcpyDtoH(out, dC, BYTES)); + for (i = 0; i < N; i++) { + int want = (base + i) * 2 + 3; + if (out[i] != want) { + if (bad < 3) printf(" MISMATCH rep=%d i=%d got=%d want=%d\n", rep, i, out[i], want); + bad++; + } + } + printf(" replay %d: base=%d out[0]=%d want=%d mismatches=%d\n", rep, base, + out[0], base * 2 + 3, bad); + if (bad) fails++; + } + + if (p_cuGraphExecDestroy) { p_cuGraphExecDestroy(e0); p_cuGraphExecDestroy(e1); p_cuGraphExecDestroy(e2); } + if (p_cuGraphDestroy) { p_cuGraphDestroy(g0); p_cuGraphDestroy(g1); p_cuGraphDestroy(g2); } + + printf("\nfails=%d\n", fails); + printf("VERDICT: %s\n", fails ? "FAILED" : "REBEGIN_HOLDS"); + return fails ? 1 : 0; +} diff --git a/scripts/probe_cudagraph_rebegin.cu b/scripts/probe_cudagraph_rebegin.cu new file mode 100644 index 000000000..963bbf0b8 --- /dev/null +++ b/scripts/probe_cudagraph_rebegin.cu @@ -0,0 +1,167 @@ +// ENG-CUDAGRAPH-BREAK (#1163) W1 exit criterion, issue #1192. +// +// Question W0 left open: does CUDA permit cudaStreamEndCapture followed by +// cudaStreamBeginCapture on the SAME stream mid-forward, with EAGER work between +// them, on OUR stream configuration? +// +// "Our stream configuration" is mirrored exactly from src/vt/cuda/cuda_backend.cu: +// BeginCapture -> cudaStreamBeginCapture(s, cudaStreamCaptureModeThreadLocal) (:204) +// EndCaptureGraph -> cudaStreamEndCapture + cudaGraphInstantiate, stores nothing (:225) +// ReplayGraph -> cudaGraphLaunch(exec, s) (:233) +// DestroyGraph -> cudaGraphExecDestroy (:288) +// The stream is a non-blocking stream, as vt queues are. +// +// The shape under test is the seam's, not a toy: segment, break, segment, break, +// segment. The break function is genuinely host-dependent (device -> host -> +// host arithmetic -> device), which is the class of operation a break point +// exists for, and it runs eagerly on the SAME stream that capture just left. +#include +#include + +#define N 1024 + +static int fails = 0; +#define CK(expr) \ + do { \ + cudaError_t _e = (expr); \ + if (_e != cudaSuccess) { \ + std::printf("FAIL %s -> %d (%s)\n", #expr, (int)_e, \ + cudaGetErrorString(_e)); \ + ++fails; \ + } \ + } while (0) + +__global__ void add_k(int* dst, const int* src, int add, int n) { + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) dst[i] = src[i] + add; +} + +static cudaStream_t S; +static int *dA, *dB, *dC; +static int hostbuf[N]; + +// The break: host-dependent work, eager, on the same stream. Mirrors what a real +// break point does (a host readback that decides something, then a writeback). +static void break_fn(int mul) { + CK(cudaMemcpyAsync(hostbuf, dB, N * sizeof(int), cudaMemcpyDeviceToHost, S)); + CK(cudaStreamSynchronize(S)); + for (int i = 0; i < N; ++i) hostbuf[i] = hostbuf[i] * mul; + CK(cudaMemcpyAsync(dB, hostbuf, N * sizeof(int), cudaMemcpyHostToDevice, S)); + CK(cudaStreamSynchronize(S)); +} + +// EndCaptureGraph, byte-for-byte the cuda_backend.cu:225-232 body. +static cudaGraphExec_t EndCaptureGraph(cudaStream_t s) { + cudaGraph_t g = nullptr; + CK(cudaStreamEndCapture(s, &g)); + cudaGraphExec_t e = nullptr; + CK(cudaGraphInstantiate(&e, g, 0)); + cudaGraphDestroy(g); + return e; +} + +int main() { + cudaDeviceProp prop{}; + int dev = 0; + CK(cudaGetDevice(&dev)); + CK(cudaGetDeviceProperties(&prop, dev)); + int rt = 0, drv = 0; + cudaRuntimeGetVersion(&rt); + cudaDriverGetVersion(&drv); + std::printf("device=%s sm_%d%d runtime=%d driver=%d integrated=%d\n", prop.name, + prop.major, prop.minor, rt, drv, prop.integrated); + + CK(cudaStreamCreateWithFlags(&S, cudaStreamNonBlocking)); + CK(cudaMalloc(&dA, N * sizeof(int))); + CK(cudaMalloc(&dB, N * sizeof(int))); + CK(cudaMalloc(&dC, N * sizeof(int))); + if (fails) { std::printf("VERDICT: SETUP_FAILED\n"); return 3; } + + int in[N], out[N]; + for (int i = 0; i < N; ++i) in[i] = i + 1; + CK(cudaMemcpy(dA, in, N * sizeof(int), cudaMemcpyHostToDevice)); + + const int blocks = (N + 255) / 256; + + // ---------------- capture phase ---------------- + std::printf("\n[capture] segment 0: BeginCapture(thread-local)\n"); + CK(cudaStreamBeginCapture(S, cudaStreamCaptureModeThreadLocal)); + cudaStreamCaptureStatus st{}; + CK(cudaStreamIsCapturing(S, &st)); + std::printf("[capture] is_capturing in segment 0 = %d\n", (int)st); + add_k<<>>(dB, dA, 1, N); // B = A + 1 + cudaGraphExec_t e0 = EndCaptureGraph(S); + CK(cudaStreamIsCapturing(S, &st)); + std::printf("[capture] segment 0 ended, exec=%p, is_capturing=%d\n", (void*)e0, (int)st); + if (fails) { std::printf("VERDICT: SEGMENT0_FAILED\n"); return 3; } + + std::printf("[capture] BREAK 0: eager host-dependent work on the SAME stream\n"); + break_fn(2); // B = B * 2 + if (fails) { std::printf("VERDICT: EAGER_BETWEEN_SEGMENTS_FAILED\n"); return 3; } + + std::printf("[capture] segment 1: RE-BEGIN capture on the same stream <-- EXIT CRITERION\n"); + { + cudaError_t r = cudaStreamBeginCapture(S, cudaStreamCaptureModeThreadLocal); + if (r != cudaSuccess) { + std::printf("FAIL re-begin -> %d (%s)\nVERDICT: REBEGIN_REFUSED\n", (int)r, + cudaGetErrorString(r)); + return 3; + } + } + add_k<<>>(dC, dB, 3, N); // C = B + 3 + cudaGraphExec_t e1 = EndCaptureGraph(S); + std::printf("[capture] segment 1 ended, exec=%p\n", (void*)e1); + if (fails) { std::printf("VERDICT: SEGMENT1_FAILED\n"); return 3; } + + // The BARE break: end then immediately re-begin with NO work in between, the + // degenerate case of breakable_cuda_graph.py:370-374. + std::printf("[capture] BREAK 1 (bare): re-begin with NO work between\n"); + { + cudaError_t r = cudaStreamBeginCapture(S, cudaStreamCaptureModeThreadLocal); + if (r != cudaSuccess) { + std::printf("FAIL bare re-begin -> %d (%s)\nVERDICT: BARE_REBEGIN_REFUSED\n", + (int)r, cudaGetErrorString(r)); + return 3; + } + } + add_k<<>>(dC, dC, 10, N); // C = C + 10 + cudaGraphExec_t e2 = EndCaptureGraph(S); + std::printf("[capture] segment 2 ended, exec=%p (3 segments, 2 breaks)\n", (void*)e2); + if (fails) { std::printf("VERDICT: BARE_SEGMENT_FAILED\n"); return 3; } + + // ---------------- replay phase: THREE replays, fresh input each ---------------- + std::printf("\n[replay] 3 replays, different input each time\n"); + for (int rep = 0; rep < 3; ++rep) { + const int base = 100 * (rep + 1); + for (int i = 0; i < N; ++i) in[i] = base + i; + CK(cudaMemcpy(dA, in, N * sizeof(int), cudaMemcpyHostToDevice)); + + CK(cudaGraphLaunch(e0, S)); // segment 0 + CK(cudaStreamSynchronize(S)); + break_fn(2); // break 0 (eager) + CK(cudaGraphLaunch(e1, S)); // segment 1 + // break 1 is bare: nothing runs + CK(cudaGraphLaunch(e2, S)); // segment 2 + CK(cudaStreamSynchronize(S)); + CK(cudaMemcpy(out, dC, N * sizeof(int), cudaMemcpyDeviceToHost)); + + int bad = 0; + for (int i = 0; i < N; ++i) { + const int want = ((base + i) + 1) * 2 + 3 + 10; + if (out[i] != want) { + if (bad < 3) + std::printf(" MISMATCH rep=%d i=%d got=%d want=%d\n", rep, i, out[i], want); + ++bad; + } + } + std::printf(" replay %d: base=%d out[0]=%d want=%d mismatches=%d\n", rep, base, + out[0], (base + 1) * 2 + 13, bad); + if (bad) ++fails; + } + + cudaGraphExecDestroy(e0); + cudaGraphExecDestroy(e1); + cudaGraphExecDestroy(e2); + std::printf("\nfails=%d\nVERDICT: %s\n", fails, fails ? "FAILED" : "REBEGIN_HOLDS"); + return fails ? 1 : 0; +} diff --git a/src/vllm/model_executor/models/qwen3.cpp b/src/vllm/model_executor/models/qwen3.cpp index 80ed8b3b8..27116710a 100644 --- a/src/vllm/model_executor/models/qwen3.cpp +++ b/src/vllm/model_executor/models/qwen3.cpp @@ -55,6 +55,8 @@ #include "vllm/model_executor/models/qwen3_5_internal.h" // detail::DeviceTokenIds seam #include "vllm/platforms/interface.h" #include "vt/backend.h" +#include "vllm/model_executor/models/dense_attn_graph_break.h" // CopyOutput(optional) +#include "vt/breakable_graph.h" // ENG-CUDAGRAPH-BREAK: the break-point seam #include "vt/ops.h" #include "vt/tenstorrent/tenstorrent_device.h" // WarmRopeCosSin (item-5 TT-only) #include "vt/recipes.h" @@ -132,14 +134,40 @@ void RunLayer(Dev d, const Qwen3DenseLayerWeights& layer, const HfConfig& cfg, vt::RmsNorm(d.q, dhn.t(), hidden.t(), w_in, vt::RmsNormArgs{eps, false}, &res.t()); } - DBuf attn = AttnBlock(d, layer.attn, cfg, dhn.t(), si, meta, kv, T, tp); + // ENG-CUDAGRAPH-BREAK (#1163) W1 (#1192): THE DENSE ATTENTION ENTRY IS A BREAK + // POINT. The boundary is vLLM's, not ours to invent: its v1 default splits at + // `splitting_ops`, defaulted to the attention family + // (`vllm/config/compilation.py:517,764-772,1145` @ pin `5559679229`). The + // registration form is SGLang's, because vLLM gets its split from Dynamo and + // FX and we have no compiler: one line at the site, exactly as + // `layers/radix_attention.py:256` @ `f63458b5be`. THE SITE IS THE + // REGISTRATION. + // + // Destination form, because `AttnBlock` returns a FRESH pooled buffer on every + // call. The destination is a `vt::BreakSlot`, not a local `std::optional`: + // the following segment bakes the destination's address, and a plain local + // dies on this function's `return` while the pooled block it named goes back + // on the `DevicePool` free list — a host use-after-scope plus D1's reuse + // hazard, at the one site that has to obey the rule. The slot hands its + // storage to the seam on the capturing path and keeps it inline on the + // pass-through path. + // + // Outside a capture scope — which is every production step today, until W2 + // migrates this model's decode driver onto the seam — `GraphBreak` moves the + // result into the slot and returns, so this is byte-identical to the + // `DBuf attn = AttnBlock(...)` it replaces, allocates nothing extra, and makes + // zero backend calls. + vt::BreakSlot> attn; + vt::GraphBreak([&] { return AttnBlock(d, layer.attn, cfg, dhn.t(), si, meta, kv, T, tp); }, + attn); + DBuf& attn_buf = attn->value(); Tensor w_post = ResidentWeight(d, layer.post_attention_layernorm, {H}); DBuf dh2(d, DType::kBF16, {T, H}); if (FusedChainAdoptEnabled()) { - vt::FusedChain(d.q, dh2.t(), attn.t(), w_post, &res.t(), vt::kFusedAddRmsNormStd, eps); + vt::FusedChain(d.q, dh2.t(), attn_buf.t(), w_post, &res.t(), vt::kFusedAddRmsNormStd, eps); } else { - vt::RmsNorm(d.q, dh2.t(), attn.t(), w_post, vt::RmsNormArgs{eps, false}, &res.t()); + vt::RmsNorm(d.q, dh2.t(), attn_buf.t(), w_post, vt::RmsNormArgs{eps, false}, &res.t()); } hidden = MlpBlock(d, layer.mlp, cfg, dh2.t(), T, tp); diff --git a/src/vt/breakable_graph.cpp b/src/vt/breakable_graph.cpp new file mode 100644 index 000000000..7341f6f9c --- /dev/null +++ b/src/vt/breakable_graph.cpp @@ -0,0 +1,247 @@ +// Implementation of the break-point capture seam. See +// `include/vt/breakable_graph.h` for the port map and the lifetime rules, and +// `.agents/specs/eng-cudagraph-break.md` for the design. Row +// ENG-CUDAGRAPH-BREAK W1, issue #1192, parent #1163. +// +// Upstream: SGLang at pin `f63458b5be`, +// `python/sglang/srt/model_executor/runner_backend_utils/breakable_cuda_graph/breakable_cuda_graph.py`; +// every bare `:N` is a line in that file. +#include "vt/breakable_graph.h" + +#include +#include +#include +#include +#include + +namespace vt { +namespace { + +// The capture scope of the CURRENT THREAD. Replaces upstream's `ContextVar` +// (`:63`); see the header for why this is thread-local and not global. +thread_local GraphCaptureScope* t_current_scope = nullptr; + +std::atomic g_break_points{0}; +std::atomic g_segments{0}; +std::atomic g_breaks{0}; +std::atomic g_replays{0}; + +} // namespace + +namespace detail { +void CountBreakPoint() { g_break_points.fetch_add(1, std::memory_order_relaxed); } +} // namespace detail + +GraphBreakStats GetGraphBreakStats() { + GraphBreakStats s; + s.break_points_reached = g_break_points.load(std::memory_order_relaxed); + s.segments_captured = g_segments.load(std::memory_order_relaxed); + s.breaks_registered = g_breaks.load(std::memory_order_relaxed); + s.replays = g_replays.load(std::memory_order_relaxed); + return s; +} + +void ResetGraphBreakStats() { + g_break_points.store(0, std::memory_order_relaxed); + g_segments.store(0, std::memory_order_relaxed); + g_breaks.store(0, std::memory_order_relaxed); + g_replays.store(0, std::memory_order_relaxed); +} + +// Read ONCE into a function-local static, so a process is in exactly one lane +// for its whole life and nothing can toggle it mid-run. +bool GraphCaptureEnabled() { + static const bool enabled = [] { + const char* e = std::getenv("VLLM_CPP_CUDAGRAPH"); + return e == nullptr || std::string(e) != "0"; + }(); + return enabled; +} + +// --------------------------------------------------------------------------- +// CopyOutput — port of `_copy_output` (`:172-201`). +// --------------------------------------------------------------------------- + +static size_t Numel(const Tensor& t) { + size_t n = 1; + for (int i = 0; i < t.rank; ++i) n *= static_cast(t.shape[i]); + return n; +} + +void CopyOutput(Backend& b, Queue& q, Tensor& dst, const Tensor& src) { + if (dst.data == nullptr || src.data == nullptr) return; + if (dst.data == src.data) return; // already the same storage: nothing to do + // The two sides must AGREE before a byte moves. Sizing the copy from `dst` + // alone turns a break function that returned the wrong shape or dtype into an + // out-of-bounds read of `src` — a fault at best and someone else's data at + // worst — where upstream's `dst.copy_(src)` raises. This is the refusal that + // makes a mis-registered break point loud instead of silently wrong. + VT_CHECK(dst.dtype == src.dtype, + "CopyOutput: break output dtype does not match its destination"); + VT_CHECK(Numel(dst) == Numel(src), + "CopyOutput: break output element count does not match its destination"); + const size_t bytes = Numel(dst) * SizeOf(dst.dtype); + // In place, into the address the FOLLOWING SEGMENT baked. `dst` is not + // rebound, which is upstream's `assertIs(result, dst)` (`:190`). + b.Copy(q, dst.data, src.data, bytes); +} + +void CopyOutput(Backend& b, Queue& q, std::vector& dst, + const std::vector& src) { + const size_t n = dst.size() < src.size() ? dst.size() : src.size(); + for (size_t i = 0; i < n; ++i) CopyOutput(b, q, dst[i], src[i]); +} + +void CopyOutput(Backend& b, Queue& q, std::map& dst, + const std::map& src) { + for (const auto& kv : src) { + auto it = dst.find(kv.first); + // A key the destination holds is copied IN PLACE; a key it does not hold is + // ASSIGNED, which is upstream's `dst[key] = src_val` branch (`:200`). + if (it != dst.end()) + CopyOutput(b, q, it->second, kv.second); + else + dst[kv.first] = kv.second; + } +} + +// --------------------------------------------------------------------------- +// BreakableGraph — port of `BreakableCUDAGraph` (`:246-274`). +// --------------------------------------------------------------------------- + +BreakableGraph::~BreakableGraph() { Reset(); } + +void BreakableGraph::Reset() { + if (backend_ != nullptr) { + for (void* g : segments_) backend_->DestroyGraph(g); + } + segments_.clear(); + break_fns_.clear(); + // The replay count describes the graph that was just released. Leaving it + // behind makes the next capture report replays it never ran, and G3's whole + // job is to be the number nobody has to trust twice. + replays_ = 0; +} + +void BreakableGraph::Replay(Queue& q) { + if (backend_ == nullptr || segments_.empty()) return; + // The invariant, asserted where violating it would be SILENT. A container with + // `break_count() == segment_count()` replays with its last break dropped, and + // the only symptom is a forward that skips one host-dependent operation. + VT_CHECK(break_fns_.size() + 1 == segments_.size(), + "BreakableGraph::Replay: segment_count() must equal break_count() + 1"); + // Interleaved, mirroring `:255-263`: segment i, then break i. + for (size_t i = 0; i < segments_.size(); ++i) { + backend_->ReplayGraph(q, segments_[i]); + if (i + 1 < segments_.size()) break_fns_[i](); + } + ++replays_; + g_replays.fetch_add(1, std::memory_order_relaxed); +} + +// --------------------------------------------------------------------------- +// GraphCaptureScope — port of `BreakableCUDAGraphCapture` (`:277-367`). +// --------------------------------------------------------------------------- + +GraphCaptureScope* GraphCaptureScope::Current() { return t_current_scope; } + +GraphCaptureScope::GraphCaptureScope(Backend& b, Queue& q, BreakableGraph& out) + : b_(&b), q_(&q), g_(&out), uncaught_on_entry_(std::uncaught_exceptions()) { + // Refused before ANY backend call, and refused in both lanes so the switch + // cannot change which programs are legal. + // + // NESTING. `prev_`-style save and restore made a nested pair look contemplated; + // it is not. Two scopes on one queue trace `Begin Begin EndCaptureGraph Begin + // EndCaptureGraph EndCaptureGraph ...`, whose second `BeginCapture` lands on an + // already-capturing stream — `CUDA_ERROR_ILLEGAL_STATE` on a real backend. + VT_CHECK(t_current_scope == nullptr, + "GraphCaptureScope: a capture scope is already open on this thread"); + // RE-ENTRY. Appending to a container that already holds a capture yields + // `break_count() == segment_count()`, whose last break Replay drops. Reset it. + VT_CHECK(!out.captured() && out.break_count() == 0, + "GraphCaptureScope: this BreakableGraph already holds a capture; Reset() it first"); + active_ = GraphCaptureEnabled() && b.SupportsGraphCapture(); + if (!active_) return; // inert: the forward runs eager, exactly as today + g_->backend_ = b_; + t_current_scope = this; + BeginSegment(); +} + +GraphCaptureScope::~GraphCaptureScope() { + if (!active_) return; + // `__exit__` closes the final segment inside a `finally` (`:322-333`), so the + // scope pointer is restored even when the close throws. A destructor that + // propagates would terminate, and a skipped close poisons the stream + // permanently — which is why three drivers hand-rolled this same drain + // (`qwen3_5.cpp:9913`, `qwen3_5.cpp:10335`, `qwen3_dflash.cpp:1106`). + // + // THE DRAIN IS ABOUT THE EXCEPTION, NOT ONLY ABOUT THE CLOSE. Catching a + // throwing `EndCaptureGraph` covers one failure. The failure that actually + // happens is a break function, or ordinary model code between two break + // points, throwing MID-CAPTURE: the close then succeeds, the stream is clean, + // and the container is left holding a PARTIAL forward that reports + // `captured() == true` and replays as half a step. Comparing the uncaught + // depth against the one recorded at entry is what tells the two apart, and the + // partial container is destroyed rather than handed back. + const bool unwinding = std::uncaught_exceptions() > uncaught_on_entry_; + try { + EndSegment(); + } catch (...) { + g_->Reset(); + t_current_scope = nullptr; + return; + } + if (unwinding) g_->Reset(); + t_current_scope = nullptr; +} + +void GraphCaptureScope::BeginSegment() { + if (!active_ || segment_open_) return; + b_->BeginCapture(*q_); + segment_open_ = true; +} + +void GraphCaptureScope::EndSegment() { + if (!active_ || !segment_open_) return; + segment_open_ = false; // cleared FIRST: a throwing end must not be retried + void* seg = b_->EndCaptureGraph(*q_); + g_->segments_.push_back(seg); + g_segments.fetch_add(1, std::memory_order_relaxed); +} + +void GraphCaptureScope::AppendBreak(std::function fn, const void* destination) { + if (!active_) return; + // The ALIASING half of lifetime rule 1, refused rather than documented. Two + // break points writing through one CELL bind both replay closures to the same + // address: the earlier writeback is overwritten on every replay, and any + // segment that baked the earlier destination reads the later break's data. + // Nothing faults and the token gate cannot see it, which is why this is a + // refusal and not a comment. The check is here, at the ONE registration + // point, so a `GraphBreak` form added later cannot opt out of it. + if (destination != nullptr) { + for (const void* d : destinations_) { + VT_CHECK(d != destination, + "GraphBreak: this destination is already the target of another break point " + "in the open capture; both replays would write through the SAME address. " + "Give each break point its own BreakSlot"); + } + destinations_.push_back(destination); + } + g_->break_fns_.push_back(std::move(fn)); + g_breaks.fetch_add(1, std::memory_order_relaxed); +} + +// The bare marker (`:370-374`). +void GraphBreak() { + GraphCaptureScope* s = GraphCaptureScope::Current(); + if (s == nullptr || !s->active()) { + detail::CountBreakPoint(); + return; + } + s->EndSegment(); + s->AppendBreak([] {}, nullptr); + s->BeginSegment(); + detail::CountBreakPoint(); +} + +} // namespace vt diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index b865d484f..959157887 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -402,6 +402,7 @@ vllm_cpp_add_test(test_fp8_block_quant vllm/model_executor/layers/test_fp8_block_quant.cpp) vllm_cpp_add_test(test_linear_method vllm/model_executor/layers/test_linear_method.cpp) target_include_directories(test_linear_method PRIVATE ${CMAKE_SOURCE_DIR}/src) +vllm_cpp_add_test(test_qwen3_break_point vllm/models/test_qwen3_break_point.cpp) vllm_cpp_add_test(test_decode_graph_sizes vllm/models/test_decode_graph_sizes.cpp) vllm_cpp_add_test(test_qwen35_paged_forward vllm/models/test_qwen35_paged_forward.cpp) # GDN-MOE-BF16-OUT (#1168): the MoE forward's GDN output-activation dtype probe @@ -1686,6 +1687,7 @@ target_include_directories(test_gdn_prefill_conv PRIVATE ${CMAKE_SOURCE_DIR}/src # resident), never freed, so a device pointer baked into a captured pure-decode CUDA # graph stays valid — the fix for the 35B c2+ online-serving illegal-memory-access. # The device proof is the DGX serving harness; this pins the never-free guarantee. +vllm_cpp_add_test(test_breakable_graph vt/test_breakable_graph.cpp) vllm_cpp_add_test(test_graph_safe_scratch vt/test_graph_safe_scratch.cpp) target_include_directories(test_graph_safe_scratch PRIVATE ${CMAKE_SOURCE_DIR}/src) # CPU threadpool (QUANT-GGUF-CPU-THREADPOOL): barrier/park-wake/chunk-steal diff --git a/tests/vllm/models/test_qwen3_break_point.cpp b/tests/vllm/models/test_qwen3_break_point.cpp new file mode 100644 index 000000000..b5584bc8f --- /dev/null +++ b/tests/vllm/models/test_qwen3_break_point.cpp @@ -0,0 +1,240 @@ +// THE G2 REACHABILITY GATE for ENG-CUDAGRAPH-BREAK W1 (#1192, parent #1163), +// per `.agents/reachability.md` and the AGENTS.md "Nothing lands dead" rule. +// +// WHAT IT MEASURES, and why a unit test could not. `tests/vt/test_breakable_graph.cpp` +// proves the seam WORKS by constructing it by hand; that measures a class. This +// file proves something else entirely: that a REAL MODEL'S FORWARD reaches it. +// It drives `Qwen3DenseModel::Forward` — the same function the registered +// forward calls under `ModelRegistry::Forward` +// (`src/vllm/v1/worker/gpu/runner.cpp:1465`) — with a capture scope open, and +// counts the segments that come out. Delete the one `vt::GraphBreak` line at the +// dense attention entry in `src/vllm/model_executor/models/qwen3.cpp` and this +// file goes RED, because the forward then produces one segment instead of +// `num_hidden_layers + 1`. +// +// The recording backend is the harness adaptation stated once in +// `tests/vt/recording_capture_backend.h`: the model still computes on the CPU +// backend, and the scope's backend only records the capture call sequence, so +// this gate needs no GPU and runs in continuous integration. +// +// It also carries the G4 half of the spec's `## Gates`: the seam changes NO +// numerics. The logits from the scoped run are compared BIT FOR BIT against the +// unscoped one, which is what makes this stage reversible. +#include + +#include +#include +#include +#include +#include +#include +#include +#include + +#include "vllm/model_executor/models/dense_attn_graph_break.h" // the writeback overload +#include "vllm/model_executor/models/qwen3.h" +#include "vllm/transformers_utils/hf_config.h" +#include "vt/backend.h" +#include "vt/breakable_graph.h" +#include "vt/dtype.h" +#include "vt/recording_capture_backend.h" +#include "vt/tensor.h" + +// THE WRITEBACK BRANCH IS SELECTED, asserted at COMPILE TIME rather than +// believed. The destination form's `static_assert` already refuses a type with +// no `CopyOutput`, so a renamed, moved or shadowed overload is now a build +// failure at the model — but only this line says WHICH branch the production +// destination type takes, and it is the branch the D9 wrong-numerics path hides +// behind. Renaming `vllm::dense_attn::CopyOutput` reds here and at the model, +// where it used to leave both suites green. +static_assert(vt::detail::HasCopyOutput>::value, + "the qwen3 break point's destination must have a CopyOutput overload; without " + "one every replay leaves the following segment reading capture-time data"); + +namespace { + +using vllm::HfConfig; +using vllm::PagedKvCache; +using vllm::Qwen3DenseWeights; +using vllm::v1::CommonAttentionMetadata; +using vt::DType; +using vt_test::RecordingCaptureBackend; + +vt::Queue Q() { return vt::Queue{vt::Device{vt::DeviceType::kCPU, 0}, nullptr}; } + +vllm::OwnedTensor MakeBf16(const std::vector& shape, bool nk, uint32_t seed, + float scale = 0.08f) { + vllm::OwnedTensor o; + o.dtype = DType::kBF16; + o.nk = nk; + o.rank = static_cast(shape.size()); + int64_t numel = 1; + for (int i = 0; i < o.rank; ++i) { + o.shape[i] = shape[static_cast(i)]; + numel *= shape[static_cast(i)]; + } + o.bytes.resize(static_cast(numel) * sizeof(uint16_t)); + auto* p = reinterpret_cast(o.bytes.data()); + std::mt19937 rng(seed); + std::uniform_real_distribution dist(-scale, scale); + for (int64_t i = 0; i < numel; ++i) p[i] = vt::F32ToBF16(dist(rng)); + return o; +} + +HfConfig TinyConfig() { + HfConfig c; + c.num_hidden_layers = 2; + c.hidden_size = 64; + c.num_attention_heads = 4; + c.num_key_value_heads = 2; + c.head_dim = 16; + c.rotary_dim = 16; // partial_rotary_factor 1.0 + c.intermediate_size = 128; + c.rms_norm_eps = 1e-6; + c.rope_theta = 1000000.0; + c.vocab_size = 100; + return c; +} + +Qwen3DenseWeights TinyWeights(const HfConfig& c) { + const int64_t H = c.hidden_size, Hq = c.num_attention_heads, Hkv = c.num_key_value_heads; + const int64_t Dh = c.head_dim, I = c.intermediate_size, V = c.vocab_size; + const int64_t qdim = Hq * Dh, kdim = Hkv * Dh; + Qwen3DenseWeights w; + w.tie_word_embeddings = true; + w.attention_bias = false; + w.embed_tokens = MakeBf16({V, H}, /*nk=*/false, 1); + w.final_norm = MakeBf16({H}, false, 2, 0.5f); + uint32_t seed = 100; + for (int64_t l = 0; l < c.num_hidden_layers; ++l) { + vllm::Qwen3DenseLayerWeights lw; + lw.input_layernorm = MakeBf16({H}, false, seed++, 0.5f); + lw.post_attention_layernorm = MakeBf16({H}, false, seed++, 0.5f); + lw.attn.qkv_proj = MakeBf16({qdim + 2 * kdim, H}, /*nk=*/true, seed++); + lw.attn.o_proj = MakeBf16({H, qdim}, /*nk=*/true, seed++); + lw.attn.q_norm = MakeBf16({Dh}, false, seed++, 0.5f); + lw.attn.k_norm = MakeBf16({Dh}, false, seed++, 0.5f); + lw.mlp.gate_up_proj = MakeBf16({2 * I, H}, /*nk=*/true, seed++); + lw.mlp.down_proj = MakeBf16({H, I}, /*nk=*/true, seed++); + w.layers.push_back(std::move(lw)); + } + return w; +} + +// Single-sequence prefill paged KV cache (all layers full-attention). +struct CachePool { + std::vector> buf; + std::vector attn_kv; + CachePool(const HfConfig& c, int64_t num_blocks, int64_t block_size) { + const int64_t Hkv = c.num_key_value_heads, Dh = c.head_dim; + for (int64_t l = 0; l < c.num_hidden_layers; ++l) + buf.emplace_back(static_cast(num_blocks * 2 * block_size * Hkv * Dh), 0.0f); + for (auto& b : buf) { + PagedKvCache kv; + kv.data = b.data(); + kv.dtype = DType::kF32; + kv.num_blocks = num_blocks; + kv.block_size = block_size; + kv.num_kv_heads = Hkv; + kv.head_size = Dh; + attn_kv.push_back(kv); + } + } +}; + +CommonAttentionMetadata PrefillMeta(int64_t T, int64_t block_size) { + CommonAttentionMetadata m; + m.num_reqs = 1; + m.num_actual_tokens = static_cast(T); + m.query_start_loc = {0, static_cast(T)}; + m.query_start_loc_cpu = m.query_start_loc; + m.seq_lens = {static_cast(T)}; + m.seq_lens_cpu = m.seq_lens; + m.max_query_len = static_cast(T); + m.max_seq_len = static_cast(T); + m.block_table_num_cols = 1; + m.block_table_tensor = {0}; + for (int64_t t = 0; t < T; ++t) m.slot_mapping.push_back(t % block_size); + m.causal = true; + return m; +} + +// One prefill forward through the production model function, optionally with a +// capture scope open over `rec`. +std::vector RunForward(const HfConfig& c, const Qwen3DenseWeights& w, + RecordingCaptureBackend* rec, vt::BreakableGraph* g) { + const int64_t T = 5; + CachePool pool(c, /*num_blocks=*/2, /*block_size=*/8); + const CommonAttentionMetadata am = PrefillMeta(T, 8); + const std::vector tokens = {3, 17, 42, 8, 61}; + const std::vector positions = {0, 1, 2, 3, 4}; + vt::Queue q = Q(); + if (rec == nullptr) + return vllm::Qwen3DenseModel::Forward(tokens, positions, am, pool.attn_kv, w, c, q); + vt::Queue rq = rec->CreateQueue(); + // The CAPTURING lane is this gate's precondition: with VLLM_CPP_CUDAGRAPH=0 + // exported the scope is inert by design, and an environment red must not read + // as a code red. + REQUIRE_MESSAGE(vt::GraphCaptureEnabled(), + "this gate needs the CAPTURING lane; VLLM_CPP_CUDAGRAPH=0 is set"); + vt::GraphCaptureScope scope(*rec, rq, *g); + REQUIRE(scope.active()); + return vllm::Qwen3DenseModel::Forward(tokens, positions, am, pool.attn_kv, w, c, q); +} + +} // namespace + +TEST_CASE("G2: the Qwen3 dense forward REACHES the break seam, once per layer") { + const HfConfig c = TinyConfig(); + const Qwen3DenseWeights w = TinyWeights(c); + REQUIRE(c.num_hidden_layers == 2); + + RecordingCaptureBackend rec; + vt::BreakableGraph g; + const std::vector scoped = RunForward(c, w, &rec, &g); + + // One break per decoder layer, so N layers yield N breaks and N+1 segments. + // This is the number that goes to 0 and 1 when the call site is deleted. + CHECK(g.break_count() == static_cast(c.num_hidden_layers)); + CHECK(g.segment_count() == static_cast(c.num_hidden_layers) + 1); + CHECK(rec.Count("Begin") == c.num_hidden_layers + 1); + CHECK(rec.Count("EndCaptureGraph") == c.num_hidden_layers + 1); + // The invariant the seam ASSERTS — `GraphCaptureScope` refuses a container that + // already holds a capture, and `Replay` refuses counts that disagree — + // restated here where a real model produced them. + CHECK(g.segment_count() == g.break_count() + 1); + + // G4: the seam changed no numerics. Bit for bit, not approximately. + const std::vector eager = RunForward(c, w, nullptr, nullptr); + REQUIRE(scoped.size() == eager.size()); + REQUIRE(!scoped.empty()); + size_t differing = 0; + for (size_t i = 0; i < scoped.size(); ++i) + if (std::memcmp(&scoped[i], &eager[i], sizeof(float)) != 0) ++differing; + CHECK(differing == 0); + MESSAGE("logits compared bit for bit: " << scoped.size() << " values, " << differing + << " differing"); + for (float v : scoped) REQUIRE(std::isfinite(v)); +} + +// Outside a scope — which is every production step today, until W2 migrates this +// model's decode driver onto the seam — the registered break point is a pure +// pass-through: the forward runs eager and makes ZERO backend capture calls. +TEST_CASE("G2: with no scope open the same forward captures nothing at all") { + const HfConfig c = TinyConfig(); + const Qwen3DenseWeights w = TinyWeights(c); + RecordingCaptureBackend rec; + + const vt::GraphBreakStats before = vt::GetGraphBreakStats(); + const std::vector out = RunForward(c, w, nullptr, nullptr); + const vt::GraphBreakStats after = vt::GetGraphBreakStats(); + + REQUIRE(!out.empty()); + CHECK(rec.Trace().empty()); + CHECK(after.segments_captured == before.segments_captured); + CHECK(after.breaks_registered == before.breaks_registered); + // The break point was still REACHED, once per layer — which is the counter + // that distinguishes "the model routes into the seam" from "the seam exists". + CHECK(after.break_points_reached - before.break_points_reached == + c.num_hidden_layers); +} diff --git a/tests/vt/recording_capture_backend.h b/tests/vt/recording_capture_backend.h new file mode 100644 index 000000000..a38897e8c --- /dev/null +++ b/tests/vt/recording_capture_backend.h @@ -0,0 +1,187 @@ +// vllm.cpp original test helper (ENG-CUDAGRAPH-BREAK W1, #1192); no upstream +// mirror, because upstream has none to mirror. +// +// THE HARNESS ADAPTATION, stated once and stated exactly. Every class in +// SGLang's BCG suite +// (`test/registered/cuda_graph/breakable/test_breakable_cuda_graph.py` @ +// `f63458b5be`) raises `unittest.SkipTest` without CUDA (`:34-36`, `:177-178`, +// `:235-236`), so upstream runs its eleven unit cases on a real device. This box +// and continuous integration have no NVIDIA GPU, so the ported cases run against +// a backend that RECORDS the capture call sequence and SIMULATES a graph. +// +// What "simulates a graph" means here, because a recording backend that only +// logged would have cost the upstream cases their whole point. `Record(fn)` +// models one captured operation: during capture the work is filed against the +// currently open segment and NOT executed, exactly as a real stream capture +// files a kernel without running it, and `ReplayGraph` runs the filed work in +// order. That is what makes upstream's post-replay arithmetic chains portable — +// `x.fill_(5); graph.replay(); assert y == 6` (`:49-63`) is the same assertion +// here — instead of degrading them into capture-time value checks, which assert +// something upstream never asserted. +// +// What is NOT modelled: real device memory, asynchrony, and any error a real +// runtime would raise. A case that needs those is a GPU case and is recorded as +// owed in the spec rather than approximated here. +// +// A recording backend is also strictly STRONGER than a real one for the two +// questions these cases ask. "How many segments did this capture produce?" and +// "in what ORDER did replay emit segments and break functions?" are answered by +// the log directly; on a real device both are inferred from a composed value, +// which a wrong order can still satisfy for a commutative chain (spec +// `## Tests to port` test 12). `Note()` is what puts break functions into the +// SAME log as segments, so one assertion covers the interleaving rather than two +// independent ones that a batched replay satisfies separately. +#pragma once + +#include +#include +#include +#include +#include +#include +#include +#include + +#include "vt/backend.h" +#include "vt/device.h" + +namespace vt_test { + +// Host-memory backend that implements the capture vocabulary +// (`include/vt/backend.h:208-222`) by appending to a call log. A "graph handle" +// is a small heap tag, which is what makes DestroyGraph observable: the seam +// must route every acquisition and release through EndCaptureGraph and +// DestroyGraph so ENG-CUDAGRAPH-DEDUP (#1162) can later interpose at the +// backend (spec `## Risks/decisions` D4). +class RecordingCaptureBackend final : public vt::Backend { + public: + explicit RecordingCaptureBackend(bool supports_capture = true) + : supports_capture_(supports_capture) {} + ~RecordingCaptureBackend() override { + for (void* p : owned_) std::free(p); + for (void* p : tags_) std::free(p); + } + + // ---- vt::Backend ---- + void* Alloc(size_t bytes) override { + void* p = std::malloc(bytes == 0 ? 1 : bytes); + owned_.push_back(p); + return p; + } + // Deliberately does not std::free: cases compare pointer identity across a + // free, and a freed pointer is not a value you may reason about. + void Free(void*) override {} + void Memset(vt::Queue&, void* p, int v, size_t bytes) override { + std::memset(p, v, bytes); + } + void Copy(vt::Queue&, void* dst, const void* src, size_t bytes) override { + std::memcpy(dst, src, bytes); + ++copies_; + } + vt::Queue CreateQueue() override { return vt::Queue{}; } + bool UnifiedMemory() const override { return true; } + + bool SupportsGraphCapture() const override { return supports_capture_; } + void BeginCapture(vt::Queue&) override { + log_.push_back("Begin"); + open_work_.clear(); + capturing_ = true; + } + void* EndCaptureGraph(vt::Queue&) override { + if (fail_next_end_) { + fail_next_end_ = false; + capturing_ = false; + throw std::runtime_error("recording backend: EndCaptureGraph refused"); + } + log_.push_back("EndCaptureGraph"); + void* tag = std::malloc(1); + tags_.push_back(tag); + live_graphs_.push_back(tag); + graph_work_[tag] = std::move(open_work_); + open_work_.clear(); + capturing_ = false; + return tag; + } + void ReplayGraph(vt::Queue&, void* graph) override { + log_.push_back("ReplayGraph"); + replayed_.push_back(graph); + auto it = graph_work_.find(graph); + if (it != graph_work_.end()) + for (const auto& fn : it->second) fn(); + } + void DestroyGraph(void* graph) override { + log_.push_back("DestroyGraph"); + destroyed_.push_back(graph); + graph_work_.erase(graph); + for (size_t i = 0; i < live_graphs_.size(); ++i) { + if (live_graphs_[i] == graph) { + live_graphs_.erase(live_graphs_.begin() + static_cast(i)); + break; + } + } + } + + // ---- the instrument ---- + + // File one captured operation against the open segment. NOT executed now: a + // real capture records the kernel and runs nothing, which is why upstream's + // value assertions all sit AFTER `graph.replay()`. + void Record(std::function fn) { + if (capturing_) + open_work_.push_back(std::move(fn)); + else + fn(); // outside a capture there is nothing to record: run it, as eager work does + } + + // Append a marker into the CALL LOG itself, so a break function's position is + // comparable with the segments around it in one sequence. Two independent + // sequences — segments in the backend log, breaks in a private vector — are + // exactly what a batched replay satisfies while interleaving nothing. + void Note(const std::string& what) { log_.push_back(what); } + + // Arm a single EndCaptureGraph refusal, for the drain case: the one failure a + // destructor cannot let propagate. + void FailNextEndCapture() { fail_next_end_ = true; } + + const std::vector& log() const { return log_; } + const std::vector& replayed() const { return replayed_; } + const std::vector& destroyed() const { return destroyed_; } + size_t live_graphs() const { return live_graphs_.size(); } + int copies() const { return copies_; } + void ClearLog() { + log_.clear(); + replayed_.clear(); + } + // The whole call sequence as one string, so a case can assert the ORDER + // rather than only the counts. + std::string Trace() const { + std::string s; + for (const std::string& e : log_) { + if (!s.empty()) s += " "; + s += e; + } + return s; + } + int Count(const char* what) const { + int n = 0; + for (const std::string& e : log_) + if (e == what) ++n; + return n; + } + + private: + bool supports_capture_; + bool capturing_ = false; + bool fail_next_end_ = false; + std::vector log_; + std::vector owned_; + std::vector tags_; + std::vector live_graphs_; + std::vector replayed_; + std::vector destroyed_; + std::vector> open_work_; + std::map>> graph_work_; + int copies_ = 0; +}; + +} // namespace vt_test diff --git a/tests/vt/test_breakable_graph.cpp b/tests/vt/test_breakable_graph.cpp new file mode 100644 index 000000000..582dd914e --- /dev/null +++ b/tests/vt/test_breakable_graph.cpp @@ -0,0 +1,885 @@ +// PORT of SGLang's breakable-CUDA-graph unit suite, +// `test/registered/cuda_graph/breakable/test_breakable_cuda_graph.py` @ +// `f63458b5be` (305 lines, four classes). AGENTS.md requires the upstream tests +// in the same change that ports the behavior; `.agents/specs/eng-cudagraph-break.md` +// `## Tests to port` maps every upstream case to the local case that owes it, +// and those numbers (T1..T11) are repeated in each case name below. +// +// Row ENG-CUDAGRAPH-BREAK W1, issue #1192, parent #1163. +// +// THE HARNESS ADAPTATION is stated in `tests/vt/recording_capture_backend.h` and +// nowhere else: upstream skips without CUDA, we record the call sequence and +// simulate the graph, so a captured operation is filed at capture and RUN at +// replay. Upstream's arithmetic chains and its post-replay assertions therefore +// port literally — `x.fill_(5); graph.replay(); y == 6` is the same assertion +// here — as do the break counts, the segment counts, the in-place-versus-assign +// split and the non-copyable fallback. +// +// TWO deliberate deviations, named rather than claimed away: +// * Upstream's `eager_on_graph` always returns a fresh tensor, so every +// upstream break is the DESTINATION form. The in-place form has no upstream +// case; it is exercised alongside T2 and by the bare marker. +// * `test_gsm8k_accuracy` (`:288`) is EXCLUDED with the reason recorded in the +// spec: it is a distributional accuracy floor (`mgsm_en >= 0.80`) on a +// PREFILL capture path, and both halves are wrong for us. Our gate polarity +// is bit-exactness against the model's own eager forward (spec `## Gates` +// G1), which is strictly stronger. +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#ifdef __linux__ +#include // readlink("/proc/self/exe"), getpid, for T4's child arm +#endif + +#include "vt/breakable_graph.h" +#include "vt/recording_capture_backend.h" + +namespace { + +using vt::BreakableGraph; +using vt::BreakSlot; +using vt::GraphBreak; +using vt::GraphCaptureScope; +using vt_test::RecordingCaptureBackend; + +// THE SUITE'S OWN PRECONDITION. Every capturing case below asserts the CAPTURING +// lane, so exporting `VLLM_CPP_CUDAGRAPH=0` around a run must not read as a code +// failure: an environment red and a code red are different verdicts and a suite +// that cannot tell them apart reports neither. The kill-switch case is the one +// that deliberately runs in the other lane, and it does so in a CHILD process. +void RequireCaptureLane() { + REQUIRE_MESSAGE(vt::GraphCaptureEnabled(), + "this suite gates the CAPTURING lane; VLLM_CPP_CUDAGRAPH=0 is set"); +} + +// A device-resident int32 vector on the recording backend, so a case can talk +// about a tensor's ADDRESS as well as its contents. The address is the whole +// point of the writeback contract (spec D9): on replay N the eager operation +// returns a FRESH allocation whose address is not the one the next segment +// baked at capture. +struct IntBuf { + vt::Tensor t; + int32_t* host = nullptr; +}; + +IntBuf MakeBuf(RecordingCaptureBackend& b, const std::vector& v) { + IntBuf out; + auto* p = static_cast(b.Alloc(v.size() * sizeof(int32_t))); + std::memcpy(p, v.data(), v.size() * sizeof(int32_t)); + out.host = p; + out.t.data = p; + out.t.dtype = vt::DType::kI32; + out.t.rank = 1; + out.t.shape[0] = static_cast(v.size()); + out.t.stride[0] = 1; + return out; +} + +// Always read THROUGH the tensor view, never through a cached pointer. The +// capture-time semantic of the destination form is that the slot BECOMES the +// first eager result (`:222-227`, upstream's +// `captured_output = _weak_ref_if_tensor(output)`), so a pointer cached before +// the break names a buffer nothing writes to afterwards. +std::vector ReadT(const vt::Tensor& t) { + std::vector v(static_cast(t.shape[0])); + std::memcpy(v.data(), t.data, v.size() * sizeof(int32_t)); + return v; +} +std::vector Read(const IntBuf& b) { return ReadT(b.t); } + +int32_t* Data(IntBuf& b) { return static_cast(b.t.data); } + +void Fill(IntBuf& b, int32_t v) { + for (int64_t i = 0; i < b.t.shape[0]; ++i) Data(b)[i] = v; +} + +// `dst = src + k`, the shape of every captured segment in the upstream cases. +void AddInto(IntBuf& dst, const vt::Tensor& src, int32_t k) { + const std::vector s = ReadT(src); + for (size_t i = 0; i < s.size(); ++i) Data(dst)[i] = s[i] + k; +} + +// An eager break function: reads its input, scales it, and returns a FRESH +// allocation — which is exactly what makes the writeback contract load-bearing. +vt::Tensor FreshScaled(RecordingCaptureBackend& b, const vt::Tensor& src, int32_t mul, + int32_t add) { + std::vector v = ReadT(src); + for (int32_t& e : v) e = e * mul + add; + return MakeBuf(b, v).t; +} + +} // namespace + +// --------------------------------------------------------------------------- +// TestBreakableCUDAGraphBasic (`:30`) — the capture and replay mechanism. +// --------------------------------------------------------------------------- + +// T1 <- test_no_break_capture_replay (`:49-63`). Zero breaks captures and +// replays exactly like a plain graph: `y = x + 1`, refill x to 5, replay, 6. +TEST_CASE("T1 breakable graph: no break yields ONE segment and replays like a plain graph") { + RequireCaptureLane(); + RecordingCaptureBackend b; + vt::Queue q = b.CreateQueue(); + BreakableGraph g; + + IntBuf x = MakeBuf(b, {0, 0, 0, 0}); + IntBuf y = MakeBuf(b, {0, 0, 0, 0}); + { + GraphCaptureScope scope(b, q, g); + REQUIRE(scope.active()); + b.Record([&] { AddInto(y, x.t, 1); }); // captured, NOT run + } + CHECK(g.segment_count() == 1); + CHECK(g.break_count() == 0); + CHECK(g.captured()); + // Exactly one Begin/EndCaptureGraph pair, in that order. + CHECK(b.Trace() == "Begin EndCaptureGraph"); + CHECK(Read(y) == std::vector{0, 0, 0, 0}); // capture executed nothing + + Fill(x, 5); // upstream's x.fill_(5.0) + b.ClearLog(); + g.Replay(q); + CHECK(b.Trace() == "ReplayGraph"); + CHECK(Read(y) == std::vector{6, 6, 6, 6}); // upstream's y == 6.0 + CHECK(g.replay_count() == 1); +} + +// T2 <- test_single_break (`:65-87`). One break splits capture into two segments +// and the chain composes: x=10 -> +1=11 -> eager *2=22 -> +3=25. +TEST_CASE("T2 breakable graph: ONE break yields two segments and the chain composes") { + RequireCaptureLane(); + RecordingCaptureBackend b; + vt::Queue q = b.CreateQueue(); + BreakableGraph g; + + IntBuf x = MakeBuf(b, {0, 0, 0, 0}); + IntBuf mid = MakeBuf(b, {0, 0, 0, 0}); + IntBuf y = MakeBuf(b, {0, 0, 0, 0}); + BreakSlot broken; + + { + GraphCaptureScope scope(b, q, g); + b.Record([&] { AddInto(mid, x.t, 1); }); // segment 0: mid = x + 1 + GraphBreak([&] { return FreshScaled(b, mid.t, 2, 0); }, broken); // the break: eager *2 + b.Record([&] { AddInto(y, *broken, 3); }); // segment 1: y = broken + 3 + } + + CHECK(g.segment_count() == 2); + CHECK(g.break_count() == 1); + CHECK(b.Trace() == "Begin EndCaptureGraph Begin EndCaptureGraph"); + + Fill(x, 10); // upstream's x.fill_(10.0) + g.Replay(q); + // x=10 -> +1=11 -> eager *2=22 -> +3=25, the upstream chain and values. + CHECK(ReadT(*broken) == std::vector{22, 22, 22, 22}); + CHECK(Read(y) == std::vector{25, 25, 25, 25}); +} + +// T3 <- test_multiple_breaks (`:89-115`). Two breaks, three segments, chained: +// x=5 -> +1=6 -> add_one=7 -> +1=8 -> double=16. +TEST_CASE("T3 breakable graph: N breaks yield N+1 segments and the chain composes (N=2)") { + RequireCaptureLane(); + RecordingCaptureBackend b; + vt::Queue q = b.CreateQueue(); + BreakableGraph g; + + IntBuf x = MakeBuf(b, {0, 0, 0, 0}); + IntBuf t1 = MakeBuf(b, {0, 0, 0, 0}); + IntBuf t3 = MakeBuf(b, {0, 0, 0, 0}); + IntBuf y = MakeBuf(b, {0, 0, 0, 0}); + BreakSlot t2, t4; + + { + GraphCaptureScope scope(b, q, g); + b.Record([&] { AddInto(t1, x.t, 1); }); // segment 0: t1 = x + 1 + GraphBreak([&] { return FreshScaled(b, t1.t, 1, 1); }, t2); // break 1: add_one + b.Record([&] { AddInto(t3, *t2, 1); }); // segment 1: t3 = t2 + 1 + GraphBreak([&] { return FreshScaled(b, t3.t, 2, 0); }, t4); // break 2: double + b.Record([&] { AddInto(y, *t4, 0); }); // segment 2: y = t4 + } + CHECK(g.segment_count() == 3); + CHECK(g.break_count() == 2); + CHECK(b.Count("Begin") == 3); + CHECK(b.Count("EndCaptureGraph") == 3); + + Fill(x, 5); // upstream's x.fill_(5.0) + g.Replay(q); + // x=5 -> +1=6 -> add_one=7 -> +1=8 -> double=16, the upstream chain and values. + CHECK(Read(y) == std::vector{16, 16, 16, 16}); +} + +// T4 <- test_eager_on_graph_disabled (`:117-129`). With the wrapper DISABLED the +// function is returned unchanged and runs normally. Our disable switch is item 1 +// of the spec's `## Our baseline`: VLLM_CPP_CUDAGRAPH=0. +// +// IT RUNS IN A CHILD PROCESS, and that is the point. The switch is read once per +// process into a function-local static, so no case can toggle it. Asserting the +// OTHER inert arm instead — a backend that cannot capture — substitutes a +// different conjunct of `active_ = GraphCaptureEnabled() && SupportsGraphCapture()`, +// and deleting `GraphCaptureEnabled()` from that expression left the whole suite +// green. `docs/USAGE.md` sells this switch to users; a switch nothing executes is +// a claim, not a feature. +TEST_CASE("T4 breakable graph: the VLLM_CPP_CUDAGRAPH kill switch makes a capture-capable " + "backend inert") { + const char* sentinel_path = std::getenv("VLLM_CPP_BREAK_KILLSWITCH_SENTINEL"); + if (sentinel_path != nullptr) { + // ---- CHILD ARM: this process was started with VLLM_CPP_CUDAGRAPH=0. ---- + REQUIRE_FALSE(vt::GraphCaptureEnabled()); + RecordingCaptureBackend b(/*supports_capture=*/true); // the backend CAN capture + vt::Queue q = b.CreateQueue(); + BreakableGraph g; + int ran = 0; + { + GraphCaptureScope scope(b, q, g); + REQUIRE_FALSE(scope.active()); // inert anyway: the switch is the only reason + GraphBreak([&] { ++ran; }); + } + REQUIRE(ran == 1); // the break function still runs + REQUIRE(g.segment_count() == 0); + REQUIRE(g.break_count() == 0); + REQUIRE(b.Trace().empty()); // zero backend calls + std::ofstream out(sentinel_path); + out << "INERT-OK"; + return; + } + + // ---- PARENT ARM. ---- +#ifdef __linux__ + char exe[4096] = {0}; + const ssize_t n = ::readlink("/proc/self/exe", exe, sizeof(exe) - 1); + REQUIRE_MESSAGE(n > 0, "cannot resolve this test binary; the child arm cannot be run"); + const std::filesystem::path sentinel = + std::filesystem::temp_directory_path() / + ("vllm_cpp_break_killswitch_" + std::to_string(::getpid()) + ".txt"); + std::filesystem::remove(sentinel); + const std::string cmd = "VLLM_CPP_CUDAGRAPH=0 VLLM_CPP_BREAK_KILLSWITCH_SENTINEL='" + + sentinel.string() + "' '" + std::string(exe) + + "' -tc='*kill switch*' >/dev/null 2>&1"; + const int rc = std::system(cmd.c_str()); + CHECK(rc == 0); + // The sentinel is what defeats "0 cases ran, SUCCESS!": a filter that matched + // nothing exits 0 and writes no file. + REQUIRE(std::filesystem::exists(sentinel)); + std::ifstream in(sentinel); + std::string got; + in >> got; + CHECK(got == "INERT-OK"); + std::filesystem::remove(sentinel); +#else + MESSAGE("kill-switch child arm needs /proc/self/exe; not run on this platform"); +#endif +} + +// T5 <- test_eager_on_graph_outside_capture (`:131-142`). Outside any capture +// the wrapper is a pass-through. +TEST_CASE("T5 breakable graph: OUTSIDE a scope GraphBreak runs fn and calls no backend") { + RecordingCaptureBackend b; + CHECK(GraphCaptureScope::Current() == nullptr); + + int ran = 0; + GraphBreak([&] { ++ran; }); + CHECK(ran == 1); + GraphBreak(); // the bare marker is also a pass-through + CHECK(b.Trace().empty()); + + // The destination form is a pass-through too, and it does NOT fabricate a + // copy: the slot takes the value the break function produced, and its storage + // is never pinned, so a non-capturing forward allocates nothing extra. + BreakSlot out; + GraphBreak([] { return 7; }, out, vt::NoWriteback{}); + CHECK(*out == 7); + CHECK_FALSE(out.pinned()); + CHECK(b.copies() == 0); +} + +// T6 <- test_replay_updates_output (`:144-169`). TWO replays with different +// inputs give different outputs: 3.0, then after x.fill_(10), 33.0. It is the +// upstream anchor for G1's "more than one replay" requirement and the one case +// that can see a break function writing to a stale address — and it asserts the +// break's POSITION in the same log as the segments, not in a second sequence. +TEST_CASE("T6 breakable graph: two replays with different inputs give different outputs") { + RequireCaptureLane(); + RecordingCaptureBackend b; + vt::Queue q = b.CreateQueue(); + BreakableGraph g; + + IntBuf x = MakeBuf(b, {0, 0, 0, 0}); + IntBuf t = MakeBuf(b, {0, 0, 0, 0}); + IntBuf y = MakeBuf(b, {0, 0, 0, 0}); + BreakSlot t2; + { + GraphCaptureScope scope(b, q, g); + b.Record([&] { AddInto(t, x.t, 1); }); // t = x + 1 + GraphBreak( + [&] { + b.Note("scale"); + return FreshScaled(b, t.t, 3, 0); + }, + t2); // eager: t * 3 + b.Record([&] { AddInto(y, *t2, 0); }); // y = t2 + } + REQUIRE(g.segment_count() == 2); + REQUIRE(g.break_count() == 1); + + // First replay: x=0 -> 0+1=1 -> 1*3=3. + b.ClearLog(); + g.Replay(q); + CHECK(b.Trace() == "ReplayGraph scale ReplayGraph"); + CHECK(Read(y) == std::vector{3, 3, 3, 3}); + + // Second replay: x=10 -> 10+1=11 -> 11*3=33. + Fill(x, 10); + b.ClearLog(); + g.Replay(q); + CHECK(b.Trace() == "ReplayGraph scale ReplayGraph"); + CHECK(Read(y) == std::vector{33, 33, 33, 33}); + CHECK(g.replay_count() == 2); +} + +// --------------------------------------------------------------------------- +// TestCopyOutput (`:172`) — the output-writeback contract, spec `## Port map` §3 +// and D9. Upstream tests `_copy_output` (`:172-201`) SEPARATELY from the capture +// machinery because it is a separate guarantee, and so does this. +// --------------------------------------------------------------------------- + +// T7 <- test_tensor_copy (`:187-192`). A tensor destination is written IN PLACE +// and the SAME object is returned (`assertIs(result, dst)`). +TEST_CASE("T7 CopyOutput: a tensor destination is written IN PLACE, address preserved") { + RecordingCaptureBackend b; + vt::Queue q = b.CreateQueue(); + IntBuf dst = MakeBuf(b, {0, 0, 0, 0}); + IntBuf src = MakeBuf(b, {5, 5, 5, 5}); + void* const dst_addr = dst.t.data; + REQUIRE(dst.t.data != src.t.data); + + vt::CopyOutput(b, q, dst.t, src.t); + + CHECK(dst.t.data == dst_addr); // assertIs(result, dst): identity preserved + CHECK(Read(dst) == std::vector{5, 5, 5, 5}); + CHECK(b.copies() == 1); // a real device copy, not a pointer assignment +} + +// The two sides have to AGREE first. Sizing the copy from the destination alone +// makes a break function that returned the wrong shape or dtype an out-of-bounds +// READ of the source, where upstream's `dst.copy_(src)` raises. This is the +// refusal that keeps a mis-registered break point loud (spec D7). +TEST_CASE("T7c CopyOutput: a destination the break output does not match is REFUSED") { + RecordingCaptureBackend b; + vt::Queue q = b.CreateQueue(); + IntBuf dst = MakeBuf(b, {0, 0, 0, 0}); + IntBuf small = MakeBuf(b, {1, 1}); + CHECK_THROWS_AS(vt::CopyOutput(b, q, dst.t, small.t), std::runtime_error); + CHECK(b.copies() == 0); // nothing moved + + IntBuf other = MakeBuf(b, {2, 2, 2, 2}); + other.t.dtype = vt::DType::kF32; // same element count, different dtype + CHECK_THROWS_AS(vt::CopyOutput(b, q, dst.t, other.t), std::runtime_error); + CHECK(b.copies() == 0); +} + +// T8 <- test_dict_copy (`:194-208`). A keyed set of destinations is copied +// field by field. This is the qwen3_5.cpp:9570-9577 shape: logits plus the +// auxiliary hidden taps. +TEST_CASE("T8 CopyOutput: a KEYED set of destinations is written field by field") { + RecordingCaptureBackend b; + vt::Queue q = b.CreateQueue(); + IntBuf da = MakeBuf(b, {0, 0, 0, 0}); + IntBuf db = MakeBuf(b, {0, 0, 0, 0}); + IntBuf sa = MakeBuf(b, {1, 1, 1, 1}); + IntBuf sb = MakeBuf(b, {2, 2, 2, 2}); + void* const da_addr = da.t.data; + void* const db_addr = db.t.data; + + std::map dst{{"a", da.t}, {"b", db.t}}; + const std::map src{{"a", sa.t}, {"b", sb.t}}; + vt::CopyOutput(b, q, dst, src); + + CHECK(dst.at("a").data == da_addr); + CHECK(dst.at("b").data == db_addr); + CHECK(Read(da) == std::vector{1, 1, 1, 1}); + CHECK(Read(db) == std::vector{2, 2, 2, 2}); + CHECK(b.copies() == 2); +} + +// T9 <- test_object_copy (`:210-223`). A struct destination copies its tensor +// fields IN PLACE and ASSIGNS its non-tensor fields (`dst.label == "new"`). +// +// Harness adaptation: upstream reaches the fields through `__dict__` +// (`:182-192`). C++ has no reflection, so a struct participates by providing its +// own CopyOutput overload, found by argument-dependent lookup. The SPLIT the +// case pins — in place for the device field, assign for the rest — is preserved +// exactly. +namespace { +struct BreakStructOut { + vt::Tensor tensor; + std::string label; +}; +void CopyOutput(vt::Backend& b, vt::Queue& q, BreakStructOut& dst, + const BreakStructOut& src) { + vt::CopyOutput(b, q, dst.tensor, src.tensor); // in place + dst.label = src.label; // assigned +} +} // namespace + +TEST_CASE("T9 CopyOutput: a struct copies tensor fields in place and ASSIGNS the rest") { + RecordingCaptureBackend b; + vt::Queue q = b.CreateQueue(); + IntBuf dt = MakeBuf(b, {0, 0, 0, 0}); + IntBuf st = MakeBuf(b, {3, 3, 3, 3}); + void* const dt_addr = dt.t.data; + + BreakStructOut dst{dt.t, "old"}; + const BreakStructOut src{st.t, "new"}; + CopyOutput(b, q, dst, src); + + CHECK(dst.tensor.data == dt_addr); + CHECK(Read(dt) == std::vector{3, 3, 3, 3}); + CHECK(dst.label == "new"); +} + +// T10 <- test_non_tensor_fallback (`:225-227`). With nothing copyable, +// `_copy_output` returns `src` — the documented fallback. Here it must be ASKED +// FOR: `vt::NoWriteback{}` is the opt-in, and without it the destination form +// refuses this type at compile time rather than degrading into the branch whose +// own comment says such a break must use the in-place form. +TEST_CASE("T10 CopyOutput: the non-copyable fallback is legal and must be asked for") { + RequireCaptureLane(); + CHECK_FALSE(vt::detail::HasCopyOutput::value); + CHECK(vt::detail::HasCopyOutput::value); + + RecordingCaptureBackend b; + vt::Queue q = b.CreateQueue(); + BreakableGraph g; + BreakSlot out; + { + GraphCaptureScope scope(b, q, g); + GraphBreak([] { return 42; }, out, vt::NoWriteback{}); + } + CHECK(*out == 42); + CHECK(g.segment_count() == 2); // it still splits + CHECK(g.break_count() == 1); + CHECK(b.copies() == 0); // and it fabricates no copy + + // On replay the break function runs, and — this is the fallback — the seam + // writes NOTHING back, because it has no way to copy into this destination. + // That is upstream's `return src` (`:201`) and it is why such a break must use + // the in-place form instead. Asserting the ABSENCE of the writeback is the + // point: it is the warning a future author needs. + b.ClearLog(); + *out = 0; + g.Replay(q); + CHECK(*out == 0); + CHECK(b.copies() == 0); +} + +// The destination form's REPLAY behavior, which is what D9 is about: on replay +// the eager operation produces a FRESH result and the seam writes it back into +// the capture-time destination, so the following segment's baked address holds +// the new data. A container that replayed the caller's raw `fn` and dropped its +// return value would leave the destination holding capture-time data forever — +// wrong numerics, not a fault. +TEST_CASE("T7b GraphBreak(fn slot): every REPLAY writes back into the capture-time address") { + RequireCaptureLane(); + RecordingCaptureBackend b; + vt::Queue q = b.CreateQueue(); + BreakableGraph g; + + BreakSlot dst; + int32_t feed = 11; + // The break returns a FRESH buffer every call, exactly as an eager operation + // allocating its own output does. + auto fresh = [&] { return MakeBuf(b, {feed, feed, feed, feed}).t; }; + + void* baked = nullptr; + { + GraphCaptureScope scope(b, q, g); + GraphBreak(fresh, dst); + // The destination is now the CAPTURE-TIME result, and this address is what + // the following segment bakes. Upstream does the same: `output = inner(...)` + // then `captured_output = _weak_ref_if_tensor(output)` (`:222-227`). + baked = dst->data; + } + CHECK(baked != nullptr); + CHECK(dst.pinned()); // the seam owns the destination, not this frame + CHECK(ReadT(*dst) == std::vector{11, 11, 11, 11}); + + for (int32_t v : {22, 33, 44}) { + feed = v; + g.Replay(q); + CHECK(dst->data == baked); // the address never moved + CHECK(ReadT(*dst) == std::vector{v, v, v, v}); + } +} + +// The destination OUTLIVES the frame that declared it. This is the rule the +// header used to state in a comment and the first production site broke: a +// `std::optional` local of the function containing the break dies on that +// function's `return`, while the following segment goes on reading its address +// on every replay. With `BreakSlot` the seam owns the storage, so the slot's +// frame can be gone and the writeback still lands. +TEST_CASE("Lifetime rule 1: the destination survives the frame that declared it") { + RequireCaptureLane(); + RecordingCaptureBackend b; + vt::Queue q = b.CreateQueue(); + BreakableGraph g; + int32_t feed = 7; + void* baked = nullptr; + std::vector seen; + + { + GraphCaptureScope scope(b, q, g); + // The slot is a local of THIS block, exactly as the qwen3 site's was a local + // of RunLayer, and it is destroyed before the replay below. + { + BreakSlot inner; + GraphBreak([&] { return MakeBuf(b, {feed, feed}).t; }, inner); + baked = inner->data; + // The following segment BAKES THE ADDRESS, which is what a real capture + // records: a pointer, not a reference to whatever frame declared it. + const vt::Tensor baked_view = *inner; + b.Record([&, baked_view] { seen = ReadT(baked_view); }); // NOT run at capture + } + } + REQUIRE(baked != nullptr); + + feed = 99; + g.Replay(q); + // The break wrote into the cell the closure owns, and the following segment + // read the same address. A destination that died with its frame would have + // been read after its storage was recycled. + CHECK(seen == std::vector{99, 99}); +} + +// The ALIASING half of the SAME rule, which `BreakSlot` did NOT close. The type +// closed the LIFETIME half: the seam owns the storage, so no caller can hand the +// break a destination that dies first. It left the other half expressible. ONE +// slot reused for TWO break points inside one capture compiled, and +// `PinForCapture` returned the cell it had already made, so BOTH replay closures +// wrote through the SAME address: break 0's writeback was overwritten by break +// 1's on every replay, and any segment that baked break 0's destination read +// break 1's data. Measured on the head before this repair, `&*slot` was byte +// identical after break 0 and after break 1, and after `Replay` the slot held +// only break 1's value. That is the shape of the defect the review already +// found once at the production site, with the lifetime half fixed and this half +// open — and W2 through W5 add nine more callers. +// +// The rule is one destination per break point, and registration now REFUSES the +// second. Refusing at `AppendBreak` rather than at the slot is deliberate: every +// form of `GraphBreak` must state its destination to register at all, so a form +// added later cannot opt out by forgetting. +TEST_CASE("Aliasing: one BreakSlot reused for TWO break points in one capture is REFUSED") { + RequireCaptureLane(); + RecordingCaptureBackend b; + vt::Queue q = b.CreateQueue(); + BreakableGraph g; + int32_t feed = 3; + + auto two_breaks_into_one_slot = [&] { + GraphCaptureScope scope(b, q, g); + BreakSlot shared; + GraphBreak([&] { return MakeBuf(b, {feed, feed}).t; }, shared); + GraphBreak([&] { return MakeBuf(b, {feed * 10, feed * 10}).t; }, shared); + }; + CHECK_THROWS_AS(two_breaks_into_one_slot(), std::runtime_error); + // The refusal leaves NO capture behind. The exception propagates out of the + // scope, so the drain sees it and destroys the half-built container rather + // than handing back a forward that replays two segments of three. + CHECK_FALSE(g.captured()); + CHECK(g.break_count() == 0); + CHECK(GraphCaptureScope::Current() == nullptr); + + // CONTROL 1: two DISTINCT slots in the same capture are accepted, and each + // keeps its own destination. A refusal that fired here would be a mute switch + // on the whole destination form. + BreakableGraph g2; + BreakSlot first, second; + { + GraphCaptureScope scope(b, q, g2); + GraphBreak([&] { return MakeBuf(b, {feed, feed}).t; }, first); + GraphBreak([&] { return MakeBuf(b, {feed * 10, feed * 10}).t; }, second); + } + CHECK(g2.segment_count() == 3); + CHECK(g2.break_count() == 2); + CHECK(first->data != second->data); + feed = 5; + g2.Replay(q); + CHECK(ReadT(*first) == std::vector{5, 5}); + CHECK(ReadT(*second) == std::vector{50, 50}); + + // CONTROL 2: the constraint is per CAPTURE, not per slot for life. Reusing one + // slot in a LATER capture is legal, because the later capture bakes its own + // addresses and there is only one closure writing through the cell. + BreakableGraph g3; + { + GraphCaptureScope scope(b, q, g3); + GraphBreak([&] { return MakeBuf(b, {feed + 1, feed + 1}).t; }, first); + } + CHECK(g3.segment_count() == 2); + CHECK(g3.break_count() == 1); + CHECK(ReadT(*first) == std::vector{6, 6}); +} + +// --------------------------------------------------------------------------- +// TestBreakGraphHelper (`:230`). +// --------------------------------------------------------------------------- + +// T11 <- test_break_graph_inserts_segment (`:249-265`). The BARE marker splits +// the segment even though its body does nothing: x=10 -> +1=11 -> break -> +2=13. +TEST_CASE("T11 GraphBreak(): the BARE marker splits the segment and runs nothing") { + RequireCaptureLane(); + RecordingCaptureBackend b; + vt::Queue q = b.CreateQueue(); + BreakableGraph g; + + IntBuf x = MakeBuf(b, {0, 0, 0, 0}); + IntBuf t = MakeBuf(b, {0, 0, 0, 0}); + IntBuf y = MakeBuf(b, {0, 0, 0, 0}); + { + GraphCaptureScope scope(b, q, g); + b.Record([&] { AddInto(t, x.t, 1); }); // segment 0: t = x + 1 + GraphBreak(); // the bare marker + b.Record([&] { AddInto(y, t.t, 2); }); // segment 1: y = t + 2 + } + CHECK(g.segment_count() == 2); + CHECK(g.break_count() == 1); + CHECK(b.Trace() == "Begin EndCaptureGraph Begin EndCaptureGraph"); + + Fill(x, 10); // upstream's x.fill_(10.0) + b.ClearLog(); + g.Replay(q); + // The bare marker's break function is empty, so replay is segments only. + CHECK(b.Trace() == "ReplayGraph ReplayGraph"); + CHECK(Read(y) == std::vector{13, 13, 13, 13}); // x=10 -> 11 -> 13 +} + +// --------------------------------------------------------------------------- +// Tests this row owes with no upstream counterpart, numbered on from the ported +// set in the spec's `## Tests to port`. +// --------------------------------------------------------------------------- + +// Test 12, REPLAY ORDER rather than replay arithmetic. Upstream asserts the +// composed VALUE, which a wrong order can still satisfy for a commutative chain. +// The break markers go into the BACKEND's own log, so segments and breaks form +// ONE sequence: a container that replayed every segment and then ran every break +// would satisfy two independent assertions and fails this one. +TEST_CASE("Test 12: replay emits segment break segment break segment IN THAT ORDER every time") { + RequireCaptureLane(); + RecordingCaptureBackend b; + vt::Queue q = b.CreateQueue(); + BreakableGraph g; + { + GraphCaptureScope scope(b, q, g); + GraphBreak([&] { b.Note("break0"); }); + GraphBreak([&] { b.Note("break1"); }); + } + REQUIRE(g.segment_count() == 3); + REQUIRE(g.break_count() == 2); + // The two break functions ran ONCE eagerly during capture, so the outputs hold + // real data (breakable_cuda_graph.py:222-223). + CHECK(b.Trace() == + "Begin EndCaptureGraph break0 Begin EndCaptureGraph break1 Begin EndCaptureGraph"); + + for (int rep = 0; rep < 3; ++rep) { + b.ClearLog(); + g.Replay(q); + CHECK(b.Trace() == "ReplayGraph break0 ReplayGraph break1 ReplayGraph"); + CHECK(g.replay_count() == rep + 1); + } + CHECK(b.replayed().size() == 3); +} + +// Test 13, the CAPTURE-FAILURE DRAIN. The spec recorded this as "the destructor +// already catches and resets; what is owed is the test", and that was not true: +// the catch only ever guarded a throwing `EndCaptureGraph`. An exception from +// ordinary user code mid-capture left a partial 2-segment graph reporting +// `captured() == true`, which is replayable as HALF A FORWARD. All three arms +// below are the behaviour, not only its test. +TEST_CASE("Test 13a: a break function that THROWS during capture leaves nothing captured") { + RequireCaptureLane(); + RecordingCaptureBackend b; + vt::Queue q = b.CreateQueue(); + BreakableGraph g; + CHECK_THROWS_AS( + [&] { + GraphCaptureScope scope(b, q, g); + GraphBreak([] { throw std::runtime_error("break blew up"); }); + }(), + std::runtime_error); + + CHECK_FALSE(g.captured()); + CHECK(g.segment_count() == 0); + CHECK(g.break_count() == 0); + CHECK(b.live_graphs() == 0); // every instantiated segment released + CHECK(b.Count("EndCaptureGraph") == 1); // and the stream taken out of capture + CHECK(b.Count("DestroyGraph") == 1); + CHECK(GraphCaptureScope::Current() == nullptr); +} + +TEST_CASE("Test 13b: ORDINARY code throwing mid-segment leaves no partial capture") { + RequireCaptureLane(); + RecordingCaptureBackend b; + vt::Queue q = b.CreateQueue(); + BreakableGraph g; + CHECK_THROWS_AS( + [&] { + GraphCaptureScope scope(b, q, g); + GraphBreak(); // one clean break: two segments are now in flight + throw std::runtime_error("model code blew up"); + }(), + std::runtime_error); + + CHECK_FALSE(g.captured()); + CHECK(g.segment_count() == 0); + CHECK(b.live_graphs() == 0); + CHECK(b.Count("EndCaptureGraph") == 2); + CHECK(GraphCaptureScope::Current() == nullptr); +} + +TEST_CASE("Test 13c: a throwing EndCaptureGraph never propagates out of the destructor") { + RequireCaptureLane(); + RecordingCaptureBackend b; + vt::Queue q = b.CreateQueue(); + BreakableGraph g; + { + GraphCaptureScope scope(b, q, g); + GraphBreak(); + b.FailNextEndCapture(); // the FINAL close refuses + } // a destructor that propagated would terminate + CHECK_FALSE(g.captured()); + CHECK(b.live_graphs() == 0); + CHECK(GraphCaptureScope::Current() == nullptr); +} + +// The invariant `segment_count() == break_count() + 1`, asserted rather than +// documented. Re-entering a scope on an already-captured container appended to +// it and produced `seg == brk`, whose last break `Replay` silently drops. +TEST_CASE("Invariant: a scope REFUSES a BreakableGraph that already holds a capture") { + RequireCaptureLane(); + RecordingCaptureBackend b; + vt::Queue q = b.CreateQueue(); + BreakableGraph g; + { + GraphCaptureScope scope(b, q, g); + GraphBreak(); + } + REQUIRE(g.segment_count() == 2); + REQUIRE(g.break_count() == 1); + + CHECK_THROWS_AS(GraphCaptureScope(b, q, g), std::runtime_error); + CHECK(g.segment_count() == 2); // untouched + CHECK(g.break_count() == 1); + CHECK(GraphCaptureScope::Current() == nullptr); + + g.Replay(q); + REQUIRE(g.replay_count() == 1); + + g.Reset(); // the documented way to capture into it again + CHECK_FALSE(g.captured()); + // A released graph's replay count goes with it. A count left behind describes + // a graph that no longer exists, and G3's whole job is to be the number nobody + // has to trust twice. + CHECK(g.replay_count() == 0); + { + GraphCaptureScope scope(b, q, g); + GraphBreak(); + } + CHECK(g.segment_count() == 2); + CHECK(g.break_count() == 1); +} + +// Nesting is REFUSED. `BeginCapture` on a stream that is already capturing is +// CUDA_ERROR_ILLEGAL_STATE on a real backend, and the trace a nested pair emits +// is not a legal capture sequence on any backend. +TEST_CASE("Nesting: a second scope on the same thread is REFUSED before any backend call") { + RequireCaptureLane(); + RecordingCaptureBackend b; + vt::Queue q = b.CreateQueue(); + BreakableGraph outer_g, inner_g; + { + GraphCaptureScope outer(b, q, outer_g); + REQUIRE(outer.active()); + CHECK_THROWS_AS(GraphCaptureScope(b, q, inner_g), std::runtime_error); + // The refusal made no backend call: still exactly one Begin. + CHECK(b.Count("Begin") == 1); + CHECK(b.Count("EndCaptureGraph") == 0); + } + CHECK(outer_g.segment_count() == 1); // the outer scope closed normally + CHECK(inner_g.segment_count() == 0); + CHECK(GraphCaptureScope::Current() == nullptr); +} + +// --------------------------------------------------------------------------- +// Ownership: the container routes every acquisition and release through +// Backend::EndCaptureGraph and Backend::DestroyGraph, treating a segment handle +// as OPAQUE, so ENG-CUDAGRAPH-DEDUP (#1162) can interpose at the backend +// without editing the container (spec `## Risks/decisions` D4). +// --------------------------------------------------------------------------- +TEST_CASE("BreakableGraph: destruction releases EVERY segment through DestroyGraph") { + RequireCaptureLane(); + RecordingCaptureBackend b; + vt::Queue q = b.CreateQueue(); + { + BreakableGraph g; + { + GraphCaptureScope scope(b, q, g); + GraphBreak(); + GraphBreak(); + } + REQUIRE(g.segment_count() == 3); + CHECK(b.live_graphs() == 3); + } + CHECK(b.destroyed().size() == 3); + CHECK(b.live_graphs() == 0); +} + +// Test 14, the NON-CAPTURING BACKEND (spec `## Tests to port`). Vulkan +// (`vulkan_backend.cpp:16`) and Metal (`metal_backend.mm:13`) are the live +// cases: the scope is inert, the forward runs eager, and GraphBreak makes ZERO +// backend calls. This is the OTHER conjunct of `active_`; T4 covers the switch. +TEST_CASE("Test 14: on a backend that cannot capture the scope is inert and the forward is eager") { + RecordingCaptureBackend no_capture(/*supports_capture=*/false); + vt::Queue q = no_capture.CreateQueue(); + BreakableGraph g; + int ran = 0; + { + GraphCaptureScope scope(no_capture, q, g); + CHECK_FALSE(scope.active()); + GraphBreak([&] { ++ran; }); + } + CHECK(ran == 1); // the break function still runs + CHECK(g.segment_count() == 0); + CHECK(g.break_count() == 0); + CHECK(no_capture.Trace().empty()); // zero backend calls +} + +// The G3 observability counters (spec `## Gates` G3): without them there is no +// way to tell a two-segment capture from a fully eager step, and "the graph ran" +// is exactly the claim a broken instrument fabricates. +TEST_CASE("G3: the seam reports segments captured, breaks registered and replays run") { + RequireCaptureLane(); + vt::ResetGraphBreakStats(); + RecordingCaptureBackend b; + vt::Queue q = b.CreateQueue(); + BreakableGraph g; + { + GraphCaptureScope scope(b, q, g); + GraphBreak(); + } + g.Replay(q); + g.Replay(q); + const vt::GraphBreakStats s = vt::GetGraphBreakStats(); + CHECK(s.break_points_reached == 1); + CHECK(s.segments_captured == 2); + CHECK(s.breaks_registered == 1); + CHECK(s.replays == 2); +}