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