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 81074c7dd..cf121531e 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 | 7 | 2 | 2 | 6 | +| Engine and scheduling | 30 | 6 | 3 | 1 | 3 | 8 | 2 | 2 | 5 | | KV cache and memory | 24 | 7 | 3 | 2 | 3 | 3 | 2 | 0 | 4 | | Parallelism | 6 | 0 | 0 | 0 | 1 | 0 | 0 | 0 | 5 | | Sampling and generation | 15 | 4 | 2 | 0 | 0 | 4 | 0 | 1 | 4 | @@ -48,7 +48,7 @@ forensics: roadmap_v1.md and the parity ledger. | LoRA and adapters | 2 | 0 | 0 | 0 | 0 | 1 | 0 | 0 | 1 | | Long context and attention | 10 | 5 | 0 | 0 | 1 | 0 | 1 | 0 | 3 | | Loading, tokenizer, config | 10 | 3 | 3 | 0 | 0 | 1 | 1 | 1 | 1 | -| **Total** | **162** | **35** | **17** | **5** | **12** | **34** | **8** | **11** | **39** | +| **Total** | **162** | **35** | **17** | **5** | **12** | **35** | **8** | **11** | **38** | ## 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 NINE drivers (count corrected 2026-08-18, [#1179](https://github.com/mudler/vllm.cpp/issues/1179); `9bc4d7f44` recorded eight, missing the DFlash draft graph `src/vllm/model_executor/models/qwen3_dflash.cpp:771,870,1038,1091,1095,1106`) | owed: same-binary A/B proving a deduped replay is byte-identical to a non-deduped one; a capture-count/exec-count log line so the ratio is observable | [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `INVENTORIED` | [#1162](https://github.com/mudler/vllm.cpp/issues/1162) | +| `ENG-CUDAGRAPH-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: bit-exactness vs eager on every migrated model over MORE than one replay, on a real GPU (W2); the host-lifetime contract of `decode-graph-scratch-uaf-2026-07-18.md` enforced AT the seam — D1's INPUT half, making the intermediates a segment reads unavailable to the `DevicePool` free list (W2); the auxiliary-stream auto-join before every segment close (`:353-361`, spec D10), live at `src/vllm/model_executor/models/qwen3_5.cpp:6254-6255,6384` and `src/vllm/model_executor/models/laguna.cpp:2572-2576,2612` (W4, W5). **Delivered by W1** ([#1192](https://github.com/mudler/vllm.cpp/issues/1192)): the reachability mutation (performed; deleting the call site reds `tests/vllm/models/test_qwen3_break_point.cpp` and leaves the unit suite green); the ported SGLang unit cases with their arithmetic chains and post-replay assertions; and the break-function OUTPUT writeback (`replay_fn`/`_copy_output` `breakable_cuda_graph.py:231-235,172-201`, spec D9), whose destination is a `vt::BreakSlot` the seam owns rather than a caller reference it cannot outlive | spec [eng-cudagraph-break.md](specs/eng-cudagraph-break.md) (W0 spike DONE 2026-08-18: the existing `vt` capture vocabulary `include/vt/backend.h:208-222` expresses a SEGMENTED capture with NO new virtual, because `EndCaptureGraph` stores nothing (`src/vt/cuda/cuda_backend.cu:225-232`); a break point is expressible with one `thread_local` capture pointer plus a free function, no compiler and no decorator); **W1 DONE 2026-08-18 ([#1192](https://github.com/mudler/vllm.cpp/issues/1192)): the seam LANDS** — `vt::BreakableGraph`, `vt::GraphCaptureScope` and `vt::GraphBreak` (`include/vt/breakable_graph.h`, `src/vt/breakable_graph.cpp`), the SGLang unit suite ported case for case (`tests/vt/test_breakable_graph.cpp`, 24 cases / 163 assertions, re-derived 2026-08-18 by `ninja test_breakable_graph && ./build/tests/test_breakable_graph`; the recorded 14/81 never re-derived at any head of this branch), and ONE break point registered at the DENSE ATTENTION ENTRY of `Qwen3ForCausalLM` (`src/vllm/model_executor/models/qwen3.cpp`, inside `RunLayer`). **The exit criterion W0 deliberately left open is ANSWERED on a leased GPU:** `cudaStreamEndCapture` then `cudaStreamBeginCapture` on the SAME stream mid-forward with EAGER work between is LEGAL under `cudaStreamCaptureModeThreadLocal` (`src/vt/cuda/cuda_backend.cu:204-206`) — `orin:gpu0` via an `rc` lease, driver 12060, 3 replays with fresh inputs, 0 mismatches, bare zero-work re-begin legal too. G2 reachability is `tests/vllm/models/test_qwen3_break_point.cpp`, which drives the production `Qwen3DenseModel::Forward` with a scope open and counts `num_hidden_layers + 1` segments (mutation: delete the call site ⇒ 1 segment ⇒ RED), and holds G4 in the same case at 500 logits / 0 differing bit for bit. STAGED SLICE, named: the scope and the container are not yet ENTERED from a production step — no driver opens a scope until W2 migrates `Qwen3DenseDecodeGraph` — and the spec's `## Owed` lists it with W2 as owner, alongside the D10 auxiliary-stream auto-join (W4/W5), G5's ROCm/Tenstorrent arms (W3) and G1 on a real GPU (W2). **The capture-failure drain is NOT among them: it landed HERE**, as behaviour (`std::uncaught_exceptions()` compared against the depth recorded at scope entry, so a break function or ordinary model code throwing mid-capture destroys the partial container instead of handing back a forward that reports `captured() == true`) and as three gated arms (tests 13a, 13b, 13c). The spec's `## Owed` strikes the item through and reads DELIVERED in W1; this cell said the opposite until 2026-08-18 because `cba969857` re-derived field 6 alone. Still NO throughput claim; analysis [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `ACTIVE` | `CLAIM-ENG-CUDAGRAPH-BREAK-W1`; [#1163](https://github.com/mudler/vllm.cpp/issues/1163), [#1192](https://github.com/mudler/vllm.cpp/issues/1192), [#1020](https://github.com/mudler/vllm.cpp/issues/1020) | | `ENG-CUDAGRAPH-DIFFUSION` | Capture the LTX-2.5 denoise loop (fixed shapes, many identical iterations — the ideal graph target). **BLOCKED, and the blocker is ours:** the render does almost no device compute to capture | T2 | SGLang enabled BCG on this shape AFTER our pin — LTX-2 H200 two-stage 10.75s->6.90s (`d4be483efb`), SANA 1024px -26% (`6c7498113f`), SANA denoise 0.73->0.457s (`56ef810cad`). Dated events, NOT pinned evidence; their win is mostly PyTorch host tax we do not pay | NO capture at all: `grep` for capture across `src/vllm/model_executor/models/ltx2*.cpp` returns nothing | blocked by [#1024](https://github.com/mudler/vllm.cpp/issues/1024) (GPU util **exactly 0 in 321 of 347 samples**, 1.00 core of 20 held for 17+ min after staging), [#1007](https://github.com/mudler/vllm.cpp/issues/1007) (VAE decode has no device arm), [#1087](https://github.com/mudler/vllm.cpp/issues/1087) (**57-66% of wall** is ONE resolution-CONSTANT serial host phase), [#1010](https://github.com/mudler/vllm.cpp/issues/1010) (no phase-boundary log). Decision point is a MEASUREMENT of GPU-busy vs wall once device-resident, not an implementation | [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `INVENTORIED` | [#1164](https://github.com/mudler/vllm.cpp/issues/1164) | | `ENG-BATCH-INVARIANT` | Opt-in deterministic execution across scheduler batch sizes (`VLLM_BATCH_INVARIANT=1`): batch-invariant matmul/norm/attention/collectives plus persistent-scheduler NVFP4; production default remains off | T1 | default/env `vllm/envs.py:89,576-578`; initialization `vllm/v1/worker/gpu_worker.py:1262`; NVFP4 dispatch `csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_sm120_kernels.cu:212-220`; suite fixture `tests/v1/determinism/conftest.py:9-12`; operator/e2e `tests/v1/determinism/test_nvfp4_batch_invariant_scaled_mm.py`, `tests/v1/determinism/test_nvfp4_batch_invariant.py` @ `702f481` | - | [W3-C3R executed contract](specs/nvfp4-persistent-plan-cache.md#w3-c3r-batch-shape-localization-and-gate-correction-2026-07-13): production-default ours and vLLM both change outputs across batch shapes; no local opt-in implementation is claimed | `planned: specs/batch-invariant-execution.md` | `INVENTORIED` | - | diff --git a/.agents/issue-index.md b/.agents/issue-index.md index 16c0a8390..5fc4b1cd8 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -366,6 +366,7 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1179](https://github.com/mudler/vllm.cpp/issues/1179) | `ENG-CUDAGRAPH-BREAK` | The hand-rolled decode-graph driver count recorded in `9bc4d7f44` is **eight** and is actually **nine**, and the row it feeds was framed as coverage-only when it is also correctness. The ninth is the DFlash draft graph, file-local with no header declaration, at `src/vllm/model_executor/models/qwen3_dflash.cpp:771,870,1038,1091,1095,1106` — its own `int g_state = 0` three-state machine (`:771`), its own `VT_DFLASH_GRAPH` kill switch (`:870`) instead of the `VLLM_CPP_CUDAGRAPH` the six batched drivers read, its own invalidate-on-block-width-change (`:1038-1047`) and its own `try { EndCaptureGraph(); } catch (...) {}` drain (`:1106`). The eight-count is stated in four places, all corrected here: [`sglang-breakable-cuda-graph.md`](specs/sglang-breakable-cuda-graph.md) §4 and `## Owed`, [`.agents/engine-matrix.md`](engine-matrix.md) rows `ENG-CUDAGRAPH-BREAK` and `ENG-CUDAGRAPH-DEDUP` ("times eight drivers", which sizes #1162's signature table), and [`.agents/roadmap_v1.md`](roadmap_v1.md) track `C12`. The reframing is the substantive half: `ENG-CUDAGRAPH-BREAK` was recorded as a COVERAGE row, and the duplication has already cost a SHIPPED model its decode graph. `src/vllm/model_executor/models/qwen3.cpp:961-986` declines the decode graph outright whenever the asynchronous device-token mirror is live, on its own measured battery — `depth-1, graph ON PASS 78/78`; `depth-2, graph OFF PASS 82/82`; `depth-2, graph ON FAIL, slots 1-3 degenerate` — because `Step()` replays against the HOST `input.token_ids` and the combine has patched the DEVICE ids. The comment names the real fix as reading the identifiers at replay time from a stable device buffer, and that fix exists, in exactly one sibling driver, as `StepDevInputs` (`src/vllm/model_executor/models/qwen3_5.cpp:3894`): `grep -c StepDevInputs` returns 41 lines there and 0 in each of `qwen3_moe.cpp`, `qwen3.cpp`, `deepseek_v2.cpp` and `voxtral.cpp`. One capability, written once, unavailable to four models, with a live mitigation standing in its place. This does NOT weaken the framing rule that `ENG-CUDAGRAPH` established: the row still makes no throughput claim, and the prefill refutation (GB10 3.8% host-idle between launches, GPU-busy >96%, 27B prefill gap 92.5% non-GEMM glue) stands unchanged. Coverage AND correctness, never speed. Fixed in flow with the [`eng-cudagraph-break.md`](specs/eng-cudagraph-break.md) review repair ([#1163](https://github.com/mudler/vllm.cpp/issues/1163)) | record | | [#1181](https://github.com/mudler/vllm.cpp/issues/1181) | `FIX-READ-F32-SCALAR-GUARD` | `ReadF32Scalar` (`src/vllm/model_executor/models/qwen3_5_weights.cpp:312-318` @ `ab6e65216`) bounds its input with `t.data != nullptr && t.nbytes >= sizeof(float)`, a LOWER bound, and then `memcpy`s four bytes into a `float`. Two silent wrong-value paths follow and neither fails: an ARRAY is reduced to element 0, so a block-wise FP8 scale grid of shape `[ceil(N/128), ceil(K/128)]` passes and stands in for the whole weight (measured under [#1166](https://github.com/mudler/vllm.cpp/issues/1166) on `Qwen/Qwen3.8-27B-FP8` @ `017b9c7af6b5689d5dd426a76e0bc077eb5ca20a`, `q_proj.weight_scale_inv` is `[96, 40]`), and ANY dtype is reinterpreted, since that same tensor is `BF16` and its four bytes are two bf16 values read as one float. Both return a finite plausible float, so the output is fluent, plausible and wrong, which is what a token gate cannot see. Upstream makes both facts structural rather than optional: a per-tensor scale is a distinct parameter TYPE that asserts `loaded_weight.shape[0] == 1` (`vllm/model_executor/parameter.py:260-272,304-309` @ `555967922`, plus the `_assert_and_load` shape assert at `:93-96`), the slot is allocated `torch.float32` so a narrow on-disk dtype is VALUE-converted rather than reinterpreted (`utils/fp8_utils.py:1276`), and the declared strategy TENSOR/CHANNEL/BLOCK picks the parameter type before a byte is read (`compressed_tensors/schemes/compressed_tensors_w8a8_fp8.py:63,128`). The AUDIT corrects the issue's own framing twice. The 27 grep hits across five files are 5 definitions, 20 call sites and 2 comment references, and both counts are short: `ReadCtF32Scalar` (`include/vllm/model_executor/models/dense_weight_loaders.h:376`) is a SIXTH copy of the same defect under another name, reached from a SIXTH model file (`src/vllm/model_executor/models/qwen3_weights.cpp:100,126-128` through `LoadCtNvfp4W4A16`). Of the six, three check nothing, `LnReadF32Scalar`/`ShReadF32Scalar` check dtype but not count, and only `nemotron_h_weights.cpp:557-573` is correct, which makes it the model the shared guard generalizes. No call site legitimately passes a multi-element or non-F32 tensor, and every existing fixture emits rank-0 or `{1}` `F32`, so nothing in the tree needed the leniency. It is NOT merely latent: `dense_weight_loaders.h:73-74` and `docs/BENCHMARKS.md:52` both record `unsloth/Qwen3.6-27B-NVFP4` @ `ccdaab7e` as FP8 W8A8 throughout with BF16 PER-OUTPUT-CHANNEL scales, and `LoadAttnDense` branches on the weight dtype alone (`qwen3_5_dense_weights.cpp:478-480`), so those projections enter the per-tensor arm and hit both defects at once under the tensor name the loader actually asked for, with no misspelling to stop them. Fixed in flow by one `dense_loaders::ReadF32Scalar(get, name)` that refuses `numel != 1` naming the shape, refuses a non-`F32` dtype naming the dtype, and requires exactly four readable bytes, with the other five copies deleted onto it and `nemotron_h`'s `Loader`-based twin kept as the one tracked exception. A narrow dtype is refused rather than converted, because a one-element BF16 scale has never been read correctly here and the BF16 layout that IS shipped is per-channel, which the count check refuses first. Per-channel FP8, block-wise FP8 and any explicit narrow-dtype conversion stay owed. Spec [`read-f32-scalar-guard.md`](specs/read-f32-scalar-guard.md) | bug | | [#1185](https://github.com/mudler/vllm.cpp/issues/1185) | `ENV-ORACLE-WHEEL-IN-LEASE` | The pinned vLLM oracle BUILDS, installs, imports and sees the GPU inside an `rc` lease on `dgx:gpu0`, measured 2026-08-18, which falsifies the `nvcc` clause four records carried. [`lease-runtime-staging.md`](specs/lease-runtime-staging.md) said the oracle "needs `nvcc`, which the worker still lacks", and `.agents/environment.md`, [`mtp-k-gt-1.md`](specs/mtp-k-gt-1.md) and [`gpu-lease-methodology.md`](specs/gpu-lease-methodology.md) each derived a blocker from it. The build job (`buildvllm.sh`, staged sha256 prefix `15e140d41f44e7c2`) asserted the checkout against the pin BEFORE compiling, printing `PIN CONFIRMED` at `5559679229bc961848b121ccdeaa8fa5d79bec98` and aborting otherwise, took `nvcc` from the toolkit row `MODEL-NEMOTRON-H-ABI-A3-E2E` staged (`NVCC_RC=0`, CUDA `release 13.3, V13.3.73`) and produced `WHEEL_RC=0`, `PERSIST_RC=0` and a 434 MiB `vllm-0.1.dev1+g555967922.cu133-cp312-cp312-linux_aarch64.whl`, sha256 `7c58b339741a288fbb313f4f5196c9c92a9e3b3c3ebe2ea970b0ff50bb9bcba4`. The identity job (`oracleenv.sh`, prefix `6119f5223f5d818c`) asserted from `cd /`, outside any source tree: `vllm.__version__ = 0.1.dev1+g555967922`, `IDENTITY_RC=0`, `cuda True NVIDIA GB10`, `CUDA_RC=0`. SCOPE, and it carries the same weight as the result: RUNNING A MODEL IS UNTESTED. Only build, install, import and `torch.cuda.is_available()` are measured, and [`mtp-k-gt-1.md`](specs/mtp-k-gt-1.md) records that the last time an oracle reached this far it consumed the host in the step AFTER `torch.compile` and REBOOTED the box, at `gpu_memory_utilization` 0.75 and again at 0.30, so the fraction is not the lever. The version string is an OPEN discrepancy: `.agents/upstream-sync.md` records `vllm_runtime_version = 0.23.1rc1.dev1511+g555967922`, the commit segment matches and satisfies the pin's binding `+g` rule, and the prefix differs because a shallow fetch stops `setuptools_scm` counting commits since the last tag, so a full-string gate needs a deeper fetch or a recorded pretend-version. The venv is NOT staged, because that job was killed at a 90-minute ceiling mid-copy and its partial tree was removed, so only the WHEEL is durable. Four staging walls, all artifacts of the NAS rather than of CUDA: `cp -a` preserves `file_mode=0664` so `nvcc` exited 126. CIFS `nounix` stores no symlink so `include` and `lib64` vanished and CMake reported `Could NOT find CUDA (missing: CUDA_INCLUDE_DIRS CUDA_CUDART_LIBRARY) (found version "13.3")`, naming the version and denying the toolkit in one line. 32 library links `libfoo.so` and `libfoo.so.MAJOR` had to be rebuilt because only the `libfoo.so.X.Y.Z` real files survived. And `markupsafe` existed as a dist-info with NO package files from a `pip --target` killed at a 35-minute ceiling, so Marlin codegen died on `ModuleNotFoundError`. The `rc` worker container is REUSED between jobs, so a repair inside a staging branch is skipped on the next run (`nvcc already in place`) and an environment repair must be unconditional and assert its postcondition. CONSEQUENCE for the rows #1129 blocked, [#1003](https://github.com/mudler/vllm.cpp/issues/1003), [#915](https://github.com/mudler/vllm.cpp/issues/915), [#821](https://github.com/mudler/vllm.cpp/issues/821) and [#81](https://github.com/mudler/vllm.cpp/issues/81): UNBLOCKED FOR THE BUILD STEP and STILL BLOCKED FOR A MODEL RUN. None can take a measurement until a model run is demonstrated. Job details, walls and non-claims in [`oracle-wheel-in-lease.md`](specs/oracle-wheel-in-lease.md) | verification | +| [#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 | | [#1192](https://github.com/mudler/vllm.cpp/issues/1192) | `ENG-CUDAGRAPH-BREAK` | W1 of the break-point capture seam: `vt::BreakableGraph`, `vt::GraphCaptureScope` and `vt::GraphBreak` land in `include/vt/breakable_graph.h` + `src/vt/breakable_graph.cpp`, with the SGLang unit suite ported case for case and ONE break point registered on ONE model. **The stage's exit criterion was answered first, on a leased GPU, because W0 deliberately did not answer it**: `cudaStreamEndCapture` followed by `cudaStreamBeginCapture` on the SAME stream mid-forward with EAGER work between them is LEGAL under `cudaStreamCaptureModeThreadLocal`, our backend's mode (`src/vt/cuda/cuda_backend.cu:204-206`) — measured on `orin:gpu0` through an `rc` lease, driver 12060, three replays with fresh inputs and 0 mismatches, plus the bare zero-work re-begin (`breakable_cuda_graph.py:370-374`). The first probe REFUSED and that refusal was the probe's, not CUDA's: `dlsym` on the bare symbol name binds the LEGACY v1 entry points, which are not capture-aware, so `cuMemcpyDtoDAsync` returned `CUDA_ERROR_STREAM_CAPTURE_UNSUPPORTED` and `cuStreamBeginCapture` (v1) silently took no capture-mode argument — a broken instrument failing toward a verdict about the code. Bound by exact versioned name, the criterion holds. The break point is the DENSE ATTENTION ENTRY of `Qwen3ForCausalLM` (`src/vllm/model_executor/models/qwen3.cpp`, inside `RunLayer`), which is vLLM's boundary and not one this row invents (`splitting_ops` defaulted to the attention family, `vllm/config/compilation.py:517,764-772,1145` @ `5559679229`) registered in SGLang's form because we have no compiler to match operation names against (`layers/radix_attention.py:256` @ `f63458b5be`). The G2 reachability gate is `tests/vllm/models/test_qwen3_break_point.cpp`: it drives the production `Qwen3DenseModel::Forward` with a scope open and counts `num_hidden_layers + 1` segments, so deleting the one call site takes it to 1 and turns the gate RED — a unit test that constructs the container by hand could not have asked that. It also holds G4 in the same case: the logits are compared BIT FOR BIT against the unscoped forward, 500 values, 0 differing. **STAGED SLICE, named rather than implied:** `GraphCaptureScope` and `BreakableGraph` are not yet ENTERED from a production step, because no driver opens a scope until W2 migrates `Qwen3DenseDecodeGraph`; the break point itself is on the production path on every forward and takes the pass-through arm. Listed under `## Owed` in [`eng-cudagraph-break.md`](specs/eng-cudagraph-break.md). The capture-failure drain is NOT owed: it landed here, as behaviour and as three gated arms. The destructor `catch` alone covered a throwing `EndCaptureGraph` and nothing else, so a break function or ordinary model code throwing mid-capture left a container reporting `captured() == true` over a PARTIAL forward; comparing `std::uncaught_exceptions()` against the depth recorded at scope entry is what tells those apart, and tests 13a, 13b and 13c gate the three arms. Still owed by later stages and NOT claimed here: the auxiliary-stream auto-join before every segment close (spec D10, live at `qwen3_5.cpp:6254-6255,6384` and `laguna.cpp:2572-2576,2612`, which W4 and W5 migrate), the non-capturing-backend arm on ROCm and Tenstorrent (G5), G1 bit-exactness on a real GPU over more than one replay, and one residual the fresh review surfaced: an exception CAUGHT INSIDE the capture scope leaves the segment closed and the rest of the forward uncaptured while `captured()` stays true, which the drain cannot see because nothing is unwinding at scope exit (spec `## Owed`, W2). No throughput claim is made anywhere in this change: the prefill refutation stands (GB10 3.8% host-idle between launches, GPU-busy above 96%, 27B prefill gap 92.5% non-GEMM glue) | feature | | [#1189](https://github.com/mudler/vllm.cpp/issues/1189) | `VT-QUANT-FP8-GROUP` | Block-wise (128x128) FP8 so `Qwen/Qwen3.8-27B-FP8` runs instead of being refused. `weight_block_size` appears nowhere in `src/` or `include/`, so `469f38395` refuses the arm by name (#1166). Six independently landable milestones; M1 lands here. M1 is `vt::QuantFp8Group`, the dynamic per-token per-group activation quant, CPU and CUDA. The numerics mirror the kernel that ACTUALLY EXECUTES on a CUDA-alike platform with a contiguous input, which is the C++ custom op at `csrc/libtorch_stable/quantization/w8a8/fp8/per_token_group_quant.cu:42-96` and NOT the Triton kernel at `fp8_utils.py:95-150`: `fp8_utils.py:635-650` calls the former and returns before the latter. The two arms differ, and the difference is measurable rather than cosmetic. The CUDA kernel divides twice, `local_absmax / max_8bit` at `:68` and `static_cast(src) / y_s` at `:85`; the Triton kernel multiplies by `(1.0 / fp8_max)` at `fp8_utils.py:145` under a comment that names the 1-ULP gap. Upstream's own test tolerates the gap with `rtol=0.15` (`test_block_fp8.py:112-114`), so a value comparison cannot tell the two apart and only a byte comparison against a spelled-out reference can. `eps` is the reduction's INITIAL value (`:47`), not a post-clamp, which is what keeps an all-zero group from dividing by zero. Scope refused here and owed to later milestones: the block-scaled GEMM (M2), `Fp8BlockWeight` and the loader (M3), `Fp8BlockLinearMethod` and the Qwen3.5 wiring (M4), the mainloop-scaled CUTLASS kernel and the column-major/TMA-aligned scale layouts (M5), merged `gate_up`/QKV (M6). M1 lands UNREACHED: no production entry point dispatches `vt::QuantFp8Group` at its merge commit, M4 owns the wiring, and `.agents/specs/vt-quant-fp8-group.md` lists it under `## Owed`. The CUDA arm compiles and its on-hardware leg is owed too, because the row took no GPU lease by design: the CPU arm is the gateable one | feature | | [#1190](https://github.com/mudler/vllm.cpp/issues/1190) | `ENV-AGNOSTIC-CAMPAIGN` | One operator's hosts, share paths and addresses are written into 227 tracked files where a `.env` placeholder belongs, so a second developer who follows the protocol documents is told to reach a box on one home network. Re-derived at `fd64c76ee` with `git grep -cIE 'dgx\.casa\|nas_share\|192\.168\.\|thor:gpu0'`: `dgx.casa` 203 files, `nas_share` 32, `192.168.` 24, `thor:gpu0` 11. The mechanism was already there and almost unused: `.env.example` is tracked, `.env` and `.agents/developer-preferences.md` are ignored, `scripts/agent-onboard.py --env-set` already records one answered value and refuses an undeclared key, and `${VLLM_SOURCE}` and `${GPU_LOCK}` already resolve in 59 and 50 files while `${DEVICE_ARCH}`, `${DEVICE_TOOLKIT_ROOT}` and `${DEVICE_COMPILER}` resolve in none. This SCOPING row lands the rule and the mechanism, not the sweep. The rule is a substitution test: replace the literal with a second developer's value and if the sentence stays true it is CONFIGURATION and becomes `${KEY}`, and if it becomes false it is PROVENANCE and stays literal. That reverses the issue's own ranking, because all 34 hits in the densest guide file `.agents/environment.md` are provenance or named-profile definition and none is configuration, so density does not predict the defect and a blind `sed` would falsify records. Landed: three new keys `GATE_CHECKOUT`, `SHARED_STORAGE_ROOT` and `GATE_DEVICE` derived from the literals that recur and map to no existing key; the create-on-first-use route in `scripts/agent-start.py`, which printed `environment: missing` as a status label and then listed next actions that never mentioned it, so the fallback in practice was a host name copied from a document; the matching obligation in `AGENTS.md`; and one worked example. The worked example `scripts/dgx-bringup.sh` found a live defect rather than a cosmetic one: it defaulted `CUTLASS_DIR` to `$HOME/cutlass_probe` while `.agents/environment.md:389` records `$HOME/cutlass-4.5.0` as mandatory on the same box, and a configure that misses CUTLASS silently drops the sm120a NVFP4 GEMM and FlashAttention-2, which that file measures as moving the SACRED `test_qwen27_paged_engine` from 235/235 to 234/235 with the source untouched, so a stale hard-coded default is a false green. Its new test also caught that `set -a; . ./.env; set +a` over an `.env.example` copy blanks a value the caller exported, so the process-environment contract is now executable. Waves `ENV-AGNOSTIC-W1-TOOLING` through `ENV-AGNOSTIC-W5-LEDGERS` own the sweep, partition all 227 files, and open their own issues. Spec [`env-agnostic.md`](specs/env-agnostic.md) | bug | diff --git a/.agents/specs/eng-cudagraph-dedup.md b/.agents/specs/eng-cudagraph-dedup.md new file mode 100644 index 000000000..04a9e6e7d --- /dev/null +++ b/.agents/specs/eng-cudagraph-dedup.md @@ -0,0 +1,413 @@ +# 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, +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 + +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 the programmatic + `CapturedCount()` / `ExecCount()` accessors behind it, 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 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. + +## 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::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 +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 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 + 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 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`, 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 +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, 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. 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. | +| `.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; `(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` | `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 | + +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`, 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 + 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. + 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 + 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. +- **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 — 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. +- **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. +- **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"` 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 + 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 + +| 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 | +| 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 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 + +- 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 bb01a3dfb..c9edb888f 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -272,6 +272,7 @@ Qwen3.6-27B NVFP4, GB10, whole serving window. | Peak `MemAvailable` drop | 68.35 GiB | 80.66 GiB | 1.180x | **PASS** | | Weight offload, resident device bytes (`ENG-WEIGHT-OFFLOAD` W6) | not measured | not measured | n/a | **BLOCKED**, not pending: unmeasurable on every host we own (GB10 shares one pool, so `cpu_offload_gb` frees nothing). Needs a discrete-GPU rig ([record](../.agents/benchmark-record.md)) | | Disk residency via `--offload-config` (`ENG-RESIDENCY-CONFIG`, [#1110](https://github.com/mudler/vllm.cpp/issues/1110)) | not measured | n/a (no disk tier upstream) | n/a | **PENDING** a GB10 run. The row changes no kernel, dtype or allocation, so it claims no throughput axis; the 370 GiB reproduction through the JSON form is owed ([spec](../.agents/specs/weight-residency-config.md)) | +| Decode-graph executables, `VT_CUDA_GRAPH_DEDUP` on vs off (`ENG-CUDAGRAPH-DEDUP`) | not measured | n/a | n/a | **PENDING**, not blocked: needs a leased CUDA box for the same-binary A/B (token-identity first, then the exec-count ratio) ([spec](../.agents/specs/eng-cudagraph-dedup.md)) | 35B steady-serving PSS is 3.53 GiB against vLLM's 13.3 GiB after the routed-expert host mirror is freed once the device Marlin resident is built. diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index 6857087fa..5ca79b2f6 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -132,6 +132,7 @@ portable/reference path. In normal operation leave them unset. | `VT_ASYNC_DEVICE_MIRROR` | on | The CUDA device-resident sampled-token mirror (ENG-ASYNC-SCHED W4): the async serving loop's sampled ids stay on the device instead of round-tripping the host. On a DISCRETE GPU this removes the host-fallback main-stream `Synchronize`. On the INTEGRATED GB10 it moves `last_sampled_tokens` off the host array `update_states`' condense reorders (a drain move, not a removal). DEFAULT ON since the 2026-08-06 correctness flip (ROW-SERVE-ASYNC-LLM): the OFF path's host-array combine writes `step.input_token_ids` on the main queue while the decode graph reads it on the CPU without a sync, so async batch-1 greedy decode nondeterministically degenerates into token-0 garbage; the mirror routes the ids into the embed on-queue and fixes it. `=0` is the rollback to the (racy) host-array path. Speed-neutral (c16 0.999x). No effect on CPU or the sync `LLMEngine` | | `VT_ASYNC_EXECUTOR` | off (opt-in) | `=1` enables Option A: the decode-graph per-step input H2D staged OUT of the captured replay (the c16/c32 overlap unlock, ENG-ASYNC-SCHED). The Qwen3.5 MoE/dense decode-graph drivers give each padded-size slot PERSISTENT device input buffers the captured graph reads and PINNED host staging; per step the input H2D is enqueued on the main queue BEFORE `ReplayGraph` and an input-staged event is recorded right after it, so the next same-slot Refresh waits only that tiny copy, never the GPU tail (the faithful vLLM `_prepare_input_ids`/`synchronize_input_prep` structure, `states.py:64`). The 2-slot parity ring is retained (the depth-2 loop enqueues sample(i-1) after forward(i), so persistent logits must double-buffer). The runner skips the depth-2 pre-forward `Synchronize` whenever the previous step's logits are a non-owning graph-slot view. Default OFF routes through the single-slot baked-H2D driver with the drain intact — byte-identical to production. No effect on CPU or the sync `LLMEngine`. `=1` roughly doubles the captured decode-graph memory (logits-dominated) and adds small pinned host + persistent device input buffers per slot | | `VLLM_CPP_CUDAGRAPH` | on (CUDA) | Eager launches instead of a captured CUDA graph | +| `VT_CUDA_GRAPH_DEDUP` | off (opt-in) | `=1` folds captured graphs that share a node topology onto ONE graph executable, re-pointing it with `cudaGraphExecUpdate` / `hipGraphExecUpdate` instead of instantiating a second (`ENG-CUDAGRAPH-DEDUP`, [#1162](https://github.com/mudler/vllm.cpp/issues/1162)). Today a model holds one executable per padded decode bucket — 7 at `max_num_seqs=32`, 11 at 64 — times eight capture drivers. A MEMORY and capture-time change, not a throughput one: a deduped replay launches the same nodes, and each candidate fold is probed with the real driver update on a throwaway executable before it is honoured, so a capture the driver refuses simply keeps its own executable. It logs `vt graph dedup: captured N graphs, deduped to G execs` per capture so the ratio is readable. **Default OFF and it stays off until measured:** a workload that alternates padded buckets every step pays one `cudaGraphExecUpdate` per switch, and the device byte-identity A/B that would price that is still owed. Unset leaves the capture path byte-identical to the pre-dedup one | | `VLLM_CPP_DENSE_DECODE_GRAPH` | on (CUDA dense) | Non-graphed dense decode | | `VLLM_CPP_QWEN3_DENSE_DECODE_GRAPH` | **on** | Routes pure-decode steps for the SHARED pure-dense forward (`Qwen3DenseModel`, i.e. Qwen3 / Llama / InternLM3 / Mistral / InternLM2 `ForCausalLM`) through the captured decode CUDA graph. **DEFAULT-ON since `QUANT-CT-MXFP4-MARLIN-STRUCT`** (parity-enabler; banks ~+1.3% TPOT@c8 by collapsing the eager inter-kernel launch gap). Token-exact with eager — `=0` opts out to the byte-identical eager decode; gated dgx SACRED `test_qwen3_paged_engine` 184/184 (graph ON==OFF, Qwen3-0.6B near-tie + 4B), async 82/82, Qwen3-8B-MXFP4 #44 smoke 3/3, Qwen3-32B-NVFP4A16 142/142. Honors `VLLM_CPP_CUDAGRAPH=0` | | `VT_MOE_FUSED_W13` | on | Runs a dense/shared MLP gate+up PAIR as ONE Marlin GEMM over the N-concatenated `[2I,H]` operand + `SiluAndMul` (vLLM's merged `gate_up_proj` structure), instead of two split GEMMs + `MoeSiluMul`. **Generalized NVFP4→MXFP4 in `QUANT-CT-MXFP4-MARLIN-STRUCT`** (drops the classic-dense Qwen3-8B-MXFP4 decode from 180→144 marlin GEMM/step = vLLM-structural parity). `=0` is the split A/B fallback. Numerically equivalent to split (the fused fp32 split-K reduce regroups by 1 bf16 ULP; token-exact vs the oracle — #44 fused==split 3/3, unit `test_linear_method` 99.9% bit-exact) | diff --git a/docs/STATUS.md b/docs/STATUS.md index fb60c528c..b88656085 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -1705,6 +1705,8 @@ Gemma4/ROCm env split: public `VT_GEMMA4_EXPERT_VRAM_MB` caps expert LRU in posi `BACKEND-TENSTORRENT-HOST-FREE-FORWARD`: `ACTIVE`: env-gated `VT_TT_HOST_FREE_DECODE` decode-graph capture. Implementer P150 run of Qwen3-0.6B, 80 tokens: 79 replays, no hang, 5.8x vs eager, 22/22 vs the per-step-copy baseline. Default path inert. Operator gate and full-engine golden still owed. A new batch after the first capture is refused. +`ENG-CUDAGRAPH-DEDUP`: `ACTIVE`: env-gated `VT_CUDA_GRAPH_DEDUP` graph-executable dedup — one `cudaGraphExec` per captured TOPOLOGY instead of one per padded decode bucket per model. The shared `vt` registry hashes a captured graph's structure and re-points a single executable with `cudaGraphExecUpdate` on a hit, falling back to a private executable when the driver rejects the update. A memory and capture-time change, NOT a throughput change. Default OFF until a leased-GPU same-binary A/B proves a deduped replay token-identical; that A/B and the ROCm compile are owed ([#1162](https://github.com/mudler/vllm.cpp/issues/1162)). + **Platform SELECTION is the one non-additive site, and is now gated.** A platform missing from `CurrentPlatform()`'s hardcoded walk registers and answers correctly but is NEVER selected, with no compiler diagnostic. `test_platform` diff --git a/scripts/check-gate-commands.py b/scripts/check-gate-commands.py index 19a2dc04c..e01af5cff 100755 --- a/scripts/check-gate-commands.py +++ b/scripts/check-gate-commands.py @@ -311,8 +311,18 @@ def audit() -> list[dict]: # reproduction through the JSON form is recorded as OWED rather than skipped # because dgx.casa was unreachable at the SSH layer. Growth, so the set is # re-pinned in the same change. +# 2026-08-18: +ENG-CUDAGRAPH-DEDUP. A NEW row arriving at ACTIVE (issue #1162), +# so it enters GATED_STATES for the first time. Its spec's Gates section names +# `ctest -R test_graph_dedup` and `scripts/agent-preflight.sh`, both of which +# genuinely fail when the row regresses -- the focused suite detected 9 of 9 +# negative mutations of the registry it gates. It also records what is NOT +# claimed and why: the device byte-identity A/B needs a leased CUDA box this +# session did not have, so it is carried under the spec's `## Owed` rather than +# reported as run, and the CUDA leg's compile rests on the `cuda-fat-build` CI +# job. Growth, so the set is re-pinned in the same change. RUNNABLE_BASELINE = frozenset({ "ENG-RESIDENCY-CONFIG", + "ENG-CUDAGRAPH-DEDUP", "SPEC-MTP-K-GT-1", "ATTN-CHUNKED-LOCAL", "SERVE-RECIPE-ARGS", 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..c6776c3a0 --- /dev/null +++ b/src/vt/graph_dedup.h @@ -0,0 +1,301 @@ +// 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 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 +// 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, 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. +// +// 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. +// +// 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_ + +#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); + // 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)); + } + 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); + // 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; + } + 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()); + // 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); + 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_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 new file mode 100644 index 000000000..a38f1e094 --- /dev/null +++ b/src/vt/graph_dedup_runtime.h @@ -0,0 +1,297 @@ +// vllm.cpp original — the CUDA / HIP binding for the graph-executable dedup registry. +// +// 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 -> 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 +// 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 "vt/graph_dedup.h" +#include "vt/graph_dedup_latch.h" +#include "vt/graph_dedup_signature.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 { + +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 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 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 + // 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 +} + +// (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; +} + +// 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; + + // 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)(); } + + 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; + } + return true; + } + + 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; + } + return true; + } + + // 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; + } + } +}; + +// 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) { + return vt::graph_dedup_sig::Signature(static_cast(raw_graph)); +} + +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"); +} + +} // namespace unguarded + +inline const GraphDedupOps& Ops() { + static const GraphDedupOps ops = + vt::graph_dedup_latch::MakeLatchGuardedOps(); + return ops; +} + +} // namespace vt::graph_dedup_rt + +#endif // VT_GRAPH_DEDUP_RUNTIME_H_ 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/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 959157887..d97303bc8 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -1690,6 +1690,14 @@ target_include_directories(test_gdn_prefill_conv PRIVATE ${CMAKE_SOURCE_DIR}/src vllm_cpp_add_test(test_breakable_graph vt/test_breakable_graph.cpp) vllm_cpp_add_test(test_graph_safe_scratch vt/test_graph_safe_scratch.cpp) target_include_directories(test_graph_safe_scratch PRIVATE ${CMAKE_SOURCE_DIR}/src) +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/scripts/test_check_gate_commands.py b/tests/scripts/test_check_gate_commands.py index 78d4cfbd3..2b88dc7ce 100644 --- a/tests/scripts/test_check_gate_commands.py +++ b/tests/scripts/test_check_gate_commands.py @@ -641,6 +641,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() diff --git a/tests/vt/test_graph_dedup.cpp b/tests/vt/test_graph_dedup.cpp new file mode 100644 index 000000000..e7b8f1d42 --- /dev/null +++ b/tests/vt/test_graph_dedup.cpp @@ -0,0 +1,504 @@ +// 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 nine 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 +#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 + // (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 + // 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 + // 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) { + 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}; +} + +// 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; + 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 || + 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; + } + 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("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("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(); + + 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")); + + // 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")); +} 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) == "[]"); +}