From fcd7a9b2f1d992e7b477c1cc11940ea665a1be89 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Tue, 18 Aug 2026 12:19:27 +0000 Subject: [PATCH] record(ENG-CUDAGRAPH-DEDUP): the device A/B ran, and the fold it was filed for NEVER HAPPENS (#1162, #1184, #1226) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The owed W4 gate ran on `dgx:gpu0` on 2026-08-18 (GB10, driver 580.173.02, nvcc 13.0.88, `rc` job f88d484b, one binary, `VT_ASYNC_RUNNER=0` and `VT_DECODE_GRAPH_STATS=1` throughout, the only variable `VT_CUDA_GRAPH_DEDUP`). It came back with two opposite verdicts, and this is the record of both. No engine code changes here. The gated commit is `72de552c8`, not the merge `2a976eb9f`. The row squashed, so the gated tree is not an ancestor of what landed; what carries the claim across that gap is that all four dedup sources are byte-identical at the two commits, which was verified rather than assumed. CORRECTNESS PASSES, and #1184 is closed by this run rather than by the CPU suite that could never see it: 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 across 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. THE BENEFIT IS REFUTED for exactly the case the row was filed for. `N == M` in every ON cell, now over 2 and 3 distinct padded buckets per process where the first attempt managed one and could prove nothing, with the registry count climbing 1→1, 2→2, 3→3. The cause was pre-registered before the run and is structural: `AppendKernelPayload` hashes the grid and block dimensions and the memcpy payload hashes the copy extent, so the padded batch dimension is in the KEY, no candidate group forms, and `cudaGraphExecUpdate` is never attempted. That contradicts the row's own header comment, and SGLang keys the same fields, so whatever folds upstream is not decode buckets either. No throughput or memory number is recorded on any axis, for two independent reasons either of which suffices: the clocks were not pinned, and the ON arm allocated exactly as many executables as OFF, so there is no delta to claim. THE MATRIX ROW STAYS `ACTIVE`, argued rather than inherited. `DONE` would owe an `## Outcome` claiming a delivered capability, and the capability measurably folds nothing on the only driver exercised; banking that as done is how a refuted premise becomes a number people quote. `PARTIAL` means the implementation is also known to omit upstream behavior, and we omit none — we key the same fields SGLang keys, faithfully. The gap is a design question about the key, which is a live hypothesis with an owner, and that is what `ACTIVE` means. The summary-table counts are therefore unchanged, because no state moved. The open question is filed as #1226 and deliberately NOT decided here: a coarser key that keeps the function addresses and the topology but drops the launch dimensions and the memcpy extents would let two padded buckets form a candidate group at all, and because the registry probes with `cudaGraphExecUpdate` BEFORE it folds, a wrong grouping costs one wasted probe and a private executable rather than a wrong replay. That makes it a cost question needing its own spec and red-first evidence. "Unreachable with THIS key" is not "unreachable", and AGENTS.md forbids declaring a ceiling. Four honest gaps ride into every surface rather than into none: per-shape replay counts are unavailable because the driver prints a total, so workload B's ~30-per-shape is arithmetic and is labelled as arithmetic; the driver's "N captured size(s)" line counts slots, not captures; the container's own cuBLASLt was never re-tested at CUDA 13.0 because the staged cu130 prefix was probed first and worked, so the shim's original reason may no longer apply; and only the Qwen3 dense decode driver was exercised. The supporting `orin:gpu0` lane is BLOCKED, cleanly: the Jetson 540.4.0 driver cannot run a CUDA 13 runtime (`cudaGetDeviceCount err=35`), and a CUDA 12.x toolkit is the untried route. Checkers run from this worktree: `check-agent-record.py` OK (ENGINE=162 unchanged), `check-public-doc-tables.py` OK, `check-issue-index-append-only.py --base origin/main --head HEAD` OK. `agent-preflight.sh` was NOT run: it has been killed repeatedly on this box, so the component checkers were run instead and are named here. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code] --- .agents/benchmark-record.md | 195 +++++++++++++++++++++++++++ .agents/engine-matrix.md | 2 +- .agents/issue-index.md | 1 + .agents/specs/eng-cudagraph-dedup.md | 128 ++++++++++++++++-- docs/BENCHMARKS.md | 2 +- docs/STATUS.md | 2 +- 6 files changed, 316 insertions(+), 14 deletions(-) diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md index 0728bc454..91b3578a5 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -22765,3 +22765,198 @@ verdict above therefore rests on the recorded reading, and the only way to re-derive it is to re-run the committed source under the recipe above on a leased GPU. Nothing here was re-run for this record: the repair pass that added it had no GPU lease. + +## ENG-CUDAGRAPH-DEDUP W4 — the device A/B: correctness PASSES, and the fold NEVER HAPPENS (2026-08-18, `row/ENG-CUDAGRAPH-DEDUP-RESULT`, gated commit `72de552c8`, GB10 sm_121a, #1162 / #1184 / #1226) + +**No throughput or memory number is recorded, on any axis, and none is owed.** Two +independent reasons and either alone is sufficient: the clocks were not pinned, and the +`VT_CUDA_GRAPH_DEDUP=1` arm allocated exactly as many graph executables as the OFF arm, +so there is no memory delta to claim. The `replay branch avg` figures in the cell logs +(0.033-0.120 ms/step) are diagnostics, not a measurement. This row was never a +throughput row; its load-bearing gate is byte-identity, and its headline number is an +executable count. + +### The commit that was gated is NOT the commit that landed + +`72de552c8` was built and run. The merge is `2a976eb9f`, and the row squashed, so the +gated tree is not an ancestor of it. What carries the claim across that gap is a content +equality that was checked rather than assumed: all four dedup sources — `src/vt/graph_dedup.h`, +`src/vt/graph_dedup_runtime.h`, `src/vt/graph_dedup_latch.h` and +`src/vt/graph_dedup_signature.h` — are **byte-identical** at the two commits. The merge +commit itself was not executed, and this record does not claim it was. + +### Environment + +Device `dgx:gpu0`, leased through `rc`, job `f88d484b`, never ssh. Worker pod +`rc-worker-4b8lj`, aarch64, root, 20 cores. boot_id +`1cf6179f-0150-4052-b507-506fd6751953` for the build and every cell. GPU NVIDIA GB10, +driver `580.173.02`, `clocks.sm` 208 MHz idle / 3003 MHz max, `applications.graphics` +2418 MHz, persistence Disabled. **Clocks NOT pinned.** loadavg `1.82 3.22 4.07` at +series start and `2.88 3.98 3.98` at series end, per-cell range 2.94-3.45; the box was +quiet at both ends. + +**The in-pod toolkit CHANGED under us between the two attempts.** +`/usr/local/cuda/bin/nvcc` measured **CUDA 13.0.88** here and **13.3.73** on the same +pod and the same boot at the first attempt, because a neighbouring session held the box +for about two hours immediately before with a job named `oracle-build130`. The toolkit +is a property of the shared pod, not of the job. 13.0.88 is the version the recorded dgx +gate stack names. The lesson generalizes past this row: a job that needs a known toolkit +must ASSERT it, not assume it. This one recorded it. + +Runtime cuBLASLt was the staged cu130 `/tmp/tsite/nvidia/cu13/lib/libcublasLt.so.13` +(629,945,952 bytes), selected by a smoke probe (rc=0, 0 cuBLAS errors, a graph +captured). **Honest gap:** the probe tried that prefix FIRST and it worked, so the +container's own cuBLASLt was never re-tested at 13.0. The reason the shim existed — a +13.6.0.2 cuBLASLt that could not capture — may no longer apply. + +### Build recipe + +```sh +cmake -S /tmp/dedup72/src -B /tmp/dedup72/build -G Ninja \ + -DCMAKE_BUILD_TYPE=Release -DVLLM_CPP_CUDA=ON -DVLLM_CPP_CUDA_ARCHITECTURES=121a \ + -DCMAKE_CUDA_COMPILER=/usr/local/cuda/bin/nvcc \ + -DVLLM_CPP_CUTLASS_DIR=/tmp/dedup72/cutlass -DVLLM_CPP_TRITON=ON \ + -DVLLM_CPP_BUILD_TESTS=OFF -DVLLM_CPP_BUILD_EXAMPLES=ON +ninja -C /tmp/dedup72/build -j 4 vllm-bench +``` + +`rc=0`. All four mandatory production-stack assertions present in the configure log: +`CUTLASS found … enabling sm120a NVFP4 cutlass GEMM`; `FlashAttention-2 +prefill/decode: ENABLED for arch(es) [121a]`; `Triton AOT: gdn_deltah_h48 <- sm_121a`; +`VLLM_CPP_CUDA_ARCHITECTURES:STRING=121a`. Built on local disk, never the NAS. + +| Artifact | sha256 | +|---|---| +| `vllm-bench` binary, copied out of the build tree | `e166ed8d7b39ff131bf832bdef0c3ddb26b0e38ea7cddf3eedb54e7d417666fb` | +| model `model.safetensors`, Qwen3ForCausalLM 0.6B bf16 | `11293257a8df593c154a8ecd5fc039f3076de35411e35f06d41b471e136f6641` | +| source tar `src-72de552c8.tar` | `5e940df5bd94edd94da19048b8c864e697b100971f3fe91c9c35edda337ce62f` | + +Evidence directory `/mnt/nas_share/rc/dedup-gate2/` — `EVIDENCE.md`, `logs-ab/`, +`out-ab/`, `build72.sh`, `run72.sh`, `probe_orin.sh`. + +### Method, and the knob the FIRST attempt got wrong + +One binary in every cell. The only variable is `VT_CUDA_GRAPH_DEDUP`. +`VT_ASYNC_RUNNER=0` in every cell, and that is THE knob: `qwen3.cpp` +`DenseDecodeGraphForward` declines the graph while `input.device_token_ids != nullptr` +(the #323 mitigation), and that pointer is owned by the async RUNNER, not by +`VT_ASYNC_SCHED`. The first attempt's runs 1-2 were VOID because they used the wrong +knob and captured no graph at all. `VT_DECODE_GRAPH_STATS=1` throughout. The compared +artifact is `--output-token-ids`, the generated identifiers per request index — a real +byte artifact, not a summary statistic and not an exit code. + +**Multi-bucket design, which the first attempt lacked.** Buckets come from +`vllm::DecodeGraphSizes(max_num_seqs)`; `examples/bench` sets `max_num_seqs = +--concurrency` and refills to `concurrency` each poll, so the shapes seen are `pad(C)` +and `pad(N mod C)`. Workload A "drain" C=24 N=24 in=64 out=40 seed=4242 captured +24/16/8; B "tail" C=16 N=21 out=32 seed=777 captured 16/8; C "tail2" C=32 N=49 out=24 +seed=31337 captured 32/24. The first attempt captured ONE size in every cell, so its 1:1 +ratio measured nothing. + +### The 12 cells + +| cell | DEDUP | exit | dedup final line | sizes | replays | ids sha256 | +|---|---|---|---|---|---|---| +| a_off_a | unset | 0 | (none) | 24 16 8 | 33 | `d3b7028b…` | +| a_off_b | unset | 0 | (none) | 24 16 8 | 33 | `d3b7028b…` | +| a_on_a | 1 | 0 | captured 3 graphs, deduped to 3 execs | 24 16 8 | 33 | `d3b7028b…` | +| a_on_b | 1 | 0 | captured 3 graphs, deduped to 3 execs | 24 16 8 | 33 | `d3b7028b…` | +| a_zero | 0 | 0 | (none) | 24 16 8 | 33 | `d3b7028b…` | +| b_off_a | unset | 0 | (none) | 16 8 | 60 | `02a1add6…` | +| b_off_b | unset | 0 | (none) | 16 8 | 60 | `02a1add6…` | +| b_on_a | 1 | 0 | captured 2 graphs, deduped to 2 execs | 16 8 | 60 | `02a1add6…` | +| b_on_b | 1 | 0 | captured 2 graphs, deduped to 2 execs | 16 8 | 60 | `02a1add6…` | +| c_off_a | unset | 0 | (none) | 32 24 | 43 | `4f8714db…` | +| c_on_a | 1 | 0 | captured 2 graphs, deduped to 2 execs | 32 24 | 43 | `4f8714db…` | +| c_off_b | unset | 0 | (none) | 32 24 | 43 | `4f8714db…` | + +Every cell: successful requests = N, `empty_rows` 0, token counts 960 / 672 / 1176. + +### Result 1 — #1184 is FIXED, on the device + +All 12 cells exit 0. `grep -c "invalid device function"` and `grep -c "engine-fatal"` +return **zero** in every cell log. On the pre-fix head `e4ce5571a` every `dedup=1` cell +died after exactly one replay. The ON arms now replay as often as the OFF arms: **60 = +60** on B, **33 = 33** on A, **43 = 43** on C. A CPU suite drives a fake runtime and +cannot observe the real latched error, so this run is what closes the issue, not the +suite that proved the guard's structure. + +### Result 2 — byte-identity, 10/10, controls first + +```text +-- OFF/OFF controls, which must pass before any OFF-vs-ON comparison means anything +IDENTICAL a_off_a == a_off_b IDENTICAL b_off_a == b_off_b +IDENTICAL c_off_a == c_off_b +-- OFF vs ON +IDENTICAL a_off_a == a_on_a IDENTICAL a_off_a == a_on_b +IDENTICAL a_on_a == a_on_b IDENTICAL a_off_a == a_zero +IDENTICAL b_off_a == b_on_a IDENTICAL b_off_a == b_on_b +IDENTICAL c_off_a == c_on_a +``` + +The three workloads hash to three DIFFERENT values, so the identity is not the vacuous +kind a constant artifact would produce. + +### Result 3 — THE NEGATIVE: the fold never happens, and the cause is structural + +`N == M` in every ON cell, now with 2 and 3 *distinct* padded buckets per process: + +```text +a_on_a / a_on_b: captured 3 graphs, deduped to 3 execs sizes=[24 16 8] +b_on_a / b_on_b: captured 2 graphs, deduped to 2 execs sizes=[16 8] +c_on_a: captured 2 graphs, deduped to 2 execs sizes=[32 24] +``` + +The registry logs one line per registration and the count CLIMBS — `1 -> 1`, `2 -> 2`, +`3 -> 3` — which is the proof that more than one capture reached it. This was +pre-registered in `EVIDENCE.md` BEFORE the run rather than reasoned backwards from it: +`AppendKernelPayload` hashes `(func, gridDim.{x,y,z}, blockDim.{x,y,z}, sharedMemBytes)` +(`src/vt/graph_dedup_runtime.h:121-128`) and the memcpy payload hashes the copy extent. +The padded batch dimension is in BOTH, so two decode buckets never share a key, no +candidate group forms, and `cudaGraphExecUpdate` is **never attempted**. + +**This refutes the row's own premise.** `graph_dedup.h`'s header comment says the fold +exists because "the decode graphs of two padded batch sizes are usually the same node +topology with different parameters", and the signature as written cannot group exactly +those. SGLang hashes the same fields (`cuda_graph_dedup_mixin.py:105-114`), so whatever +folds upstream is not decode buckets either. The machinery is correct and does nothing +on the workload it was built for. `VT_CUDA_GRAPH_DEDUP` therefore stays OFF: a default +is a measurement, and this measurement does not support one. + +**The next traceable hypothesis, filed rather than decided** ([#1226](https://github.com/mudler/vllm.cpp/issues/1226)): +a coarser key that keeps the function addresses and the topology but drops the launch +dimensions and the memcpy extents would let two padded buckets form a candidate group at +all. The probe-before-fold design means such a key costs one wasted `cudaGraphExecUpdate` +probe and a private executable when the driver refuses, rather than a wrong replay, so it +is a cost question and not an obviously unsafe one. "Unreachable with THIS key" is not +"unreachable", and nothing here declares a ceiling. + +### Honest gaps a reviewer must weigh + +- **Per-shape replay counts are unavailable.** The driver prints a TOTAL only. Workload + B's ~30 replays per shape is arithmetic over that total (60 replays, 2 captured + shapes, a 16-request wave then a 5-request wave each decoding 31 steps). Stated as + arithmetic, not as a measurement. The ON and OFF arms report the SAME totals. +- **The driver's "N captured size(s)" figure counts SLOTS, not captures.** Workload A + reports 6 sizes and emits only 3 `captured dense decode graph for padded size S=` + lines, because the smallest buckets appeared for too few steps to pass warm-up. +- **The container's own cuBLASLt was never re-tested at CUDA 13.0**, as recorded above. +- **Only the Qwen3 dense decode driver was exercised.** Whether any other capture site, + or two models sharing the process-singleton registry, can produce a fold is untested. +- **The model is a Qwen3ForCausalLM 0.6B derivative from the NAS**, not a SACRED gate + checkpoint. +- **The merge commit `2a976eb9f` was not run**; `72de552c8` was, with the four dedup + sources verified byte-identical between them. + +### The supporting `orin:gpu0` lane — BLOCKED, cleanly + +Two `rc run` jobs on `orin:gpu0`. Job 1 exited `FATAL_NO_CUDA_PKG` because the script +hardcoded the ubuntu2404 suite without checking. Job 2 measured the worker properly: +Ubuntu 24.04.4, root, 12 cores, 29 GiB, `gcc`/`g++`/`cmake`/`ninja`/`git`/`python3`/`curl` +present, `/workspace` a LOCAL 1.8T disk and not the NAS, driver "NVIDIA UNIX Open Kernel +Module for aarch64 540.4.0", `/dev/nvidia0` present, `nvcc` 13.0 V13.0.88 present. The +CUDA smoke then failed: `cudaGetDeviceCount err=35 CUDA driver version is insufficient +for CUDA runtime version`. The Jetson 540.4.0 driver cannot run a CUDA 13 runtime. The +untried route is a CUDA 12.x toolkit for that driver; it was not pursued, because the +dgx gate had already answered the question orin was there to support. No lease held; +`orin:gpu0` returned to ready. diff --git a/.agents/engine-matrix.md b/.agents/engine-matrix.md index cf121531e..d1b668b78 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. STILL OWED: the device same-binary A/B proving a deduped replay token-identical AND that #1184 is gone on a device, the exec-count ratio, device-tier signature stability/discrimination tests, probing `current_raw` instead of `raws.front()` to retire the update-transitivity assumption, the ROCm compile, and reaching the feature from the default serving path at all — the async runner captures no decode graph, so dedup engages only under `VT_ASYNC_RUNNER=0` | [eng-cudagraph-dedup.md](specs/eng-cudagraph-dedup.md); analysis [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `ACTIVE` | `CLAIM-ENG-CUDAGRAPH-DEDUP` ([#1162](https://github.com/mudler/vllm.cpp/issues/1162)) | +| `ENG-CUDAGRAPH-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, so dedup engages only under `VT_ASYNC_RUNNER=0` | [eng-cudagraph-dedup.md](specs/eng-cudagraph-dedup.md); analysis [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `ACTIVE` | `CLAIM-ENG-CUDAGRAPH-DEDUP` ([#1162](https://github.com/mudler/vllm.cpp/issues/1162)) | | `ENG-CUDAGRAPH-BREAK` | One shared `vt` capture seam that accepts BREAK POINTS, so a forward containing a host-dependent op is still graphed instead of falling out entirely — and so the NINE hand-rolled drivers become one (count corrected 2026-08-18, [#1179](https://github.com/mudler/vllm.cpp/issues/1179); `9bc4d7f44` recorded eight). **Coverage AND CORRECTNESS row, not a throughput row** | T1 | mirror vLLM `CUDAGraphMode.PIECEWISE` splitting at `splitting_ops` (`vllm/config/compilation.py:60-63,517,615,630` @ `555967922`); construction from SGLang BCG `python/sglang/srt/model_executor/runner_backend_utils/breakable_cuda_graph/breakable_cuda_graph.py:204-243,246-274,309-333,335-367` @ `f63458b5be` (decorator + runtime stream capture, no compiler); its unit suite `test/registered/cuda_graph/breakable/test_breakable_cuda_graph.py:30,172,230` (305 lines, 11 unit cases) is mapped case for case in the spec's `## Tests to port` | all-or-nothing today: `src/vllm/v1/worker/gpu/runner.cpp:1338-1341` routes only `pure_decode`; drivers `qwen3_5.h:275`, `qwen3_5_dense.h:391`, `qwen3_moe.h:117`, `qwen3.h:243`, `deepseek_v2.h:324`, `voxtral.h:126`, plus `deepseek_v4.cpp`, `laguna.cpp` — and the spike found the NINTH already written, `src/vllm/model_executor/models/qwen3_dflash.cpp:771,1091`. The re-derivation is measured, not asserted: `StepDevInputs` (`src/vllm/model_executor/models/qwen3_5.cpp:3894`, the persistent DEVICE input path) exists in ONE driver and `grep -c` returns 0 in `qwen3_moe.cpp`, `qwen3.cpp`, `deepseek_v2.cpp` and `voxtral.cpp`, which is why `src/vllm/model_executor/models/qwen3.cpp:961-986` DECLINES the graph outright when the async device-token mirror is live. **That decline is why this is also a CORRECTNESS row** ([#1179](https://github.com/mudler/vllm.cpp/issues/1179)): a SHIPPED model has already lost its decode graph to the duplication, on the driver's own measurement (`depth-1, graph ON PASS 78/78`; `depth-2, graph OFF PASS 82/82`; `depth-2, graph ON FAIL, slots 1-3 degenerate`), and the fix its comment names is the sibling's `StepDevInputs`. The row still makes NO throughput claim: the prefill refutation on the `ENG-CUDAGRAPH` row (3.8% host idle, >96% GPU-busy, 92.5% glue) stands unchanged | owed: bit-exactness vs eager on every migrated model over MORE than one replay, on a real GPU (W2); the host-lifetime contract of `decode-graph-scratch-uaf-2026-07-18.md` enforced AT the seam — D1's INPUT half, making the intermediates a segment reads unavailable to the `DevicePool` free list (W2); the auxiliary-stream auto-join before every segment close (`:353-361`, spec D10), live at `src/vllm/model_executor/models/qwen3_5.cpp:6254-6255,6384` and `src/vllm/model_executor/models/laguna.cpp:2572-2576,2612` (W4, W5). **Delivered by W1** ([#1192](https://github.com/mudler/vllm.cpp/issues/1192)): the reachability mutation (performed; deleting the call site reds `tests/vllm/models/test_qwen3_break_point.cpp` and leaves the unit suite green); the ported SGLang unit cases with their arithmetic chains and post-replay assertions; and the break-function OUTPUT writeback (`replay_fn`/`_copy_output` `breakable_cuda_graph.py:231-235,172-201`, spec D9), whose destination is a `vt::BreakSlot` the seam owns rather than a caller reference it cannot outlive | spec [eng-cudagraph-break.md](specs/eng-cudagraph-break.md) (W0 spike DONE 2026-08-18: the existing `vt` capture vocabulary `include/vt/backend.h:208-222` expresses a SEGMENTED capture with NO new virtual, because `EndCaptureGraph` stores nothing (`src/vt/cuda/cuda_backend.cu:225-232`); a break point is expressible with one `thread_local` capture pointer plus a free function, no compiler and no decorator); **W1 DONE 2026-08-18 ([#1192](https://github.com/mudler/vllm.cpp/issues/1192)): the seam LANDS** — `vt::BreakableGraph`, `vt::GraphCaptureScope` and `vt::GraphBreak` (`include/vt/breakable_graph.h`, `src/vt/breakable_graph.cpp`), the SGLang unit suite ported case for case (`tests/vt/test_breakable_graph.cpp`, 24 cases / 163 assertions, re-derived 2026-08-18 by `ninja test_breakable_graph && ./build/tests/test_breakable_graph`; the recorded 14/81 never re-derived at any head of this branch), and ONE break point registered at the DENSE ATTENTION ENTRY of `Qwen3ForCausalLM` (`src/vllm/model_executor/models/qwen3.cpp`, inside `RunLayer`). **The exit criterion W0 deliberately left open is ANSWERED on a leased GPU:** `cudaStreamEndCapture` then `cudaStreamBeginCapture` on the SAME stream mid-forward with EAGER work between is LEGAL under `cudaStreamCaptureModeThreadLocal` (`src/vt/cuda/cuda_backend.cu:204-206`) — `orin:gpu0` via an `rc` lease, driver 12060, 3 replays with fresh inputs, 0 mismatches, bare zero-work re-begin legal too. G2 reachability is `tests/vllm/models/test_qwen3_break_point.cpp`, which drives the production `Qwen3DenseModel::Forward` with a scope open and counts `num_hidden_layers + 1` segments (mutation: delete the call site ⇒ 1 segment ⇒ RED), and holds G4 in the same case at 500 logits / 0 differing bit for bit. STAGED SLICE, named: the scope and the container are not yet ENTERED from a production step — no driver opens a scope until W2 migrates `Qwen3DenseDecodeGraph` — and the spec's `## Owed` lists it with W2 as owner, alongside the D10 auxiliary-stream auto-join (W4/W5), G5's ROCm/Tenstorrent arms (W3) and G1 on a real GPU (W2). **The capture-failure drain is NOT among them: it landed HERE**, as behaviour (`std::uncaught_exceptions()` compared against the depth recorded at scope entry, so a break function or ordinary model code throwing mid-capture destroys the partial container instead of handing back a forward that reports `captured() == true`) and as three gated arms (tests 13a, 13b, 13c). The spec's `## Owed` strikes the item through and reads DELIVERED in W1; this cell said the opposite until 2026-08-18 because `cba969857` re-derived field 6 alone. Still NO throughput claim; analysis [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `ACTIVE` | `CLAIM-ENG-CUDAGRAPH-BREAK-W1`; [#1163](https://github.com/mudler/vllm.cpp/issues/1163), [#1192](https://github.com/mudler/vllm.cpp/issues/1192), [#1020](https://github.com/mudler/vllm.cpp/issues/1020) | | `ENG-CUDAGRAPH-DIFFUSION` | Capture the LTX-2.5 denoise loop (fixed shapes, many identical iterations — the ideal graph target). **BLOCKED, and the blocker is ours:** the render does almost no device compute to capture | T2 | SGLang enabled BCG on this shape AFTER our pin — LTX-2 H200 two-stage 10.75s->6.90s (`d4be483efb`), SANA 1024px -26% (`6c7498113f`), SANA denoise 0.73->0.457s (`56ef810cad`). Dated events, NOT pinned evidence; their win is mostly PyTorch host tax we do not pay | NO capture at all: `grep` for capture across `src/vllm/model_executor/models/ltx2*.cpp` returns nothing | blocked by [#1024](https://github.com/mudler/vllm.cpp/issues/1024) (GPU util **exactly 0 in 321 of 347 samples**, 1.00 core of 20 held for 17+ min after staging), [#1007](https://github.com/mudler/vllm.cpp/issues/1007) (VAE decode has no device arm), [#1087](https://github.com/mudler/vllm.cpp/issues/1087) (**57-66% of wall** is ONE resolution-CONSTANT serial host phase), [#1010](https://github.com/mudler/vllm.cpp/issues/1010) (no phase-boundary log). Decision point is a MEASUREMENT of GPU-busy vs wall once device-resident, not an implementation | [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `INVENTORIED` | [#1164](https://github.com/mudler/vllm.cpp/issues/1164) | | `ENG-BATCH-INVARIANT` | Opt-in deterministic execution across scheduler batch sizes (`VLLM_BATCH_INVARIANT=1`): batch-invariant matmul/norm/attention/collectives plus persistent-scheduler NVFP4; production default remains off | T1 | default/env `vllm/envs.py:89,576-578`; initialization `vllm/v1/worker/gpu_worker.py:1262`; NVFP4 dispatch `csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_sm120_kernels.cu:212-220`; suite fixture `tests/v1/determinism/conftest.py:9-12`; operator/e2e `tests/v1/determinism/test_nvfp4_batch_invariant_scaled_mm.py`, `tests/v1/determinism/test_nvfp4_batch_invariant.py` @ `702f481` | - | [W3-C3R executed contract](specs/nvfp4-persistent-plan-cache.md#w3-c3r-batch-shape-localization-and-gate-correction-2026-07-13): production-default ours and vLLM both change outputs across batch shapes; no local opt-in implementation is claimed | `planned: specs/batch-invariant-execution.md` | `INVENTORIED` | - | diff --git a/.agents/issue-index.md b/.agents/issue-index.md index 5fc4b1cd8..4fef4a09b 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -372,3 +372,4 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1190](https://github.com/mudler/vllm.cpp/issues/1190) | `ENV-AGNOSTIC-CAMPAIGN` | One operator's hosts, share paths and addresses are written into 227 tracked files where a `.env` placeholder belongs, so a second developer who follows the protocol documents is told to reach a box on one home network. Re-derived at `fd64c76ee` with `git grep -cIE 'dgx\.casa\|nas_share\|192\.168\.\|thor:gpu0'`: `dgx.casa` 203 files, `nas_share` 32, `192.168.` 24, `thor:gpu0` 11. The mechanism was already there and almost unused: `.env.example` is tracked, `.env` and `.agents/developer-preferences.md` are ignored, `scripts/agent-onboard.py --env-set` already records one answered value and refuses an undeclared key, and `${VLLM_SOURCE}` and `${GPU_LOCK}` already resolve in 59 and 50 files while `${DEVICE_ARCH}`, `${DEVICE_TOOLKIT_ROOT}` and `${DEVICE_COMPILER}` resolve in none. This SCOPING row lands the rule and the mechanism, not the sweep. The rule is a substitution test: replace the literal with a second developer's value and if the sentence stays true it is CONFIGURATION and becomes `${KEY}`, and if it becomes false it is PROVENANCE and stays literal. That reverses the issue's own ranking, because all 34 hits in the densest guide file `.agents/environment.md` are provenance or named-profile definition and none is configuration, so density does not predict the defect and a blind `sed` would falsify records. Landed: three new keys `GATE_CHECKOUT`, `SHARED_STORAGE_ROOT` and `GATE_DEVICE` derived from the literals that recur and map to no existing key; the create-on-first-use route in `scripts/agent-start.py`, which printed `environment: missing` as a status label and then listed next actions that never mentioned it, so the fallback in practice was a host name copied from a document; the matching obligation in `AGENTS.md`; and one worked example. The worked example `scripts/dgx-bringup.sh` found a live defect rather than a cosmetic one: it defaulted `CUTLASS_DIR` to `$HOME/cutlass_probe` while `.agents/environment.md:389` records `$HOME/cutlass-4.5.0` as mandatory on the same box, and a configure that misses CUTLASS silently drops the sm120a NVFP4 GEMM and FlashAttention-2, which that file measures as moving the SACRED `test_qwen27_paged_engine` from 235/235 to 234/235 with the source untouched, so a stale hard-coded default is a false green. Its new test also caught that `set -a; . ./.env; set +a` over an `.env.example` copy blanks a value the caller exported, so the process-environment contract is now executable. Waves `ENV-AGNOSTIC-W1-TOOLING` through `ENV-AGNOSTIC-W5-LEDGERS` own the sweep, partition all 227 files, and open their own issues. Spec [`env-agnostic.md`](specs/env-agnostic.md) | bug | | [#1193](https://github.com/mudler/vllm.cpp/issues/1193) | `SPEC-DSPARK-QWEN3-ROUTING` | A Qwen3 DSpark draft declaring `architectures=["DSparkDraftModel"]` with `model_type` `qwen3` has no route. The pin forces every DSpark draft that is not `Qwen3DSparkModel` or `Gemma4DSparkModel` onto `model_type` `deepseek_v4` (`vllm/config/speculative.py:934-944` @ `555967922`), and vLLM PR 52197 (merged 2026-08-17 at `7075ddac`) replaced that with a leading branch normalizing the pair to `Qwen3DSparkModel`. We diverge from BOTH: the forced rewrite was never ported, so nothing in `src/vllm/entrypoints/model_loader.cpp` reads a draft config's `architectures` key at all, and `SpeculativeConfig::IsDsparkDraft` (`include/vllm/config/speculative.h:120-136`) has no production caller — every reference outside its header is in `tests/vllm/config/test_speculative_dspark.cpp:132-140`, and `ResolveSpecConfig` branches on `cli.method` alone. The checkpoint is real and gateable here: `RadixArk/Qwen3.8-27B-DSpark` at revision `85ef153be924f17ce4bf62726954eeaa4a73e854` carries exactly that config shape in one 2718576122-byte shard, drafting five layers for a 64-layer Qwen3.8-27B target | bug | | [#1213](https://github.com/mudler/vllm.cpp/issues/1213) | `ENV-LEASE-RUNTIME-STAGING` | `AGENTS.md` stated that a leased worker "has no compiler, no downloader and no Python, so it cannot produce a runtime in place", and `.agents/environment.md` carried the matching clause twice for `dgx:gpu0`. All three negatives are false. `rc describe dgx:gpu0` states that a job runs as root in an Ubuntu 24.04 container carrying `git`, `curl`, `wget`, `ssh`, `gcc`, `g++`, `make`, `cmake`, `ninja`, `pkg-config`, `python3`, `pip` and `venv`, and it instructs the reader to install anything missing; the one limit it names is the absent CUDA toolkit. Two jobs then compiled inside a lease on 2026-08-18: `claude/mudler-ubuntu-box/qwen38-gate` apt-installed `cuda-nvcc-13-0` from the `ubuntu2404/sbsa` lane and built this tree 1791/1791 to `BUILD_RC=0` (`/mnt/nas_share/rc/qwen38-gate/out-main/cfg.log` records `nvcc` 13.0.88 and `CUDA feature cutlass-fp8: ENABLED for [121a]`), and `/mnt/nas_share/rc/mtp_test/build.sh` cloned `github.com/mudler/llama.cpp` from inside a job and left a 97 MB `libggml-cuda.so` on the share. **Why it matters:** "the lease cannot produce a runtime in place" is the stated basis for treating the pinned vLLM oracle as unreachable from a lease, and that oracle is the denominator for every speed-parity number the project owes, so the premise needs re-testing rather than inheriting. This claims nothing about a model run; #1185 owns that and stays open. FIXED IN FLOW: the `AGENTS.md` paragraph and both `.agents/environment.md` clauses now say what the measurement supports, and the four real limits (no preinstalled CUDA toolkit, global installs leak until the pod restarts, CIFS `/workspace` holds no symlink so build in `/tmp` and `cp -rL`, and `-j 4` because unconstrained parallelism OOM-reboots the box) plus the host-versus-container egress distinction ride with the correction. | record | +| [#1226](https://github.com/mudler/vllm.cpp/issues/1226) | `ENG-CUDAGRAPH-DEDUP` | The owed device A/B for #1162 RAN on `dgx:gpu0` 2026-08-18 (GB10, driver 580.173.02, nvcc 13.0.88, `rc` job f88d484b, gated commit `72de552c8`, whose four dedup sources are byte-identical to the merged `2a976eb9f`) and it SPLIT. CORRECTNESS PASSED and closes [#1184](https://github.com/mudler/vllm.cpp/issues/1184): 12/12 cells exit 0, zero `invalid device function` and zero `engine-fatal` 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 three workloads hashing to three DIFFERENT values. THE BENEFIT IS REFUTED for the case the row was filed for: `N == M` in every `VT_CUDA_GRAPH_DEDUP=1` 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 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 of the first attempt. The cause is structural and was pre-registered before the run: `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 is IN THE KEY, no candidate group forms and `cudaGraphExecUpdate` is never attempted. That contradicts the row's own premise, 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. `VT_CUDA_GRAPH_DEDUP` therefore stays OFF and the default flip is NOT JUSTIFIED on this evidence. THE OPEN HYPOTHESIS this issue owns, deliberately not decided by the record that filed it: a COARSER key keeping the function addresses and topology but dropping the launch dimensions and memcpy extents would let two padded buckets form a candidate group at all, and the probe-before-fold design means a wrong grouping costs one wasted `cudaGraphExecUpdate` probe and a private executable rather than a wrong replay, so it is a cost question needing its own spec, a red-first discrimination test and a device probe-refusal rate. Also owed here: whether ANY other capture site or two models sharing the process-singleton registry can fold (only the Qwen3 dense decode driver was exercised), and a supporting `orin:gpu0` leg, BLOCKED because the Jetson 540.4.0 driver cannot run a CUDA 13 runtime (`cudaGetDeviceCount err=35`; a CUDA 12.x toolkit is the untried route). Honest gaps carried into the record: per-shape replay counts are unavailable, so B's ~30-per-shape is arithmetic over a printed TOTAL; the driver's "N captured size(s)" counts SLOTS not captures; the container's own cuBLASLt was never re-tested at CUDA 13.0. Full evidence [`benchmark-record.md`](benchmark-record.md) entry `ENG-CUDAGRAPH-DEDUP W4`; spec [`eng-cudagraph-dedup.md`](specs/eng-cudagraph-dedup.md) `## Outcome` | perf | diff --git a/.agents/specs/eng-cudagraph-dedup.md b/.agents/specs/eng-cudagraph-dedup.md index 04a9e6e7d..eebe438b1 100644 --- a/.agents/specs/eng-cudagraph-dedup.md +++ b/.agents/specs/eng-cudagraph-dedup.md @@ -10,12 +10,17 @@ is why the load-bearing gate below is byte-identity rather than a speed ratio. ## Now -`ACTIVE`. The shared registry, its CPU contract suite and the CUDA wiring land here, -together with the [#1184](https://github.com/mudler/vllm.cpp/issues/1184) repair: the -runtime binding is allowed to see calls fail and never consumed the runtime's latched -error, so the first unrelated kernel after a capture reported our refusal as its own -failure and every `VT_CUDA_GRAPH_DEDUP=1` run died. The device-level byte-identity A/B -and the ROCm leg are named under [`## Owed`](#owed) with the issue that owns each. +`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 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. The row is not `DONE`, because its stated saving +does not occur; it is not `PARTIAL`, because nothing upstream is omitted — SGLang keys +the same fields. It stays `ACTIVE` on one open, traceable hypothesis +([#1226](https://github.com/mudler/vllm.cpp/issues/1226)) and the items under +[`## Owed`](#owed). Full result: [`## Outcome`](#outcome). ## Scope @@ -218,11 +223,14 @@ raises rather than latching, so the whole error-latch class does not exist there with a negative control proving the instrument can fail. A proxy is not the gate: the CUDA 13 `cudaGraphGetEdges` break this file already took was invisible to exactly this kind of local check and only `cuda-fat-build` reported it. -6. **Device byte-identity A/B (owed, see below).** Same binary, `VT_CUDA_GRAPH_DEDUP` +6. **Device byte-identity A/B — RAN 2026-08-18, PASS; see [`## Outcome`](#outcome).** Same binary, `VT_CUDA_GRAPH_DEDUP` off then on, identical prompts and sampling, token-exact equality, with **more than one replay per padded bucket** — a first replay can be correct by accident, and the whole hazard of a shared executable is that the *second* visit to a shape is the one - that has to re-point it. + that has to re-point it. **Stated limit on how that clause was satisfied:** the driver + prints a replay TOTAL and no per-shape breakdown, so "more than one replay per bucket" + rests on the total over the captured shapes (workload B, 60 replays over 2 shapes, + ~30 each) and is arithmetic rather than a direct measurement. 7. **Not gated, deliberately:** throughput. This row must not be sold as a speed change. The reportable numbers are executable count and capture wall time. @@ -390,17 +398,115 @@ raises rather than latching, so the whole error-latch class does not exist there general lesson for this file: a header-shape check is only as current as the toolkit it ran against, and the CI job is the gate, not the local proxy. +## Outcome + +The device gate ran on 2026-08-18 on `dgx:gpu0` through an `rc` lease (job `f88d484b`, +pod `rc-worker-4b8lj`, boot_id `1cf6179f-0150-4052-b507-506fd6751953`), GB10, driver +`580.173.02`, nvcc `13.0.88`, runtime cuBLASLt the staged cu130 +`/tmp/tsite/nvidia/cu13/lib/libcublasLt.so.13`. Twelve cells of one binary +(sha256 `e166ed8d…7666fb`), `VT_ASYNC_RUNNER=0` and `VT_DECODE_GRAPH_STATS=1` +throughout, the only variable `VT_CUDA_GRAPH_DEDUP`. Recipe, hashes, per-cell table and +caveats: [`.agents/benchmark-record.md`](../benchmark-record.md), entry +`ENG-CUDAGRAPH-DEDUP W4`. + +**The commit gated is `72de552c8`, not the merge `2a976eb9f`.** The row squashed, so the +gated tree is not an ancestor of what landed. What carries the claim instead is a +content equality that was checked rather than assumed: all four dedup sources — +`src/vt/graph_dedup.h`, `graph_dedup_runtime.h`, `graph_dedup_latch.h` and +`graph_dedup_signature.h` — are byte-identical at the two commits. + +### What was measured + +1. **[#1184](https://github.com/mudler/vllm.cpp/issues/1184) is fixed on the device.** + All 12 cells exit 0. `grep -c "invalid device function"` and `grep -c "engine-fatal"` + return zero in every cell log. On the pre-fix head `e4ce5571a` every `dedup=1` cell + died after exactly one replay. The ON arms now replay as often as the OFF arms: + 60 = 60 on workload B, 33 = 33 on A, 43 = 43 on C. +2. **Byte-identity holds, 10/10.** The OFF/OFF controls passed first — `a_off_a == + a_off_b`, `b_off_a == b_off_b`, `c_off_a == c_off_b` — and then every OFF-vs-ON pair + was identical. The artifact is `--output-token-ids`, real generated identifiers, 0 + empty rows, 960 / 672 / 1176 tokens. The three workloads hash to three DIFFERENT + values (`d3b7028b…`, `02a1add6…`, `4f8714db…`), so the identity is not vacuous. +3. **The fold never happens.** `N == M` in every ON cell, now with two and three + *distinct* padded buckets per process, where the first attempt managed only one and + could therefore prove nothing: + +```text +a_on_a / a_on_b: captured 3 graphs, deduped to 3 execs sizes=[24 16 8] +b_on_a / b_on_b: captured 2 graphs, deduped to 2 execs sizes=[16 8] +c_on_a: captured 2 graphs, deduped to 2 execs sizes=[32 24] +``` + +The registry's own line climbs `1 -> 1`, `2 -> 2`, `3 -> 3`, which is the proof that +more than one capture reached it. + +### What was refuted, and why the row's premise was wrong + +The cause is structural rather than a tuning miss. `AppendKernelPayload` hashes +`(func, gridDim.{x,y,z}, blockDim.{x,y,z}, sharedMemBytes)` +(`src/vt/graph_dedup_runtime.h:121-128`) and the memcpy payload hashes the copy extent. +The padded batch dimension is in both, so two decode buckets never share a key, no +candidate group ever forms, and `cudaGraphExecUpdate` is **never attempted**. + +This refutes the row's own premise. `graph_dedup.h`'s header comment says the fold +exists because "the decode graphs of two padded batch sizes are usually the same node +topology with different parameters", and the signature as written cannot group exactly +those. SGLang hashes the same fields (`cuda_graph_dedup_mixin.py:105-114`), so whatever +folds upstream is not decode buckets either. The machinery is correct and does nothing +on the workload it was built for. + +### Why the defaults have their values + +`VT_CUDA_GRAPH_DEDUP` stays **OFF**. The flip was gated on this A/B, and the A/B says +the ON arm allocates exactly as many executables as the OFF arm on the only driver +exercised. A default is a measurement, and this measurement does not support one. + +**No throughput or memory number is recorded, on any axis.** Two independent reasons, +and either alone is sufficient: the clocks were not pinned (208 MHz idle, 3003 MHz max, +2418 MHz applications), and dedup allocated the same number of executables as OFF, so +there is no memory delta to claim. The `replay branch avg` figures in the logs +(0.033-0.120 ms/step) are diagnostics, not a measurement. + +### Honest gaps in this run + +- **Per-shape replay counts are unavailable.** The driver prints a total only. + Workload B's ~30 replays per shape is arithmetic over that total (60 replays, 2 + shapes), stated as arithmetic and not as a measurement. +- **The driver's "N captured size(s)" line counts SLOTS, not captures.** Workload A + reports 6 sizes and emits only 3 `captured … padded size S=` lines. +- **The container's own cuBLASLt was never re-tested at CUDA 13.0.** The smoke probe + tried the staged cu130 prefix first and it worked, so the original reason for the shim + — a 13.6.0.2 cuBLASLt that could not capture — may no longer apply. +- **Only the Qwen3 dense decode driver was exercised.** Whether any other capture site, + or two models sharing the process-singleton registry, can produce a fold is untested. +- **The in-pod toolkit changed under us**, from 13.3.73 at the first attempt to 13.0.88 + here, on the same pod and the same boot, because a neighbouring session held the box + first. A job that needs a known toolkit must assert it rather than assume it. +- **The supporting `orin:gpu0` lane is BLOCKED**, cleanly: the Jetson `540.4.0` driver + cannot run a CUDA 13 runtime (`cudaGetDeviceCount err=35`). A CUDA 12.x toolkit is the + untried route. It is recorded, not pursued, because the dgx gate answered the question + orin was there to support. + ## Owed +**DELIVERED by W4, 2026-08-18 — see [`## Outcome`](#outcome).** The device +byte-identity A/B: **PASS**, 10/10, OFF/OFF controls first. The `#1184` re-run +confirming that `VT_CUDA_GRAPH_DEDUP=1` completes a decode step on a device: **PASS**, +12/12 cells exit 0, zero `invalid device function`, zero `engine-fatal`. The +executable-count ratio: **DELIVERED AND NEGATIVE**, `N == M` in every ON cell over two +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. + | Item | Issue | Why not here | |---|---|---| -| The device byte-identity A/B and the executable-count measurement (W4) | [#1162](https://github.com/mudler/vllm.cpp/issues/1162) | needs a leased CUDA box; the CPU tier proves the registry's launch-sequence identity, not the device's | -| Flipping `VT_CUDA_GRAPH_DEDUP` on by default | [#1162](https://github.com/mudler/vllm.cpp/issues/1162) | gated on W4. A default is a measurement, not a preference | +| Flipping `VT_CUDA_GRAPH_DEDUP` on by default | [#1162](https://github.com/mudler/vllm.cpp/issues/1162) | **NOT JUSTIFIED on this evidence.** W4 delivered and the ON arm allocated exactly as many executables as OFF, so the flip would ship a probe cost for no measured saving. It is gated now on a key that can actually group two decode buckets ([#1226](https://github.com/mudler/vllm.cpp/issues/1226)), not on a rerun of the same A/B | | Probing `group.current_raw` rather than `raws.front()`, retiring the transitivity assumption above | [#1162](https://github.com/mudler/vllm.cpp/issues/1162) | it changes probe behaviour, and the device A/B is measuring the current one. Land it with the A/B rerun, not before | | Executable coverage for the signature builder's DEVICE half — stability and discrimination on a real `cudaGraph_t`, plus the five node-payload cases and their query escapes | [#1162](https://github.com/mudler/vllm.cpp/issues/1162) | the device-free half is now covered by `tests/vt/test_graph_dedup_runtime.cpp`; what remains needs a real `cudaGraph_t`, so it rides with the leased box the A/B already needs | -| Re-running the device A/B after the [#1184](https://github.com/mudler/vllm.cpp/issues/1184) repair, and confirming that `VT_CUDA_GRAPH_DEDUP=1` now completes a decode step | [#1184](https://github.com/mudler/vllm.cpp/issues/1184) | the CPU suite proves the guard's structure over a fake runtime; only a device can prove that `cudaGetLastError` clears the real latch and that the run survives. #1184 stays open until that run exists | | Reaching the feature from the DEFAULT serving path. Dedup engages only under `VT_ASYNC_RUNNER=0`, because the async path captures no decode graph at all ([#323](https://github.com/mudler/vllm.cpp/issues/323) mitigation) | [#1179](https://github.com/mudler/vllm.cpp/issues/1179) | the repair is the `StepDevInputs`-shaped one `ENG-CUDAGRAPH-BREAK` already owns; until it lands, this row's saving is unreachable on the configuration users serve with | | The ROCm leg's compile and run verification | [#41](https://github.com/mudler/vllm.cpp/issues/41) | no ROCm hardware or `hipcc` is reachable from this session and CI has no ROCm job, so the HIP wiring is written against the same shared header but is compile-unverified | +| **THE NEXT TRACEABLE HYPOTHESIS, not decided here.** A COARSER signature key — one that keeps the function addresses and the topology but drops the launch dimensions and the memcpy extents — would let two padded decode buckets form a candidate group at all. It is not obviously unsafe: the registry probes with `cudaGraphExecUpdate` BEFORE it folds, so a key that groups two graphs the driver then refuses costs one wasted probe and a private executable rather than a wrong replay. That makes it a cost question. It needs its own spec, a red-first discrimination test proving the coarser key still separates two genuinely different topologies, and a device run measuring the probe-refusal rate | [#1226](https://github.com/mudler/vllm.cpp/issues/1226) | it changes what the row optimizes, so it is a design decision with its own evidence, not a repair to fold into the record of the run that found it. AGENTS.md forbids declaring a ceiling, and "unreachable with THIS key" is not "unreachable" | +| Whether ANY capture site can fold. W4 exercised only the Qwen3 dense decode driver. The other hand-rolled drivers, and two models sharing the process-singleton registry, are untested | [#1226](https://github.com/mudler/vllm.cpp/issues/1226) | the same lease that answers the key question answers this one, and answering it before the key is settled measures the wrong thing | +| A supporting device leg on `orin:gpu0` | [#1226](https://github.com/mudler/vllm.cpp/issues/1226) | BLOCKED, cleanly: the Jetson `540.4.0` driver cannot run a CUDA 13 runtime (`cudaGetDeviceCount err=35`). A CUDA 12.x toolkit is the untried route. Not needed now that the dgx gate answered | ## Stop conditions diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index c9edb888f..4139eee47 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -272,7 +272,7 @@ Qwen3.6-27B NVFP4, GB10, whole serving window. | Peak `MemAvailable` drop | 68.35 GiB | 80.66 GiB | 1.180x | **PASS** | | Weight offload, resident device bytes (`ENG-WEIGHT-OFFLOAD` W6) | not measured | not measured | n/a | **BLOCKED**, not pending: unmeasurable on every host we own (GB10 shares one pool, so `cpu_offload_gb` frees nothing). Needs a discrete-GPU rig ([record](../.agents/benchmark-record.md)) | | Disk residency via `--offload-config` (`ENG-RESIDENCY-CONFIG`, [#1110](https://github.com/mudler/vllm.cpp/issues/1110)) | not measured | n/a (no disk tier upstream) | n/a | **PENDING** a GB10 run. The row changes no kernel, dtype or allocation, so it claims no throughput axis; the 370 GiB reproduction through the JSON form is owed ([spec](../.agents/specs/weight-residency-config.md)) | -| Decode-graph executables, `VT_CUDA_GRAPH_DEDUP` on vs off (`ENG-CUDAGRAPH-DEDUP`) | not measured | n/a | n/a | **PENDING**, not blocked: needs a leased CUDA box for the same-binary A/B (token-identity first, then the exec-count ratio) ([spec](../.agents/specs/eng-cudagraph-dedup.md)) | +| Decode-graph executables, `VT_CUDA_GRAPH_DEDUP` on vs off (`ENG-CUDAGRAPH-DEDUP`, [#1226](https://github.com/mudler/vllm.cpp/issues/1226)) | 3 execs / 3 captures; 2/2; 2/2 | n/a | 1.00x | **MEASURED, NEGATIVE**: GB10, 10/10 token-byte-identical, #1184 gone; N==M, the key carries the padded batch dim so buckets never fold. No speed/memory number ([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 b88656085..4a1a9e5d9 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -1705,7 +1705,7 @@ Gemma4/ROCm env split: public `VT_GEMMA4_EXPERT_VRAM_MB` caps expert LRU in posi `BACKEND-TENSTORRENT-HOST-FREE-FORWARD`: `ACTIVE`: env-gated `VT_TT_HOST_FREE_DECODE` decode-graph capture. Implementer P150 run of Qwen3-0.6B, 80 tokens: 79 replays, no hang, 5.8x vs eager, 22/22 vs the per-step-copy baseline. Default path inert. Operator gate and full-engine golden still owed. A new batch after the first capture is refused. -`ENG-CUDAGRAPH-DEDUP`: `ACTIVE`: env-gated `VT_CUDA_GRAPH_DEDUP` graph-executable dedup — one `cudaGraphExec` per captured TOPOLOGY instead of one per padded decode bucket per model. The shared `vt` registry hashes a captured graph's structure and re-points a single executable with `cudaGraphExecUpdate` on a hit, falling back to a private executable when the driver rejects the update. A memory and capture-time change, NOT a throughput change. Default OFF until a leased-GPU same-binary A/B proves a deduped replay token-identical; that A/B and the ROCm compile are owed ([#1162](https://github.com/mudler/vllm.cpp/issues/1162)). +`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 same-binary A/B ran on 2026-08-18 and split: replays are byte-identical 10/10 and [#1184](https://github.com/mudler/vllm.cpp/issues/1184) is gone, but the fold NEVER engages — the signature carries the padded batch dimension, so two decode buckets never group and ON allocates as many executables as OFF. Default stays OFF; the flip is unjustified on this evidence, and a coarser key is the open hypothesis ([#1226](https://github.com/mudler/vllm.cpp/issues/1226)). **Platform SELECTION is the one non-additive site, and is now gated.** A platform missing from `CurrentPlatform()`'s hardcoded walk registers and answers