feat(ENG-CUDAGRAPH-DEDUP): fold captures that share a topology onto one graph executable

.agents/claims/CLAIM-ENG-CUDAGRAPH-DEDUP.md

@@ -0,0 +1,5 @@
+# CLAIM-ENG-CUDAGRAPH-DEDUP
+
+| Claim | Row IDs | Agent | Worktree / remote dir | Branch | Owned scope | State | Last update |
+|---|---|---|---|---|---|---|---|
+| `CLAIM-ENG-CUDAGRAPH-DEDUP` | `ENG-CUDAGRAPH-DEDUP` (`ACTIVE`) | Claude Code (claude-opus-5), fresh implementer | worktree `wt-dedup`, no GPU on this host | `row/ENG-CUDAGRAPH-DEDUP`, issue [#1162](https://github.com/mudler/vllm.cpp/issues/1162), derived by the analysis in [sglang-breakable-cuda-graph.md](../specs/sglang-breakable-cuda-graph.md) | Owns ONLY: the new `src/vt/graph_dedup.h` and `src/vt/graph_dedup_runtime.h`; the graph-handle trio (`EndCaptureGraph`, `ReplayGraph`, `DestroyGraph`) in `src/vt/cuda/cuda_backend.cu` and `src/vt/rocm/rocm_backend.hip`; the new `tests/vt/test_graph_dedup.cpp` and its `tests/CMakeLists.txt` registration; the `VT_CUDA_GRAPH_DEDUP` entry in `docs/ENVIRONMENT.md`; `.agents/specs/eng-cudagraph-dedup.md`, this claim file, and the `ENG-CUDAGRAPH-DEDUP` row. **NON-COLLISION:** disjoint from `ENG-CUDAGRAPH-BREAK` ([#1163](https://github.com/mudler/vllm.cpp/issues/1163)), which a parallel agent owns and which works ABOVE this seam on the eight decode drivers and the runner routing; this claim touches no driver internals and no bucket selection. EXCLUDED: flipping `VT_CUDA_GRAPH_DEDUP` on by default (owes the device A/B), break points and piecewise capture, and the diffusion path ([#1164](https://github.com/mudler/vllm.cpp/issues/1164), blocked) | `ACTIVE` | 2026-08-18 — spec committed before implementation. W1 (the shared registry and its CPU contract suite) and W2 (the CUDA/HIP ops table and the backend wiring) land on this branch, default OFF. Still owed and named in the spec's `## Owed`: the leased-GPU same-binary byte-identity A/B, the executable-count ratio, the default flip, and the ROCm compile (no ROCm hardware or `hipcc` is reachable from this session and CI has no ROCm job) |

docs/BENCHMARKS.md

@@ -272,6 +272,7 @@ Qwen3.6-27B NVFP4, GB10, whole serving window.
 | Peak `MemAvailable` drop | 68.35 GiB | 80.66 GiB | 1.180x | **PASS** |
 | Weight offload, resident device bytes (`ENG-WEIGHT-OFFLOAD` W6) | not measured | not measured | n/a | **BLOCKED**, not pending: unmeasurable on every host we own (GB10 shares one pool, so `cpu_offload_gb` frees nothing). Needs a discrete-GPU rig ([record](../.agents/benchmark-record.md)) |
 | Disk residency via `--offload-config` (`ENG-RESIDENCY-CONFIG`, [#1110](https://github.com/mudler/vllm.cpp/issues/1110)) | not measured | n/a (no disk tier upstream) | n/a | **PENDING** a GB10 run. The row changes no kernel, dtype or allocation, so it claims no throughput axis; the 370 GiB reproduction through the JSON form is owed ([spec](../.agents/specs/weight-residency-config.md)) |
+| Decode-graph executables, `VT_CUDA_GRAPH_DEDUP` on vs off (`ENG-CUDAGRAPH-DEDUP`) | not measured | n/a | n/a | **PENDING**, not blocked: needs a leased CUDA box for the same-binary A/B (token-identity first, then the exec-count ratio) ([spec](../.agents/specs/eng-cudagraph-dedup.md)) |
 
 35B steady-serving PSS is 3.53 GiB against vLLM's 13.3 GiB after the routed-expert
 host mirror is freed once the device Marlin resident is built.

docs/ENVIRONMENT.md

@@ -132,6 +132,7 @@ portable/reference path. In normal operation leave them unset.
 | `VT_ASYNC_DEVICE_MIRROR` | on | The CUDA device-resident sampled-token mirror (ENG-ASYNC-SCHED W4): the async serving loop's sampled ids stay on the device instead of round-tripping the host. On a DISCRETE GPU this removes the host-fallback main-stream `Synchronize`. On the INTEGRATED GB10 it moves `last_sampled_tokens` off the host array `update_states`' condense reorders (a drain move, not a removal). DEFAULT ON since the 2026-08-06 correctness flip (ROW-SERVE-ASYNC-LLM): the OFF path's host-array combine writes `step.input_token_ids` on the main queue while the decode graph reads it on the CPU without a sync, so async batch-1 greedy decode nondeterministically degenerates into token-0 garbage; the mirror routes the ids into the embed on-queue and fixes it. `=0` is the rollback to the (racy) host-array path. Speed-neutral (c16 0.999x). No effect on CPU or the sync `LLMEngine` |
 | `VT_ASYNC_EXECUTOR` | off (opt-in) | `=1` enables Option A: the decode-graph per-step input H2D staged OUT of the captured replay (the c16/c32 overlap unlock, ENG-ASYNC-SCHED). The Qwen3.5 MoE/dense decode-graph drivers give each padded-size slot PERSISTENT device input buffers the captured graph reads and PINNED host staging; per step the input H2D is enqueued on the main queue BEFORE `ReplayGraph` and an input-staged event is recorded right after it, so the next same-slot Refresh waits only that tiny copy, never the GPU tail (the faithful vLLM `_prepare_input_ids`/`synchronize_input_prep` structure, `states.py:64`). The 2-slot parity ring is retained (the depth-2 loop enqueues sample(i-1) after forward(i), so persistent logits must double-buffer). The runner skips the depth-2 pre-forward `Synchronize` whenever the previous step's logits are a non-owning graph-slot view. Default OFF routes through the single-slot baked-H2D driver with the drain intact — byte-identical to production. No effect on CPU or the sync `LLMEngine`. `=1` roughly doubles the captured decode-graph memory (logits-dominated) and adds small pinned host + persistent device input buffers per slot |
 | `VLLM_CPP_CUDAGRAPH` | on (CUDA) | Eager launches instead of a captured CUDA graph |
+| `VT_CUDA_GRAPH_DEDUP` | off (opt-in) | `=1` folds captured graphs that share a node topology onto ONE graph executable, re-pointing it with `cudaGraphExecUpdate` / `hipGraphExecUpdate` instead of instantiating a second (`ENG-CUDAGRAPH-DEDUP`, [#1162](https://github.com/mudler/vllm.cpp/issues/1162)). Today a model holds one executable per padded decode bucket — 7 at `max_num_seqs=32`, 11 at 64 — times eight capture drivers. A MEMORY and capture-time change, not a throughput one: a deduped replay launches the same nodes, and each candidate fold is probed with the real driver update on a throwaway executable before it is honoured, so a capture the driver refuses simply keeps its own executable. It logs `vt graph dedup: captured N graphs, deduped to G execs` per capture so the ratio is readable. **Default OFF and it stays off until measured:** a workload that alternates padded buckets every step pays one `cudaGraphExecUpdate` per switch, and the device byte-identity A/B that would price that is still owed. Unset leaves the capture path byte-identical to the pre-dedup one |
 | `VLLM_CPP_DENSE_DECODE_GRAPH` | on (CUDA dense) | Non-graphed dense decode |
 | `VLLM_CPP_QWEN3_DENSE_DECODE_GRAPH` | **on** | Routes pure-decode steps for the SHARED pure-dense forward (`Qwen3DenseModel`, i.e. Qwen3 / Llama / InternLM3 / Mistral / InternLM2 `ForCausalLM`) through the captured decode CUDA graph. **DEFAULT-ON since `QUANT-CT-MXFP4-MARLIN-STRUCT`** (parity-enabler; banks ~+1.3% TPOT@c8 by collapsing the eager inter-kernel launch gap). Token-exact with eager — `=0` opts out to the byte-identical eager decode; gated dgx SACRED `test_qwen3_paged_engine` 184/184 (graph ON==OFF, Qwen3-0.6B near-tie + 4B), async 82/82, Qwen3-8B-MXFP4 #44 smoke 3/3, Qwen3-32B-NVFP4A16 142/142. Honors `VLLM_CPP_CUDAGRAPH=0` |
 | `VT_MOE_FUSED_W13` | on | Runs a dense/shared MLP gate+up PAIR as ONE Marlin GEMM over the N-concatenated `[2I,H]` operand + `SiluAndMul` (vLLM's merged `gate_up_proj` structure), instead of two split GEMMs + `MoeSiluMul`. **Generalized NVFP4→MXFP4 in `QUANT-CT-MXFP4-MARLIN-STRUCT`** (drops the classic-dense Qwen3-8B-MXFP4 decode from 180→144 marlin GEMM/step = vLLM-structural parity). `=0` is the split A/B fallback. Numerically equivalent to split (the fused fp32 split-K reduce regroups by 1 bf16 ULP; token-exact vs the oracle — #44 fused==split 3/3, unit `test_linear_method` 99.9% bit-exact) |

docs/STATUS.md

@@ -1705,6 +1705,8 @@ Gemma4/ROCm env split: public `VT_GEMMA4_EXPERT_VRAM_MB` caps expert LRU in posi
 
 `BACKEND-TENSTORRENT-HOST-FREE-FORWARD`: `ACTIVE`: env-gated `VT_TT_HOST_FREE_DECODE` decode-graph capture. Implementer P150 run of Qwen3-0.6B, 80 tokens: 79 replays, no hang, 5.8x vs eager, 22/22 vs the per-step-copy baseline. Default path inert. Operator gate and full-engine golden still owed. A new batch after the first capture is refused.
 
+`ENG-CUDAGRAPH-DEDUP`: `ACTIVE`: env-gated `VT_CUDA_GRAPH_DEDUP` graph-executable dedup — one `cudaGraphExec` per captured TOPOLOGY instead of one per padded decode bucket per model. The shared `vt` registry hashes a captured graph's structure and re-points a single executable with `cudaGraphExecUpdate` on a hit, falling back to a private executable when the driver rejects the update. A memory and capture-time change, NOT a throughput change. Default OFF until a leased-GPU same-binary A/B proves a deduped replay token-identical; that A/B and the ROCm compile are owed ([#1162](https://github.com/mudler/vllm.cpp/issues/1162)).
+
 **Platform SELECTION is the one non-additive site, and is now gated.** A
 platform missing from `CurrentPlatform()`'s hardcoded walk registers and answers
 correctly but is NEVER selected, with no compiler diagnostic. `test_platform`

scripts/check-gate-commands.py

@@ -311,8 +311,18 @@ def audit() -> list[dict]:
 # reproduction through the JSON form is recorded as OWED rather than skipped
 # because dgx.casa was unreachable at the SSH layer. Growth, so the set is
 # re-pinned in the same change.
+# 2026-08-18: +ENG-CUDAGRAPH-DEDUP. A NEW row arriving at ACTIVE (issue #1162),
+# so it enters GATED_STATES for the first time. Its spec's Gates section names
+# `ctest -R test_graph_dedup` and `scripts/agent-preflight.sh`, both of which
+# genuinely fail when the row regresses -- the focused suite detected 9 of 9
+# negative mutations of the registry it gates. It also records what is NOT
+# claimed and why: the device byte-identity A/B needs a leased CUDA box this
+# session did not have, so it is carried under the spec's `## Owed` rather than
+# reported as run, and the CUDA leg's compile rests on the `cuda-fat-build` CI
+# job. Growth, so the set is re-pinned in the same change.
 RUNNABLE_BASELINE = frozenset({
     "ENG-RESIDENCY-CONFIG",
+    "ENG-CUDAGRAPH-DEDUP",
     "SPEC-MTP-K-GT-1",
     "ATTN-CHUNKED-LOCAL",
     "SERVE-RECIPE-ARGS",

src/vt/cuda/cuda_backend.cu

@@ -13,6 +13,8 @@
 
 #include "vt/backend.h"
 #include "vt/cuda/cuda_device_caps.h"
+#include "vt/graph_dedup.h"
+#include "vt/graph_dedup_runtime.h"
 #ifdef VT_BENCH_PROFILE_CONTROL
 #include "vt/cuda/cuda_profiler_control.h"
 #endif
@@ -225,6 +227,18 @@ class CudaBackend final : public Backend {
   void* EndCaptureGraph(Queue& q) override {
     cudaGraph_t graph = nullptr;
     Check(cudaStreamEndCapture(AsStream(q), &graph), "cudaStreamEndCapture");
+    // ENG-CUDAGRAPH-DEDUP (#1162): with VT_CUDA_GRAPH_DEDUP set, hand the RAW graph to
+    // the dedup registry, which keys it by topology and folds it onto an existing
+    // executable when the driver accepts the update. The raw graph is retained for the
+    // handle's lifetime because both the key and cudaGraphExecUpdate need it — SGLang
+    // buys the same retention with torch.cuda.CUDAGraph(keep_graph=True). Default OFF,
+    // in which case this stays the pre-dedup path byte for byte.
+    if (vt::GraphDedupEnabled()) {
+      if (dedup_ == nullptr) {
+        dedup_ = std::make_unique<vt::GraphDedupRegistry>(vt::graph_dedup_rt::Ops());
+      }
+      return dedup_->Register(reinterpret_cast<void*>(graph));
+    }
     cudaGraphExec_t exec = nullptr;
     Check(cudaGraphInstantiate(&exec, graph, 0), "cudaGraphInstantiate");
     cudaGraphDestroy(graph);
@@ -264,8 +278,14 @@ class CudaBackend final : public Backend {
       }
     }
 #endif
-    Check(cudaGraphLaunch(reinterpret_cast<cudaGraphExec_t>(graph), AsStream(q)),
-          "cudaGraphLaunch");
+    if (dedup_ != nullptr && dedup_->Owns(graph)) {
+      // Owns() rather than a mode flag: a handle minted before the registry existed is
+      // still a plain executable, and the seam must never guess at a pointer's origin.
+      dedup_->Replay(graph, reinterpret_cast<void*>(AsStream(q)));
+    } else {
+      Check(cudaGraphLaunch(reinterpret_cast<cudaGraphExec_t>(graph), AsStream(q)),
+            "cudaGraphLaunch");
+    }
 #ifdef VT_BENCH_PROFILE_CONTROL
     if (g_cuda_profile_active) {
       if (g_cuda_profile_remaining_replays == 0) {
@@ -286,7 +306,12 @@ class CudaBackend final : public Backend {
 #endif
   }
   void DestroyGraph(void* graph) override {
-    if (graph != nullptr) cudaGraphExecDestroy(reinterpret_cast<cudaGraphExec_t>(graph));
+    if (graph == nullptr) return;
+    if (dedup_ != nullptr && dedup_->Owns(graph)) {
+      dedup_->Destroy(graph);
+      return;
+    }
+    cudaGraphExecDestroy(reinterpret_cast<cudaGraphExec_t>(graph));
   }
 
  private:
@@ -295,6 +320,9 @@ class CudaBackend final : public Backend {
   int sm_major_ = 0;
   int sm_minor_ = 0;
   cudaGraphExec_t exec_ = nullptr;  // last instantiated captured graph
+  // ENG-CUDAGRAPH-DEDUP (#1162): built on the first capture, and only when
+  // VT_CUDA_GRAPH_DEDUP asked for it, so an unset environment allocates nothing.
+  std::unique_ptr<vt::GraphDedupRegistry> dedup_;
 };
 
 // Registers kCUDA during static init (registration must complete before