From 6639f410db9596b8c5db2a3ccff4602197cdbe61 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Tue, 18 Aug 2026 00:43:44 +0000 Subject: [PATCH 1/9] spec(ENG-CUDAGRAPH-DEDUP): one graph executable per topology, not one per bucket MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit We instantiate a fresh `cudaGraphExec` for every capture and throw the raw graph away (`src/vt/cuda/cuda_backend.cu:222-232`), so the executables multiply along two axes: 7 padded decode buckets at `max_num_seqs=32` or 11 at 64, times eight hand-rolled drivers. `grep -rn "cudaGraphExecUpdate" src include` returns nothing. This spec lands before its implementation, as the protocol requires. It ports the construction from SGLang's dedup mixin at the pinned `f63458b5be`, because vLLM gets its executables from `torch.compile` and therefore has no analogue to mirror, and it records the three places where our seam forces a divergence. The load-bearing one: SGLang holds a second probe executable per group for the whole capture phase and frees it in `seal()`, but our drivers capture lazily and interleaved with replay, so there is no capture-phase end and a persistent probe would hold two executables per group forever — worse than the baseline this row exists to reduce. The probe is therefore transient, and a probe the driver rejects falls back to a private executable instead of aborting, which is what makes an under-specified signature safe rather than wrong. The row moves `INVENTORIED` to `ACTIVE`, which owes `docs/STATUS.md`, `docs/BENCHMARKS.md` and the spec's own `## Now`. The benchmark entry is recorded `PENDING` rather than left absent: the device byte-identity A/B needs a leased CUDA box, and this session has none. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code] --- .agents/claims/CLAIM-ENG-CUDAGRAPH-DEDUP.md | 5 + .agents/engine-matrix.md | 6 +- .agents/specs/eng-cudagraph-dedup.md | 254 ++++++++++++++++++++ docs/BENCHMARKS.md | 1 + docs/STATUS.md | 2 + 5 files changed, 265 insertions(+), 3 deletions(-) create mode 100644 .agents/claims/CLAIM-ENG-CUDAGRAPH-DEDUP.md create mode 100644 .agents/specs/eng-cudagraph-dedup.md diff --git a/.agents/claims/CLAIM-ENG-CUDAGRAPH-DEDUP.md b/.agents/claims/CLAIM-ENG-CUDAGRAPH-DEDUP.md new file mode 100644 index 000000000..cdcdcaf8e --- /dev/null +++ b/.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) | diff --git a/.agents/engine-matrix.md b/.agents/engine-matrix.md index 0b049156d..3e3e9fcb9 100644 --- a/.agents/engine-matrix.md +++ b/.agents/engine-matrix.md @@ -38,7 +38,7 @@ forensics: roadmap_v1.md and the parity ledger. | Area | Rows | `ANCHOR-BACKFILL` | `PARTIAL` | `SPIKE` | `READY` | `ACTIVE` | `GATING` | `DONE` | `INVENTORIED` | |---|---:|---:|---:|---:|---:|---:|---:|---:|---:| -| Engine and scheduling | 30 | 6 | 3 | 1 | 3 | 6 | 2 | 2 | 7 | +| Engine and scheduling | 30 | 6 | 3 | 1 | 3 | 7 | 2 | 2 | 6 | | KV cache and memory | 23 | 7 | 3 | 2 | 3 | 2 | 2 | 0 | 4 | | Parallelism | 6 | 0 | 0 | 0 | 1 | 0 | 0 | 0 | 5 | | Sampling and generation | 15 | 4 | 2 | 0 | 0 | 4 | 0 | 1 | 4 | @@ -48,7 +48,7 @@ forensics: roadmap_v1.md and the parity ledger. | LoRA and adapters | 2 | 0 | 0 | 0 | 0 | 1 | 0 | 0 | 1 | | Long context and attention | 10 | 5 | 0 | 0 | 1 | 0 | 1 | 0 | 3 | | Loading, tokenizer, config | 10 | 3 | 3 | 0 | 0 | 1 | 1 | 1 | 1 | -| **Total** | **160** | **35** | **17** | **5** | **11** | **32** | **8** | **11** | **40** | +| **Total** | **160** | **35** | **17** | **5** | **11** | **33** | **8** | **11** | **39** | ## Engine core and scheduling @@ -60,7 +60,7 @@ forensics: roadmap_v1.md and the parity ledger. | `KV-PREFIX-MATCH-UNIT` | `--prefix-match-unit` (config `prefix_match_unit`): the finest token boundary a prefix-cache hit can land on == the `hash_block_size`/"prefix match unit" the block hasher uses. NEW in 0.26 (absent at the prior `e24d1b24`/0.25.0 pin). For a HYBRID/multi-group model the resolver `resolve_kv_cache_block_sizes` computes `hash_block_size = prefix_match_unit if set else gcd(group_block_sizes)` (scheduler block size = `lcm`), letting matching land FINER than a physical block (e.g. 16/32 tokens inside a 1024-token block) provided every group block size is divisible by it; single-group (dense) models ignore the knob. Backs off to the scheduler block size when no prefix-cache/connector consumer is active or a mamba group diverges from `cache_block_size` (mamba_cache_mode != "align"); throws on a non-divisible unit. **W0 spike + W1 resolver LANDED 2026-07-28 (`CLAIM-PREFIX-MATCH-UNIT`, NOT pushed):** `resolve_kv_cache_block_sizes` ported 1:1 (explicit-parameter signature vs upstream's `VllmConfig`, our config surface is threaded), RED-first unit-gated (default gcd `!=` `=16` override). `PARTIAL`: the config/CLI/ABI field (W2), the scheduler threading of a resolved `hash_block_size != block_size` + mamba partial-tail stop (W3, needs the `KV-BLOCK-POOL` align path that still throws), and the benchmark (W4) are deferred. Default path byte-identical (single-group inert; scheduler still passes `block_size`). | T1 | `vllm/engine/arg_utils.py:696,1222,1940`; `vllm/config/cache.py:56-67`; resolver `vllm/v1/core/kv_cache_utils.py:626-688`; hasher `:691-748`; call site `vllm/v1/engine/core.py:154`; scheduler `vllm/v1/core/sched/scheduler.py:76,268-270,282,312-318`; fine-grained view `vllm/v1/core/single_type_kv_cache_manager.py:683,697` | resolver `src/vllm/v1/core/kv_cache_utils.cpp:638` (`resolve_kv_cache_block_sizes`), decl `include/vllm/v1/core/kv_cache_utils.h`; hash_block_size already plumbed `get_request_block_hasher` `src/vllm/v1/core/kv_cache_utils.cpp:577`; DEFERRED align path throws `src/vllm/v1/core/block_pool.cpp:93,220` (shared with `KV-BLOCK-POOL`) | `tests/vllm/v1/test_prefix_match_unit.cpp:64,88,99,119,129,145,164,186` 8/8 (29 assertions): single-group inert + DCP scale, multi-group default=gcd, `=16` override finer-than-default (RED), finer-than-1024-block, non-divisible throws, no-consumer back-off + connector-alone re-enable, mamba non-align back-off vs align gcd, hasher-granularity RED (coarse 2 vs fine 4 hashes); [parity-ledger.md](parity-ledger.md) | [prefix-match-unit.md](specs/prefix-match-unit.md) | `PARTIAL` | `CLAIM-PREFIX-MATCH-UNIT` | | `ENG-PREEMPT-RECOMPUTE` | FCFS tail preemption with recompute | T0 | `vllm/v1/core/sched/scheduler.py:1142`; `tests/v1/core/test_scheduler.py:930` | `src/vllm/v1/core/sched/scheduler.cpp:102,157`; `src/vllm/v1/core/sched/request_queue.cpp:36` | `tests/vllm/v1/test_scheduler.cpp:247,295`; `tests/vllm/v1/test_request_queue.cpp:91` | `planned: specs/preemption.md` | `ANCHOR-BACKFILL` | - | | `ENG-CUDAGRAPH` | Decode graph capture/replay modes (host-cluster cleanup: capture-size set derived from `max_num_seqs` mirroring vLLM `_set_cudagraph_sizes`; 2026-07-18 graph-baked-scratch use-after-free fix — the 35B c2+ online-serving IMA blocker) | T0 | `vllm/config/compilation.py:53,1319,683-684,1438-1444`; `vllm/config/vllm.py:1667-1770`; `vllm/v1/worker/gpu/cudagraph_utils.py:116`; `tests/compile/test_config.py:122,229` | `src/vt/cuda/cuda_backend.cu:76,97,105`; `include/vllm/model_executor/models/decode_graph_sizes.h`; `src/vllm/model_executor/models/qwen3_5.cpp:3754,3952`; `src/vllm/v1/worker/gpu/runner.cpp:577,597`; graph-safe scratch (retire-on-grow so graph-baked scratch pointers stay valid) `src/vt/cuda/graph_safe_scratch.h`, `src/vt/cuda/cuda_moe_marlin.cu:75`, `src/vt/cuda/cuda_matmul_nvfp4.cu:766`, `src/vt/cuda/cuda_matmul_nvfp4_cutlass.cu:105`, `src/vt/cuda/cuda_matmul_fp8_cutlass.cu:95` | `tests/vt/test_cuda_backend.cpp:98`; `tests/vllm/models/test_decode_graph_sizes.cpp`; `tests/vt/test_graph_safe_scratch.cpp`; explicit 35B gate `tests/parity/test_qwen36_paged_engine.cpp:140` | [blocktable-host-cluster-cleanup.md](specs/blocktable-host-cluster-cleanup.md); [decode-graph-scratch-uaf-2026-07-18.md](specs/decode-graph-scratch-uaf-2026-07-18.md) | `PARTIAL` | **PREFILL capture REFUTED as a lever (2026-08-17, [#1161](https://github.com/mudler/vllm.cpp/issues/1161)).** vLLM's v1 default already captures prefill piecewise (`vllm/config/compilation.py:60-63,615,630` @ `555967922`) and it is in our denominator; SGLang reached the same coverage without `torch.compile` via BCG (`SGLANG-BCG` in [sglang-matrix.md](sglang-matrix.md)). Neither helps us: GB10 2026-07-09 measured prefill GPU-idle-between-launches at **3.8%** with GPU-busy >96% on both arms, and the 27B prefill gap at **92.5% non-GEMM glue GPU work** with the dominant GEMM at +0.17% and attention AHEAD. There are no launch bubbles in our prefill to collapse. Row stays `PARTIAL`; the real residuals are exec dedup ([#1162](https://github.com/mudler/vllm.cpp/issues/1162)) and the break-point seam ([#1163](https://github.com/mudler/vllm.cpp/issues/1163)). Spec [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | -| `ENG-CUDAGRAPH-DEDUP` | Graph-executable dedup: hash each captured graph's topology and re-point ONE `cudaGraphExec` with `cudaGraphExecUpdate` on a signature hit, instead of instantiating one exec per padded bucket per model. A memory and capture-time change, NOT a throughput change — a deduped replay launches the same nodes | T2 | vLLM has no analogue (its execs come from `torch.compile`); SGLang `python/sglang/srt/model_executor/runner_backend/cuda_graph_dedup_mixin.py:219-242,358` @ `f63458b5be` | UNPORTED — `cudaGraphExecUpdate` appears nowhere in `src/` or `include/`; today `src/vt/cuda/cuda_backend.cu:222-232` instantiates a fresh exec per capture, over the 7 (`max_num_seqs=32`) or 11 (64) buckets of `include/vllm/model_executor/models/decode_graph_sizes.h:32-41`, times eight drivers | owed: same-binary A/B proving a deduped replay is byte-identical to a non-deduped one; a capture-count/exec-count log line so the ratio is observable | [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `INVENTORIED` | [#1162](https://github.com/mudler/vllm.cpp/issues/1162) | +| `ENG-CUDAGRAPH-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 eight drivers | CPU-tier contract suite over a fake ops table, RED-first, so every guarantee is gated on every platform and not only on a CUDA box; the fake's launch log makes "the right nodes ran" an observable sequence over MORE than one replay per shape. CUDA leg compile-gated by the `cuda-fat-build` CI job. STILL OWED: the device same-binary A/B proving a deduped replay is token-identical, the exec-count ratio, and the ROCm compile | [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 eight hand-rolled drivers become one. **Coverage 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 `breakable_cuda_graph.py:204-241,244-260` @ `f63458b5be` (decorator + runtime stream capture, no compiler) | 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` | owed: reachability mutation (delete the production call site, rerun the focused gate); bit-exactness vs eager on every migrated model over MORE than one replay; the host-lifetime contract of [decode-graph-scratch-uaf-2026-07-18.md](specs/decode-graph-scratch-uaf-2026-07-18.md) enforced AT the seam | [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `INVENTORIED` | [#1163](https://github.com/mudler/vllm.cpp/issues/1163), [#1020](https://github.com/mudler/vllm.cpp/issues/1020) | | `ENG-CUDAGRAPH-DIFFUSION` | Capture the LTX-2.5 denoise loop (fixed shapes, many identical iterations — the ideal graph target). **BLOCKED, and the blocker is ours:** the render does almost no device compute to capture | T2 | SGLang enabled BCG on this shape AFTER our pin — LTX-2 H200 two-stage 10.75s->6.90s (`d4be483efb`), SANA 1024px -26% (`6c7498113f`), SANA denoise 0.73->0.457s (`56ef810cad`). Dated events, NOT pinned evidence; their win is mostly PyTorch host tax we do not pay | NO capture at all: `grep` for capture across `src/vllm/model_executor/models/ltx2*.cpp` returns nothing | blocked by [#1024](https://github.com/mudler/vllm.cpp/issues/1024) (GPU util **exactly 0 in 321 of 347 samples**, 1.00 core of 20 held for 17+ min after staging), [#1007](https://github.com/mudler/vllm.cpp/issues/1007) (VAE decode has no device arm), [#1087](https://github.com/mudler/vllm.cpp/issues/1087) (**57-66% of wall** is ONE resolution-CONSTANT serial host phase), [#1010](https://github.com/mudler/vllm.cpp/issues/1010) (no phase-boundary log). Decision point is a MEASUREMENT of GPU-busy vs wall once device-resident, not an implementation | [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `INVENTORIED` | [#1164](https://github.com/mudler/vllm.cpp/issues/1164) | | `ENG-BATCH-INVARIANT` | Opt-in deterministic execution across scheduler batch sizes (`VLLM_BATCH_INVARIANT=1`): batch-invariant matmul/norm/attention/collectives plus persistent-scheduler NVFP4; production default remains off | T1 | default/env `vllm/envs.py:89,576-578`; initialization `vllm/v1/worker/gpu_worker.py:1262`; NVFP4 dispatch `csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_sm120_kernels.cu:212-220`; suite fixture `tests/v1/determinism/conftest.py:9-12`; operator/e2e `tests/v1/determinism/test_nvfp4_batch_invariant_scaled_mm.py`, `tests/v1/determinism/test_nvfp4_batch_invariant.py` @ `702f481` | - | [W3-C3R executed contract](specs/nvfp4-persistent-plan-cache.md#w3-c3r-batch-shape-localization-and-gate-correction-2026-07-13): production-default ours and vLLM both change outputs across batch shapes; no local opt-in implementation is claimed | `planned: specs/batch-invariant-execution.md` | `INVENTORIED` | - | diff --git a/.agents/specs/eng-cudagraph-dedup.md b/.agents/specs/eng-cudagraph-dedup.md new file mode 100644 index 000000000..21d486f89 --- /dev/null +++ b/.agents/specs/eng-cudagraph-dedup.md @@ -0,0 +1,254 @@ +# ENG-CUDAGRAPH-DEDUP: one graph executable per topology, not one per bucket + +**Row:** `ENG-CUDAGRAPH-DEDUP` in [engine-matrix.md](../engine-matrix.md). +**Issue:** [#1162](https://github.com/mudler/vllm.cpp/issues/1162). +**Analysis that derived this row:** [sglang-breakable-cuda-graph.md](sglang-breakable-cuda-graph.md), `## Owed`. +**Kind:** memory and capture-time optimization behind the `vt` graph seam. **Not a +throughput row.** A deduped replay launches exactly the same nodes with exactly the +same parameters; if it launched anything else it would be a correctness defect, which +is why the load-bearing gate below is byte-identity rather than a speed ratio. + +## Now + +`ACTIVE`. The shared registry, its CPU contract suite and the CUDA wiring land here. +The device-level byte-identity A/B and the ROCm leg are named under +[`## Owed`](#owed) with the issue that owns each. + +## Scope + +In scope: + +1. A device-agnostic graph-executable dedup registry behind the existing `vt` + `Backend::{EndCaptureGraph,ReplayGraph,DestroyGraph}` seam, keyed on a structural + signature of the captured graph, folding compatible captures onto one executable + with `cudaGraphExecUpdate` / `hipGraphExecUpdate`. +2. Retention of the raw captured graph, which the current seam destroys immediately + after instantiate. Both the signature and the update need it. +3. The CUDA ops table and the CUDA backend wiring, off by default behind + `VT_CUDA_GRAPH_DEDUP`. +4. An observable capture-count / exec-count log line and a programmatic `Stats()` + accessor, so the ratio is a measurement rather than a claim. + +Out of scope, each owned elsewhere: + +- Break points and piecewise capture: `ENG-CUDAGRAPH-BREAK`, + [#1163](https://github.com/mudler/vllm.cpp/issues/1163). +- The diffusion / denoise-loop capture: `ENG-CUDAGRAPH-DIFFUSION`, + [#1164](https://github.com/mudler/vllm.cpp/issues/1164), blocked. +- Retiring the eight hand-rolled decode drivers. This row does not touch their + internals; it changes only what the seam they already call hands back. +- Any change to bucket selection, padding, or the `pure_decode` predicate. + +## Upstream chain + +vLLM has **no analogue**. Its graph executables come from `torch.compile` +(`vllm/config/compilation.py:60-63,517,615,630` @ `555967922`), which owns +instantiation, so there is no vLLM file to mirror. Per AGENTS.md `## When vLLM has no +implementation`, the secondary oracle is SGLang, pinned in +[oracles/sglang.md](../oracles/sglang.md) at `v0.5.15` / `f63458b5beaceabbd9d749b9fc956370e1b649e6`. +Paths below are relative to `python/sglang/srt/`. + +| Mechanism | Upstream anchor | Ported here | +|---|---|---| +| `cudaGraphExecUpdate` wrapper returning (ok, detail) | `model_executor/runner_backend/cuda_graph_dedup_mixin.py:27-37` | `GraphDedupOps::update` | +| Structural signature over node payloads in a deterministic topological order plus the edge set | `:139-179` | `GraphSignature()` in `src/vt/graph_dedup_runtime.h` | +| Per-node payload by node type (kernel / memcpy / memset / child graph / empty) | `:105-136` | same function, same five cases | +| Group holding one live exec plus the raw graph it currently reflects | `:182-188` (`GraphExecGroup`) | `GraphDedupRegistry::Group` | +| Register: signature hit re-points, signature miss instantiates a new group | `:219-242` | `GraphDedupRegistry::Register` | +| Replay: update only when the group's current raw graph differs, then launch | `:258-275` | `GraphDedupRegistry::Replay` | +| Compatibility probe on a separate executable so a failed probe cannot leave the live one in an undefined state | `:236-238` (`compat_exec`), destroyed by `seal()` at `:244-251` | `GraphDedupRegistry::Register`, **transient** probe — see the divergence below | +| `"captured %d CUDA graphs, deduped to %d execs"` | `:353-358` | `GraphDedupRegistry::Stats()` plus the log line | +| Retained raw graph (`torch.cuda.CUDAGraph(keep_graph=True)`) | `:326` | `EndCaptureGraph` stops calling `cudaGraphDestroy` when dedup is on | + +Three deliberate divergences from the upstream construction, each forced by a +difference in our seam rather than chosen: + +- **The probe executable is transient, not persistent.** SGLang holds a second + `compat_exec` per group for the whole capture phase and frees it in `seal()`, + which its backend calls at `end_cuda_graph_capture()` (`:352-358`). **Our seam has + no capture-phase end.** The eight drivers capture lazily, the first time a padded + bucket is seen, interleaved with replays of buckets already captured + (`src/vllm/model_executor/models/qwen3_5.cpp:9918`, + `src/vllm/model_executor/models/qwen3.cpp:888`). A persistent probe would therefore + never be sealed, would live for the process, and would hold **two** executables per + group — which is worse than the one-per-bucket baseline this row exists to reduce. + So the probe is instantiated from the group's first raw graph, updated to the + candidate, and destroyed inside `Register`. Peak executables during a registration + is 2, the same as today's steady state per capture, and steady state afterwards is + 1 per signature group. +- **A failed probe falls back; it does not abort.** SGLang asserts (`:229`). Here a + signature hit whose update the driver rejects gives that graph its own group, so the + worst case degrades exactly to today's behaviour instead of killing the process. This + is what makes an under-specified signature safe rather than wrong: the signature is + only a candidate filter, and `cudaGraphExecUpdate` is the authority. +- **Kernel identity is the host function pointer, not the kernel name.** SGLang reads + names through the CUDA driver API (`cuKernelGetName` / `cuFuncGetName`, `:44-55`) + because it holds driver handles. The runtime API hands us + `cudaKernelNodeParams::func` directly, and pointer identity is strictly stronger than + name identity for the same reason two `__global__` templates can share a demangled + prefix. SGLang's launch-attribute probing (`:58-102`) is driver-API-only and is + dropped for the same reason; dropping it can only produce a false signature *hit*, + which the probe then rejects, and never a false miss. + +## Our baseline + +`grep -rn "cudaGraphExecUpdate" src include` returns nothing at the base of this +branch (`9bc4d7f44`). Today every capture instantiates its own executable and throws +the raw graph away: + +- `src/vt/cuda/cuda_backend.cu:222-232` — `EndCaptureGraph` does + `cudaStreamEndCapture` → `cudaGraphInstantiate` → `cudaGraphDestroy(graph)` and + returns the exec as the opaque `void*` handle. +- `src/vt/rocm/rocm_backend.hip:271-283` — the same shape on `hipGraph`. +- `src/vt/backend.cpp:36-38` — the default implementations `VT_CHECK(false)`; + Vulkan, Metal and CPU do not support capture. Tenstorrent implements the trio over + `MeshTraceId` (`src/vt/tenstorrent/tenstorrent_backend.cpp:79-81`) and is not in + scope. + +The executables multiply along two axes. `include/vllm/model_executor/models/decode_graph_sizes.h:32-41` +yields 7 padded decode buckets at `max_num_seqs=32` and 11 at 64, and eight drivers +each build their own set: `Qwen3_5DecodeGraph`, `Qwen3_5DenseDecodeGraph`, +`Qwen3MoeDecodeGraph`, `Qwen3DenseDecodeGraph`, `DeepseekV2DecodeGraph`, +`VoxtralDecodeGraph`, plus the graph code in +`src/vllm/model_executor/models/deepseek_v4.cpp:1900` and +`src/vllm/model_executor/models/laguna.cpp:2691`. + +Why it matters here specifically, restated from the issue so this spec stands alone: +on GB10 unified memory an out-of-memory event reboots the box, capture time is startup +latency and startup latency is a recorded gate axis, and bucket count is exactly the +quantity that `ENG-CUDAGRAPH-BREAK` would raise if it widened coverage. + +The standing hazard on this path is recorded in +[decode-graph-scratch-uaf-2026-07-18.md](decode-graph-scratch-uaf-2026-07-18.md): +capture bakes the pointers live at capture time, so a later free of a captured buffer +is a use-after-free on replay, and a clean `compute-sanitizer` run is **not** evidence +that a capture path is safe, because its serialization hides exactly this bug. Dedup +interacts with that hazard in one direction only and it is the safe one: retaining the +raw graph extends a host-side object's lifetime and frees nothing earlier than before. +Nothing in this change frees a device buffer, and the executables it does not +instantiate were never a source of baked pointers. + +## Port map + +| New / changed | What | +|---|---| +| `src/vt/graph_dedup.h` (new) | `GraphDedupOps` (six function pointers: signature, instantiate, update, destroy exec, destroy graph, launch) and `GraphDedupRegistry`. Header-only, no device dependency, CPU-unit-testable — the same shape `src/vt/cuda/graph_safe_scratch.h` used for the graph-safe scratch bookkeeping. Also `GraphDedupEnabled()`, the `VT_CUDA_GRAPH_DEDUP` read. | +| `src/vt/graph_dedup_runtime.h` (new) | The CUDA / HIP ops table, written **once** and bound to either runtime by a macro alias block. `GraphSignature()` lives here. This is the file that must not become two hand-written copies. | +| `src/vt/cuda/cuda_backend.cu` | `EndCaptureGraph` routes through the registry when enabled and retains the raw graph; `ReplayGraph` and `DestroyGraph` dispatch on `registry->Owns(handle)`. | +| `src/vt/rocm/rocm_backend.hip` | The same three edits against the same shared header. | +| `docs/ENVIRONMENT.md` | `VT_CUDA_GRAPH_DEDUP`. Required: `scripts/check-env-doc.py` fails a `VT_*` name that is neither documented nor allowlisted, and this is a user-facing behaviour knob rather than kernel-internal tuning. | +| `.agents/engine-matrix.md` | The row moves `INVENTORIED` → `ACTIVE`; the summary table's "Engine and scheduling" line moves one row between the two state columns. No row is added or removed, so `ENGINE_ROWS` stays 160. | + +The `void*` handle contract does not change. With dedup off, `EndCaptureGraph` returns +the raw executable exactly as today and the path is byte-identical. With dedup on it +returns a registry-owned handle, and `Owns()` — an O(1) set membership test on handles +the registry itself minted — is what lets one backend serve both kinds without +guessing at a pointer's provenance. + +## Tests to port + +SGLang ships no unit test for the mixin, so there is nothing to port parameter-for- +parameter; what is portable is the guarantee set, and each case below names the +upstream line it pins. New suite `tests/vt/test_graph_dedup.cpp`, driven by a fake ops +table over a toy graph model so every guarantee is gated on **every** platform rather +than only on a CUDA box. The fake's `launch` appends the id of the graph the executable +currently reflects, which makes "the right nodes ran" an observable sequence instead of +an assertion about intent. + +| Case | Guarantee | Upstream | +|---|---|---| +| `folds compatible captures onto one executable` | K same-signature captures leave exactly 1 live exec; `Stats()` is `(K, 1)` | `:219-242`, `:253-256` | +| `keeps incompatible captures apart` | different signatures never share | `:225-241` | +| `replays the graph the caller asked for` | over an interleaved multi-replay sequence, the deduped launch log equals the non-deduped control log, element for element | `:258-275` | +| `does not re-update for a repeated shape` | a second consecutive replay of one handle issues no update | `:267` | +| `falls back to a private executable when the driver rejects the update` | a probe failure yields its own exec, the existing group's live exec is untouched, and replay stays correct | diverges from `:229`, stated above | +| `frees the shared executable only with its last graph` | destroying one sibling keeps the exec; destroying the last frees the exec and every retained raw graph, with all counters balanced | `:277-292` | +| `does not claim a handle it did not mint` | `Owns()` is false for a foreign pointer | ours; the seam has no upstream twin | +| `is off unless the environment asks for it` | `GraphDedupEnabled()` polarity | `:321` | + +## Gates + +1. **Red first.** Every case above is written and run against the un-implemented + registry, and the red output is captured, before the implementation exists. +2. **Focused green.** `ctest -R test_graph_dedup`. +3. **Mutation.** Each guarantee is deleted or inverted in a scratch copy, the focused + suite is proven to fail, and the tree is restored byte-for-byte. `git diff --stat` + and the compile status are printed for every mutation, because a mutation that fails + to build and a mutation that never applied both read as a passing test. +4. **Full gate.** `./scripts/agent-preflight.sh`. +5. **CUDA compile.** The `cuda-fat-build` CI job (`.github/workflows/ci.yml:669-712`, + container `nvidia/cuda:13.3.0-devel-ubuntu24.04`) compiles the CUDA leg on the PR. +6. **Device byte-identity A/B (owed, see below).** Same binary, `VT_CUDA_GRAPH_DEDUP` + off then on, identical prompts and sampling, token-exact equality, with **more than + one replay per padded bucket** — a first replay can be correct by accident, and the + whole hazard of a shared executable is that the *second* visit to a shape is the one + that has to re-point it. +7. **Not gated, deliberately:** throughput. This row must not be sold as a speed + change. The reportable numbers are executable count and capture wall time. + +## Dependencies + +- The `vt` graph seam `include/vt/backend.h:196-209` and its two accelerator + implementations. No other row is editing them: `ENG-CUDAGRAPH-BREAK` is a parallel + agent on [#1163](https://github.com/mudler/vllm.cpp/issues/1163) and works above the + seam, on the drivers and the runner routing. +- `scripts/check-env-doc.py` for the new environment name. +- No dependency on the diffusion path, which captures nothing today. + +## Work breakdown + +- **W1 — the shared registry.** `src/vt/graph_dedup.h`, `tests/vt/test_graph_dedup.cpp`, + red first, then green, then the mutations. Device-free; gated on every platform. +- **W2 — the runtime ops table.** `src/vt/graph_dedup_runtime.h`, one source bound to + both runtimes, plus the CUDA backend wiring and the environment documentation. + Compile-gated by `cuda-fat-build`. +- **W3 — the ROCm wiring.** The same three edits in `rocm_backend.hip`. +- **W4 — the device A/B.** A leased GPU, both arms of the same binary, token-exact + comparison and the executable-count ratio. + +## Risks/decisions + +- **A shared executable is shared state, and that is the whole risk.** Two handles that + resolve to one executable cannot be replayed concurrently on two streams, because the + second replay's update would re-point the executable out from under the first. Today's + drivers replay decode graphs from one runner thread on one queue + (`src/vllm/model_executor/models/qwen3_5.cpp:9840,10270`), so the constraint holds at + the base, but it is a constraint the seam did not previously carry. It is documented at + the registry, and `ENG-CUDAGRAPH-BREAK` must honour it if it widens who replays. +- **Alternating shapes cost one update per switch.** The registry re-points only when + the group's current raw graph differs, mirroring `:267`, so a run that stays in one + bucket pays nothing. A workload that alternates buckets every step pays one + `cudaGraphExecUpdate` per step. This is the reason the row is `T2` and the reason the + default is off until W4 measures it: shipping it on by default without that number + would be trading an unmeasured latency for an unmeasured memory saving. +- **Retaining the raw graph costs host memory.** A `cudaGraph_t` holds the node + descriptions, not the device-side executable image; it is retained only while dedup is + on, and it is released with its handle. The direction is safe with respect to + [decode-graph-scratch-uaf-2026-07-18.md](decode-graph-scratch-uaf-2026-07-18.md), + which is a premature-**free** hazard. +- **The signature can only be too coarse, never too fine, and coarse is safe.** Every + hit is probed with the real driver call before it is honoured. This is deliberate: + a signature that had to be exhaustive would be a correctness surface, and it is + instead a lookup key. +- **`cudaGraphExecUpdate` changed signature at CUDA 12.** The 3-argument + `cudaGraphExecUpdateResultInfo` form is used at `CUDART_VERSION >= 12000` and the + legacy 4-argument form below it, so the file compiles on both. HIP has only the + 4-argument form. + +## Owed + +| Item | Issue | Why not here | +|---|---|---| +| The device byte-identity A/B and the executable-count measurement (W4) | [#1162](https://github.com/mudler/vllm.cpp/issues/1162) | needs a leased CUDA box; the CPU tier proves the registry's launch-sequence identity, not the device's | +| Flipping `VT_CUDA_GRAPH_DEDUP` on by default | [#1162](https://github.com/mudler/vllm.cpp/issues/1162) | gated on W4. A default is a measurement, not a preference | +| The ROCm leg's compile and run verification | [#41](https://github.com/mudler/vllm.cpp/issues/41) | no ROCm hardware or `hipcc` is reachable from this session and CI has no ROCm job, so the HIP wiring is written against the same shared header but is compile-unverified | + +## Stop conditions + +- Stop and report if the device A/B shows any token divergence between the arms. Byte + identity is the premise of the whole row, not a nice property of it. +- Stop if a measurement shows the per-switch update cost is material at the bucket + churn a real serving grid produces. The row is then a memory-versus-latency trade to + decide, not an optimization to enable. +- Do not extend this row to break points or coverage. That is + [#1163](https://github.com/mudler/vllm.cpp/issues/1163) and it has an owner. diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index 0a7606c3d..35cbec05e 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -270,6 +270,7 @@ Qwen3.6-27B NVFP4, GB10, whole serving window. | Peak GPU memory | 40,996 MiB | 70,531 MiB | 1.720x | **PASS** | | 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)) | +| 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. diff --git a/docs/STATUS.md b/docs/STATUS.md index a589bb4c4..3c4a163c6 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -1626,6 +1626,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` From c1ca9ac4bcb84db348c890894017ad6d42fb62ae Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Tue, 18 Aug 2026 00:59:21 +0000 Subject: [PATCH 2/9] feat(ENG-CUDAGRAPH-DEDUP): fold captures that share a topology onto one graph executable MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Every capture used to instantiate its own executable and throw the raw graph away, so a model held one executable per padded decode bucket — 7 at `max_num_seqs=32`, 11 at 64 — and eight capture drivers each held their own set. The decode graphs of two padded batch sizes are usually the same node topology with different parameters, which is the case `cudaGraphExecUpdate` exists for. `src/vt/graph_dedup.h` is the registry: it keys captures by a structural signature and re-points one executable instead of instantiating a second. `src/vt/graph_dedup_runtime.h` binds it to CUDA and to HIP from one source, because the two graph APIs differ only by symbol prefix and by two call shapes, and a second hand-written copy is the parallel path the protocol forbids. Both accelerator backends route `EndCaptureGraph`, `ReplayGraph` and `DestroyGraph` through it, dispatching on `Owns()` rather than a mode flag so a handle minted before the registry existed is still replayed as the plain executable it is. The signature is a lookup key and never an authority. Every candidate fold is probed with the real driver update on a throwaway executable first, and a probe the driver refuses gives that capture its own executable, so a signature that is too coarse costs a wasted probe and can never make a replay wrong. That is also why the probe is transient where SGLang's is persistent: its backend seals the capture phase and frees the probes, and ours has no capture phase to seal, so a persistent probe would hold two executables per group forever. Default OFF behind `VT_CUDA_GRAPH_DEDUP`, and it stays off until measured. A workload that alternates padded buckets every step pays one update per switch, and the device byte-identity A/B that would price that needs a leased CUDA box this session does not have. It is owed under the spec's `## Owed`, together with the ROCm compile: no ROCm hardware or `hipcc` is reachable here and CI has no ROCm job, so the HIP leg is written against the shared header but unverified. Gated RED-first: `tests/vt/test_graph_dedup.cpp` was written and run against an absent header, then 10 cases / 43 assertions green. The suite drives the registry through a fake device whose launch call records which graph the executable actually reflects, so "the right nodes ran" is an observable sequence rather than an assertion about intent, and every case replays more than once per shape — the risk of a shared executable is that the second visit to a shape is the one that has to re-point it. One case had to be added after the fact: the original sibling-destroy case could not see a capture landing on a freed graph's address, so the fake now hands that storage back and the case reuses it deliberately. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code] --- .agents/engine-matrix.md | 2 +- docs/ENVIRONMENT.md | 1 + src/vt/cuda/cuda_backend.cu | 34 +++- src/vt/graph_dedup.h | 260 +++++++++++++++++++++++ src/vt/graph_dedup_runtime.h | 333 ++++++++++++++++++++++++++++++ src/vt/rocm/rocm_backend.hip | 29 ++- tests/CMakeLists.txt | 2 + tests/vt/test_graph_dedup.cpp | 374 ++++++++++++++++++++++++++++++++++ 8 files changed, 1028 insertions(+), 7 deletions(-) create mode 100644 src/vt/graph_dedup.h create mode 100644 src/vt/graph_dedup_runtime.h create mode 100644 tests/vt/test_graph_dedup.cpp diff --git a/.agents/engine-matrix.md b/.agents/engine-matrix.md index 3e3e9fcb9..c0b3f10c1 100644 --- a/.agents/engine-matrix.md +++ b/.agents/engine-matrix.md @@ -60,7 +60,7 @@ forensics: roadmap_v1.md and the parity ledger. | `KV-PREFIX-MATCH-UNIT` | `--prefix-match-unit` (config `prefix_match_unit`): the finest token boundary a prefix-cache hit can land on == the `hash_block_size`/"prefix match unit" the block hasher uses. NEW in 0.26 (absent at the prior `e24d1b24`/0.25.0 pin). For a HYBRID/multi-group model the resolver `resolve_kv_cache_block_sizes` computes `hash_block_size = prefix_match_unit if set else gcd(group_block_sizes)` (scheduler block size = `lcm`), letting matching land FINER than a physical block (e.g. 16/32 tokens inside a 1024-token block) provided every group block size is divisible by it; single-group (dense) models ignore the knob. Backs off to the scheduler block size when no prefix-cache/connector consumer is active or a mamba group diverges from `cache_block_size` (mamba_cache_mode != "align"); throws on a non-divisible unit. **W0 spike + W1 resolver LANDED 2026-07-28 (`CLAIM-PREFIX-MATCH-UNIT`, NOT pushed):** `resolve_kv_cache_block_sizes` ported 1:1 (explicit-parameter signature vs upstream's `VllmConfig`, our config surface is threaded), RED-first unit-gated (default gcd `!=` `=16` override). `PARTIAL`: the config/CLI/ABI field (W2), the scheduler threading of a resolved `hash_block_size != block_size` + mamba partial-tail stop (W3, needs the `KV-BLOCK-POOL` align path that still throws), and the benchmark (W4) are deferred. Default path byte-identical (single-group inert; scheduler still passes `block_size`). | T1 | `vllm/engine/arg_utils.py:696,1222,1940`; `vllm/config/cache.py:56-67`; resolver `vllm/v1/core/kv_cache_utils.py:626-688`; hasher `:691-748`; call site `vllm/v1/engine/core.py:154`; scheduler `vllm/v1/core/sched/scheduler.py:76,268-270,282,312-318`; fine-grained view `vllm/v1/core/single_type_kv_cache_manager.py:683,697` | resolver `src/vllm/v1/core/kv_cache_utils.cpp:638` (`resolve_kv_cache_block_sizes`), decl `include/vllm/v1/core/kv_cache_utils.h`; hash_block_size already plumbed `get_request_block_hasher` `src/vllm/v1/core/kv_cache_utils.cpp:577`; DEFERRED align path throws `src/vllm/v1/core/block_pool.cpp:93,220` (shared with `KV-BLOCK-POOL`) | `tests/vllm/v1/test_prefix_match_unit.cpp:64,88,99,119,129,145,164,186` 8/8 (29 assertions): single-group inert + DCP scale, multi-group default=gcd, `=16` override finer-than-default (RED), finer-than-1024-block, non-divisible throws, no-consumer back-off + connector-alone re-enable, mamba non-align back-off vs align gcd, hasher-granularity RED (coarse 2 vs fine 4 hashes); [parity-ledger.md](parity-ledger.md) | [prefix-match-unit.md](specs/prefix-match-unit.md) | `PARTIAL` | `CLAIM-PREFIX-MATCH-UNIT` | | `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 eight drivers | CPU-tier contract suite over a fake ops table, RED-first, so every guarantee is gated on every platform and not only on a CUDA box; the fake's launch log makes "the right nodes ran" an observable sequence over MORE than one replay per shape. CUDA leg compile-gated by the `cuda-fat-build` CI job. STILL OWED: the device same-binary A/B proving a deduped replay is token-identical, the exec-count ratio, and the ROCm compile | [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-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 eight drivers | `tests/vt/test_graph_dedup.cpp` 10/10 cases, 43 assertions, RED-first (written and run against an absent header) 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; 9/9 negative mutations detected. CUDA leg compile-gated by `.github/workflows/ci.yml:669`. STILL OWED: the device same-binary A/B proving a deduped replay token-identical, the exec-count ratio, and the ROCm compile | [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 eight hand-rolled drivers become one. **Coverage 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 `breakable_cuda_graph.py:204-241,244-260` @ `f63458b5be` (decorator + runtime stream capture, no compiler) | 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` | owed: reachability mutation (delete the production call site, rerun the focused gate); bit-exactness vs eager on every migrated model over MORE than one replay; the host-lifetime contract of [decode-graph-scratch-uaf-2026-07-18.md](specs/decode-graph-scratch-uaf-2026-07-18.md) enforced AT the seam | [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `INVENTORIED` | [#1163](https://github.com/mudler/vllm.cpp/issues/1163), [#1020](https://github.com/mudler/vllm.cpp/issues/1020) | | `ENG-CUDAGRAPH-DIFFUSION` | Capture the LTX-2.5 denoise loop (fixed shapes, many identical iterations — the ideal graph target). **BLOCKED, and the blocker is ours:** the render does almost no device compute to capture | T2 | SGLang enabled BCG on this shape AFTER our pin — LTX-2 H200 two-stage 10.75s->6.90s (`d4be483efb`), SANA 1024px -26% (`6c7498113f`), SANA denoise 0.73->0.457s (`56ef810cad`). Dated events, NOT pinned evidence; their win is mostly PyTorch host tax we do not pay | NO capture at all: `grep` for capture across `src/vllm/model_executor/models/ltx2*.cpp` returns nothing | blocked by [#1024](https://github.com/mudler/vllm.cpp/issues/1024) (GPU util **exactly 0 in 321 of 347 samples**, 1.00 core of 20 held for 17+ min after staging), [#1007](https://github.com/mudler/vllm.cpp/issues/1007) (VAE decode has no device arm), [#1087](https://github.com/mudler/vllm.cpp/issues/1087) (**57-66% of wall** is ONE resolution-CONSTANT serial host phase), [#1010](https://github.com/mudler/vllm.cpp/issues/1010) (no phase-boundary log). Decision point is a MEASUREMENT of GPU-busy vs wall once device-resident, not an implementation | [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `INVENTORIED` | [#1164](https://github.com/mudler/vllm.cpp/issues/1164) | | `ENG-BATCH-INVARIANT` | Opt-in deterministic execution across scheduler batch sizes (`VLLM_BATCH_INVARIANT=1`): batch-invariant matmul/norm/attention/collectives plus persistent-scheduler NVFP4; production default remains off | T1 | default/env `vllm/envs.py:89,576-578`; initialization `vllm/v1/worker/gpu_worker.py:1262`; NVFP4 dispatch `csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_sm120_kernels.cu:212-220`; suite fixture `tests/v1/determinism/conftest.py:9-12`; operator/e2e `tests/v1/determinism/test_nvfp4_batch_invariant_scaled_mm.py`, `tests/v1/determinism/test_nvfp4_batch_invariant.py` @ `702f481` | - | [W3-C3R executed contract](specs/nvfp4-persistent-plan-cache.md#w3-c3r-batch-shape-localization-and-gate-correction-2026-07-13): production-default ours and vLLM both change outputs across batch shapes; no local opt-in implementation is claimed | `planned: specs/batch-invariant-execution.md` | `INVENTORIED` | - | diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index d18bd64fa..7674681f4 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -97,6 +97,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) | diff --git a/src/vt/cuda/cuda_backend.cu b/src/vt/cuda/cuda_backend.cu index 75c5f63b1..d654cfa65 100644 --- a/src/vt/cuda/cuda_backend.cu +++ b/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::graph_dedup_rt::Ops()); + } + return dedup_->Register(reinterpret_cast(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(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(AsStream(q))); + } else { + Check(cudaGraphLaunch(reinterpret_cast(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(graph)); + if (graph == nullptr) return; + if (dedup_ != nullptr && dedup_->Owns(graph)) { + dedup_->Destroy(graph); + return; + } + cudaGraphExecDestroy(reinterpret_cast(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 dedup_; }; // Registers kCUDA during static init (registration must complete before diff --git a/src/vt/graph_dedup.h b/src/vt/graph_dedup.h new file mode 100644 index 000000000..f816ed265 --- /dev/null +++ b/src/vt/graph_dedup.h @@ -0,0 +1,260 @@ +// vllm.cpp original — graph-executable dedup behind the vt capture seam. +// +// Row ENG-CUDAGRAPH-DEDUP, issue #1162, spec .agents/specs/eng-cudagraph-dedup.md. +// +// WHAT THIS REPLACES. Backend::EndCaptureGraph instantiates one graph executable per +// capture and destroys the raw graph immediately (src/vt/cuda/cuda_backend.cu, and the +// same shape on hipGraph). A model therefore holds one executable per padded decode +// bucket — 7 of them at max_num_seqs=32 and 11 at 64 — and eight hand-rolled drivers +// each hold their own set. The decode graphs of two padded batch sizes are usually the +// same node topology with different parameters, which is exactly the case +// cudaGraphExecUpdate exists for: re-point ONE executable instead of instantiating a +// second. +// +// This is a MEMORY and CAPTURE-TIME change, not a throughput change. A deduped replay +// launches the same nodes with the same parameters; if it launched anything else that +// would be a correctness defect, which is why the gate on this file is byte-identity +// and not a ratio. +// +// WHY THE SIGNATURE IS ONLY A LOOKUP KEY. Register groups captures by a structural +// signature, but never trusts it: every candidate match is PROBED with the real driver +// update before it is honoured, on a throwaway executable, and a probe the driver +// refuses gives that capture its own executable. So a signature that is too coarse +// costs a wasted probe and nothing else, and a signature can never make a replay wrong. +// The alternative — a signature exhaustive enough to be trusted — would make this file +// a correctness surface that has to track every node-parameter class the driver knows. +// +// WHY THE PROBE IS TRANSIENT. SGLang keeps a second `compat_exec` per group alive for +// the whole capture phase and frees it in seal() at end_cuda_graph_capture +// (cuda_graph_dedup_mixin.py:236-238,244-251 @ f63458b5be). OUR SEAM HAS NO CAPTURE +// PHASE: the drivers capture lazily, the first time a padded bucket is seen, interleaved +// with replays of buckets already captured. A persistent probe would therefore never be +// sealed and would hold TWO executables per group for the process, which is worse than +// the one-per-bucket baseline this file exists to reduce. Instantiating the probe inside +// Register and destroying it there keeps steady state at one executable per group, and +// costs exactly what the pre-dedup path already paid: one instantiate per capture. +// +// THREADING. A group's executable is shared state. Two handles that resolve to one +// executable must not be replayed concurrently on two streams, because the second +// replay's update would re-point the executable out from under the first. Today's +// drivers replay decode graphs from one runner thread on one queue, so the constraint +// holds; it is stated here because it is a constraint the seam did not previously carry. +#ifndef VT_GRAPH_DEDUP_H_ +#define VT_GRAPH_DEDUP_H_ + +#include +#include +#include +#include +#include +#include +#include +#include + +#include "vt/dtype.h" // VT_CHECK + +namespace vt { + +// The device operations the registry needs, supplied by whichever accelerator backend +// owns the capture. Free function pointers rather than virtuals so this header stays +// device-free and unit-testable on every platform, the way +// src/vt/cuda/graph_safe_scratch.h keeps its retire bookkeeping portable. +struct GraphDedupOps { + // A structural signature of a raw captured graph. Equal signatures are CANDIDATES + // for sharing one executable; the driver update below is the authority. + std::string (*signature)(void* raw_graph) = nullptr; + // cudaGraphInstantiate / hipGraphInstantiate. + void* (*instantiate)(void* raw_graph) = nullptr; + // cudaGraphExecUpdate / hipGraphExecUpdate. False on any non-success, with a + // human-readable reason in `detail` when it is not null. + bool (*update)(void* exec, void* raw_graph, std::string* detail) = nullptr; + void (*destroy_exec)(void* exec) = nullptr; + void (*destroy_graph)(void* raw_graph) = nullptr; + void (*launch)(void* exec, void* stream) = nullptr; +}; + +// The polarity of VT_CUDA_GRAPH_DEDUP, split out so it is testable without touching the +// process environment. Anything that is not exactly "1" leaves the capture path +// byte-identical to the pre-dedup one. +inline bool GraphDedupEnabledFor(const char* value) { + return value != nullptr && value[0] == '1' && value[1] == '\0'; +} + +// Read once: the drivers capture lazily and repeatedly, and a knob that could change +// between two captures of the same model would produce a mixed set of handles for no +// stated reason. +inline bool GraphDedupEnabled() { + static const bool enabled = GraphDedupEnabledFor(std::getenv("VT_CUDA_GRAPH_DEDUP")); + return enabled; +} + +// Owns every captured raw graph handed to it and every executable it instantiates. +// Register returns the opaque handle the vt seam's void* contract already promises; the +// backend routes ReplayGraph / DestroyGraph back here for the handles Owns() claims, and +// down the plain path for the ones it does not. +class GraphDedupRegistry { + public: + // `log` receives the capture-count / executable-count line after each registration. + // Null silences it, which is what the unit suite wants and what a caller that only + // reads Stats() programmatically can ask for. + explicit GraphDedupRegistry(const GraphDedupOps& ops, std::FILE* log = stderr) + : ops_(ops), log_(log) { + VT_CHECK(ops_.signature != nullptr && ops_.instantiate != nullptr && + ops_.update != nullptr && ops_.destroy_exec != nullptr && + ops_.destroy_graph != nullptr && ops_.launch != nullptr, + "graph dedup: incomplete device op table"); + } + ~GraphDedupRegistry() { Close(); } + + GraphDedupRegistry(const GraphDedupRegistry&) = delete; + GraphDedupRegistry& operator=(const GraphDedupRegistry&) = delete; + + // Takes ownership of `raw_graph` and returns the handle to replay and destroy it by. + void* Register(void* raw_graph) { + VT_CHECK(raw_graph != nullptr, "graph dedup: cannot register a null graph"); + const std::string signature = ops_.signature(raw_graph); + std::vector>& candidates = groups_[signature]; + + Group* group = nullptr; + for (const std::unique_ptr& candidate : candidates) { + if (ProbeAccepts(*candidate, raw_graph)) { + group = candidate.get(); + break; + } + } + if (group == nullptr) { + auto created = std::make_unique(); + created->exec = ops_.instantiate(raw_graph); + created->current_raw = raw_graph; + group = created.get(); + candidates.push_back(std::move(created)); + } + group->raws.push_back(raw_graph); + + auto handle = std::make_unique(); + handle->group = group; + handle->raw = raw_graph; + handle->signature = signature; + void* key = handle.get(); + handles_.emplace(key, std::move(handle)); + + if (log_ != nullptr) { + // Mirrors SGLang's "captured %d CUDA graphs, deduped to %d execs" + // (cuda_graph_dedup_mixin.py:358). Emitted per registration rather than once at + // the end because there is no end: see the capture-phase note at the top. + std::fprintf(log_, "vt graph dedup: captured %zu graphs, deduped to %zu execs\n", + CapturedCount(), ExecCount()); + } + return key; + } + + bool Owns(void* handle) const { return handles_.find(handle) != handles_.end(); } + + void Replay(void* handle, void* stream) { + Handle& entry = Lookup(handle); + Group& group = *entry.group; + if (group.current_raw != entry.raw) { + std::string detail; + const bool ok = ops_.update(group.exec, entry.raw, &detail); + // Register already probed this exact fold, so a refusal here is an invariant + // violation rather than an unsupported capture, and silently launching the + // executable's PREVIOUS contents would be the worst possible response. + VT_CHECK(ok, std::string("graph dedup: replay update refused (") + detail + ")"); + group.current_raw = entry.raw; + } + ops_.launch(group.exec, stream); + } + + void Destroy(void* handle) { + auto found = handles_.find(handle); + if (found == handles_.end()) return; + Handle& entry = *found->second; + Group& group = *entry.group; + + group.raws.erase(std::remove(group.raws.begin(), group.raws.end(), entry.raw), + group.raws.end()); + // The executable outlives the graph it was last pointed at, and a later capture can + // land on the freed address. Forgetting the address is what forces the next replay + // to re-point rather than assume. + if (group.current_raw == entry.raw) group.current_raw = nullptr; + ops_.destroy_graph(entry.raw); + + if (group.raws.empty()) { + ops_.destroy_exec(group.exec); + auto bucket = groups_.find(entry.signature); + if (bucket != groups_.end()) { + std::vector>& candidates = bucket->second; + candidates.erase(std::remove_if(candidates.begin(), candidates.end(), + [&group](const std::unique_ptr& owned) { + return owned.get() == &group; + }), + candidates.end()); + if (candidates.empty()) groups_.erase(bucket); + } + } + handles_.erase(found); + } + + // Live captures, and the executables they share between them. The pre-dedup path + // would report these equal. + std::size_t CapturedCount() const { return handles_.size(); } + std::size_t ExecCount() const { + std::size_t count = 0; + for (const auto& bucket : groups_) count += bucket.second.size(); + return count; + } + + void Close() { + for (auto& bucket : groups_) { + for (std::unique_ptr& group : bucket.second) { + for (void* raw : group->raws) ops_.destroy_graph(raw); + group->raws.clear(); + ops_.destroy_exec(group->exec); + group->exec = nullptr; + } + } + groups_.clear(); + handles_.clear(); + } + + private: + struct Group { + void* exec = nullptr; + // The raw graph the executable currently reflects, or null when that graph has been + // destroyed. Never dereferenced here — only compared, and only against a live raw. + void* current_raw = nullptr; + std::vector raws; + }; + struct Handle { + Group* group = nullptr; + void* raw = nullptr; + std::string signature; + }; + + // Can `candidate`'s executable be re-pointed onto `raw_graph`? Answered on a + // throwaway executable so a refusal cannot leave the live one in a state the driver + // documentation does not define. + bool ProbeAccepts(const Group& candidate, void* raw_graph) const { + if (candidate.raws.empty()) return false; + void* probe = ops_.instantiate(candidate.raws.front()); + std::string detail; + const bool ok = ops_.update(probe, raw_graph, &detail); + ops_.destroy_exec(probe); + return ok; + } + + Handle& Lookup(void* handle) { + auto found = handles_.find(handle); + VT_CHECK(found != handles_.end(), "graph dedup: handle does not belong to this registry"); + return *found->second; + } + + GraphDedupOps ops_; + std::FILE* log_ = nullptr; + std::unordered_map>> groups_; + std::unordered_map> handles_; +}; + +} // namespace vt + +#endif // VT_GRAPH_DEDUP_H_ diff --git a/src/vt/graph_dedup_runtime.h b/src/vt/graph_dedup_runtime.h new file mode 100644 index 000000000..3f76c8e23 --- /dev/null +++ b/src/vt/graph_dedup_runtime.h @@ -0,0 +1,333 @@ +// vllm.cpp original — the CUDA / HIP binding for the graph-executable dedup registry. +// +// Row ENG-CUDAGRAPH-DEDUP, issue #1162, spec .agents/specs/eng-cudagraph-dedup.md. +// Ported from SGLang's dedup mixin at pin f63458b5be: +// python/sglang/srt/model_executor/runner_backend/cuda_graph_dedup_mixin.py +// :27-37 dedup_update -> Update +// :105-114 kernel_node_payload -> AppendKernelPayload +// :117-136 graph_node_payload -> AppendNodePayload +// :139-179 graph_signature -> AppendGraphSignature +// +// ONE source, two runtimes. The CUDA and HIP graph APIs differ only by symbol prefix +// and by two call shapes, so binding them with an alias block keeps this a single path +// instead of the hand-written parallel path AGENTS.md forbids. Include it from a .cu +// for the CUDA leg, or define VT_GRAPH_DEDUP_HIP first and include it from a .hip. +// +// DELIBERATE ADAPTATIONS, both recorded in the spec's `## Upstream chain`: +// * kernel identity is the host function POINTER (cudaKernelNodeParams::func) rather +// than the demangled name SGLang reads through cuKernelGetName / cuFuncGetName. The +// runtime API hands us the pointer directly, no driver-API dependency is added, and +// pointer identity is strictly stronger than name identity — two instantiations of +// one template can share a name prefix but never an address. +// * SGLang's launch-attribute probing (cuGraphKernelNodeGetAttribute, :58-102) is +// driver-API-only and is dropped. Dropping a signature component can only produce a +// false HIT, which Register probes and rejects, never a false miss that silently +// shares an executable. +#ifndef VT_GRAPH_DEDUP_RUNTIME_H_ +#define VT_GRAPH_DEDUP_RUNTIME_H_ + +#include +#include +#include +#include +#include +#include +#include +#include + +#include "vt/graph_dedup.h" + +#if defined(VT_GRAPH_DEDUP_HIP) +#include +#define VTGD_FN(name) hip##name +#define VTGD_SUCCESS hipSuccess +#define VTGD_NODE_KERNEL hipGraphNodeTypeKernel +#define VTGD_NODE_MEMCPY hipGraphNodeTypeMemcpy +#define VTGD_NODE_MEMSET hipGraphNodeTypeMemset +#define VTGD_NODE_GRAPH hipGraphNodeTypeGraph +namespace vt::graph_dedup_rt { +using Graph = hipGraph_t; +using GraphExec = hipGraphExec_t; +using Node = hipGraphNode_t; +using NodeType = hipGraphNodeType; +using KernelNodeParams = hipKernelNodeParams; +using MemcpyNodeParams = hipMemcpy3DParms; +using MemsetNodeParams = hipMemsetParams; +using Stream = hipStream_t; +} // namespace vt::graph_dedup_rt +#else +#include +#define VTGD_FN(name) cuda##name +#define VTGD_SUCCESS cudaSuccess +#define VTGD_NODE_KERNEL cudaGraphNodeTypeKernel +#define VTGD_NODE_MEMCPY cudaGraphNodeTypeMemcpy +#define VTGD_NODE_MEMSET cudaGraphNodeTypeMemset +#define VTGD_NODE_GRAPH cudaGraphNodeTypeGraph +namespace vt::graph_dedup_rt { +using Graph = cudaGraph_t; +using GraphExec = cudaGraphExec_t; +using Node = cudaGraphNode_t; +using NodeType = cudaGraphNodeType; +using KernelNodeParams = cudaKernelNodeParams; +using MemcpyNodeParams = cudaMemcpy3DParms; +using MemsetNodeParams = cudaMemsetParams; +using Stream = cudaStream_t; +} // namespace vt::graph_dedup_rt +#endif + +namespace vt::graph_dedup_rt { + +inline void AppendNumber(std::string* out, long long value) { + out->append(std::to_string(value)); + out->push_back(','); +} + +inline void AppendGraphSignature(Graph graph, std::string* out, int depth); + +// (func, gridDim, blockDim, sharedMemBytes) — cuda_graph_dedup_mixin.py:105-114 minus +// the driver-API attribute tuple, per the header note. +inline bool AppendKernelPayload(Node node, std::string* out) { + KernelNodeParams params{}; + if (VTGD_FN(GraphKernelNodeGetParams)(node, ¶ms) != VTGD_SUCCESS) return false; + AppendNumber(out, static_cast(reinterpret_cast(params.func))); + AppendNumber(out, params.gridDim.x); + AppendNumber(out, params.gridDim.y); + AppendNumber(out, params.gridDim.z); + AppendNumber(out, params.blockDim.x); + AppendNumber(out, params.blockDim.y); + AppendNumber(out, params.blockDim.z); + AppendNumber(out, params.sharedMemBytes); + return true; +} + +// cuda_graph_dedup_mixin.py:117-136, the same five cases in the same order. +inline bool AppendNodePayload(Node node, std::string* out, int depth) { + NodeType type{}; + if (VTGD_FN(GraphNodeGetType)(node, &type) != VTGD_SUCCESS) return false; + AppendNumber(out, static_cast(type)); + switch (type) { + case VTGD_NODE_KERNEL: + return AppendKernelPayload(node, out); + case VTGD_NODE_MEMCPY: { + MemcpyNodeParams params{}; + if (VTGD_FN(GraphMemcpyNodeGetParams)(node, ¶ms) != VTGD_SUCCESS) return false; + AppendNumber(out, static_cast(params.kind)); + AppendNumber(out, static_cast(params.extent.width)); + AppendNumber(out, static_cast(params.extent.height)); + AppendNumber(out, static_cast(params.extent.depth)); + return true; + } + case VTGD_NODE_MEMSET: { + MemsetNodeParams params{}; + if (VTGD_FN(GraphMemsetNodeGetParams)(node, ¶ms) != VTGD_SUCCESS) return false; + AppendNumber(out, static_cast(params.elementSize)); + AppendNumber(out, static_cast(params.width)); + AppendNumber(out, static_cast(params.height)); + return true; + } + case VTGD_NODE_GRAPH: { + Graph child = nullptr; + if (VTGD_FN(GraphChildGraphNodeGetGraph)(node, &child) != VTGD_SUCCESS) return false; + // Bounded: a capture nested past this depth degrades to a coarser key, which the + // probe still guards. Unbounded recursion over driver-owned structure is not a + // risk worth taking inside a capture path. + if (depth >= 4) return true; + AppendGraphSignature(child, out, depth + 1); + return true; + } + default: + return true; + } +} + +// cuda_graph_dedup_mixin.py:139-179. Node order as the runtime reports it is not a +// contract, so the nodes are re-indexed by a deterministic topological order (Kahn, +// always taking the lowest available index, exactly as the upstream heapq does) and the +// edge set is emitted in that order's terms. +inline void AppendGraphSignature(Graph graph, std::string* out, int depth) { + std::size_t num_nodes = 0; + if (VTGD_FN(GraphGetNodes)(graph, nullptr, &num_nodes) != VTGD_SUCCESS) { + out->append("nodes?;"); + return; + } + std::vector nodes(num_nodes); + if (num_nodes > 0 && + VTGD_FN(GraphGetNodes)(graph, nodes.data(), &num_nodes) != VTGD_SUCCESS) { + out->append("nodes?;"); + return; + } + + std::size_t num_edges = 0; + if (VTGD_FN(GraphGetEdges)(graph, nullptr, nullptr, &num_edges) != VTGD_SUCCESS) { + out->append("edges?;"); + return; + } + std::vector from(num_edges); + std::vector to(num_edges); + if (num_edges > 0 && VTGD_FN(GraphGetEdges)(graph, from.data(), to.data(), + &num_edges) != VTGD_SUCCESS) { + out->append("edges?;"); + return; + } + + std::unordered_map index; + index.reserve(num_nodes * 2); + for (std::size_t i = 0; i < num_nodes; ++i) index[nodes[i]] = static_cast(i); + + std::vector> children(num_nodes); + std::vector indegree(num_nodes, 0); + std::vector> edges; + edges.reserve(num_edges); + for (std::size_t e = 0; e < num_edges; ++e) { + const auto src = index.find(from[e]); + const auto dst = index.find(to[e]); + if (src == index.end() || dst == index.end()) { + out->append("edge?;"); + return; + } + children[static_cast(src->second)].push_back(dst->second); + ++indegree[static_cast(dst->second)]; + edges.emplace_back(src->second, dst->second); + } + + std::priority_queue, std::greater> ready; + for (std::size_t i = 0; i < num_nodes; ++i) { + if (indegree[i] == 0) ready.push(static_cast(i)); + } + std::vector order; + order.reserve(num_nodes); + while (!ready.empty()) { + const int current = ready.top(); + ready.pop(); + order.push_back(current); + for (int child : children[static_cast(current)]) { + if (--indegree[static_cast(child)] == 0) ready.push(child); + } + } + if (order.size() != num_nodes) { + // A cycle is impossible in a captured graph; if the runtime ever reports one, the + // key degrades rather than the process aborting inside capture. + out->append("cycle?;"); + return; + } + + std::vector topo(num_nodes, 0); + for (std::size_t i = 0; i < order.size(); ++i) { + topo[static_cast(order[i])] = static_cast(i); + } + + out->push_back('['); + for (int node_index : order) { + if (!AppendNodePayload(nodes[static_cast(node_index)], out, depth)) { + out->append("node?;"); + return; + } + out->push_back(';'); + } + out->push_back(']'); + + std::vector> topo_edges; + topo_edges.reserve(edges.size()); + for (const auto& edge : edges) { + topo_edges.emplace_back(topo[static_cast(edge.first)], + topo[static_cast(edge.second)]); + } + std::sort(topo_edges.begin(), topo_edges.end()); + for (const auto& edge : topo_edges) { + AppendNumber(out, edge.first); + AppendNumber(out, edge.second); + out->push_back(';'); + } +} + +inline std::string Signature(void* raw_graph) { + std::string out; + out.reserve(4096); + AppendGraphSignature(static_cast(raw_graph), &out, 0); + return out; +} + +inline void* Instantiate(void* raw_graph) { + GraphExec exec = nullptr; +#if defined(VT_GRAPH_DEDUP_HIP) + if (hipGraphInstantiate(&exec, static_cast(raw_graph), nullptr, nullptr, 0) != + hipSuccess) { + return nullptr; + } +#else + if (cudaGraphInstantiate(&exec, static_cast(raw_graph), 0) != cudaSuccess) { + return nullptr; + } +#endif + return reinterpret_cast(exec); +} + +// cuda_graph_dedup_mixin.py:27-37. False on any non-success, including the case where +// the call succeeds but the result is not "updated", which is the outcome that matters +// and the one a bare error check would miss. +inline bool Update(void* exec, void* raw_graph, std::string* detail) { + auto handle = reinterpret_cast(exec); + auto graph = static_cast(raw_graph); +#if defined(VT_GRAPH_DEDUP_HIP) + hipGraphNode_t error_node = nullptr; + hipGraphExecUpdateResult result = hipGraphExecUpdateSuccess; + const hipError_t status = hipGraphExecUpdate(handle, graph, &error_node, &result); + const bool ok = status == hipSuccess && result == hipGraphExecUpdateSuccess; + if (!ok && detail != nullptr) { + *detail = "err=" + std::to_string(static_cast(status)) + + " result=" + std::to_string(static_cast(result)); + } + return ok; +#elif CUDART_VERSION >= 12000 + cudaGraphExecUpdateResultInfo info{}; + const cudaError_t status = cudaGraphExecUpdate(handle, graph, &info); + const bool ok = status == cudaSuccess && info.result == cudaGraphExecUpdateSuccess; + if (!ok && detail != nullptr) { + *detail = "err=" + std::to_string(static_cast(status)) + + " result=" + std::to_string(static_cast(info.result)); + } + return ok; +#else + cudaGraphNode_t error_node = nullptr; + cudaGraphExecUpdateResult result = cudaGraphExecUpdateSuccess; + const cudaError_t status = cudaGraphExecUpdate(handle, graph, &error_node, &result); + const bool ok = status == cudaSuccess && result == cudaGraphExecUpdateSuccess; + if (!ok && detail != nullptr) { + *detail = "err=" + std::to_string(static_cast(status)) + + " result=" + std::to_string(static_cast(result)); + } + return ok; +#endif +} + +inline void DestroyExec(void* exec) { + if (exec != nullptr) VTGD_FN(GraphExecDestroy)(reinterpret_cast(exec)); +} + +inline void DestroyGraph(void* raw_graph) { + if (raw_graph != nullptr) VTGD_FN(GraphDestroy)(static_cast(raw_graph)); +} + +inline void Launch(void* exec, void* stream) { + VT_CHECK(VTGD_FN(GraphLaunch)(reinterpret_cast(exec), + static_cast(stream)) == VTGD_SUCCESS, + "graph dedup: graph launch failed"); +} + +inline const GraphDedupOps& Ops() { + static const GraphDedupOps ops = [] { + GraphDedupOps table; + table.signature = &Signature; + table.instantiate = &Instantiate; + table.update = &Update; + table.destroy_exec = &DestroyExec; + table.destroy_graph = &DestroyGraph; + table.launch = &Launch; + return table; + }(); + return ops; +} + +} // namespace vt::graph_dedup_rt + +#endif // VT_GRAPH_DEDUP_RUNTIME_H_ diff --git a/src/vt/rocm/rocm_backend.hip b/src/vt/rocm/rocm_backend.hip index 0ee462aff..f539aa6c0 100644 --- a/src/vt/rocm/rocm_backend.hip +++ b/src/vt/rocm/rocm_backend.hip @@ -43,6 +43,9 @@ #include #include "vt/backend.h" +#define VT_GRAPH_DEDUP_HIP 1 +#include "vt/graph_dedup.h" +#include "vt/graph_dedup_runtime.h" #include "vt/rocm/rocm_arch.h" namespace vt::rocm { @@ -271,6 +274,16 @@ class RocmBackend final : public Backend { void* EndCaptureGraph(Queue& q) override { hipGraph_t graph = nullptr; Check(hipStreamEndCapture(AsStream(q), &graph), "hipStreamEndCapture"); + // ENG-CUDAGRAPH-DEDUP (#1162), the same wiring as the CUDA leg against the same + // shared registry and the same shared ops table: with VT_CUDA_GRAPH_DEDUP set, the + // RAW graph is retained and keyed by topology so compatible captures share one + // executable. Default OFF, in which case this is the pre-dedup path byte for byte. + if (vt::GraphDedupEnabled()) { + if (dedup_ == nullptr) { + dedup_ = std::make_unique(vt::graph_dedup_rt::Ops()); + } + return dedup_->Register(reinterpret_cast(graph)); + } hipGraphExec_t exec = nullptr; Check(hipGraphInstantiate(&exec, graph, nullptr, nullptr, 0), "hipGraphInstantiate"); @@ -281,6 +294,10 @@ class RocmBackend final : public Backend { // cudaProfilerStart/Stop around a chosen replay. A rocprofiler equivalent is // later work and deliberately out of W1's scope (spec §2). void ReplayGraph(Queue& q, void* graph) override { + if (dedup_ != nullptr && dedup_->Owns(graph)) { + dedup_->Replay(graph, reinterpret_cast(AsStream(q))); + return; + } Check(hipGraphLaunch(reinterpret_cast(graph), AsStream(q)), "hipGraphLaunch"); } @@ -299,10 +316,13 @@ class RocmBackend final : public Backend { // leak this backend would rather report than swallow, independent of what // the build flags happen to enforce. void DestroyGraph(void* graph) override { - if (graph != nullptr) { - Check(hipGraphExecDestroy(reinterpret_cast(graph)), - "hipGraphExecDestroy"); + if (graph == nullptr) return; + if (dedup_ != nullptr && dedup_->Owns(graph)) { + dedup_->Destroy(graph); + return; } + Check(hipGraphExecDestroy(reinterpret_cast(graph)), + "hipGraphExecDestroy"); } // THE LOAD-BEARING BOOL. This is what decides whether the portable CPU @@ -353,6 +373,9 @@ class RocmBackend final : public Backend { // Single-graph slot for the EndCapture/Replay pair. The handle API // (EndCaptureGraph) stores nothing here — its caller owns the exec. hipGraphExec_t exec_ = nullptr; + // 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 dedup_; }; // Registers every visible AMD GPU at its own Device{kROCM, i} slot, mirroring diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 63706934f..69f6bc6f4 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -1640,6 +1640,8 @@ target_include_directories(test_gdn_prefill_conv PRIVATE ${CMAKE_SOURCE_DIR}/src # The device proof is the DGX serving harness; this pins the never-free guarantee. vllm_cpp_add_test(test_graph_safe_scratch vt/test_graph_safe_scratch.cpp) target_include_directories(test_graph_safe_scratch PRIVATE ${CMAKE_SOURCE_DIR}/src) +vllm_cpp_add_test(test_graph_dedup vt/test_graph_dedup.cpp) +target_include_directories(test_graph_dedup PRIVATE ${CMAKE_SOURCE_DIR}/src) # CPU threadpool (QUANT-GGUF-CPU-THREADPOOL): barrier/park-wake/chunk-steal # stress + the n_threads 1/3/20 determinism battery. Reaches the internal # header under src/ (the pool is not public API). diff --git a/tests/vt/test_graph_dedup.cpp b/tests/vt/test_graph_dedup.cpp new file mode 100644 index 000000000..cfb625976 --- /dev/null +++ b/tests/vt/test_graph_dedup.cpp @@ -0,0 +1,374 @@ +// vllm.cpp original — CPU-tier contract for the graph-executable dedup registry +// (src/vt/graph_dedup.h), row ENG-CUDAGRAPH-DEDUP, issue #1162. +// +// The production hazard is accelerator-only: today every capture instantiates its own +// graph executable (src/vt/cuda/cuda_backend.cu EndCaptureGraph, and the same shape on +// hipGraph), so a model holds one executable per padded decode bucket and eight drivers +// each hold their own set. The registry folds captures whose topology matches onto ONE +// executable and re-points it with cudaGraphExecUpdate / hipGraphExecUpdate. +// +// A shared executable is shared state, so the guarantee that matters is not "fewer +// executables" — it is that a deduped replay launches EXACTLY what the caller captured. +// This suite pins that on every platform by driving the registry through a fake device. +// The fake's launch() records which graph the executable currently reflects, so "the +// right nodes ran" is an observable sequence rather than an assertion about intent, and +// every case replays MORE THAN ONCE per shape: the whole risk of a shared executable is +// that the SECOND visit to a shape is the one that has to re-point it, so a suite that +// replayed each shape once would pass while measuring nothing. +// +// The device-side proof (a same-binary VT_CUDA_GRAPH_DEDUP off/on A/B that is +// token-identical over a real decode) needs a leased CUDA box and is owed by #1162; it +// is recorded under `## Owed` in .agents/specs/eng-cudagraph-dedup.md, not claimed here. +#include + +#include +#include +#include +#include + +#include "vt/graph_dedup.h" + +using vt::GraphDedupOps; +using vt::GraphDedupRegistry; + +namespace { + +// --- the fake device ------------------------------------------------------------- +// A "graph" is a signature plus an identity; an "executable" is a mutable pointer to +// the graph it currently reflects. That is precisely the CUDA model at the level this +// registry works: instantiate binds an executable to a graph, update re-binds it, and +// launch runs whatever it is currently bound to. +struct FakeGraph { + std::string sig; + int id = 0; +}; +// An executable holds a COPY of what it was last bound to, never a pointer into the +// graph. That is the real contract: cudaGraphInstantiate and cudaGraphExecUpdate copy +// the node parameters out, so an executable stays launchable after its source graph is +// destroyed. A fake that stored the pointer would read freed memory in the +// destroy-a-sibling case and would be modelling a hazard the driver does not have. +struct FakeExec { + std::string sig; + int id = 0; +}; + +struct FakeDevice { + int instantiated = 0; + int exec_destroyed = 0; + int graph_destroyed = 0; + int updates = 0; + std::set reject_update_for; // graph ids the driver refuses to update onto + std::vector launch_log; // the id of the graph each launch actually ran + // With `recycle` set, a destroyed graph's STORAGE is handed back to the test instead + // of being freed, so a later capture can be made to land on the same address on + // purpose. That is the only way to test the freed-address hazard deterministically: + // relying on the allocator to reuse the block would make the case a coin flip. + bool recycle = false; + std::vector recycled; + + int live_execs() const { return instantiated - exec_destroyed; } +}; + +FakeDevice* g_dev = nullptr; + +std::string FakeSignature(void* graph) { + return static_cast(graph)->sig; +} + +void* FakeInstantiate(void* graph) { + ++g_dev->instantiated; + auto* source = static_cast(graph); + return new FakeExec{source->sig, source->id}; +} + +// Mirrors cudaGraphExecUpdate's contract: it can only re-point an executable onto a +// graph whose topology matches the one the executable currently holds, and it may +// refuse for a reason the caller cannot see from the topology alone — which is exactly +// what reject_update_for models. +bool FakeUpdate(void* exec_handle, void* graph_handle, std::string* detail) { + ++g_dev->updates; + auto* exec = static_cast(exec_handle); + auto* graph = static_cast(graph_handle); + if (g_dev->reject_update_for.count(graph->id) != 0 || exec->sig != graph->sig) { + if (detail != nullptr) *detail = "fake driver refused the update"; + return false; + } + exec->sig = graph->sig; + exec->id = graph->id; + return true; +} + +void FakeDestroyExec(void* exec_handle) { + ++g_dev->exec_destroyed; + delete static_cast(exec_handle); +} + +void FakeDestroyGraph(void* graph_handle) { + ++g_dev->graph_destroyed; + if (g_dev->recycle) { + g_dev->recycled.push_back(graph_handle); + return; + } + delete static_cast(graph_handle); +} + +void FakeLaunch(void* exec_handle, void* /*stream*/) { + g_dev->launch_log.push_back(static_cast(exec_handle)->id); +} + +GraphDedupOps FakeOps() { + GraphDedupOps ops; + ops.signature = &FakeSignature; + ops.instantiate = &FakeInstantiate; + ops.update = &FakeUpdate; + ops.destroy_exec = &FakeDestroyExec; + ops.destroy_graph = &FakeDestroyGraph; + ops.launch = &FakeLaunch; + return ops; +} + +FakeGraph* MakeGraph(const char* sig, int id) { return new FakeGraph{sig, id}; } + +// A registry that never writes to a log stream, so the suite stays quiet. +GraphDedupRegistry MakeRegistry() { return GraphDedupRegistry(FakeOps(), nullptr); } + +struct DeviceScope { + FakeDevice dev; + DeviceScope() { g_dev = &dev; } + ~DeviceScope() { + for (void* graph : dev.recycled) delete static_cast(graph); + g_dev = nullptr; + } +}; + +} // namespace + +TEST_CASE("dedup folds compatible captures onto one executable") { + DeviceScope scope; + auto registry = MakeRegistry(); + + std::vector handles; + for (int i = 1; i <= 5; ++i) handles.push_back(registry.Register(MakeGraph("decode", i))); + + CHECK(registry.CapturedCount() == 5); + CHECK(registry.ExecCount() == 1); + // The saving is real only if the extra executables are actually GONE, not merely + // unreferenced: a transient probe that leaked would leave the count unchanged while + // ExecCount() reported 1. + CHECK(scope.dev.live_execs() == 1); + + registry.Close(); + CHECK(scope.dev.live_execs() == 0); + CHECK(scope.dev.graph_destroyed == 5); +} + +TEST_CASE("dedup keeps incompatible captures apart") { + DeviceScope scope; + auto registry = MakeRegistry(); + + void* a1 = registry.Register(MakeGraph("prefill", 1)); + void* b1 = registry.Register(MakeGraph("decode", 2)); + void* a2 = registry.Register(MakeGraph("prefill", 3)); + + CHECK(registry.CapturedCount() == 3); + CHECK(registry.ExecCount() == 2); + CHECK(scope.dev.live_execs() == 2); + + // Two shapes, alternating, twice around: neither group may answer for the other. + for (int round = 0; round < 2; ++round) { + registry.Replay(a1, nullptr); + registry.Replay(b1, nullptr); + registry.Replay(a2, nullptr); + } + CHECK(scope.dev.launch_log == std::vector{1, 2, 3, 1, 2, 3}); + + registry.Close(); +} + +TEST_CASE("a deduped replay launches the graph the caller asked for") { + DeviceScope scope; + + // The control arm: what the SAME sequence launches with no dedup at all, which is + // trivially the graph each handle was captured from. Comparing against a recomputed + // expectation rather than a hand-written literal is what makes this an A/B and not a + // restatement of the implementation. + const std::vector order = {1, 1, 2, 3, 2, 2, 1, 3, 3, 1, 2}; + + auto registry = MakeRegistry(); + std::vector handles; + for (int i = 1; i <= 3; ++i) handles.push_back(registry.Register(MakeGraph("decode", i))); + REQUIRE(registry.ExecCount() == 1); + + for (int id : order) registry.Replay(handles[static_cast(id - 1)], nullptr); + + CHECK(scope.dev.launch_log == order); + registry.Close(); +} + +TEST_CASE("dedup does not re-point the executable for a repeated shape") { + DeviceScope scope; + auto registry = MakeRegistry(); + + void* a = registry.Register(MakeGraph("decode", 1)); + void* b = registry.Register(MakeGraph("decode", 2)); + const int after_register = scope.dev.updates; + + registry.Replay(a, nullptr); + const int after_first = scope.dev.updates; + registry.Replay(a, nullptr); + registry.Replay(a, nullptr); + // Three replays of one shape, and only the first may cost an update. This is the + // whole reason a steady decode workload pays nothing for dedup. + CHECK(scope.dev.updates == after_first); + + registry.Replay(b, nullptr); + CHECK(scope.dev.updates == after_first + 1); + CHECK(after_first >= after_register); + CHECK(scope.dev.launch_log == std::vector{1, 1, 1, 2}); + + registry.Close(); +} + +TEST_CASE("a capture the driver refuses to fold gets its own executable") { + DeviceScope scope; + scope.dev.reject_update_for.insert(2); + auto registry = MakeRegistry(); + + void* g1 = registry.Register(MakeGraph("decode", 1)); + void* g2 = registry.Register(MakeGraph("decode", 2)); // same signature, refused + void* g3 = registry.Register(MakeGraph("decode", 3)); // same signature, accepted + + // Degrades to today's behaviour for the refused capture only. It must never abort, + // and it must never leave the refused graph sharing an executable it cannot drive. + CHECK(registry.CapturedCount() == 3); + CHECK(registry.ExecCount() == 2); + + for (int round = 0; round < 2; ++round) { + registry.Replay(g1, nullptr); + registry.Replay(g2, nullptr); + registry.Replay(g3, nullptr); + } + CHECK(scope.dev.launch_log == std::vector{1, 2, 3, 1, 2, 3}); + + registry.Close(); +} + +TEST_CASE("destroying one capture keeps its siblings replayable") { + DeviceScope scope; + auto registry = MakeRegistry(); + + void* g1 = registry.Register(MakeGraph("decode", 1)); + void* g2 = registry.Register(MakeGraph("decode", 2)); + void* g3 = registry.Register(MakeGraph("decode", 3)); + REQUIRE(registry.ExecCount() == 1); + + registry.Replay(g2, nullptr); // the executable now reflects graph 2 + registry.Destroy(g2); // and that graph is about to be freed + CHECK(registry.CapturedCount() == 2); + CHECK(registry.ExecCount() == 1); + CHECK(scope.dev.live_execs() == 1); + CHECK(scope.dev.graph_destroyed == 1); + + // The surviving siblings stay replayable through the executable their destroyed + // sibling was last pointed at. (The address-reuse half of this is a separate case + // below; this one cannot see it, because re-pointing onto a DIFFERENT address happens + // either way.) + registry.Replay(g1, nullptr); + registry.Replay(g3, nullptr); + registry.Replay(g1, nullptr); + CHECK(scope.dev.launch_log == std::vector{2, 1, 3, 1}); + + registry.Destroy(g1); + CHECK(scope.dev.live_execs() == 1); + registry.Destroy(g3); + CHECK(scope.dev.live_execs() == 0); + CHECK(registry.CapturedCount() == 0); + CHECK(registry.ExecCount() == 0); + CHECK(scope.dev.graph_destroyed == 3); + + registry.Close(); +} + +TEST_CASE("a capture reusing a freed graph address does not inherit its replay state") { + DeviceScope scope; + scope.dev.recycle = true; + auto registry = MakeRegistry(); + + registry.Register(MakeGraph("decode", 1)); + void* second = registry.Register(MakeGraph("decode", 2)); + + registry.Replay(second, nullptr); // the shared executable now reflects graph 2 + registry.Destroy(second); // and graph 2 is freed + REQUIRE(scope.dev.recycled.size() == 1); + + // A later capture lands on the freed address. The registry remembers which raw graph + // its executable reflects by ADDRESS, so if it did not forget the destroyed one it + // would compare equal here, skip the re-point, and replay graph 2's nodes under + // graph 7's handle. That is a wrong answer, not a lost optimisation. + auto* reused = static_cast(scope.dev.recycled.back()); + scope.dev.recycled.pop_back(); + *reused = FakeGraph{"decode", 7}; + + void* seventh = registry.Register(reused); + registry.Replay(seventh, nullptr); + registry.Replay(seventh, nullptr); + CHECK(scope.dev.launch_log == std::vector{2, 7, 7}); + + registry.Close(); +} + +TEST_CASE("the registry does not claim a handle it did not mint") { + DeviceScope scope; + auto registry = MakeRegistry(); + + void* mine = registry.Register(MakeGraph("decode", 1)); + int not_a_handle = 0; + + // This is what lets one backend serve a deduped and a plain executable through the + // same void* seam without guessing at a pointer's provenance. + CHECK(registry.Owns(mine)); + CHECK_FALSE(registry.Owns(¬_a_handle)); + CHECK_FALSE(registry.Owns(nullptr)); + + registry.Destroy(mine); + CHECK_FALSE(registry.Owns(mine)); + + registry.Close(); +} + +TEST_CASE("the dedup log line reports the running capture and executable counts") { + DeviceScope scope; + std::FILE* log = std::tmpfile(); + REQUIRE(log != nullptr); + + { + GraphDedupRegistry registry(FakeOps(), log); + registry.Register(MakeGraph("decode", 1)); + registry.Register(MakeGraph("decode", 2)); + registry.Register(MakeGraph("prefill", 3)); + registry.Close(); + } + + std::fflush(log); + std::rewind(log); + std::string text; + char buffer[512]; + while (std::fgets(buffer, sizeof(buffer), log) != nullptr) text += buffer; + std::fclose(log); + + // Mirrors SGLang's "captured %d CUDA graphs, deduped to %d execs" + // (cuda_graph_dedup_mixin.py:358). The ratio has to be readable off a log, because a + // dedup that silently folded nothing would otherwise look identical to one that works. + CHECK(text.find("captured 1 graphs, deduped to 1 execs") != std::string::npos); + CHECK(text.find("captured 2 graphs, deduped to 1 execs") != std::string::npos); + CHECK(text.find("captured 3 graphs, deduped to 2 execs") != std::string::npos); +} + +TEST_CASE("dedup is off unless the environment asks for it") { + // The registry only exists when this is true, so its polarity decides whether the + // production path is byte-identical to the pre-dedup one. + CHECK_FALSE(vt::GraphDedupEnabledFor(nullptr)); + CHECK_FALSE(vt::GraphDedupEnabledFor("")); + CHECK_FALSE(vt::GraphDedupEnabledFor("0")); + CHECK(vt::GraphDedupEnabledFor("1")); +} From ff8237977ae195965a161791abd4e3b2ef7c17d8 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Tue, 18 Aug 2026 01:08:32 +0000 Subject: [PATCH 3/9] chore(ENG-CUDAGRAPH-DEDUP): re-pin RUNNABLE_BASELINE for the row that just entered it `ENG-CUDAGRAPH-DEDUP` reached `ACTIVE` with a spec whose `## Gates` section names commands that can actually fail, so the row entered the runnable gated population and `test_check_gate_commands` went red on the equality that makes the ratchet work in both directions. Growth, re-pinned in the same change and named the way the checker's own message demands. The credit is not inherited: `ctest -R test_graph_dedup` detected 9 of 9 negative mutations of the registry the row adds, which is the strongest kind of credit in this set rather than the weakest. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code] --- scripts/check-gate-commands.py | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/scripts/check-gate-commands.py b/scripts/check-gate-commands.py index 4818e49c1..cf213e66e 100755 --- a/scripts/check-gate-commands.py +++ b/scripts/check-gate-commands.py @@ -300,7 +300,17 @@ def audit() -> list[dict]: # no CUDA or SACRED gate is claimed here, and the DGX three-way at k=2..4 is # recorded as OWED rather than skipped. 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-CUDAGRAPH-DEDUP", "SPEC-MTP-K-GT-1", "ATTN-CHUNKED-LOCAL", "SERVE-RECIPE-ARGS", From 89635e4707930f4d55ceb6c120028dbfc62218cf Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Tue, 18 Aug 2026 01:27:13 +0000 Subject: [PATCH 4/9] test(ENG-CUDAGRAPH-DEDUP): pin the new RUNNABLE_BASELINE entry with its own mutation `check-pr-size` refuses a change to a governance checker that arrives without semantic mutation evidence in the paired suite, and it is right to: a re-pin that only added a row to a frozenset would be indistinguishable from one that quieted the gate. Two cases, mirroring what `SPEC-MTP-K-GT-1` and `SERVE-RECIPE-ARGS` already carry. The first pins both halves of the credit -- the row is in the exact pin, AND its `## Gates` section really does yield a command that can fail -- and asserts on `ctest -R test_graph_dedup` rather than on the `git diff --stat` the extractor also finds, because that one exits 0 unconditionally in a repo and is one of the weak credits this checker's own header names. It also asserts the spec still says the device leg is owed, so the credit cannot quietly come to rest on the CPU tier while the record goes silent about the arm nobody ran. The second removes the entry and proves the set equality goes red. Verified by mutation: deleting `ENG-CUDAGRAPH-DEDUP` from `RUNNABLE_BASELINE` fails both new cases (7 of 39 in total), and the file restores byte-for-byte. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code] --- tests/scripts/test_check_gate_commands.py | 40 +++++++++++++++++++++++ 1 file changed, 40 insertions(+) diff --git a/tests/scripts/test_check_gate_commands.py b/tests/scripts/test_check_gate_commands.py index 737d62cba..574602999 100644 --- a/tests/scripts/test_check_gate_commands.py +++ b/tests/scripts/test_check_gate_commands.py @@ -589,6 +589,46 @@ def test_the_checker_is_wired_into_preflight_and_ci(self): self.assertIn("scripts/check-gate-commands.py --check", ci) self.assertIn("tests/scripts/test_check_gate_commands.py", ci) + def test_cudagraph_dedup_is_credited_for_real_commands(self): + # ENG-CUDAGRAPH-DEDUP (#1162) is a NEW row arriving at ACTIVE, which is + # the first state that puts it in GATED_STATES at all, so it joins the + # runnable population on arrival. TWO things are pinned, not one: the + # row is in the exact pin, AND its Gates section really does yield a + # command that can fail. A row pinned without the second half is a + # certificate for nothing, which is the failure the checker's own header + # admits to for five older credits. + self.assertIn("ENG-CUDAGRAPH-DEDUP", gates.RUNNABLE_BASELINE) + verdicts = {r["id"]: r["verdict"] for r in gates.audit()} + self.assertEqual(verdicts.get("ENG-CUDAGRAPH-DEDUP"), "runnable") + + spec = (ROOT / ".agents/specs/eng-cudagraph-dedup.md").read_text(encoding="utf-8") + section = gates.gates_section(spec) + self.assertIsNotNone(section) + commands = gates.runnable_commands(section) + # The credit rests on the focused suite, not on the preflight line: that + # suite detected 9 of 9 negative mutations of the registry this row adds, + # so it genuinely goes red when the row regresses. `git diff --stat` is + # also extracted from the mutation bullet and is one of the weak credits + # this checker's header names; it is deliberately not what is asserted. + self.assertIn("ctest -R test_graph_dedup", commands) + self.assertIn("./scripts/agent-preflight.sh", commands) + # The row's DEVICE leg is OWED, not skipped, and the spec has to say so. + # Without this the credit could rest on the CPU tier while the record + # stayed silent about the arm nobody ran, which reads as coverage. + self.assertIn("Device byte-identity A/B (owed", spec) + + def test_dropping_cudagraph_dedup_from_the_pin_breaks_it(self): + # MUTATION, in the direction this re-pin actually moved: the entry added + # for #1162 must be what keeps the exact pin agreeing with the audit. + # Remove it and set equality has to go red, which is what proves the row + # was pinned because it entered the population and not to quiet a gate. + reduced = set(gates.RUNNABLE_BASELINE) - {"ENG-CUDAGRAPH-DEDUP"} + self.assertNotEqual(reduced, set(gates.RUNNABLE_BASELINE)) + runnable = {r["id"] for r in gates.audit() if r["verdict"] == "runnable"} + self.assertNotEqual(runnable, reduced) + self.assertEqual(runnable - reduced, {"ENG-CUDAGRAPH-DEDUP"}) + self.assertEqual(runnable, set(gates.RUNNABLE_BASELINE)) + if __name__ == "__main__": unittest.main() From e4ce5571a7b27a81a6b96d13b569bbd409282451 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Tue, 18 Aug 2026 01:46:04 +0000 Subject: [PATCH 5/9] fix(ENG-CUDAGRAPH-DEDUP): bind both shapes of cudaGraphGetEdges, which changed at CUDA 13 `cuda-fat-build` on `nvidia/cuda:13.3.0` failed with four errors at `graph_dedup_runtime.h:161,168`: CUDA 13's `cudaGraphGetEdges` takes a fifth `cudaGraphEdgeData*` parameter, where CUDA 12 and HIP take four. The local syntax check could not have caught this and did not. It ran against CUDA 12.9 runtime headers and passed, which is the honest limit of a proxy: a header-shape check is only as current as the toolkit it ran against. Recorded in the spec's `## Risks/decisions` beside the `cudaGraphExecUpdate` split, which was foreseen, so the difference between the two is on the record. Both shapes are bound in one `GetEdges` helper rather than at the two call sites, so the topology walk stays one piece of code. The CUDA 13 arm fills a real edge data buffer instead of passing null, because null edge data on the filling call is not a shape the documentation promises; the data is discarded, since an edge annotation is not part of the identity this key needs. Verified locally against CUDA 12.9's `cudaGraphGetEdges_v2`, whose signature is identical to CUDA 13's: both arms compile clean under `-Wall -Wextra -Werror`. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code] --- .agents/specs/eng-cudagraph-dedup.md | 15 ++++++++++---- src/vt/graph_dedup_runtime.h | 29 +++++++++++++++++++++++++--- 2 files changed, 37 insertions(+), 7 deletions(-) diff --git a/.agents/specs/eng-cudagraph-dedup.md b/.agents/specs/eng-cudagraph-dedup.md index 21d486f89..3eb365d6c 100644 --- a/.agents/specs/eng-cudagraph-dedup.md +++ b/.agents/specs/eng-cudagraph-dedup.md @@ -230,10 +230,17 @@ an assertion about intent. hit is probed with the real driver call before it is honoured. This is deliberate: a signature that had to be exhaustive would be a correctness surface, and it is instead a lookup key. -- **`cudaGraphExecUpdate` changed signature at CUDA 12.** The 3-argument - `cudaGraphExecUpdateResultInfo` form is used at `CUDART_VERSION >= 12000` and the - legacy 4-argument form below it, so the file compiles on both. HIP has only the - 4-argument form. +- **Two runtime-API calls changed shape across CUDA major versions, and only one was + foreseen.** `cudaGraphExecUpdate` took the 3-argument `cudaGraphExecUpdateResultInfo` + form at CUDA 12, so both it and the legacy 4-argument form are bound; HIP has only + the 4-argument form. `cudaGraphGetEdges` gained a fifth `cudaGraphEdgeData*` + parameter at **CUDA 13**, which was NOT foreseen: the local syntax check ran against + CUDA 12.9 headers and passed, and `cuda-fat-build` on `nvidia/cuda:13.3.0` is what + reported it (4 errors, `graph_dedup_runtime.h:161,168`). Both shapes are now bound in + one `GetEdges` helper so the topology walk stays one piece of code, and the 5-argument + shape is verified locally against CUDA 12.9's identical `cudaGraphGetEdges_v2`. The + general lesson for this file: a header-shape check is only as current as the toolkit + it ran against, and the CI job is the gate, not the local proxy. ## Owed diff --git a/src/vt/graph_dedup_runtime.h b/src/vt/graph_dedup_runtime.h index 3f76c8e23..89c053d08 100644 --- a/src/vt/graph_dedup_runtime.h +++ b/src/vt/graph_dedup_runtime.h @@ -77,6 +77,30 @@ using Stream = cudaStream_t; namespace vt::graph_dedup_rt { +// cudaGraphGetEdges gained a fifth `cudaGraphEdgeData*` parameter in CUDA 13; the +// four-argument form is what CUDA 12 and HIP ship. Both shapes are bound here rather +// than at the call sites so the topology walk below stays one piece of code. This is +// the one place a local compile against CUDA 12 headers could not have caught, and did +// not: the `cuda-fat-build` job on nvidia/cuda:13.3.0 is what reported it. +inline bool GetEdges(Graph graph, Node* from, Node* to, std::size_t* num_edges) { +#if defined(VT_GRAPH_DEDUP_HIP) + return hipGraphGetEdges(graph, from, to, num_edges) == hipSuccess; +#elif CUDART_VERSION >= 13000 + // Queried with every pointer null; filled with a real buffer, because passing null + // edge data to the filling call is not a shape the documentation promises. The data + // is discarded: an edge's annotation is not part of the identity this key needs. + std::vector edge_data; + cudaGraphEdgeData* data = nullptr; + if (from != nullptr) { + edge_data.resize(*num_edges); + data = edge_data.data(); + } + return cudaGraphGetEdges(graph, from, to, data, num_edges) == cudaSuccess; +#else + return cudaGraphGetEdges(graph, from, to, num_edges) == cudaSuccess; +#endif +} + inline void AppendNumber(std::string* out, long long value) { out->append(std::to_string(value)); out->push_back(','); @@ -158,14 +182,13 @@ inline void AppendGraphSignature(Graph graph, std::string* out, int depth) { } std::size_t num_edges = 0; - if (VTGD_FN(GraphGetEdges)(graph, nullptr, nullptr, &num_edges) != VTGD_SUCCESS) { + if (!GetEdges(graph, nullptr, nullptr, &num_edges)) { out->append("edges?;"); return; } std::vector from(num_edges); std::vector to(num_edges); - if (num_edges > 0 && VTGD_FN(GraphGetEdges)(graph, from.data(), to.data(), - &num_edges) != VTGD_SUCCESS) { + if (num_edges > 0 && !GetEdges(graph, from.data(), to.data(), &num_edges)) { out->append("edges?;"); return; } From a0b8ff3cf5b52e3ec3a67a30e99138d1a6e64b3a Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Tue, 18 Aug 2026 03:56:40 +0000 Subject: [PATCH 6/9] fix(ENG-CUDAGRAPH-DEDUP): a failed instantiate now throws where the capture happened, not at the first replay Repairs the findings of the fresh review of #1178. Issue #1162. `Register` assigned `ops_.instantiate(raw_graph)` and moved on, while `graph_dedup_runtime.h` returns null on a driver failure. The pre-dedup path wrapped `cudaGraphInstantiate` in `Check()` and threw at the capture site with the driver's code; with dedup on the same failure minted a valid-looking handle over a null executable, counted it live in `ExecCount()`, and surfaced at the first replay as "graph launch failed" -- the wrong site, the wrong message, and the driver's reason already discarded. `Register` now fails there, releasing the capture whose ownership it had taken so the throw does not also leak it. The probe instantiates a SECOND executable, so it can fail independently, and `cudaGraphExecUpdate` has no defined behaviour for a null one. `ProbeAccepts` now answers "cannot fold" instead of asking the driver about nothing, which degrades to exactly today's one-executable-per-capture path. Two guarantees were unpinned and are now tests. A capture the driver cannot instantiate must throw and count nothing; a probe that cannot instantiate must not drive a null executable. Both were run RED against the unfixed source (6 failing assertions) before the fix. The environment polarity gained the values only the terminator check rejects: deleting `&& value[1] == '\0'` left the whole suite green, so "10" enabled a default-off correctness-sensitive path undetected. 12/12 cases, 61 assertions; three negative mutations applied, detected and byte-restored. Three findings are records rather than code, because fixing them would change behaviour a device A/B is measuring on this commit or is not code at all. `Replay` does not issue the fold `Register` probed: the probe tests `(raws.front(), raw_graph)` and the replay issues `(current_raw, entry.raw)`, which diverge from a group's third member onwards. Honouring the probe therefore treats `cudaGraphExecUpdate` compatibility as transitive across a group, which no documentation states and nothing here asserts. The comment claiming the fold was already probed is corrected to say what is actually assumed; the stronger fix -- probe `current_raw` -- is owed rather than taken, because it changes probe behaviour mid-measurement. `src/vt/graph_dedup_runtime.h` has zero executable coverage on any tier. The mutation count is a statement about `graph_dedup.h` and must not be read as covering the signature builder, whose one silent mode is a signature unstable for some topology: it folds nothing while every CPU test stays green. Recorded in the spec and in the matrix row, with the two device-tier tests that would cover it named and owed. The registry lives on a process-singleton per-device backend, so one unsynchronised registry serves every model on a device. That is wider than the "one runner thread" note it carried, and it is now stated where a future capture-widening row will read it. Also merged up: the spec's `Stats()` never existed and is corrected to the shipped `CapturedCount()` / `ExecCount()`, and the decode-driver count moves from eight to nine here too, after #1179 corrected it on main. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code] --- .agents/engine-matrix.md | 2 +- .agents/specs/eng-cudagraph-dedup.md | 69 ++++++++++++++++++--- src/vt/graph_dedup.h | 53 ++++++++++++++-- tests/vt/test_graph_dedup.cpp | 93 +++++++++++++++++++++++++++- 4 files changed, 198 insertions(+), 19 deletions(-) diff --git a/.agents/engine-matrix.md b/.agents/engine-matrix.md index 3ef5e7f4b..74d5b88e8 100644 --- a/.agents/engine-matrix.md +++ b/.agents/engine-matrix.md @@ -60,7 +60,7 @@ forensics: roadmap_v1.md and the parity ledger. | `KV-PREFIX-MATCH-UNIT` | `--prefix-match-unit` (config `prefix_match_unit`): the finest token boundary a prefix-cache hit can land on == the `hash_block_size`/"prefix match unit" the block hasher uses. NEW in 0.26 (absent at the prior `e24d1b24`/0.25.0 pin). For a HYBRID/multi-group model the resolver `resolve_kv_cache_block_sizes` computes `hash_block_size = prefix_match_unit if set else gcd(group_block_sizes)` (scheduler block size = `lcm`), letting matching land FINER than a physical block (e.g. 16/32 tokens inside a 1024-token block) provided every group block size is divisible by it; single-group (dense) models ignore the knob. Backs off to the scheduler block size when no prefix-cache/connector consumer is active or a mamba group diverges from `cache_block_size` (mamba_cache_mode != "align"); throws on a non-divisible unit. **W0 spike + W1 resolver LANDED 2026-07-28 (`CLAIM-PREFIX-MATCH-UNIT`, NOT pushed):** `resolve_kv_cache_block_sizes` ported 1:1 (explicit-parameter signature vs upstream's `VllmConfig`, our config surface is threaded), RED-first unit-gated (default gcd `!=` `=16` override). `PARTIAL`: the config/CLI/ABI field (W2), the scheduler threading of a resolved `hash_block_size != block_size` + mamba partial-tail stop (W3, needs the `KV-BLOCK-POOL` align path that still throws), and the benchmark (W4) are deferred. Default path byte-identical (single-group inert; scheduler still passes `block_size`). | T1 | `vllm/engine/arg_utils.py:696,1222,1940`; `vllm/config/cache.py:56-67`; resolver `vllm/v1/core/kv_cache_utils.py:626-688`; hasher `:691-748`; call site `vllm/v1/engine/core.py:154`; scheduler `vllm/v1/core/sched/scheduler.py:76,268-270,282,312-318`; fine-grained view `vllm/v1/core/single_type_kv_cache_manager.py:683,697` | resolver `src/vllm/v1/core/kv_cache_utils.cpp:638` (`resolve_kv_cache_block_sizes`), decl `include/vllm/v1/core/kv_cache_utils.h`; hash_block_size already plumbed `get_request_block_hasher` `src/vllm/v1/core/kv_cache_utils.cpp:577`; DEFERRED align path throws `src/vllm/v1/core/block_pool.cpp:93,220` (shared with `KV-BLOCK-POOL`) | `tests/vllm/v1/test_prefix_match_unit.cpp:64,88,99,119,129,145,164,186` 8/8 (29 assertions): single-group inert + DCP scale, multi-group default=gcd, `=16` override finer-than-default (RED), finer-than-1024-block, non-divisible throws, no-consumer back-off + connector-alone re-enable, mamba non-align back-off vs align gcd, hasher-granularity RED (coarse 2 vs fine 4 hashes); [parity-ledger.md](parity-ledger.md) | [prefix-match-unit.md](specs/prefix-match-unit.md) | `PARTIAL` | `CLAIM-PREFIX-MATCH-UNIT` | | `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` 10/10 cases, 43 assertions, RED-first (written and run against an absent header) 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; 9/9 negative mutations detected. CUDA leg compile-gated by `.github/workflows/ci.yml:669`. STILL OWED: the device same-binary A/B proving a deduped replay token-identical, the exec-count ratio, and the ROCm compile | [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-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` 12/12 cases, 61 assertions, RED-first (written and run against an absent header, and the three cases added by the fresh review of #1178 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; 12/12 negative mutations detected (9 at implementation, 3 at review repair). That count covers `src/vt/graph_dedup.h` ONLY: `src/vt/graph_dedup_runtime.h`, which builds the signature, has NO executable coverage on any tier and is compile-gated by `.github/workflows/ci.yml:669` alone — its one silent mode is a signature unstable for some topology, which folds nothing while every CPU test stays green. STILL OWED: the device same-binary A/B proving a deduped replay token-identical, the exec-count ratio, device-tier signature stability/discrimination tests, probing `current_raw` instead of `raws.front()` to retire the update-transitivity assumption, and the ROCm compile | [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:961-986` DECLINES the graph outright when the async device-token mirror is live. **That decline is why this is also a CORRECTNESS row** ([#1179](https://github.com/mudler/vllm.cpp/issues/1179)): a SHIPPED model has already lost its decode graph to the duplication, on the driver's own measurement (`depth-1, graph ON PASS 78/78`; `depth-2, graph OFF PASS 82/82`; `depth-2, graph ON FAIL, slots 1-3 degenerate`), and the fix its comment names is the sibling's `StepDevInputs`. The row still makes NO throughput claim: the prefill refutation on the `ENG-CUDAGRAPH` row (3.8% host idle, >96% GPU-busy, 92.5% glue) stands unchanged | owed: reachability mutation (delete the production call site, rerun the focused gate); bit-exactness vs eager on every migrated model over MORE than one replay; the host-lifetime contract of [decode-graph-scratch-uaf-2026-07-18.md](specs/decode-graph-scratch-uaf-2026-07-18.md) enforced AT the seam; the 11 ported SGLang unit cases; the break-function OUTPUT writeback (`replay_fn`/`_copy_output` `breakable_cuda_graph.py:231-235,172-201`, spec D9) without which a replay leaves the next segment reading capture-time data; the auxiliary-stream auto-join before every segment close (`:353-361`, spec D10), live at `src/vllm/model_executor/models/qwen3_5.cpp:6254-6255,6384` and `src/vllm/model_executor/models/laguna.cpp:2572-2576,2612` | spec [eng-cudagraph-break.md](specs/eng-cudagraph-break.md) (W0 spike DONE 2026-08-18: the existing `vt` capture vocabulary `include/vt/backend.h:208-222` expresses a SEGMENTED capture with NO new virtual, because `EndCaptureGraph` stores nothing (`src/vt/cuda/cuda_backend.cu:225-232`); a break point is expressible with one `thread_local` capture pointer plus a free function, no compiler and no decorator); analysis [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `READY` | [#1163](https://github.com/mudler/vllm.cpp/issues/1163), [#1020](https://github.com/mudler/vllm.cpp/issues/1020) | | `ENG-CUDAGRAPH-DIFFUSION` | Capture the LTX-2.5 denoise loop (fixed shapes, many identical iterations — the ideal graph target). **BLOCKED, and the blocker is ours:** the render does almost no device compute to capture | T2 | SGLang enabled BCG on this shape AFTER our pin — LTX-2 H200 two-stage 10.75s->6.90s (`d4be483efb`), SANA 1024px -26% (`6c7498113f`), SANA denoise 0.73->0.457s (`56ef810cad`). Dated events, NOT pinned evidence; their win is mostly PyTorch host tax we do not pay | NO capture at all: `grep` for capture across `src/vllm/model_executor/models/ltx2*.cpp` returns nothing | blocked by [#1024](https://github.com/mudler/vllm.cpp/issues/1024) (GPU util **exactly 0 in 321 of 347 samples**, 1.00 core of 20 held for 17+ min after staging), [#1007](https://github.com/mudler/vllm.cpp/issues/1007) (VAE decode has no device arm), [#1087](https://github.com/mudler/vllm.cpp/issues/1087) (**57-66% of wall** is ONE resolution-CONSTANT serial host phase), [#1010](https://github.com/mudler/vllm.cpp/issues/1010) (no phase-boundary log). Decision point is a MEASUREMENT of GPU-busy vs wall once device-resident, not an implementation | [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `INVENTORIED` | [#1164](https://github.com/mudler/vllm.cpp/issues/1164) | | `ENG-BATCH-INVARIANT` | Opt-in deterministic execution across scheduler batch sizes (`VLLM_BATCH_INVARIANT=1`): batch-invariant matmul/norm/attention/collectives plus persistent-scheduler NVFP4; production default remains off | T1 | default/env `vllm/envs.py:89,576-578`; initialization `vllm/v1/worker/gpu_worker.py:1262`; NVFP4 dispatch `csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_sm120_kernels.cu:212-220`; suite fixture `tests/v1/determinism/conftest.py:9-12`; operator/e2e `tests/v1/determinism/test_nvfp4_batch_invariant_scaled_mm.py`, `tests/v1/determinism/test_nvfp4_batch_invariant.py` @ `702f481` | - | [W3-C3R executed contract](specs/nvfp4-persistent-plan-cache.md#w3-c3r-batch-shape-localization-and-gate-correction-2026-07-13): production-default ours and vLLM both change outputs across batch shapes; no local opt-in implementation is claimed | `planned: specs/batch-invariant-execution.md` | `INVENTORIED` | - | diff --git a/.agents/specs/eng-cudagraph-dedup.md b/.agents/specs/eng-cudagraph-dedup.md index 3eb365d6c..7dd7165b9 100644 --- a/.agents/specs/eng-cudagraph-dedup.md +++ b/.agents/specs/eng-cudagraph-dedup.md @@ -26,8 +26,9 @@ In scope: after instantiate. Both the signature and the update need it. 3. The CUDA ops table and the CUDA backend wiring, off by default behind `VT_CUDA_GRAPH_DEDUP`. -4. An observable capture-count / exec-count log line and a programmatic `Stats()` - accessor, so the ratio is a measurement rather than a claim. +4. An observable capture-count / exec-count log line and the programmatic + `CapturedCount()` / `ExecCount()` accessors behind it, so the ratio is a measurement + rather than a claim. Out of scope, each owned elsewhere: @@ -35,7 +36,7 @@ Out of scope, each owned elsewhere: [#1163](https://github.com/mudler/vllm.cpp/issues/1163). - The diffusion / denoise-loop capture: `ENG-CUDAGRAPH-DIFFUSION`, [#1164](https://github.com/mudler/vllm.cpp/issues/1164), blocked. -- Retiring the eight hand-rolled decode drivers. This row does not touch their +- Retiring the nine hand-rolled decode drivers. This row does not touch their internals; it changes only what the seam they already call hands back. - Any change to bucket selection, padding, or the `pure_decode` predicate. @@ -57,7 +58,7 @@ Paths below are relative to `python/sglang/srt/`. | Register: signature hit re-points, signature miss instantiates a new group | `:219-242` | `GraphDedupRegistry::Register` | | Replay: update only when the group's current raw graph differs, then launch | `:258-275` | `GraphDedupRegistry::Replay` | | Compatibility probe on a separate executable so a failed probe cannot leave the live one in an undefined state | `:236-238` (`compat_exec`), destroyed by `seal()` at `:244-251` | `GraphDedupRegistry::Register`, **transient** probe — see the divergence below | -| `"captured %d CUDA graphs, deduped to %d execs"` | `:353-358` | `GraphDedupRegistry::Stats()` plus the log line | +| `"captured %d CUDA graphs, deduped to %d execs"` | `:353-358` | `GraphDedupRegistry::CapturedCount()` / `ExecCount()` plus the log line | | Retained raw graph (`torch.cuda.CUDAGraph(keep_graph=True)`) | `:326` | `EndCaptureGraph` stops calling `cudaGraphDestroy` when dedup is on | Three deliberate divergences from the upstream construction, each forced by a @@ -66,7 +67,7 @@ difference in our seam rather than chosen: - **The probe executable is transient, not persistent.** SGLang holds a second `compat_exec` per group for the whole capture phase and frees it in `seal()`, which its backend calls at `end_cuda_graph_capture()` (`:352-358`). **Our seam has - no capture-phase end.** The eight drivers capture lazily, the first time a padded + no capture-phase end.** The nine drivers capture lazily, the first time a padded bucket is seen, interleaved with replays of buckets already captured (`src/vllm/model_executor/models/qwen3_5.cpp:9918`, `src/vllm/model_executor/models/qwen3.cpp:888`). A persistent probe would therefore @@ -106,12 +107,16 @@ the raw graph away: scope. The executables multiply along two axes. `include/vllm/model_executor/models/decode_graph_sizes.h:32-41` -yields 7 padded decode buckets at `max_num_seqs=32` and 11 at 64, and eight drivers +yields 7 padded decode buckets at `max_num_seqs=32` and 11 at 64, and nine drivers each build their own set: `Qwen3_5DecodeGraph`, `Qwen3_5DenseDecodeGraph`, `Qwen3MoeDecodeGraph`, `Qwen3DenseDecodeGraph`, `DeepseekV2DecodeGraph`, `VoxtralDecodeGraph`, plus the graph code in `src/vllm/model_executor/models/deepseek_v4.cpp:1900` and -`src/vllm/model_executor/models/laguna.cpp:2691`. +`src/vllm/model_executor/models/laguna.cpp:2691`, and the DFlash draft graph in +`src/vllm/model_executor/models/qwen3_dflash.cpp:771,870,1038,1091,1095,1106`. This +spec said **eight** until 2026-08-18; the ninth was found and the count corrected by +[#1179](https://github.com/mudler/vllm.cpp/issues/1179) on `ENG-CUDAGRAPH-BREAK`, and +carried here so the two rows describe the same tree. Why it matters here specifically, restated from the issue so this spec stands alone: on GB10 unified memory an out-of-memory event reboots the box, capture time is startup @@ -132,7 +137,7 @@ instantiate were never a source of baked pointers. | New / changed | What | |---|---| -| `src/vt/graph_dedup.h` (new) | `GraphDedupOps` (six function pointers: signature, instantiate, update, destroy exec, destroy graph, launch) and `GraphDedupRegistry`. Header-only, no device dependency, CPU-unit-testable — the same shape `src/vt/cuda/graph_safe_scratch.h` used for the graph-safe scratch bookkeeping. Also `GraphDedupEnabled()`, the `VT_CUDA_GRAPH_DEDUP` read. | +| `src/vt/graph_dedup.h` (new) | `GraphDedupOps` (six function pointers: signature, instantiate, update, destroy exec, destroy graph, launch) and `GraphDedupRegistry`. Header-only, no device dependency, CPU-unit-testable — the same shape `src/vt/cuda/graph_safe_scratch.h` used for the graph-safe scratch bookkeeping. Also `GraphDedupEnabled()`, the `VT_CUDA_GRAPH_DEDUP` read, and the `CapturedCount()` / `ExecCount()` accessors. There is no single `Stats()` method: two accessors read at the call site beat one struct nobody stores. | | `src/vt/graph_dedup_runtime.h` (new) | The CUDA / HIP ops table, written **once** and bound to either runtime by a macro alias block. `GraphSignature()` lives here. This is the file that must not become two hand-written copies. | | `src/vt/cuda/cuda_backend.cu` | `EndCaptureGraph` routes through the registry when enabled and retains the raw graph; `ReplayGraph` and `DestroyGraph` dispatch on `registry->Owns(handle)`. | | `src/vt/rocm/rocm_backend.hip` | The same three edits against the same shared header. | @@ -157,14 +162,16 @@ an assertion about intent. | Case | Guarantee | Upstream | |---|---|---| -| `folds compatible captures onto one executable` | K same-signature captures leave exactly 1 live exec; `Stats()` is `(K, 1)` | `:219-242`, `:253-256` | +| `folds compatible captures onto one executable` | K same-signature captures leave exactly 1 live exec; `(CapturedCount(), ExecCount())` is `(K, 1)` | `:219-242`, `:253-256` | | `keeps incompatible captures apart` | different signatures never share | `:225-241` | | `replays the graph the caller asked for` | over an interleaved multi-replay sequence, the deduped launch log equals the non-deduped control log, element for element | `:258-275` | | `does not re-update for a repeated shape` | a second consecutive replay of one handle issues no update | `:267` | | `falls back to a private executable when the driver rejects the update` | a probe failure yields its own exec, the existing group's live exec is untouched, and replay stays correct | diverges from `:229`, stated above | | `frees the shared executable only with its last graph` | destroying one sibling keeps the exec; destroying the last frees the exec and every retained raw graph, with all counters balanced | `:277-292` | | `does not claim a handle it did not mint` | `Owns()` is false for a foreign pointer | ours; the seam has no upstream twin | -| `is off unless the environment asks for it` | `GraphDedupEnabled()` polarity | `:321` | +| `is off unless the environment asks for it` | `GraphDedupEnabledFor()` polarity, including the values that only the terminator check rejects (`"10"`, `"11"`, `"1 "`, `"1x"`) | `:321` | +| `a capture the driver cannot instantiate fails at the capture site` | a null `instantiate` throws inside `Register`, mints no handle, counts nothing, and releases the capture it took ownership of | ours; added by the fresh review of #1178 | +| `a probe the driver cannot instantiate degrades instead of driving a null exec` | a failed probe answers "cannot fold" rather than passing a null executable to `cudaGraphExecUpdate` | ours; added by the fresh review of #1178 | ## Gates @@ -230,6 +237,46 @@ an assertion about intent. hit is probed with the real driver call before it is honoured. This is deliberate: a signature that had to be exhaustive would be a correctness surface, and it is instead a lookup key. +- **The fold that is probed is not always the fold that is replayed, and the gap is an + unasserted transitivity assumption.** `Register` probes the pair + `(group.raws.front(), raw_graph)`; `Replay` issues `(group.current_raw, entry.raw)`. + For a group of one or two members those coincide. From the third member onwards they + do not, so honouring the probe treats `cudaGraphExecUpdate` compatibility as + **transitive** across a group's members — if the driver re-points A onto B and A onto + C, then it re-points B onto C. Neither the CUDA nor the HIP documentation states this, + and nothing here asserts it. The fresh review of [#1178](https://github.com/mudler/vllm.cpp/pull/1178) + demonstrated it with a driver refusal keyed on the `(current, target)` pair: every + `Register` probe succeeded and the **second** `Replay` threw. The failure polarity is + what makes it survivable — a refusal lands on the `VT_CHECK` in `Replay`, loudly, and + never on a silent launch of the executable's previous contents. The stronger fix is to + probe `group.current_raw` instead of `raws.front()`, which removes the assumption + entirely; it is **not** taken here because it changes probe behaviour while the device + A/B (below) is measuring this exact commit. It is owed. +- **One registry serves every model on a device, and it is unsynchronised.** + `CudaBackend::dedup_` is a member of a process-singleton per-device backend + (`src/vt/cuda/cuda_backend.cu:355-365`), so the registry — and its two plain + `unordered_map`s — is shared by every model loaded onto that device. The constraint is + therefore wider than the "one runner thread" note above implies: a second model + capturing or destroying a graph concurrently with the first races the containers + themselves, before any question about a shared executable arises. Correct today + because the engine drives one device from one thread. A lock belongs at the registry + the moment that stops being true, and `ENG-CUDAGRAPH-BREAK` must not widen who + captures without adding one. +- **The signature builder has no executable coverage; it is compile-gated only.** The + 12 cases in `tests/vt/test_graph_dedup.cpp` drive `GraphDedupRegistry` through a fake + ops table, so they gate the registry's launch-sequence identity and nothing else. + `src/vt/graph_dedup_runtime.h` — the Kahn ordering, the topological re-index, the + sorted edge emission, the five node-payload cases, the depth-4 child-graph bound and + the five degradation escapes — is reached by no test on any tier; `cuda-fat-build` + proves only that it compiles. "12/12 negative mutations detected" is a statement about + `graph_dedup.h`, and must not be read as coverage of that file. The probe caps the + blast radius, so an unstable signature can never produce a wrong replay — but it has + exactly one silent mode: a signature that is unstable for some topology folds nothing, + every CPU test stays green, and the only observable is the device-side + `"captured N graphs, deduped to N execs"` log line. What would cover it: a CUDA-tier + test that captures the same trivial graph twice and asserts the two signatures are + byte-equal, and a second that captures two deliberately different topologies and + asserts they differ. Both need a device, so both are owed with the A/B. - **Two runtime-API calls changed shape across CUDA major versions, and only one was foreseen.** `cudaGraphExecUpdate` took the 3-argument `cudaGraphExecUpdateResultInfo` form at CUDA 12, so both it and the legacy 4-argument form are bound; HIP has only @@ -248,6 +295,8 @@ an assertion about intent. |---|---|---| | The device byte-identity A/B and the executable-count measurement (W4) | [#1162](https://github.com/mudler/vllm.cpp/issues/1162) | needs a leased CUDA box; the CPU tier proves the registry's launch-sequence identity, not the device's | | Flipping `VT_CUDA_GRAPH_DEDUP` on by default | [#1162](https://github.com/mudler/vllm.cpp/issues/1162) | gated on W4. A default is a measurement, not a preference | +| Probing `group.current_raw` rather than `raws.front()`, retiring the transitivity assumption above | [#1162](https://github.com/mudler/vllm.cpp/issues/1162) | it changes probe behaviour, and the device A/B is measuring the current one. Land it with the A/B rerun, not before | +| Executable coverage for `src/vt/graph_dedup_runtime.h`'s signature builder (stability and discrimination, both device-tier) | [#1162](https://github.com/mudler/vllm.cpp/issues/1162) | it needs a real `cudaGraph_t`, so it rides with the leased box the A/B already needs | | The ROCm leg's compile and run verification | [#41](https://github.com/mudler/vllm.cpp/issues/41) | no ROCm hardware or `hipcc` is reachable from this session and CI has no ROCm job, so the HIP wiring is written against the same shared header but is compile-unverified | ## Stop conditions diff --git a/src/vt/graph_dedup.h b/src/vt/graph_dedup.h index f816ed265..827573f59 100644 --- a/src/vt/graph_dedup.h +++ b/src/vt/graph_dedup.h @@ -5,8 +5,8 @@ // WHAT THIS REPLACES. Backend::EndCaptureGraph instantiates one graph executable per // capture and destroys the raw graph immediately (src/vt/cuda/cuda_backend.cu, and the // same shape on hipGraph). A model therefore holds one executable per padded decode -// bucket — 7 of them at max_num_seqs=32 and 11 at 64 — and eight hand-rolled drivers -// each hold their own set. The decode graphs of two padded batch sizes are usually the +// bucket — 7 of them at max_num_seqs=32 and 11 at 64 — and nine hand-rolled drivers +// each hold their own set (count corrected 2026-08-18, issue #1179). The decode graphs of two padded batch sizes are usually the // same node topology with different parameters, which is exactly the case // cudaGraphExecUpdate exists for: re-point ONE executable instead of instantiating a // second. @@ -20,7 +20,9 @@ // signature, but never trusts it: every candidate match is PROBED with the real driver // update before it is honoured, on a throwaway executable, and a probe the driver // refuses gives that capture its own executable. So a signature that is too coarse -// costs a wasted probe and nothing else, and a signature can never make a replay wrong. +// costs a wasted probe, or at worst the loud Replay-time VT_CHECK described there when +// update compatibility turns out not to be transitive across a group -- never a wrong +// replay. A signature can only cost a fold, it can never make a replay wrong. // The alternative — a signature exhaustive enough to be trusted — would make this file // a correctness surface that has to track every node-parameter class the driver knows. // @@ -39,6 +41,15 @@ // replay's update would re-point the executable out from under the first. Today's // drivers replay decode graphs from one runner thread on one queue, so the constraint // holds; it is stated here because it is a constraint the seam did not previously carry. +// +// The constraint is wider than "one runner thread", though, and the reason is where the +// registry LIVES: `CudaBackend::dedup_` belongs to a process-singleton per-device +// backend (src/vt/cuda/cuda_backend.cu:355-365), so ONE registry serves every model on +// a device. Its two `unordered_map`s carry no synchronisation, so a second model +// capturing or destroying a graph concurrently with the first is a data race on the +// containers themselves, before any question about a shared executable arises. Correct +// today because the engine drives one device from one thread; a lock belongs here the +// moment that stops being true. #ifndef VT_GRAPH_DEDUP_H_ #define VT_GRAPH_DEDUP_H_ @@ -125,6 +136,20 @@ class GraphDedupRegistry { if (group == nullptr) { auto created = std::make_unique(); created->exec = ops_.instantiate(raw_graph); + // FAIL HERE, at the capture site, exactly where the pre-dedup path did: + // EndCaptureGraph wrapped cudaGraphInstantiate in Check(), so an instantiate + // failure threw with the driver's error code while the caller was still inside + // the capture it belongs to. Accepting a null executable instead would mint a + // valid-looking handle over nothing, count it live in ExecCount(), and defer the + // failure to the first Replay as "graph launch failed" -- the wrong site, the + // wrong message, and the driver's reason already discarded. The capture's + // ownership transferred to us at the top of Register, so release it before + // unwinding rather than leaking it on the way out. + const bool instantiated = created->exec != nullptr; + if (!instantiated) ops_.destroy_graph(raw_graph); + VT_CHECK(instantiated, + "graph dedup: instantiate failed for the graph being registered at this " + "capture site; the capture cannot be replayed"); created->current_raw = raw_graph; group = created.get(); candidates.push_back(std::move(created)); @@ -156,9 +181,19 @@ class GraphDedupRegistry { if (group.current_raw != entry.raw) { std::string detail; const bool ok = ops_.update(group.exec, entry.raw, &detail); - // Register already probed this exact fold, so a refusal here is an invariant - // violation rather than an unsupported capture, and silently launching the - // executable's PREVIOUS contents would be the worst possible response. + // THIS IS NOT THE FOLD Register PROBED. Register probes (group.raws.front(), + // raw_graph); this update is (group.current_raw, entry.raw), and from the third + // member of a group onwards those pairs differ. Accepting the fold on the probe's + // word therefore treats cudaGraphExecUpdate compatibility as TRANSITIVE across a + // group's members -- if the driver can re-point A onto B and A onto C, then it can + // re-point B onto C. Neither the CUDA nor the HIP documentation says so, and this + // file does not assert it; it is an assumption, recorded as such under + // `## Risks/decisions` in .agents/specs/eng-cudagraph-dedup.md and owed there. + // The reason it is safe to hold it here anyway is the polarity of the failure: a + // refusal is either an invariant violation or that transitivity not holding, and + // both land on this VT_CHECK. What must never happen is the alternative -- + // launching the executable's PREVIOUS contents under this handle, which is a wrong + // answer rather than a loud one. VT_CHECK(ok, std::string("graph dedup: replay update refused (") + detail + ")"); group.current_raw = entry.raw; } @@ -237,6 +272,12 @@ class GraphDedupRegistry { bool ProbeAccepts(const Group& candidate, void* raw_graph) const { if (candidate.raws.empty()) return false; void* probe = ops_.instantiate(candidate.raws.front()); + // The probe instantiates a SECOND executable, so it can fail on its own -- most + // plausibly under the memory pressure this row exists to relieve. cudaGraphExecUpdate + // has no defined behaviour for a null executable, so answer "cannot fold" instead of + // asking the driver about nothing. Unlike the instantiate above this is not fatal: + // declining to fold degrades to exactly today's one-executable-per-capture path. + if (probe == nullptr) return false; std::string detail; const bool ok = ops_.update(probe, raw_graph, &detail); ops_.destroy_exec(probe); diff --git a/tests/vt/test_graph_dedup.cpp b/tests/vt/test_graph_dedup.cpp index cfb625976..bd66b104a 100644 --- a/tests/vt/test_graph_dedup.cpp +++ b/tests/vt/test_graph_dedup.cpp @@ -3,7 +3,7 @@ // // The production hazard is accelerator-only: today every capture instantiates its own // graph executable (src/vt/cuda/cuda_backend.cu EndCaptureGraph, and the same shape on -// hipGraph), so a model holds one executable per padded decode bucket and eight drivers +// hipGraph), so a model holds one executable per padded decode bucket and nine drivers // each hold their own set. The registry folds captures whose topology matches onto ONE // executable and re-points it with cudaGraphExecUpdate / hipGraphExecUpdate. // @@ -23,6 +23,7 @@ #include #include +#include #include #include @@ -58,6 +59,15 @@ struct FakeDevice { int graph_destroyed = 0; int updates = 0; std::set reject_update_for; // graph ids the driver refuses to update onto + // Graph ids the driver refuses to INSTANTIATE. cudaGraphInstantiate fails for reasons + // that have nothing to do with the topology -- out of memory is the common one on the + // unified-memory box this row exists for -- so it is a distinct outcome from a refused + // update and it has to have its own case. + std::set fail_instantiate_for; + // Times the registry handed the driver a NULL executable. cudaGraphExecUpdate has no + // defined behaviour for one; counting it here rather than dereferencing it keeps the + // red observable instead of a crash. + int null_exec_updates = 0; std::vector launch_log; // the id of the graph each launch actually ran // With `recycle` set, a destroyed graph's STORAGE is handed back to the test instead // of being freed, so a later capture can be made to land on the same address on @@ -76,8 +86,12 @@ std::string FakeSignature(void* graph) { } void* FakeInstantiate(void* graph) { - ++g_dev->instantiated; auto* source = static_cast(graph); + // Mirrors the driver: on failure nothing is allocated and the out-parameter is left + // null, which is what src/vt/graph_dedup_runtime.h Instantiate() turns into a null + // return. Counting only the successes keeps live_execs() meaningful. + if (g_dev->fail_instantiate_for.count(source->id) != 0) return nullptr; + ++g_dev->instantiated; return new FakeExec{source->sig, source->id}; } @@ -87,6 +101,11 @@ void* FakeInstantiate(void* graph) { // what reject_update_for models. bool FakeUpdate(void* exec_handle, void* graph_handle, std::string* detail) { ++g_dev->updates; + if (exec_handle == nullptr) { + ++g_dev->null_exec_updates; + if (detail != nullptr) *detail = "fake driver got a null executable"; + return false; + } auto* exec = static_cast(exec_handle); auto* graph = static_cast(graph_handle); if (g_dev->reject_update_for.count(graph->id) != 0 || exec->sig != graph->sig) { @@ -317,6 +336,62 @@ TEST_CASE("a capture reusing a freed graph address does not inherit its replay s registry.Close(); } +TEST_CASE("a capture the driver cannot instantiate fails at the capture site") { + DeviceScope scope; + scope.dev.fail_instantiate_for.insert(1); + auto registry = MakeRegistry(); + + // The pre-dedup path called Check(cudaGraphInstantiate(...), "cudaGraphInstantiate") + // inside EndCaptureGraph, so an instantiate failure threw where the capture happened. + // Accepting a null executable instead would mint a valid-looking handle over nothing, + // count it live, and defer the failure to the first replay -- a different site, a + // different message, and the driver's reason already thrown away. + CHECK_THROWS_AS(registry.Register(MakeGraph("decode", 1)), std::runtime_error); + + // And it must not be counted. A registry that reported a capture it can never replay + // would make the exec-count ratio this row is measured by a fiction. + CHECK(registry.CapturedCount() == 0); + CHECK(registry.ExecCount() == 0); + CHECK(scope.dev.live_execs() == 0); + // The capture's ownership had already transferred, so unwinding has to release it. + CHECK(scope.dev.graph_destroyed == 1); + + // The registry is still usable for a capture the driver will accept. + void* ok = registry.Register(MakeGraph("decode", 2)); + registry.Replay(ok, nullptr); + registry.Replay(ok, nullptr); + CHECK(registry.ExecCount() == 1); + CHECK(scope.dev.launch_log == std::vector{2, 2}); + + registry.Close(); +} + +TEST_CASE("a probe the driver cannot instantiate degrades instead of driving a null exec") { + DeviceScope scope; + auto registry = MakeRegistry(); + + void* first = registry.Register(MakeGraph("decode", 1)); + // The probe re-instantiates an EXISTING group member, so its instantiate can fail + // independently of the capture being registered -- most plausibly under the memory + // pressure that made this row worth doing at all. + scope.dev.fail_instantiate_for.insert(1); + void* second = registry.Register(MakeGraph("decode", 2)); + + // cudaGraphExecUpdate has no defined behaviour for a null executable, so a failed + // probe must answer "cannot fold" rather than ask the driver about nothing. + CHECK(scope.dev.null_exec_updates == 0); + CHECK(registry.CapturedCount() == 2); + CHECK(registry.ExecCount() == 2); + + for (int round = 0; round < 2; ++round) { + registry.Replay(first, nullptr); + registry.Replay(second, nullptr); + } + CHECK(scope.dev.launch_log == std::vector{1, 2, 1, 2}); + + registry.Close(); +} + TEST_CASE("the registry does not claim a handle it did not mint") { DeviceScope scope; auto registry = MakeRegistry(); @@ -371,4 +446,18 @@ TEST_CASE("dedup is off unless the environment asks for it") { CHECK_FALSE(vt::GraphDedupEnabledFor("")); CHECK_FALSE(vt::GraphDedupEnabledFor("0")); CHECK(vt::GraphDedupEnabledFor("1")); + + // Exactly "1", not "starts with 1". Without the terminator check these all enable + // dedup, and every case above still passes, because none of them ever asks about a + // value whose FIRST character is the one the polarity accepts. "10" is the one that + // matters in practice: it is what a hand-edited "1" plus a stray keystroke looks + // like, and turning a default-off correctness-sensitive path on by accident is the + // failure this polarity exists to prevent. + CHECK_FALSE(vt::GraphDedupEnabledFor("10")); + CHECK_FALSE(vt::GraphDedupEnabledFor("11")); + CHECK_FALSE(vt::GraphDedupEnabledFor("1 ")); + CHECK_FALSE(vt::GraphDedupEnabledFor("1x")); + CHECK_FALSE(vt::GraphDedupEnabledFor("01")); + CHECK_FALSE(vt::GraphDedupEnabledFor("true")); + CHECK_FALSE(vt::GraphDedupEnabledFor("on")); } From 5791c8b92648f6121a6b7b110f89f771ba46be9a Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Tue, 18 Aug 2026 04:04:08 +0000 Subject: [PATCH 7/9] style(ENG-CUDAGRAPH-DEDUP): reflow the header paragraph the driver-count correction broke The nine-driver correction was spliced into a wrapped comment paragraph and left one 129-column line in a file wrapped at 88. Reflowed the paragraph; no other change. The signature note also said "described there" of a call site 150 lines below it and then restated its own conclusion, so it now points at the call site and says the thing once. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code] --- src/vt/graph_dedup.h | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/src/vt/graph_dedup.h b/src/vt/graph_dedup.h index 827573f59..c6776c3a0 100644 --- a/src/vt/graph_dedup.h +++ b/src/vt/graph_dedup.h @@ -5,11 +5,11 @@ // WHAT THIS REPLACES. Backend::EndCaptureGraph instantiates one graph executable per // capture and destroys the raw graph immediately (src/vt/cuda/cuda_backend.cu, and the // same shape on hipGraph). A model therefore holds one executable per padded decode -// bucket — 7 of them at max_num_seqs=32 and 11 at 64 — and nine hand-rolled drivers -// each hold their own set (count corrected 2026-08-18, issue #1179). The decode graphs of two padded batch sizes are usually the -// same node topology with different parameters, which is exactly the case -// cudaGraphExecUpdate exists for: re-point ONE executable instead of instantiating a -// second. +// bucket — 7 of them at max_num_seqs=32 and 11 at 64 — and nine hand-rolled drivers each +// hold their own set (count corrected 2026-08-18, issue #1179). The decode graphs of two +// padded batch sizes are usually the same node topology with different parameters, which +// is exactly the case cudaGraphExecUpdate exists for: re-point ONE executable instead of +// instantiating a second. // // This is a MEMORY and CAPTURE-TIME change, not a throughput change. A deduped replay // launches the same nodes with the same parameters; if it launched anything else that @@ -20,9 +20,9 @@ // signature, but never trusts it: every candidate match is PROBED with the real driver // update before it is honoured, on a throwaway executable, and a probe the driver // refuses gives that capture its own executable. So a signature that is too coarse -// costs a wasted probe, or at worst the loud Replay-time VT_CHECK described there when -// update compatibility turns out not to be transitive across a group -- never a wrong -// replay. A signature can only cost a fold, it can never make a replay wrong. +// costs a wasted probe, or at worst the loud Replay-time VT_CHECK explained at that +// call site when update compatibility turns out not to be transitive across a group. +// A signature can cost a fold; it can never make a replay wrong. // The alternative — a signature exhaustive enough to be trusted — would make this file // a correctness surface that has to track every node-parameter class the driver knows. // From 3ed6253c6b94ce7e5b77424f7c8008601f5fb40b Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Tue, 18 Aug 2026 04:12:25 +0000 Subject: [PATCH 8/9] test(ENG-CUDAGRAPH-DEDUP): gate the loud-failure claim the transitivity note now rests on The previous commit corrected the Replay comment to say that honouring Register's probe assumes cudaGraphExecUpdate compatibility is transitive across a group's members, and argued the assumption is survivable because a refusal lands on a VT_CHECK rather than on a silent launch of the executable's previous contents. Nothing gated that argument, and an argument for why a defect is tolerable is exactly the kind of claim that has to be executable. The fake driver can now refuse a (currently-reflected, target) PAIR, which is the shape the target-only rejection set cannot express and the shape that matters. Refusing exactly the pair Register never asks about reproduces the case: all three probes succeed, the group folds to one executable, and the second replay is where the assumption is tested for real. The case asserts the throw AND that the launch log is unchanged. Mutation, detected: replacing the VT_CHECK with `(void)ok` makes the replay launch graph 2's nodes under graph 3's handle, which is the silent wrong answer the check exists to prevent, and fails three assertions. Non-empty `git diff --stat`, clean compile, byte-for-byte restore. 13/13 cases, 65 assertions. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code] --- .agents/engine-matrix.md | 2 +- .agents/specs/eng-cudagraph-dedup.md | 9 ++++-- tests/vt/test_graph_dedup.cpp | 43 +++++++++++++++++++++++++++- 3 files changed, 49 insertions(+), 5 deletions(-) diff --git a/.agents/engine-matrix.md b/.agents/engine-matrix.md index 74d5b88e8..678b635fe 100644 --- a/.agents/engine-matrix.md +++ b/.agents/engine-matrix.md @@ -60,7 +60,7 @@ forensics: roadmap_v1.md and the parity ledger. | `KV-PREFIX-MATCH-UNIT` | `--prefix-match-unit` (config `prefix_match_unit`): the finest token boundary a prefix-cache hit can land on == the `hash_block_size`/"prefix match unit" the block hasher uses. NEW in 0.26 (absent at the prior `e24d1b24`/0.25.0 pin). For a HYBRID/multi-group model the resolver `resolve_kv_cache_block_sizes` computes `hash_block_size = prefix_match_unit if set else gcd(group_block_sizes)` (scheduler block size = `lcm`), letting matching land FINER than a physical block (e.g. 16/32 tokens inside a 1024-token block) provided every group block size is divisible by it; single-group (dense) models ignore the knob. Backs off to the scheduler block size when no prefix-cache/connector consumer is active or a mamba group diverges from `cache_block_size` (mamba_cache_mode != "align"); throws on a non-divisible unit. **W0 spike + W1 resolver LANDED 2026-07-28 (`CLAIM-PREFIX-MATCH-UNIT`, NOT pushed):** `resolve_kv_cache_block_sizes` ported 1:1 (explicit-parameter signature vs upstream's `VllmConfig`, our config surface is threaded), RED-first unit-gated (default gcd `!=` `=16` override). `PARTIAL`: the config/CLI/ABI field (W2), the scheduler threading of a resolved `hash_block_size != block_size` + mamba partial-tail stop (W3, needs the `KV-BLOCK-POOL` align path that still throws), and the benchmark (W4) are deferred. Default path byte-identical (single-group inert; scheduler still passes `block_size`). | T1 | `vllm/engine/arg_utils.py:696,1222,1940`; `vllm/config/cache.py:56-67`; resolver `vllm/v1/core/kv_cache_utils.py:626-688`; hasher `:691-748`; call site `vllm/v1/engine/core.py:154`; scheduler `vllm/v1/core/sched/scheduler.py:76,268-270,282,312-318`; fine-grained view `vllm/v1/core/single_type_kv_cache_manager.py:683,697` | resolver `src/vllm/v1/core/kv_cache_utils.cpp:638` (`resolve_kv_cache_block_sizes`), decl `include/vllm/v1/core/kv_cache_utils.h`; hash_block_size already plumbed `get_request_block_hasher` `src/vllm/v1/core/kv_cache_utils.cpp:577`; DEFERRED align path throws `src/vllm/v1/core/block_pool.cpp:93,220` (shared with `KV-BLOCK-POOL`) | `tests/vllm/v1/test_prefix_match_unit.cpp:64,88,99,119,129,145,164,186` 8/8 (29 assertions): single-group inert + DCP scale, multi-group default=gcd, `=16` override finer-than-default (RED), finer-than-1024-block, non-divisible throws, no-consumer back-off + connector-alone re-enable, mamba non-align back-off vs align gcd, hasher-granularity RED (coarse 2 vs fine 4 hashes); [parity-ledger.md](parity-ledger.md) | [prefix-match-unit.md](specs/prefix-match-unit.md) | `PARTIAL` | `CLAIM-PREFIX-MATCH-UNIT` | | `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` 12/12 cases, 61 assertions, RED-first (written and run against an absent header, and the three cases added by the fresh review of #1178 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; 12/12 negative mutations detected (9 at implementation, 3 at review repair). That count covers `src/vt/graph_dedup.h` ONLY: `src/vt/graph_dedup_runtime.h`, which builds the signature, has NO executable coverage on any tier and is compile-gated by `.github/workflows/ci.yml:669` alone — its one silent mode is a signature unstable for some topology, which folds nothing while every CPU test stays green. STILL OWED: the device same-binary A/B proving a deduped replay token-identical, the exec-count ratio, device-tier signature stability/discrimination tests, probing `current_raw` instead of `raws.front()` to retire the update-transitivity assumption, and the ROCm compile | [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-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`, which builds the signature, has NO executable coverage on any tier and is compile-gated by `.github/workflows/ci.yml:669` alone — its one silent mode is a signature unstable for some topology, which folds nothing while every CPU test stays green. STILL OWED: the device same-binary A/B proving a deduped replay token-identical, the exec-count ratio, device-tier signature stability/discrimination tests, probing `current_raw` instead of `raws.front()` to retire the update-transitivity assumption, and the ROCm compile | [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:961-986` DECLINES the graph outright when the async device-token mirror is live. **That decline is why this is also a CORRECTNESS row** ([#1179](https://github.com/mudler/vllm.cpp/issues/1179)): a SHIPPED model has already lost its decode graph to the duplication, on the driver's own measurement (`depth-1, graph ON PASS 78/78`; `depth-2, graph OFF PASS 82/82`; `depth-2, graph ON FAIL, slots 1-3 degenerate`), and the fix its comment names is the sibling's `StepDevInputs`. The row still makes NO throughput claim: the prefill refutation on the `ENG-CUDAGRAPH` row (3.8% host idle, >96% GPU-busy, 92.5% glue) stands unchanged | owed: reachability mutation (delete the production call site, rerun the focused gate); bit-exactness vs eager on every migrated model over MORE than one replay; the host-lifetime contract of [decode-graph-scratch-uaf-2026-07-18.md](specs/decode-graph-scratch-uaf-2026-07-18.md) enforced AT the seam; the 11 ported SGLang unit cases; the break-function OUTPUT writeback (`replay_fn`/`_copy_output` `breakable_cuda_graph.py:231-235,172-201`, spec D9) without which a replay leaves the next segment reading capture-time data; the auxiliary-stream auto-join before every segment close (`:353-361`, spec D10), live at `src/vllm/model_executor/models/qwen3_5.cpp:6254-6255,6384` and `src/vllm/model_executor/models/laguna.cpp:2572-2576,2612` | spec [eng-cudagraph-break.md](specs/eng-cudagraph-break.md) (W0 spike DONE 2026-08-18: the existing `vt` capture vocabulary `include/vt/backend.h:208-222` expresses a SEGMENTED capture with NO new virtual, because `EndCaptureGraph` stores nothing (`src/vt/cuda/cuda_backend.cu:225-232`); a break point is expressible with one `thread_local` capture pointer plus a free function, no compiler and no decorator); analysis [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `READY` | [#1163](https://github.com/mudler/vllm.cpp/issues/1163), [#1020](https://github.com/mudler/vllm.cpp/issues/1020) | | `ENG-CUDAGRAPH-DIFFUSION` | Capture the LTX-2.5 denoise loop (fixed shapes, many identical iterations — the ideal graph target). **BLOCKED, and the blocker is ours:** the render does almost no device compute to capture | T2 | SGLang enabled BCG on this shape AFTER our pin — LTX-2 H200 two-stage 10.75s->6.90s (`d4be483efb`), SANA 1024px -26% (`6c7498113f`), SANA denoise 0.73->0.457s (`56ef810cad`). Dated events, NOT pinned evidence; their win is mostly PyTorch host tax we do not pay | NO capture at all: `grep` for capture across `src/vllm/model_executor/models/ltx2*.cpp` returns nothing | blocked by [#1024](https://github.com/mudler/vllm.cpp/issues/1024) (GPU util **exactly 0 in 321 of 347 samples**, 1.00 core of 20 held for 17+ min after staging), [#1007](https://github.com/mudler/vllm.cpp/issues/1007) (VAE decode has no device arm), [#1087](https://github.com/mudler/vllm.cpp/issues/1087) (**57-66% of wall** is ONE resolution-CONSTANT serial host phase), [#1010](https://github.com/mudler/vllm.cpp/issues/1010) (no phase-boundary log). Decision point is a MEASUREMENT of GPU-busy vs wall once device-resident, not an implementation | [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `INVENTORIED` | [#1164](https://github.com/mudler/vllm.cpp/issues/1164) | | `ENG-BATCH-INVARIANT` | Opt-in deterministic execution across scheduler batch sizes (`VLLM_BATCH_INVARIANT=1`): batch-invariant matmul/norm/attention/collectives plus persistent-scheduler NVFP4; production default remains off | T1 | default/env `vllm/envs.py:89,576-578`; initialization `vllm/v1/worker/gpu_worker.py:1262`; NVFP4 dispatch `csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_sm120_kernels.cu:212-220`; suite fixture `tests/v1/determinism/conftest.py:9-12`; operator/e2e `tests/v1/determinism/test_nvfp4_batch_invariant_scaled_mm.py`, `tests/v1/determinism/test_nvfp4_batch_invariant.py` @ `702f481` | - | [W3-C3R executed contract](specs/nvfp4-persistent-plan-cache.md#w3-c3r-batch-shape-localization-and-gate-correction-2026-07-13): production-default ours and vLLM both change outputs across batch shapes; no local opt-in implementation is claimed | `planned: specs/batch-invariant-execution.md` | `INVENTORIED` | - | diff --git a/.agents/specs/eng-cudagraph-dedup.md b/.agents/specs/eng-cudagraph-dedup.md index 7dd7165b9..164522252 100644 --- a/.agents/specs/eng-cudagraph-dedup.md +++ b/.agents/specs/eng-cudagraph-dedup.md @@ -172,6 +172,7 @@ an assertion about intent. | `is off unless the environment asks for it` | `GraphDedupEnabledFor()` polarity, including the values that only the terminator check rejects (`"10"`, `"11"`, `"1 "`, `"1x"`) | `:321` | | `a capture the driver cannot instantiate fails at the capture site` | a null `instantiate` throws inside `Register`, mints no handle, counts nothing, and releases the capture it took ownership of | ours; added by the fresh review of #1178 | | `a probe the driver cannot instantiate degrades instead of driving a null exec` | a failed probe answers "cannot fold" rather than passing a null executable to `cudaGraphExecUpdate` | ours; added by the fresh review of #1178 | +| `a replay update the driver refuses fails loudly rather than launching stale nodes` | a refusal on the fold `Register` never probed throws, and the executable's previous contents are NOT launched under the asking handle | ours; pins the safety claim that makes the transitivity assumption below survivable | ## Gates @@ -248,7 +249,9 @@ an assertion about intent. demonstrated it with a driver refusal keyed on the `(current, target)` pair: every `Register` probe succeeded and the **second** `Replay` threw. The failure polarity is what makes it survivable — a refusal lands on the `VT_CHECK` in `Replay`, loudly, and - never on a silent launch of the executable's previous contents. The stronger fix is to + never on a silent launch of the executable's previous contents — which the case + `a replay update the driver refuses fails loudly rather than launching stale nodes` + now gates, by refusing exactly the pair `Register` never asks about. The stronger fix is to probe `group.current_raw` instead of `raws.front()`, which removes the assumption entirely; it is **not** taken here because it changes probe behaviour while the device A/B (below) is measuring this exact commit. It is owed. @@ -263,12 +266,12 @@ an assertion about intent. the moment that stops being true, and `ENG-CUDAGRAPH-BREAK` must not widen who captures without adding one. - **The signature builder has no executable coverage; it is compile-gated only.** The - 12 cases in `tests/vt/test_graph_dedup.cpp` drive `GraphDedupRegistry` through a fake + 13 cases in `tests/vt/test_graph_dedup.cpp` drive `GraphDedupRegistry` through a fake ops table, so they gate the registry's launch-sequence identity and nothing else. `src/vt/graph_dedup_runtime.h` — the Kahn ordering, the topological re-index, the sorted edge emission, the five node-payload cases, the depth-4 child-graph bound and the five degradation escapes — is reached by no test on any tier; `cuda-fat-build` - proves only that it compiles. "12/12 negative mutations detected" is a statement about + proves only that it compiles. "13/13 negative mutations detected" is a statement about `graph_dedup.h`, and must not be read as coverage of that file. The probe caps the blast radius, so an unstable signature can never produce a wrong replay — but it has exactly one silent mode: a signature that is unstable for some topology folds nothing, diff --git a/tests/vt/test_graph_dedup.cpp b/tests/vt/test_graph_dedup.cpp index bd66b104a..e7b8f1d42 100644 --- a/tests/vt/test_graph_dedup.cpp +++ b/tests/vt/test_graph_dedup.cpp @@ -24,6 +24,7 @@ #include #include #include +#include #include #include @@ -59,6 +60,11 @@ struct FakeDevice { int graph_destroyed = 0; int updates = 0; std::set reject_update_for; // graph ids the driver refuses to update onto + // (currently-reflected graph id, target graph id) pairs the driver refuses. This is + // the shape reject_update_for cannot express and the shape that matters: whether an + // update is accepted can depend on WHICH graph the executable currently holds, not + // only on the one it is being pointed at. + std::set> reject_update_pair; // Graph ids the driver refuses to INSTANTIATE. cudaGraphInstantiate fails for reasons // that have nothing to do with the topology -- out of memory is the common one on the // unified-memory box this row exists for -- so it is a distinct outcome from a refused @@ -108,7 +114,9 @@ bool FakeUpdate(void* exec_handle, void* graph_handle, std::string* detail) { } auto* exec = static_cast(exec_handle); auto* graph = static_cast(graph_handle); - if (g_dev->reject_update_for.count(graph->id) != 0 || exec->sig != graph->sig) { + if (g_dev->reject_update_for.count(graph->id) != 0 || + g_dev->reject_update_pair.count({exec->id, graph->id}) != 0 || + exec->sig != graph->sig) { if (detail != nullptr) *detail = "fake driver refused the update"; return false; } @@ -392,6 +400,39 @@ TEST_CASE("a probe the driver cannot instantiate degrades instead of driving a n registry.Close(); } +TEST_CASE("a replay update the driver refuses fails loudly rather than launching stale nodes") { + DeviceScope scope; + // Register probes (group.raws.front(), candidate) but Replay issues + // (group.current_raw, target), and from the third group member onwards those pairs + // differ -- so honouring the probe assumes update compatibility is TRANSITIVE across a + // group. Nothing asserts that. Refusing exactly the pair Register never asks about + // reproduces the case: every probe succeeds, the group folds to one executable, and + // the SECOND replay is where the assumption is tested for real. + scope.dev.reject_update_pair.insert({2, 3}); + auto registry = MakeRegistry(); + + void* g1 = registry.Register(MakeGraph("decode", 1)); + void* g2 = registry.Register(MakeGraph("decode", 2)); + void* g3 = registry.Register(MakeGraph("decode", 3)); + // Every probe was (graph 1, candidate), and the driver accepts those. + REQUIRE(registry.ExecCount() == 1); + + registry.Replay(g2, nullptr); // (1 -> 2), accepted; the executable now reflects 2 + + // And now the fold nobody probed. What must NOT happen is the silent alternative: + // leaving the executable pointing at graph 2 and launching it under g3's handle, which + // is a wrong answer rather than a loud one. This VT_CHECK is the entire reason the + // transitivity assumption is survivable, so it is gated rather than asserted in prose. + CHECK_THROWS_AS(registry.Replay(g3, nullptr), std::runtime_error); + CHECK(scope.dev.launch_log == std::vector{2}); + + // g1 is unaffected: its own fold is one the driver still accepts. + registry.Replay(g1, nullptr); + CHECK(scope.dev.launch_log == std::vector{2, 1}); + + registry.Close(); +} + TEST_CASE("the registry does not claim a handle it did not mint") { DeviceScope scope; auto registry = MakeRegistry(); From 72de552c8107694a1ab88eaf380088ae19f7544d Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Tue, 18 Aug 2026 07:54:55 +0000 Subject: [PATCH 9/9] fix(ENG-CUDAGRAPH-DEDUP): a file allowed to see the runtime FAIL must consume the latched error, and this one never did (#1184, #1162) `VT_CUDA_GRAPH_DEDUP=1` could not complete one decode step once a graph was actually captured. 6/6 deterministic on GB10: vt graph dedup: captured 1 graphs, deduped to 1 execs [Qwen3DenseDecodeGraph] captured dense decode graph for padded size S=8 engine-fatal: EngineCore busy loop threw: vt cuda: greedy_argmax launch: invalid device function from a `greedy_argmax` launch that had SUCCEEDED. The OFF and `=0` arms were clean and byte-identical across 7 runs on the same binary and libraries, which is what makes the asymmetry arm-attributable. `greedy_argmax` was never at fault. The whole safety argument for this row is that the signature is a lookup key and `cudaGraphExecUpdate` is the authority, so a probe the driver REFUSES is normal operation, not an exception; the topology walk has five more escapes of the same kind. A CUDA call that fails also latches its code in the runtime's sticky per-thread slot, and none of the twelve fallible calls in `graph_dedup_runtime.h` consumed it -- `grep -n cudaGetLastError` on that file returned nothing. The next unrelated kernel, launched with the ordinary `kernel<<<>>>(); Check(cudaGetLastError())` pattern, then read our routine refusal and reported it as its own failure. Every symptom follows from that and none from the kernel named: it needs BOTH `dedup=1` and a real capture, `CUDA_LAUNCH_BLOCKING=1` does not move it because the latch is host-side and synchronous rather than a deferred async error, and `cudaGraphLaunch` returns success because a return value does not consume the latch. The fix is one type, not twelve clears. Twelve hand-placed `cudaGetLastError()` calls are a fix the thirteenth fallible call silently misses, and a file whose design is "these calls are allowed to fail" will grow a thirteenth. The clear lives in `ScopedLatchClear`'s destructor and is installed at the binding's ENTRY POINTS, which are exactly the six `GraphDedupOps` members the backends reach through `Ops()`. Every exit runs it: a plain return, a degradation escape, and the `VT_CHECK` unwinding out of `Launch`. `MakeLatchGuardedOps` is the table's only constructor and takes the raw functions as template arguments, so no raw address reaches a field; a seventh operation wired anywhere else leaves its field null and `GraphDedupRegistry` refuses an incomplete table. One line covers both arms because there is one source: `VTGD_FN(GetLastError)` resolves to `cudaGetLastError` or `hipGetLastError`. The fold and probe DECISION logic is untouched. This is error-state hygiene. MEDIUM-3, the coverage gap that hid this for a whole review cycle, closes in the same change. `graph_dedup_runtime.h` was reached by no test on any tier, so its Kahn ordering, topological re-index, sorted edge emission, depth-4 child bound and degradation escapes were compile-gated only. The device-free half moves to `src/vt/graph_dedup_signature.h` behind an `Rt` policy and is driven by a fake runtime in the new `tests/vt/test_graph_dedup_runtime.cpp`: 13 cases, 51 assertions, RED-first against the pre-fix guard, where the suite reports 22 failed assertions including the production message reproduced from the mechanism alone. 7/7 negative mutations detected, each with its `git diff --stat`, compile status and exit status; one first attempt failed to build under `-Werror` and was re-run in a compiling form, because a mutation that fails to build reads as a passing test. BE HONEST ABOUT THE LIMIT. A CPU test drives a FAKE runtime. It cannot observe the CUDA runtime's real latched-error state, so it proves the guard's structure and not that #1184 is gone on a device. The device A/B re-run stays owed under #1184, and the five node-payload cases stay device-only. The spec also now records what the device gate found: the shipped async serving path captures no decode graph at all (the #323 mitigation in `DenseDecodeGraphForward`), so dedup engages only under `VT_ASYNC_RUNNER=0`, which materially bounds this row's worth and is owned by #1179. No local CUDA toolkit was reachable. A scratch header-shape stub instantiated `Ops()` at CUDART_VERSION 12090 and 13030 and on the HIP arm under `-Werror`, all three clean, with a negative control proving the instrument can fail. That is a proxy; `cuda-fat-build` remains the gate, exactly as it was for the CUDA 13 `cudaGraphGetEdges` break this file already took. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code] --- .agents/engine-matrix.md | 2 +- .agents/issue-index.md | 1 + .agents/specs/eng-cudagraph-dedup.md | 140 ++++++-- src/vt/graph_dedup_latch.h | 120 +++++++ src/vt/graph_dedup_runtime.h | 255 ++++++-------- src/vt/graph_dedup_signature.h | 162 +++++++++ tests/CMakeLists.txt | 6 + tests/vt/test_graph_dedup_runtime.cpp | 479 ++++++++++++++++++++++++++ 8 files changed, 987 insertions(+), 178 deletions(-) create mode 100644 src/vt/graph_dedup_latch.h create mode 100644 src/vt/graph_dedup_signature.h create mode 100644 tests/vt/test_graph_dedup_runtime.cpp diff --git a/.agents/engine-matrix.md b/.agents/engine-matrix.md index fb89daafc..db2ba959a 100644 --- a/.agents/engine-matrix.md +++ b/.agents/engine-matrix.md @@ -60,7 +60,7 @@ forensics: roadmap_v1.md and the parity ledger. | `KV-PREFIX-MATCH-UNIT` | `--prefix-match-unit` (config `prefix_match_unit`): the finest token boundary a prefix-cache hit can land on == the `hash_block_size`/"prefix match unit" the block hasher uses. NEW in 0.26 (absent at the prior `e24d1b24`/0.25.0 pin). For a HYBRID/multi-group model the resolver `resolve_kv_cache_block_sizes` computes `hash_block_size = prefix_match_unit if set else gcd(group_block_sizes)` (scheduler block size = `lcm`), letting matching land FINER than a physical block (e.g. 16/32 tokens inside a 1024-token block) provided every group block size is divisible by it; single-group (dense) models ignore the knob. Backs off to the scheduler block size when no prefix-cache/connector consumer is active or a mamba group diverges from `cache_block_size` (mamba_cache_mode != "align"); throws on a non-divisible unit. **W0 spike + W1 resolver LANDED 2026-07-28 (`CLAIM-PREFIX-MATCH-UNIT`, NOT pushed):** `resolve_kv_cache_block_sizes` ported 1:1 (explicit-parameter signature vs upstream's `VllmConfig`, our config surface is threaded), RED-first unit-gated (default gcd `!=` `=16` override). `PARTIAL`: the config/CLI/ABI field (W2), the scheduler threading of a resolved `hash_block_size != block_size` + mamba partial-tail stop (W3, needs the `KV-BLOCK-POOL` align path that still throws), and the benchmark (W4) are deferred. Default path byte-identical (single-group inert; scheduler still passes `block_size`). | T1 | `vllm/engine/arg_utils.py:696,1222,1940`; `vllm/config/cache.py:56-67`; resolver `vllm/v1/core/kv_cache_utils.py:626-688`; hasher `:691-748`; call site `vllm/v1/engine/core.py:154`; scheduler `vllm/v1/core/sched/scheduler.py:76,268-270,282,312-318`; fine-grained view `vllm/v1/core/single_type_kv_cache_manager.py:683,697` | resolver `src/vllm/v1/core/kv_cache_utils.cpp:638` (`resolve_kv_cache_block_sizes`), decl `include/vllm/v1/core/kv_cache_utils.h`; hash_block_size already plumbed `get_request_block_hasher` `src/vllm/v1/core/kv_cache_utils.cpp:577`; DEFERRED align path throws `src/vllm/v1/core/block_pool.cpp:93,220` (shared with `KV-BLOCK-POOL`) | `tests/vllm/v1/test_prefix_match_unit.cpp:64,88,99,119,129,145,164,186` 8/8 (29 assertions): single-group inert + DCP scale, multi-group default=gcd, `=16` override finer-than-default (RED), finer-than-1024-block, non-divisible throws, no-consumer back-off + connector-alone re-enable, mamba non-align back-off vs align gcd, hasher-granularity RED (coarse 2 vs fine 4 hashes); [parity-ledger.md](parity-ledger.md) | [prefix-match-unit.md](specs/prefix-match-unit.md) | `PARTIAL` | `CLAIM-PREFIX-MATCH-UNIT` | | `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`, which builds the signature, has NO executable coverage on any tier and is compile-gated by `.github/workflows/ci.yml:669` alone — its one silent mode is a signature unstable for some topology, which folds nothing while every CPU test stays green. STILL OWED: the device same-binary A/B proving a deduped replay token-identical, the exec-count ratio, device-tier signature stability/discrimination tests, probing `current_raw` instead of `raws.front()` to retire the update-transitivity assumption, and the ROCm compile | [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-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. STILL OWED: the device same-binary A/B proving a deduped replay token-identical AND that #1184 is gone on a device, the exec-count ratio, device-tier signature stability/discrimination tests, probing `current_raw` instead of `raws.front()` to retire the update-transitivity assumption, the ROCm compile, and reaching the feature from the default serving path at all — the async runner captures no decode graph, so dedup engages only under `VT_ASYNC_RUNNER=0` | [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:961-986` DECLINES the graph outright when the async device-token mirror is live. **That decline is why this is also a CORRECTNESS row** ([#1179](https://github.com/mudler/vllm.cpp/issues/1179)): a SHIPPED model has already lost its decode graph to the duplication, on the driver's own measurement (`depth-1, graph ON PASS 78/78`; `depth-2, graph OFF PASS 82/82`; `depth-2, graph ON FAIL, slots 1-3 degenerate`), and the fix its comment names is the sibling's `StepDevInputs`. The row still makes NO throughput claim: the prefill refutation on the `ENG-CUDAGRAPH` row (3.8% host idle, >96% GPU-busy, 92.5% glue) stands unchanged | owed: reachability mutation (delete the production call site, rerun the focused gate); bit-exactness vs eager on every migrated model over MORE than one replay; the host-lifetime contract of [decode-graph-scratch-uaf-2026-07-18.md](specs/decode-graph-scratch-uaf-2026-07-18.md) enforced AT the seam; the 11 ported SGLang unit cases; the break-function OUTPUT writeback (`replay_fn`/`_copy_output` `breakable_cuda_graph.py:231-235,172-201`, spec D9) without which a replay leaves the next segment reading capture-time data; the auxiliary-stream auto-join before every segment close (`:353-361`, spec D10), live at `src/vllm/model_executor/models/qwen3_5.cpp:6254-6255,6384` and `src/vllm/model_executor/models/laguna.cpp:2572-2576,2612` | spec [eng-cudagraph-break.md](specs/eng-cudagraph-break.md) (W0 spike DONE 2026-08-18: the existing `vt` capture vocabulary `include/vt/backend.h:208-222` expresses a SEGMENTED capture with NO new virtual, because `EndCaptureGraph` stores nothing (`src/vt/cuda/cuda_backend.cu:225-232`); a break point is expressible with one `thread_local` capture pointer plus a free function, no compiler and no decorator); analysis [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `READY` | [#1163](https://github.com/mudler/vllm.cpp/issues/1163), [#1020](https://github.com/mudler/vllm.cpp/issues/1020) | | `ENG-CUDAGRAPH-DIFFUSION` | Capture the LTX-2.5 denoise loop (fixed shapes, many identical iterations — the ideal graph target). **BLOCKED, and the blocker is ours:** the render does almost no device compute to capture | T2 | SGLang enabled BCG on this shape AFTER our pin — LTX-2 H200 two-stage 10.75s->6.90s (`d4be483efb`), SANA 1024px -26% (`6c7498113f`), SANA denoise 0.73->0.457s (`56ef810cad`). Dated events, NOT pinned evidence; their win is mostly PyTorch host tax we do not pay | NO capture at all: `grep` for capture across `src/vllm/model_executor/models/ltx2*.cpp` returns nothing | blocked by [#1024](https://github.com/mudler/vllm.cpp/issues/1024) (GPU util **exactly 0 in 321 of 347 samples**, 1.00 core of 20 held for 17+ min after staging), [#1007](https://github.com/mudler/vllm.cpp/issues/1007) (VAE decode has no device arm), [#1087](https://github.com/mudler/vllm.cpp/issues/1087) (**57-66% of wall** is ONE resolution-CONSTANT serial host phase), [#1010](https://github.com/mudler/vllm.cpp/issues/1010) (no phase-boundary log). Decision point is a MEASUREMENT of GPU-busy vs wall once device-resident, not an implementation | [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `INVENTORIED` | [#1164](https://github.com/mudler/vllm.cpp/issues/1164) | | `ENG-BATCH-INVARIANT` | Opt-in deterministic execution across scheduler batch sizes (`VLLM_BATCH_INVARIANT=1`): batch-invariant matmul/norm/attention/collectives plus persistent-scheduler NVFP4; production default remains off | T1 | default/env `vllm/envs.py:89,576-578`; initialization `vllm/v1/worker/gpu_worker.py:1262`; NVFP4 dispatch `csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_sm120_kernels.cu:212-220`; suite fixture `tests/v1/determinism/conftest.py:9-12`; operator/e2e `tests/v1/determinism/test_nvfp4_batch_invariant_scaled_mm.py`, `tests/v1/determinism/test_nvfp4_batch_invariant.py` @ `702f481` | - | [W3-C3R executed contract](specs/nvfp4-persistent-plan-cache.md#w3-c3r-batch-shape-localization-and-gate-correction-2026-07-13): production-default ours and vLLM both change outputs across batch shapes; no local opt-in implementation is claimed | `planned: specs/batch-invariant-execution.md` | `INVENTORIED` | - | diff --git a/.agents/issue-index.md b/.agents/issue-index.md index 83876d612..2153caedc 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -366,3 +366,4 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1179](https://github.com/mudler/vllm.cpp/issues/1179) | `ENG-CUDAGRAPH-BREAK` | The hand-rolled decode-graph driver count recorded in `9bc4d7f44` is **eight** and is actually **nine**, and the row it feeds was framed as coverage-only when it is also correctness. The ninth is the DFlash draft graph, file-local with no header declaration, at `src/vllm/model_executor/models/qwen3_dflash.cpp:771,870,1038,1091,1095,1106` — its own `int g_state = 0` three-state machine (`:771`), its own `VT_DFLASH_GRAPH` kill switch (`:870`) instead of the `VLLM_CPP_CUDAGRAPH` the six batched drivers read, its own invalidate-on-block-width-change (`:1038-1047`) and its own `try { EndCaptureGraph(); } catch (...) {}` drain (`:1106`). The eight-count is stated in four places, all corrected here: [`sglang-breakable-cuda-graph.md`](specs/sglang-breakable-cuda-graph.md) §4 and `## Owed`, [`.agents/engine-matrix.md`](engine-matrix.md) rows `ENG-CUDAGRAPH-BREAK` and `ENG-CUDAGRAPH-DEDUP` ("times eight drivers", which sizes #1162's signature table), and [`.agents/roadmap_v1.md`](roadmap_v1.md) track `C12`. The reframing is the substantive half: `ENG-CUDAGRAPH-BREAK` was recorded as a COVERAGE row, and the duplication has already cost a SHIPPED model its decode graph. `src/vllm/model_executor/models/qwen3.cpp:961-986` declines the decode graph outright whenever the asynchronous device-token mirror is live, on its own measured battery — `depth-1, graph ON PASS 78/78`; `depth-2, graph OFF PASS 82/82`; `depth-2, graph ON FAIL, slots 1-3 degenerate` — because `Step()` replays against the HOST `input.token_ids` and the combine has patched the DEVICE ids. The comment names the real fix as reading the identifiers at replay time from a stable device buffer, and that fix exists, in exactly one sibling driver, as `StepDevInputs` (`src/vllm/model_executor/models/qwen3_5.cpp:3894`): `grep -c StepDevInputs` returns 41 lines there and 0 in each of `qwen3_moe.cpp`, `qwen3.cpp`, `deepseek_v2.cpp` and `voxtral.cpp`. One capability, written once, unavailable to four models, with a live mitigation standing in its place. This does NOT weaken the framing rule that `ENG-CUDAGRAPH` established: the row still makes no throughput claim, and the prefill refutation (GB10 3.8% host-idle between launches, GPU-busy >96%, 27B prefill gap 92.5% non-GEMM glue) stands unchanged. Coverage AND correctness, never speed. Fixed in flow with the [`eng-cudagraph-break.md`](specs/eng-cudagraph-break.md) review repair ([#1163](https://github.com/mudler/vllm.cpp/issues/1163)) | record | | [#1181](https://github.com/mudler/vllm.cpp/issues/1181) | `FIX-READ-F32-SCALAR-GUARD` | `ReadF32Scalar` (`src/vllm/model_executor/models/qwen3_5_weights.cpp:312-318` @ `ab6e65216`) bounds its input with `t.data != nullptr && t.nbytes >= sizeof(float)`, a LOWER bound, and then `memcpy`s four bytes into a `float`. Two silent wrong-value paths follow and neither fails: an ARRAY is reduced to element 0, so a block-wise FP8 scale grid of shape `[ceil(N/128), ceil(K/128)]` passes and stands in for the whole weight (measured under [#1166](https://github.com/mudler/vllm.cpp/issues/1166) on `Qwen/Qwen3.8-27B-FP8` @ `017b9c7af6b5689d5dd426a76e0bc077eb5ca20a`, `q_proj.weight_scale_inv` is `[96, 40]`), and ANY dtype is reinterpreted, since that same tensor is `BF16` and its four bytes are two bf16 values read as one float. Both return a finite plausible float, so the output is fluent, plausible and wrong, which is what a token gate cannot see. Upstream makes both facts structural rather than optional: a per-tensor scale is a distinct parameter TYPE that asserts `loaded_weight.shape[0] == 1` (`vllm/model_executor/parameter.py:260-272,304-309` @ `555967922`, plus the `_assert_and_load` shape assert at `:93-96`), the slot is allocated `torch.float32` so a narrow on-disk dtype is VALUE-converted rather than reinterpreted (`utils/fp8_utils.py:1276`), and the declared strategy TENSOR/CHANNEL/BLOCK picks the parameter type before a byte is read (`compressed_tensors/schemes/compressed_tensors_w8a8_fp8.py:63,128`). The AUDIT corrects the issue's own framing twice. The 27 grep hits across five files are 5 definitions, 20 call sites and 2 comment references, and both counts are short: `ReadCtF32Scalar` (`include/vllm/model_executor/models/dense_weight_loaders.h:376`) is a SIXTH copy of the same defect under another name, reached from a SIXTH model file (`src/vllm/model_executor/models/qwen3_weights.cpp:100,126-128` through `LoadCtNvfp4W4A16`). Of the six, three check nothing, `LnReadF32Scalar`/`ShReadF32Scalar` check dtype but not count, and only `nemotron_h_weights.cpp:557-573` is correct, which makes it the model the shared guard generalizes. No call site legitimately passes a multi-element or non-F32 tensor, and every existing fixture emits rank-0 or `{1}` `F32`, so nothing in the tree needed the leniency. It is NOT merely latent: `dense_weight_loaders.h:73-74` and `docs/BENCHMARKS.md:52` both record `unsloth/Qwen3.6-27B-NVFP4` @ `ccdaab7e` as FP8 W8A8 throughout with BF16 PER-OUTPUT-CHANNEL scales, and `LoadAttnDense` branches on the weight dtype alone (`qwen3_5_dense_weights.cpp:478-480`), so those projections enter the per-tensor arm and hit both defects at once under the tensor name the loader actually asked for, with no misspelling to stop them. Fixed in flow by one `dense_loaders::ReadF32Scalar(get, name)` that refuses `numel != 1` naming the shape, refuses a non-`F32` dtype naming the dtype, and requires exactly four readable bytes, with the other five copies deleted onto it and `nemotron_h`'s `Loader`-based twin kept as the one tracked exception. A narrow dtype is refused rather than converted, because a one-element BF16 scale has never been read correctly here and the BF16 layout that IS shipped is per-channel, which the count check refuses first. Per-channel FP8, block-wise FP8 and any explicit narrow-dtype conversion stay owed. Spec [`read-f32-scalar-guard.md`](specs/read-f32-scalar-guard.md) | bug | | [#1185](https://github.com/mudler/vllm.cpp/issues/1185) | `ENV-ORACLE-WHEEL-IN-LEASE` | The pinned vLLM oracle BUILDS, installs, imports and sees the GPU inside an `rc` lease on `dgx:gpu0`, measured 2026-08-18, which falsifies the `nvcc` clause four records carried. [`lease-runtime-staging.md`](specs/lease-runtime-staging.md) said the oracle "needs `nvcc`, which the worker still lacks", and `.agents/environment.md`, [`mtp-k-gt-1.md`](specs/mtp-k-gt-1.md) and [`gpu-lease-methodology.md`](specs/gpu-lease-methodology.md) each derived a blocker from it. The build job (`buildvllm.sh`, staged sha256 prefix `15e140d41f44e7c2`) asserted the checkout against the pin BEFORE compiling, printing `PIN CONFIRMED` at `5559679229bc961848b121ccdeaa8fa5d79bec98` and aborting otherwise, took `nvcc` from the toolkit row `MODEL-NEMOTRON-H-ABI-A3-E2E` staged (`NVCC_RC=0`, CUDA `release 13.3, V13.3.73`) and produced `WHEEL_RC=0`, `PERSIST_RC=0` and a 434 MiB `vllm-0.1.dev1+g555967922.cu133-cp312-cp312-linux_aarch64.whl`, sha256 `7c58b339741a288fbb313f4f5196c9c92a9e3b3c3ebe2ea970b0ff50bb9bcba4`. The identity job (`oracleenv.sh`, prefix `6119f5223f5d818c`) asserted from `cd /`, outside any source tree: `vllm.__version__ = 0.1.dev1+g555967922`, `IDENTITY_RC=0`, `cuda True NVIDIA GB10`, `CUDA_RC=0`. SCOPE, and it carries the same weight as the result: RUNNING A MODEL IS UNTESTED. Only build, install, import and `torch.cuda.is_available()` are measured, and [`mtp-k-gt-1.md`](specs/mtp-k-gt-1.md) records that the last time an oracle reached this far it consumed the host in the step AFTER `torch.compile` and REBOOTED the box, at `gpu_memory_utilization` 0.75 and again at 0.30, so the fraction is not the lever. The version string is an OPEN discrepancy: `.agents/upstream-sync.md` records `vllm_runtime_version = 0.23.1rc1.dev1511+g555967922`, the commit segment matches and satisfies the pin's binding `+g` rule, and the prefix differs because a shallow fetch stops `setuptools_scm` counting commits since the last tag, so a full-string gate needs a deeper fetch or a recorded pretend-version. The venv is NOT staged, because that job was killed at a 90-minute ceiling mid-copy and its partial tree was removed, so only the WHEEL is durable. Four staging walls, all artifacts of the NAS rather than of CUDA: `cp -a` preserves `file_mode=0664` so `nvcc` exited 126. CIFS `nounix` stores no symlink so `include` and `lib64` vanished and CMake reported `Could NOT find CUDA (missing: CUDA_INCLUDE_DIRS CUDA_CUDART_LIBRARY) (found version "13.3")`, naming the version and denying the toolkit in one line. 32 library links `libfoo.so` and `libfoo.so.MAJOR` had to be rebuilt because only the `libfoo.so.X.Y.Z` real files survived. And `markupsafe` existed as a dist-info with NO package files from a `pip --target` killed at a 35-minute ceiling, so Marlin codegen died on `ModuleNotFoundError`. The `rc` worker container is REUSED between jobs, so a repair inside a staging branch is skipped on the next run (`nvcc already in place`) and an environment repair must be unconditional and assert its postcondition. CONSEQUENCE for the rows #1129 blocked, [#1003](https://github.com/mudler/vllm.cpp/issues/1003), [#915](https://github.com/mudler/vllm.cpp/issues/915), [#821](https://github.com/mudler/vllm.cpp/issues/821) and [#81](https://github.com/mudler/vllm.cpp/issues/81): UNBLOCKED FOR THE BUILD STEP and STILL BLOCKED FOR A MODEL RUN. None can take a measurement until a model run is demonstrated. Job details, walls and non-claims in [`oracle-wheel-in-lease.md`](specs/oracle-wheel-in-lease.md) | verification | +| [#1184](https://github.com/mudler/vllm.cpp/issues/1184) | `ENG-CUDAGRAPH-DEDUP` | `VT_CUDA_GRAPH_DEDUP=1` cannot complete a single decode step once a CUDA graph is actually captured. Measured 6/6 deterministic on GB10: `vt graph dedup: captured 1 graphs, deduped to 1 execs`, then `[Qwen3DenseDecodeGraph] captured dense decode graph for padded size S=8`, then `engine-fatal: EngineCore busy loop threw: vt cuda: greedy_argmax launch: invalid device function` -- from a `greedy_argmax` launch that had SUCCEEDED. The OFF and `=0` arms are clean and byte-identical across 7 runs on the identical binary and libraries, which is what makes the asymmetry arm-attributable. `greedy_argmax` is not at fault. `src/vt/graph_dedup_runtime.h` @ `3ed6253c6` walks the retained raw graph with CUDA runtime APIs and is DESIGNED to see several of them fail -- the `cudaGraphExecUpdate` probe refusing a fold is the feature working, not an exception, and the topology walk has five more escapes that degrade the key rather than abort inside a capture -- but it never consumed the runtime's sticky per-thread error: `grep -n 'cudaGetLastError\|hipGetLastError' src/vt/graph_dedup_runtime.h` returned nothing. Twelve sites swallowed a non-success return without clearing (`:98,100,115,130,137,146,154,173,179,281,307,317`); the two that mattered are the probe (`:307,:317`), whose refusal is routine, and `cudaGraphInstantiate` (`:281`). The next unrelated kernel, launched with the ordinary `kernel<<<>>>(); Check(cudaGetLastError())` pattern, then reported our refusal as its own failure. Every symptom follows: it needs BOTH `dedup=1` and a real capture, `CUDA_LAUNCH_BLOCKING=1` does not move it because the latch is host-side and synchronous rather than a deferred async error, and `cudaGraphLaunch` returns success because a return value does not consume the latch. Fixed in flow, structurally rather than site-by-site: twelve hand-placed clears are a fix the thirteenth fallible call misses, so the clear lives in `ScopedLatchClear`'s destructor (`src/vt/graph_dedup_latch.h`, new) installed at the binding's entry points, which are exactly the six `GraphDedupOps` members; `MakeLatchGuardedOps` is the table's only constructor and takes the raw functions as template arguments, so no raw address reaches a field and a seventh operation wired elsewhere leaves its field null, which the registry constructor refuses. One line covers CUDA and HIP because there is one source. The coverage gap that hid it is closed in the same change: the device-free half of the signature walk moves to `src/vt/graph_dedup_signature.h` and is gated by the new `tests/vt/test_graph_dedup_runtime.cpp` (13 cases, 51 assertions, 7/7 negative mutations detected) -- Kahn ordering, topological re-index, sorted edge emission, the depth-4 child bound and the four graph-level escapes, none of which was reached by any test on any tier before. STATED LIMIT: a CPU test drives a fake runtime and cannot observe the real latched error, so it proves the guard's structure and not the device outcome; the device A/B re-run is owed under `## Owed` in [`eng-cudagraph-dedup.md`](specs/eng-cudagraph-dedup.md), which also records that the shipped async serving path never reaches this feature at all | bug | diff --git a/.agents/specs/eng-cudagraph-dedup.md b/.agents/specs/eng-cudagraph-dedup.md index 164522252..04a9e6e7d 100644 --- a/.agents/specs/eng-cudagraph-dedup.md +++ b/.agents/specs/eng-cudagraph-dedup.md @@ -10,9 +10,12 @@ is why the load-bearing gate below is byte-identity rather than a speed ratio. ## Now -`ACTIVE`. The shared registry, its CPU contract suite and the CUDA wiring land here. -The device-level byte-identity A/B and the ROCm leg are named under -[`## Owed`](#owed) with the issue that owns each. +`ACTIVE`. The shared registry, its CPU contract suite and the CUDA wiring land here, +together with the [#1184](https://github.com/mudler/vllm.cpp/issues/1184) repair: the +runtime binding is allowed to see calls fail and never consumed the runtime's latched +error, so the first unrelated kernel after a capture reported our refusal as its own +failure and every `VT_CUDA_GRAPH_DEDUP=1` run died. The device-level byte-identity A/B +and the ROCm leg are named under [`## Owed`](#owed) with the issue that owns each. ## Scope @@ -138,7 +141,9 @@ instantiate were never a source of baked pointers. | New / changed | What | |---|---| | `src/vt/graph_dedup.h` (new) | `GraphDedupOps` (six function pointers: signature, instantiate, update, destroy exec, destroy graph, launch) and `GraphDedupRegistry`. Header-only, no device dependency, CPU-unit-testable — the same shape `src/vt/cuda/graph_safe_scratch.h` used for the graph-safe scratch bookkeeping. Also `GraphDedupEnabled()`, the `VT_CUDA_GRAPH_DEDUP` read, and the `CapturedCount()` / `ExecCount()` accessors. There is no single `Stats()` method: two accessors read at the call site beat one struct nobody stores. | -| `src/vt/graph_dedup_runtime.h` (new) | The CUDA / HIP ops table, written **once** and bound to either runtime by a macro alias block. `GraphSignature()` lives here. This is the file that must not become two hand-written copies. | +| `src/vt/graph_dedup_runtime.h` (new) | The CUDA / HIP ops table, written **once** and bound to either runtime by a macro alias block. This is the file that must not become two hand-written copies. After [#1184](https://github.com/mudler/vllm.cpp/issues/1184) it holds only what is device-shaped: the `GetEdges` version binding, the five node-payload cases, the six raw operations, and the `Runtime` policy that carries `ClearLatchedError`. `Ops()` is its only export. | +| `src/vt/graph_dedup_latch.h` (new, #1184) | `ScopedLatchClear` and `MakeLatchGuardedOps`. Device-free. Every `GraphDedupOps` field is a wrapper address, so a runtime error latched anywhere inside the binding is consumed before control leaves it — on a plain return, on a degradation escape, and on a `VT_CHECK` unwinding. | +| `src/vt/graph_dedup_signature.h` (new, #1184) | The device-free half of the structural signature: Kahn ordering, the topological re-index, the sorted edge emission, the child-graph depth bound and the four degradation escapes, templated on the `Rt` policy above. Split out so a CPU test can drive it; it was previously reachable by no test on any tier. | | `src/vt/cuda/cuda_backend.cu` | `EndCaptureGraph` routes through the registry when enabled and retains the raw graph; `ReplayGraph` and `DestroyGraph` dispatch on `registry->Owns(handle)`. | | `src/vt/rocm/rocm_backend.hip` | The same three edits against the same shared header. | | `docs/ENVIRONMENT.md` | `VT_CUDA_GRAPH_DEDUP`. Required: `scripts/check-env-doc.py` fails a `VT_*` name that is neither documented nor allowlisted, and this is a user-facing behaviour knob rather than kernel-internal tuning. | @@ -174,11 +179,32 @@ an assertion about intent. | `a probe the driver cannot instantiate degrades instead of driving a null exec` | a failed probe answers "cannot fold" rather than passing a null executable to `cudaGraphExecUpdate` | ours; added by the fresh review of #1178 | | `a replay update the driver refuses fails loudly rather than launching stale nodes` | a refusal on the fold `Register` never probed throws, and the executable's previous contents are NOT launched under the asking handle | ours; pins the safety claim that makes the transitivity assumption below survivable | +A second suite, `tests/vt/test_graph_dedup_runtime.cpp`, was added by +[#1184](https://github.com/mudler/vllm.cpp/issues/1184) over the two device-free halves +of the runtime binding. It has no upstream twin: SGLang runs on a Python runtime that +raises rather than latching, so the whole error-latch class does not exist there. + +| Case | Guarantee | +|---|---| +| `every guarded operation consumes the latch its own failure set` | five arms, one per fallible operation; after each the latch is clean and a stand-in for the next unrelated kernel launch does not throw. The refusal still carries the driver's `detail` string, so clearing does not swallow the reason the caller acts on | +| `the clear survives an operation that throws` | `Launch` reports a failed `cudaGraphLaunch` by throwing. This case is what decides between a clear placed after the call and a clear placed in a destructor: only the destructor runs on the unwinding path | +| `a refused fold leaves nothing latched for the next kernel` | #1184 in the shape a CPU tier can hold: a registry whose probe the driver REFUSES — the feature's normal operation — leaves nothing for the next caller to misread | +| `no table field holds an address that skips the guard` | each of the six `GraphDedupOps` fields differs from the raw function's address | +| `an operation the guarded builder does not wire cannot reach the registry` | a `GraphDedupOps` member that `MakeLatchGuardedOps` does not set stays null, and the registry's constructor refuses an incomplete table | +| `the signature is independent of the order the runtime reports nodes in` | the load-bearing property of the walk: `cudaGraphGetNodes` promises no order, so without the Kahn re-index two captures of one topology key differently and fold nothing, observable only in the device-side log line | +| `the signature is independent of the order the runtime reports edges in` | a diamond with its edges scrambled keys identically | +| `the signature separates topologies that differ only in their edges` / `only in a node payload` | discrimination in both directions | +| `a child graph contributes its own signature, bounded at depth four` | the bound is exact on both sides: level 4's payload is emitted and its child is not walked | +| `a child graph is walked, not merely noted` | the nested signature appears inside the parent's, in the exact byte form | +| `each runtime failure degrades the key instead of aborting inside a capture` | five escapes, five exact strings: `nodes?;`, `edges?;`, `edge?;`, `cycle?;`, `[A,;node?;` | +| `an empty graph still produces a signature` | `[]` rather than an escape | + ## Gates 1. **Red first.** Every case above is written and run against the un-implemented registry, and the red output is captured, before the implementation exists. -2. **Focused green.** `ctest -R test_graph_dedup`. +2. **Focused green.** `ctest -R test_graph_dedup`, which now selects both + `test_graph_dedup` and `test_graph_dedup_runtime`. 3. **Mutation.** Each guarantee is deleted or inverted in a scratch copy, the focused suite is proven to fail, and the tree is restored byte-for-byte. `git diff --stat` and the compile status are printed for every mutation, because a mutation that fails @@ -186,6 +212,12 @@ an assertion about intent. 4. **Full gate.** `./scripts/agent-preflight.sh`. 5. **CUDA compile.** The `cuda-fat-build` CI job (`.github/workflows/ci.yml:669-712`, container `nvidia/cuda:13.3.0-devel-ubuntu24.04`) compiles the CUDA leg on the PR. + This is the gate. The #1184 session had no local toolkit and used a scratch + header-shape stub as a proxy — `graph_dedup_runtime.h` instantiated through `Ops()` + at `CUDART_VERSION` 12090 and 13030 and on the HIP arm, `-Werror`, all three clean, + with a negative control proving the instrument can fail. A proxy is not the gate: the + CUDA 13 `cudaGraphGetEdges` break this file already took was invisible to exactly + this kind of local check and only `cuda-fat-build` reported it. 6. **Device byte-identity A/B (owed, see below).** Same binary, `VT_CUDA_GRAPH_DEDUP` off then on, identical prompts and sampling, token-exact equality, with **more than one replay per padded bucket** — a first replay can be correct by accident, and the @@ -265,21 +297,87 @@ an assertion about intent. because the engine drives one device from one thread. A lock belongs at the registry the moment that stops being true, and `ENG-CUDAGRAPH-BREAK` must not widen who captures without adding one. -- **The signature builder has no executable coverage; it is compile-gated only.** The - 13 cases in `tests/vt/test_graph_dedup.cpp` drive `GraphDedupRegistry` through a fake - ops table, so they gate the registry's launch-sequence identity and nothing else. - `src/vt/graph_dedup_runtime.h` — the Kahn ordering, the topological re-index, the - sorted edge emission, the five node-payload cases, the depth-4 child-graph bound and - the five degradation escapes — is reached by no test on any tier; `cuda-fat-build` - proves only that it compiles. "13/13 negative mutations detected" is a statement about - `graph_dedup.h`, and must not be read as coverage of that file. The probe caps the - blast radius, so an unstable signature can never produce a wrong replay — but it has - exactly one silent mode: a signature that is unstable for some topology folds nothing, +- **A file allowed to see the runtime fail must consume the runtime's latched error, + and this one did not.** [#1184](https://github.com/mudler/vllm.cpp/issues/1184). The + whole safety argument above rests on a refused `cudaGraphExecUpdate` being NORMAL — + the signature is a lookup key, the driver is the authority, and a refusal means "do + not fold". A CUDA call that fails also LATCHES its code in the runtime's sticky + per-thread slot, and none of the twelve fallible calls in `graph_dedup_runtime.h` + consumed it. The rest of this tree launches kernels with the ordinary + `kernel<<<>>>(...); Check(cudaGetLastError())` pattern, so the next unrelated kernel + read our routine refusal and reported it as its own failure. On GB10 that was 6/6 + deterministic: + + vt graph dedup: captured 1 graphs, deduped to 1 execs + [Qwen3DenseDecodeGraph] captured dense decode graph for padded size S=8 (real B=8) + engine-fatal: EngineCore busy loop threw: + vt cuda: greedy_argmax launch: invalid device function + + from a `greedy_argmax` launch that had SUCCEEDED. Every symptom follows from the + mechanism and none of them from `greedy_argmax`: it needs both `dedup=1` and a real + capture (no capture, no probe, no latch); `CUDA_LAUNCH_BLOCKING=1` does not move it, + because the latch is host-side and synchronous rather than a deferred async error; and + `cudaGraphLaunch` itself returns success, because reading a return value does not + consume the latch. The two sites that mattered most were the ones whose failure is by + design — the `cudaGraphExecUpdate` probe, and the `cudaGraphInstantiate` the earlier + review made fatal — which is the general shape: the more deliberate a "this may fail" + path is, the more reliably it poisons the next caller. + + **The repair is a scope guard, not twelve clears.** Twelve hand-placed + `cudaGetLastError()` calls are a fix the thirteenth fallible call silently misses, and + a file whose design is "these calls are allowed to fail" will grow a thirteenth. The + clear therefore lives in `ScopedLatchClear`'s destructor + (`src/vt/graph_dedup_latch.h`), installed at the binding's entry points, which are + exactly the six `GraphDedupOps` members. Every exit runs it: a plain return, a + degradation escape, and the `VT_CHECK` unwinding out of `Launch`. `MakeLatchGuardedOps` + is the only constructor of the table and takes the raw functions as template + arguments, so no raw address reaches a field; a seventh operation wired anywhere else + leaves its field null and `GraphDedupRegistry`'s constructor refuses the table. One + line covers both arms because there is one source: `VTGD_FN(GetLastError)` resolves to + `cudaGetLastError` or `hipGetLastError`. + + **What the CPU suite proves and what it does not.** A CPU test drives a fake runtime, + so it cannot observe the CUDA runtime's real latched-error state and it cannot prove + #1184 is gone on a device. It proves the structure the fix rests on, red-first: with + the destructor emptied to the pre-fix state the suite reports 22 failed assertions + including `CHECK_NOTHROW(NextUnrelatedKernelLaunch()) THREW exception: "greedy_argmax + launch: invalid device function"`, which is the production message reproduced from the + mechanism alone. The device half stays owed. + +- **The signature builder had no executable coverage, and now has most of one.** + Previously the 13 cases in `tests/vt/test_graph_dedup.cpp` gated `GraphDedupRegistry` + and nothing else, and all of `src/vt/graph_dedup_runtime.h` was reached by no test on + any tier — the gap that hid #1184 for a whole review cycle. The device-free half is now + `src/vt/graph_dedup_signature.h` and is driven by `tests/vt/test_graph_dedup_runtime.cpp` + over a fake `Rt`: the Kahn ordering, the topological re-index, the sorted edge + emission, the child-graph depth bound and the four graph-level degradation escapes. + Seven negative mutations were detected, each with its `git diff --stat`, compile status + and exit status recorded; one first attempt failed to build under `-Werror` and was + re-run in a compiling form, because a mutation that fails to build reads as a passing + test. **Still uncovered on any tier:** the five node-payload cases (kernel, memcpy, + memset, child graph, default) and their four query escapes, which read CUDA parameter + structs and stay behind the policy. And the silent mode the original note named is + unchanged in kind: a signature that is unstable for some real topology folds nothing, every CPU test stays green, and the only observable is the device-side - `"captured N graphs, deduped to N execs"` log line. What would cover it: a CUDA-tier - test that captures the same trivial graph twice and asserts the two signatures are - byte-equal, and a second that captures two deliberately different topologies and - asserts they differ. Both need a device, so both are owed with the A/B. + `"captured N graphs, deduped to N execs"` line. What would close it is still + device-tier — capture one trivial graph twice and assert the signatures are byte-equal, + capture two different topologies and assert they differ — so both remain owed with the + A/B. + +- **The shipped async serving path never reaches this feature, which bounds the row's + value.** Measured during the #1184 device gate: `vllm-bench` on the async path captures + NO decode graph at all, because `DenseDecodeGraphForward` returns `nullopt` whenever + `input.device_token_ids != nullptr` (`src/vllm/model_executor/models/qwen3.cpp`, the + [#323](https://github.com/mudler/vllm.cpp/issues/323) mitigation that + [#1179](https://github.com/mudler/vllm.cpp/issues/1179) re-derived). No capture means + no registration, so dedup engages only under `VT_ASYNC_RUNNER=0` — which is not the + default and not what a user serves with. This does not change the row's correctness + argument and it does change its worth: the executables this row folds are, on the + default path, executables that are never instantiated in the first place. The saving is + real only for the drivers and configurations that still capture. Recorded here rather + than filed separately because the fix is `StepDevInputs`-shaped and already owned by + `ENG-CUDAGRAPH-BREAK` / #1179; this row must not be sold on a number the default path + cannot produce. - **Two runtime-API calls changed shape across CUDA major versions, and only one was foreseen.** `cudaGraphExecUpdate` took the 3-argument `cudaGraphExecUpdateResultInfo` form at CUDA 12, so both it and the legacy 4-argument form are bound; HIP has only @@ -299,7 +397,9 @@ an assertion about intent. | The device byte-identity A/B and the executable-count measurement (W4) | [#1162](https://github.com/mudler/vllm.cpp/issues/1162) | needs a leased CUDA box; the CPU tier proves the registry's launch-sequence identity, not the device's | | Flipping `VT_CUDA_GRAPH_DEDUP` on by default | [#1162](https://github.com/mudler/vllm.cpp/issues/1162) | gated on W4. A default is a measurement, not a preference | | Probing `group.current_raw` rather than `raws.front()`, retiring the transitivity assumption above | [#1162](https://github.com/mudler/vllm.cpp/issues/1162) | it changes probe behaviour, and the device A/B is measuring the current one. Land it with the A/B rerun, not before | -| Executable coverage for `src/vt/graph_dedup_runtime.h`'s signature builder (stability and discrimination, both device-tier) | [#1162](https://github.com/mudler/vllm.cpp/issues/1162) | it needs a real `cudaGraph_t`, so it rides with the leased box the A/B already needs | +| Executable coverage for the signature builder's DEVICE half — stability and discrimination on a real `cudaGraph_t`, plus the five node-payload cases and their query escapes | [#1162](https://github.com/mudler/vllm.cpp/issues/1162) | the device-free half is now covered by `tests/vt/test_graph_dedup_runtime.cpp`; what remains needs a real `cudaGraph_t`, so it rides with the leased box the A/B already needs | +| Re-running the device A/B after the [#1184](https://github.com/mudler/vllm.cpp/issues/1184) repair, and confirming that `VT_CUDA_GRAPH_DEDUP=1` now completes a decode step | [#1184](https://github.com/mudler/vllm.cpp/issues/1184) | the CPU suite proves the guard's structure over a fake runtime; only a device can prove that `cudaGetLastError` clears the real latch and that the run survives. #1184 stays open until that run exists | +| Reaching the feature from the DEFAULT serving path. Dedup engages only under `VT_ASYNC_RUNNER=0`, because the async path captures no decode graph at all ([#323](https://github.com/mudler/vllm.cpp/issues/323) mitigation) | [#1179](https://github.com/mudler/vllm.cpp/issues/1179) | the repair is the `StepDevInputs`-shaped one `ENG-CUDAGRAPH-BREAK` already owns; until it lands, this row's saving is unreachable on the configuration users serve with | | The ROCm leg's compile and run verification | [#41](https://github.com/mudler/vllm.cpp/issues/41) | no ROCm hardware or `hipcc` is reachable from this session and CI has no ROCm job, so the HIP wiring is written against the same shared header but is compile-unverified | ## Stop conditions diff --git a/src/vt/graph_dedup_latch.h b/src/vt/graph_dedup_latch.h new file mode 100644 index 000000000..80533bcf9 --- /dev/null +++ b/src/vt/graph_dedup_latch.h @@ -0,0 +1,120 @@ +// vllm.cpp original — the error-latch discipline for the graph dedup ops table. +// +// Row ENG-CUDAGRAPH-DEDUP, issues #1162 and #1184, spec +// .agents/specs/eng-cudagraph-dedup.md. +// +// WHAT THIS EXISTS FOR. `src/vt/graph_dedup_runtime.h` calls the CUDA / HIP runtime and +// is ALLOWED to see several of those calls fail. That is not an accident of the port, +// it is the feature: the whole safety argument for dedup is that the signature is only +// a lookup key and `cudaGraphExecUpdate` is the authority, so a probe the driver +// REFUSES is the normal, expected outcome that means "do not fold". The topology walk +// has five more escapes of the same kind, each of which degrades the key rather than +// aborting inside a capture. +// +// A CUDA runtime call that fails LATCHES its code in the runtime's per-thread sticky +// slot, and nothing in that file consumed the latch. The rest of this tree launches +// kernels with the ordinary `kernel<<<>>>(...); Check(cudaGetLastError())` pattern, so +// the NEXT unrelated kernel read our routine refusal and reported it as its own +// failure. That is #1184, observed 6/6 on GB10 as +// +// vt graph dedup: captured 1 graphs, deduped to 1 execs +// engine-fatal: EngineCore busy loop threw: +// vt cuda: greedy_argmax launch: invalid device function +// +// from a `greedy_argmax` launch that had SUCCEEDED. Every symptom follows: it needs a +// real capture to exist (no capture, no probe, no latch), `CUDA_LAUNCH_BLOCKING=1` does +// not move it (the latch is host-side and synchronous), and `cudaGraphLaunch` itself +// returns success (a return value does not consume the latch). +// +// WHY THIS IS A TYPE AND NOT TWELVE `cudaGetLastError()` CALLS. Twelve hand-placed +// clears are a fix the thirteenth fallible call silently misses, and a file whose whole +// design is "these calls are allowed to fail" will grow a thirteenth. The clear +// therefore lives in a destructor at the file's ENTRY POINTS, and the entry points are +// exactly the six members of `GraphDedupOps`, which the backends reach only through +// `graph_dedup_rt::Ops()`. Every path out of the file — a plain return, a degradation +// escape, or a `VT_CHECK` unwinding — runs it. +// +// `MakeLatchGuardedOps` below is the only constructor of that table, and it takes the +// raw functions as template arguments, so their addresses never reach a `GraphDedupOps` +// field. A seventh operation added to `GraphDedupOps` and wired anywhere else leaves the +// field null, and `GraphDedupRegistry`'s constructor refuses an incomplete table. So the +// bypass is not merely discouraged, it does not construct. +// +// This header is device-free on purpose, the same way `vt/graph_dedup.h` is: the guard +// discipline is then gated by a CPU test on every platform instead of only where a +// driver exists. See the honesty note in `tests/vt/test_graph_dedup_runtime.cpp` about +// what such a test can and cannot prove. +#ifndef VT_GRAPH_DEDUP_LATCH_H_ +#define VT_GRAPH_DEDUP_LATCH_H_ + +#include + +#include "vt/graph_dedup.h" + +namespace vt::graph_dedup_latch { + +// Runs `Rt::ClearLatchedError()` on every exit from the scope it is declared in, +// including an exception unwinding through it. `Rt` is the runtime policy — +// `graph_dedup_rt::CudaRuntime` in production, a counting fake in the suite. +template +struct ScopedLatchClear { + ScopedLatchClear() = default; + // Destructors are implicitly noexcept, and cudaGetLastError / hipGetLastError cannot + // throw, so this is safe on the unwinding path that `Launch`'s VT_CHECK takes. + ~ScopedLatchClear() { Rt::ClearLatchedError(); } + ScopedLatchClear(const ScopedLatchClear&) = delete; + ScopedLatchClear& operator=(const ScopedLatchClear&) = delete; +}; + +// One wrapper per `GraphDedupOps` member. The signatures are spelled out rather than +// deduced through an `auto` non-type parameter partial specialisation, because this +// header compiles under MSVC as well and the explicit form needs nothing past C++14. +template +inline std::string GuardedSignature(void* raw_graph) { + ScopedLatchClear clear; + return Fn(raw_graph); +} + +template +inline void* GuardedInstantiate(void* raw_graph) { + ScopedLatchClear clear; + return Fn(raw_graph); +} + +template +inline bool GuardedUpdate(void* exec, void* raw_graph, std::string* detail) { + ScopedLatchClear clear; + return Fn(exec, raw_graph, detail); +} + +template +inline void GuardedUnary(void* handle) { + ScopedLatchClear clear; + Fn(handle); +} + +template +inline void GuardedLaunch(void* exec, void* stream) { + ScopedLatchClear clear; + Fn(exec, stream); +} + +// The ONLY way a GraphDedupOps table is built for a real runtime. Every field is a +// wrapper address; none is a raw function address. +template +inline GraphDedupOps MakeLatchGuardedOps() { + GraphDedupOps table; + table.signature = &GuardedSignature; + table.instantiate = &GuardedInstantiate; + table.update = &GuardedUpdate; + table.destroy_exec = &GuardedUnary; + table.destroy_graph = &GuardedUnary; + table.launch = &GuardedLaunch; + return table; +} + +} // namespace vt::graph_dedup_latch + +#endif // VT_GRAPH_DEDUP_LATCH_H_ diff --git a/src/vt/graph_dedup_runtime.h b/src/vt/graph_dedup_runtime.h index 89c053d08..a38f1e094 100644 --- a/src/vt/graph_dedup_runtime.h +++ b/src/vt/graph_dedup_runtime.h @@ -1,18 +1,30 @@ // vllm.cpp original — the CUDA / HIP binding for the graph-executable dedup registry. // -// Row ENG-CUDAGRAPH-DEDUP, issue #1162, spec .agents/specs/eng-cudagraph-dedup.md. +// Row ENG-CUDAGRAPH-DEDUP, issues #1162 and #1184, spec +// .agents/specs/eng-cudagraph-dedup.md. // Ported from SGLang's dedup mixin at pin f63458b5be: // python/sglang/srt/model_executor/runner_backend/cuda_graph_dedup_mixin.py // :27-37 dedup_update -> Update // :105-114 kernel_node_payload -> AppendKernelPayload -// :117-136 graph_node_payload -> AppendNodePayload -// :139-179 graph_signature -> AppendGraphSignature +// :117-136 graph_node_payload -> Runtime::AppendNodePayload +// :139-179 graph_signature -> graph_dedup_sig::AppendGraphSignature // // ONE source, two runtimes. The CUDA and HIP graph APIs differ only by symbol prefix // and by two call shapes, so binding them with an alias block keeps this a single path // instead of the hand-written parallel path AGENTS.md forbids. Include it from a .cu // for the CUDA leg, or define VT_GRAPH_DEDUP_HIP first and include it from a .hip. // +// THIS FILE IS ALLOWED TO SEE THE RUNTIME FAIL, AND THAT IS WHY IT NEEDS A LATCH GUARD. +// The cudaGraphExecUpdate probe refusing a fold is the feature working, not an +// exception, and the topology walk has four more escapes that degrade the key rather +// than abort inside a capture. Every one of those latches an error code in the runtime's +// sticky per-thread slot, and until #1184 nothing here consumed it, so the next +// unrelated kernel reported OUR refusal as its own failure. The clear is therefore not +// placed beside each fallible call — a thirteenth fallible call would miss it — but in a +// destructor wrapped around every entry point, by `vt/graph_dedup_latch.h`. Ops() below +// is the file's ONLY export, and it is built solely through MakeLatchGuardedOps, so no +// raw function address in this file ever reaches a GraphDedupOps field. +// // DELIBERATE ADAPTATIONS, both recorded in the spec's `## Upstream chain`: // * kernel identity is the host function POINTER (cudaKernelNodeParams::func) rather // than the demangled name SGLang reads through cuKernelGetName / cuFuncGetName. The @@ -26,16 +38,14 @@ #ifndef VT_GRAPH_DEDUP_RUNTIME_H_ #define VT_GRAPH_DEDUP_RUNTIME_H_ -#include #include #include -#include -#include #include -#include #include #include "vt/graph_dedup.h" +#include "vt/graph_dedup_latch.h" +#include "vt/graph_dedup_signature.h" #if defined(VT_GRAPH_DEDUP_HIP) #include @@ -77,12 +87,14 @@ using Stream = cudaStream_t; namespace vt::graph_dedup_rt { +using vt::graph_dedup_sig::AppendNumber; + // cudaGraphGetEdges gained a fifth `cudaGraphEdgeData*` parameter in CUDA 13; the // four-argument form is what CUDA 12 and HIP ship. Both shapes are bound here rather -// than at the call sites so the topology walk below stays one piece of code. This is +// than at the call sites so the topology walk stays one piece of code. This is // the one place a local compile against CUDA 12 headers could not have caught, and did // not: the `cuda-fat-build` job on nvidia/cuda:13.3.0 is what reported it. -inline bool GetEdges(Graph graph, Node* from, Node* to, std::size_t* num_edges) { +inline bool GetEdgesRaw(Graph graph, Node* from, Node* to, std::size_t* num_edges) { #if defined(VT_GRAPH_DEDUP_HIP) return hipGraphGetEdges(graph, from, to, num_edges) == hipSuccess; #elif CUDART_VERSION >= 13000 @@ -101,13 +113,6 @@ inline bool GetEdges(Graph graph, Node* from, Node* to, std::size_t* num_edges) #endif } -inline void AppendNumber(std::string* out, long long value) { - out->append(std::to_string(value)); - out->push_back(','); -} - -inline void AppendGraphSignature(Graph graph, std::string* out, int depth); - // (func, gridDim, blockDim, sharedMemBytes) — cuda_graph_dedup_mixin.py:105-114 minus // the driver-API attribute tuple, per the header note. inline bool AppendKernelPayload(Node node, std::string* out) { @@ -124,150 +129,87 @@ inline bool AppendKernelPayload(Node node, std::string* out) { return true; } -// cuda_graph_dedup_mixin.py:117-136, the same five cases in the same order. -inline bool AppendNodePayload(Node node, std::string* out, int depth) { - NodeType type{}; - if (VTGD_FN(GraphNodeGetType)(node, &type) != VTGD_SUCCESS) return false; - AppendNumber(out, static_cast(type)); - switch (type) { - case VTGD_NODE_KERNEL: - return AppendKernelPayload(node, out); - case VTGD_NODE_MEMCPY: { - MemcpyNodeParams params{}; - if (VTGD_FN(GraphMemcpyNodeGetParams)(node, ¶ms) != VTGD_SUCCESS) return false; - AppendNumber(out, static_cast(params.kind)); - AppendNumber(out, static_cast(params.extent.width)); - AppendNumber(out, static_cast(params.extent.height)); - AppendNumber(out, static_cast(params.extent.depth)); - return true; - } - case VTGD_NODE_MEMSET: { - MemsetNodeParams params{}; - if (VTGD_FN(GraphMemsetNodeGetParams)(node, ¶ms) != VTGD_SUCCESS) return false; - AppendNumber(out, static_cast(params.elementSize)); - AppendNumber(out, static_cast(params.width)); - AppendNumber(out, static_cast(params.height)); - return true; - } - case VTGD_NODE_GRAPH: { - Graph child = nullptr; - if (VTGD_FN(GraphChildGraphNodeGetGraph)(node, &child) != VTGD_SUCCESS) return false; - // Bounded: a capture nested past this depth degrades to a coarser key, which the - // probe still guards. Unbounded recursion over driver-owned structure is not a - // risk worth taking inside a capture path. - if (depth >= 4) return true; - AppendGraphSignature(child, out, depth + 1); - return true; - } - default: - return true; - } -} - -// cuda_graph_dedup_mixin.py:139-179. Node order as the runtime reports it is not a -// contract, so the nodes are re-indexed by a deterministic topological order (Kahn, -// always taking the lowest available index, exactly as the upstream heapq does) and the -// edge set is emitted in that order's terms. -inline void AppendGraphSignature(Graph graph, std::string* out, int depth) { - std::size_t num_nodes = 0; - if (VTGD_FN(GraphGetNodes)(graph, nullptr, &num_nodes) != VTGD_SUCCESS) { - out->append("nodes?;"); - return; - } - std::vector nodes(num_nodes); - if (num_nodes > 0 && - VTGD_FN(GraphGetNodes)(graph, nodes.data(), &num_nodes) != VTGD_SUCCESS) { - out->append("nodes?;"); - return; - } - - std::size_t num_edges = 0; - if (!GetEdges(graph, nullptr, nullptr, &num_edges)) { - out->append("edges?;"); - return; - } - std::vector from(num_edges); - std::vector to(num_edges); - if (num_edges > 0 && !GetEdges(graph, from.data(), to.data(), &num_edges)) { - out->append("edges?;"); - return; - } +// The `Rt` policy vt/graph_dedup_signature.h walks a graph through. Everything +// device-shaped lives here; the ordering, re-indexing and edge emission do not. +struct Runtime { + using Graph = vt::graph_dedup_rt::Graph; + using Node = vt::graph_dedup_rt::Node; - std::unordered_map index; - index.reserve(num_nodes * 2); - for (std::size_t i = 0; i < num_nodes; ++i) index[nodes[i]] = static_cast(i); + // Consume the runtime's sticky per-thread error. Called from a destructor at every + // entry point (vt/graph_dedup_latch.h) rather than beside each fallible call. On the + // HIP leg VTGD_FN resolves this to hipGetLastError, which has the same latch + // semantics; one line covers both arms because there is one source. + static void ClearLatchedError() { VTGD_FN(GetLastError)(); } - std::vector> children(num_nodes); - std::vector indegree(num_nodes, 0); - std::vector> edges; - edges.reserve(num_edges); - for (std::size_t e = 0; e < num_edges; ++e) { - const auto src = index.find(from[e]); - const auto dst = index.find(to[e]); - if (src == index.end() || dst == index.end()) { - out->append("edge?;"); - return; + static bool GetNodes(Graph graph, std::vector* out) { + std::size_t num_nodes = 0; + if (VTGD_FN(GraphGetNodes)(graph, nullptr, &num_nodes) != VTGD_SUCCESS) return false; + out->assign(num_nodes, Node{}); + if (num_nodes > 0 && + VTGD_FN(GraphGetNodes)(graph, out->data(), &num_nodes) != VTGD_SUCCESS) { + return false; } - children[static_cast(src->second)].push_back(dst->second); - ++indegree[static_cast(dst->second)]; - edges.emplace_back(src->second, dst->second); + return true; } - std::priority_queue, std::greater> ready; - for (std::size_t i = 0; i < num_nodes; ++i) { - if (indegree[i] == 0) ready.push(static_cast(i)); - } - std::vector order; - order.reserve(num_nodes); - while (!ready.empty()) { - const int current = ready.top(); - ready.pop(); - order.push_back(current); - for (int child : children[static_cast(current)]) { - if (--indegree[static_cast(child)] == 0) ready.push(child); + static bool GetEdges(Graph graph, std::vector* from, std::vector* to) { + std::size_t num_edges = 0; + if (!GetEdgesRaw(graph, nullptr, nullptr, &num_edges)) return false; + from->assign(num_edges, Node{}); + to->assign(num_edges, Node{}); + if (num_edges > 0 && !GetEdgesRaw(graph, from->data(), to->data(), &num_edges)) { + return false; } - } - if (order.size() != num_nodes) { - // A cycle is impossible in a captured graph; if the runtime ever reports one, the - // key degrades rather than the process aborting inside capture. - out->append("cycle?;"); - return; + return true; } - std::vector topo(num_nodes, 0); - for (std::size_t i = 0; i < order.size(); ++i) { - topo[static_cast(order[i])] = static_cast(i); - } - - out->push_back('['); - for (int node_index : order) { - if (!AppendNodePayload(nodes[static_cast(node_index)], out, depth)) { - out->append("node?;"); - return; + // cuda_graph_dedup_mixin.py:117-136, the same five cases in the same order. `*child` + // is set only for a child-graph node; the depth bound that governs recursing into it + // belongs to the walk, not to this switch. + static bool AppendNodePayload(Node node, std::string* out, Graph* child) { + NodeType type{}; + if (VTGD_FN(GraphNodeGetType)(node, &type) != VTGD_SUCCESS) return false; + AppendNumber(out, static_cast(type)); + switch (type) { + case VTGD_NODE_KERNEL: + return AppendKernelPayload(node, out); + case VTGD_NODE_MEMCPY: { + MemcpyNodeParams params{}; + if (VTGD_FN(GraphMemcpyNodeGetParams)(node, ¶ms) != VTGD_SUCCESS) return false; + AppendNumber(out, static_cast(params.kind)); + AppendNumber(out, static_cast(params.extent.width)); + AppendNumber(out, static_cast(params.extent.height)); + AppendNumber(out, static_cast(params.extent.depth)); + return true; + } + case VTGD_NODE_MEMSET: { + MemsetNodeParams params{}; + if (VTGD_FN(GraphMemsetNodeGetParams)(node, ¶ms) != VTGD_SUCCESS) return false; + AppendNumber(out, static_cast(params.elementSize)); + AppendNumber(out, static_cast(params.width)); + AppendNumber(out, static_cast(params.height)); + return true; + } + case VTGD_NODE_GRAPH: { + Graph nested = nullptr; + if (VTGD_FN(GraphChildGraphNodeGetGraph)(node, &nested) != VTGD_SUCCESS) { + return false; + } + *child = nested; + return true; + } + default: + return true; } - out->push_back(';'); } - out->push_back(']'); +}; - std::vector> topo_edges; - topo_edges.reserve(edges.size()); - for (const auto& edge : edges) { - topo_edges.emplace_back(topo[static_cast(edge.first)], - topo[static_cast(edge.second)]); - } - std::sort(topo_edges.begin(), topo_edges.end()); - for (const auto& edge : topo_edges) { - AppendNumber(out, edge.first); - AppendNumber(out, edge.second); - out->push_back(';'); - } -} +// The six operations, BEFORE the latch guard. Ops() is what the backends see, and it +// wraps every one of these; nothing outside this namespace takes their addresses. +namespace unguarded { inline std::string Signature(void* raw_graph) { - std::string out; - out.reserve(4096); - AppendGraphSignature(static_cast(raw_graph), &out, 0); - return out; + return vt::graph_dedup_sig::Signature(static_cast(raw_graph)); } inline void* Instantiate(void* raw_graph) { @@ -337,17 +279,16 @@ inline void Launch(void* exec, void* stream) { "graph dedup: graph launch failed"); } +} // namespace unguarded + inline const GraphDedupOps& Ops() { - static const GraphDedupOps ops = [] { - GraphDedupOps table; - table.signature = &Signature; - table.instantiate = &Instantiate; - table.update = &Update; - table.destroy_exec = &DestroyExec; - table.destroy_graph = &DestroyGraph; - table.launch = &Launch; - return table; - }(); + static const GraphDedupOps ops = + vt::graph_dedup_latch::MakeLatchGuardedOps(); return ops; } diff --git a/src/vt/graph_dedup_signature.h b/src/vt/graph_dedup_signature.h new file mode 100644 index 000000000..e32bcdc14 --- /dev/null +++ b/src/vt/graph_dedup_signature.h @@ -0,0 +1,162 @@ +// vllm.cpp original — the device-free half of the graph dedup structural signature. +// +// Row ENG-CUDAGRAPH-DEDUP, issues #1162 and #1184, spec +// .agents/specs/eng-cudagraph-dedup.md. +// Ported from SGLang's dedup mixin at pin f63458b5be: +// python/sglang/srt/model_executor/runner_backend/cuda_graph_dedup_mixin.py:139-179. +// +// WHY THIS IS SPLIT OUT OF graph_dedup_runtime.h. The topology walk — Kahn ordering, +// the topological re-index, the sorted edge emission, the child-graph depth bound and +// the four degradation escapes — is ordinary graph code with no CUDA in it. Leaving it +// inside a header that includes made it reachable by NO test on any +// tier: `cuda-fat-build` proved it compiled and nothing proved it was right. That gap is +// what the spec recorded under `## Risks/decisions` and what this file closes; the +// device-shaped part (node types, kernel and memcpy and memset parameters, the runtime's +// own node and edge queries) stays behind the `Rt` policy and is still device-only. +// +// THE POLICY CONTRACT. `Rt` supplies: +// +// using Graph = ...; // a handle type comparable against nullptr +// using Node = ...; // a handle type usable as an unordered_map key +// static bool GetNodes(Graph, std::vector* out); +// static bool GetEdges(Graph, std::vector* from, std::vector* to); +// static bool AppendNodePayload(Node, std::string* out, Graph* child); +// +// Each returns false on a runtime failure the walk is allowed to degrade around, and +// `AppendNodePayload` sets `*child` non-null exactly for a child-graph node, which the +// walk recurses into subject to `kMaxChildDepth`. Nothing here interprets a payload; the +// signature is a LOOKUP KEY and cudaGraphExecUpdate is the authority, so a key that is +// too coarse costs a wasted probe and can never make a replay wrong. +#ifndef VT_GRAPH_DEDUP_SIGNATURE_H_ +#define VT_GRAPH_DEDUP_SIGNATURE_H_ + +#include +#include +#include +#include +#include +#include +#include +#include + +namespace vt::graph_dedup_sig { + +// Bounded: a capture nested past this depth degrades to a coarser key, which the probe +// still guards. Unbounded recursion over driver-owned structure is not a risk worth +// taking inside a capture path. +inline constexpr int kMaxChildDepth = 4; + +inline void AppendNumber(std::string* out, long long value) { + out->append(std::to_string(value)); + out->push_back(','); +} + +// cuda_graph_dedup_mixin.py:139-179. Node order as the runtime reports it is not a +// contract, so the nodes are re-indexed by a deterministic topological order (Kahn, +// always taking the lowest available index, exactly as the upstream heapq does) and the +// edge set is emitted in that order's terms. +template +inline void AppendGraphSignature(typename Rt::Graph graph, std::string* out, int depth) { + using Node = typename Rt::Node; + using Graph = typename Rt::Graph; + + std::vector nodes; + if (!Rt::GetNodes(graph, &nodes)) { + out->append("nodes?;"); + return; + } + const std::size_t num_nodes = nodes.size(); + + std::vector from; + std::vector to; + if (!Rt::GetEdges(graph, &from, &to)) { + out->append("edges?;"); + return; + } + const std::size_t num_edges = std::min(from.size(), to.size()); + + std::unordered_map index; + index.reserve(num_nodes * 2); + for (std::size_t i = 0; i < num_nodes; ++i) index[nodes[i]] = static_cast(i); + + std::vector> children(num_nodes); + std::vector indegree(num_nodes, 0); + std::vector> edges; + edges.reserve(num_edges); + for (std::size_t e = 0; e < num_edges; ++e) { + const auto src = index.find(from[e]); + const auto dst = index.find(to[e]); + if (src == index.end() || dst == index.end()) { + out->append("edge?;"); + return; + } + children[static_cast(src->second)].push_back(dst->second); + ++indegree[static_cast(dst->second)]; + edges.emplace_back(src->second, dst->second); + } + + std::priority_queue, std::greater> ready; + for (std::size_t i = 0; i < num_nodes; ++i) { + if (indegree[i] == 0) ready.push(static_cast(i)); + } + std::vector order; + order.reserve(num_nodes); + while (!ready.empty()) { + const int current = ready.top(); + ready.pop(); + order.push_back(current); + for (int child : children[static_cast(current)]) { + if (--indegree[static_cast(child)] == 0) ready.push(child); + } + } + if (order.size() != num_nodes) { + // A cycle is impossible in a captured graph; if the runtime ever reports one, the + // key degrades rather than the process aborting inside capture. + out->append("cycle?;"); + return; + } + + std::vector topo(num_nodes, 0); + for (std::size_t i = 0; i < order.size(); ++i) { + topo[static_cast(order[i])] = static_cast(i); + } + + out->push_back('['); + for (int node_index : order) { + Graph child = nullptr; + if (!Rt::AppendNodePayload(nodes[static_cast(node_index)], out, &child)) { + out->append("node?;"); + return; + } + if (child != nullptr && depth < kMaxChildDepth) { + AppendGraphSignature(child, out, depth + 1); + } + out->push_back(';'); + } + out->push_back(']'); + + std::vector> topo_edges; + topo_edges.reserve(edges.size()); + for (const auto& edge : edges) { + topo_edges.emplace_back(topo[static_cast(edge.first)], + topo[static_cast(edge.second)]); + } + std::sort(topo_edges.begin(), topo_edges.end()); + for (const auto& edge : topo_edges) { + AppendNumber(out, edge.first); + AppendNumber(out, edge.second); + out->push_back(';'); + } +} + +template +inline std::string Signature(typename Rt::Graph graph) { + std::string out; + out.reserve(4096); + AppendGraphSignature(graph, &out, 0); + return out; +} + +} // namespace vt::graph_dedup_sig + +#endif // VT_GRAPH_DEDUP_SIGNATURE_H_ diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 7799ab2d7..0af193dee 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -1674,6 +1674,12 @@ vllm_cpp_add_test(test_graph_safe_scratch vt/test_graph_safe_scratch.cpp) target_include_directories(test_graph_safe_scratch PRIVATE ${CMAKE_SOURCE_DIR}/src) vllm_cpp_add_test(test_graph_dedup vt/test_graph_dedup.cpp) target_include_directories(test_graph_dedup PRIVATE ${CMAKE_SOURCE_DIR}/src) +# ENG-CUDAGRAPH-DEDUP / #1184: the two device-free halves of the CUDA/HIP binding -- +# the error-latch guard every entry point installs, and the structural-signature walk. +# Both lived inside a header that includes and were therefore reached +# by NO test on any tier; the latch defect hid in exactly that gap. +vllm_cpp_add_test(test_graph_dedup_runtime vt/test_graph_dedup_runtime.cpp) +target_include_directories(test_graph_dedup_runtime PRIVATE ${CMAKE_SOURCE_DIR}/src) # CPU threadpool (QUANT-GGUF-CPU-THREADPOOL): barrier/park-wake/chunk-steal # stress + the n_threads 1/3/20 determinism battery. Reaches the internal # header under src/ (the pool is not public API). diff --git a/tests/vt/test_graph_dedup_runtime.cpp b/tests/vt/test_graph_dedup_runtime.cpp new file mode 100644 index 000000000..67adb7675 --- /dev/null +++ b/tests/vt/test_graph_dedup_runtime.cpp @@ -0,0 +1,479 @@ +// vllm.cpp original — CPU-tier contract for the two device-free halves of the graph +// dedup runtime binding: the error-latch discipline (src/vt/graph_dedup_latch.h) and the +// structural-signature walk (src/vt/graph_dedup_signature.h). +// Row ENG-CUDAGRAPH-DEDUP, issues #1162 and #1184. +// +// WHY THIS SUITE EXISTS. Until #1184 the whole of src/vt/graph_dedup_runtime.h was +// reached by NO test on any tier: `grep -rn graph_dedup_runtime tests/` returned +// nothing, and `cuda-fat-build` proved only that it compiled. Two things hid inside that +// gap. The first is the defect: twelve runtime calls that this file is DESIGNED to see +// fail — the cudaGraphExecUpdate probe refusing a fold is the feature working, not an +// exception — none of which consumed the runtime's sticky per-thread error, so the next +// unrelated kernel launched with the ordinary `Check(cudaGetLastError())` pattern +// reported OUR refusal as its own failure. On GB10 that presented 6/6 as +// `greedy_argmax launch: invalid device function` from a launch that had succeeded. The +// second is the signature walk itself: Kahn ordering, the topological re-index, the +// sorted edge emission and the depth bound were all unexecuted by anything. +// +// WHAT THIS SUITE CANNOT PROVE, STATED PLAINLY. A CPU test drives a FAKE runtime. It +// cannot observe the CUDA runtime's real latched-error state, so it cannot prove that +// cudaGetLastError is the right call, that hipGetLastError has the same semantics, or +// that #1184 is gone on a device. What it does prove is the STRUCTURE the fix rests on: +// that every entry point clears on every exit path including an unwinding one, that a +// refusal inside the registry leaves nothing latched for the next caller, and that no +// table field holds an address that skips the guard. The device half — a same-binary +// VT_CUDA_GRAPH_DEDUP off/on A/B that survives a real capture — stays owed under #1162 +// and is recorded in .agents/specs/eng-cudagraph-dedup.md, not claimed here. +#include + +#include +#include +#include +#include + +#include "vt/graph_dedup.h" +#include "vt/graph_dedup_latch.h" +#include "vt/graph_dedup_signature.h" + +namespace { + +// =================================================================================== +// Part 1 — the error-latch discipline. +// =================================================================================== + +// A stand-in for the CUDA runtime's sticky per-thread error slot. A failing fake call +// SETS it; only a clear consumes it. That is the whole of the mechanism #1184 turned on: +// a return value does not consume the latch, which is why cudaGraphLaunch could return +// success while a stale code was still waiting for the next kernel to read it. +struct FakeLatch { + bool latched = false; + int clears = 0; + // Arming flags: which fake operation is made to fail on its next call. + bool fail_signature = false; + bool fail_instantiate = false; + bool refuse_update = false; + bool fail_destroy = false; + bool fail_launch = false; +}; + +FakeLatch* g_latch = nullptr; + +void Latch() { g_latch->latched = true; } + +// The policy vt/graph_dedup_latch.h installs at every entry point. +struct CountingRuntime { + static void ClearLatchedError() { + ++g_latch->clears; + g_latch->latched = false; + } +}; + +// Stands in for the very next `kernel<<<>>>(...); Check(cudaGetLastError())` the engine +// runs after the capture path returns. This is the caller that #1184 blamed. +void NextUnrelatedKernelLaunch() { + if (g_latch->latched) { + throw std::runtime_error("greedy_argmax launch: invalid device function"); + } +} + +// --- the fake device, one that latches on every failure ----------------------------- +struct LatchGraph { + std::string sig; + int id = 0; +}; +struct LatchExec { + int id = 0; +}; + +std::vector g_execs; + +std::string RawSignature(void* raw_graph) { + if (g_latch->fail_signature) Latch(); + return static_cast(raw_graph)->sig; +} + +void* RawInstantiate(void* raw_graph) { + if (g_latch->fail_instantiate) { + Latch(); + return nullptr; + } + auto* exec = new LatchExec{static_cast(raw_graph)->id}; + g_execs.push_back(exec); + return exec; +} + +bool RawUpdate(void* exec, void* raw_graph, std::string* detail) { + if (g_latch->refuse_update) { + Latch(); + if (detail != nullptr) *detail = "err=1 result=1"; + return false; + } + static_cast(exec)->id = static_cast(raw_graph)->id; + return true; +} + +void RawDestroyExec(void* exec) { + if (g_latch->fail_destroy) Latch(); + delete static_cast(exec); +} + +void RawDestroyGraph(void* raw_graph) { + if (g_latch->fail_destroy) Latch(); + (void)raw_graph; +} + +void RawLaunch(void* exec, void* stream) { + (void)stream; + if (g_latch->fail_launch) { + Latch(); + // Exactly the shape of the production Launch: the failure is reported by throwing, + // which means the clear has to survive an unwinding stack. + VT_CHECK(false, "graph dedup: graph launch failed"); + } + (void)exec; +} + +vt::GraphDedupOps GuardedTable() { + return vt::graph_dedup_latch::MakeLatchGuardedOps(); +} + +// =================================================================================== +// Part 2 — the structural signature walk. +// =================================================================================== + +struct SigGraph; + +struct SigNode { + std::string payload; + const SigGraph* child = nullptr; + bool payload_fails = false; +}; + +// An edge endpoint that is deliberately NOT a member of the graph reporting it, so the +// `edge?` degradation escape has something to fire on. +const SigNode kAlien{"alien", nullptr, false}; + +struct SigGraph { + std::vector nodes; + // Index pairs into `nodes`; -1 on either side means "report the alien node instead". + std::vector> edges; + bool nodes_fail = false; + bool edges_fail = false; +}; + +struct FakeSigRuntime { + using Graph = const SigGraph*; + using Node = const SigNode*; + + static bool GetNodes(Graph graph, std::vector* out) { + if (graph->nodes_fail) return false; + *out = graph->nodes; + return true; + } + + static bool GetEdges(Graph graph, std::vector* from, std::vector* to) { + if (graph->edges_fail) return false; + from->clear(); + to->clear(); + for (const auto& edge : graph->edges) { + from->push_back(edge.first < 0 ? &kAlien : graph->nodes[static_cast(edge.first)]); + to->push_back(edge.second < 0 ? &kAlien : graph->nodes[static_cast(edge.second)]); + } + return true; + } + + static bool AppendNodePayload(Node node, std::string* out, Graph* child) { + if (node->payload_fails) return false; + out->append(node->payload); + *child = node->child; + return true; + } +}; + +std::string SignatureOf(const SigGraph& graph) { + return vt::graph_dedup_sig::Signature(&graph); +} + +} // namespace + +// =================================================================================== + +TEST_CASE("every guarded operation consumes the latch its own failure set") { + FakeLatch latch; + g_latch = &latch; + const vt::GraphDedupOps ops = GuardedTable(); + LatchGraph graph{"sig", 1}; + + // Each arm: arm the failure, drive the operation through the TABLE (never the raw + // function), then stand in the shoes of the next unrelated kernel. + SUBCASE("signature") { + latch.fail_signature = true; + ops.signature(&graph); + } + SUBCASE("instantiate") { + latch.fail_instantiate = true; + CHECK(ops.instantiate(&graph) == nullptr); + } + SUBCASE("update") { + latch.fail_instantiate = false; + void* exec = ops.instantiate(&graph); + latch.refuse_update = true; + std::string detail; + CHECK_FALSE(ops.update(exec, &graph, &detail)); + // The refusal still reports its reason: clearing the latch must not swallow the + // driver's own explanation, which is what the caller acts on. + CHECK(detail == "err=1 result=1"); + ops.destroy_exec(exec); + } + SUBCASE("destroy_exec") { + void* exec = ops.instantiate(&graph); + latch.fail_destroy = true; + ops.destroy_exec(exec); + } + SUBCASE("destroy_graph") { + latch.fail_destroy = true; + ops.destroy_graph(&graph); + } + + CHECK_FALSE(latch.latched); + CHECK_NOTHROW(NextUnrelatedKernelLaunch()); + CHECK(latch.clears >= 1); + g_latch = nullptr; +} + +TEST_CASE("the clear survives an operation that throws") { + // The production Launch reports a failed cudaGraphLaunch by throwing through + // VT_CHECK. A clear placed after the call in the function body would never run on that + // path; a clear in a destructor does. This is the case that decides between the two. + FakeLatch latch; + g_latch = &latch; + const vt::GraphDedupOps ops = GuardedTable(); + LatchGraph graph{"sig", 1}; + void* exec = ops.instantiate(&graph); + const int clears_before = latch.clears; + + latch.fail_launch = true; + CHECK_THROWS_AS(ops.launch(exec, nullptr), std::runtime_error); + + CHECK(latch.clears == clears_before + 1); + CHECK_FALSE(latch.latched); + CHECK_NOTHROW(NextUnrelatedKernelLaunch()); + + latch.fail_launch = false; + ops.destroy_exec(exec); + g_latch = nullptr; +} + +TEST_CASE("a refused fold leaves nothing latched for the next kernel") { + // #1184 in the shape a CPU tier can hold. The registry's REFUSED probe is normal + // operation -- it is what "the signature is only a lookup key" means -- so the whole + // safety argument for this feature depends on a refusal being free of side effects + // outside the registry. Before the fix the refusal latched, nothing consumed it, and + // the next unrelated kernel launch reported it as its own failure. + FakeLatch latch; + g_latch = &latch; + vt::GraphDedupRegistry registry(GuardedTable(), nullptr); + + LatchGraph first{"same", 1}; + LatchGraph second{"same", 2}; + + void* handle_a = registry.Register(&first); + CHECK_NOTHROW(NextUnrelatedKernelLaunch()); + + latch.refuse_update = true; // the driver declines to fold the second capture + void* handle_b = registry.Register(&second); + CHECK_NOTHROW(NextUnrelatedKernelLaunch()); + // The refusal did what it is supposed to do: two groups, not one. + CHECK(registry.ExecCount() == 2); + latch.refuse_update = false; + + registry.Replay(handle_a, nullptr); + CHECK_NOTHROW(NextUnrelatedKernelLaunch()); + registry.Replay(handle_b, nullptr); + CHECK_NOTHROW(NextUnrelatedKernelLaunch()); + + registry.Close(); + CHECK_NOTHROW(NextUnrelatedKernelLaunch()); + g_latch = nullptr; +} + +TEST_CASE("no table field holds an address that skips the guard") { + // The guard is only structural if the raw addresses cannot reach the table. Assigning + // `table.signature = &RawSignature` anywhere would compile and would silently restore + // the defect for that one operation, so the distinctness is asserted rather than + // trusted to review. + FakeLatch latch; + g_latch = &latch; + const vt::GraphDedupOps ops = GuardedTable(); + CHECK(ops.signature != &RawSignature); + CHECK(ops.instantiate != &RawInstantiate); + CHECK(ops.update != &RawUpdate); + CHECK(ops.destroy_exec != &RawDestroyExec); + CHECK(ops.destroy_graph != &RawDestroyGraph); + CHECK(ops.launch != &RawLaunch); + g_latch = nullptr; +} + +TEST_CASE("an operation the guarded builder does not wire cannot reach the registry") { + // The backstop for the seventh operation. If GraphDedupOps grows a member and + // MakeLatchGuardedOps is not extended, the field stays null and the registry refuses + // to construct -- so the bypass is not merely discouraged, it does not run. + FakeLatch latch; + g_latch = &latch; + vt::GraphDedupOps partial = GuardedTable(); + partial.update = nullptr; + CHECK_THROWS_AS(vt::GraphDedupRegistry(partial, nullptr), std::runtime_error); + g_latch = nullptr; +} + +// --- the signature walk ------------------------------------------------------------ + +TEST_CASE("the signature is independent of the order the runtime reports nodes in") { + // THE load-bearing property of the walk. cudaGraphGetNodes promises no order, so two + // captures of one topology can arrive with their nodes permuted. Without the Kahn + // re-index they would key differently, fold nothing, and the only observable would be + // the device-side "captured N graphs, deduped to N execs" line -- exactly the silent + // mode the spec records. + const SigNode a{"A,"}; + const SigNode b{"B,"}; + const SigNode c{"C,"}; + + SigGraph forward; + forward.nodes = {&a, &b, &c}; + forward.edges = {{0, 1}, {1, 2}}; + + SigGraph reversed; + reversed.nodes = {&c, &b, &a}; + reversed.edges = {{2, 1}, {1, 0}}; // the same A->B->C chain, reported back to front + + CHECK(SignatureOf(forward) == SignatureOf(reversed)); + CHECK(SignatureOf(forward) == "[A,;B,;C,;]0,1,;1,2,;"); +} + +TEST_CASE("the signature is independent of the order the runtime reports edges in") { + const SigNode a{"A,"}; + const SigNode b{"B,"}; + const SigNode c{"C,"}; + const SigNode d{"D,"}; + + SigGraph one; + one.nodes = {&a, &b, &c, &d}; + one.edges = {{0, 1}, {0, 2}, {1, 3}, {2, 3}}; // a diamond + + SigGraph other; + other.nodes = {&a, &b, &c, &d}; + other.edges = {{2, 3}, {1, 3}, {0, 2}, {0, 1}}; // the same diamond, edges scrambled + + CHECK(SignatureOf(one) == SignatureOf(other)); + CHECK(SignatureOf(one) == "[A,;B,;C,;D,;]0,1,;0,2,;1,3,;2,3,;"); +} + +TEST_CASE("the signature separates topologies that differ only in their edges") { + const SigNode a{"A,"}; + const SigNode b{"B,"}; + const SigNode c{"C,"}; + + SigGraph chain; + chain.nodes = {&a, &b, &c}; + chain.edges = {{0, 1}, {1, 2}}; + + SigGraph fan; + fan.nodes = {&a, &b, &c}; + fan.edges = {{0, 1}, {0, 2}}; + + CHECK(SignatureOf(chain) != SignatureOf(fan)); +} + +TEST_CASE("the signature separates topologies that differ only in a node payload") { + const SigNode a{"A,"}; + const SigNode b{"B,"}; + const SigNode b_other{"B',"}; + + SigGraph one; + one.nodes = {&a, &b}; + one.edges = {{0, 1}}; + + SigGraph other; + other.nodes = {&a, &b_other}; + other.edges = {{0, 1}}; + + CHECK(SignatureOf(one) != SignatureOf(other)); +} + +TEST_CASE("a child graph contributes its own signature, bounded at depth four") { + // The bound is exact on both sides: level 4's payload is emitted and its child is not + // walked. A bound of three would drop "L4," and a bound of five would admit "L5,". + SigGraph level[6]; + SigNode node[6]; + for (int i = 5; i >= 0; --i) { + node[i].payload = "L" + std::to_string(i) + ","; + if (i < 5) node[i].child = &level[i + 1]; + level[i].nodes = {&node[i]}; + } + + const std::string signature = SignatureOf(level[0]); + CHECK(signature.find("L0,") != std::string::npos); + CHECK(signature.find("L4,") != std::string::npos); + CHECK(signature.find("L5,") == std::string::npos); +} + +TEST_CASE("a child graph is walked, not merely noted") { + const SigNode leaf{"LEAF,"}; + SigGraph child; + child.nodes = {&leaf}; + + const SigNode parent_node{"P,", &child}; + SigGraph parent; + parent.nodes = {&parent_node}; + + CHECK(SignatureOf(parent) == "[P,[LEAF,;];]"); +} + +TEST_CASE("each runtime failure degrades the key instead of aborting inside a capture") { + // Five escapes, five exact strings. Degradation is deliberate: a coarse key costs a + // wasted probe, and the probe is the authority, so the walk must never throw from + // inside a stream capture. + const SigNode a{"A,"}; + const SigNode b{"B,"}; + + SUBCASE("the node query fails") { + SigGraph graph; + graph.nodes = {&a}; + graph.nodes_fail = true; + CHECK(SignatureOf(graph) == "nodes?;"); + } + SUBCASE("the edge query fails") { + SigGraph graph; + graph.nodes = {&a, &b}; + graph.edges = {{0, 1}}; + graph.edges_fail = true; + CHECK(SignatureOf(graph) == "edges?;"); + } + SUBCASE("an edge names a node the graph did not report") { + SigGraph graph; + graph.nodes = {&a, &b}; + graph.edges = {{0, -1}}; + CHECK(SignatureOf(graph) == "edge?;"); + } + SUBCASE("the reported edges describe a cycle") { + SigGraph graph; + graph.nodes = {&a, &b}; + graph.edges = {{0, 1}, {1, 0}}; + CHECK(SignatureOf(graph) == "cycle?;"); + } + SUBCASE("a node payload query fails") { + const SigNode broken{"X,", nullptr, true}; + SigGraph graph; + graph.nodes = {&a, &broken}; + graph.edges = {{0, 1}}; + CHECK(SignatureOf(graph) == "[A,;node?;"); + } +} + +TEST_CASE("an empty graph still produces a signature") { + SigGraph graph; + CHECK(SignatureOf(graph) == "[]"); +}