diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md index aa81bd159..3ce9df22d 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -23876,3 +23876,255 @@ second driver and a second CUDA release for the `refused=0` rate; landing #1232 any of that, because the key it would default to is still a draft; and [#1283](https://github.com/mudler/vllm.cpp/issues/1283), without which no future A/B on this row can use a ragged-tail workload shape. + +## ENG-CUDAGRAPH-DEDUP W6 — the DEVICE-BYTE measurement: the fold ENGAGES at the shipped bucket set, the saving DOES NOT SURVIVE its own null control, and the driver refuses 73% of probes on TOPOLOGY (2026-08-19, `row/ENG-CUDAGRAPH-DEDUP-BYTES-RESULT`, tested `origin/main` `2c8f53d93`, GB10 sm_121a, #1162 / #1226) + +**This supersedes nothing. It delivers the one measurement W4 and W5 both named as +owed** — the row's saving is a MEMORY saving, every earlier number counted +executables, and nobody had measured bytes. It also removes W5's largest caveat: +[PR #1232](https://github.com/mudler/vllm.cpp/pull/1232) LANDED as `2c8f53d93`, so the +coarse key is on `main` (default OFF inside a default-OFF flag) and this entry measures +a configuration that ships rather than a draft. + +**The decision this run was taken to make is DELIVERED and NEGATIVE: leave +`VT_CUDA_GRAPH_DEDUP` default OFF.** + +### Recipe + +`dgx:gpu0` through an `rc` lease, job `93f783de-228f-47d5-806d-c5b56aa72c3a`, pod +`rc-worker-4b8lj`, `### BYTES START 2026-08-19T04:30:07Z` to +`### DONE_BYTES 2026-08-19T04:57:19Z`. GB10, driver `580.173.02`, persistence Enabled, +boot_id `3fd9745a-d25a-426c-ba3c-97c958a85515` — **the same value at both ends**, so the +box did not reboot mid-series. `loadavg` start `2.16 2.18 2.33`, end `3.17 3.93 3.80`. + +Source: `git archive` of `origin/main` `2c8f53d93`, tar sha256 +`8fc4109b10b5070ae6b531f8ba1c2f3ec227e2d78dad5a7d1d9e41c5b89d0cc9`, asserted by the +harness before extraction. Binary `vllm-bench` sha256 +`be6972682fa7ba2dacbd5d03166314c45e46960875d0b0b61690715a0ce657a7`, copied out of the +build tree before anything ran it. Release build, `VLLM_CPP_CUDA_ARCHITECTURES=121a`, +CUTLASS 4.5.0 sm120a NVFP4 ENABLED, FlashAttention-2 ENABLED, vendored sm_121a Triton +AOT. `nvcc` **13.0.88**. Model Qwen3-0.6B bf16, 28 layers, from +`/workspace/dedup-gate/model`. + +`VT_ASYNC_RUNNER=0` and `VT_DECODE_GRAPH_STATS=1` throughout; the only variables are +`VT_CUDA_GRAPH_DEDUP` and `VT_CUDA_GRAPH_DEDUP_COARSE_KEY`. Twelve cells (two +calibration, ten A/B), all exit 0, zero VOID markers. + +| Workload | `--concurrency` | `--num-prompts` | `--input-len` | `--output-len` | `--max-num-batched-tokens` | captured padded sizes | +|---|---:|---:|---:|---:|---|---| +| W32 | 32 | 32 | 128 | 96 | 64 | 1, 2, 4, 8, 16, 24, 32 — **7 of 7** | +| W64 | 64 | 64 | 256 | 256 | 128 | 1, 2, 4, 8, 16, 24, 32, 40, 48, 56, 64 — **11 of 11** | + +`vllm-bench` sets `max_num_seqs = concurrency`, so these are exactly the shipped bucket +sets of `include/vllm/model_executor/models/decode_graph_sizes.h:32-41`. **This is the +churn W5 could not produce**: W5's three workloads reached 2-3 buckets, and every W4 and +W5 conclusion about probe behaviour was drawn from a process that presented the driver +with one pair. + +### Instruments, and the assertion each one had to pass first + +**A — `nvidia-smi --query-compute-apps=used_memory`**, sampled against the benchmark pid +every 0.3 s. `--query-gpu=memory.used` returns `[N/A]` on this box and was NOT used. + +**B — an `LD_PRELOAD` shim** reading `cudaMemGetInfo` immediately before and after every +`cudaGraphInstantiate`, `cudaGraphExecUpdate` and `cudaGraphExecDestroy`, plus +`cudaGraphGetNodes` for a node count per instantiate. Interposition was asserted at +build time (`vllm-bench: undefined cudaGraphInstantiate refs = 1`) and again by a +preflight cell that had to exit 0, capture a graph AND write its token artifact. + +### Result 1 — THE FOLD ENGAGES, at the bucket set that ships + +```text +w32_coarse_a vt graph dedup: captured 7 graphs, deduped to 3 execs (probes=7 refused=3) +w32_coarse_b vt graph dedup: captured 7 graphs, deduped to 3 execs (probes=7 refused=3) +w32_exact_a vt graph dedup: captured 7 graphs, deduped to 7 execs (probes=0 refused=0) +w64_coarse_a vt graph dedup: captured 11 graphs, deduped to 5 execs (probes=22 refused=16) +w64_coarse_b vt graph dedup: captured 11 graphs, deduped to 5 execs (probes=22 refused=16) +w64_exact_a vt graph dedup: captured 11 graphs, deduped to 11 execs (probes=0 refused=0) +``` + +Both OFF cells of each workload emit no dedup line at all and instantiate 7 and 11 +executables respectively. Token ids are byte-identical across every cell of a workload, +**including both OFF/OFF controls**: W32 all +`ff0db6c6f4cf0eaa9cba7278880b2c24ec5f6602df6e47d16f9f0cea0b15be9d` (11720 B), W64 all +`e1cbf5fcd3843a07787b39fd167fa713084e84a7b6e5a80836e044912298e5d0` (57620 B). Neither +workload reproduces [#1283](https://github.com/mudler/vllm.cpp/issues/1283)'s +ragged-tail shape, and neither hit it. + +### Result 2 — THE DRIVER REFUSES, and it says TOPOLOGY + +| buckets | probes | refused | rate | reason, verbatim | +|---:|---:|---:|---:|---| +| 7 | 7 | 3 | 43% | `vt graph dedup: probe refused a fold (err=910 result=2)` x3 | +| 11 | 22 | 16 | **73%** | `vt graph dedup: probe refused a fold (err=910 result=2)` x16 | + +`err=910` is `cudaErrorGraphExecUpdateFailure`; `result=2` is +`cudaGraphExecUpdateErrorTopologyChanged`. **Every refusal in this run gave that one +reason.** W5's `refused=0` was an artefact of a workload whose buckets only ever shrank, +so only one pair was ever presented. + +The shim's `cudaGraphGetNodes` reading says why false candidates form at all: the decode +graphs are **not one topology**. They come in two node counts, **376 and 404**, mixed +across the bucket set — `w32_off_a` instantiated, in capture order, +`404 404 376 376 404 404 404`, and `w64_coarse_a` saw +`376 376 376 376 376 376 404 404 376 …`. + +**This inverts the hypothesis the coarse key was built on.** Every refusal is about +topology, never about a parameter. A key that cannot see a topology difference therefore +produces MORE false candidates, not more folds, and coarsening it further makes the +refusal rate worse rather than better. A refusal is a cost and never a wrong replay — +the probe is the authority and a refused capture gets its own executable — but it is the +cost that decides the flip. + +### Result 3 — THE BYTES, and why the saving is NOT ESTABLISHED + +**Instrument A, per-process device footprint, tail median (MiB):** + +| workload | OFF a | OFF b | COARSE a | COARSE b | EXACT a | +|---|---:|---:|---:|---:|---:| +| W32 (7 buckets) | 3252 | 3262 | 3262 | 3275 | 3262 | +| W64 (11 buckets) | 9737 | 9737 | 9737 | 9737 | 9737 | + +**No saving is visible at MiB resolution.** W64 is identical to the megabyte in all five +cells. W32's coarse arm reads 10-23 MiB *higher* than its OFF arm. + +**Instrument B, signed sum of `cudaMemGetInfo` free-byte deltas over every +`cudaGraphInstantiate` in the cell:** + +| cell | retained execs | instantiate calls | bytes | MiB | +|---|---:|---:|---:|---:| +| w32_off_a | 7 | 7 | 29,356,032 | 28.00 | +| w32_off_b | 7 | 7 | 31,952,896 | 30.47 | +| **w32_exact_a** (null control) | **7** | **7** | **18,206,720** | **17.36** | +| w32_coarse_a | 3 | 10 | 16,654,336 | 15.88 | +| w32_coarse_b | 3 | 10 | 15,654,912 | 14.93 | +| w64_off_a | 11 | 11 | 43,134,976 | 41.14 | +| w64_off_b | 11 | 11 | 24,014,848 | 22.90 | +| **w64_exact_a** (null control) | **11** | **11** | **43,855,872** | **41.82** | +| w64_coarse_a | 5 | 27 | 35,045,376 | 33.42 | +| w64_coarse_b | 5 | 27 | 33,673,216 | 32.11 | + +**The nominal effect.** W32: OFF mean 29.24 MiB against COARSE mean 15.40 MiB, a nominal +**13.83 MiB**, which is **0.42%** of the 3262 MiB process footprint instrument A +measured. W64: OFF mean 32.02 MiB against COARSE mean 32.77 MiB, i.e. **−0.75 MiB, +nothing, in the wrong direction**, on a 9737 MiB process. + +**And the nominal effect is not established, on four independent grounds:** + +1. **The null control disagrees by as much as the candidate effect.** `EXACT` retains + exactly as many executables as `OFF` (7 and 11) at `probes=0`, so it allocates the + same executables and its true difference from OFF is ZERO. It measured 17.36 MiB + where OFF measured 28.00 and 30.47 — a **10.6-13.1 MiB** disagreement, against a + candidate effect of 13.83 MiB. +2. **An OFF/OFF pair disagrees with itself by more.** W64 OFF a/b are 41.14 and 22.90 + MiB: **18.2 MiB** apart on two runs of one binary on one workload. +3. **One instantiate recorded a NEGATIVE delta.** `w32_exact_a` recorded, in order, + `10,055,680 / -5,165,056 / 2,326,528 / 2,727,936 / 4,153,344 / 2,273,280 / + 1,835,008` bytes: one instantiate left MORE device memory free than it found, which + is only possible if something else in the process released memory inside the same + window. +4. **`cudaGraphExecDestroy` reclaimed nothing, in every cell.** `reclaimed_bytes=0` + everywhere. A destroy that returns no measurable memory says the reading is pool + growth, not per-object cost. + +The mechanism is the same one in all four. Per-instantiate deltas for byte-identical +404-node graphs range from **0 to 10,514,432 bytes**, and 17 of 27 instantiates in +`w64_coarse_a` read exactly `delta_bytes=0`. The driver serves these allocations from a +pool that grows in chunks and does not shrink, so a per-cell total measures POOL GROWTH, +and the coarse arm's throwaway probe executables grow that pool exactly like retained +ones do. That is how an 11→5 fold arrives at no measurable saving. + +**What CAN be priced.** Over the four OFF cells, bytes ÷ retained execs gives **2.08 to +4.35 MiB for one ~390-node decode-graph executable** (4.00, 4.35, 3.74, 2.08). The +harness also printed bytes per node for a reader sizing the same fold on a deeper +checkpoint: `w32_off_a mean_nodes=396.0 mean_bytes_per_node=10590.2`, +`w64_off_a mean_nodes=391.3 mean_bytes_per_node=10022.1`. + +### Result 4 — THE PROBE COST, exact and clock-independent + +| workload | arm | instantiate calls | exec destroys | `cudaGraphExecUpdate` calls | +|---|---|---:|---:|---:| +| W32 | OFF | 7 | 7 | 0 | +| W32 | COARSE | **10** (3 retained + 7 probes) | 10 | 11 | +| W64 | OFF | 11 | 11 | 0 | +| W64 | COARSE | **27** (5 retained + 22 probes) | 27 | 28 | + +At 11 buckets the coarse key runs **2.45x the instantiate calls** and adds 28 update +calls on the capture path, to retain 6 fewer executables. + +**Peak transient memory did NOT double.** In every ON cell the shim's live-bytes trace +peaks at its end value (W32 COARSE `live_bytes_peak=16654336 live_bytes_end=16654336`; +W64 COARSE `35045376 = 35045376`), because `Register` destroys the probe executable +before it returns. The feared "double the peak to save the steady state" trade did not +occur. + +**A replay-time re-point DID occur, and it is arithmetic, not a counter.** `probes=7` +against 11 update calls, and `probes=22` against 28, leaves 4 and 6 updates that are not +probes — the reverse re-point W5 recorded as untested on a device. Those cells exit 0 +with byte-identical ids, so the transitivity assumption in `Replay` did not abort and did +not change a token here. It is stated as ARITHMETIC over two printed totals, because the +registry still counts probe refusals only and a replay-time refusal would abort on +`Replay`'s `VT_CHECK` rather than increment anything. + +### What this run does NOT establish + +- **NO time-based figure is attributable.** The clock pin was REFUSED inside the lease: + `The current user does not have permission to change clocks for GPU 0000000F:01:00.0`, + and the log records `clocks_pinned=0`. The instantiate-wall and update-wall figures in + `bytes.log` are diagnostics and are deliberately not quoted as results anywhere in + this record. +- **nvcc was 13.0.88 here and 13.3.73 for the W5 fold run.** The OFF-vs-ON and + EXACT-vs-COARSE comparisons WITHIN this one binary are valid; this run and W5 are + **not directly comparable to each other**. 13.0.88 is the toolkit the completed W4 + baseline gate ran. +- **`result=2` is one driver (`580.173.02`), one GB10, one toolkit.** A different + refusal rate elsewhere is a different cost, not a different correctness. +- **Only the Qwen3 dense decode driver was exercised**, as in W4 and W5. The other eight + capture drivers, and two models sharing the process-singleton registry, are untested. +- **`VT_ASYNC_RUNNER=0` throughout.** Dedup engages only there, because the async path + captures no decode graph, so the feature remains unreachable on the DEFAULT serving + path ([#1179](https://github.com/mudler/vllm.cpp/issues/1179)). +- **`cudaMemGetInfo` cannot separate an executable's own cost from the pool chunk that + satisfied it.** Every byte figure above inherits that limit. +- **Two summaries of the same rows differ, and neither is wrong.** `run.sh`'s per-cell + `shim_inst … sum_bytes` reducer sums only POSITIVE deltas and therefore reads higher in + the two cells containing a negative one (`w32_exact_a` 23,371,776 against the signed + 18,206,720). The tables above use the SIGNED sum, which agrees with the shim's own + in-process `SUMMARY` counter in every cell. + +### Verdict + +1. **`VT_CUDA_GRAPH_DEDUP` stays default OFF — supported by these numbers, not merely + ungated.** Best nominal saving 13.83 MiB (0.42% of process) at 7 buckets, smaller + than its own null control's disagreement; nothing at 11 buckets; cost is 16 extra + instantiate/destroy pairs and 28 update calls on the capture path. +2. **`VT_CUDA_GRAPH_DEDUP_COARSE_KEY` alone is a NO-OP, not merely unsupported.** + `GraphDedupCoarseKeyEnabled()` (`src/vt/graph_dedup.h:114`) is read only by the + signature builder (`src/vt/graph_dedup_runtime.h:177`), which runs only from + `Register`, which `src/vt/cuda/cuda_backend.cu:237` calls only under + `GraphDedupEnabled()`. With dedup off its sole observable is one stderr line. +3. **Both on — not supported.** Same numbers plus a 73% probe-refusal rate. + +### Next traceable hypotheses — NO CEILING is declared + +- **Find where the 376/404 node split comes from.** The FA-2 split-KV grid is the first + suspect. If a capture exists that fixes the node set across buckets, every refusal in + this run disappears and the fold ratio is the whole bucket set rather than 3/7 and + 5/11. Until then a coarser key is the wrong direction. +- **An instrument that resolves a single 2-4 MiB executable against driver pool + granularity.** `cuMemGetAllocationGranularity`, a driver-pool statistics query, or a + build with pool instrumentation would price the fold directly instead of through a + total whose noise exceeds the effect. +- **A 60-80 layer checkpoint.** Bytes scale with node count and this graph is 376-404 + nodes on a 28-layer 0.6B model. The measured 10.0-10.6 KB per node is the number to + re-run against before the flip is refused permanently for large models. + +### Evidence + +`/mnt/nas_share/rc/dedup-bytes/` — `RESULT.md` and `STATUS.md` (the pre-registered +"what must be true before any number here is quoted" list and the result against it), +`run.sh`, `build.sh`, `memshim.c`, `src-2c8f53d93.tar`, `logs-bytes/bytes.log`, +`logs-bytes/clockpin.log` (the refusal), `logs-bytes/run_.log` for all 12 cells, +`mem/mem_.csv` (the per-call `cudaMemGetInfo` trace with node counts), +`mem/smp_.csv` (the `nvidia-smi` and RSS sampler) and `out-bytes/ids_.json` +(the token artifacts). diff --git a/.agents/engine-matrix.md b/.agents/engine-matrix.md index 9bacf8396..fc3d5b1f8 100644 --- a/.agents/engine-matrix.md +++ b/.agents/engine-matrix.md @@ -60,7 +60,7 @@ forensics: roadmap_v1.md and the parity ledger. | `KV-PREFIX-MATCH-UNIT` | `--prefix-match-unit` (config `prefix_match_unit`): the finest token boundary a prefix-cache hit can land on == the `hash_block_size`/"prefix match unit" the block hasher uses. NEW in 0.26 (absent at the prior `e24d1b24`/0.25.0 pin). For a HYBRID/multi-group model the resolver `resolve_kv_cache_block_sizes` computes `hash_block_size = prefix_match_unit if set else gcd(group_block_sizes)` (scheduler block size = `lcm`), letting matching land FINER than a physical block (e.g. 16/32 tokens inside a 1024-token block) provided every group block size is divisible by it; single-group (dense) models ignore the knob. Backs off to the scheduler block size when no prefix-cache/connector consumer is active or a mamba group diverges from `cache_block_size` (mamba_cache_mode != "align"); throws on a non-divisible unit. **W0 spike + W1 resolver LANDED 2026-07-28 (`CLAIM-PREFIX-MATCH-UNIT`, NOT pushed):** `resolve_kv_cache_block_sizes` ported 1:1 (explicit-parameter signature vs upstream's `VllmConfig`, our config surface is threaded), RED-first unit-gated (default gcd `!=` `=16` override). `PARTIAL`: the config/CLI/ABI field (W2), the scheduler threading of a resolved `hash_block_size != block_size` + mamba partial-tail stop (W3, needs the `KV-BLOCK-POOL` align path that still throws), and the benchmark (W4) are deferred. Default path byte-identical (single-group inert; scheduler still passes `block_size`). | T1 | `vllm/engine/arg_utils.py:696,1222,1940`; `vllm/config/cache.py:56-67`; resolver `vllm/v1/core/kv_cache_utils.py:626-688`; hasher `:691-748`; call site `vllm/v1/engine/core.py:154`; scheduler `vllm/v1/core/sched/scheduler.py:76,268-270,282,312-318`; fine-grained view `vllm/v1/core/single_type_kv_cache_manager.py:683,697` | resolver `src/vllm/v1/core/kv_cache_utils.cpp:638` (`resolve_kv_cache_block_sizes`), decl `include/vllm/v1/core/kv_cache_utils.h`; hash_block_size already plumbed `get_request_block_hasher` `src/vllm/v1/core/kv_cache_utils.cpp:577`; DEFERRED align path throws `src/vllm/v1/core/block_pool.cpp:93,220` (shared with `KV-BLOCK-POOL`) | `tests/vllm/v1/test_prefix_match_unit.cpp:64,88,99,119,129,145,164,186` 8/8 (29 assertions): single-group inert + DCP scale, multi-group default=gcd, `=16` override finer-than-default (RED), finer-than-1024-block, non-divisible throws, no-consumer back-off + connector-alone re-enable, mamba non-align back-off vs align gcd, hasher-granularity RED (coarse 2 vs fine 4 hashes); [parity-ledger.md](parity-ledger.md) | [prefix-match-unit.md](specs/prefix-match-unit.md) | `PARTIAL` | `CLAIM-PREFIX-MATCH-UNIT` | | `ENG-PREEMPT-RECOMPUTE` | FCFS tail preemption with recompute | T0 | `vllm/v1/core/sched/scheduler.py:1142`; `tests/v1/core/test_scheduler.py:930` | `src/vllm/v1/core/sched/scheduler.cpp:102,157`; `src/vllm/v1/core/sched/request_queue.cpp:36` | `tests/vllm/v1/test_scheduler.cpp:247,295`; `tests/vllm/v1/test_request_queue.cpp:91` | `planned: specs/preemption.md` | `ANCHOR-BACKFILL` | - | | `ENG-CUDAGRAPH` | Decode graph capture/replay modes (host-cluster cleanup: capture-size set derived from `max_num_seqs` mirroring vLLM `_set_cudagraph_sizes`; 2026-07-18 graph-baked-scratch use-after-free fix — the 35B c2+ online-serving IMA blocker) | T0 | `vllm/config/compilation.py:53,1319,683-684,1438-1444`; `vllm/config/vllm.py:1667-1770`; `vllm/v1/worker/gpu/cudagraph_utils.py:116`; `tests/compile/test_config.py:122,229` | `src/vt/cuda/cuda_backend.cu:76,97,105`; `include/vllm/model_executor/models/decode_graph_sizes.h`; `src/vllm/model_executor/models/qwen3_5.cpp:3754,3952`; `src/vllm/v1/worker/gpu/runner.cpp:577,597`; graph-safe scratch (retire-on-grow so graph-baked scratch pointers stay valid) `src/vt/cuda/graph_safe_scratch.h`, `src/vt/cuda/cuda_moe_marlin.cu:75`, `src/vt/cuda/cuda_matmul_nvfp4.cu:766`, `src/vt/cuda/cuda_matmul_nvfp4_cutlass.cu:105`, `src/vt/cuda/cuda_matmul_fp8_cutlass.cu:95` | `tests/vt/test_cuda_backend.cpp:98`; `tests/vllm/models/test_decode_graph_sizes.cpp`; `tests/vt/test_graph_safe_scratch.cpp`; explicit 35B gate `tests/parity/test_qwen36_paged_engine.cpp:140` | [blocktable-host-cluster-cleanup.md](specs/blocktable-host-cluster-cleanup.md); [decode-graph-scratch-uaf-2026-07-18.md](specs/decode-graph-scratch-uaf-2026-07-18.md) | `PARTIAL` | **PREFILL capture REFUTED as a lever (2026-08-17, [#1161](https://github.com/mudler/vllm.cpp/issues/1161)).** vLLM's v1 default already captures prefill piecewise (`vllm/config/compilation.py:60-63,615,630` @ `555967922`) and it is in our denominator; SGLang reached the same coverage without `torch.compile` via BCG (`SGLANG-BCG` in [sglang-matrix.md](sglang-matrix.md)). Neither helps us: GB10 2026-07-09 measured prefill GPU-idle-between-launches at **3.8%** with GPU-busy >96% on both arms, and the 27B prefill gap at **92.5% non-GEMM glue GPU work** with the dominant GEMM at +0.17% and attention AHEAD. There are no launch bubbles in our prefill to collapse. Row stays `PARTIAL`; the real residuals are exec dedup ([#1162](https://github.com/mudler/vllm.cpp/issues/1162)) and the break-point seam ([#1163](https://github.com/mudler/vllm.cpp/issues/1163)). Spec [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | -| `ENG-CUDAGRAPH-DEDUP` | Graph-executable dedup: hash each captured graph's topology and re-point ONE `cudaGraphExec` with `cudaGraphExecUpdate` on a signature hit, instead of instantiating one exec per padded bucket per model. A memory and capture-time change, NOT a throughput change — a deduped replay launches the same nodes, and the load-bearing gate is byte-identity rather than a ratio | T2 | vLLM has no analogue (its execs come from `torch.compile`, `vllm/config/compilation.py:60-63,517,615,630` @ `555967922`); secondary oracle SGLang `python/sglang/srt/model_executor/runner_backend/cuda_graph_dedup_mixin.py:27-37,105-179,219-242,258-275,353-358` @ `f63458b5be` ([oracles/sglang.md](oracles/sglang.md)) | W1+W2 landing here behind `VT_CUDA_GRAPH_DEDUP`, default OFF until the device A/B measures the per-switch update cost: a device-agnostic dedup registry shared by both accelerator backends plus one CUDA/HIP ops table written once, wired into `EndCaptureGraph`/`ReplayGraph`/`DestroyGraph`. Baseline it replaces: `src/vt/cuda/cuda_backend.cu:222-232` instantiates a fresh exec per capture and destroys the raw graph, over the 7 (`max_num_seqs=32`) or 11 (64) buckets of `include/vllm/model_executor/models/decode_graph_sizes.h:32-41`, times NINE drivers (count corrected 2026-08-18, [#1179](https://github.com/mudler/vllm.cpp/issues/1179); `9bc4d7f44` recorded eight, missing the DFlash draft graph `src/vllm/model_executor/models/qwen3_dflash.cpp:771,870,1038,1091,1095,1106`) | `tests/vt/test_graph_dedup.cpp` 13/13 cases, 65 assertions, RED-first (written and run against an absent header, and the four cases added by the fresh review of #1178, three of them run against the unfixed source) and gated on every platform via a fake ops table whose launch log makes "the right nodes ran" an observable sequence over MORE than one replay per shape; 13/13 negative mutations detected (9 at implementation, 4 at review repair). That count covers `src/vt/graph_dedup.h` ONLY. `src/vt/graph_dedup_runtime.h` had NO executable coverage on any tier, and [#1184](https://github.com/mudler/vllm.cpp/issues/1184) is what hid in that gap: the file is DESIGNED to see runtime calls fail — a refused `cudaGraphExecUpdate` probe is the feature working — and never consumed the runtime's latched error, so the next unrelated kernel reported the refusal as its own failure and every `VT_CUDA_GRAPH_DEDUP=1` run died 6/6 on GB10 as `greedy_argmax launch: invalid device function` from a launch that had succeeded. Repaired structurally rather than at twelve sites: the clear lives in `ScopedLatchClear`'s destructor (`src/vt/graph_dedup_latch.h`) installed at the six `GraphDedupOps` entry points by `MakeLatchGuardedOps`, the table's only constructor, so no raw function address reaches a field and an unwired seventh operation leaves a null the registry refuses; one line covers CUDA and HIP. The device-free half of the signature walk moved to `src/vt/graph_dedup_signature.h` and is gated by `tests/vt/test_graph_dedup_runtime.cpp` 13/13 cases, 51 assertions, RED-first against the pre-fix guard (22 failed assertions reproducing the production message), 7/7 negative mutations detected — Kahn ordering, topological re-index, sorted edge emission, the depth-4 child bound and the four graph-level escapes. STILL compile-gated only: the five node-payload cases behind the device policy. **DEVICE A/B DELIVERED 2026-08-18 on `dgx:gpu0` (GB10, driver 580.173.02, nvcc 13.0.88, `rc` job f88d484b), and it SPLIT.** Gated commit `72de552c8`, whose four dedup sources are byte-identical to the merged `2a976eb9f` — the row squashed, so the gated tree is not an ancestor of the merge and that sha equality is what carries the claim. CORRECTNESS PASSES: 12/12 cells exit 0, zero `invalid device function` and zero `engine-fatal` in every cell log where the pre-fix head `e4ce5571a` died after exactly one replay, ON replays as often as OFF (60=60, 33=33, 43=43), and `--output-token-ids` is IDENTICAL over 10/10 comparisons with the three OFF/OFF controls passing FIRST and the three workloads hashing to three DIFFERENT values, so the identity is not vacuous. #1184 is closed by this run, because a CPU suite drives a fake runtime and cannot observe the real latched error. THE BENEFIT IS REFUTED for the case this row was filed for: `N == M` in every ON cell — 3 graphs to 3 execs on sizes [24 16 8], 2 to 2 on [16 8], 2 to 2 on [32 24] — with the registry's count CLIMBING 1→1, 2→2, 3→3, so more than one capture reached it and the 1:1 is a measurement rather than the single-capture artefact the first attempt produced. Cause pre-registered before the run and then confirmed, structural rather than a tuning miss: `AppendKernelPayload` hashes (`func`, `gridDim.{x,y,z}`, `blockDim.{x,y,z}`, `sharedMemBytes`) at `src/vt/graph_dedup_runtime.h:121-128` and the memcpy payload hashes the copy extent, so the padded batch dimension sits in the KEY, no candidate group ever forms and `cudaGraphExecUpdate` is NEVER ATTEMPTED. That contradicts this row's own premise — `graph_dedup.h`'s header says the fold is for "two padded batch sizes … the same node topology with different parameters" — and SGLang keys the same fields (`cuda_graph_dedup_mixin.py:105-114`), so whatever folds upstream is not decode buckets either. NO throughput or memory number is recorded: clocks unpinned AND the ON arm allocated exactly as many executables as OFF. Honest gaps: per-shape replay counts are unavailable (the driver prints a TOTAL, so B's ~30-per-shape is arithmetic); the driver's "N captured size(s)" counts SLOTS not captures (A reports 6, emits 3); the container's own cuBLASLt was never re-tested at CUDA 13.0 because the staged cu130 prefix was probed first and worked; only the Qwen3 dense decode driver was exercised. STILL OWED: the default flip, now NOT JUSTIFIED on this evidence rather than merely ungated; a COARSER key that could group two decode buckets at all, which the probe-before-fold design makes a cost question rather than an obviously unsafe one ([#1226](https://github.com/mudler/vllm.cpp/issues/1226), the next traceable hypothesis, deliberately NOT decided by this record); device-tier signature stability/discrimination tests; probing `current_raw` instead of `raws.front()` to retire the update-transitivity assumption; the ROCm compile; a supporting `orin:gpu0` leg, BLOCKED because the Jetson 540.4.0 driver cannot run a CUDA 13 runtime (`cudaGetDeviceCount err=35`); and reaching the feature from the default serving path at all — the async runner captures no decode graph, **W5, THE SAME DAY, CONFIRMED THE HYPOTHESIS THAT NEGATIVE PRODUCED ([#1226](https://github.com/mudler/vllm.cpp/issues/1226) DELIVERED).** Same box, `rc-worker-4b8lj`, boot_id `3fd9745a-d25a-426c-ba3c-97c958a85515` at both ends, GB10, driver `580.173.02`, `### DONE_AB_KEY 2026-08-18T20:58:46Z`, binary sha256 `ca114abb…c772ad` from `b48b51df1` (tar sha256 asserted before extraction). Drop the launch dimensions and the memcpy extents from the key and every bucket folds: `a_coarse` 3 graphs to 2 execs, `b_coarse` 2 to 1, `c_coarse` 2 to 1, each `probes=1 refused=0`, against `probes=0 refused=0` in every EXACT cell. **`probes=0` in the EXACT cells is the direct process-level proof of W4's source-level diagnosis** — with the launch dimensions in the key no candidate group forms and `cudaGraphExecUpdate` is never asked; drop them and it is asked once per fold and ACCEPTED EVERY TIME. The saving W4 recorded as unreachable is reachable via the key. Byte-identity holds on A (five cells, `59ebff4a…`) and C (four cells, `ff205260…`). **Workload B is VOID rather than a pass, and its cause is a NEW DEFECT that is not this row's:** the two `VT_CUDA_GRAPH_DEDUP`-unset control cells DISAGREED (`5973c5a1…` 2638 bytes vs `4cf79230…` 2650 bytes) on one binary, one workload, greedy `--temperature 0 --seed 777` at `--concurrency 16`, 23 s apart — 672 tokens both, so the byte delta is JSON width and not a length; exactly rows 17 and 18 of 21 differ, both mid-decode, both in the ragged tail `21 % 16` leaves. B's `b_off_a == b_exact` and `b_off_a == b_coarse_a` therefore compare against a baseline that does not reproduce itself and are WORTHLESS; only the OFF/OFF control made that visible, and without it B would have read as three more confirmations. Filed [#1283](https://github.com/mudler/vllm.cpp/issues/1283). **Caveats that bound this result:** nvcc was `13.3.73` here and `13.0.88` for the W4 baseline the recorded dgx gate stack names, so the OFF-vs-ON and EXACT-vs-COARSE comparisons WITHIN this binary are valid while this run and that baseline are NOT directly comparable; clocks unpinned (2405 MHz current, 3003 max, 2418 applications) and nothing measured bytes, so NO throughput and NO memory number is claimed or implied; only the Qwen3 dense decode driver was exercised; `refused=0` is ONE driver on ONE hardware and toolkit pair, which is no more a floor than W4's negative was a ceiling; and the coarse key is behind `VT_CUDA_GRAPH_DEDUP_COARSE_KEY`, default OFF, inside a default-OFF flag, on **PR [#1232](https://github.com/mudler/vllm.cpp/pull/1232) which is STILL A DRAFT — nothing on `main` folds today.** **Row stays `ACTIVE`, argued:** not `DONE`, because the fold is unreachable on every shipping configuration and the row's stated MEMORY saving has never been measured in bytes on either key; not `PARTIAL`, because nothing upstream is omitted — the coarse key is our own extension past SGLang, which keys the fields we started from; not `BLOCKED`, because nothing external stops the next step. What is owed is now a DECISION about the default plus the byte measurement and the probe-cost-at-real-churn measurement it needs, and landing #1232 first | [eng-cudagraph-dedup.md](specs/eng-cudagraph-dedup.md); analysis [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `ACTIVE` | `CLAIM-ENG-CUDAGRAPH-DEDUP` ([#1162](https://github.com/mudler/vllm.cpp/issues/1162)) | +| `ENG-CUDAGRAPH-DEDUP` | Graph-executable dedup: hash each captured graph's topology and re-point ONE `cudaGraphExec` with `cudaGraphExecUpdate` on a signature hit, instead of instantiating one exec per padded bucket per model. A memory and capture-time change, NOT a throughput change — a deduped replay launches the same nodes, and the load-bearing gate is byte-identity rather than a ratio | T2 | vLLM has no analogue (its execs come from `torch.compile`, `vllm/config/compilation.py:60-63,517,615,630` @ `555967922`); secondary oracle SGLang `python/sglang/srt/model_executor/runner_backend/cuda_graph_dedup_mixin.py:27-37,105-179,219-242,258-275,353-358` @ `f63458b5be` ([oracles/sglang.md](oracles/sglang.md)) | W1+W2 landing here behind `VT_CUDA_GRAPH_DEDUP`, default OFF until the device A/B measures the per-switch update cost: a device-agnostic dedup registry shared by both accelerator backends plus one CUDA/HIP ops table written once, wired into `EndCaptureGraph`/`ReplayGraph`/`DestroyGraph`. Baseline it replaces: `src/vt/cuda/cuda_backend.cu:222-232` instantiates a fresh exec per capture and destroys the raw graph, over the 7 (`max_num_seqs=32`) or 11 (64) buckets of `include/vllm/model_executor/models/decode_graph_sizes.h:32-41`, times NINE drivers (count corrected 2026-08-18, [#1179](https://github.com/mudler/vllm.cpp/issues/1179); `9bc4d7f44` recorded eight, missing the DFlash draft graph `src/vllm/model_executor/models/qwen3_dflash.cpp:771,870,1038,1091,1095,1106`) | `tests/vt/test_graph_dedup.cpp` 13/13 cases, 65 assertions, RED-first (written and run against an absent header, and the four cases added by the fresh review of #1178, three of them run against the unfixed source) and gated on every platform via a fake ops table whose launch log makes "the right nodes ran" an observable sequence over MORE than one replay per shape; 13/13 negative mutations detected (9 at implementation, 4 at review repair). That count covers `src/vt/graph_dedup.h` ONLY. `src/vt/graph_dedup_runtime.h` had NO executable coverage on any tier, and [#1184](https://github.com/mudler/vllm.cpp/issues/1184) is what hid in that gap: the file is DESIGNED to see runtime calls fail — a refused `cudaGraphExecUpdate` probe is the feature working — and never consumed the runtime's latched error, so the next unrelated kernel reported the refusal as its own failure and every `VT_CUDA_GRAPH_DEDUP=1` run died 6/6 on GB10 as `greedy_argmax launch: invalid device function` from a launch that had succeeded. Repaired structurally rather than at twelve sites: the clear lives in `ScopedLatchClear`'s destructor (`src/vt/graph_dedup_latch.h`) installed at the six `GraphDedupOps` entry points by `MakeLatchGuardedOps`, the table's only constructor, so no raw function address reaches a field and an unwired seventh operation leaves a null the registry refuses; one line covers CUDA and HIP. The device-free half of the signature walk moved to `src/vt/graph_dedup_signature.h` and is gated by `tests/vt/test_graph_dedup_runtime.cpp` 13/13 cases, 51 assertions, RED-first against the pre-fix guard (22 failed assertions reproducing the production message), 7/7 negative mutations detected — Kahn ordering, topological re-index, sorted edge emission, the depth-4 child bound and the four graph-level escapes. STILL compile-gated only: the five node-payload cases behind the device policy. **DEVICE A/B DELIVERED 2026-08-18 on `dgx:gpu0` (GB10, driver 580.173.02, nvcc 13.0.88, `rc` job f88d484b), and it SPLIT.** Gated commit `72de552c8`, whose four dedup sources are byte-identical to the merged `2a976eb9f` — the row squashed, so the gated tree is not an ancestor of the merge and that sha equality is what carries the claim. CORRECTNESS PASSES: 12/12 cells exit 0, zero `invalid device function` and zero `engine-fatal` in every cell log where the pre-fix head `e4ce5571a` died after exactly one replay, ON replays as often as OFF (60=60, 33=33, 43=43), and `--output-token-ids` is IDENTICAL over 10/10 comparisons with the three OFF/OFF controls passing FIRST and the three workloads hashing to three DIFFERENT values, so the identity is not vacuous. #1184 is closed by this run, because a CPU suite drives a fake runtime and cannot observe the real latched error. THE BENEFIT IS REFUTED for the case this row was filed for: `N == M` in every ON cell — 3 graphs to 3 execs on sizes [24 16 8], 2 to 2 on [16 8], 2 to 2 on [32 24] — with the registry's count CLIMBING 1→1, 2→2, 3→3, so more than one capture reached it and the 1:1 is a measurement rather than the single-capture artefact the first attempt produced. Cause pre-registered before the run and then confirmed, structural rather than a tuning miss: `AppendKernelPayload` hashes (`func`, `gridDim.{x,y,z}`, `blockDim.{x,y,z}`, `sharedMemBytes`) at `src/vt/graph_dedup_runtime.h:121-128` and the memcpy payload hashes the copy extent, so the padded batch dimension sits in the KEY, no candidate group ever forms and `cudaGraphExecUpdate` is NEVER ATTEMPTED. That contradicts this row's own premise — `graph_dedup.h`'s header says the fold is for "two padded batch sizes … the same node topology with different parameters" — and SGLang keys the same fields (`cuda_graph_dedup_mixin.py:105-114`), so whatever folds upstream is not decode buckets either. NO throughput or memory number is recorded: clocks unpinned AND the ON arm allocated exactly as many executables as OFF. Honest gaps: per-shape replay counts are unavailable (the driver prints a TOTAL, so B's ~30-per-shape is arithmetic); the driver's "N captured size(s)" counts SLOTS not captures (A reports 6, emits 3); the container's own cuBLASLt was never re-tested at CUDA 13.0 because the staged cu130 prefix was probed first and worked; only the Qwen3 dense decode driver was exercised. STILL OWED: the default flip, now NOT JUSTIFIED on this evidence rather than merely ungated; a COARSER key that could group two decode buckets at all, which the probe-before-fold design makes a cost question rather than an obviously unsafe one ([#1226](https://github.com/mudler/vllm.cpp/issues/1226), the next traceable hypothesis, deliberately NOT decided by this record); device-tier signature stability/discrimination tests; probing `current_raw` instead of `raws.front()` to retire the update-transitivity assumption; the ROCm compile; a supporting `orin:gpu0` leg, BLOCKED because the Jetson 540.4.0 driver cannot run a CUDA 13 runtime (`cudaGetDeviceCount err=35`); and reaching the feature from the default serving path at all — the async runner captures no decode graph, **W5, THE SAME DAY, CONFIRMED THE HYPOTHESIS THAT NEGATIVE PRODUCED ([#1226](https://github.com/mudler/vllm.cpp/issues/1226) DELIVERED).** Same box, `rc-worker-4b8lj`, boot_id `3fd9745a-d25a-426c-ba3c-97c958a85515` at both ends, GB10, driver `580.173.02`, `### DONE_AB_KEY 2026-08-18T20:58:46Z`, binary sha256 `ca114abb…c772ad` from `b48b51df1` (tar sha256 asserted before extraction). Drop the launch dimensions and the memcpy extents from the key and every bucket folds: `a_coarse` 3 graphs to 2 execs, `b_coarse` 2 to 1, `c_coarse` 2 to 1, each `probes=1 refused=0`, against `probes=0 refused=0` in every EXACT cell. **`probes=0` in the EXACT cells is the direct process-level proof of W4's source-level diagnosis** — with the launch dimensions in the key no candidate group forms and `cudaGraphExecUpdate` is never asked; drop them and it is asked once per fold and ACCEPTED EVERY TIME. The saving W4 recorded as unreachable is reachable via the key. Byte-identity holds on A (five cells, `59ebff4a…`) and C (four cells, `ff205260…`). **Workload B is VOID rather than a pass, and its cause is a NEW DEFECT that is not this row's:** the two `VT_CUDA_GRAPH_DEDUP`-unset control cells DISAGREED (`5973c5a1…` 2638 bytes vs `4cf79230…` 2650 bytes) on one binary, one workload, greedy `--temperature 0 --seed 777` at `--concurrency 16`, 23 s apart — 672 tokens both, so the byte delta is JSON width and not a length; exactly rows 17 and 18 of 21 differ, both mid-decode, both in the ragged tail `21 % 16` leaves. B's `b_off_a == b_exact` and `b_off_a == b_coarse_a` therefore compare against a baseline that does not reproduce itself and are WORTHLESS; only the OFF/OFF control made that visible, and without it B would have read as three more confirmations. Filed [#1283](https://github.com/mudler/vllm.cpp/issues/1283). **Caveats that bound this result:** nvcc was `13.3.73` here and `13.0.88` for the W4 baseline the recorded dgx gate stack names, so the OFF-vs-ON and EXACT-vs-COARSE comparisons WITHIN this binary are valid while this run and that baseline are NOT directly comparable; clocks unpinned (2405 MHz current, 3003 max, 2418 applications) and nothing measured bytes, so NO throughput and NO memory number is claimed or implied; only the Qwen3 dense decode driver was exercised; `refused=0` is ONE driver on ONE hardware and toolkit pair, which is no more a floor than W4's negative was a ceiling; and the coarse key is behind `VT_CUDA_GRAPH_DEDUP_COARSE_KEY`, default OFF, inside a default-OFF flag, on **PR [#1232](https://github.com/mudler/vllm.cpp/pull/1232) which is STILL A DRAFT — nothing on `main` folds today.** **Row stays `ACTIVE`, argued:** not `DONE`, because the fold is unreachable on every shipping configuration and the row's stated MEMORY saving has never been measured in bytes on either key; not `PARTIAL`, because nothing upstream is omitted — the coarse key is our own extension past SGLang, which keys the fields we started from; not `BLOCKED`, because nothing external stops the next step. What is owed is now a DECISION about the default plus the byte measurement and the probe-cost-at-real-churn measurement it needs, and landing #1232 first **W6, 2026-08-19, THE DEVICE-BYTE MEASUREMENT — THE BENEFIT QUESTION IS NOW CLOSED AND THE ANSWER IS NEGATIVE.** Tested `origin/main` `2c8f53d93`, which is PR #1232 LANDED, so the "nothing on `main` folds today" caveat every earlier record carried is RETIRED and this measures a configuration that ships. Same box, `rc` job `93f783de`, pod `rc-worker-4b8lj`, boot_id `3fd9745a-…` at BOTH ends, GB10, driver `580.173.02`, nvcc **13.0.88** (the W4 baseline toolkit; W5 ran 13.3.73, so W6 and W5 are NOT directly comparable while comparisons WITHIN this one binary are valid), binary sha256 `be697268…0ce657a7`, `### DONE_BYTES 2026-08-19T04:57:19Z`, 12/12 cells exit 0, zero VOID markers. **THE FOLD ENGAGES AT THE SHIPPED BUCKET SET**, which is the churn W5 could not produce: `vllm-bench` sets `max_num_seqs = concurrency`, so W32 captured `[1 2 4 8 16 24 32]` 7-of-7 and W64 captured `[1 … 64]` 11-of-11, exactly `decode_graph_sizes.h:32-41`, against the 2-3 buckets every earlier conclusion was drawn from. COARSE folds 7 graphs to 3 execs (`probes=7 refused=3`) and 11 to 5 (`probes=22 refused=16`); EXACT folds NOTHING at `probes=0`, reproducing W4 at four times the bucket count. Token ids byte-identical across every cell of a workload INCLUDING both OFF/OFF controls (`ff0db6c6…be9d` 11720 B; `e1cbf5fc…e5d0` 57620 B) — neither workload has #1283's ragged-tail shape and neither hit it. **THE SAVING DOES NOT SURVIVE ITS OWN NULL CONTROL.** `nvidia-smi --query-compute-apps` tail median (the `--query-gpu=memory.used` axis returns `[N/A]` on this box) shows W64 IDENTICAL to the megabyte in all five cells (9737) and W32's coarse arm reading 10-23 MiB HIGHER than OFF (3252/3262 vs 3262/3275). A `cudaMemGetInfo` shim summed over every instantiate gives a nominal 13.83 MiB at 7 buckets — **0.42% of a 3.25 GiB process** — and **−0.75 MiB, i.e. NOTHING, at 11**. That nominal effect is NOT ESTABLISHED on four independent grounds: `EXACT` is a TRUE NULL (same 7 and 11 retained execs, `probes=0`, so it allocates what OFF allocates) and disagrees with OFF by 10.6-13.1 MiB against a 13.83 MiB candidate; the W64 OFF/OFF pair disagrees with ITSELF by 18.2 MiB; one instantiate recorded a NEGATIVE delta (`-5,165,056` B); and `cudaGraphExecDestroy` reclaimed `0` in EVERY cell. Per-instantiate deltas for byte-identical 404-node graphs span 0 to 10,514,432 B and 17 of 27 instantiates in one cell read exactly zero, so these are POOL-GRANULAR readings and the coarse arm's throwaway probes grow that pool exactly like retained execs do. What CAN be priced: one ~390-node executable at **2.08-4.35 MiB**, 10.0-10.6 KB per node — the figure to re-run on a deep checkpoint. **THE MECHANISM INVERTS THIS ROW'S PREMISE.** The driver refuses **43% of probes at 7 buckets and 73% at 11**, every one of them `probe refused a fold (err=910 result=2)` = `cudaErrorGraphExecUpdateFailure` / `cudaGraphExecUpdateErrorTopologyChanged`. The shim's `cudaGraphGetNodes` reading says why false candidates form: the decode graphs are **TWO topologies, 376 and 404 nodes**, mixed across the buckets (`w32_off_a` captured `404 404 376 376 404 404 404`). Every refusal is about TOPOLOGY, never a parameter, so a COARSER key produces MORE false hits rather than more folds — the opposite of what W5's 2-bucket A/B suggested, and W5's `refused=0` is now explained as an artefact of workloads whose buckets only ever SHRANK, so exactly one pair was ever presented. **COST:** W32 OFF 7 instantiates / 0 updates vs COARSE 10 (3 retained + 7 probes) / 11 updates; W64 OFF 11 / 0 vs COARSE **27** (5 retained + 22 probes) / 28 updates — **2.45x the instantiate calls** to retain 6 fewer executables. **Peak transient did NOT double** — in every ON cell live-bytes peak == end, because `Register` destroys the probe before returning, so the feared "double the peak to save the steady state" trade did not occur. **A replay-time re-point DID occur** — 4 and 6 non-probe updates over 88 and 244 replays, ARITHMETIC over two printed totals and not a counter — with every cell exiting 0 and byte-identical, so `Replay`'s transitivity assumption neither aborted nor changed a token; W5 recorded that case as untested. **CAVEATS THAT BOUND THIS RESULT:** the clock pin was **REFUSED inside the lease** (`The current user does not have permission to change clocks for GPU 0000000F:01:00.0`, `clocks_pinned=0`), so **NO time-based figure is attributable** and the instantiate-wall and update-wall figures in `bytes.log` are diagnostics quoted nowhere as a result; `result=2` is ONE driver, ONE GB10, ONE toolkit; only the Qwen3 dense decode driver was exercised, as in W4 and W5; `VT_ASYNC_RUNNER=0` throughout, so the feature is STILL unreachable on the DEFAULT serving path (#1179); and `cudaMemGetInfo` cannot separate an executable's own cost from the pool chunk that satisfied it. **VERDICT, DELIVERED AND NEGATIVE:** `VT_CUDA_GRAPH_DEDUP` stays default OFF, now on MEASUREMENT rather than on silence; `VT_CUDA_GRAPH_DEDUP_COARSE_KEY` alone is a **NO-OP, not merely unsupported** — `GraphDedupCoarseKeyEnabled()` (`src/vt/graph_dedup.h:114`) is read only by the signature builder (`src/vt/graph_dedup_runtime.h:177`), only from `Register`, only under `GraphDedupEnabled()` (`src/vt/cuda/cuda_backend.cu:237`), so with dedup off its sole observable is one stderr line; both on is unsupported. **NOT A CEILING.** Three things would change it and each is traceable: find where the 376/404 split comes from (the FA-2 split-KV grid is the first suspect — a capture that fixes the node set across buckets removes every refusal); an instrument that resolves a single 2-4 MiB executable against driver pool granularity (`cuMemGetAllocationGranularity` or a pool-statistics query); and the same measurement on a 60-80 layer checkpoint, where bytes scale with node count. **Row STAYS `ACTIVE`, argued, and the argument is now narrow.** The MEASUREMENT obligations are discharged and the DECISION is delivered, which is the `DONE` case and it is a real one. Three things stop the flip and none is a checker technicality: the feature is unreachable on the DEFAULT serving path, owned by `ENG-CUDAGRAPH-BREAK` (#1179) and the "nothing lands dead" half of this row; two items still sit under #1162 itself — the device-tier signature stability/discrimination tests and probing `group.current_raw` instead of `raws.front()` to retire the transitivity assumption; and the `DONE` record surface owes a `.agents/parity-ledger.md` entry, a closing-commit owner in place of the claim, an exact test anchor and the RELEASE of `CLAIM-ENG-CUDAGRAPH-DEDUP`, which is an operator act and which this record-only branch does not own. Not `PARTIAL` — nothing upstream is omitted. Not `BLOCKED` — nothing external stops the next step. Full evidence: [benchmark-record.md](benchmark-record.md) entry `ENG-CUDAGRAPH-DEDUP W6`, raw at `/mnt/nas_share/rc/dedup-bytes/` | [eng-cudagraph-dedup.md](specs/eng-cudagraph-dedup.md); analysis [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `ACTIVE` | `CLAIM-ENG-CUDAGRAPH-DEDUP` ([#1162](https://github.com/mudler/vllm.cpp/issues/1162)) | | `ENG-CUDAGRAPH-BREAK` | One shared `vt` capture seam that accepts BREAK POINTS, so a forward containing a host-dependent op is still graphed instead of falling out entirely — and so the NINE hand-rolled drivers become one (count corrected 2026-08-18, [#1179](https://github.com/mudler/vllm.cpp/issues/1179); `9bc4d7f44` recorded eight). **Coverage AND CORRECTNESS row, not a throughput row** | T1 | mirror vLLM `CUDAGraphMode.PIECEWISE` splitting at `splitting_ops` (`vllm/config/compilation.py:60-63,517,615,630` @ `555967922`); construction from SGLang BCG `python/sglang/srt/model_executor/runner_backend_utils/breakable_cuda_graph/breakable_cuda_graph.py:204-243,246-274,309-333,335-367` @ `f63458b5be` (decorator + runtime stream capture, no compiler); its unit suite `test/registered/cuda_graph/breakable/test_breakable_cuda_graph.py:30,172,230` (305 lines, 11 unit cases) is mapped case for case in the spec's `## Tests to port` | all-or-nothing today: `src/vllm/v1/worker/gpu/runner.cpp:1338-1341` routes only `pure_decode`; drivers `qwen3_5.h:275`, `qwen3_5_dense.h:391`, `qwen3_moe.h:117`, `qwen3.h:243`, `deepseek_v2.h:324`, `voxtral.h:126`, plus `deepseek_v4.cpp`, `laguna.cpp` — and the spike found the NINTH already written, `src/vllm/model_executor/models/qwen3_dflash.cpp:771,1091`. The re-derivation is measured, not asserted: `StepDevInputs` (`src/vllm/model_executor/models/qwen3_5.cpp:3894`, the persistent DEVICE input path) exists in ONE driver and `grep -c` returns 0 in `qwen3_moe.cpp`, `qwen3.cpp`, `deepseek_v2.cpp` and `voxtral.cpp`, which is why `src/vllm/model_executor/models/qwen3.cpp: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 did NOT meet it and says so**: no `rc` lease was obtainable in its window and a CPU harness cannot replay a captured segment, so it moves to W3 with the three drivers of the same shape (G1); the host-lifetime contract of `decode-graph-scratch-uaf-2026-07-18.md` enforced AT the seam — D1's INPUT half, making the intermediates a segment reads unavailable to the `DevicePool` free list, which becomes live only for the first PIECEWISE production capture (W4); the auxiliary-stream auto-join before every segment close (`:353-361`, spec D10), live at `src/vllm/model_executor/models/qwen3_5.cpp:6254-6255,6384` and `src/vllm/model_executor/models/laguna.cpp:2572-2576,2612` (W4, W5). **Delivered by W1** ([#1192](https://github.com/mudler/vllm.cpp/issues/1192)): the reachability mutation (performed; deleting the call site reds `tests/vllm/models/test_qwen3_break_point.cpp` and leaves the unit suite green); the ported SGLang unit cases with their arithmetic chains and post-replay assertions; and the break-function OUTPUT writeback (`replay_fn`/`_copy_output` `breakable_cuda_graph.py:231-235,172-201`, spec D9), whose destination is a `vt::BreakSlot` the seam owns rather than a caller reference it cannot outlive | spec [eng-cudagraph-break.md](specs/eng-cudagraph-break.md) (W0 spike DONE 2026-08-18: the existing `vt` capture vocabulary `include/vt/backend.h:208-222` expresses a SEGMENTED capture with NO new virtual, because `EndCaptureGraph` stores nothing (`src/vt/cuda/cuda_backend.cu:225-232`); a break point is expressible with one `thread_local` capture pointer plus a free function, no compiler and no decorator); **W1 DONE 2026-08-18 ([#1192](https://github.com/mudler/vllm.cpp/issues/1192)): the seam LANDS** — `vt::BreakableGraph`, `vt::GraphCaptureScope` and `vt::GraphBreak` (`include/vt/breakable_graph.h`, `src/vt/breakable_graph.cpp`), the SGLang unit suite ported case for case (`tests/vt/test_breakable_graph.cpp`, 24 cases / 163 assertions, re-derived 2026-08-18 by `ninja test_breakable_graph && ./build/tests/test_breakable_graph`; the recorded 14/81 never re-derived at any head of this branch), and ONE break point registered at the DENSE ATTENTION ENTRY of `Qwen3ForCausalLM` (`src/vllm/model_executor/models/qwen3.cpp`, inside `RunLayer`). **The exit criterion W0 deliberately left open is ANSWERED on a leased GPU:** `cudaStreamEndCapture` then `cudaStreamBeginCapture` on the SAME stream mid-forward with EAGER work between is LEGAL under `cudaStreamCaptureModeThreadLocal` (`src/vt/cuda/cuda_backend.cu:204-206`) — `orin:gpu0` via an `rc` lease, driver 12060, 3 replays with fresh inputs, 0 mismatches, bare zero-work re-begin legal too. G2 reachability is `tests/vllm/models/test_qwen3_break_point.cpp`, which drives the production `Qwen3DenseModel::Forward` with a scope open and counts `num_hidden_layers + 1` segments (mutation: delete the call site ⇒ 1 segment ⇒ RED), and holds G4 in the same case at 500 logits / 0 differing bit for bit. STAGED SLICE, named: the scope and the container are not yet ENTERED from a production step — no driver opens a scope until W2 migrates `Qwen3DenseDecodeGraph` — and the spec's `## Owed` lists it with W2 as owner, alongside the D10 auxiliary-stream auto-join (W4/W5), G5's ROCm/Tenstorrent arms (W3) and G1 on a real GPU (W2). **The capture-failure drain is NOT among them: it landed HERE**, as behaviour (`std::uncaught_exceptions()` compared against the depth recorded at scope entry, so a break function or ordinary model code throwing mid-capture destroys the partial container instead of handing back a forward that reports `captured() == true`) and as three gated arms (tests 13a, 13b, 13c). The spec's `## Owed` strikes the item through and reads DELIVERED in W1; this cell said the opposite until 2026-08-18 because `cba969857` re-derived field 6 alone. **W2 DONE 2026-08-18 ([#1261](https://github.com/mudler/vllm.cpp/issues/1261)): `Qwen3DenseDecodeGraph` MIGRATED and the seam is ENTERED from a production step**, which retires W1's staged slice. `Qwen3DenseDecodeGraph::Step` opens a `vt::GraphCaptureScope` over a per-slot `vt::BreakableGraph` and replays through `BreakableGraph::Replay`; the hand-rolled `BeginCapture`/`EndCaptureGraph` pair, the raw `void*` handle, the `bool captured` flag, the `DestroyGraph` loop and the driver's own `VLLM_CPP_CUDAGRAPH` read are gone (re-derivation items 1, 2, 5, 6). The migration ADDED `vt::GraphCaptureMode`, mirroring vLLM's `CUDAGraphMode` (`vllm/config/compilation.py:59-63`), whose v1 default `FULL_AND_PIECEWISE` (`:63`) is documented at `:630-632` as a FULL graph for DECODE batches and a piecewise one for prefill/mixed, with `decode_mode()` (`:65-66`) selecting the full half and the runtime reading it at `vllm/v1/worker/gpu/cudagraph_utils.py:185-186`. A decode driver opened `kPiecewise` would have turned a fully graphed decode step into ONE EAGER ATTENTION CALL PER LAYER between graph replays — not vLLM's decode behaviour, and invisible to every token gate here. `GraphBreak` in a `kFull` scope takes the pass-through arm and `AppendBreak` REFUSES a registration in that mode. G2 is `tests/vllm/models/test_qwen3_decode_graph_seam.cpp` (3 cases / 124 assertions), which asserts the SEAM's counters because a driver calling `Backend::ReplayGraph` directly leaves an identical backend log; the mutation restoring the pre-W2 raw pair (18 lines, compiled clean) left `test_breakable_graph` 27/27, `test_qwen3_break_point` 2/2 and `test_qwen3_forward` 10/10 GREEN and reddened only this file. G4 in the same file: capture step vs `Qwen3DenseModel::Forward`, 100 logits, 0 differing. **The async decline at `qwen3.cpp` STANDS and is now GATED in both arms**: migrating the capture does not move the INPUTS, so the depth-2 race is untouched, and the fix is `StepDevInputs` as a SEAM capability, which is W4. **G1 is NOT met by W2** and is recorded owed rather than implied. Still NO throughput claim; analysis [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `ACTIVE` | `CLAIM-ENG-CUDAGRAPH-BREAK-W2`; [#1163](https://github.com/mudler/vllm.cpp/issues/1163), [#1192](https://github.com/mudler/vllm.cpp/issues/1192), [#1261](https://github.com/mudler/vllm.cpp/issues/1261), [#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. **The unblock order now has an owning row:** `LTX25-DEVICE-RESIDENCY` ([#1264](https://github.com/mudler/vllm.cpp/issues/1264), [ltx25-device-residency.md](specs/ltx25-device-residency.md)) stages those defects W0-W6 and carries this decision point as its W7 — if the loop comes back GPU-bound, #1164 closes as a refutation the way [#1161](https://github.com/mudler/vllm.cpp/issues/1161) closed prefill capture | [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `INVENTORIED` | [#1164](https://github.com/mudler/vllm.cpp/issues/1164) | | `ENG-BATCH-INVARIANT` | Opt-in deterministic execution across scheduler batch sizes (`VLLM_BATCH_INVARIANT=1`): batch-invariant matmul/norm/attention/collectives plus persistent-scheduler NVFP4; production default remains off | T1 | default/env `vllm/envs.py:89,576-578`; initialization `vllm/v1/worker/gpu_worker.py:1262`; NVFP4 dispatch `csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_sm120_kernels.cu:212-220`; suite fixture `tests/v1/determinism/conftest.py:9-12`; operator/e2e `tests/v1/determinism/test_nvfp4_batch_invariant_scaled_mm.py`, `tests/v1/determinism/test_nvfp4_batch_invariant.py` @ `702f481` | - | [W3-C3R executed contract](specs/nvfp4-persistent-plan-cache.md#w3-c3r-batch-shape-localization-and-gate-correction-2026-07-13): production-default ours and vLLM both change outputs across batch shapes; no local opt-in implementation is claimed | `planned: specs/batch-invariant-execution.md` | `INVENTORIED` | - | diff --git a/.agents/issue-index.md b/.agents/issue-index.md index e1efded2b..d2d98bafc 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -412,3 +412,4 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1247](https://github.com/mudler/vllm.cpp/issues/1247) | `MODEL-MUSIC-minimax-music3-mini-max-music3-for-conditional-generation` | MiniMax-Music3 depth-stage RECORD, four defects, no measured number affected: (a) `test_minimax_music3_ar.cpp`, spec §15.5 and PR #1238's body all say the committed goldens are "8-wide with ONE head" so a head stride is invisible — `minimax_music3_ar_goldens.inc` sets `kMusic3DepthHeads = 2` over `kMusic3DepthHidden = 8`, so they are 2 heads of 4 and dropping the head stride from the cached KEY index reds 4 cases / 32 assertions, TWO of them at the goldens' own geometry; (b) "12 timed rounds per arm" and "all 20 runs printed `f0cfeed6eee4f55d`" cannot be reconciled with the record's own table — 5 pairs x 1 round + 3 pairs x 4 rounds is 17 rounds per arm, and 8 pairs x 2 arms is 16 processes, each printing ONE fingerprint after its round loop; (c) the heading named base `origin/main` `727163997` while the body named `fc163f62b`, and only the latter was built and timed (the delta is #1231's profiler, which the driver never enters); (d) "70 rows a frame to read 16" conflates the 14 rows the OLD arm read with the AFTER arm's 16. FIXED IN FLOW, owned by `MUSIC3-DEPTH-SPEED`, spec §15.2, §15.3, §15.5, §15.6. | bug | | [#1283](https://github.com/mudler/vllm.cpp/issues/1283) | — | **Greedy decode is NOT reproducible at concurrency 16.** Two runs of ONE binary (sha256 `ca114abb…c772ad`) on ONE workload with `--temperature 0 --seed 777 --concurrency 16 --num-prompts 21 --output-len 32`, 23 seconds apart on GB10 / driver `580.173.02` / `VT_ASYNC_RUNNER=0`, emitted DIFFERENT token ids: `5973c5a10a6210085417fb25a29edbd0dc15fe61d7d4f774dd8ff3883dae1d64` (2638 bytes) vs `4cf7923080db6aa29759537f2192f3d9500db11c0e8d72bbb2b4ac6e4614af7c` (2650 bytes). **Not a dedup defect** — `VT_CUDA_GRAPH_DEDUP` is UNSET in both cells; these are the OFF/OFF control of the `ENG-CUDAGRAPH-DEDUP` W5 coarse-key A/B ([#1226](https://github.com/mudler/vllm.cpp/issues/1226)), and the control is the only reason it was seen. Both runs emit `ids_requests=21 ids_total_tokens=672 empty_rows=0`, so the byte delta is JSON decimal width and NOT a length difference, not a truncation and not an early stop; exactly rows 17 and 18 of 21 differ and both diverge MID-DECODE (token index 11 and 5), not at the first token, and both sit in the ragged tail `21 % 16` leaves. Workloads A (conc 24) and C (conc 32) in the same series on the same binary in the same minutes each reproduced themselves exactly, so it is one configuration of three at one repetition each. **The cost: it VOIDS workload B of that run** — `b_off_a == b_exact` and `b_off_a == b_coarse_a` compare against a baseline that does not reproduce itself, and without the OFF/OFF control they would have read as three more byte-identity confirmations. NOT diagnosed beyond the evidence; the issue carries an isolation plan (establish the rate over N repetitions; re-run with `VLLM_CPP_CUDAGRAPH=0` to separate the scheduler from the graph path; a non-ragged `--num-prompts` multiple of `--concurrency`; per-step batch composition for the two diverging requests; the top-2 logit margin at the divergence step, because a near-tie a reduction order can flip is a different defect from a wrong value). Evidence `/mnt/nas_share/rc/dedup-key/logs-ab/ab.log` and `out-ab/ids_b_off_*.json`; owed under `## Owed` of [eng-cudagraph-dedup.md](specs/eng-cudagraph-dedup.md) | bug | | [#1298](https://github.com/mudler/vllm.cpp/issues/1298) | — | `scripts/agent-integration.py` CANNOT RUN. `AGENTS.md` §Commands names it as the pre-merge gate, but `main()` calls `cutover_oid()` before it reads anything and that function loads `.agents/policy-cutover`, DELETED by `0f3e44eee` ("policy: the code is the state, git is the history", 2026-08-09) and never restored -- so every invocation exits `INTEGRATION FAILED: missing policy cutover` whatever the tree, the pull request or the trailers hold. Measured at `origin/main` `489a9a4c0`. `scripts/agent-integration.py:107` is the tree's ONLY `--cutover` caller, and it sits behind that raise; `.agents/specs/fix-trailer-lane-cutover.md` already recorded the file as absent when it rejected `--cutover` as an instrument, but not that the absence bricks the command. NO SUITE IS RED, because `tests/scripts/test_agent_gates.py:205` exercises `cutover_oid` against a SYNTHETIC repository it writes the anchor into: the function is tested, the command is not. Found while `GATE-PR-BODY-TRAILERS` ([#1263](https://github.com/mudler/vllm.cpp/issues/1263)) evaluated it as the home for the pre-merge body check and rejected it on this ground -- wiring a new check there would land it behind a permanent refusal. NOT fixed in flow: the repair is a decision (drop `--cutover` from the command, or restore an anchor and say what value it should hold and why), so it owes its own spec, red-before and reviewer. Owed under `## Owed` in [`gate-pr-body-trailers.md`](specs/gate-pr-body-trailers.md) | bug | +| [#1303](https://github.com/mudler/vllm.cpp/issues/1303) | `ENG-CUDAGRAPH-DEDUP` | **The decode graphs are TWO topologies, not one: 376 and 404 nodes, mixed across the shipped padded bucket set.** Found by the `ENG-CUDAGRAPH-DEDUP` W6 device-byte run (`origin/main` `2c8f53d93`, GB10, driver `580.173.02`, nvcc 13.0.88, `rc` job `93f783de`, 2026-08-19) with a `cudaGraphGetNodes` shim on every `cudaGraphInstantiate`: `w32_off_a` captured `404 404 376 376 404 404 404` and `w64_coarse_a` saw `376 376 376 376 376 376 404 404 376 …` over 11 buckets. **The measured consequence:** with the coarse key the driver refuses **43% of probes at 7 buckets and 73% at 11**, and EVERY refusal in the run gave one reason — `probe refused a fold (err=910 result=2)`, `cudaErrorGraphExecUpdateFailure` / `cudaGraphExecUpdateErrorTopologyChanged`. Not one was about a parameter. **This INVERTS the premise of [#1226](https://github.com/mudler/vllm.cpp/issues/1226) (now closed):** `cudaGraphExecUpdate` permits parameter changes and requires topology to match, so dropping the launch dimensions from the key was expected to fold more; it produces more FALSE CANDIDATES instead, because the key cannot see the difference the driver rejects on. W5's `refused=0` is explained as an artefact of workloads whose buckets only ever SHRANK, so exactly one pair was ever presented. **Why it matters beyond the key:** at 11 buckets the coarse arm ran 27 instantiate calls against OFF's 11 (2.45x) plus 28 update calls, to retain 6 fewer executables for no measurable byte saving; if the split has a removable cause, every refusal disappears and the fold ratio becomes the whole bucket set rather than 3/7 and 5/11, which is the only route by which the default flip could become supportable on this model class. First suspect, NOT verified: the FA-2 split-KV grid. Owed: attribute the split to a node or subgraph by DUMPING and diffing the two node sets (`cudaGraphNodeGetType`, `cudaGraphDebugDotPrint`) rather than reading source; say whether a capture that fixes the node set exists and what it costs; re-run the W6 A/B if it does. Evidence `/mnt/nas_share/rc/dedup-bytes/` (`RESULT.md`, `logs-bytes/bytes.log` `### PER-INSTANTIATE BYTE DISTRIBUTION`, `mem/mem_.csv` column 8); recorded in [`benchmark-record.md`](benchmark-record.md) entry `ENG-CUDAGRAPH-DEDUP W6` and [`eng-cudagraph-dedup.md`](specs/eng-cudagraph-dedup.md) `## Outcome` (W6) and `## Owed` | perf | diff --git a/.agents/specs/eng-cudagraph-dedup.md b/.agents/specs/eng-cudagraph-dedup.md index 458dea25c..62d22b183 100644 --- a/.agents/specs/eng-cudagraph-dedup.md +++ b/.agents/specs/eng-cudagraph-dedup.md @@ -11,28 +11,39 @@ is why the load-bearing gate below is byte-identity rather than a speed ratio. ## Now `ACTIVE`, and the state is argued rather than inherited. The shared registry, its CPU -contract suite, the CUDA wiring and the -[#1184](https://github.com/mudler/vllm.cpp/issues/1184) repair all landed. The owed -device A/B ran on 2026-08-18 (W4) and split in two: the correctness half PASSED, and the -benefit half was REFUTED for exactly the case this row was filed for. `N == M` in every -`VT_CUDA_GRAPH_DEDUP=1` cell, because the signature carries the padded batch dimension, -so no two decode buckets ever group. - -**W5, the same day, tested the hypothesis that negative produced and CONFIRMED it.** -A COARSE key that drops the launch dimensions and the memcpy extents folds every bucket -the exact key left alone — 3 graphs to 2 execs, 2 to 1, 2 to 1 — with -`cudaGraphExecUpdate` probed once per fold and **accepted every time**, `refused=0`. The -benefit W4 recorded as unreachable is reachable via the key -([#1226](https://github.com/mudler/vllm.cpp/issues/1226) DELIVERED). - -The row stays `ACTIVE`, and the argument moved rather than stayed still. It is still not -`DONE`: the coarse key is behind its own opt-in in **PR #1232, which has not landed**, -so nothing on `main` folds today, and the row's stated saving is a MEMORY saving that -**nobody has measured in bytes on either key**. It is still not `PARTIAL`, because -nothing upstream is omitted; the coarse key is our own extension past SGLang, which keys -the same fields we started from. It is not `BLOCKED`, because nothing external stops the -next step. What is owed is now a DECISION about the default, and that decision needs a -measurement that does not exist. Full result: [`## Outcome`](#outcome). +contract suite, the CUDA wiring, the +[#1184](https://github.com/mudler/vllm.cpp/issues/1184) repair and the coarse key of +[#1226](https://github.com/mudler/vllm.cpp/issues/1226) have all landed — the last of +them as `2c8f53d93`, so the "PR #1232 is a draft" caveat every earlier record carried is +retired and the fold is reachable on `main` today, behind two default-OFF flags. + +**The row's benefit question is now CLOSED, and the answer is NEGATIVE.** W6 ran the +device-byte measurement W4 and W5 both named as owed, on the SHIPPED bucket sets rather +than the 2-3 buckets earlier runs reached. The fold is real there: 7 captures fold to 3 +executables at `max_num_seqs=32` and 11 fold to 5 at 64. The saving is not: 13.83 MiB +nominal at 7 buckets, which is 0.42% of a 3.25 GiB process and SMALLER than the +disagreement of its own null control, and nothing at all at 11 buckets. The driver +refuses 43% of probes at 7 buckets and 73% at 11, every one of them +`cudaGraphExecUpdateErrorTopologyChanged`, because the decode graphs are two topologies +of 376 and 404 nodes rather than one. **That inverts the premise the coarse key was +built on**: a key that cannot see a topology difference produces more false candidates, +not more folds. Both flags therefore stay OFF, on measurement rather than on silence. +Full result: [`## Outcome`](#outcome), W6. + +**Why the row is still `ACTIVE` and not `DONE`.** The measurement obligations are +discharged and the decision is delivered, which is the DONE case and it is a real one. +Three things stop the flip here, and none of them is a checker technicality. The +feature is unreachable on the DEFAULT serving path, because the async runner captures no +decode graph — that reach is owned by `ENG-CUDAGRAPH-BREAK` +([#1179](https://github.com/mudler/vllm.cpp/issues/1179)) and is the "nothing lands +dead" half of this row. Two items still sit under this row's own issue: the device-tier +signature stability and discrimination tests, and probing `group.current_raw` instead of +`raws.front()` to retire the transitivity assumption in `Replay`. And the `DONE` record +surface itself owes a `.agents/parity-ledger.md` entry, a closing-commit owner in place +of `CLAIM-ENG-CUDAGRAPH-DEDUP`, an exact test anchor, and the release of that claim — +files this record-only branch does not own. It is not `PARTIAL`, because nothing +upstream is omitted; the coarse key is our own extension past SGLang. It is not +`BLOCKED`, because nothing external stops the next step. ## Scope @@ -720,6 +731,140 @@ measured bytes saved — the fold count is a count of executables, not of memory **Only the Qwen3 dense decode driver was exercised**, as in W4. Whether any other capture site folds under either key is untested. +### W6, 2026-08-19: the DEVICE-BYTE measurement, and the topology finding that inverts the premise + +This is the measurement W4 and W5 both recorded as owed. Every number before it counted +executables; the row's stated saving is a MEMORY saving. It also removes W5's largest +caveat: **PR #1232 landed as `2c8f53d93`**, so this run measures a configuration that +ships rather than a draft. + +Ran on `dgx:gpu0` through an `rc` lease, job `93f783de-228f-47d5-806d-c5b56aa72c3a`, pod +`rc-worker-4b8lj`, boot_id `3fd9745a-d25a-426c-ba3c-97c958a85515` at BOTH ends, GB10, +driver `580.173.02`, nvcc **13.0.88**, `### DONE_BYTES 2026-08-19T04:57:19Z`. Twelve +cells of one binary (sha256 `be697268…0ce657a7`) built from `origin/main` `2c8f53d93` +(tar sha256 asserted before extraction), `VT_ASYNC_RUNNER=0` and +`VT_DECODE_GRAPH_STATS=1` throughout, the only variables `VT_CUDA_GRAPH_DEDUP` and +`VT_CUDA_GRAPH_DEDUP_COARSE_KEY`. All 12 cells exit 0 with zero VOID markers. Recipe, +instruments, cell tables and every caveat: +[`.agents/benchmark-record.md`](../benchmark-record.md), entry `ENG-CUDAGRAPH-DEDUP W6`. + +**The bucket churn is the shipped one.** `vllm-bench` sets +`max_num_seqs = concurrency`, so W32 captured `[1 2 4 8 16 24 32]`, **7 of 7**, and W64 +captured `[1 2 4 8 16 24 32 40 48 56 64]`, **11 of 11** — exactly +`include/vllm/model_executor/models/decode_graph_sizes.h:32-41`. W5 reached 2-3 buckets +and every earlier conclusion about probe behaviour came from a process that presented +the driver with one pair. + +**The fold engages.** + +```text +w32_coarse_a/b captured 7 graphs, deduped to 3 execs (probes=7 refused=3) +w32_exact_a captured 7 graphs, deduped to 7 execs (probes=0 refused=0) +w64_coarse_a/b captured 11 graphs, deduped to 5 execs (probes=22 refused=16) +w64_exact_a captured 11 graphs, deduped to 11 execs (probes=0 refused=0) +``` + +Token ids are byte-identical across every cell of a workload, including both OFF/OFF +controls (W32 `ff0db6c6…be9d`, 11720 B; W64 `e1cbf5fc…e5d0`, 57620 B). Neither workload +has [#1283](https://github.com/mudler/vllm.cpp/issues/1283)'s ragged-tail shape, and +neither hit it. + +**The saving does not survive its own null control.** Two instruments, and they agree. +`nvidia-smi --query-compute-apps` per-process tail median shows W64 identical to the +megabyte in all five cells (9737 MiB) and W32's coarse arm reading 10-23 MiB *higher* +than OFF (3252/3262 OFF against 3262/3275 COARSE). A `cudaMemGetInfo` shim summed over +every `cudaGraphInstantiate` gives a nominal **13.83 MiB** at 7 buckets — **0.42%** of +the 3.25 GiB process — and **−0.75 MiB**, i.e. nothing, at 11. + +That nominal effect is **NOT ESTABLISHED**, on four grounds that are independent of each +other. `EXACT` is a true null control — it retains exactly as many executables as OFF (7 +and 11) at `probes=0` — and it disagrees with OFF by 10.6-13.1 MiB, against a candidate +effect of 13.83 MiB. The W64 OFF/OFF pair disagrees with ITSELF by 18.2 MiB. One +instantiate recorded a NEGATIVE delta (`-5,165,056` B in `w32_exact_a`), which is only +possible if something else released memory inside the same window. And +`cudaGraphExecDestroy` returned `reclaimed_bytes=0` in every cell. Per-instantiate +deltas for byte-identical 404-node graphs range from 0 to 10,514,432 B, and 17 of 27 +instantiates in `w64_coarse_a` read exactly zero: these are POOL-GRANULAR readings, and +the coarse arm's throwaway probe executables grow that pool exactly like retained ones +do. What can be priced is one ~390-node executable at **2.08-4.35 MiB**, or 10.0-10.6 KB +per node. + +**THE MECHANISM, and why it inverts this row's own hypothesis.** The driver refuses +**43% of probes at 7 buckets and 73% at 11**, and every single refusal is +`vt graph dedup: probe refused a fold (err=910 result=2)` — +`cudaErrorGraphExecUpdateFailure` with `cudaGraphExecUpdateErrorTopologyChanged`. The +shim's `cudaGraphGetNodes` reading says why false candidates form at all: the decode +graphs are **two topologies, 376 and 404 nodes**, mixed across the bucket set +(`w32_off_a` captured `404 404 376 376 404 404 404`). + +Every refusal is about TOPOLOGY, never about a parameter. **So a coarser key produces +more false hits, not more folds** — the opposite of what W5's 2-bucket A/B suggested. +W5's "the driver accepts every time, `refused=0`" was an artefact of workloads whose +buckets only ever shrank, so exactly one pair was ever presented. This is the correction +W5's own `refused=0` caveat asked for, and it arrived from the same box. + +**The probe cost, exact and clock-independent.** W32: OFF 7 instantiates / 0 updates +against COARSE **10** (3 retained + 7 probes) / 11 updates. W64: OFF 11 / 0 against +COARSE **27** (5 retained + 22 probes) / 28 updates — **2.45x the instantiate calls** to +retain 6 fewer executables. **Peak transient did NOT double**: in every ON cell the +live-bytes peak equals its end value, because `Register` destroys the probe executable +before it returns. The feared "double the peak to save the steady state" trade did not +occur. + +**A replay-time re-point occurred, as ARITHMETIC over two printed totals.** `probes=7` +against 11 update calls and `probes=22` against 28 leaves 4 and 6 updates that are not +probes, over 88 and 244 replays — the reverse re-point W5 recorded as untested on a +device. Those cells exit 0 with byte-identical ids, so `Replay`'s transitivity +assumption neither aborted nor changed a token here. It is stated as arithmetic and NOT +as a measurement, because the registry counts probe refusals only and a replay-time +refusal would abort on `VT_CHECK` rather than increment anything. + +**Times: NONE is attributable.** The clock pin was REFUSED inside the lease — `The +current user does not have permission to change clocks for GPU 0000000F:01:00.0`, and +the log records `clocks_pinned=0`. The instantiate-wall and update-wall figures in +`bytes.log` are diagnostics and are deliberately quoted nowhere as a result. + +### W6: why the defaults keep their values + +1. **`VT_CUDA_GRAPH_DEDUP` stays OFF — supported by numbers, not merely ungated.** Best + nominal saving 13.83 MiB (0.42% of process) at the 7-bucket set, smaller than its own + null control's disagreement; nothing at 11 buckets; cost is 16 extra + instantiate/destroy pairs and 28 update calls on the capture path. +2. **`VT_CUDA_GRAPH_DEDUP_COARSE_KEY` alone is a NO-OP, not merely unsupported.** + `GraphDedupCoarseKeyEnabled()` (`src/vt/graph_dedup.h:114`) is read only by the + signature builder (`src/vt/graph_dedup_runtime.h:177`), which runs only from + `Register`, which `src/vt/cuda/cuda_backend.cu:237` calls only under + `GraphDedupEnabled()`. With dedup off, its sole observable is one stderr line. +3. **Both on — not supported.** The same numbers plus a 73% probe-refusal rate. + +### W6: what this run does NOT establish + +- **No time-based figure is attributable**, because the clock pin was refused. +- **nvcc was 13.0.88 here and 13.3.73 for the W5 fold run.** Comparisons WITHIN this one + binary are valid; this run and W5 are NOT directly comparable to each other. +- **`result=2` is one driver (`580.173.02`), one GB10, one toolkit.** +- **Only the Qwen3 dense decode driver was exercised**, as in W4 and W5. +- **`VT_ASYNC_RUNNER=0` throughout**, so the feature remains unreachable on the DEFAULT + serving path ([#1179](https://github.com/mudler/vllm.cpp/issues/1179)). +- **`cudaMemGetInfo` cannot separate an executable's own cost from the pool chunk that + satisfied it.** Every byte figure inherits that limit. +- `run.sh`'s per-cell `shim_inst … sum_bytes` reducer sums only POSITIVE deltas and + reads higher in the two cells holding a negative one. The tables use the SIGNED sum, + which matches the shim's in-process `SUMMARY` counter in every cell. + +### W6: no ceiling — what would change the answer + +- **Where the 376/404 node split comes from.** The FA-2 split-KV grid is the first + suspect. A capture that fixes the node set across buckets removes every refusal in this + run and turns the fold ratio into the whole bucket set. Until then, coarsening the key + is the wrong direction. +- **An instrument that resolves a single 2-4 MiB executable** against driver pool + granularity — `cuMemGetAllocationGranularity`, a pool-statistics query, or a build with + pool instrumentation. +- **A 60-80 layer checkpoint.** Bytes scale with node count and this graph is 376-404 + nodes on a 28-layer 0.6B model. 10.0-10.6 KB per node is the figure to re-run against + before the flip is refused permanently for large models. + ## Owed **DELIVERED by W4, 2026-08-18 — see [`## Outcome`](#outcome).** The device @@ -730,6 +875,36 @@ executable-count ratio: **DELIVERED AND NEGATIVE**, `N == M` in every ON cell ov and three distinct padded buckets, so the row's headline saving is measured NOT to occur with the current signature. What remains owed is below. +**DELIVERED by W6, 2026-08-19 — THE DECISION, and it is NEGATIVE.** The default flip +this row has owed since W4 is no longer owed: it is DECIDED, against, on measurement. +`VT_CUDA_GRAPH_DEDUP` stays OFF, `VT_CUDA_GRAPH_DEDUP_COARSE_KEY` stays OFF and is a +no-op on its own, and both together are unsupported. The two numbers W5 named as +missing were both taken. **Bytes:** 13.83 MiB nominal at the shipped 7-bucket set, 0.42% +of a 3.25 GiB process, smaller than the disagreement of its own null control, and +nothing at 11 buckets. **The probe cost at real bucket churn:** 43% of probes refused at +7 buckets and 73% at 11, 2.45x the instantiate calls, every refusal +`cudaGraphExecUpdateErrorTopologyChanged`. Two W5 items are closed as a side effect: PR +#1232 LANDED as `2c8f53d93`, and the replay-time re-point did occur (4 and 6 of them, +arithmetic over the update and probe totals) without an abort or a token change. W6 also +INVERTED this row's premise — the decode graphs are two topologies of 376 and 404 nodes, +so a coarser key produces more false candidates rather than more folds. See +[`## Outcome`](#outcome), W6. **What remains owed is below, and it is now three +measurements plus the residual engineering — never another attempt at the same flip.** + +**[#1226](https://github.com/mudler/vllm.cpp/issues/1226) is CLOSED.** It owned the +coarse-key hypothesis, and W5 confirmed it while W6 priced it and inverted it. The rows +in the older tables below that still name it are history and are kept as such; the live +residue moved on 2026-08-19 to +[#1303](https://github.com/mudler/vllm.cpp/issues/1303), which owns the 376/404 topology +split, and to [#1162](https://github.com/mudler/vllm.cpp/issues/1162), which owns this +row and everything else still open on it. + +| Item | Issue | Why not here | +|---|---|---| +| **Where the 376/404 node split comes from.** Every refusal in W6 was `cudaGraphExecUpdateErrorTopologyChanged`, and the node counts say the decode buckets genuinely differ in STRUCTURE, not only in parameters. The FA-2 split-KV grid is the first suspect. A capture that fixes the node set across buckets would remove every refusal and turn the fold ratio into the whole bucket set; until then, coarsening the key is measured to be the wrong direction | [#1303](https://github.com/mudler/vllm.cpp/issues/1303) | it is a NEW hypothesis about the capture path rather than about the key, needs its own spec and its own red-first evidence, and W6's job was to price the key that shipped | +| **An instrument that resolves a single 2-4 MiB executable against driver pool granularity.** Neither W6 instrument can: `cudaMemGetInfo` measures pool growth, 17 of 27 instantiates in one cell read exactly zero, one read NEGATIVE, and `cudaGraphExecDestroy` reclaimed nothing anywhere. `cuMemGetAllocationGranularity`, a driver-pool statistics query, or a build with pool instrumentation would price the fold directly instead of through a total whose noise exceeds the effect | [#1162](https://github.com/mudler/vllm.cpp/issues/1162) | it is an instrument-building task, and W6 deliberately reported the noisy total WITH its null control rather than presenting a 13.83 MiB number the control cannot support | +| **A 60-80 layer checkpoint.** Bytes scale with node count and W6's graph is 376-404 nodes on a 28-layer 0.6B model, measured at 10.0-10.6 KB per node. The flip is refused for THIS model class on THIS evidence; it is not refused permanently for large models until the same measurement runs on a deep one | [#1162](https://github.com/mudler/vllm.cpp/issues/1162) | it needs another lease and a large checkpoint, and refusing the flip for the measured case does not depend on it | + **DELIVERED by W5, 2026-08-18 — the [#1226](https://github.com/mudler/vllm.cpp/issues/1226) hypothesis: CONFIRMED.** A coarse key folds every decode bucket the exact key left alone (3→2, 2→1, 2→1) and the driver accepted every probe, `refused=0`. That closes the @@ -739,10 +914,10 @@ nobody has taken.** Every number so far counts executables. Nothing has measured | Item | Issue | Why not here | |---|---|---| -| Flipping `VT_CUDA_GRAPH_DEDUP` on by default | [#1162](https://github.com/mudler/vllm.cpp/issues/1162) | **STILL NOT JUSTIFIED, and the reason CHANGED.** W4's reason was that the ON arm allocated as many executables as OFF. W5 removed that reason: with the coarse key the ON arm allocates FEWER. The reason now is that nobody has priced the saving. A default is a measurement, and "2 execs instead of 3" is a count, not bytes. See the row below | -| **THE DECISION, and the measurement it needs: does the COARSE key become the default?** Two numbers are missing and neither is optional. (1) **Bytes.** What does one `cudaGraphExec_t` cost on this model at this bucket set, and what does the fold actually return — device allocation at steady state, ON versus OFF, on a pinned-clock idle box. Every W4 and W5 figure counts executables. (2) **The probe cost at real bucket churn.** W5 probed once per fold on a 2-3 bucket process. A serving grid with 7 or 11 buckets across nine capture drivers churns differently, and the probe runs on the capture path. Only then is the flip a measurement rather than a preference | [#1226](https://github.com/mudler/vllm.cpp/issues/1226) | it needs a leased GPU, pinned clocks, and PR #1232 landed first — the key it would default to is still a draft. Deliberately NOT decided by this record, exactly as W5's own hypothesis was not decided by W4's | -| Landing PR [#1232](https://github.com/mudler/vllm.cpp/pull/1232), the coarse-key mode itself. It is a DRAFT, so nothing on `main` folds today and this row's saving is unreachable on every configuration that ships | [#1226](https://github.com/mudler/vllm.cpp/issues/1226) | the experiment measured whether the key works before asking whether it should land, which is the correct order. It needs its own fresh review, and its red-first discrimination test — that the coarse key still SEPARATES two genuinely different topologies — is the load-bearing one | -| **An ALTERNATING workload, and a counter for replay-time re-points.** W5's three workloads captured their padded buckets in strictly descending order, so no group was ever replayed back to an earlier member and the only pair any coarse group issued was the one the probe tested. The reverse re-point — the pair that is untested from group size two onwards — is therefore unmeasured on a device, and `refused=0` does not cover it. The registry counts PROBE refusals only, so the ask is not even observable today: a replay-time refusal aborts on `Replay`'s `VT_CHECK` instead of incrementing anything | [#1226](https://github.com/mudler/vllm.cpp/issues/1226) | it needs a workload whose decode batch returns to a bucket it already captured, which the fixed-pool bench harness does not produce, plus a replay-side counter. Both ride with the lease the memory measurement above already needs | +| ~~Flipping `VT_CUDA_GRAPH_DEDUP` on by default~~ **DECIDED AGAINST by W6, 2026-08-19.** It is no longer owed, and no longer merely unjustified: 13.83 MiB nominal (0.42% of process) at 7 buckets, nothing at 11, smaller than its own null control's disagreement, against 16 extra instantiate/destroy pairs and 28 update calls. History below | [#1162](https://github.com/mudler/vllm.cpp/issues/1162) | **STILL NOT JUSTIFIED, and the reason CHANGED.** W4's reason was that the ON arm allocated as many executables as OFF. W5 removed that reason: with the coarse key the ON arm allocates FEWER. The reason now is that nobody has priced the saving. A default is a measurement, and "2 execs instead of 3" is a count, not bytes. See the row below | +| ~~**THE DECISION, and the measurement it needs: does the COARSE key become the default?**~~ **DELIVERED AND NEGATIVE by W6, 2026-08-19.** Both numbers were taken at the SHIPPED bucket sets: bytes (0.42% at 7 buckets, none at 11, inside the null control's own noise) and the probe cost at real churn (43% and 73% refused, 2.45x the instantiate calls). The original wording follows.** Two numbers are missing and neither is optional. (1) **Bytes.** What does one `cudaGraphExec_t` cost on this model at this bucket set, and what does the fold actually return — device allocation at steady state, ON versus OFF, on a pinned-clock idle box. Every W4 and W5 figure counts executables. (2) **The probe cost at real bucket churn.** W5 probed once per fold on a 2-3 bucket process. A serving grid with 7 or 11 buckets across nine capture drivers churns differently, and the probe runs on the capture path. Only then is the flip a measurement rather than a preference | [#1226](https://github.com/mudler/vllm.cpp/issues/1226) | it needs a leased GPU, pinned clocks, and PR #1232 landed first — the key it would default to is still a draft. Deliberately NOT decided by this record, exactly as W5's own hypothesis was not decided by W4's | +| ~~Landing PR [#1232](https://github.com/mudler/vllm.cpp/pull/1232)~~ **LANDED as `2c8f53d93`, and W6 measured that exact commit.** The original wording follows. It is a DRAFT, so nothing on `main` folds today and this row's saving is unreachable on every configuration that ships | [#1226](https://github.com/mudler/vllm.cpp/issues/1226) | the experiment measured whether the key works before asking whether it should land, which is the correct order. It needs its own fresh review, and its red-first discrimination test — that the coarse key still SEPARATES two genuinely different topologies — is the load-bearing one | +| **PARTLY ANSWERED by W6, 2026-08-19:** 4 and 6 replay-time re-points occurred over 88 and 244 replays — arithmetic over the update and probe totals, not a counter — with every cell exiting 0 and byte-identical, so `Replay`'s transitivity assumption neither aborted nor changed a token. The counter is STILL owed, and so is a workload that alternates by design rather than by accident. Original wording: **An ALTERNATING workload, and a counter for replay-time re-points.** W5's three workloads captured their padded buckets in strictly descending order, so no group was ever replayed back to an earlier member and the only pair any coarse group issued was the one the probe tested. The reverse re-point — the pair that is untested from group size two onwards — is therefore unmeasured on a device, and `refused=0` does not cover it. The registry counts PROBE refusals only, so the ask is not even observable today: a replay-time refusal aborts on `Replay`'s `VT_CHECK` instead of incrementing anything | [#1226](https://github.com/mudler/vllm.cpp/issues/1226) | it needs a workload whose decode batch returns to a bucket it already captured, which the fixed-pool bench harness does not produce, plus a replay-side counter. Both ride with the lease the memory measurement above already needs | | A second driver, a second architecture, a second CUDA release. `refused=0` is one driver (`580.173.02`) on one GB10 under nvcc `13.3.73`. A refusal elsewhere costs a wasted probe rather than a wrong replay, but the rate is unknown and the default decision depends on it | [#1226](https://github.com/mudler/vllm.cpp/issues/1226) | no second CUDA box with a different driver is reachable from this session, and `orin:gpu0` stays BLOCKED on the CUDA 13 runtime | | **[#1283](https://github.com/mudler/vllm.cpp/issues/1283): greedy decode is not reproducible at concurrency 16.** W5's OFF/OFF control caught it and it VOIDED workload B. NOT this row's defect — `VT_CUDA_GRAPH_DEDUP` is unset in both cells that disagree — but this row's future A/Bs cannot use that workload shape until it is understood | [#1283](https://github.com/mudler/vllm.cpp/issues/1283) | it is a decode-path defect with its own issue and its own isolation plan. Folding it into this row would hide it behind a graph feature it has nothing to do with | | 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 | diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index 7f965fc7e..a0eaaae32 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -282,7 +282,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`, [#1226](https://github.com/mudler/vllm.cpp/issues/1226)) | EXACT key 3/3, 2/2, 2/2; COARSE key 2/3, 1/2, 1/2 | n/a | 1.00x EXACT; 0.67x / 0.50x COARSE | **SPLIT.** EXACT key folds NOTHING (`probes=0`); COARSE folds every bucket, `refused=0`, but is opt-in and **PR #1232 is UNLANDED**. No memory or speed number: clocks unpinned ([record](../.agents/benchmark-record.md)) | +| Decode-graph executables and device bytes, `VT_CUDA_GRAPH_DEDUP` (`ENG-CUDAGRAPH-DEDUP`, [#1162](https://github.com/mudler/vllm.cpp/issues/1162)) | COARSE key 3/7 and 5/11 execs; 15.40 vs 29.24 MiB nominal at 7 buckets | n/a | 0.43x / 0.45x execs; bytes NOT ESTABLISHED | **NEGATIVE, decided.** The fold engages; the saving fails its null control -- 0.42% of process at 7 buckets, none at 11. Default stays OFF. No time figure, clocks unpinned ([record](../.agents/benchmark-record.md)) | 35B steady-serving PSS is 3.53 GiB against vLLM's 13.3 GiB after the routed-expert host mirror is freed once the device Marlin resident is built. diff --git a/docs/STATUS.md b/docs/STATUS.md index c5c461c39..c0fe12221 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -1756,9 +1756,11 @@ 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 GB10 A/B of 2026-08-18 split: byte-identical 10/10, [#1184](https://github.com/mudler/vllm.cpp/issues/1184) gone, and the fold never engaged because the shipped key carries the padded batch dimension. +`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. Byte-identical on GB10 over 10/10 comparisons, [#1184](https://github.com/mudler/vllm.cpp/issues/1184) gone. The exact key folds nothing, because it carries the padded batch dimension; the coarse key folds and is measured below. Default OFF, and the flag stays OFF on measurement rather than on silence. -`ENG-CUDAGRAPH-DEDUP` coarse key: the [#1226](https://github.com/mudler/vllm.cpp/issues/1226) hypothesis is CONFIRMED on GB10. Drop the launch dimensions and every bucket folds — 3 graphs to 2 execs, 2 to 1, 2 to 1 — with `cudaGraphExecUpdate` probed once per fold and `refused=0`. It is opt-in, default OFF, and PR #1232 has NOT landed, so nothing on `main` folds today. No memory or speed number: clocks were unpinned and nobody measured bytes saved. +`ENG-CUDAGRAPH-DEDUP` coarse key ([#1226](https://github.com/mudler/vllm.cpp/issues/1226)): landed and MEASURED, and the benefit is NEGATIVE. At the shipped bucket sets it folds 7 captures to 3 execs and 11 to 5, so the fold is real. The device-byte saving is not: 13.83 MiB nominal at 7 buckets (0.42% of a 3.25 GiB process), nothing at 11, and smaller than the disagreement of its own null control. The driver refuses 43% of probes at 7 buckets and 73% at 11, always on TOPOLOGY — the decode graphs are 376 and 404 nodes, so a coarser key makes false candidates, not folds. + +`ENG-CUDAGRAPH-DEDUP` defaults: both flags stay OFF. `VT_CUDA_GRAPH_DEDUP_COARSE_KEY` alone is a no-op, because the coarse key is read only inside the dedup path. **No time figure is attributable** — the clock pin was refused inside the lease. The feature is also still unreachable on the DEFAULT serving path, which captures no decode graph ([#1179](https://github.com/mudler/vllm.cpp/issues/1179)). **Platform SELECTION is the one non-additive site, and is now gated.** A platform missing from `CurrentPlatform()`'s hardcoded walk registers and answers