diff --git a/.agents/claims/CLAIM-ENG-CUDAGRAPH-BREAK-W5.md b/.agents/claims/CLAIM-ENG-CUDAGRAPH-BREAK-W5.md
new file mode 100644
index 000000000..e3517d778
--- /dev/null
+++ b/.agents/claims/CLAIM-ENG-CUDAGRAPH-BREAK-W5.md
@@ -0,0 +1,5 @@
+# CLAIM-ENG-CUDAGRAPH-BREAK-W5
+
+| Claim | Row IDs | Agent | Worktree / remote dir | Branch | Owned scope | State | Last update |
+|---|---|---|---|---|---|---|---|
+| `CLAIM-ENG-CUDAGRAPH-BREAK-W5` | `ENG-CUDAGRAPH-BREAK` (`ACTIVE`) | Claude Code (opus-5), fresh implementer | isolated worktree; one queued `rc` job on `thor:gpu0` for the CUDA compile arm and G1; no registry, no publication | `row/ENG-CUDAGRAPH-BREAK-W5`, issue [#1335](https://github.com/mudler/vllm.cpp/issues/1335), parent [#1163](https://github.com/mudler/vllm.cpp/issues/1163), predecessors [#1192](https://github.com/mudler/vllm.cpp/issues/1192) (W1), [#1261](https://github.com/mudler/vllm.cpp/issues/1261) (W2), [#1291](https://github.com/mudler/vllm.cpp/issues/1291) (W3) and [#1307](https://github.com/mudler/vllm.cpp/issues/1307) (W4) | Owns ONLY W5 of [`eng-cudagraph-break.md`](../specs/eng-cudagraph-break.md): migrating the three SINGLE-SHAPE drivers onto `vt::GraphCaptureScope` + `vt::BreakableGraph` in `kFull` — the DFlash draft graph (`src/vllm/model_executor/models/qwen3_dflash.cpp`), the DeepSeek V4 decode graph (`src/vllm/model_executor/models/deepseek_v4.cpp`) and the Laguna decode graph (`src/vllm/model_executor/models/laguna.cpp`); D10, the auxiliary-stream fork/join, as `GraphCaptureScope::NoteFork`/`NoteJoin`, the free `vt::GraphNoteFork`/`vt::GraphNoteJoin`, the auto-join inside `EndSegment`, and the `forks_tracked`/`forks_auto_joined` counters in `include/vt/breakable_graph.h` and `src/vt/breakable_graph.cpp`; routing Laguna's aux fork and join through those hooks; the aux-stream logging and `SupportsAuxStream()` in `tests/vt/recording_capture_backend.h`; test 15 in `tests/vt/test_breakable_graph.cpp`; the gate `tests/vllm/models/test_qwen3_dflash_decode_graph_seam.cpp` and its CMake registration. EXCLUDES the `pure_decode` predicate at `src/vllm/v1/worker/gpu/runner.cpp:1341` and the PIECEWISE arm (W6), the async device-token decline and [#1305](https://github.com/mudler/vllm.cpp/issues/1305) (row-owned, needs a `dgx` window with checkpoints), and graph-executable dedup ([#1162](https://github.com/mudler/vllm.cpp/issues/1162), a parallel branch whose files this claim never touches) | `ACTIVE` | 2026-08-19 (fresh implementer) — all three single-shape drivers migrated, so NINE OF NINE drivers are on the seam and no `BeginCapture`/`EndCaptureGraph`/`ReplayGraph`/`DestroyGraph` CALL survives anywhere in `src/vllm/`. D10 discharged with a five-arm gate and two mutations, and given its production caller in `laguna.cpp`. DFlash gated red-first on CPU (3 cases, 18 assertions) with a G2 mutation that reds ONLY its own file. **G1 and G2 for DeepSeek V4 and Laguna are OWED and the reason is the admission predicate and the build flag, not effort**: V4 refuses a CPU queue and needs the four CUDA-registered kernel families, and Laguna's capture class exists only under `VT_MARLIN_NVFP4`. G5's ROCm and Tenstorrent arms stay BLOCKED on hardware the fleet does not carry. GPU work done through one `rc` job on `thor:gpu0`: G1 RE-RUN at this head (5 cases, 2066 assertions, 0 differing) because D10 changed the shared segment-close path, and a COMPILE-REACHABILITY mutation proving both `#ifdef`-guarded migrated regions are really built (`laguna.cpp:2735` and `deepseek_v4.cpp:1921` each fail the object build under `-Werror`, restoring to an empty diff). Also found, filed and FIXED IN FLOW: [#1352](https://github.com/mudler/vllm.cpp/issues/1352), where W5's own migration made `VLLM_CPP_CUDAGRAPH=0` route the DFlash draft step into an INERT capture lane that ran the forward twice per propose; gated by LANE IDENTITY against a live control in a child process, after the first instrument (allocation counts) turned out to report the ORDER of the runs rather than their shape, because the `DevicePool` free list absorbs the second forward's allocations |
diff --git a/.agents/engine-matrix.md b/.agents/engine-matrix.md
index d20b0eb92..0a67e38b2 100644
--- a/.agents/engine-matrix.md
+++ b/.agents/engine-matrix.md
@@ -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, and the load-bearing gate is byte-identity rather than a ratio | T2 | vLLM has no analogue (its execs come from `torch.compile`, `vllm/config/compilation.py:60-63,517,615,630` @ `555967922`); secondary oracle SGLang `python/sglang/srt/model_executor/runner_backend/cuda_graph_dedup_mixin.py:27-37,105-179,219-242,258-275,353-358` @ `f63458b5be` ([oracles/sglang.md](oracles/sglang.md)) | W1+W2 landing here behind `VT_CUDA_GRAPH_DEDUP`, default OFF until the device A/B measures the per-switch update cost: a device-agnostic dedup registry shared by both accelerator backends plus one CUDA/HIP ops table written once, wired into `EndCaptureGraph`/`ReplayGraph`/`DestroyGraph`. Baseline it replaces: `src/vt/cuda/cuda_backend.cu:222-232` instantiates a fresh exec per capture and destroys the raw graph, 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`) | `tests/vt/test_graph_dedup.cpp` 13/13 cases, 65 assertions, RED-first (written and run against an absent header, and the four cases added by the fresh review of #1178, three of them run against the unfixed source) and gated on every platform via a fake ops table whose launch log makes "the right nodes ran" an observable sequence over MORE than one replay per shape; 13/13 negative mutations detected (9 at implementation, 4 at review repair). That count covers `src/vt/graph_dedup.h` ONLY. `src/vt/graph_dedup_runtime.h` had NO executable coverage on any tier, and [#1184](https://github.com/mudler/vllm.cpp/issues/1184) is what hid in that gap: the file is DESIGNED to see runtime calls fail — a refused `cudaGraphExecUpdate` probe is the feature working — and never consumed the runtime's latched error, so the next unrelated kernel reported the refusal as its own failure and every `VT_CUDA_GRAPH_DEDUP=1` run died 6/6 on GB10 as `greedy_argmax launch: invalid device function` from a launch that had succeeded. Repaired structurally rather than at twelve sites: the clear lives in `ScopedLatchClear`'s destructor (`src/vt/graph_dedup_latch.h`) installed at the six `GraphDedupOps` entry points by `MakeLatchGuardedOps`, the table's only constructor, so no raw function address reaches a field and an unwired seventh operation leaves a null the registry refuses; one line covers CUDA and HIP. The device-free half of the signature walk moved to `src/vt/graph_dedup_signature.h` and is gated by `tests/vt/test_graph_dedup_runtime.cpp` 13/13 cases, 51 assertions, RED-first against the pre-fix guard (22 failed assertions reproducing the production message), 7/7 negative mutations detected — Kahn ordering, topological re-index, sorted edge emission, the depth-4 child bound and the four graph-level escapes. STILL compile-gated only: the five node-payload cases behind the device policy. **DEVICE A/B DELIVERED 2026-08-18 on `dgx:gpu0` (GB10, driver 580.173.02, nvcc 13.0.88, `rc` job f88d484b), and it SPLIT.** Gated commit `72de552c8`, whose four dedup sources are byte-identical to the merged `2a976eb9f` — the row squashed, so the gated tree is not an ancestor of the merge and that sha equality is what carries the claim. CORRECTNESS PASSES: 12/12 cells exit 0, zero `invalid device function` and zero `engine-fatal` in every cell log where the pre-fix head `e4ce5571a` died after exactly one replay, ON replays as often as OFF (60=60, 33=33, 43=43), and `--output-token-ids` is IDENTICAL over 10/10 comparisons with the three OFF/OFF controls passing FIRST and the three workloads hashing to three DIFFERENT values, so the identity is not vacuous. #1184 is closed by this run, because a CPU suite drives a fake runtime and cannot observe the real latched error. THE BENEFIT IS REFUTED for the case this row was filed for: `N == M` in every ON cell — 3 graphs to 3 execs on sizes [24 16 8], 2 to 2 on [16 8], 2 to 2 on [32 24] — with the registry's count CLIMBING 1→1, 2→2, 3→3, so more than one capture reached it and the 1:1 is a measurement rather than the single-capture artefact the first attempt produced. Cause pre-registered before the run and then confirmed, structural rather than a tuning miss: `AppendKernelPayload` hashes (`func`, `gridDim.{x,y,z}`, `blockDim.{x,y,z}`, `sharedMemBytes`) at `src/vt/graph_dedup_runtime.h:121-128` and the memcpy payload hashes the copy extent, so the padded batch dimension sits in the KEY, no candidate group ever forms and `cudaGraphExecUpdate` is NEVER ATTEMPTED. That contradicts this row's own premise — `graph_dedup.h`'s header says the fold is for "two padded batch sizes … the same node topology with different parameters" — and SGLang keys the same fields (`cuda_graph_dedup_mixin.py:105-114`), so whatever folds upstream is not decode buckets either. NO throughput or memory number is recorded: clocks unpinned AND the ON arm allocated exactly as many executables as OFF. Honest gaps: per-shape replay counts are unavailable (the driver prints a TOTAL, so B's ~30-per-shape is arithmetic); the driver's "N captured size(s)" counts SLOTS not captures (A reports 6, emits 3); the container's own cuBLASLt was never re-tested at CUDA 13.0 because the staged cu130 prefix was probed first and worked; only the Qwen3 dense decode driver was exercised. STILL OWED: the default flip, now NOT JUSTIFIED on this evidence rather than merely ungated; a COARSER key that could group two decode buckets at all, which the probe-before-fold design makes a cost question rather than an obviously unsafe one ([#1226](https://github.com/mudler/vllm.cpp/issues/1226), the next traceable hypothesis, deliberately NOT decided by this record); device-tier signature stability/discrimination tests; probing `current_raw` instead of `raws.front()` to retire the update-transitivity assumption; the ROCm compile; a supporting `orin:gpu0` leg, BLOCKED because the Jetson 540.4.0 driver cannot run a CUDA 13 runtime (`cudaGetDeviceCount err=35`); and reaching the feature from the default serving path at all — the async runner captures no decode graph,  **W5, THE SAME DAY, CONFIRMED THE HYPOTHESIS THAT NEGATIVE PRODUCED ([#1226](https://github.com/mudler/vllm.cpp/issues/1226) DELIVERED).** Same box, `rc-worker-4b8lj`, boot_id `3fd9745a-d25a-426c-ba3c-97c958a85515` at both ends, GB10, driver `580.173.02`, `### DONE_AB_KEY 2026-08-18T20:58:46Z`, binary sha256 `ca114abb…c772ad` from `b48b51df1` (tar sha256 asserted before extraction). Drop the launch dimensions and the memcpy extents from the key and every bucket folds: `a_coarse` 3 graphs to 2 execs, `b_coarse` 2 to 1, `c_coarse` 2 to 1, each `probes=1 refused=0`, against `probes=0 refused=0` in every EXACT cell. **`probes=0` in the EXACT cells is the direct process-level proof of W4's source-level diagnosis** — with the launch dimensions in the key no candidate group forms and `cudaGraphExecUpdate` is never asked; drop them and it is asked once per fold and ACCEPTED EVERY TIME. The saving W4 recorded as unreachable is reachable via the key. Byte-identity holds on A (five cells, `59ebff4a…`) and C (four cells, `ff205260…`). **Workload B is VOID rather than a pass, and its cause is a NEW DEFECT that is not this row's:** the two `VT_CUDA_GRAPH_DEDUP`-unset control cells DISAGREED (`5973c5a1…` 2638 bytes vs `4cf79230…` 2650 bytes) on one binary, one workload, greedy `--temperature 0 --seed 777` at `--concurrency 16`, 23 s apart — 672 tokens both, so the byte delta is JSON width and not a length; exactly rows 17 and 18 of 21 differ, both mid-decode, both in the ragged tail `21 % 16` leaves. B's `b_off_a == b_exact` and `b_off_a == b_coarse_a` therefore compare against a baseline that does not reproduce itself and are WORTHLESS; only the OFF/OFF control made that visible, and without it B would have read as three more confirmations. Filed [#1283](https://github.com/mudler/vllm.cpp/issues/1283). **Caveats that bound this result:** nvcc was `13.3.73` here and `13.0.88` for the W4 baseline the recorded dgx gate stack names, so the OFF-vs-ON and EXACT-vs-COARSE comparisons WITHIN this binary are valid while this run and that baseline are NOT directly comparable; clocks unpinned (2405 MHz current, 3003 max, 2418 applications) and nothing measured bytes, so NO throughput and NO memory number is claimed or implied; only the Qwen3 dense decode driver was exercised; `refused=0` is ONE driver on ONE hardware and toolkit pair, which is no more a floor than W4's negative was a ceiling; and the coarse key is behind `VT_CUDA_GRAPH_DEDUP_COARSE_KEY`, default OFF, inside a default-OFF flag, on **PR [#1232](https://github.com/mudler/vllm.cpp/pull/1232) which is STILL A DRAFT — nothing on `main` folds today.** **Row stays `ACTIVE`, argued:** not `DONE`, because the fold is unreachable on every shipping configuration and the row's stated MEMORY saving has never been measured in bytes on either key; not `PARTIAL`, because nothing upstream is omitted — the coarse key is our own extension past SGLang, which keys the fields we started from; not `BLOCKED`, because nothing external stops the next step. What is owed is now a DECISION about the default plus the byte measurement and the probe-cost-at-real-churn measurement it needs, and landing #1232 first **W6, 2026-08-19, THE DEVICE-BYTE MEASUREMENT — THE BENEFIT QUESTION IS NOW CLOSED AND THE ANSWER IS NEGATIVE.** Tested `origin/main` `2c8f53d93`, which is PR #1232 LANDED, so the "nothing on `main` folds today" caveat every earlier record carried is RETIRED and this measures a configuration that ships. Same box, `rc` job `93f783de`, pod `rc-worker-4b8lj`, boot_id `3fd9745a-…` at BOTH ends, GB10, driver `580.173.02`, nvcc **13.0.88** (the W4 baseline toolkit; W5 ran 13.3.73, so W6 and W5 are NOT directly comparable while comparisons WITHIN this one binary are valid), binary sha256 `be697268…0ce657a7`, `### DONE_BYTES 2026-08-19T04:57:19Z`, 12/12 cells exit 0, zero VOID markers. **THE FOLD ENGAGES AT THE SHIPPED BUCKET SET**, which is the churn W5 could not produce: `vllm-bench` sets `max_num_seqs = concurrency`, so W32 captured `[1 2 4 8 16 24 32]` 7-of-7 and W64 captured `[1 … 64]` 11-of-11, exactly `decode_graph_sizes.h:32-41`, against the 2-3 buckets every earlier conclusion was drawn from. COARSE folds 7 graphs to 3 execs (`probes=7 refused=3`) and 11 to 5 (`probes=22 refused=16`); EXACT folds NOTHING at `probes=0`, reproducing W4 at four times the bucket count. Token ids byte-identical across every cell of a workload INCLUDING both OFF/OFF controls (`ff0db6c6…be9d` 11720 B; `e1cbf5fc…e5d0` 57620 B) — neither workload has #1283's ragged-tail shape and neither hit it. **THE SAVING DOES NOT SURVIVE ITS OWN NULL CONTROL.** `nvidia-smi --query-compute-apps` tail median (the `--query-gpu=memory.used` axis returns `[N/A]` on this box) shows W64 IDENTICAL to the megabyte in all five cells (9737) and W32's coarse arm reading 10-23 MiB HIGHER than OFF (3252/3262 vs 3262/3275). A `cudaMemGetInfo` shim summed over every instantiate gives a nominal 13.83 MiB at 7 buckets — **0.42% of a 3.25 GiB process** — and **−0.75 MiB, i.e. NOTHING, at 11**. That nominal effect is NOT ESTABLISHED on four independent grounds: `EXACT` is a TRUE NULL (same 7 and 11 retained execs, `probes=0`, so it allocates what OFF allocates) and disagrees with OFF by 10.6-13.1 MiB against a 13.83 MiB candidate; the W64 OFF/OFF pair disagrees with ITSELF by 18.2 MiB; one instantiate recorded a NEGATIVE delta (`-5,165,056` B); and `cudaGraphExecDestroy` reclaimed `0` in EVERY cell. Per-instantiate deltas for byte-identical 404-node graphs span 0 to 10,514,432 B and 17 of 27 instantiates in one cell read exactly zero, so these are POOL-GRANULAR readings and the coarse arm's throwaway probes grow that pool exactly like retained execs do. What CAN be priced: one ~390-node executable at **2.08-4.35 MiB**, 10.0-10.6 KB per node — the figure to re-run on a deep checkpoint. **THE MECHANISM INVERTS THIS ROW'S PREMISE.** The driver refuses **43% of probes at 7 buckets and 73% at 11**, every one of them `probe refused a fold (err=910 result=2)` = `cudaErrorGraphExecUpdateFailure` / `cudaGraphExecUpdateErrorTopologyChanged`. The shim's `cudaGraphGetNodes` reading says why false candidates form: the decode graphs are **TWO topologies, 376 and 404 nodes**, mixed across the buckets (`w32_off_a` captured `404 404 376 376 404 404 404`). Every refusal is about TOPOLOGY, never a parameter, so a COARSER key produces MORE false hits rather than more folds — the opposite of what W5's 2-bucket A/B suggested, and W5's `refused=0` is now explained as an artefact of workloads whose buckets only ever SHRANK, so exactly one pair was ever presented. **COST:** W32 OFF 7 instantiates / 0 updates vs COARSE 10 (3 retained + 7 probes) / 11 updates; W64 OFF 11 / 0 vs COARSE **27** (5 retained + 22 probes) / 28 updates — **2.45x the instantiate calls** to retain 6 fewer executables. **Peak transient did NOT double** — in every ON cell live-bytes peak == end, because `Register` destroys the probe before returning, so the feared "double the peak to save the steady state" trade did not occur. **A replay-time re-point DID occur** — 4 and 6 non-probe updates over 88 and 244 replays, ARITHMETIC over two printed totals and not a counter — with every cell exiting 0 and byte-identical, so `Replay`'s transitivity assumption neither aborted nor changed a token; W5 recorded that case as untested. **CAVEATS THAT BOUND THIS RESULT:** the clock pin was **REFUSED inside the lease** (`The current user does not have permission to change clocks for GPU 0000000F:01:00.0`, `clocks_pinned=0`), so **NO time-based figure is attributable** and the instantiate-wall and update-wall figures in `bytes.log` are diagnostics quoted nowhere as a result; `result=2` is ONE driver, ONE GB10, ONE toolkit; only the Qwen3 dense decode driver was exercised, as in W4 and W5; `VT_ASYNC_RUNNER=0` throughout, so the feature is STILL unreachable on the DEFAULT serving path (#1179); and `cudaMemGetInfo` cannot separate an executable's own cost from the pool chunk that satisfied it. **VERDICT, DELIVERED AND NEGATIVE:** `VT_CUDA_GRAPH_DEDUP` stays default OFF, now on MEASUREMENT rather than on silence; `VT_CUDA_GRAPH_DEDUP_COARSE_KEY` alone is a **NO-OP, not merely unsupported** — `GraphDedupCoarseKeyEnabled()` (`src/vt/graph_dedup.h:114`) is read only by the signature builder (`src/vt/graph_dedup_runtime.h:177`), only from `Register`, only under `GraphDedupEnabled()` (`src/vt/cuda/cuda_backend.cu:237`), so with dedup off its sole observable is one stderr line; both on is unsupported. **NOT A CEILING.** Three things would change it and each is traceable: find where the 376/404 split comes from (the FA-2 split-KV grid is the first suspect — a capture that fixes the node set across buckets removes every refusal); an instrument that resolves a single 2-4 MiB executable against driver pool granularity (`cuMemGetAllocationGranularity` or a pool-statistics query); and the same measurement on a 60-80 layer checkpoint, where bytes scale with node count. **Row STAYS `ACTIVE`, argued, and the argument is now narrow.** The MEASUREMENT obligations are discharged and the DECISION is delivered, which is the `DONE` case and it is a real one. Three things stop the flip and none is a checker technicality: the feature is unreachable on the DEFAULT serving path, owned by `ENG-CUDAGRAPH-BREAK` (#1179) and the "nothing lands dead" half of this row; two items still sit under #1162 itself — the device-tier signature stability/discrimination tests and probing `group.current_raw` instead of `raws.front()` to retire the transitivity assumption; and the `DONE` record surface owes a `.agents/parity-ledger.md` entry, a closing-commit owner in place of the claim, an exact test anchor and the RELEASE of `CLAIM-ENG-CUDAGRAPH-DEDUP`, which is an operator act and which this record-only branch does not own. Not `PARTIAL` — nothing upstream is omitted. Not `BLOCKED` — nothing external stops the next step. Full evidence: [benchmark-record.md](benchmark-record.md) entry `ENG-CUDAGRAPH-DEDUP W6`, raw at `/mnt/nas_share/rc/dedup-bytes/` | [eng-cudagraph-dedup.md](specs/eng-cudagraph-dedup.md); analysis [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `ACTIVE` | `CLAIM-ENG-CUDAGRAPH-DEDUP` ([#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`'s `DenseDecodeGraphForward` 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 did NOT meet it and says so**: no `rc` lease was obtainable in its window and a CPU harness cannot replay a captured segment, so it moves to W3 with the three drivers of the same shape (G1); 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, which becomes live only for the first PIECEWISE production capture (W4); 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. **W2 DONE 2026-08-18 ([#1261](https://github.com/mudler/vllm.cpp/issues/1261)): `Qwen3DenseDecodeGraph` MIGRATED and the seam is ENTERED from a production step**, which retires W1's staged slice. `Qwen3DenseDecodeGraph::Step` opens a `vt::GraphCaptureScope` over a per-slot `vt::BreakableGraph` and replays through `BreakableGraph::Replay`; the hand-rolled `BeginCapture`/`EndCaptureGraph` pair, the raw `void*` handle, the `bool captured` flag, the `DestroyGraph` loop and the driver's own `VLLM_CPP_CUDAGRAPH` read are gone (re-derivation items 1, 2, 5, 6). The migration ADDED `vt::GraphCaptureMode`, mirroring vLLM's `CUDAGraphMode` (`vllm/config/compilation.py:59-63`), whose v1 default `FULL_AND_PIECEWISE` (`:63`) is documented at `:630-632` as a FULL graph for DECODE batches and a piecewise one for prefill/mixed, with `decode_mode()` (`:65-66`) selecting the full half and the runtime reading it at `vllm/v1/worker/gpu/cudagraph_utils.py:185-186`. A decode driver opened `kPiecewise` would have turned a fully graphed decode step into ONE EAGER ATTENTION CALL PER LAYER between graph replays — not vLLM's decode behaviour, and invisible to every token gate here. `GraphBreak` in a `kFull` scope takes the pass-through arm and `AppendBreak` REFUSES a registration in that mode. G2 is `tests/vllm/models/test_qwen3_decode_graph_seam.cpp` (3 cases / 124 assertions), which asserts the SEAM's counters because a driver calling `Backend::ReplayGraph` directly leaves an identical backend log; the mutation restoring the pre-W2 raw pair (18 lines, compiled clean) left `test_breakable_graph` 27/27, `test_qwen3_break_point` 2/2 and `test_qwen3_forward` 10/10 GREEN and reddened only this file. G4 in the same file: capture step vs `Qwen3DenseModel::Forward`, 100 logits, 0 differing. **The async decline at `qwen3.cpp` STANDS and is now GATED in both arms**: migrating the capture does not move the INPUTS, so the depth-2 race is untouched, and the fix is `StepDevInputs` as a SEAM capability, which is W4. **G1 is NOT met by W2** and is recorded owed rather than implied. **W3 DONE 2026-08-19 ([#1291](https://github.com/mudler/vllm.cpp/issues/1291)): the three remaining PLAIN BATCHED drivers migrate — `Qwen3MoeDecodeGraph`, `VoxtralDecodeGraph`, `DeepseekV2DecodeGraph` — one commit each, each with its own RED-first G2 gate.** Four of the nine drivers are now on the seam, and the six batched-driver `VLLM_CPP_CUDAGRAPH` reads `## Our baseline` item 1 counted are down to TWO, both in `qwen3_5.cpp` (W4). Each gate asserts the SEAM's counters and not the backend log, because a driver that kept its raw pair produces identical logits, an identical backend log and an identical `replay_count()`; red-first on four assertions each (`test_qwen3_moe_decode_graph_seam` 222/226, `test_voxtral_decode_graph_seam` 224/228, `test_deepseek_v2_decode_graph_seam` 224/228, all exit 1), green 3/3 each after. The G2 mutation — restoring each pre-W3 driver file, 25/102, 23/92 and 25/94 lines, each compiled clean — reddens ONLY its own gate and leaves `test_breakable_graph` 216/216 and W2's `test_qwen3_decode_graph_seam` 231/231 green. The gate harness is now SHARED (`tests/vllm/models/decode_graph_seam_harness.h`); three more copies inside `tests/` would have reproduced the duplication this row removes from `src/`. **G1 IS DELIVERED and is no longer owed** — the item W1 and W2 both carried. `tests/vllm/models/test_decode_graph_seam_g1_cuda.cpp` runs each driver COLD, CAPTURE and THREE consecutive replays against its own eager arm (selected by `max_num_reqs == 0`, so both arms are one binary on one device rather than two builds, each with its OWN device KV cache) on `thor:gpu0` through an `rc` lease — NVIDIA Thor sm_110, driver 595.78, nvcc 13.0.88, source `c905bb536`, 32 `.cu.o` objects, binary resolving `libcudart.so.13`/`libcublasLt.so.13`: **3 cases, 1600 assertions, exit 0, `5 steps x 100 logits, 0 differing, 4 replays` per driver.** The COUNT carries that claim, not the status line: with no CUDA backend the same file prints `SUCCESS!` over `assertions: 0`. Bounded honestly — synthetic tiny models rather than a checkpoint, and W2's driver shares the seam by argument rather than by measurement. **W3 also found a gate that could not fail.** The three gates' `breaks_registered == 0` mode guard is a TAUTOLOGY for any model with no registered break point, and the one production `vt::GraphBreak` in the tree is W1's in `qwen3.cpp`: flipping `kFull` to `kPiecewise` in `qwen3_moe.cpp`, one token, compiled clean and left that gate GREEN at 226/226. The mode was UNOBSERVABLE from outside a driver, so `vt::GraphBreakStats` gains `full_scopes`/`piecewise_scopes`, counted in `GraphCaptureScope`'s constructor on the ACTIVE path only, with an inert-scope control; the same flip now reds all three gates on exactly those two assertions. **NO break point is registered in these three models, deliberately**: under `kFull` it would be pass-through machinery no gate can exercise, and the break-point set is what the PIECEWISE arm needs (W4/W6). **The async decline, per driver:** Voxtral needs none (its only construction site is `VoxtralGenerateGreedy`, unreachable from the runner); Qwen3-Coder and DeepSeek carry a NEW FINDING instead — `qwen3_moe_registry.cpp:107`, `deepseek_v2_registry.cpp:106` and `glm4_moe_lite_registry.cpp:125` route an async step into a host-vector replay with no `device_token_ids` check at all, filed [#1305](https://github.com/mudler/vllm.cpp/issues/1305) with W4 as owner rather than mitigated on a measurement W3 cannot make. G5's ROCm/Tenstorrent arm is NOT discharged and moves to W5: the fleet carries no such device, so it is blocked on hardware rather than unattempted. **W4 DONE 2026-08-19 ([#1307](https://github.com/mudler/vllm.cpp/issues/1307)): the persistent device input path becomes a SEAM CAPABILITY, and the two Qwen3.5 drivers migrate.** `vt::PersistentStepInput` (`include/vt/persistent_step_input.h`, `src/vt/persistent_step_input.cpp`) binds a capture-stable device destination the DRIVER owns together with its pinned host staging block, and refreshes it in place from a host source or a DEVICE one; it owns the address-stability rule as a REFUSAL, the staging block, and the refreshing ARM as an observable (`last_source()`, `vt::StepInputStats`), and deliberately NOT the device allocation, because `Qwen3_5DecodeGraph` draws its retained inputs from a DEDICATED `DevicePool` so they never pop a block the captured forward's scratch then needs (D3). RED-first against a stub with the declared API and no guarantees: `tests/vt/test_persistent_step_input.cpp` 9 cases / 0 passed / 59 assertions / 32 failed / exit 1, GREEN after at 9/9 and 59/59; three mutations (delete the capacity refusal, make a null device source a silent no-op, collapse the host arm out of staging) each compiled clean and each reds exactly one case. `Qwen3_5DecodeGraph` and `Qwen3_5DenseDecodeGraph` open a `vt::GraphCaptureScope` over a per-slot `vt::BreakableGraph` in `kFull` and replay through it, and their `PinnedStepInputs`/`StageStepInputs` staging now runs THROUGH the capability, which is what makes it reachable rather than a class with a unit test. **Six of the nine drivers are on the seam** and `grep -rn 'std::getenv("VLLM_CPP_CUDAGRAPH")' src/` returns exactly ONE line, `src/vt/breakable_graph.cpp:61` — one switch, at last. Gate `tests/vllm/models/test_qwen3_5_decode_graph_seam.cpp` RED-first on the MoE driver's five seam assertions (3 cases / 62 assertions / 5 failed / exit 1) and GREEN after at 7/7 and 129, G4 reading `40 values, 0 differing` per driver; G2 mutations: the whole pre-W4 file restored reds BOTH drivers (296 lines, 10 assertions), the MoE replay bypassing the container reds ONLY the MoE case (7 lines), the MoE `kFull`->`kPiecewise` flip reds ONLY its mode counters (3 lines), and deleting the `StageStepInputs` call site reds ONLY the reachability case while `test_persistent_step_input` stays 59/59 green — the difference between a class that works and a capability something reaches. **W4 FALSIFIED THIS ROW'S OWN PREMISE, which is its most important result.** This record and the spec both said the fix `qwen3.cpp`'s `DenseDecodeGraphForward`'s decline names already existed as `StepDevInputs`. It does not: `StepDevInputs` has NO token-id member, and its pinned sibling `PinnedStepInputs::token_ids` was allocated at capture, filled every step, zeroed by the poison hook, and NEVER uploaded or read — the embed runs OUTSIDE the captured region from the HOST vector in every batched driver, so **the decode graph carries no token ids to the device in ANY driver**. The dead block is removed. Consequently the DECLINE STANDS and [#1305](https://github.com/mudler/vllm.cpp/issues/1305) STAYS OPEN: W4 also read the decline's recorded cause against the tree at its own parent and found it falsified (the `DeviceTokenIdsScope` WAS live on the graph path, consumed by `EmbedInto` on all three arms at `qwen3.cpp:610,621,644 @ 338cbbfd1^`), so the measured failure is real and its mechanism is unidentified — not a state from which a refactor may retire a mitigation. The async battery was NOT run and W4 says so plainly: it needs `dgx` WITH the Qwen3-0.6B/4B checkpoints, `dgx:gpu0` was held by another session for W4's whole window, and W4's lease was `thor:gpu0`. Still NO throughput claim. Still NO throughput claim; analysis [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `ACTIVE` | `CLAIM-ENG-CUDAGRAPH-BREAK-W4`; [#1163](https://github.com/mudler/vllm.cpp/issues/1163), [#1192](https://github.com/mudler/vllm.cpp/issues/1192), [#1261](https://github.com/mudler/vllm.cpp/issues/1261), [#1291](https://github.com/mudler/vllm.cpp/issues/1291), [#1307](https://github.com/mudler/vllm.cpp/issues/1307), [#1305](https://github.com/mudler/vllm.cpp/issues/1305), [#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`'s `DenseDecodeGraphForward` 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 did NOT meet it and says so**: no `rc` lease was obtainable in its window and a CPU harness cannot replay a captured segment, so it moves to W3 with the three drivers of the same shape (G1); 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, which becomes live only for the first PIECEWISE production capture (W4); 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. **W2 DONE 2026-08-18 ([#1261](https://github.com/mudler/vllm.cpp/issues/1261)): `Qwen3DenseDecodeGraph` MIGRATED and the seam is ENTERED from a production step**, which retires W1's staged slice. `Qwen3DenseDecodeGraph::Step` opens a `vt::GraphCaptureScope` over a per-slot `vt::BreakableGraph` and replays through `BreakableGraph::Replay`; the hand-rolled `BeginCapture`/`EndCaptureGraph` pair, the raw `void*` handle, the `bool captured` flag, the `DestroyGraph` loop and the driver's own `VLLM_CPP_CUDAGRAPH` read are gone (re-derivation items 1, 2, 5, 6). The migration ADDED `vt::GraphCaptureMode`, mirroring vLLM's `CUDAGraphMode` (`vllm/config/compilation.py:59-63`), whose v1 default `FULL_AND_PIECEWISE` (`:63`) is documented at `:630-632` as a FULL graph for DECODE batches and a piecewise one for prefill/mixed, with `decode_mode()` (`:65-66`) selecting the full half and the runtime reading it at `vllm/v1/worker/gpu/cudagraph_utils.py:185-186`. A decode driver opened `kPiecewise` would have turned a fully graphed decode step into ONE EAGER ATTENTION CALL PER LAYER between graph replays — not vLLM's decode behaviour, and invisible to every token gate here. `GraphBreak` in a `kFull` scope takes the pass-through arm and `AppendBreak` REFUSES a registration in that mode. G2 is `tests/vllm/models/test_qwen3_decode_graph_seam.cpp` (3 cases / 124 assertions), which asserts the SEAM's counters because a driver calling `Backend::ReplayGraph` directly leaves an identical backend log; the mutation restoring the pre-W2 raw pair (18 lines, compiled clean) left `test_breakable_graph` 27/27, `test_qwen3_break_point` 2/2 and `test_qwen3_forward` 10/10 GREEN and reddened only this file. G4 in the same file: capture step vs `Qwen3DenseModel::Forward`, 100 logits, 0 differing. **The async decline at `qwen3.cpp` STANDS and is now GATED in both arms**: migrating the capture does not move the INPUTS, so the depth-2 race is untouched, and the fix is `StepDevInputs` as a SEAM capability, which is W4. **G1 is NOT met by W2** and is recorded owed rather than implied. **W3 DONE 2026-08-19 ([#1291](https://github.com/mudler/vllm.cpp/issues/1291)): the three remaining PLAIN BATCHED drivers migrate — `Qwen3MoeDecodeGraph`, `VoxtralDecodeGraph`, `DeepseekV2DecodeGraph` — one commit each, each with its own RED-first G2 gate.** Four of the nine drivers are now on the seam, and the six batched-driver `VLLM_CPP_CUDAGRAPH` reads `## Our baseline` item 1 counted are down to TWO, both in `qwen3_5.cpp` (W4). Each gate asserts the SEAM's counters and not the backend log, because a driver that kept its raw pair produces identical logits, an identical backend log and an identical `replay_count()`; red-first on four assertions each (`test_qwen3_moe_decode_graph_seam` 222/226, `test_voxtral_decode_graph_seam` 224/228, `test_deepseek_v2_decode_graph_seam` 224/228, all exit 1), green 3/3 each after. The G2 mutation — restoring each pre-W3 driver file, 25/102, 23/92 and 25/94 lines, each compiled clean — reddens ONLY its own gate and leaves `test_breakable_graph` 216/216 and W2's `test_qwen3_decode_graph_seam` 231/231 green. The gate harness is now SHARED (`tests/vllm/models/decode_graph_seam_harness.h`); three more copies inside `tests/` would have reproduced the duplication this row removes from `src/`. **G1 IS DELIVERED and is no longer owed** — the item W1 and W2 both carried. `tests/vllm/models/test_decode_graph_seam_g1_cuda.cpp` runs each driver COLD, CAPTURE and THREE consecutive replays against its own eager arm (selected by `max_num_reqs == 0`, so both arms are one binary on one device rather than two builds, each with its OWN device KV cache) on `thor:gpu0` through an `rc` lease — NVIDIA Thor sm_110, driver 595.78, nvcc 13.0.88, source `c905bb536`, 32 `.cu.o` objects, binary resolving `libcudart.so.13`/`libcublasLt.so.13`: **3 cases, 1600 assertions, exit 0, `5 steps x 100 logits, 0 differing, 4 replays` per driver.** The COUNT carries that claim, not the status line: with no CUDA backend the same file prints `SUCCESS!` over `assertions: 0`. Bounded honestly — synthetic tiny models rather than a checkpoint, and W2's driver shares the seam by argument rather than by measurement. **W3 also found a gate that could not fail.** The three gates' `breaks_registered == 0` mode guard is a TAUTOLOGY for any model with no registered break point, and the one production `vt::GraphBreak` in the tree is W1's in `qwen3.cpp`: flipping `kFull` to `kPiecewise` in `qwen3_moe.cpp`, one token, compiled clean and left that gate GREEN at 226/226. The mode was UNOBSERVABLE from outside a driver, so `vt::GraphBreakStats` gains `full_scopes`/`piecewise_scopes`, counted in `GraphCaptureScope`'s constructor on the ACTIVE path only, with an inert-scope control; the same flip now reds all three gates on exactly those two assertions. **NO break point is registered in these three models, deliberately**: under `kFull` it would be pass-through machinery no gate can exercise, and the break-point set is what the PIECEWISE arm needs (W4/W6). **The async decline, per driver:** Voxtral needs none (its only construction site is `VoxtralGenerateGreedy`, unreachable from the runner); Qwen3-Coder and DeepSeek carry a NEW FINDING instead — `qwen3_moe_registry.cpp:107`, `deepseek_v2_registry.cpp:106` and `glm4_moe_lite_registry.cpp:125` route an async step into a host-vector replay with no `device_token_ids` check at all, filed [#1305](https://github.com/mudler/vllm.cpp/issues/1305) with W4 as owner rather than mitigated on a measurement W3 cannot make. G5's ROCm/Tenstorrent arm is NOT discharged and moves to W5: the fleet carries no such device, so it is blocked on hardware rather than unattempted. **W4 DONE 2026-08-19 ([#1307](https://github.com/mudler/vllm.cpp/issues/1307)): the persistent device input path becomes a SEAM CAPABILITY, and the two Qwen3.5 drivers migrate.** `vt::PersistentStepInput` (`include/vt/persistent_step_input.h`, `src/vt/persistent_step_input.cpp`) binds a capture-stable device destination the DRIVER owns together with its pinned host staging block, and refreshes it in place from a host source or a DEVICE one; it owns the address-stability rule as a REFUSAL, the staging block, and the refreshing ARM as an observable (`last_source()`, `vt::StepInputStats`), and deliberately NOT the device allocation, because `Qwen3_5DecodeGraph` draws its retained inputs from a DEDICATED `DevicePool` so they never pop a block the captured forward's scratch then needs (D3). RED-first against a stub with the declared API and no guarantees: `tests/vt/test_persistent_step_input.cpp` 9 cases / 0 passed / 59 assertions / 32 failed / exit 1, GREEN after at 9/9 and 59/59; three mutations (delete the capacity refusal, make a null device source a silent no-op, collapse the host arm out of staging) each compiled clean and each reds exactly one case. `Qwen3_5DecodeGraph` and `Qwen3_5DenseDecodeGraph` open a `vt::GraphCaptureScope` over a per-slot `vt::BreakableGraph` in `kFull` and replay through it, and their `PinnedStepInputs`/`StageStepInputs` staging now runs THROUGH the capability, which is what makes it reachable rather than a class with a unit test. **Six of the nine drivers are on the seam** and `grep -rn 'std::getenv("VLLM_CPP_CUDAGRAPH")' src/` returns exactly ONE line, `src/vt/breakable_graph.cpp:61` — one switch, at last. Gate `tests/vllm/models/test_qwen3_5_decode_graph_seam.cpp` RED-first on the MoE driver's five seam assertions (3 cases / 62 assertions / 5 failed / exit 1) and GREEN after at 7/7 and 129, G4 reading `40 values, 0 differing` per driver; G2 mutations: the whole pre-W4 file restored reds BOTH drivers (296 lines, 10 assertions), the MoE replay bypassing the container reds ONLY the MoE case (7 lines), the MoE `kFull`->`kPiecewise` flip reds ONLY its mode counters (3 lines), and deleting the `StageStepInputs` call site reds ONLY the reachability case while `test_persistent_step_input` stays 59/59 green — the difference between a class that works and a capability something reaches. **W4 FALSIFIED THIS ROW'S OWN PREMISE, which is its most important result.** This record and the spec both said the fix `qwen3.cpp`'s `DenseDecodeGraphForward`'s decline names already existed as `StepDevInputs`. It does not: `StepDevInputs` has NO token-id member, and its pinned sibling `PinnedStepInputs::token_ids` was allocated at capture, filled every step, zeroed by the poison hook, and NEVER uploaded or read — the embed runs OUTSIDE the captured region from the HOST vector in every batched driver, so **the decode graph carries no token ids to the device in ANY driver**. The dead block is removed. Consequently the DECLINE STANDS and [#1305](https://github.com/mudler/vllm.cpp/issues/1305) STAYS OPEN: W4 also read the decline's recorded cause against the tree at its own parent and found it falsified (the `DeviceTokenIdsScope` WAS live on the graph path, consumed by `EmbedInto` on all three arms at `qwen3.cpp:610,621,644 @ 338cbbfd1^`), so the measured failure is real and its mechanism is unidentified — not a state from which a refactor may retire a mitigation. The async battery was NOT run and W4 says so plainly: it needs `dgx` WITH the Qwen3-0.6B/4B checkpoints, `dgx:gpu0` was held by another session for W4's whole window, and W4's lease was `thor:gpu0`. Still NO throughput claim. W5 DONE 2026-08-19 ([#1335](https://github.com/mudler/vllm.cpp/issues/1335)): the THREE SINGLE-SHAPE drivers migrate — the DFlash draft graph, the DeepSeek V4 decode graph and the Laguna decode graph, whose own note at `laguna.cpp:2116-2119` asked for this seam by name and named V4's as the sibling that moves with it. **NINE OF NINE DRIVERS ARE ON THE SEAM and the migration is COMPLETE**: a call-shaped grep over `src/vllm/` for `BeginCapture`, `EndCaptureGraph`, `ReplayGraph` and `DestroyGraph`, with comment lines excluded, returns NOTHING. The three per-model rollback switches stay (each an A/B lever for one driver); `VLLM_CPP_CUDAGRAPH` reaches all three for the first time. **D10, the auxiliary-stream fork/join, is DISCHARGED and REACHED** — `GraphCaptureScope` owns the outstanding-fork set and joins it before `EndCaptureGraph` (port of `breakable_cuda_graph.py:353-361` plus the `wait_stream` hook `:101-153`), registered by `vt::GraphNoteFork`/`GraphNoteJoin` from `laguna.cpp:2572-2576,2612`, the only fork inside a captured region by construction. Every prior stage opened `kFull`, which has ONE segment and so no between-segments window, so the rule could not be exercised before W5 and untested machinery was not landed for it. Gated as a COUNTER and an ORDER out of one backend trace, five arms including the control where the model joins first, and two mutations (deleting the join reds only the new case on 5 assertions; making it over-fire reds it on 8). DFlash is the ONE single-shape driver gateable without a GPU, because its admission predicate names neither a device type nor a kernel registry: `test_qwen3_dflash_decode_graph_seam.cpp` RED-first 3 cases/0 passed/16 assertions/7 failed exit 1, GREEN after 3/18, and the G2 mutation reds ONLY that file while seven other suites — the driver's own `test_dflash_propose` included — stay green. **G1 RE-RUN at W5's head on `thor:gpu0`** (sm_110, driver 595.78, nvcc 13.0.88, 32 `.cu.o`, source `79dc6b5bd`) because D10 put a join on the path of EVERY segment close, so the seam changed underneath the five measured drivers: `test_decode_graph_seam_g1_cuda` 5 cases / 2066 assertions / 0 failed, each reading `0 differing, 4 replays`, plus `test_breakable_graph` 265 on the same device. **And the one thing a green build could NOT have told us was measured separately**: Laguna's capture class sits behind `#ifdef VT_MARLIN_NVFP4`, so a passing build is the SAME OBSERVATION as one that compiled the region out. `-DVT_MARLIN_NVFP4=1` is on `laguna.cpp`'s own compile command, and an undeclared identifier injected immediately after its `GraphCaptureScope` line FAILED the object build under `-Werror` (`laguna.cpp:2735`) against an rc-0 baseline, restoring to an empty diff; the identical mutation on V4 failed at `deepseek_v4.cpp:1921`. Both migrated regions are COMPILED, which retires the could-not-even-be-built half. **G1 for all three and G2 for V4 and Laguna are OWED on hardware**, per driver and per reason: V4's `CanRunResidentDecode` refuses `kCPU` and needs the four CUDA-registered kernel families, Laguna's capture class exists only under `VT_MARLIN_NVFP4`. G5's ROCm/Tenstorrent arm stays BLOCKED — the fleet is all NVIDIA — and its owner moves from W5 to the ROW. Still NO throughput claim; analysis [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `ACTIVE` | `CLAIM-ENG-CUDAGRAPH-BREAK-W5`; [#1163](https://github.com/mudler/vllm.cpp/issues/1163), [#1192](https://github.com/mudler/vllm.cpp/issues/1192), [#1261](https://github.com/mudler/vllm.cpp/issues/1261), [#1291](https://github.com/mudler/vllm.cpp/issues/1291), [#1307](https://github.com/mudler/vllm.cpp/issues/1307), [#1305](https://github.com/mudler/vllm.cpp/issues/1305), [#1020](https://github.com/mudler/vllm.cpp/issues/1020), [#1335](https://github.com/mudler/vllm.cpp/issues/1335)|
 | `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. **The unblock order now has an owning row:** `LTX25-DEVICE-RESIDENCY` ([#1264](https://github.com/mudler/vllm.cpp/issues/1264), [ltx25-device-residency.md](specs/ltx25-device-residency.md)) stages those defects W0-W6 and carries this decision point as its W7 — if the loop comes back GPU-bound, #1164 closes as a refutation the way [#1161](https://github.com/mudler/vllm.cpp/issues/1161) closed prefill capture | [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 a3e019058..773a21ec0 100644
--- a/.agents/issue-index.md
+++ b/.agents/issue-index.md
@@ -429,6 +429,9 @@ rather than merged. `scripts/check-agent-record.py` gates both.
 | [#1325](https://github.com/mudler/vllm.cpp/issues/1325) | `ENG-RECORD-ANCHOR-RATCHET` | `scripts/record-anchor-baseline.json` stores a `"total"` that no code reads. `load_record_anchor_baseline` in `scripts/check-agent-record.py` returns `{bucket: int(data["buckets"][bucket]) for bucket in RECORD_ANCHOR_BUCKETS}` with `RECORD_ANCHOR_BUCKETS = ("stale", "broken")`, and it is the file's only reader; `check_record_anchors` iterates those two buckets, and `write_record_anchor_baseline`'s refusal compares `result.total > sum(previous.values())` — the buckets, not the stored `total`. Measured at `af87251c5`: mutating `"total": 38` to `39` (a file whose total disagrees with `32 + 6`) leaves `check-agent-record.py` exit 0, tree restored byte-for-byte by sha256. So this row's own budget file carries the exact shape the row exists to name: a recorded figure no gate reads, sitting beside the figures that are read and presenting as if it were checked. `--write-baseline` compounds it by printing `-> 38`, which reads as the value it stored and is the one no later run consults. TWO candidate resolutions, deliberately not chosen here because choosing belongs to the fixing row: read it and assert `total == stale + broken` on load, or drop the field and derive it at read time, which is the shape AGENTS.md §Records prefers. Either is a semantic change to `check-agent-record.py` owing a spec and a red-before case in `tests/scripts/test_agent_record.py` `RecordAnchorRatchet`. Distinct from [#1287](https://github.com/mudler/vllm.cpp/issues/1287) and [#1270](https://github.com/mudler/vllm.cpp/issues/1270), neither of which reaches the unread field. Also under `## Owed` in [`record-anchor-ratchet.md`](specs/record-anchor-ratchet.md) | bug |
 | [#1316](https://github.com/mudler/vllm.cpp/issues/1316) | — | `scripts/main-baseline.py` renders a scheduled run that executed ZERO jobs as `RED` with all 11 covered jobs `missing`, so `NEWEST BASELINE: RED at <sha>` names a tree the run never checked out. Measured at `origin/main` `250db75a2`: runs `32206456661` and `32140419182` both return `startedAt: null` for every job, because GitHub cancelled them while they were pending in the single `ci-schedule-refs/heads/main-mudler/vllm.cpp` group, whose queue holds one run ([#274](https://github.com/mudler/vllm.cpp/issues/274)). Fail-closed, and the `missing (expected, never ran)` line is accurate about the jobs; the defect is the verdict word, because a run that executed nothing is NOT RUN rather than RED, and the newest verdict should fall through to the newest run that actually ran. NOT fixed in flow: `test_an_expected_job_the_payload_never_mentions_is_red` and `test_a_narrowed_run_reports_red_and_names_what_never_ran` deliberately assert missing-is-red so a narrowed run cannot pass, and separating "narrowed" from "never started" changes what the verdict means, which owes its own spec, red-before evidence and a fresh reviewer. Owed under `## Owed` of [baseline-lane-eviction.md](specs/baseline-lane-eviction.md), which removes the only observed producer of a zero-job run | bug |
 | [#1314](https://github.com/mudler/vllm.cpp/issues/1314) | `SPEC-DFLASH2` | **DFlash2 (`DFlash2DraftModel`) is unported**, and one config rule would run the published checkpoint wrong in silence. Upstream carries DFlash2 as a SECOND architecture beside DFlash rather than as a change to it ([vllm#52816](https://github.com/vllm-project/vllm/pull/52816), OPEN at head `19c9351904df4c63042671bc67a866ca48dc7d6f`, base `9842d701`, 755+/5-, 11 files, plus the stacked guard fix [vllm#52883](https://github.com/vllm-project/vllm/pull/52883)): DFlash1 gains two subclass seams and keeps every behaviour, and the new architecture adds a GROUPED DYNAMIC DEPTHWISE CONVOLUTION around each attention and each MLP sublayer plus a CANDIDATE SELECTOR that replaces the independent per-slot argmax with a scored path walk over the target head's top-K. Shapes taken from the published checkpoint rather than from the diff: `z-lab/Qwen3.8-27B-DFlash2`, safetensors header range-read 2026-08-19, 81 tensors -- DFlash1's set plus `layers.N.{attention,mlp}_conv.{base_kernel (2,2,5120), kernel_projection.weight (1280,5120)}` x5 and `candidate_selector.{hidden_projection.weight (256,5120), predecessor_codebook, successor_codebook}` at `(248320,256)` bf16 each, ~254 MB resident the DFlash1 lane never allocates; `conv_kernel_size 2`, `conv_group_size 16`, `selector_rank 256`, `selector_top_k 16`, `block_size 8`. **The silent one:** that config declares all five layers `sliding_attention` AND `is_causal false`, while our resolution mirrors the OLD upstream rule (causal iff SWA, unless `dflash_config.causal`, `include/vllm/model_executor/models/qwen3_dflash.h:22-24`), so every layer would run CAUSAL -- plausible tokens, a token gate against our own output sees nothing, and only ACCEPTANCE moves, which the lossless verify hides. Upstream changes `_dflash_layer_causal` to read `is_causal` first, in the same commit. Three further things are owed and none of them is silent: no route for the `DFlash2DraftModel` architecture string (the same classification code as the open `DSparkDraftModel` gap, [#1193](https://github.com/mudler/vllm.cpp/issues/1193)); no top-k that EMITS the surviving (id, value) pairs, where the decision is to extend the sort-free pivot-bracket threshold search already ported from the same FlashInfer approach at `src/vt/cuda/cuda_sample.cu:297-506` rather than port FlashInfer's 3380-line general radix kernel; and the path walk must run ON DEVICE from the first landing, because the identical sequential shape in DSpark shipped host-side and measured 28% of the 27B draft step ([#436](https://github.com/mudler/vllm.cpp/issues/436)) before `SampleSequentialDevice` moved it. Already reusable unchanged: `vt::DFlashBlockAttention`, the DFlash runner/rejection/GDN-rollback lane, and the loader's target-shared `embed_tokens`+`lm_head`, which is already what a DFlash2 checkpoint needs. BEYOND-PIN by developer decision 2026-08-19 (mirror the open PR now, reconcile if review moves it), in the same posture `SPEC-DSPARK-QWEN3-ROUTING` takes toward vllm#52197; the parity pin `555967922` does not carry the architecture and is NOT advanced. Gate arm is bf16 27B plus the GGUF drafter arm in the same wave, oracle = vLLM built at the PR head, acceptance measured SAME-TRAJECTORY because `SPEC-DFLASH` D8 spent a campaign on a divergent-trajectory confound that D9 refuted. Spec [dflash2-spec-decode.md](specs/dflash2-spec-decode.md) | feature |
+| [#1335](https://github.com/mudler/vllm.cpp/issues/1335) | `ENG-CUDAGRAPH-BREAK` | W5, the LAST migration stage of the break-point capture seam: the three SINGLE-SHAPE drivers move onto `vt::GraphCaptureScope` + `vt::BreakableGraph` in `kFull` — the DFlash draft graph (`qwen3_dflash.cpp`), the DeepSeek V4 decode graph (`deepseek_v4.cpp`) and the Laguna decode graph (`laguna.cpp`), whose own note at `:2116-2119` asked for this seam by name and named V4's driver as the sibling that moves with it. **NINE OF NINE DRIVERS ARE NOW ON THE SEAM**, and a call-shaped grep over `src/vllm/` for `BeginCapture`, `EndCaptureGraph`, `ReplayGraph` and `DestroyGraph`, with comment lines excluded, returns NOTHING — every surviving textual hit is prose: the last three hand-rolled `BeginCapture`/`EndCaptureGraph` pairs, their raw `void*` handles, two hand-written `DestroyGraph` destructors and DFlash's private `try`/drain are gone. The three per-model rollback switches STAY (`VT_V4_DECODE_GRAPH`, `VT_DFLASH_GRAPH`, `VT_LAGUNA_DECODE_GRAPH`) because each is a same-binary A/B lever for exactly one driver; what changes is that `VLLM_CPP_CUDAGRAPH` reaches all three for the first time, through `vt::GraphCaptureEnabled()` inside the scope. **D10, the auxiliary-stream fork/join, is DISCHARGED here and could not have been discharged earlier**: `GraphCaptureScope` now owns the set of side queues forked since the current segment opened and joins every outstanding one before `Backend::EndCaptureGraph`, which is the port of `_end_current_segment` (`breakable_cuda_graph.py:353-361` @ SGLang pin `f63458b5be`) plus the `wait_stream` hook (`:101-153`) whose only purpose is to populate that set. W1's break point was on a model that forks nothing and W2/W3/W4 all opened `kFull`, which has ONE segment and so no segment CLOSE inside a fork window; W5 owns the only driver whose fork is inside the captured region by construction (`laguna.cpp:2572-2576` fork, `:2612` join), and that driver is the rule's production caller through `vt::GraphNoteFork`/`vt::GraphNoteJoin` — no monkey-patch, because our fork is an explicit `RecordEvent`/`QueueWaitEvent` pair rather than an implicit torch call. **The rule is a COUNTER and an ORDER, not prose**, because W3 measured a mode guard that stayed green at 226/226 under the mutation it named: `vt::GraphBreakStats` gains `forks_tracked` and `forks_auto_joined`, and the assertion is the single backend trace `Begin RecordEvent QueueWaitEvent EndCaptureGraph …`, since closing a capture with an unjoined fork FAILS at `cudaStreamEndCapture` and "the join happened BEFORE the close" is the whole claim. Five arms including the CONTROL where the model joins first and the scope must do nothing; two mutations, both compiled clean — deleting `JoinOutstandingForks()` reds only the new case on 5 assertions, and making it over-fire reds it on 8. **DFlash is the one single-shape driver gateable without a GPU, and the reason is its admission predicate**: it names neither a device type nor a kernel registry, while V4's `CanRunResidentDecode` refuses `kCPU` outright and again without the four CUDA-registered V4 kernel families, and Laguna's whole capture class is behind `#ifdef VT_MARLIN_NVFP4` (a CUDA build on a marlin-nvfp4 arch — `sm_12xa` or `sm_110`). So `tests/vllm/models/test_qwen3_dflash_decode_graph_seam.cpp` lands RED-first (3 cases, 0 passed, 16 assertions, 7 failed, exit 1 against the unmigrated driver) and GREEN after (3 cases, 18 assertions), with the G2 mutation restoring the pre-W5 raw pair reddening ONLY that file while `test_breakable_graph` (265), `test_qwen3_decode_graph_seam` (231), `test_qwen3_moe_decode_graph_seam` (228), `test_voxtral_decode_graph_seam` (230), `test_deepseek_v2_decode_graph_seam` (230), `test_qwen3_5_decode_graph_seam` (129) and the driver's OWN `test_dflash_propose` (31) all stay green — that last one being the point, since a driver's bit-exactness suite cannot see which capture machinery ran. **G1 and G2 for V4 and Laguna are OWED on hardware and named rather than implied**, as is G5's ROCm/Tenstorrent arm, which the fleet carries no device for. Still NO throughput claim anywhere: prefill is 3.8% host-idle at >96% GPU-busy, the 27B prefill gap is 92.5% non-GEMM glue, decode is already captured, and the sibling row `ENG-CUDAGRAPH-DEDUP` closed its benefit question negative after measurement | feature |
+| [#1352](https://github.com/mudler/vllm.cpp/issues/1352) | `ENG-CUDAGRAPH-BREAK` | The DFlash draft step ran TWO full forwards per propose under `VLLM_CPP_CUDAGRAPH=0`. CAUSED by W5's migration ([#1335](https://github.com/mudler/vllm.cpp/issues/1335)) and FIXED IN THE SAME FLOW. `ForwardBlockLogitsWithDeviceKV`'s paged predicate was `UseDflashGraph() && d.b.SupportsGraphCapture() && support_static_graph_mode()`, which was complete while the driver's capture was its own `BeginCapture` pair — nothing outside the file could turn that off. W5 routes the capture through `vt::GraphCaptureScope`, which reads `VLLM_CPP_CUDAGRAPH` itself, so the switch gained a state the predicate did not know about: with it set to `0` on a capture-capable backend the driver still chose the CAPTURE lane, ran its eager warm pass, opened an INERT scope and ran the whole `ForwardPagedBody` a SECOND time inside it — two full draft forwards per propose, forever, because `g_state` could never reach 2. NOT a correctness defect (an inert scope executes the forward eagerly, so the logits are real and identical) and confined to a non-default diagnostic lane, which is precisely why nothing would have found it later. Fix: `vt::GraphCaptureEnabled()` becomes the third conjunct, so the switch selects the driver's existing single-forward eager path, which is what it means for the other eight drivers. **The gate is unusual because the defect is invisible to every ordinary one**: both versions emit IDENTICAL logits and both leave `segments_captured` at 0, so neither a token gate nor the seam's counters can separate them. The assertion is therefore LANE IDENTITY against a live control — a backend that CANNOT capture takes the eager path by construction, so a capture-capable backend with the switch OFF must do exactly the same work — and it runs in a CHILD PROCESS because `vt::GraphCaptureEnabled()` is read once per process into a function-local static. **The first instrument was WRONG and the harness now records why**: counting `Backend::Alloc` read `control=11, switch-off=0`, and would have read those same two numbers whichever lane ran second, because the `DevicePool` serves the second of two identical forwards out of its free list without calling `Alloc` at all — an instrument reporting the ORDER of the runs rather than their shape, which would have failed the CORRECT program. `CaptureCapableCpuBackend` counts `Copy` instead, per-operation work the pool cannot absorb, and gains a `supports_capture` flag so `StaticGraphCpu(false)` produces the control lane. Red-first on the unfixed driver: `DIFFERENT-LANE`, `control copies=4` against `switch-off copies=5`, exit 1, one assertion, compiled clean at 22 insertions and 3 deletions — and that same run reported the logits `IDENTICAL`, which is the measurement showing why nothing else could have caught it. Green after at 4 cases / 23 assertions, with all seven sibling seam gates unchanged | bug |
+| [#1361](https://github.com/mudler/vllm.cpp/issues/1361) | `ENG-CUDAGRAPH-BREAK` | [`eng-cudagraph-break.md`](specs/eng-cudagraph-break.md) credited the W1 exit criterion to GB10 `sm_121a`, and the same file says it ran on `orin:gpu0`. Found by W5 ([#1335](https://github.com/mudler/vllm.cpp/issues/1335)) while reading `## Owed` to record what that stage did and did not discharge; record-only, no product behaviour. The G5 entry closed with "the seam's CUDA arm now runs on TWO architectures rather than one, sm_110 here and sm_121a on GB10 for the W1 exit criterion", while `## Work breakdown` W1 in the same file records the criterion as "measured on `orin:gpu0` through an `rc` lease, driver `12060`" — a Jetson AGX Orin, which is neither a GB10 nor `sm_121a`. Wrong in both halves: it named a device the measurement did not run on and an architecture nothing in this row has measured the criterion against. The two-architecture claim itself SURVIVES and only the attribution was wrong: the exit criterion (`cudaStreamEndCapture` then `cudaStreamBeginCapture` mid-forward with eager work between) ran on `orin:gpu0`, and G1 plus the unit suite ran on `thor:gpu0` at sm_110 for W3, W4 and W5. The criterion has NOT been re-measured on `thor`, for a structural reason rather than an omission: every migrated driver opens `kFull`, so nothing in the tree re-begins a capture mid-forward, and G1 exercises capture and replay rather than the re-begin. `sm_121a` on GB10 is OWED and not done — W5 could not take it because `dgx:gpu0` was held by another session for its whole window. This is the shape where a number quoted often starts being treated as measured, which is why it is filed rather than quietly reworded. FIXED IN FLOW by the W5 pull request, which states what ran where and moves `sm_121a` into `## Owed` | bug |
 | [#1357](https://github.com/mudler/vllm.cpp/issues/1357) | `BACKEND-CUDA-COMP-FA` | `src/vllm/platforms/cuda.cpp::CudaPlatform::supports_fa2_attention` returns `true` unconditionally for every CUDA device, and is consumed at `src/vllm/model_executor/models/qwen3_5.cpp:5163` to gate `fa2_prefill`/`fa2_decode` and therefore `attn_dt` — bf16 FlashAttention-2 versus the f32 graph-captured fallback, on the default decode path. `CMakeLists.txt:169` defaults `VLLM_CPP_CUDA_ARCHITECTURES` to `121a` alone and `cmake/CudaArchFeatures.cmake:349` narrows FA2 to the intersection with the `fa2` row `8.0,8.6,8.7,8.9,12.0a,12.1a`, so a default build carries FA2 cubins for ONE architecture while the predicate promises FA2 to all of them; run it on an sm_86 card and the model takes a path with no SASS for the device. Same class of claim that made vLLM select an unrunnable FlashAttention on a GB10 ([#1332](https://github.com/mudler/vllm.cpp/issues/1332)), where requesting `FLASHINFER` generates text and exits 0 while the default resolves `FLASH_ATTN` and dies at the first attention call with `cudaErrorUnsupportedPtxVersion`. Fixable here in a way upstream cannot manage, because the answer already exists in the build: `vt_cuda_feature_archs(VT_FA2_ARCHS "fa2")` (`CMakeLists.txt:499`) computes the compiled FA2 arch set and `CMakeLists.txt:2279` passes that SAME variable to `vt_cuda_set_source_gencode` to emit the `-gencode` flags, so a `configure_file` manifest is DERIVED from the flag computation rather than tracking it. A hand-written list would be `CUDA_SUPPORTED_ARCHS` again and is refused. Spec [`cuda-compiled-arch-manifest.md`](specs/cuda-compiled-arch-manifest.md) | bug |
 | [#1354](https://github.com/mudler/vllm.cpp/issues/1354) | — | Clock pinning is UNAVAILABLE inside an `rc` lease. `nvidia-smi -lgc 2190` returns `LGC_RC=4`, "The current user does not have permission to change clocks for GPU 0000000F:01:00.0", in three separate `rc run` jobs on `dgx:gpu0` on 2026-08-19, each running as **root** in the worker pod. `.agents/benchmarking.md` instructs "Pin the clocks before measuring, under the lock" and ships `sudo nvidia-smi -lgc 2100` as the recipe, and **every clock-pinned figure in this repository was taken over the host + `ssh` + `flock` path** that AGENTS.md now forbids for a fleet device — so the migration to leases silently removed clock pinning and no record said so. Same class as [#1265](https://github.com/mudler/vllm.cpp/issues/1265): a capability the records assume, which the current access path does not provide. Measured cost: nine timed windows across two arms recorded within-run SM-clock spreads of 12.92% to 26.36% against the 5% ceiling, `SwThermalSlowdown` active in every one and `HwSlowdown+HwThermal` in one, so `gpu_clock_state compare` returned `PAIRING_VERDICT=DISCARD` on all three Qwen3.8-27B c1 pairings **even though the cross-arm rule passed perfectly** (same boot, both arms 2489 MHz median, 0.0% offset). The cell therefore has two clean complete absolutes and no ratio. NOT fixed in flow: the fix is either an `rc` worker capability this row has no authority over, or a demonstrated settle-and-hold procedure, or a ratified different clock rule for lease-measured pairs — each its own spec, and none of them a widening of the assertion to turn a red green. Records updated meanwhile in `.agents/environment.md`, `.agents/benchmarking.md` and `.agents/benchmark-record.md`. Owed under `## Owed` in [bench-qwen38-27b-four-way.md](specs/bench-qwen38-27b-four-way.md) | bug |
 | [#1355](https://github.com/mudler/vllm.cpp/issues/1355) | — | Our server reports 5,942 prompt tokens where vLLM reports 6,144 for the IDENTICAL generated prompts. Found 2026-08-19 in the raw `vllm bench serve --save-detailed` files of the Qwen3.8-27B bf16 re-measure ([#915](https://github.com/mudler/vllm.cpp/issues/915)), both arms driven by the byte-identical client invocation from the same pinned wheel `0.1.dev1+g555967922`, same dataset, same seed. `input_lens` is the SERVER-reported length — `vllm/benchmarks/lib/endpoint_request_func.py:247` overwrites `output.prompt_len` from the streamed `usage.prompt_tokens` — and reads `[915, 931, 1024, 1024, 1024, 1024]` for us against `[1024] x 6` for vLLM at c1, with 19 of 48 short (877-941) at c8, byte-identical across all three reps of each leg. Not the client: `_align_prompts_to_server_tokenizer` (`vllm/benchmarks/serve.py:74,2041-2044`) re-aligns against the server's own `/tokenize` and prints `WARNING: tokenizer mismatch` when it disagrees, and NEITHER arm printed it, so our `/tokenize` agreed on 1024 while our `usage.prompt_tokens` reported 915 for the same request. `output_lens` is `[128]xN` on both arms in every leg, so the campaign's output-throughput, TPOT and ITL figures are unaffected; `total_token_throughput` is affected, our c8 196.10 tok/s being computed over 47,072 input tokens where the intended workload is 49,152. TWO causes and the artifacts cannot separate them: under-reported usage, or a genuinely truncated prompt — and the second would mean the two arms did not run the same workload. A greedy token gate cannot see either, which is why it survived the gate on this checkpoint. NOT fixed in flow: the finding row writes no product code and holds no GPU. Owed under `## Owed` in [qwen38-27b-bf16-gate.md](specs/qwen38-27b-bf16-gate.md) | bug |
diff --git a/.agents/specs/eng-cudagraph-break.md b/.agents/specs/eng-cudagraph-break.md
index f664cc7cb..65a679dbf 100644
--- a/.agents/specs/eng-cudagraph-break.md
+++ b/.agents/specs/eng-cudagraph-break.md
@@ -604,13 +604,53 @@ G1 Qwen3_5DecodeGraph     on CUDA: 5 steps x  40 logits, 0 differing, 4 replays
 G1 Qwen3_5DenseDecodeGraph on CUDA: 5 steps x 40 logits, 0 differing, 4 replays
 ```
 
+**G1 WAS RE-RUN AT W5's HEAD on `thor:gpu0` through an `rc` lease, and it is a
+re-run rather than a new case.** The seam changed UNDER the five measured
+drivers: D10 put a `JoinOutstandingForks()` call on the path of EVERY segment
+close, so each migrated driver's capture now executes seam code W4's run did not,
+and a change to the shared close path is exactly what a CPU harness cannot clear.
+Same box and same provenance as W3 and W4 — NVIDIA Thor, sm_110, driver 595.78,
+nvcc 13.0.88, `-DVLLM_CPP_CUDA=ON -DVLLM_CPP_CUDA_ARCHITECTURES=110`, 32 `.cu.o`
+objects — at source `79dc6b5bd`:
+`tests/vllm/models/test_decode_graph_seam_g1_cuda.cpp` ran **5 cases, 2066
+assertions, 0 failed, exit 0**, every driver again reading `0 differing, 4
+replays`. On the same device and in the same job: the full CUDA library built
+clean (`lib_rc=0`, 742 targets), `tests/vt/test_breakable_graph.cpp` ran 30 cases
+and 265 assertions, and `tests/vllm/models/test_qwen3_dflash_decode_graph_seam.cpp`
+ran **4 cases, 23 assertions, exit 0** — so W5's own driver gate, including the
+#1352 child-process arm, is green on a CUDA build and not only on this box's CPU
+one. That last line is a build-and-routing result rather than a replay one: the
+DFlash case still runs against the CPU harness inside that binary, which is why
+G1 for that driver stays owed below.
+
+**THE ONE THING A GREEN BUILD COULD NOT HAVE TOLD US was measured separately.**
+Laguna's capture class is inside `#ifdef VT_MARLIN_NVFP4`, so "the CUDA build
+compiled `laguna.cpp`" is satisfied just as well by a build that compiled the
+migrated region OUT — success and failure are the SAME OBSERVATION. Two facts
+settle it. `-DVT_MARLIN_NVFP4=1` is on `laguna.cpp`'s own compile command in
+`compile_commands.json`, and CMake reported `CUDA feature marlin-nvfp4: ENABLED
+for [110]`. And the region itself was MUTATED: an undeclared identifier inserted
+immediately after `vt::GraphCaptureScope scope(b, q, graph, kFull)` FAILED the
+object build under `-Werror` (`laguna.cpp:2735`,
+`'VT_W5_MUTATION_THIS_MUST_NOT_COMPILE' was not declared in this scope`) against
+a baseline object build of rc 0, and the tree restored byte for byte to an empty
+`git diff`. The identical mutation on DeepSeek V4 failed at
+`deepseek_v4.cpp:1921`. **Both migrated regions are compiled.** That is the half
+of their coverage answerable without the models' own kernels and checkpoints; G1
+and G2 for those two remain owed, and `## Owed` states what each needs.
+
+**W5 ADDED NO G1 CASE OF ITS OWN, and says so rather than leaving it to be
+inferred from a count.** Its three drivers are single-shape and each refuses the
+harness for its own reason.
+
 The two Qwen3.5 cases needed a cache pool the other three did not:
 `CudaGdnCachePool` allocates the RECURRENT ssm and conv state on device beside
 the paged KV, and each arm gets its own — the GDN recurrence advances its state
 every step, so two arms sharing one state would step each other's recurrence and
-the agreement would measure nothing. **Five of the six migrated drivers are now
-covered by measurement**; W2's `Qwen3DenseDecodeGraph` still is not, and still
-shares the seam by argument rather than by measurement.
+the agreement would measure nothing. **Five of the NINE migrated drivers are
+covered by measurement**; W2's `Qwen3DenseDecodeGraph` and W5's three
+single-shape drivers are not, and share the seam by argument rather than by
+measurement.
 
 **The W3 run this extends, kept because it is the earlier evidence** (W3, #1291,
 2026-08-19, the same `thor:gpu0`) — NVIDIA Thor, sm_110, driver 595.78, nvcc
@@ -1032,10 +1072,101 @@ slot 0 captures on step THREE.
 **The async decline STANDS, and the reason is sharper than "the fix is not built
 yet".** See `## Owed`.
 
-**W5, migrate the three single-shape drivers.** DeepSeek V4, Laguna and DFlash.
-`laguna.cpp:2116-2119` already carries the note that its capture class is waiting
-for exactly this seam, and it names DeepSeek V4's driver as the sibling that moves
-with it.
+**W5, migrate the three single-shape drivers. DONE 2026-08-19,
+[#1335](https://github.com/mudler/vllm.cpp/issues/1335).** DeepSeek V4, Laguna
+and DFlash. `laguna.cpp:2116-2119` carried the note that its capture class was
+waiting for exactly this seam and named DeepSeek V4's driver as the sibling that
+moves with it; both moved in this stage and the note now says what is actually
+left device-coupled, which is the captured CHAIN and not the capture.
+
+Each driver's `Step` opens a `vt::GraphCaptureScope` over its own
+`vt::BreakableGraph` in `kFull` and replays through `BreakableGraph::Replay`.
+What is gone from all three: the hand-rolled `BeginCapture`/`EndCaptureGraph`
+pair, the raw `void*` handle, two hand-written destructors that released it, and
+DFlash's private `try`/drain. **NINE OF NINE DRIVERS ARE ON THE SEAM**, and
+`grep -rnE '\.(BeginCapture|EndCaptureGraph|ReplayGraph|DestroyGraph)\s*\('`
+over `src/vllm/`, with comment lines excluded, returns NOTHING — every surviving
+textual hit is prose about what was removed.
+
+The three per-model rollback switches STAY (`VT_V4_DECODE_GRAPH`,
+`VT_DFLASH_GRAPH`, `VT_LAGUNA_DECODE_GRAPH`): each is a same-binary A/B lever for
+exactly one driver, not a copy of the shared one. What changes is that
+`VLLM_CPP_CUDAGRAPH` reaches all three for the first time, through
+`vt::GraphCaptureEnabled()` inside the scope. `gstate`/`g_state` stays too, and
+it is not a duplicate of `captured()`: it is each driver's cold/warm/captured
+ladder, and the seam has no notion of the eager warm-run that grows the pool so
+the capture that follows can allocate nothing.
+
+**D10 IS DISCHARGED HERE, and W5 is the first stage where it could be.**
+`GraphCaptureScope` owns the set of side queues forked since the current segment
+opened and joins every outstanding one before `Backend::EndCaptureGraph` — the
+port of `_end_current_segment` (`:353-361`) plus the `wait_stream` hook
+(`:101-153`) whose only purpose is to populate that set. W1 registered its break
+point on a model that forks no auxiliary queue; W2, W3 and W4 all opened `kFull`,
+which has ONE segment and therefore no segment CLOSE inside a fork window for the
+rule to govern, so the machinery would have landed unexercised. W5 owns the only
+driver in the tree whose fork is inside the captured region by construction
+(`laguna.cpp:2572-2576` fork, `:2612` join), and that driver is the rule's
+PRODUCTION CALLER through `vt::GraphNoteFork` and `vt::GraphNoteJoin`.
+
+We need no monkey-patch to populate the set, because our fork and join are
+explicit `Backend::RecordEvent` and `Backend::QueueWaitEvent` calls rather than
+an implicit torch API — the model TELLS the scope. The retirement call is not
+politeness: without it the scope would issue a REDUNDANT second join before every
+segment close, which is exactly what the gate's CONTROL arm pins. Both hooks are
+no-ops outside an active scope and make ZERO backend calls there, which is the
+same pass-through guarantee `GraphBreak` gives.
+
+**The rule is gated as a COUNTER and an ORDER, because W3 proved prose is not a
+gate.** W3 measured a mode guard that stayed green at 226/226 under the exact
+mutation it named. So `vt::GraphBreakStats` gains `forks_tracked` and
+`forks_auto_joined`, and the load-bearing one is the second: it is 0 for a model
+that joins its own fork inside the segment and non-zero exactly when the seam did
+the work. And the claim is an ORDER rather than an event, because closing a
+capture with an unjoined fork FAILS at `cudaStreamEndCapture` — so both ends are
+asserted out of ONE backend trace
+(`Begin RecordEvent QueueWaitEvent EndCaptureGraph Begin EndCaptureGraph`), for
+the same reason W1 had to move break markers into the backend's own log. Test 15
+of `## Tests to port` carries five arms: the rule, the CONTROL where the model
+joins first and the scope must do nothing, `kFull` (the arm every migrated driver
+actually takes), the inert scope, and re-registering one queue. Two mutations,
+both compiled clean: deleting `JoinOutstandingForks()` reds ONLY the new case on
+5 assertions at 29 of 30 still passing, and making `NoteJoin` fail to retire the
+entry reds it on 8 — so neither the rule nor its control is vacuous. Green after:
+30 cases, 265 assertions, exit 0.
+
+**ONE of the three drivers is gateable without a GPU, and the reason is each
+driver's own admission predicate rather than a choice.** The DFlash draft graph
+admits on `VT_DFLASH_GRAPH` plus `Backend::SupportsGraphCapture()` plus
+`Platform::support_static_graph_mode()` — it names neither a device type nor a
+kernel registry — so the shared harness's two swapped registries reach it.
+DeepSeek V4's `CanRunResidentDecode` (`deepseek_v4.cpp:1481-1487`) refuses a CPU
+queue outright and refuses again unless `V4DeviceKernelsAvailable()`, the four V4
+kernel families registered under `kCUDA` by `cuda_deepseek_v4.cu`. Laguna's whole
+capture class is behind `#ifdef VT_MARLIN_NVFP4`, which CMake sets only for a
+CUDA build on a marlin-nvfp4 architecture (`sm_12xa` and, since the Thor
+bring-up, `sm_110`), so on a box with no `nvcc` the migrated region is not in any
+reachable binary at all.
+
+`tests/vllm/models/test_qwen3_dflash_decode_graph_seam.cpp` is therefore the
+stage's driver gate. RED FIRST against the unmigrated driver: 3 cases, 0 passed,
+16 assertions, 7 failed, exit 1, on `segments_captured`, `full_scopes` and
+`replays`. GREEN after: 3 cases, 18 assertions, exit 0. G2's mutation — replacing
+the scope and `Replay` with the pre-W5 raw pair, compiled clean at 79 insertions
+and 29 deletions — reds ONLY that file (7 assertions) and leaves
+`test_breakable_graph` 265, `test_qwen3_decode_graph_seam` 231,
+`test_qwen3_moe_decode_graph_seam` 228, `test_voxtral_decode_graph_seam` 230,
+`test_deepseek_v2_decode_graph_seam` 230, `test_qwen3_5_decode_graph_seam` 129
+and the driver's OWN `test_dflash_propose` 31 all GREEN. That last one is the
+whole argument restated as a measurement: a driver's bit-exactness suite cannot
+see which capture machinery ran. G4 holds in the same file at 24 values, 0
+differing, against the driver's own eager paged arm.
+
+**What W5 does NOT deliver, named rather than implied.** G1 for all three
+drivers, and G2 for DeepSeek V4 and Laguna, are OWED on hardware for the
+predicate and build-flag reasons above. G5's ROCm and Tenstorrent arms stay
+BLOCKED on hardware the fleet does not carry. `## Owed` carries each with its
+resource.
 
 **W6, close the coverage gap the row exists for.** Only after W1 through W5 does
 the predicate at `runner.cpp:1341` move from `pure_decode` to "eligible except at
@@ -1197,12 +1328,23 @@ point registered inside an unjoined fork window without this rule fails at
 
 `ACTIVE`. W0 (spike), W1 (the seam, its ported unit gate, and one registered
 break point), W2 (`Qwen3DenseDecodeGraph`), W3 (`Qwen3MoeDecodeGraph`,
-`VoxtralDecodeGraph`, `DeepseekV2DecodeGraph`) and W4 (the persistent device
-input path as `vt::PersistentStepInput`, plus `Qwen3_5DecodeGraph` and
-`Qwen3_5DenseDecodeGraph`) have landed. **Six of the nine drivers are on the
-seam**, and one `std::getenv("VLLM_CPP_CUDAGRAPH")` remains in `src/`, the
-seam's own. W5 and W6 remain, and `## Work breakdown` states each. Owner:
-`.agents/claims/CLAIM-ENG-CUDAGRAPH-BREAK-W4.md`.
+`VoxtralDecodeGraph`, `DeepseekV2DecodeGraph`), W4 (the persistent device input
+path as `vt::PersistentStepInput`, plus `Qwen3_5DecodeGraph` and
+`Qwen3_5DenseDecodeGraph`) and W5 (the three single-shape drivers, plus D10) have
+landed. **NINE OF THE NINE DRIVERS ARE ON THE SEAM**, the migration is complete,
+and one `std::getenv("VLLM_CPP_CUDAGRAPH")` remains in `src/` — the seam's own.
+`grep -rnE '\.(BeginCapture|EndCaptureGraph|ReplayGraph|DestroyGraph)\s*\('`
+over `src/vllm/`, with comment lines excluded, returns NOTHING. **Only W6
+remains**, and `## Work breakdown` states it. Owner:
+`.agents/claims/CLAIM-ENG-CUDAGRAPH-BREAK-W5.md`.
+
+**D10, the auxiliary-stream auto-join, is discharged and is REACHED.** Every
+prior stage opened `kFull`, which has one segment and so no between-segments
+window, so the rule could not be exercised and untested machinery was not landed
+for it. W5 owns `laguna.cpp`, the only driver whose fork is inside the captured
+region by construction, and that driver is the production caller. Gated as a
+counter and an ORDER out of one backend trace, with two mutations proving neither
+the rule nor its control arm is vacuous.
 
 **W4 corrected a premise this spec had asserted three times.** The decode graph
 carries NO token ids to the device in any driver, `StepDevInputs` included, so
@@ -1214,13 +1356,14 @@ Every migrated step opens its scope in `kFull`, mirroring vLLM's decode arm. The
 PIECEWISE arm still has no production driver and `## Owed` names what has to be
 true before one exists.
 
-**G1 is MET for five of the six migrated drivers and is no longer owed.**
+**G1 is MET for five of the NINE migrated drivers and is owed for four.**
 Bit-exactness against the eager arm over the capture step plus THREE consecutive
 replays, on `thor:gpu0` through an `rc` lease: **5 cases, 2066 assertions, 0
 differing**, the two Qwen3.5 drivers added by W4. W2's `Qwen3DenseDecodeGraph`
 is still covered by argument rather than by measurement, which `## Gates` G1
-says out loud. This is the only gate in this row a CPU harness could never have
-answered.
+says out loud, and W5's three single-shape drivers join it there for reasons
+`## Owed` names per driver. This is the only gate in this row a CPU harness could
+never have answered.
 
 W1's exit criterion — that CUDA permits `cudaStreamEndCapture` followed by
 `cudaStreamBeginCapture` mid-forward on our stream configuration — is
@@ -1237,18 +1380,24 @@ Each item names the stage that owns it. Nothing here is claimed by W1.
   `tests/vllm/models/test_qwen3_decode_graph_seam.cpp` holds it through the
   seam's own counters. This was the staged slice AGENTS.md allows; it is closed
   rather than carried.
-- **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. **W4 migrated `qwen3_5.cpp` and did NOT
-  discharge this**, and the reason is the mode rather than the effort: its scope
-  is `kFull`, so there is exactly one segment and no segment CLOSE inside the
-  fork window at `:6254-6255,6384` for the rule to govern, and the machinery
-  would have landed unexercised. Owners: **W5** (`laguna.cpp:2572-2576,2612`,
-  whose fork is inside the captured region by construction) and the first stage
-  that captures PIECEWISE.
+- ~~**The auxiliary-stream auto-join before every segment close** (D10, the port
+  of `breakable_cuda_graph.py:353-361`).~~ RETIRED by **W5**
+  ([#1335](https://github.com/mudler/vllm.cpp/issues/1335)).
+  `GraphCaptureScope` owns the outstanding-fork set and joins every entry before
+  `Backend::EndCaptureGraph`; `vt::GraphNoteFork` and `vt::GraphNoteJoin` are the
+  registration, and `laguna.cpp:2572-2576,2612` — the only fork inside a captured
+  region by construction — is the production caller. W1 through W4 could not
+  discharge it and the reason was the MODE rather than the effort: every stage
+  before this one opened `kFull`, which has one segment and therefore no segment
+  CLOSE inside a fork window for the rule to govern, so landing the machinery
+  would have landed it unexercised. Gated as test 15 of `## Tests to port` with
+  five arms and proven non-vacuous by two mutations. **One residual, and it is
+  named rather than folded in:** the seam joins a fork the model REGISTERED. A
+  model that forks without calling `GraphNoteFork` is exactly as exposed as
+  before, and no checker can see that, because the fork is an ordinary pair of
+  backend calls. Laguna is the only such site today; the next one is the
+  obligation of whoever adds it, and the `SupportsAuxStream()` capability is
+  where a reader lands.
 - ~~**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
@@ -1264,11 +1413,90 @@ Each item names the stage that owns it. Nothing here is claimed by W1.
   The lease W3 obtained was `thor:gpu0`, which is CUDA; the fleet
   (`rc devices`) carries no ROCm device and no Tenstorrent device, so this is
   BLOCKED on hardware rather than unattempted. What W3 can say is what it
-  measured: the seam's CUDA arm now runs on TWO architectures rather than one,
-  sm_110 here and sm_121a on GB10 for the W1 exit criterion. Owner: **W5**,
-  which migrates the driver family whose Tenstorrent recapture path
-  (`qwen3.cpp`'s `VT_TT_RECAPTURE_EVERY` branch) is the only place a ttnn mesh
-  trace meets this seam today.
+  measured: the seam's CUDA arm now runs on TWO architectures rather than one.
+
+  **THAT SENTENCE NAMED THE WRONG DEVICE AND THE WRONG ARCHITECTURE, and W5
+  corrected it ([#1361](https://github.com/mudler/vllm.cpp/issues/1361)).** It
+  read "sm_110 here and sm_121a on GB10 for the W1 exit criterion", while this
+  same file records at `## Work breakdown` W1 that the exit criterion was
+  measured on `orin:gpu0`, driver `12060` — a Jetson AGX Orin, which is neither
+  a GB10 nor `sm_121a`. The two architectures are real and the claim survives;
+  the attribution did not. What was measured, and where: the W1 exit criterion
+  (`cudaStreamEndCapture` then `cudaStreamBeginCapture` mid-forward with eager
+  work between) on `orin:gpu0`, and G1 plus the unit suite on `thor:gpu0` at
+  sm_110 for W3, W4 and W5. The exit criterion has NOT been re-measured on
+  `thor`, and the reason is structural rather than an omission: every migrated
+  driver opens `kFull`, so nothing in the tree re-begins a capture mid-forward,
+  and G1 exercises capture and replay rather than the re-begin. **`sm_121a` on
+  GB10 is OWED, not done** — W5 could not take it because `dgx:gpu0` was held by
+  another session for that stage's whole window. This is the shape where a
+  number quoted often starts being treated as measured.
+  **W5 did NOT discharge it either, and the fleet answer is unchanged**:
+  `rc devices` lists `dgx:gpu0`, `orin:gpu0` and `thor:gpu0`, all NVIDIA. This is
+  the second stage to inherit the item and find the same wall, which is the
+  signal that it is not a stage-sized problem: it needs a ROCm or Tenstorrent
+  device on the fleet, and no ordering of the remaining work produces one. Owner:
+  row **`ENG-CUDAGRAPH-BREAK`**, discharged by whichever stage first has the
+  hardware. The Tenstorrent half additionally has a named entry point when that
+  day comes — `qwen3.cpp`'s `VT_TT_RECAPTURE_EVERY` branch is the only place a
+  ttnn mesh trace meets this seam.
+
+- **G1 and G2 for the three SINGLE-SHAPE drivers** (W5,
+  [#1335](https://github.com/mudler/vllm.cpp/issues/1335)). Each has its own
+  reason and none of them is effort, so they are listed per driver rather than as
+  one line:
+
+  1. **DFlash** has G2 and G4 on CPU
+     (`tests/vllm/models/test_qwen3_dflash_decode_graph_seam.cpp`, red-first,
+     3 cases / 18 assertions) and owes only G1, which every CPU-gated driver in
+     this row owes for the same reason: a CPU "replay" recomputes nothing, so a
+     CPU harness cannot hold that a REPLAYED segment reproduces the eager
+     forward.
+  2. **DeepSeek V4** owes G1 AND G2. `CanRunResidentDecode`
+     (`deepseek_v4.cpp:1481-1487`) refuses `device.type == kCPU` outright and
+     refuses again unless `V4DeviceKernelsAvailable()` — the four V4 kernel
+     families registered under `kCUDA` by `cuda_deepseek_v4.cu`. The shared
+     harness swaps a backend and a platform; it cannot manufacture a device type,
+     a CUDA op registry and a whole vt kernel layer, and a stub that did would
+     measure the stub. The migration is covered by CONSTRUCTION — same seam, same
+     mode, same call shape as the seven gated drivers — which is an argument and
+     not a gate, and this entry exists so nobody reads it as one.
+  3. **Laguna** owes G1 and G2 for a build reason rather than a predicate one.
+     `LagunaGraph` is inside `#ifdef VT_MARLIN_NVFP4`, which CMake sets only for
+     a CUDA build on a marlin-nvfp4 architecture, so on a host with no `nvcc` the
+     migrated region is in no reachable binary. `laguna.cpp` compiling clean in
+     the CPU configuration verifies the FILE and not the region, and
+     `test_laguna_scaffold` (167 assertions) and `test_laguna_nvfp4_loader` (63)
+     staying green says nothing about the capture at all.
+
+  **ONE HALF OF 2 AND 3 IS NOW CLOSED and the rest is not, so the split is worth
+  stating.** W5 measured on `thor:gpu0` that both migrated regions are actually
+  COMPILED, by mutating each one and requiring the object build to fail — see
+  `## Gates` G1. So "the code might not even be built" is no longer part of what
+  these two owe. What they still owe is behavioural: that the routing reaches the
+  seam (G2) and that a REPLAYED segment reproduces the eager forward (G1), and
+  both need each model's own device kernels — V4's four `kCUDA`-registered
+  families, Laguna's NVFP4 Marlin arm — rather than only a compiler.
+
+- **The W1 exit criterion on a THIRD architecture, `sm_121a` on GB10.** Filed as
+  its own item because #1361 found the record already claiming it. What has
+  actually been measured: the criterion on `orin:gpu0`, and G1 plus the unit
+  suite on `thor:gpu0` at sm_110. `dgx:gpu0` is the only `sm_121a` device on the
+  fleet and it was held by another session for the whole of W5's window
+  (`rc devices`: `busy`, past one hour). **It cannot be discharged by re-running
+  G1 there**, and that is the part worth writing down: every migrated driver
+  opens `kFull`, so nothing in this tree re-begins a capture mid-forward, and G1
+  therefore exercises capture and replay rather than the re-begin the criterion
+  is about. Discharging it needs the W1 probe itself run on GB10, or the first
+  PIECEWISE production driver — which is W6. Owner: row
+  **`ENG-CUDAGRAPH-BREAK`**.
+
+  What settles all three: a CUDA build on a marlin-nvfp4 architecture — the same
+  `thor:gpu0` (sm_110) shape W3 and W4 used — which compiles the Laguna and V4
+  regions and can extend
+  `tests/vllm/models/test_decode_graph_seam_g1_cuda.cpp` with a DFlash case over
+  the synthetic draft model this stage's CPU gate already builds. Owner: row
+  **`ENG-CUDAGRAPH-BREAK`**, the stage that gets that window.
 - ~~**G1, bit-exactness against eager on a real GPU over MORE than one
   replay.**~~ RETIRED by **W3**
   ([#1291](https://github.com/mudler/vllm.cpp/issues/1291)).
diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md
index 6c3d38c89..e5ce172d8 100644
--- a/docs/ENVIRONMENT.md
+++ b/docs/ENVIRONMENT.md
@@ -202,7 +202,7 @@ portable/reference path. In normal operation leave them unset.
 | `VT_DFLASH_PAGED` | on (CUDA, DFlash spec-decode) | The materialized `[context;block]` draft forward instead of the fixed-capacity paged draft-KV store read through `vt::DFlashPagedBlockAttention` (bit-identical; only the DFlash single-request propose path) |
 | `VT_SPEC_DECODE_GRAPH` | on (CUDA, spec-decode) | `0` runs the T=1+k speculative VERIFY eagerly instead of replaying a captured graph of that shape. Mirrors vLLM's uniform-decode dispatch, whose captured length is `1 + num_speculative_tokens`; output is byte-identical either way, so this is a cost switch for A/B |
 | `VT_DSPARK_DEVICE_SAMPLE` | on (DSpark spec-decode) | `0` restores the HOST sequential Markov loop. The device path keeps the per-step bias and argmax on the GPU and is token-identical to the host loop (same lowest-index tie-break, same base row, same d2t-mapped chain), so this is a cost switch for A/B only |
-| `VT_DFLASH_GRAPH` | on (CUDA, DFlash spec-decode) | The eager paged draft step instead of the captured/replayed draft-step CUDA graph (replayed==eager bit-identical; only the DFlash single-request propose path) |
+| `VT_DFLASH_GRAPH` | on (CUDA, DFlash spec-decode) | The eager paged draft step instead of the captured/replayed draft-step CUDA graph (replayed==eager bit-identical; only the DFlash single-request propose path). Since `ENG-CUDAGRAPH-BREAK` W5 ([#1335](https://github.com/mudler/vllm.cpp/issues/1335)) the capture runs through the shared break-point seam, so `VLLM_CPP_CUDAGRAPH=0` also turns it off; this name stays as the same-binary A/B lever for THIS driver alone |
 | `VT_DFLASH_ATTN_BLOCK` | off (CUDA, DFlash spec-decode) | `=1` selects the D12/D13 block-per-(query,head) draft paged-attention kernel instead of the D14 default warp-scoped online-softmax kernel (same f32-softmax math within the bf16 envelope; the D14 warp kernel is ~3x faster and closed the ~2% speed residual; only the DFlash single-request propose path) |
 | `VT_DFLASH_ATTN_MMA` | on where the built arch set and the running device both have bf16 `mma.sync` (Ampere and later) | `=0` disables the bf16 tensor-core block-attention path and falls back to the CUDA-core form. The guard is both compile-time (`__CUDA_ARCH_LIST__` all >= 800) and runtime (queried compute capability), so on a pre-Ampere board the path is already off and this flag changes nothing; it exists so the tensor-core and CUDA-core forms can be A/B'd on ONE binary, which is this project's benchmark protocol |
 | `VT_DFLASH_ATTN_WARP` | off (CUDA) | `=1` falls back to the older per-key warp-reduction block-attention kernel instead of the default chunked reduce-scatter form. Kept for the same-binary A/B that recorded the verdict |
diff --git a/docs/STATUS.md b/docs/STATUS.md
index 48812b95c..04a4df406 100644
--- a/docs/STATUS.md
+++ b/docs/STATUS.md
@@ -239,6 +239,38 @@ uploaded. So the classic-dense decline that costs a shipped model its decode
 graph under asynchronous serving is not one refactor away from removable. It
 stands, and the work it needs is now named rather than assigned.
 
+W5 (2026-08-19, #1335) migrates the last three, so ALL NINE drivers are on the
+seam and the migration is complete. These are the single-shape drivers: the
+DFlash draft graph, the DeepSeek V4 decode graph, and the Laguna decode graph,
+whose own source note asked for this seam by name. No hand-rolled capture call
+survives anywhere under `src/vllm/`.
+
+W5 also gives the seam the auxiliary-stream rule. Closing a graph segment while
+a side stream forked inside it is still recording is illegal, so the capture
+scope now tracks the forks opened since the segment began and joins any that are
+still outstanding before it closes. Every earlier stage captured in FULL mode,
+which has one segment and therefore no window for the rule to govern, so this is
+the first stage that could exercise it; the Laguna decode graph, whose fork sits
+inside the captured region by construction, is what reaches it.
+
+Bit-exactness against a replayed capture is still owed for four of the nine
+drivers, and for two of them so is the routing gate. The reasons are per driver
+and are recorded: DeepSeek V4 refuses a CPU queue before it reaches its capture,
+and Laguna's capture class only exists in a CUDA build with the Marlin NVFP4
+kernels. The DFlash driver's own gate landed red first.
+
+One half of that is closed. Laguna's capture class is behind a build flag, so a
+green CUDA build is equally consistent with the migrated code having been
+compiled OUT. Injecting an error into each migrated region and requiring the
+build to fail settles it: both regions are really compiled. What the two still
+owe is behavioural, and needs each model's own device kernels rather than only a
+compiler.
+
+Bit-exactness for the five drivers that HAVE it was re-measured at this stage's
+head on a leased GPU, because the shared capture-close path changed underneath
+them: five drivers, 2066 assertions, nothing differing over three consecutive
+replays each.
+
 W3 also closed a gate that could not fail. The mode a driver captures in was
 unobservable from outside it, so a one-token FULL-to-PIECEWISE flip left a whole
 driver gate green. The seam now counts the mode, and that flip reds each gate.
diff --git a/docs/USAGE.md b/docs/USAGE.md
index 36fa28d18..4a0dfecff 100644
--- a/docs/USAGE.md
+++ b/docs/USAGE.md
@@ -4586,15 +4586,41 @@ 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`. **Production steps now open a capture
-scope**, six of them as of W4: `Qwen3DenseDecodeGraph` (W2, #1261),
-`Qwen3MoeDecodeGraph`, `VoxtralDecodeGraph` and `DeepseekV2DecodeGraph` (W3,
-#1291), and `Qwen3_5DecodeGraph` with `Qwen3_5DenseDecodeGraph` (W4, #1307).
-Every one of them opens the scope in FULL mode, mirroring the decode half of
-vLLM's v1 default `CUDAGraphMode.FULL_AND_PIECEWISE`, and a `vt::GraphBreak`
-inside a FULL scope takes its pass-through arm — so the switch still changes
-nothing about the break point beyond what it already changed about the decode
-graphs. This paragraph asserted the opposite until W4: it was written at W1,
-when it was true, and W2 falsified it without rewriting it here.
+scope, and as of W5 (#1335) ALL NINE decode and draft graphs do**:
+`Qwen3DenseDecodeGraph` (W2, #1261), `Qwen3MoeDecodeGraph`, `VoxtralDecodeGraph`
+and `DeepseekV2DecodeGraph` (W3, #1291), `Qwen3_5DecodeGraph` with
+`Qwen3_5DenseDecodeGraph` (W4, #1307), and the DFlash draft graph, the DeepSeek
+V4 decode graph and the Laguna decode graph (W5, #1335). Every one of them opens
+the scope in FULL mode, mirroring the decode half of vLLM's v1 default
+`CUDAGraphMode.FULL_AND_PIECEWISE`, and a `vt::GraphBreak` inside a FULL scope
+takes its pass-through arm — so the switch still changes nothing about the break
+point beyond what it already changed about the decode graphs. This paragraph
+asserted the opposite until W4: it was written at W1, when it was true, and W2
+falsified it without rewriting it here.
+
+**W5 WIDENS WHAT THE SWITCH REACHES, and that is a user-visible change rather
+than an internal one.** The three single-shape drivers never read
+`VLLM_CPP_CUDAGRAPH` at all: each invented its own name — `VT_V4_DECODE_GRAPH`,
+`VT_DFLASH_GRAPH` and `VT_LAGUNA_DECODE_GRAPH` — so before W5 there was no single
+setting that turned capture off everywhere. There is now, and the three
+per-driver names STAY, because each is a same-binary A/B lever for exactly one
+driver rather than a copy of the shared one. Either turns its driver's capture
+off; `VLLM_CPP_CUDAGRAPH=0` turns all nine off at once.
+
+Turning capture off on those three does NOT return uncomputed memory, and the
+distinction is worth stating because it is invisible to a token gate. An INERT
+scope runs the forward eagerly, so the driver's buffers hold real values. A
+capture that FAILS is the opposite: under stream capture nothing between the
+begin and the failure executed, so those same buffers hold whatever the
+allocator last left there. The seam reports the two states apart and every
+migrated driver propagates the failure instead of returning the buffer.
+
+**The seam also owns the auxiliary-stream rule as of W5.** A model that forks a
+side stream inside a capture — the Laguna decode graph runs its FP4 shared
+expert that way — registers the fork with the capture scope, and the scope joins
+any fork still outstanding before it closes a segment, because ending a capture
+with an unjoined fork fails. There is nothing to configure: registration is part
+of the model's fork, and outside a capture both hooks do nothing at all.
 
 Building it needs no option. `src/vt/breakable_graph.cpp` and, since W4,
 `src/vt/persistent_step_input.cpp` — the capture-stable per-step device input
diff --git a/include/vt/breakable_graph.h b/include/vt/breakable_graph.h
index 2dd7a87de..11ccb7bdb 100644
--- a/include/vt/breakable_graph.h
+++ b/include/vt/breakable_graph.h
@@ -180,6 +180,17 @@ struct GraphBreakStats {
   // would make the number report a mode that never reached a backend.
   int64_t full_scopes = 0;
   int64_t piecewise_scopes = 0;
+  // D10, the auxiliary-stream fork set. ENG-CUDAGRAPH-BREAK W5 (#1335) added
+  // these for the same reason W3 added the mode counters: the rule is invisible
+  // from outside the scope, and a guard nothing can observe is a guard nothing
+  // can prove. `forks_tracked` counts `GraphNoteFork` registrations that reached
+  // an ACTIVE scope; `forks_auto_joined` counts the forks the scope itself had
+  // to join because they were still outstanding when a segment closed. The
+  // second is the one that matters: it is ZERO for a model that joins its own
+  // fork inside the segment, and non-zero exactly when the seam did the work
+  // upstream's `_end_current_segment` (`:353-361`) does.
+  int64_t forks_tracked = 0;
+  int64_t forks_auto_joined = 0;
 };
 GraphBreakStats GetGraphBreakStats();
 void ResetGraphBreakStats();
@@ -408,6 +419,37 @@ class GraphCaptureScope {
   // share an address — and the cell is what actually aliases.
   void AppendBreak(std::function<void()> fn, const void* destination);
 
+  // ---------------------------------------------------------------------
+  // D10 — THE OUTSTANDING-FORK SET. Port of upstream's `wait_stream` hook
+  // (`:101-153`, installed `:310`, removed `:332`) and of the auto-join inside
+  // `_end_current_segment` (`:353-361`).
+  // ---------------------------------------------------------------------
+  //
+  // WHY THIS EXISTS. Closing a segment while a side stream forked inside it is
+  // still participating in the capture is illegal: `cudaStreamEndCapture` fails
+  // on a capture graph with an unjoined fork. Upstream keeps the whole
+  // `torch.cuda.Stream.wait_stream` monkey-patch for no other purpose than to
+  // know WHICH streams are outstanding, then joins them itself before
+  // `capture_end()`.
+  //
+  // WHY OURS IS CHEAPER. Our fork and join are explicit `Backend::RecordEvent`
+  // and `Backend::QueueWaitEvent` calls on a seam we control, not an implicit
+  // torch API, so the model TELLS the scope instead of the scope patching the
+  // runtime. `NoteFork` registers; `NoteJoin` retires the registration when the
+  // model joined the queue itself, which is the common case and the one that
+  // must cost nothing.
+  //
+  // `join_event` is the event the scope RECORDS on `aux` and then WAITS on the
+  // capture queue if it still has to join. It must outlive this scope — the
+  // natural owner is the capture driver that created the queue, which is what
+  // `laguna.cpp`'s `LagunaGraph` does with `aux_done` (created in the
+  // constructor, destroyed in the destructor). Registering an event that dies
+  // first is lifetime rule 2 at the `GraphBreak` declaration, applied to an
+  // event instead of a buffer.
+  void NoteFork(Queue& aux, Event& join_event);
+  void NoteJoin(Queue& aux);
+  size_t outstanding_forks() const { return forks_.size(); }
+
  private:
   Backend* b_;
   Queue* q_;
@@ -420,8 +462,37 @@ class GraphCaptureScope {
   // 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<const void*> destinations_;
+  // The forks registered since this segment opened and not yet retired. One
+  // entry per outstanding side queue; a model forks as many queues as it has
+  // overlap paths, so a linear scan is cheaper than the map that would replace
+  // it — the same argument `destinations_` makes.
+  struct ForkedQueue {
+    Queue* q;
+    Event* e;
+  };
+  std::vector<ForkedQueue> forks_;
+  // Joins every outstanding fork onto the capture queue and clears the set.
+  // Called by `EndSegment` BEFORE `Backend::EndCaptureGraph`.
+  void JoinOutstandingForks();
 };
 
+// ---------------------------------------------------------------------------
+// vt::GraphNoteFork / vt::GraphNoteJoin — the D10 registration at the SITE.
+// ---------------------------------------------------------------------------
+//
+// Same shape as `GraphBreak`: OUTSIDE a capture scope both are no-ops making
+// ZERO backend calls, so a non-capturing forward is byte-identical to today.
+// Inside an ACTIVE scope `GraphNoteFork` adds the queue to the outstanding set
+// and `GraphNoteJoin` retires it.
+//
+// THE MODEL STILL ISSUES ITS OWN FORK AND JOIN. These do not replace
+// `Backend::RecordEvent` / `Backend::QueueWaitEvent`; they tell the scope that
+// a fork is open, so the scope can close the window if a segment ends before
+// the model's own join runs. A model that always joins inside its segment sees
+// `forks_auto_joined == 0` and pays two vector operations.
+void GraphNoteFork(Queue& aux, Event& join_event);
+void GraphNoteJoin(Queue& aux);
+
 // ---------------------------------------------------------------------------
 // vt::GraphBreak — the break point. Port of `eager_on_graph` (`:204-243`),
 // whose wrapper body is `:209-241`.
diff --git a/src/vllm/model_executor/models/deepseek_v4.cpp b/src/vllm/model_executor/models/deepseek_v4.cpp
index 92ccd768f..5828bf576 100644
--- a/src/vllm/model_executor/models/deepseek_v4.cpp
+++ b/src/vllm/model_executor/models/deepseek_v4.cpp
@@ -45,6 +45,7 @@
 #include <cstdint>
 #include <cstdio>
 #include <cstdlib>
+#include <exception>
 #include <limits>
 #include <memory>
 #include <numbers>
@@ -67,6 +68,7 @@
 #include "vt/ops.h"      // vt::MatmulBT (auto-dispatches kMatmulBTQuant on block weights)
 #include "vt/tensor.h"   // vt::Tensor::Contiguous
 #include "vt/backend.h"  // vt::GetBackend / Backend::Synchronize (device GEMM drain)
+#include "vt/breakable_graph.h"  // ENG-CUDAGRAPH-BREAK W5: the shared capture seam
 
 namespace vllm {
 namespace {
@@ -1715,7 +1717,15 @@ struct V4Graph {
   float* res_cur = nullptr;
   float* res_nxt = nullptr;
   int64_t kv_base = 0;
-  void* graph = nullptr;
+  // ENG-CUDAGRAPH-BREAK W5 (#1335): the instantiated graph, its handle ownership
+  // and its release live in the SHARED SEAM instead of in a raw `void*` this
+  // class destroyed by hand. `vt::BreakableGraph` releases every segment it
+  // holds through `Backend::DestroyGraph`, which is the routing that lets
+  // ENG-CUDAGRAPH-DEDUP (#1162) interpose at the backend later without editing
+  // this file. `gstate` STAYS: it is this driver's cold/warm/captured ladder,
+  // not a duplicate of `captured()` — the seam has no notion of the eager
+  // warm-run that grows the pool before a capture may allocate nothing.
+  vt::BreakableGraph graph;
   int gstate = 0;  // 0 cold (eager warm-run), 1 warm (capture+replay), 2 captured (replay)
   vt::Queue* qu = nullptr;
 
@@ -1766,9 +1776,8 @@ struct V4Graph {
       std::copy(pref.begin(), pref.end(), cache[static_cast<size_t>(l)].begin());
     }
   }
-  ~V4Graph() {
-    if (graph != nullptr && qu != nullptr) vt::GetBackend(qu->device).DestroyGraph(graph);
-  }
+  // No destructor: `graph` releases its own segments through
+  // `Backend::DestroyGraph`, so the hand-rolled one this replaced is gone.
 
   // The per-layer resident chain over the PERSISTENT buffers (the capture region).
   void RunChain(const V4Backend& be) {
@@ -1895,13 +1904,55 @@ struct V4Graph {
       RunChain(be);
       gstate = 1;
     } else if (gstate == 1) {   // warm: capture the region once, then replay it
-      b.BeginCapture(*be.q);
-      RunChain(be);
-      graph = b.EndCaptureGraph(*be.q);
-      b.ReplayGraph(*be.q, graph);
-      gstate = 2;
+      // ENG-CUDAGRAPH-BREAK W5 (#1335): the capture is the SHARED SEAM's, not
+      // this driver's hand-rolled `BeginCapture`/`EndCaptureGraph` pair. The
+      // scope owns the segment, the handle, its release, the drain a mid-capture
+      // throw needs, and the G3 counters.
+      //
+      // kFULL, INHERITED FROM W2 AND NOT RE-ARGUED. vLLM's v1 default
+      // `FULL_AND_PIECEWISE` (`vllm/config/compilation.py:63` @ pin
+      // `5559679229`) is documented at `:630-632` as a FULL graph for DECODE
+      // batches and a piecewise one for prefill and mixed batches, and
+      // `decode_mode()` (`:65-66`) returns the full half. This is the T=1
+      // resident decode step, so its capture is ONE segment with the attention
+      // calls INSIDE it — byte-identical in shape to the region this replaces.
+      {
+        vt::GraphCaptureScope scope(b, *be.q, graph, vt::GraphCaptureMode::kFull);
+        RunChain(be);
+      }  // ~GraphCaptureScope closes the segment and files it on `graph`
+      // NOT CAPTURED covers TWO states and they mean opposite things.
+      //
+      //   * FAILED (`capture_failed() == true`): `Backend::EndCaptureGraph`
+      //     threw. Under stream capture NOTHING between `BeginCapture` and the
+      //     throw executed — every kernel was RECORDED — so `logits` holds
+      //     whatever the pool last left there, and returning it would hand this
+      //     step uncomputed device memory as its logits: no fault, and a token
+      //     gate cannot see it. It PROPAGATES, carrying the runtime's own
+      //     exception, which is what the pre-W5 driver's unguarded
+      //     `EndCaptureGraph` did.
+      //   * INERT (`capture_failed() == false`): capture is unsupported here, or
+      //     `VLLM_CPP_CUDAGRAPH=0`. The scope made no backend call, `RunChain`
+      //     ran EAGERLY, and `logits` is a real result — so this step returns
+      //     normally and the driver stays in `gstate == 1`, running eager every
+      //     step rather than pretending to hold a graph.
+      if (!graph.captured()) {
+        if (graph.capture_failed()) {
+          const std::exception_ptr err = graph.capture_error();
+          graph.Reset();  // clear the failure with the graph it described
+          if (err) std::rethrow_exception(err);
+          VT_CHECK(false,
+                   "deepseek-v4 decode graph: the capture was ABANDONED and its logits "
+                   "were never computed; refusing to return uncaptured device memory");
+        }
+      } else {
+        graph.Replay(*be.q);
+        gstate = 2;
+      }
     } else {                    // captured: one cudaGraphLaunch
-      b.ReplayGraph(*be.q, graph);
+      // Through the seam's container, never `Backend::ReplayGraph` directly: it
+      // replays its segments in order (one, here, because this capture is kFull)
+      // and owns the G3 replay counter.
+      graph.Replay(*be.q);
     }
     // append this token's deck_new → cache[kv_base] (on the stream, AFTER the step's
     // graph produced deck_new, BEFORE the next replay reads it) — the growing KV.
diff --git a/src/vllm/model_executor/models/laguna.cpp b/src/vllm/model_executor/models/laguna.cpp
index dbf1f9de2..f1e53882d 100644
--- a/src/vllm/model_executor/models/laguna.cpp
+++ b/src/vllm/model_executor/models/laguna.cpp
@@ -35,6 +35,7 @@
 #include <cstdio>
 #include <cstdlib>
 #include <cstring>
+#include <exception>
 #include <limits>
 #include <memory>
 #include <string>
@@ -48,6 +49,7 @@
 #include "vllm/model_executor/models/laguna_ops.h"
 #include "vllm/model_executor/models/qwen3_5_common.h"  // HostLogits
 #include "vt/backend.h"  // vt::GetBackend (device drain for the keep-quant GEMMs)
+#include "vt/breakable_graph.h"  // ENG-CUDAGRAPH-BREAK W5: the shared capture seam
 #include "vt/dtype.h"    // vt::IsBlockQuant / RowSizeBytes
 #include "vt/ops.h"      // vt::MatmulBT (dispatches kMatmulBTQuant on block weights)
 #include "vt/recipes.h"  // vt::kFusedAddRmsNormStd (L4 residual-add + RMSNorm fusion)
@@ -2113,10 +2115,13 @@ std::vector<float> LagunaForwardResidentDecode(const LagunaWeights& weights, vt:
 }
 
 // The decode CUDA-GRAPH capture class is irreducibly device-coupled at build time —
-// graph capture/replay is a CUDA driver concept with no portable vt op today. REPAIR
-// OWED: a portable `vt` capture/replay seam (the same one deepseek_v4.cpp's V4Graph would
-// move behind), after which both graphs leave the shared layer together.
-// DSR-ALLOW(S1): decode CUDA-graph capture class; no portable vt capture seam yet.
+// the CHAIN it captures is `VT_MARLIN_NVFP4` CUDA code, so the class compiles only in
+// that build. The REPAIR this note owed has LANDED: the portable `vt` capture/replay
+// seam is `vt::GraphCaptureScope` + `vt::BreakableGraph`
+// (`include/vt/breakable_graph.h`, row ENG-CUDAGRAPH-BREAK, #1163), and this class and
+// `deepseek_v4.cpp`'s `V4Graph` — the sibling this note named — moved onto it together
+// in W5 (#1335). What is left device-coupled is the captured CHAIN, not the capture.
+// DSR-ALLOW(S1): decode CUDA-graph capture class; the CHAIN is CUDA/Marlin, not the seam.
 #ifdef VT_MARLIN_NVFP4
 namespace {
 // ─── Brick A2: the DECODE CUDA GRAPH (mirror of deepseek_v4.cpp V4Graph) ──────────
@@ -2224,7 +2229,15 @@ struct LagunaGraph {
   vt::Event aux_done{};   // recorded on aux after the shared MLP; main waits it (join)
   std::vector<int64_t> argmax_id;
   int32_t last_sampled = -1;
-  void* graph = nullptr;
+  // ENG-CUDAGRAPH-BREAK W5 (#1335): the instantiated graph, its handle ownership
+  // and its release live in the SHARED SEAM instead of in a raw `void*` this
+  // class destroyed by hand. Every segment is released through
+  // `Backend::DestroyGraph`, which is the routing that lets ENG-CUDAGRAPH-DEDUP
+  // (#1162) interpose at the backend later without editing this file. `gstate`
+  // STAYS: it is this driver's cold/warm/captured ladder, not a duplicate of
+  // `captured()` — the seam has no notion of the eager warm-run that grows the
+  // pool so a capture may allocate nothing.
+  vt::BreakableGraph graph;
   int gstate = 0;  // 0 cold (eager warm-run), 1 warm (capture+replay), 2 captured (replay)
 
   LagunaGraph(const LagunaWeights& w_, vt::Queue& q, LagunaKvCache& cache)
@@ -2370,7 +2383,9 @@ struct LagunaGraph {
   ~LagunaGraph() {
     if (qp == nullptr) return;
     vt::Backend& b = vt::GetBackend(qp->device);
-    if (graph != nullptr) b.DestroyGraph(graph);
+    // No `DestroyGraph` here: `graph` releases its own segments. The aux queue
+    // and its two events are still this class's, because the SEAM tracks a fork
+    // and never owns one — see `vt::GraphNoteFork`.
     if (shared_aux) {
       b.DestroyEvent(aux_fork);
       b.DestroyEvent(aux_done);
@@ -2571,6 +2586,17 @@ struct LagunaGraph {
           vt::Backend& b = vt::GetBackend(dev);
           b.RecordEvent(aux_fork, q);          // event0.record() on the main stream (hn ready)
           b.QueueWaitEvent(aux_q, aux_fork);   // aux waits event0 before reading hn
+          // D10 (ENG-CUDAGRAPH-BREAK W5, #1335): tell the capture scope that a
+          // side queue is now participating in this capture. Closing a segment
+          // while a fork is outstanding FAILS at `cudaStreamEndCapture`, and the
+          // scope is the only thing that can see a segment about to close.
+          // Upstream keeps its whole `torch.cuda.Stream.wait_stream` monkey-patch
+          // (`breakable_cuda_graph.py:101-153`) for no other purpose than to know
+          // this; ours is one call, because our fork is explicit. OUTSIDE a
+          // capture it is a no-op making zero backend calls, so the eager arm is
+          // byte-identical. `aux_done` is the join event, and it outlives every
+          // scope because this class owns it.
+          vt::GraphNoteFork(aux_q, aux_done);
           ActivePoolScope guard(&AuxPool(b));  // shared scratch from AuxPool (see device_pool.h)
           LagunaSharedExpertMarlinInto(aux_q, lw.moe, hn.data(), H, so.data());  // fp4 shared on aux
           b.RecordEvent(aux_done, aux_q);      // event1.record() on the aux stream (join target)
@@ -2609,7 +2635,13 @@ struct LagunaGraph {
         // Make the main stream wait for the aux shared MLP (event1.wait) so the combine
         // below reads a fully-computed `so`. Both the routed path (main) and the shared
         // path (aux) are now complete → the combine result is byte-identical to serial.
-        if (do_aux) vt::GetBackend(dev).QueueWaitEvent(q, aux_done);
+        if (do_aux) {
+          vt::GetBackend(dev).QueueWaitEvent(q, aux_done);
+          // D10: the model joined its own fork, so retire the registration. This
+          // is the common arm and the one that must cost nothing — without it the
+          // scope would issue a REDUNDANT second join before every segment close.
+          vt::GraphNoteJoin(aux_q);
+        }
         // LEVER A: hidden += routed + shared; hn = rms_norm(hidden)*next_norm in ONE node
         // (byte-exact vs AddInto(doutb) + FusedAddNorm(so)). Only in the glue-fused regime
         // (the split path is what it replaces); =0 restores the two-node split.
@@ -2686,13 +2718,56 @@ struct LagunaGraph {
       RunChain();
       gstate = 1;
     } else if (gstate == 1) {  // warm: capture the region once, then replay it
-      b.BeginCapture(q);
-      RunChain();
-      graph = b.EndCaptureGraph(q);
-      b.ReplayGraph(q, graph);
-      gstate = 2;
+      // ENG-CUDAGRAPH-BREAK W5 (#1335): the capture is the SHARED SEAM's, not
+      // this class's hand-rolled `BeginCapture`/`EndCaptureGraph` pair. The scope
+      // owns the segment, the handle, its release, the drain a mid-capture throw
+      // needs, the D10 fork set, and the G3 counters.
+      //
+      // kFULL, INHERITED FROM W2 AND NOT RE-ARGUED. vLLM's v1 default
+      // `FULL_AND_PIECEWISE` (`vllm/config/compilation.py:63` @ pin
+      // `5559679229`) is documented at `:630-632` as a FULL graph for DECODE
+      // batches and a piecewise one for prefill and mixed batches, and
+      // `decode_mode()` (`:65-66`) returns the full half. This is the T=1
+      // resident decode step, so its capture is ONE segment with the attention
+      // calls INSIDE it — byte-identical in shape to the region this replaces.
+      {
+        vt::GraphCaptureScope scope(b, q, graph, vt::GraphCaptureMode::kFull);
+        RunChain();
+      }  // ~GraphCaptureScope joins any outstanding fork, closes the segment,
+         // and files it on `graph`
+      // NOT CAPTURED covers TWO states and they mean opposite things.
+      //
+      //   * FAILED (`capture_failed() == true`): `Backend::EndCaptureGraph`
+      //     threw. Under stream capture NOTHING between `BeginCapture` and the
+      //     throw executed — every kernel was RECORDED — so `logits` (and
+      //     `argmax_id` on the on-device-sample arm) hold whatever the pool last
+      //     left there, and returning them would hand this step uncomputed
+      //     device memory as its logits: no fault, and a token gate cannot see
+      //     it. It PROPAGATES, carrying the runtime's own exception, which is
+      //     what the pre-W5 unguarded `EndCaptureGraph` did.
+      //   * INERT (`capture_failed() == false`): capture is unsupported here, or
+      //     `VLLM_CPP_CUDAGRAPH=0`. The scope made no backend call, `RunChain`
+      //     ran EAGERLY, and the buffers hold a real result — so this step
+      //     returns normally and the driver stays in `gstate == 1`, running
+      //     eager every step rather than pretending to hold a graph.
+      if (!graph.captured()) {
+        if (graph.capture_failed()) {
+          const std::exception_ptr err = graph.capture_error();
+          graph.Reset();  // clear the failure with the graph it described
+          if (err) std::rethrow_exception(err);
+          VT_CHECK(false,
+                   "laguna decode graph: the capture was ABANDONED and its logits were "
+                   "never computed; refusing to return uncaptured device memory");
+        }
+      } else {
+        graph.Replay(q);
+        gstate = 2;
+      }
     } else {                   // captured: one cudaGraphLaunch
-      b.ReplayGraph(q, graph);
+      // Through the seam's container, never `Backend::ReplayGraph` directly: it
+      // replays its segments in order (one, here, because this capture is kFull)
+      // and owns the G3 replay counter.
+      graph.Replay(q);
     }
     // LEVER A: the per-layer K/V append is now IN-GRAPH (append_kv_row writes the new row
     // at the device-read slot *len_buf inside RunChain), so the old between-replay host
diff --git a/src/vllm/model_executor/models/qwen3_dflash.cpp b/src/vllm/model_executor/models/qwen3_dflash.cpp
index 6fd6ba806..b16db7ec6 100644
--- a/src/vllm/model_executor/models/qwen3_dflash.cpp
+++ b/src/vllm/model_executor/models/qwen3_dflash.cpp
@@ -11,13 +11,16 @@
 #include <algorithm>
 #include <cmath>
 #include <cstdlib>
+#include <exception>
 #include <memory>
+#include <optional>
 #include <vector>
 
 #include "vllm/model_executor/layers/linear.h"             // UnquantizedMlpGateUpMethod seam
 #include "vllm/model_executor/models/dense_attn_block.h"  // Dev/DBuf/ResidentWeight/Reshape/MakeRopeArgs
 #include "vllm/platforms/interface.h"                     // platforms::GetPlatform (static-graph gate)
 #include "vt/backend.h"
+#include "vt/breakable_graph.h"  // ENG-CUDAGRAPH-BREAK W5: the shared capture seam
 #include "vt/ops.h"
 
 namespace vllm {
@@ -766,14 +769,16 @@ struct DflashDeviceKVStore {
   std::unique_ptr<DBuf> g_dpos;     // [Tq] i32
   std::unique_ptr<DBuf> g_cu;       // [2] i32 {0, Tq}
   std::unique_ptr<DBuf> g_logits;   // [Tq, vocab] f32 (persistent graph output)
-  void* g_graph = nullptr;          // cudaGraphExec (opaque)
+  // ENG-CUDAGRAPH-BREAK W5 (#1335): the instantiated graph, the ownership of its
+  // handle, its release and its `captured()` state live in the SHARED SEAM
+  // instead of in a raw `void*` plus a `Backend*` this store kept alive only so
+  // its destructor could call `DestroyGraph`. `vt::BreakableGraph` releases every
+  // segment it holds through `Backend::DestroyGraph`, which is the routing that
+  // lets ENG-CUDAGRAPH-DEDUP (#1162) interpose at the backend later without
+  // editing this file.
+  vt::BreakableGraph g_graph;
   int64_t g_tq = -1;                // captured (1+k); -1 = not yet
   int g_state = 0;                  // 0 cold, 1 warm (pool warmed, capture next), 2 captured
-  Backend* g_backend = nullptr;     // for DestroyGraph in the dtor
-
-  ~DflashDeviceKVStore() {
-    if (g_graph != nullptr && g_backend != nullptr) g_backend->DestroyGraph(g_graph);
-  }
 };
 
 std::shared_ptr<DflashDeviceKVStore> Qwen3DFlashModel::MakeDeviceKVStore(
@@ -783,7 +788,6 @@ std::shared_ptr<DflashDeviceKVStore> Qwen3DFlashModel::MakeDeviceKVStore(
   const int64_t Dh = config.head_dim;
   const int64_t L = config.num_hidden_layers;
   auto s = std::make_shared<DflashDeviceKVStore>();
-  s->g_backend = &d.b;
   s->num_layers = L;
   s->block_size = kDflashPageSize;
   s->max_pages = kDflashMaxCtxSlots / kDflashPageSize;
@@ -1008,11 +1012,24 @@ std::vector<float> Qwen3DFlashModel::ForwardBlockLogitsWithDeviceKV(
     VT_CHECK(ctx_cu.back() == static_cast<int32_t>(st.num_ctx),
              "ForwardBlockLogitsWithDeviceKV(paged): ctx_cu.back() must equal store num_ctx");
 
+    // `vt::GraphCaptureEnabled()` is the THIRD conjunct and it is not decoration
+    // (#1352, found and fixed while landing #1335). Before W5 this driver's
+    // capture was its own `BeginCapture` pair, so `VLLM_CPP_CUDAGRAPH` could not
+    // reach it and the two conjuncts below were the whole predicate. The capture
+    // is now the seam's, and the seam reads that switch itself — so without this
+    // conjunct `VLLM_CPP_CUDAGRAPH=0` would still route into the CAPTURE lane,
+    // run the eager warm pass, open an INERT scope, and run the whole
+    // `ForwardPagedBody` a SECOND time inside it. Two full draft forwards per
+    // propose, forever, because the driver would never reach `g_state == 2`.
+    // Not wrong, just wasteful, which is exactly the kind of defect that
+    // survives a token gate. Asking here makes the switch select this driver's
+    // existing single-forward eager path, which is what it means everywhere else.
     const bool graph_ok =
-        UseDflashGraph() && d.b.SupportsGraphCapture() &&
+        UseDflashGraph() && vt::GraphCaptureEnabled() && d.b.SupportsGraphCapture() &&
         platforms::GetPlatform(queue.device.type).support_static_graph_mode();
 
-    // --- Eager paged path (VT_DFLASH_GRAPH=0, or capture unsupported) ---------
+    // --- Eager paged path (VT_DFLASH_GRAPH=0, VLLM_CPP_CUDAGRAPH=0, or capture
+    //     unsupported) --------------------------------------------------------
     if (!graph_ok) {
       DBuf hidden(d, DType::kBF16, {Tq, H});
       {
@@ -1034,10 +1051,12 @@ std::vector<float> Qwen3DFlashModel::ForwardBlockLogitsWithDeviceKV(
     // per config, so this fires once per request lifetime). A width change invalidates a
     // prior graph.
     if (st.g_tq != Tq) {
-      if (st.g_graph != nullptr) {
-        d.b.DestroyGraph(st.g_graph);
-        st.g_graph = nullptr;
-      }
+      // Reset() releases every segment through Backend::DestroyGraph and returns
+      // the container to its as-constructed state, which is also what lets the
+      // next capture open a scope on it: the scope REFUSES a container that
+      // already holds one, because appending to it would leave
+      // `break_count() == segment_count()` and Replay would drop the last break.
+      st.g_graph.Reset();
       st.g_hidden = std::make_unique<DBuf>(d, DType::kBF16, std::vector<int64_t>{Tq, H});
       st.g_dpos = std::make_unique<DBuf>(d, DType::kI32, std::vector<int64_t>{Tq});
       const std::vector<int32_t> cus = {0, static_cast<int32_t>(Tq)};
@@ -1062,7 +1081,10 @@ std::vector<float> Qwen3DFlashModel::ForwardBlockLogitsWithDeviceKV(
     if (st.g_state == 2) {
       // Captured: relaunch the graph over the refreshed persistent inputs + grown context
       // (which enters purely via the in-place seq_lens value + paged store), then download.
-      d.b.ReplayGraph(queue, st.g_graph);
+      // Through the seam's container, never `Backend::ReplayGraph` directly: the
+      // container replays its segments in order (one, here, because this capture
+      // is kFull) and owns the G3 replay counter the reachability gate reads.
+      st.g_graph.Replay(queue);
       st.g_logits->Download(d, out.data());
       if (DflashGraphStats()) {
         ++g_dflash_replays;
@@ -1088,24 +1110,68 @@ std::vector<float> Qwen3DFlashModel::ForwardBlockLogitsWithDeviceKV(
                                       weights, config);
       warm_lg.Download(d, out.data());
     }  // warm_lg + all ForwardPagedBody scratch freed to the pool free-list here.
-    d.b.BeginCapture(queue);
-    try {
-      DBuf lg = ForwardPagedBody(d, st, st.g_hidden->t(), st.g_dpos->t(), st.g_cu->t(),
-                                 weights, config);
-      st.g_graph = d.b.EndCaptureGraph(queue);
-      st.g_logits = std::make_unique<DBuf>(std::move(lg));
-      if (DflashGraphStats()) {
-        ++g_dflash_captures;
-        std::fprintf(stderr, "[DFLASH-GRAPH] captured #%lld Tq=%lld C=%lld\n",
-                     static_cast<long long>(g_dflash_captures), static_cast<long long>(Tq),
-                     static_cast<long long>(st.num_ctx));
+    // ENG-CUDAGRAPH-BREAK W5 (#1335): the capture is the SHARED SEAM's, not this
+    // driver's hand-rolled `BeginCapture`/`EndCaptureGraph` pair with its own
+    // `try`/drain. The scope owns the segment, the handle, its release, the drain
+    // a mid-capture throw needs, and the G3 counters.
+    //
+    // kFULL, INHERITED FROM W2 AND NOT RE-ARGUED. vLLM's v1 default
+    // `FULL_AND_PIECEWISE` (`vllm/config/compilation.py:63` @ pin `5559679229`)
+    // is documented at `:630-632` as a FULL graph for DECODE batches and a
+    // piecewise one for prefill and mixed batches, and `decode_mode()` (`:65-66`)
+    // returns the full half. This is the (1+k) DRAFT step of a speculative
+    // decode, which is a decode batch, so its capture is ONE segment with the
+    // attention calls INSIDE it — byte-identical in shape to the region this
+    // replaces. Opening it kPiecewise would turn every draft layer's attention
+    // into an eager call between two graph replays, which is not vLLM's decode
+    // behaviour and which nothing in this row's record supports.
+    std::optional<DBuf> lg;
+    {
+      vt::GraphCaptureScope scope(d.b, queue, st.g_graph, vt::GraphCaptureMode::kFull);
+      lg = ForwardPagedBody(d, st, st.g_hidden->t(), st.g_dpos->t(), st.g_cu->t(),
+                            weights, config);
+    }  // ~GraphCaptureScope closes the segment and files it on st.g_graph
+    // NOT CAPTURED covers TWO states, and only one of them may continue.
+    //
+    //   * INERT (`capture_failed() == false`): unreachable here, because
+    //     `graph_ok` above already required `SupportsGraphCapture()`; the
+    //     remaining inert cause is `VLLM_CPP_CUDAGRAPH=0`, which the seam reads
+    //     and this driver no longer does. The region ran EAGERLY, `*lg` is a real
+    //     result, and the step falls back to the eager lane for good.
+    //   * FAILED (`capture_failed() == true`): `Backend::EndCaptureGraph` threw.
+    //     Under stream capture NOTHING between `BeginCapture` and the throw
+    //     executed — every kernel was RECORDED — so `*lg` is pool-recycled memory
+    //     and downloading it would hand this draft step uncomputed device memory
+    //     as its logits. No fault, and a token gate cannot see it, because a
+    //     draft the target rejects is indistinguishable from a bad draft.
+    //
+    // The pre-W5 driver rethrew after draining, and so does this.
+    if (!st.g_graph.captured()) {
+      if (st.g_graph.capture_failed()) {
+        const std::exception_ptr err = st.g_graph.capture_error();
+        st.g_graph.Reset();  // clear the failure with the graph it described
+        std::fprintf(stderr, "[DFLASH-GRAPH] capture FAILED\n");
+        if (err) std::rethrow_exception(err);
+        VT_CHECK(false,
+                 "DFlash draft graph: the capture was ABANDONED and its logits were "
+                 "never computed; refusing to return uncaptured device memory");
       }
-    } catch (const std::exception& e) {
-      std::fprintf(stderr, "[DFLASH-GRAPH] capture FAILED: %s\n", e.what());
-      void* g = nullptr;
-      try { g = d.b.EndCaptureGraph(queue); } catch (...) {}  // drain, unstick the stream
-      if (g != nullptr) d.b.DestroyGraph(g);
-      throw;
+      // INERT is now UNREACHABLE from here: `graph_ok` above required both
+      // `vt::GraphCaptureEnabled()` and `SupportsGraphCapture()`, which are the
+      // only two things that make a scope inert. Kept as a total branch rather
+      // than an assertion, because a future inert cause added to the seam must
+      // degrade to a correct eager step rather than to undefined behaviour —
+      // the region DID run eagerly, so `*lg` holds real values.
+      st.g_state = 0;  // stay eager, and re-warm rather than re-capture
+      lg->Download(d, out.data());
+      return out;
+    }
+    st.g_logits = std::make_unique<DBuf>(std::move(*lg));
+    if (DflashGraphStats()) {
+      ++g_dflash_captures;
+      std::fprintf(stderr, "[DFLASH-GRAPH] captured #%lld Tq=%lld C=%lld\n",
+                   static_cast<long long>(g_dflash_captures), static_cast<long long>(Tq),
+                   static_cast<long long>(st.num_ctx));
     }
     st.g_state = 2;  // subsequent steps replay
     return out;      // this step's output is the eager warm pass (bit-identical to the graph)
diff --git a/src/vt/breakable_graph.cpp b/src/vt/breakable_graph.cpp
index 3a791c493..71b0520b5 100644
--- a/src/vt/breakable_graph.cpp
+++ b/src/vt/breakable_graph.cpp
@@ -27,6 +27,8 @@ std::atomic<int64_t> g_breaks{0};
 std::atomic<int64_t> g_replays{0};
 std::atomic<int64_t> g_full_scopes{0};
 std::atomic<int64_t> g_piecewise_scopes{0};
+std::atomic<int64_t> g_forks_tracked{0};
+std::atomic<int64_t> g_forks_auto_joined{0};
 
 }  // namespace
 
@@ -42,6 +44,8 @@ GraphBreakStats GetGraphBreakStats() {
   s.replays = g_replays.load(std::memory_order_relaxed);
   s.full_scopes = g_full_scopes.load(std::memory_order_relaxed);
   s.piecewise_scopes = g_piecewise_scopes.load(std::memory_order_relaxed);
+  s.forks_tracked = g_forks_tracked.load(std::memory_order_relaxed);
+  s.forks_auto_joined = g_forks_auto_joined.load(std::memory_order_relaxed);
   return s;
 }
 
@@ -52,6 +56,8 @@ void ResetGraphBreakStats() {
   g_replays.store(0, std::memory_order_relaxed);
   g_full_scopes.store(0, std::memory_order_relaxed);
   g_piecewise_scopes.store(0, std::memory_order_relaxed);
+  g_forks_tracked.store(0, std::memory_order_relaxed);
+  g_forks_auto_joined.store(0, std::memory_order_relaxed);
 }
 
 // Read ONCE into a function-local static, so a process is in exactly one lane
@@ -252,6 +258,11 @@ void GraphCaptureScope::BeginSegment() {
 
 void GraphCaptureScope::EndSegment() {
   if (!active_ || !segment_open_) return;
+  // D10, and it runs BEFORE the close rather than after it, because after is too
+  // late: `cudaStreamEndCapture` on a capture graph that still has an unjoined
+  // fork FAILS. Upstream does exactly this inside `_end_current_segment`
+  // (`:353-361`), immediately before `capture_end()`.
+  JoinOutstandingForks();
   segment_open_ = false;  // cleared FIRST: a throwing end must not be retried
   void* seg = b_->EndCaptureGraph(*q_);
   g_->segments_.push_back(seg);
@@ -290,6 +301,63 @@ void GraphCaptureScope::AppendBreak(std::function<void()> fn, const void* destin
   g_breaks.fetch_add(1, std::memory_order_relaxed);
 }
 
+// ---------------------------------------------------------------------------
+// D10 — the outstanding-fork set. Port of the `wait_stream` hook (`:101-153`)
+// and of the auto-join in `_end_current_segment` (`:353-361`).
+// ---------------------------------------------------------------------------
+
+void GraphCaptureScope::NoteFork(Queue& aux, Event& join_event) {
+  if (!active_) return;
+  // Registering the SAME queue twice would make the auto-join issue two waits
+  // for one fork, and — worse — would leave one entry behind after `NoteJoin`
+  // retired the other, so a joined queue would still look outstanding. The
+  // identity is the QUEUE, because that is what a join targets.
+  for (ForkedQueue& f : forks_) {
+    if (f.q == &aux) {
+      f.e = &join_event;  // a re-fork of the same queue: keep ONE entry, newest event
+      return;
+    }
+  }
+  forks_.push_back(ForkedQueue{&aux, &join_event});
+  g_forks_tracked.fetch_add(1, std::memory_order_relaxed);
+}
+
+void GraphCaptureScope::NoteJoin(Queue& aux) {
+  if (!active_) return;
+  for (size_t i = 0; i < forks_.size(); ++i) {
+    if (forks_[i].q == &aux) {
+      forks_.erase(forks_.begin() + static_cast<long>(i));
+      return;
+    }
+  }
+}
+
+void GraphCaptureScope::JoinOutstandingForks() {
+  if (!active_ || forks_.empty()) return;
+  for (ForkedQueue& f : forks_) {
+    // The join upstream spells `self.stream.wait_stream(s)` (`:359`): make the
+    // CAPTURE queue wait for everything submitted to the side queue. Ours is the
+    // explicit two-call form — record on the side queue, wait on the capture
+    // queue — which is the same pair the model itself issues when it joins.
+    b_->RecordEvent(*f.e, *f.q);
+    b_->QueueWaitEvent(*q_, *f.e);
+    g_forks_auto_joined.fetch_add(1, std::memory_order_relaxed);
+  }
+  forks_.clear();
+}
+
+void GraphNoteFork(Queue& aux, Event& join_event) {
+  GraphCaptureScope* s = GraphCaptureScope::Current();
+  if (s == nullptr) return;
+  s->NoteFork(aux, join_event);
+}
+
+void GraphNoteJoin(Queue& aux) {
+  GraphCaptureScope* s = GraphCaptureScope::Current();
+  if (s == nullptr) return;
+  s->NoteJoin(aux);
+}
+
 // The bare marker (`:370-374`).
 void GraphBreak() {
   GraphCaptureScope* s = GraphCaptureScope::Current();
diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt
index 68b42bb3d..de8a9ac35 100644
--- a/tests/CMakeLists.txt
+++ b/tests/CMakeLists.txt
@@ -432,6 +432,7 @@ vllm_cpp_add_test(test_qwen3_moe_decode_graph_seam vllm/models/test_qwen3_moe_de
 vllm_cpp_add_test(test_voxtral_decode_graph_seam vllm/models/test_voxtral_decode_graph_seam.cpp)
 vllm_cpp_add_test(test_deepseek_v2_decode_graph_seam vllm/models/test_deepseek_v2_decode_graph_seam.cpp)
 vllm_cpp_add_test(test_qwen3_5_decode_graph_seam vllm/models/test_qwen3_5_decode_graph_seam.cpp)
+vllm_cpp_add_test(test_qwen3_dflash_decode_graph_seam vllm/models/test_qwen3_dflash_decode_graph_seam.cpp)
 # G1 for the row: bit-exactness against eager over THREE replays, on a real
 # device. Inert without a CUDA backend (every case skips), because a CPU
 # "replay" recomputes nothing.
diff --git a/tests/vllm/models/decode_graph_seam_harness.h b/tests/vllm/models/decode_graph_seam_harness.h
index 5be7f4055..f3325a62e 100644
--- a/tests/vllm/models/decode_graph_seam_harness.h
+++ b/tests/vllm/models/decode_graph_seam_harness.h
@@ -49,14 +49,34 @@ namespace vllm_test {
 // vocabulary by logging.
 class CaptureCapableCpuBackend final : public vt::Backend {
  public:
-  explicit CaptureCapableCpuBackend(vt::Backend& inner) : inner_(inner) {}
-
-  void* Alloc(size_t bytes) override { return inner_.Alloc(bytes); }
+  // `supports_capture` is the CONTROL LANE switch. Passing false makes this the
+  // same delegating backend that answers `SupportsGraphCapture()` FALSE, so a
+  // driver's own predicate sends it down its eager path — which is the lane a
+  // gate has to compare against to tell "the switch chose eager" apart from
+  // "the switch made the capture inert and the forward ran twice" (#1352).
+  explicit CaptureCapableCpuBackend(vt::Backend& inner, bool supports_capture = true)
+      : inner_(inner), supports_capture_(supports_capture) {}
+
+  // ALLOCATIONS ARE COUNTED, and the reason is LANE IDENTITY rather than memory.
+  // ENG-CUDAGRAPH-BREAK W5 (#1352) needed to tell apart two versions of a driver
+  // that produce IDENTICAL logits and BOTH report `segments_captured == 0` — the
+  // one that takes its eager path under `VLLM_CPP_CUDAGRAPH=0`, and the one that
+  // takes the CAPTURE path and runs the forward twice inside an inert scope.
+  // Neither a token gate nor the seam's counters can see that difference. What
+  // can is the comparison against a backend that CANNOT capture: with the switch
+  // off the two must allocate exactly the same, because they must be running the
+  // same code. The count is an EQUALITY against a control, never a magic number,
+  // so it does not go stale when the driver's allocation pattern changes.
+  void* Alloc(size_t bytes) override {
+    ++allocs_;
+    return inner_.Alloc(bytes);
+  }
   void Free(void* p) override { inner_.Free(p); }
   void Memset(vt::Queue& q, void* p, int v, size_t bytes) override {
     inner_.Memset(q, p, v, bytes);
   }
   void Copy(vt::Queue& q, void* dst, const void* src, size_t bytes) override {
+    ++copies_;
     inner_.Copy(q, dst, src, bytes);
   }
   vt::Queue CreateQueue() override { return inner_.CreateQueue(); }
@@ -78,7 +98,7 @@ class CaptureCapableCpuBackend final : public vt::Backend {
   }
   bool SupportsAuxStream() const override { return inner_.SupportsAuxStream(); }
 
-  bool SupportsGraphCapture() const override { return true; }
+  bool SupportsGraphCapture() const override { return supports_capture_; }
   void BeginCapture(vt::Queue&) override { log_.push_back("Begin"); }
   void* EndCaptureGraph(vt::Queue&) override {
     // Arm-once refusal, the shape of a real `cudaStreamEndCapture` returning
@@ -106,6 +126,20 @@ class CaptureCapableCpuBackend final : public vt::Backend {
 
   void FailNextEndCapture() { fail_next_end_ = true; }
 
+  // COPIES, not allocations, and the difference is the DevicePool. A pooled
+  // allocator serves the second of two identical forwards out of its free list
+  // without calling `Alloc` at all, so an allocation count reads 11 for the
+  // first lane measured and 0 for the second REGARDLESS of which lanes they
+  // were — an instrument that reports the ORDER of the runs rather than their
+  // shape. Copies are per-operation work the pool cannot absorb, so they count
+  // the forward rather than the allocator's history (#1352).
+  int64_t copies() const { return copies_; }
+  int64_t allocs() const { return allocs_; }
+  void ResetCounters() {
+    copies_ = 0;
+    allocs_ = 0;
+  }
+
   size_t Count(std::string_view what) const {
     size_t n = 0;
     for (const auto& e : log_)
@@ -115,7 +149,10 @@ class CaptureCapableCpuBackend final : public vt::Backend {
 
  private:
   vt::Backend& inner_;
+  bool supports_capture_ = true;
   bool fail_next_end_ = false;
+  int64_t allocs_ = 0;
+  int64_t copies_ = 0;
   std::vector<std::string> log_;
   std::vector<void*> tags_;
 };
@@ -168,10 +205,10 @@ class StaticGraphCpuPlatform final : public vllm::platforms::Platform {
 // was there.
 class StaticGraphCpu {
  public:
-  StaticGraphCpu()
+  explicit StaticGraphCpu(bool supports_capture = true)
       : prev_backend_(&vt::GetBackend(vt::DeviceType::kCPU)),
         prev_platform_(&vllm::platforms::GetPlatform(vt::DeviceType::kCPU)),
-        backend_(*prev_backend_),
+        backend_(*prev_backend_, supports_capture),
         platform_(*prev_platform_, backend_) {
     vt::RegisterBackend(vt::DeviceType::kCPU, &backend_);
     vllm::platforms::RegisterPlatform(vt::DeviceType::kCPU, &platform_);
diff --git a/tests/vllm/models/test_qwen3_dflash_decode_graph_seam.cpp b/tests/vllm/models/test_qwen3_dflash_decode_graph_seam.cpp
new file mode 100644
index 000000000..274f5edc2
--- /dev/null
+++ b/tests/vllm/models/test_qwen3_dflash_decode_graph_seam.cpp
@@ -0,0 +1,415 @@
+// THE G2 REACHABILITY GATE for the DFlash DRAFT graph on the break-point capture
+// seam. Row `ENG-CUDAGRAPH-BREAK` W5, issue #1335, parent #1163; spec
+// `.agents/specs/eng-cudagraph-break.md` `## Gates` G2 and `## Work breakdown` W5.
+//
+// WHY THIS FILE HAS TO EXIST, and it is the same argument W3 made for each of its
+// three drivers. A driver that kept its hand-rolled `BeginCapture` /
+// `EndCaptureGraph` pair produces IDENTICAL logits, an IDENTICAL backend call log
+// and an identical private replay count. `vt::GraphBreakStats::segments_captured`
+// moves only when a `vt::GraphCaptureScope` closes a segment, and `replays` only
+// inside `vt::BreakableGraph::Replay`, so those two counters are the ONLY
+// observables that separate "captured a graph" from "captured a graph THROUGH THE
+// SEAM". Nothing else in this tree can see the difference, which is exactly why a
+// migration without its own gate is a refactor nobody can prove happened.
+//
+// WHY THE DFLASH DRIVER IS THE ONE SINGLE-SHAPE DRIVER THAT CAN BE GATED HERE.
+// The other two W5 drivers refuse a CPU queue before they reach their capture:
+// DeepSeek V4's `CanRunResidentDecode` returns false for
+// `device.type == kCPU` and again unless the four CUDA-registered V4 kernel
+// families are present, and Laguna's whole capture class is compiled only under
+// `VT_MARLIN_NVFP4`, which needs a CUDA build on a marlin-nvfp4 architecture.
+// The DFlash draft graph's admission predicate names neither a device type nor a
+// kernel registry — it is `VT_DFLASH_GRAPH` plus
+// `Backend::SupportsGraphCapture()` plus `Platform::support_static_graph_mode()`
+// — so the shared harness's two swapped registries are enough to reach it. The
+// spec's `## Owed` records the other two as owed on hardware rather than
+// implying this file covers them.
+//
+// WHAT THIS HARNESS CANNOT SEE, named rather than claimed away: CPU kernels are
+// direct calls, so a "captured" region EXECUTES eagerly here and a "replay"
+// recomputes nothing. This gate holds the ROUTING and the capture step's
+// numerics; that a replayed segment reproduces the eager forward is G1, it needs
+// a real device, and `## Owed` records it.
+#include <doctest/doctest.h>
+
+#include <cmath>
+#include <cstdint>
+#include <cstdlib>
+#include <filesystem>
+#include <fstream>
+#include <memory>
+#include <string>
+#include <vector>
+
+#ifdef __linux__
+#include <unistd.h>
+#endif
+
+#include "vllm/model_executor/models/qwen3_dflash.h"
+#include "vllm/transformers_utils/hf_config.h"
+#include "vt/backend.h"
+#include "vt/breakable_graph.h"
+#include "vt/dtype.h"
+
+#include "decode_graph_seam_harness.h"
+
+using namespace vllm;
+
+namespace {
+
+vt::Queue Cpu() { return vt::Queue{vt::Device{vt::DeviceType::kCPU, 0}, nullptr}; }
+
+// The same deterministic synthetic weights `tests/vllm/v1/spec_decode/
+// test_dflash_propose.cpp` builds, so this gate's model is the one the existing
+// device-KV bit-exactness case already pins rather than a second invention.
+OwnedTensor MkBf16(const std::vector<int64_t>& shape, double seed, double amp, bool nk) {
+  OwnedTensor t;
+  t.dtype = vt::DType::kBF16;
+  t.rank = static_cast<int>(shape.size());
+  t.nk = nk;
+  int64_t n = 1;
+  for (int i = 0; i < t.rank; ++i) {
+    t.shape[i] = shape[static_cast<size_t>(i)];
+    n *= t.shape[i];
+  }
+  t.bytes.resize(static_cast<size_t>(n) * sizeof(uint16_t));
+  auto* p = reinterpret_cast<uint16_t*>(t.bytes.data());
+  for (int64_t i = 0; i < n; ++i)
+    p[i] = vt::F32ToBF16(static_cast<float>(amp * std::sin(seed + 0.7 * static_cast<double>(i))));
+  return t;
+}
+
+struct Dims {
+  int64_t H = 4, Hq = 2, Hkv = 1, Dh = 2, I = 6, vocab = 8, layers = 2, taps = 2;
+};
+
+HfConfig MakeConfig(const Dims& dm) {
+  HfConfig c;
+  c.hidden_size = dm.H;
+  c.num_attention_heads = dm.Hq;
+  c.num_key_value_heads = dm.Hkv;
+  c.head_dim = dm.Dh;
+  c.rotary_dim = dm.Dh;
+  c.rope_theta = 10000.0;
+  c.intermediate_size = dm.I;
+  c.vocab_size = dm.vocab;
+  c.num_hidden_layers = dm.layers;
+  c.rms_norm_eps = 1e-6;
+  c.sliding_window = 64;
+  c.layer_types = {"sliding_attention", "full_attention"};
+  c.raw = nlohmann::json::object();
+  c.raw["dflash_config"] = {{"mask_token_id", 7}};
+  return c;
+}
+
+Qwen3DFlashWeights MakeWeights(const Dims& dm) {
+  Qwen3DFlashWeights w;
+  w.num_taps = dm.taps;
+  w.mask_token_id = 7;
+  w.draft_vocab_size = dm.vocab;
+  const int64_t qdim = dm.Hq * dm.Dh, kdim = dm.Hkv * dm.Dh;
+  w.embed_tokens = MkBf16({dm.vocab, dm.H}, 0.1, 0.3, false);
+  w.fc = MkBf16({dm.H, dm.H * dm.taps}, 0.2, 0.2, true);
+  w.hidden_norm = MkBf16({dm.H}, 0.3, 0.5, false);
+  w.final_norm = MkBf16({dm.H}, 0.4, 0.5, false);
+  w.lm_head = MkBf16({dm.vocab, dm.H}, 0.5, 0.3, true);
+  const std::vector<Qwen3DFlashLayerAttnMode> modes = {{true, 64}, {false, 0}};
+  for (int64_t l = 0; l < dm.layers; ++l) {
+    Qwen3DFlashLayerWeights lw;
+    const double s = 1.0 + static_cast<double>(l);
+    lw.input_layernorm = MkBf16({dm.H}, s + 0.01, 0.5, false);
+    lw.post_attention_layernorm = MkBf16({dm.H}, s + 0.02, 0.5, false);
+    lw.qkv_proj = MkBf16({qdim + 2 * kdim, dm.H}, s + 0.03, 0.25, true);
+    lw.o_proj = MkBf16({dm.H, qdim}, s + 0.04, 0.25, true);
+    lw.q_norm = MkBf16({dm.Dh}, s + 0.05, 0.5, false);
+    lw.k_norm = MkBf16({dm.Dh}, s + 0.06, 0.5, false);
+    lw.gate_up_proj = MkBf16({2 * dm.I, dm.H}, s + 0.07, 0.2, true);
+    lw.down_proj = MkBf16({dm.H, dm.I}, s + 0.08, 0.2, true);
+    lw.attn_mode = modes[static_cast<size_t>(l)];
+    w.layers.push_back(std::move(lw));
+  }
+  return w;
+}
+
+// One context feature row per position, deterministic in the position.
+std::vector<float> Ctx(int64_t rows, int64_t H) {
+  std::vector<float> c(static_cast<size_t>(rows * H));
+  for (size_t i = 0; i < c.size(); ++i)
+    c[i] = 0.2f * static_cast<float>(std::sin(0.13 * static_cast<double>(i) + 0.4));
+  return c;
+}
+
+}  // namespace
+
+// ---------------------------------------------------------------------------
+// G2 — the DFlash draft graph is captured and replayed THROUGH THE SEAM.
+// ---------------------------------------------------------------------------
+//
+// The step ladder mirrors the production propose loop: the store is appended to
+// once per accepted token and `ForwardBlockLogitsWithDeviceKV` runs the (1+k)
+// block per propose. The driver's own state machine is cold-with-capture on the
+// FIRST propose (its warm pass and its capture happen in one step, because a
+// full target verify runs between two draft steps and perturbs the shared pool),
+// then replay on every later one. So after N proposes the seam must report ONE
+// closed segment and N-1 replays.
+TEST_CASE("dflash draft graph: the capture and every replay go through vt::BreakableGraph") {
+  vllm_test::StaticGraphCpu graph_cpu;  // capture-capable CPU backend + static-graph platform
+
+  Dims dm;
+  HfConfig cfg = MakeConfig(dm);
+  Qwen3DFlashWeights w = MakeWeights(dm);
+  vt::Queue q = Cpu();
+  const int64_t H = dm.H;
+
+  auto store = Qwen3DFlashModel::MakeDeviceKVStore(cfg, q);
+  const std::vector<float> ctx = Ctx(3, H);
+  Qwen3DFlashModel::AppendContextKVDevice(*store, ctx, {0, 1, 2}, w, cfg, q);
+  REQUIRE(Qwen3DFlashModel::DeviceKVNumCtx(*store) == 3);
+
+  std::vector<DflashDeviceKVStore*> stores = {store.get()};
+  const std::vector<int32_t> ids = {2, 7, 7};   // anchor + two MASK rows: Tq = 1 + k
+  const std::vector<int32_t> pos = {3, 4, 5};
+  const std::vector<int32_t> block_cu = {0, 3};
+  const std::vector<int32_t> ctx_cu = {0, 3};
+
+  vt::ResetGraphBreakStats();
+
+  // PROPOSE 1 — the driver's cold step: it runs the eager warm pass, then opens
+  // the seam's scope over the identical forward and files the segment.
+  const std::vector<float> first = Qwen3DFlashModel::ForwardBlockLogitsWithDeviceKV(
+      stores, ctx_cu, ids, pos, block_cu, w, cfg, q);
+  REQUIRE(first.size() == static_cast<size_t>(3) * dm.vocab);
+  {
+    const vt::GraphBreakStats s = vt::GetGraphBreakStats();
+    // ONE segment closed by a scope, and the mode is FULL. Both are asserted,
+    // and `full_scopes` is not decoration: W3 measured that a `kFull` ->
+    // `kPiecewise` flip in a migrated driver compiled clean and left the whole
+    // gate green, because the mode was unobservable from outside the driver.
+    CHECK(s.segments_captured == 1);
+    CHECK(s.full_scopes == 1);
+    CHECK(s.piecewise_scopes == 0);
+    // NO BREAK POINT IS REGISTERED, and that is a decision rather than an
+    // omission. Under `kFull` a `vt::GraphBreak` takes the same pass-through arm
+    // it takes outside a scope, so registering one here would land machinery no
+    // gate could exercise. The DFlash draft forward reaches no `GraphBreak` site
+    // at all today, which is what this reads.
+    CHECK(s.breaks_registered == 0);
+    CHECK(s.replays == 0);  // the cold step's own output is its eager warm pass
+  }
+
+  // PROPOSE 2 and 3 — replays, and they must go through the CONTAINER. A driver
+  // that called `Backend::ReplayGraph` on a raw handle would leave the backend
+  // log identical and this counter at zero.
+  for (int i = 0; i < 2; ++i) {
+    const std::vector<float> again = Qwen3DFlashModel::ForwardBlockLogitsWithDeviceKV(
+        stores, ctx_cu, ids, pos, block_cu, w, cfg, q);
+    REQUIRE(again.size() == first.size());
+  }
+  {
+    const vt::GraphBreakStats s = vt::GetGraphBreakStats();
+    CHECK(s.segments_captured == 1);  // captured ONCE for this block width
+    CHECK(s.replays == 2);
+    CHECK(s.full_scopes == 1);
+  }
+}
+
+// ---------------------------------------------------------------------------
+// G4 — the capture step's logits are bit-identical to the eager arm.
+// ---------------------------------------------------------------------------
+//
+// This is what makes the migration reversible. `VT_DFLASH_GRAPH=0` selects the
+// driver's OWN eager paged path, over the same store and the same weights, and
+// the two must agree to the bit. Exact equality, because the graph lane is not
+// an approximation of the eager one: it is the same kernels over the same
+// persistent buffers.
+TEST_CASE("dflash draft graph: the seam capture step is bit-identical to the eager arm") {
+  Dims dm;
+  HfConfig cfg = MakeConfig(dm);
+  Qwen3DFlashWeights w = MakeWeights(dm);
+  vt::Queue q = Cpu();
+  const int64_t H = dm.H;
+  const std::vector<float> ctx = Ctx(3, H);
+  const std::vector<int32_t> ids = {2, 7, 7};
+  const std::vector<int32_t> pos = {3, 4, 5};
+  const std::vector<int32_t> block_cu = {0, 3};
+  const std::vector<int32_t> ctx_cu = {0, 3};
+
+  // EAGER: no capture-capable backend, so `graph_ok` is false and the driver
+  // takes its own eager paged lane.
+  auto eager_store = Qwen3DFlashModel::MakeDeviceKVStore(cfg, q);
+  Qwen3DFlashModel::AppendContextKVDevice(*eager_store, ctx, {0, 1, 2}, w, cfg, q);
+  std::vector<DflashDeviceKVStore*> eager_stores = {eager_store.get()};
+  const std::vector<float> eager = Qwen3DFlashModel::ForwardBlockLogitsWithDeviceKV(
+      eager_stores, ctx_cu, ids, pos, block_cu, w, cfg, q);
+
+  // SEAM: the identical call under the capture-capable registries.
+  std::vector<float> captured;
+  {
+    vllm_test::StaticGraphCpu graph_cpu;
+    vt::ResetGraphBreakStats();
+    auto store = Qwen3DFlashModel::MakeDeviceKVStore(cfg, q);
+    Qwen3DFlashModel::AppendContextKVDevice(*store, ctx, {0, 1, 2}, w, cfg, q);
+    std::vector<DflashDeviceKVStore*> stores = {store.get()};
+    captured = Qwen3DFlashModel::ForwardBlockLogitsWithDeviceKV(stores, ctx_cu, ids, pos,
+                                                               block_cu, w, cfg, q);
+    REQUIRE(vt::GetGraphBreakStats().segments_captured == 1);
+  }
+
+  REQUIRE(captured.size() == eager.size());
+  size_t differing = 0;
+  for (size_t i = 0; i < captured.size(); ++i)
+    if (captured[i] != eager[i]) ++differing;
+  CHECK(differing == 0);
+}
+
+// ---------------------------------------------------------------------------
+// THE FAILED-CAPTURE ARM — a capture that was abandoned must NOT return its
+// buffer.
+// ---------------------------------------------------------------------------
+//
+// Under stream capture nothing between `BeginCapture` and a throwing
+// `EndCaptureGraph` executed: the kernels were RECORDED, not run. So the logits
+// buffer the abandoned region produced holds whatever the pool last left there,
+// and downloading it would hand the speculator uncomputed device memory as its
+// draft. No fault, and a token gate cannot see it, because a draft the target
+// rejects looks exactly like a bad draft. The pre-W5 driver rethrew after
+// draining the stream and so must this one.
+TEST_CASE("dflash draft graph: an abandoned capture propagates instead of returning garbage") {
+  vllm_test::StaticGraphCpu graph_cpu;
+
+  Dims dm;
+  HfConfig cfg = MakeConfig(dm);
+  Qwen3DFlashWeights w = MakeWeights(dm);
+  vt::Queue q = Cpu();
+  const std::vector<float> ctx = Ctx(3, dm.H);
+  auto store = Qwen3DFlashModel::MakeDeviceKVStore(cfg, q);
+  Qwen3DFlashModel::AppendContextKVDevice(*store, ctx, {0, 1, 2}, w, cfg, q);
+  std::vector<DflashDeviceKVStore*> stores = {store.get()};
+  const std::vector<int32_t> ids = {2, 7, 7};
+  const std::vector<int32_t> pos = {3, 4, 5};
+  const std::vector<int32_t> block_cu = {0, 3};
+  const std::vector<int32_t> ctx_cu = {0, 3};
+
+  graph_cpu.backend().FailNextEndCapture();
+  CHECK_THROWS(Qwen3DFlashModel::ForwardBlockLogitsWithDeviceKV(stores, ctx_cu, ids, pos,
+                                                                block_cu, w, cfg, q));
+
+  // AND THE STORE RECOVERS. The drain reset the container, so the very next
+  // propose captures cleanly rather than inheriting a poisoned state — which is
+  // the half a `throw` alone would not give.
+  vt::ResetGraphBreakStats();
+  const std::vector<float> after = Qwen3DFlashModel::ForwardBlockLogitsWithDeviceKV(
+      stores, ctx_cu, ids, pos, block_cu, w, cfg, q);
+  CHECK(after.size() == static_cast<size_t>(3) * dm.vocab);
+  CHECK(vt::GetGraphBreakStats().segments_captured == 1);
+}
+
+
+// ---------------------------------------------------------------------------
+// #1352 — THE KILL SWITCH MUST SELECT THE EAGER LANE, not an inert capture lane.
+// ---------------------------------------------------------------------------
+//
+// WHY THIS NEEDS AN UNUSUAL INSTRUMENT, stated before the code so the shape is
+// not mistaken for cleverness. The defect it gates produces IDENTICAL logits and
+// leaves `segments_captured` at 0 in BOTH versions, so a token gate is blind to
+// it and so are the seam's own counters. What differs is which CODE PATH ran:
+// with the fix, `VLLM_CPP_CUDAGRAPH=0` sends the driver down its existing eager
+// paged path; without it, the driver takes the CAPTURE path, runs its eager warm
+// pass, opens an INERT scope and runs the whole forward a SECOND time inside it.
+//
+// So the assertion is LANE IDENTITY against a control that cannot be argued
+// with: a backend that CANNOT capture takes the eager path by construction, and
+// a capture-capable backend with the switch OFF must then allocate EXACTLY the
+// same amount, because it must be running exactly the same code. An equality
+// against a live control, never a magic number, so it does not go stale when the
+// driver's allocation pattern changes.
+//
+// IT RUNS IN A CHILD PROCESS for the reason `test_breakable_graph`'s T4 does:
+// `vt::GraphCaptureEnabled()` is read once per process into a function-local
+// static, so no case inside a running process can toggle it.
+TEST_CASE("dflash draft graph: VLLM_CPP_CUDAGRAPH=0 takes the EAGER lane, not an inert capture") {
+  const char* sentinel_path = std::getenv("VLLM_CPP_DFLASH_KILLSWITCH_SENTINEL");
+  if (sentinel_path != nullptr) {
+    // ---- CHILD ARM: this process was started with VLLM_CPP_CUDAGRAPH=0. ----
+    REQUIRE_FALSE(vt::GraphCaptureEnabled());
+
+    Dims dm;
+    HfConfig cfg = MakeConfig(dm);
+    Qwen3DFlashWeights w = MakeWeights(dm);
+    vt::Queue q = Cpu();
+    const std::vector<float> ctx = Ctx(3, dm.H);
+    const std::vector<int32_t> ids = {2, 7, 7};
+    const std::vector<int32_t> pos = {3, 4, 5};
+    const std::vector<int32_t> block_cu = {0, 3};
+    const std::vector<int32_t> ctx_cu = {0, 3};
+
+    // Measures one propose on a given lane: the allocation count and the logits.
+    const auto run = [&](bool supports_capture, int64_t* copies) {
+      vllm_test::StaticGraphCpu harness(supports_capture);
+      auto store = Qwen3DFlashModel::MakeDeviceKVStore(cfg, q);
+      Qwen3DFlashModel::AppendContextKVDevice(*store, ctx, {0, 1, 2}, w, cfg, q);
+      std::vector<DflashDeviceKVStore*> stores = {store.get()};
+      harness.backend().ResetCounters();
+      std::vector<float> out = Qwen3DFlashModel::ForwardBlockLogitsWithDeviceKV(
+          stores, ctx_cu, ids, pos, block_cu, w, cfg, q);
+      *copies = harness.backend().copies();
+      return out;
+    };
+
+    vt::ResetGraphBreakStats();
+    int64_t cannot_capture_copies = 0, switch_off_copies = 0;
+    const std::vector<float> control = run(/*supports_capture=*/false, &cannot_capture_copies);
+    const std::vector<float> switched_off = run(/*supports_capture=*/true, &switch_off_copies);
+
+    // Nothing captured in either lane, which is the part BOTH versions satisfy
+    // and which is therefore NOT the discriminator. It is asserted anyway,
+    // because a version that captured here would be a different defect.
+    REQUIRE(vt::GetGraphBreakStats().segments_captured == 0);
+    REQUIRE(cannot_capture_copies > 0);  // the control actually did work
+
+    // THE DISCRIMINATOR. Same lane => same allocations, exactly.
+    const bool same_lane = (switch_off_copies == cannot_capture_copies);
+    // ... and the same numbers, which holds in both versions and is here so a
+    // failure of the line above cannot be dismissed as the two lanes computing
+    // different things.
+    bool identical = control.size() == switched_off.size();
+    for (size_t i = 0; identical && i < control.size(); ++i)
+      if (control[i] != switched_off[i]) identical = false;
+
+    std::ofstream out(sentinel_path);
+    out << (same_lane ? "SAME-LANE" : "DIFFERENT-LANE") << " " << cannot_capture_copies << " "
+        << switch_off_copies << " " << (identical ? "IDENTICAL" : "DIVERGED");
+    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_dflash_killswitch_" + std::to_string(::getpid()) + ".txt");
+  std::filesystem::remove(sentinel);
+  // The filter must match THIS case only, and doctest splits `-tc` on COMMAS,
+  // so the pattern carries none.
+  const std::string cmd = "VLLM_CPP_CUDAGRAPH=0 VLLM_CPP_DFLASH_KILLSWITCH_SENTINEL='" +
+                          sentinel.string() + "' '" + std::string(exe) +
+                          "' -tc='*EAGER lane*' >/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 lane, control_copies, off_copies, values;
+  in >> lane >> control_copies >> off_copies >> values;
+  INFO("control(no-capture-support) copies=" << control_copies
+                                             << " switch-off copies=" << off_copies);
+  CHECK(lane == "SAME-LANE");
+  CHECK(values == "IDENTICAL");
+  std::filesystem::remove(sentinel);
+#else
+  MESSAGE("kill-switch child arm needs /proc/self/exe; not run on this platform");
+#endif
+}
diff --git a/tests/vt/recording_capture_backend.h b/tests/vt/recording_capture_backend.h
index a9d182989..c354ac149 100644
--- a/tests/vt/recording_capture_backend.h
+++ b/tests/vt/recording_capture_backend.h
@@ -97,6 +97,19 @@ class RecordingCaptureBackend final : public vt::Backend {
   vt::Queue CreateQueue() override { return vt::Queue{}; }
   bool UnifiedMemory() const override { return true; }
 
+  // THE AUX-STREAM VOCABULARY, LOGGED. ENG-CUDAGRAPH-BREAK W5 (#1335) needs the
+  // D10 auto-join to be visible IN THE SAME SEQUENCE as the segment close, for
+  // the reason W1 had to put break markers into this log rather than into a
+  // private vector: two independently asserted sequences are satisfied by an
+  // implementation that interleaves nothing. The rule under test is an ORDER —
+  // "join before `EndCaptureGraph`" — so it is only assertable if both ends are
+  // in one trace. `SupportsAuxStream()` is true here because the CUDA backend
+  // reports true (`src/vt/cuda/cuda_backend.cu:200`) and D10 exists only for
+  // backends that do.
+  bool SupportsAuxStream() const override { return true; }
+  void RecordEvent(vt::Event&, vt::Queue&) override { log_.push_back("RecordEvent"); }
+  void QueueWaitEvent(vt::Queue&, vt::Event&) override { log_.push_back("QueueWaitEvent"); }
+
   bool SupportsGraphCapture() const override { return supports_capture_; }
   void BeginCapture(vt::Queue&) override {
     log_.push_back("Begin");
diff --git a/tests/vt/test_breakable_graph.cpp b/tests/vt/test_breakable_graph.cpp
index 1baf78b85..fb8779c8f 100644
--- a/tests/vt/test_breakable_graph.cpp
+++ b/tests/vt/test_breakable_graph.cpp
@@ -1188,3 +1188,178 @@ TEST_CASE("Test 13d: a DRAINED capture records that it drained; an INERT scope d
     CHECK(g.capture_error() == nullptr);
   }
 }
+
+// ---------------------------------------------------------------------------
+// TEST 15 of `## Tests to port` — THE AUXILIARY-STREAM AUTO-JOIN (D10).
+// Port of `_end_current_segment` (`:353-361`) and of the `wait_stream` hook
+// (`:101-153`) that populates the set it walks. ENG-CUDAGRAPH-BREAK W5 (#1335).
+// ---------------------------------------------------------------------------
+//
+// WHY THIS COULD NOT BE WRITTEN BEFORE W5, and why the record says so rather
+// than implying the rule was always covered. W1 registered its break point on a
+// model that forks no auxiliary queue. W2, W3 and W4 all migrated drivers that
+// open `kFull`, which has exactly ONE segment and therefore no segment CLOSE
+// between two pieces of a forward for an outstanding fork to straddle. W5 is the
+// first stage that owns a driver whose fork is inside the captured region by
+// construction (`laguna.cpp:2572-2576` fork, `:2612` join).
+//
+// WHAT IT ASSERTS AND WHY IT IS AN ORDER. Closing a capture with an unjoined
+// fork FAILS at `cudaStreamEndCapture`, so "the join happened" is not the claim
+// — "the join happened BEFORE the close" is. Both ends are therefore asserted
+// out of ONE backend trace, for the same reason W1 had to move break markers
+// into the backend's own log: two independently asserted sequences are satisfied
+// by an implementation that interleaves nothing.
+TEST_CASE("BreakableGraph: an outstanding fork is joined BEFORE the segment closes") {
+  using vt::BreakableGraph;
+  using vt::GraphCaptureMode;
+  using vt::GraphCaptureScope;
+
+  // ARM 1, THE RULE. A fork registered inside a piecewise capture and never
+  // joined by the model is joined by the SCOPE, at the break point that closes
+  // the segment and again at the scope's own close.
+  {
+    RecordingCaptureBackend b;
+    vt::Queue q = b.CreateQueue();
+    vt::Queue aux = b.CreateQueue();
+    vt::Event done = b.CreateEvent();
+    vt::ResetGraphBreakStats();
+    BreakableGraph g;
+    {
+      GraphCaptureScope scope(b, q, g, GraphCaptureMode::kPiecewise);
+      REQUIRE(scope.active());
+      // The model forks and tells the scope, exactly as `laguna.cpp` does.
+      vt::GraphNoteFork(aux, done);
+      CHECK(scope.outstanding_forks() == 1);
+      // ... and then hits a break point WITHOUT joining. Upstream's hazard.
+      vt::GraphBreak();
+      // The scope joined it, so nothing is outstanding for the next segment.
+      CHECK(scope.outstanding_forks() == 0);
+    }
+    REQUIRE(g.captured());
+    CHECK(g.segment_count() == 2);
+    CHECK(g.break_count() == 1);
+    // THE ORDER, out of one trace. `QueueWaitEvent` sits between the fork's
+    // `RecordEvent` and the `EndCaptureGraph` it protects, on the FIRST segment.
+    CHECK(b.Trace() ==
+          "Begin RecordEvent QueueWaitEvent EndCaptureGraph Begin EndCaptureGraph");
+    const vt::GraphBreakStats s = vt::GetGraphBreakStats();
+    CHECK(s.forks_tracked == 1);
+    CHECK(s.forks_auto_joined == 1);
+  }
+
+  // ARM 2, THE CONTROL that makes arm 1 a distinction rather than a flag that is
+  // always set. The model joins its OWN fork before the break, which is what
+  // every shipped fork site actually does. The scope must then do NOTHING: no
+  // second wait, no extra event record, and `forks_auto_joined` stays 0. Without
+  // this arm an implementation that joined unconditionally on every segment
+  // close would pass arm 1 while issuing a redundant wait on every real step.
+  {
+    RecordingCaptureBackend b;
+    vt::Queue q = b.CreateQueue();
+    vt::Queue aux = b.CreateQueue();
+    vt::Event done = b.CreateEvent();
+    vt::ResetGraphBreakStats();
+    BreakableGraph g;
+    {
+      GraphCaptureScope scope(b, q, g, GraphCaptureMode::kPiecewise);
+      REQUIRE(scope.active());
+      vt::GraphNoteFork(aux, done);
+      CHECK(scope.outstanding_forks() == 1);
+      b.RecordEvent(done, aux);      // the model's own join, first half
+      b.QueueWaitEvent(q, done);     // ... and its second half
+      vt::GraphNoteJoin(aux);        // the registration is retired
+      CHECK(scope.outstanding_forks() == 0);
+      vt::GraphBreak();
+      CHECK(scope.outstanding_forks() == 0);
+    }
+    REQUIRE(g.captured());
+    // ONE RecordEvent and ONE QueueWaitEvent in the whole trace — the model's.
+    CHECK(b.Count("RecordEvent") == 1);
+    CHECK(b.Count("QueueWaitEvent") == 1);
+    CHECK(b.Trace() ==
+          "Begin RecordEvent QueueWaitEvent EndCaptureGraph Begin EndCaptureGraph");
+    const vt::GraphBreakStats s = vt::GetGraphBreakStats();
+    CHECK(s.forks_tracked == 1);
+    CHECK(s.forks_auto_joined == 0);
+  }
+
+  // ARM 3, `kFull`. The mode every migrated driver opens. A fork registered here
+  // is still tracked and is still joined before the ONE segment closes — which
+  // is the arm `laguna.cpp` actually takes, because its fork and join both sit
+  // inside `RunChain`. It is asserted rather than assumed, because "the set is
+  // empty so the rule is vacuous" is a claim about the MODEL, not about the seam,
+  // and a driver that returns early between its fork and its join would make it
+  // false without changing a line of this file.
+  {
+    RecordingCaptureBackend b;
+    vt::Queue q = b.CreateQueue();
+    vt::Queue aux = b.CreateQueue();
+    vt::Event done = b.CreateEvent();
+    vt::ResetGraphBreakStats();
+    BreakableGraph g;
+    {
+      GraphCaptureScope scope(b, q, g, GraphCaptureMode::kFull);
+      REQUIRE(scope.active());
+      vt::GraphNoteFork(aux, done);
+      CHECK(scope.outstanding_forks() == 1);
+    }
+    REQUIRE(g.captured());
+    CHECK(g.segment_count() == 1);
+    CHECK(b.Trace() == "Begin RecordEvent QueueWaitEvent EndCaptureGraph");
+    const vt::GraphBreakStats s = vt::GetGraphBreakStats();
+    CHECK(s.forks_auto_joined == 1);
+  }
+
+  // ARM 4, THE INERT SCOPE. Outside a capture, and inside a scope the backend
+  // cannot honour, both calls make ZERO backend calls and move no counter — the
+  // same pass-through guarantee `GraphBreak` gives, applied to the fork hooks.
+  // A model calls them unconditionally, so a version that tracked in the inert
+  // lane would issue joins on a forward that never captured.
+  {
+    RecordingCaptureBackend b(/*supports_capture=*/false);
+    vt::Queue q = b.CreateQueue();
+    vt::Queue aux = b.CreateQueue();
+    vt::Event done = b.CreateEvent();
+    vt::ResetGraphBreakStats();
+    vt::GraphNoteFork(aux, done);  // no scope at all
+    vt::GraphNoteJoin(aux);
+    BreakableGraph g;
+    {
+      GraphCaptureScope scope(b, q, g);
+      REQUIRE_FALSE(scope.active());
+      vt::GraphNoteFork(aux, done);
+      CHECK(scope.outstanding_forks() == 0);
+    }
+    CHECK(b.Trace().empty());
+    const vt::GraphBreakStats s = vt::GetGraphBreakStats();
+    CHECK(s.forks_tracked == 0);
+    CHECK(s.forks_auto_joined == 0);
+  }
+
+  // ARM 5, RE-REGISTERING ONE QUEUE. A model that forks the same auxiliary queue
+  // in two layers of one capture must leave ONE entry, not two: two entries make
+  // the auto-join issue two waits for one fork, and — the half that actually
+  // corrupts state — a single `GraphNoteJoin` would then retire only one of them
+  // and leave a joined queue looking outstanding forever.
+  {
+    RecordingCaptureBackend b;
+    vt::Queue q = b.CreateQueue();
+    vt::Queue aux = b.CreateQueue();
+    vt::Event done = b.CreateEvent();
+    vt::ResetGraphBreakStats();
+    BreakableGraph g;
+    {
+      GraphCaptureScope scope(b, q, g, GraphCaptureMode::kFull);
+      REQUIRE(scope.active());
+      vt::GraphNoteFork(aux, done);
+      vt::GraphNoteFork(aux, done);
+      CHECK(scope.outstanding_forks() == 1);
+      vt::GraphNoteJoin(aux);
+      CHECK(scope.outstanding_forks() == 0);
+    }
+    CHECK(b.Count("QueueWaitEvent") == 0);
+    const vt::GraphBreakStats s = vt::GetGraphBreakStats();
+    CHECK(s.forks_tracked == 1);
+    CHECK(s.forks_auto_joined == 0);
+  }
+}