feat(ENG-CUDAGRAPH-BREAK): migrate the three single-shape drivers onto the seam, and join an outstanding aux-stream fork before a segment closes (#1335, #1163)

.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 |

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 |

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.

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

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`.