From 0c3788b9cff4ac236a76c776058f270e7e525962 Mon Sep 17 00:00:00 2001 From: Luca Barbato Date: Thu, 20 Aug 2026 19:55:36 +0200 Subject: [PATCH 1/5] fix(BACKEND-TENSTORRENT-HOST-FREE-FORWARD): captured decode degenerate at the first KV block boundary (#1476) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The operator gate found captured replay deterministic-degenerate (word salad from ~generated token 30) while host-free eager stayed coherent, reproduced at the landed SHA in clean builds. Two root causes, both fixed: **The RAC page_table was `[C,1]`.** The tt-metal dataflow reader resolves `page_table_ptr[update_idx / block_size]` (`reader_update_cache_interleaved_start_id.cpp:110-111`) — it walks the STICK — so the moment `cur_pos` crossed `block_size` (32) the kernel indexed past the one-column tensor and wrote KV into a garbage physical block. The device tensor now carries the user's whole block-table row `[C, block_table_cols]`; ANY width change (`!=`, block-boundary growth or the multi-request shrink when the longest request finishes — the old `>` let the else-branch copy_to_device TT_FATAL on a shape mismatch) reallocates and RETIRES (keeps alive) the superseded tensor, because a freed device buffer can hand its address to a new allocation while a recorded trace still addresses it; steady state refreshes content only on change (`ptv != e.pt_host`), so there are zero copies inside a block. This lands the "Phase 2 full" refresh the old comment owed and never implemented. **`WarmDecodePos` keyed its skip on `GraphCapturesDone()`.** That counter is process-global and `Reset()` never clears it, so the cold eager step after a boundary reset ran no plus_one and the RE-captured trace read `cur_pos` one position behind. The regime flag now comes from the driver (`s.graph.captured()`); every cold/warm/capture step re-seeds `cur_pos = seq_lens - 1`, and only true replays leave it to the captured plus_one. | Check | Result | |---|---| | captured vs host-free eager, 80 tokens | **NOT byte-identical** — 284B md5 `3b5a579d…` vs 286B md5 `f5ffdf6a…`, first diff byte 174, both arms coherent. Adjudicated per-step top-2 (`VT_TT_DUMP_KV`): 45/80 steps argmax-identical with top-2 values agreeing to ≤0.5 logits; the FIRST divergence (step 46) is a swapped top-2 near-tie — gaps 0.25 vs 0.125 nats (1-2 bf16 ULP), the band this repo tracks for Qwen3-0.6B on TT (#1488 owes the teacher-forced re-adjudication); the 34 later differences are prefix divergence, not numeric evidence | | captured, 160 tokens | coherent, the 80-token answer a strict byte-prefix, 5 boundaries crossed, exit 0 | | captured + `VT_TT_RECAPTURE_EVERY=8` | 9 captures / 71 replays (8 mid-generation re-captures): byte-identical to the plain captured arm, same single step-46 near-tie vs eager. Without this arm the fixed-width gate never fires a `Reset()` (engine preallocates bt_cols=256), so the re-seed guarantee was previously undetected | | mutation reds (/tmp clone, same build config) | G2 (`GraphCapturesDone()>0` early-return restored): steps 1-10 agree, RED at step 11 — non-tie divergence (argmax 220 gap 1.125 vs eager 13), text degrades. G1 (page_table `[C,1]`, no steady refresh): 32 steps agree, RED at step 33 (first step past block_size) with the word salad. Both restores sha256-verified, rebuilt, rerun green with answers byte-identical to the gate runs | | `test_tenstorrent_backend` | 23/23 cases, 831/831 assertions, with AND without an ambient `VT_TT_HOST_FREE_DECODE` | | `test_qwen3_paged_engine` | unchanged from pre-fix (default path; stale golden #1488) | An earlier draft of this commit claimed captured==eager byte-identical over 80 tokens; that was a `grep -m1 '^ Answer'` first-line artifact and is corrected here in the message, the spec, STATUS, and the #1476 index row. - The RAC width-realloc condition is `!=` (shrink included), matching the driver's `cols_changed !=` reset. - `CopyDeviceDeviceIfCapture` / `MemsetDeviceIfCapture` read `VT_TT_HOST_FREE_DECODE` live instead of caching it in a static: the inertness-guard case unsets the env mid-process and must observe the decline, and a suite run under an ambient flag must not pin armed behavior for cases that unset it. The kRopeNeox bit-exact case now unsets the flag itself (mirroring the inertness case), so the suite is green under an ambient flag (was 21/23). - `WarmRacIdx` builds its host page-table tensor only on the copy path and the never-read `RacIdxEntry::idx_host` field/write is gone. - `[TT-DUMP-LOGITS]` carries a top-2 (id, value, gap) suffix under `VT_TT_DUMP_KV` — the near-tie adjudication instrument. - `test_tenstorrent_backend` exits 139 AFTER a green doctest summary (static `optional` destruction after device close). Proven pre-existing by A/B stash/build/run at `origin/main`: #1486. - `test_release_metadata` reds on every aarch64 host (fixture stages the host `/bin/true` into an x86_64-named archive), so preflight cannot go green on the TT dev fleet: #1487. It is the one gate that stays red in this change's preflight. - The TT `test_qwen3_paged_engine` golden is stale (anchor drift prompt[1] tok=10, identical before and after this fix — that test runs the default path): #1488 owes the `VT_DUMP_IDS` re-adjudication. Records: spec `## Owed`/`## Now` (including the `retired_pts` defense-in-depth reachability caveat), the STATUS cell, and the corrected #1476 issue-index row. Closes #1476. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:zai-glm-5.3 [Maki] --- .agents/benchmark-record.md | 31 +++++ .agents/issue-index.md | 5 + .../specs/tenstorrent-host-free-forward.md | 123 ++++++++++++++++-- docs/BENCHMARKS.md | 2 +- docs/STATUS.md | 2 +- src/vllm/model_executor/models/qwen3.cpp | 66 ++++++++-- src/vt/tenstorrent/tenstorrent_device.h | 19 ++- src/vt/tenstorrent/tenstorrent_ops.cpp | 109 ++++++++++------ tests/vt/test_tenstorrent_backend.cpp | 6 + 9 files changed, 291 insertions(+), 72 deletions(-) diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md index 1e7f4eae0..cec16bf2d 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -25565,3 +25565,34 @@ decision. And still owed, now that the alias is excluded on the target silicon: WHAT the step-7 divergence is. Excluding one cause is not identifying another, and the next traceable hypothesis is a per-operation two-arm comparison of the step-7 forward, first differing tensor named. + +--- + +## TT host-free decode: the 22/22 argmax is withdrawn; the operator gate records a step-46 near-tie (#1476, #1488) (2026-08-20) + +Row `BACKEND-TENSTORRENT-HOST-FREE-FORWARD`, evidence moved out of its +`docs/BENCHMARKS.md` cell by the 220-char entry budget. + +The 2026-08-16 implementer figure "22/22 argmax vs the per-step-copy baseline" +predates the final on-device `cur_pos` plus_one integration and does not +reproduce on the landed tree. The operator gate at `206afb63` (2026-08-20) found +captured replay deterministic-degenerate (word salad from ~generated token 30 = +the first KV block boundary) while host-free eager stayed coherent; issue #1476 +carries the full derivation. + +After the fix (operator rerun at `2b06f98a`, identical bytes to the implementer +and re-review runs): captured `Hello`/80-tok answer 284 bytes md5 +`3b5a579d82d58396fe4e344826946403`, eager 286 bytes md5 +`f5ffdf6aa290e11fd187673c2f3c52bb`, first divergence at byte 174 = decode step +46, a swapped top-2 near-tie (captured gap 0.25 nats, eager 0.125 nats; the +0.5-nat bar of `scripts/qwen3-neartie-gap.py`). 45/80 steps argmax-identical; +the 34 later differences are prefix divergence. `VT_TT_RECAPTURE_EVERY=8` (9 +captures / mid-generation re-captures) is byte-identical to the plain captured +arm. The #1476 degeneration is gone — G1 (`[C,1]` page_table) and G4 +(steady-state refresh suppressed) mutations each regenerate the exact word +salad; G2 (regime early-return) reds at step 11 under the recapture arm. + +Whether the step-46 flip is teacher-forced-benign stays UNADJUDICATED on the +full engine: that is #1488's `VT_DUMP_IDS` re-adjudication, the same band as +the paged-engine anchor drift it owns. No speed number is quoted from any of +these runs; a same-binary A/B on an idle host precedes any future figure. diff --git a/.agents/issue-index.md b/.agents/issue-index.md index 1904eab8b..bab9e421e 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -505,3 +505,8 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1493](https://github.com/mudler/vllm.cpp/issues/1493) | `LTX25-DFR-ROUNDS` | **DFR's unclamped `2**round_idx` tile count is mirrored and gated by nothing, because every fixture canvas in this tree has ONE keyframe segment.** `tile_ranges` clamps to `min(num_tiles, n_segments)` (`dfr_layout.py:171`), and the 9-frame fixture pads to a 25-frame canvas with a single segment, so round 1 asks for 2 windows and gets 1 and round 2 asks for 4 and gets 2. A port computing `round_idx + 1`, or `2 * round_idx`, or capping at 2 returns the SAME tile counts on every test here, with every downstream shape, frame count and exit status identical. NOT FIXED IN FLOW and the judgement is recorded rather than reversed: reaching 4 segments needs a materially longer canvas, so round 2 would denoise 4 tiles on a canvas already doubled twice - a new fixture and a substantially longer CPU run in a file that already carries 102 cases, not an assertion added to the existing render. The bound is stated in the test body and in `docs/USAGE.md`, so it was disclosed before it was owned; this row is the ownership. Closing it needs one render whose canvas carries at least 4 segments plus an assertion that `round_tile_counts` reads the unclamped `2**round_idx` for at least one round, which is the only shape that separates `2**round` from every expression agreeing with it at 1 and 2. Listed under `## Owed` in [`ltx25-dfr-rounds.md`](specs/ltx25-dfr-rounds.md) | bug | | [#1494](https://github.com/mudler/vllm.cpp/issues/1494) | `LTX25-DEVICE-RESIDENCY` | **A SECOND LTX-2.5 phase-coverage ratio decides by box load, and it is NOT the one [#1439](https://github.com/mudler/vllm.cpp/issues/1439) tracks.** `ltx2 video: the three carrying phases contain their work and the load keeps its order` asserts `CHECK_MESSAGE(covered >= c.min_coverage * leaf_seconds, ...)` at `tests/vllm/multimodal/test_ltx2_video.cpp:3696`; #1439 is `CHECK(leaves >= 0.95 * wall)` in a DIFFERENT case, now at `:3259`. Closing one does not close the other. Measured 2026-08-20, x86_64 `Release` `VLLM_CPP_CUDA=OFF`, three consecutive full-suite runs of ONE binary (`sha256 8fdbc31d...`) with no source change: loadavg 10.45 gives 94.6039% RED, a quieter run gives 96.8506% green, loadavg 16.53 gives 94.6039% RED. The comparison prints as `CHECK( 0.00414483 >= 0.00416218 )` - `denoise` is 0.00438124 s, its eight named sub-scopes cover 0.00414483 s, so the un-named residue is 0.00023641 s and **the margin is a quarter of a millisecond**. Same scheduling polarity #1439 recorded, which is the tell that this is the instrument and not the code: the run that PASSED is the run where `denoise` took 0.00940481 s, more than twice the failing runs', because the residue grows more slowly than the leaf it is divided by. NOT FIXED IN FLOW: bounding the residue in SECONDS beside the ratio, so the assertion says the same thing at fixture and production scale, changes a gate's semantics and needs its own row, spec and red-first evidence per `AGENTS.md` `## Changing the rules or a checker` - the same conclusion #1439 reached, and the two should be repaired together because one seconds bound would serve both. Found by the fresh implementer repairing the review findings of [#1481](https://github.com/mudler/vllm.cpp/pull/1481); pre-existing and not that PR's defect | bug | | [#1474](https://github.com/mudler/vllm.cpp/issues/1474) | `VT-CONV1D-F32-ACC` | `vt::Conv1d` and `vt::ConvTranspose1d` accumulate every output cell in **f64** against a reference that accumulates **f32**, and the provenance recorded for that width is wrong in two independent ways. **vLLM owns no such op** at the parity pin `555967922` — no `Conv1dLayer`, no `ConvTranspose*Layer`, and the only `ConvTranspose` string in the tree names `torch.nn` classes inside an NVTX annotation hook (`vllm/utils/nvtx_pytorch_hooks.py:47-49`); vLLM deliberately drops the vocoder it would otherwise own (`qwen3_omni_moe_thinker.py:1975`, `skip_prefixes=["talker.", "code2wav."]`), and where it DOES own a convolution it states this polarity itself (`csrc/cpu/mamba_kernels.hpp` — `// Accumulate in float32 for precision`). So the reference is torch, and torch was MEASURED rather than read: a 27-tap `[+1e8, 0.1 x 25, -1e8]` probe over a uniform input, where an f32 accumulator lands on exactly `0.0` in ANY order and an f64 one on ~2.5, returns **0.0** for `F.conv1d` at f32 AND at bf16, **0.0** for `F.conv_transpose1d` at f32, and 2.4999998510 at f64 (torch 2.11.0+cu130). **The recorded justification fails twice.** `.agents/specs/minimax-music3.md` §13.2 and §18.3 say f64 "is what every committed golden for all four consumers was taken with", but all three generators run torch in f32 — `gen-bigvgan-goldens.py:48` builds f64 then `.float()`, `gen-ltx2-vae-goldens.py:223,234` and `gen-minimax-music3-acoustic-goldens.py:81,134` cast every parameter and input with `astype(np.float32)` — so the goldens are the output of an **f32-accumulating** reference and the f64 arm has been wider than its own oracle since the op landed. And `include/vt/ops.h` clause (1) cites "every committed golden under `tests/parity/goldens/`", a directory whose 101 entries contain **no** vocoder, BigVGAN, LTX-2.5 VAE, FVQ or general-conv1d golden at all; those live in `tests/vllm/models/*_goldens.inc`. An uncheckable citation is how the first claim survived. **The precedent is landed inside one of the same models**: [#1008](https://github.com/mudler/vllm.cpp/issues/1008) narrowed LTX-2.5's conv VIDEO VAE f64->f32 on this exact argument (`d1b0ea3a8`, [PR #1036](https://github.com/mudler/vllm.cpp/pull/1036)), leaving LTX-2.5 shipping f32 video decode and f64 audio decode in ONE model. The measured risk is the summation ORDER, not the width: in #1008 naive-serial f32 pushed one arm to 5.00679e-06 against a 5e-06 tolerance, repaired by per-input-channel BLOCKED partials — which is what torch's blocked-GEMM convolution does and what `vt::Conv3d` already declares as its contract. Spec [`vt-conv1d-f32-accumulator.md`](specs/vt-conv1d-f32-accumulator.md); corrects both false records in the same flow | bug | +| [#1476](https://github.com/mudler/vllm.cpp/issues/1476) | `BACKEND-TENSTORRENT-HOST-FREE-FORWARD` | **TT host-free captured decode went degenerate at the first KV block boundary while host-free eager stayed coherent.** Found by the operator gate at `206afb63`, reproduced at the landed SHA `79ff8f31` in fresh clean builds — not a main regression; the recorded 22/22 argmax predated the final on-device `cur_pos` integration and did not reproduce on the landed tree. TWO root causes, both read at the tt-metal source and fixed together: (1) the RAC `page_table` was allocated `[C,1]`, but `reader_update_cache_interleaved_start_id.cpp:110-111` resolves `page_table_ptr[update_idx / block_size]` — the STICK — so the first write past `cur_pos >= block_size` (32) indexed past the one-column tensor and landed KV in a garbage physical block; the device tensor now carries the user's whole block-table row `[C, block_table_cols]`, grows by reallocation on width change, and RETIRES (keeps alive) the superseded tensor because a freed buffer can hand its address to a new allocation while a recorded trace still addresses it. (2) `WarmDecodePos` skipped its re-seed whenever the process-global `GraphCapturesDone()` was nonzero — a counter `Reset()` never clears — so the cold eager step after a boundary reset ran no plus_one and the RE-captured trace read `cur_pos` one position behind; the regime flag now comes from the driver (`s.graph.captured()`), and every cold/warm/capture step re-seeds. Implementer verification on the P150 (full-answer compares, never a first-line `grep -m1`): the degeneration is GONE — a 160-token captured run is coherent with the 80-token answer as a strict byte-prefix across 5 block boundaries, and reverting either root cause in a scratch build regenerates the word salad (the review's G1/G4 mutations). Captured vs host-free eager is NOT byte-identical — the earlier identical claim was a first-line compare artifact: full answers 284B md5 `3b5a579d82d58396fe4e344826946403` vs 286B md5 `f5ffdf6aa290e11fd187673c2f3c52bb`, first diff at byte 174, both arms coherent. Per-step top-2 adjudication (`VT_TT_DUMP_KV`, the `qwen3-neartie-gap.py` 0.5-nat bar): argmax identical 45/80 steps with top-2 values agreeing to ≤0.5 logits (≤4 bf16 ULP); the first divergence, decode step 46, is a swapped top-2 near-tie — captured `[11:19.75, 311:19.50]` gap 0.25 nats vs eager `[311:19.625, 11:19.50]` gap 0.125 nats, cross-arm deltas 0.125/0.25 logits — inside the near-tie band this repo already tracks for Qwen3-0.6B on TT (#1488 owes the teacher-forced golden re-adjudication); the 34 later argmax differences are prefix divergence, not numeric evidence. The re-seed regime is gate-covered by a new arm: `VT_TT_RECAPTURE_EVERY=8` (9 captures, 71 replays) is byte-identical to the plain captured arm and restoring the old `GraphCapturesDone()>0` early-return reds it; `test_tenstorrent_backend` 23/23 green with the fix, with and without an ambient flag (exit-time segfault pre-existing, [#1486](https://github.com/mudler/vllm.cpp/issues/1486)). Steady state copies the page_table only on content change — zero copies inside a block — so the fix lands the "Phase 2 full" refresh the old comment owed. Spec [`tenstorrent-host-free-forward.md`](specs/tenstorrent-host-free-forward.md) | bug | +| [#1486](https://github.com/mudler/vllm.cpp/issues/1486) | `BACKEND-TENSTORRENT-HOST-FREE-FORWARD` | **`test_tenstorrent_backend` SIGSEGVs after a fully green doctest summary (23/23 cases, 831/831 assertions): exit 139 from static destruction.** The chain is `std::optional` destruction into `ttnn::Tensor::deallocate_impl` into `tt::tt_metal::GraphTracker::is_enabled` on a tracker already torn down — the UMD close lines print first. The holders are process-lifetime cache-entry fields at `src/vt/tenstorrent/tenstorrent_ops.cpp:211,396,707,1020,1206`. PROVEN PRE-EXISTING by A/B in one build directory during the #1476 gate: stash the 3-file fix, rebuild, run → 23/23 + exit 139; restore, rebuild, run → identical. Deterministic per run on thalia (aarch64 P150 host). Impact: ctest marks the suite failed while every test passed — the same signature #1476's gate noted once on `test_qwen3_paged_engine`. Repair direction: drain the static caches before device teardown rather than rely on exit order. Listed under `## Owed` in [`tenstorrent-host-free-forward.md`](specs/tenstorrent-host-free-forward.md) | bug | +| [#1487](https://github.com/mudler/vllm.cpp/issues/1487) | `BACKEND-TENSTORRENT-HOST-FREE-FORWARD` | **`test_release_metadata` is red on every aarch64 host, so `agent-preflight` cannot go green on the TT dev fleet.** `tests/scripts/test_release_metadata.py:55` stages the HOST arch — `shutil.copy2("/bin/true", ...)` — into the hardcoded `vllm.cpp-0.0.1-linux-x86_64-glibc-cpu.tar.gz` (`:143`); on x86_64 the ELF matches the manifest by accident, on aarch64 the validator correctly refuses: `ELF host architecture does not match manifest: ARM aarch64`. Measured at `206afb63` == `origin/main` with zero local commits and an uncommitted src-only diff; NOT the #1353 disk-full mode (902G free, no scratch-repo failures, the message is the mismatch itself). Repair direction: make the fixture self-consistent (stage an ELF matching the declared arch, or derive the declared arch from the staged ELF) so the validator is exercised for a manifest that lies. Listed under `## Owed` in [`tenstorrent-host-free-forward.md`](specs/tenstorrent-host-free-forward.md) | bug | +| [#1488](https://github.com/mudler/vllm.cpp/issues/1488) | `BACKEND-TENSTORRENT-HOST-FREE-FORWARD` | **The TT `test_qwen3_paged_engine` golden is stale: `REQUIRE(anchor_ok)` fails with anchor drift prompt[1] tok=10 (engine=14126, committed anchor=62901).** Measured IDENTICALLY before and after the #1476 fix (same engine token both arms, ctest exit 8 both arms) — that test runs the DEFAULT path with no `VT_TT_HOST_FREE_DECODE`, which the fix does not touch, so the drift is not the captured-decode defect. The golden was committed `971d5506` (2026-08-09) and 13 TT ops commits landed Aug 9-18 with no on-card golden re-run. Owed: the `VT_DUMP_IDS` + `qwen3-neartie-gap.py` re-adjudication, then either a re-derived golden or an escalation to a real divergence. Split out of #1476 so the fidelity fix can close without orphaning the finding. Listed under `## Owed` in [`tenstorrent-host-free-forward.md`](specs/tenstorrent-host-free-forward.md) | bug | +| [#1504](https://github.com/mudler/vllm.cpp/issues/1504) | `SPEC-DFLASH2` | **`test_audit_live_rows` is red on `main`: `KERNEL-DFLASH2-SELECTOR-EDGES` and `KERNEL-TOPK-PAIRS` remain `ACTIVE` after their implementation landed.** Measured at `origin/main` `cffe59b0` in a clean detached worktree: `scripts/audit-live-rows.py` reports `248 live rows; 2 abandoned ACTIVE` and `test_audit_live_rows` fails `test_shipped_record_has_no_abandoned_active_row` with exactly those two IDs. Both kernel-matrix rows (.agents/kernel-matrix.md:144-145) record their implementation as landed and gated on 2026-08-20 via SPEC-DFLASH2 W3 (PR #1497) yet still carry `ACTIVE`, so every preflight on current main inherits the red. Fourth instance of the #1469 class. NOT FIXED IN FLOW: the repair is a lifecycle move on `SPEC-DFLASH2`-owned rows and owes `docs/STATUS.md`, `docs/BENCHMARKS.md` and the moved spec's `## Now`; `KERNEL-TOPK-PAIRS` also carries the open CUDA NaN-tie disagreement #1489 that the move should reconcile or name as its open gate. Found while rebasing `BACKEND-TENSTORRENT-HOST-FREE-FORWARD` onto main; filed in flow | bug | diff --git a/.agents/specs/tenstorrent-host-free-forward.md b/.agents/specs/tenstorrent-host-free-forward.md index 1cbc0f59e..36abbb8c8 100644 --- a/.agents/specs/tenstorrent-host-free-forward.md +++ b/.agents/specs/tenstorrent-host-free-forward.md @@ -259,14 +259,22 @@ investigation row but MUST be addressed by the item-5 port: That is silently wrong rather than a refusal, and it is a candidate explanation for `test_qwen3_paged_engine` still timing out under the flag. Owned by [#1105](https://github.com/mudler/vllm.cpp/issues/1105). -- **`VT_TT_RECAPTURE_EVERY` lags `cur_pos` by one per recapture cycle.** - `GraphCapturesCounter` is only ever `fetch_add`ed (`tenstorrent_ops.cpp:3225`, called - once from `:3297`); nothing resets it, and `DestroyGraph` does not. So the - recapture-triggered eager step re-seeds nothing and skips every `copy_to_device`, - and the following capture step captures a `cur_pos` one behind. The eager PA - consistency check at `:2410-2412` cannot see it: it compares `e.cp_host[0]` against - the same `seq_lens` that rewrote `e.cp_host` at `:3982`, so it validates the host - mirror against itself. Owned by [#1105](https://github.com/mudler/vllm.cpp/issues/1105). +- **`VT_TT_RECAPTURE_EVERY` lagged `cur_pos` by one per recapture cycle; the + mechanism is fixed by [#1476](https://github.com/mudler/vllm.cpp/issues/1476).** + The regime flag no longer comes from the process-global `GraphCapturesDone()` + (never cleared by `Reset()`); the driver passes `s.graph.captured()`, so the + cold eager step after ANY reset — a block boundary or a + `VT_TT_RECAPTURE_EVERY` cycle — re-seeds `cur_pos`, and the RAC page_table + refreshes on content change. The deeper defect stays open on + [#1105](https://github.com/mudler/vllm.cpp/issues/1105): `DecodePosCache` is + keyed on bare `num_reqs` with no engine or device identity, and the real fix + is a per-cache-entry seed / generation field aliased on every warm call, not + any process-global counter. The eager PA consistency check remains + self-validating (`e.cp_host` against the `seq_lens` that wrote it). The + re-seed regime now has a gate arm: `VT_TT_RECAPTURE_EVERY=8` forces + mid-generation re-captures (9 captures over 80 tokens), byte-identical to + the plain captured arm, and the `GraphCapturesDone()`-early-return mutation + reds it (see `## Now`). The seven constraints above remain. A new batch size after the first capture is now refused (`VT_CHECK` in `WarmDecodePos` / `WarmPaMeta` / @@ -275,6 +283,37 @@ per-cache-entry seed / generation field and aliasing on every warm call, not a process-global `GraphCapturesCounter`. Tracked on [#1105](https://github.com/mudler/vllm.cpp/issues/1105). +- **The `retired_pts` keep-alive is unguarded defense-in-depth.** No gate + exercises it and none can on the current fixed-width setup: the engine + preallocates `block_table_num_cols=256`, so after the first allocation no + width change (growth or shrink) ever occurs and the retire branch in + `WarmRacIdx` is structurally unreachable there. A width change needs a + multi-request run whose longest request finishes (shrink) or a driver that + grows cols per block (growth) — neither is reachable through the current + decode-graph gate. It stays because freeing a device buffer a recorded + trace addresses is the worse failure; treat it as defense-in-depth until a + gate can reach it, and do not cite it as covered. The allocator itself now + handles ANY width change (`!=`, growth or the multi-request shrink when the + longest request finishes — the old `>` let the else-branch + `copy_to_device` TT_FATAL on a shape mismatch), mirroring the driver's + `cols_changed !=` reset. +- **`test_tenstorrent_backend` exits 139 after a fully green doctest summary** + (23/23 cases, 831/831 assertions): static `std::optional` cache + fields are destroyed after the UMD device closes and `deallocate_impl` reaches + a torn-down `GraphTracker`. Proven pre-existing at `origin/main` by an A/B + stash/build/run during the [#1476](https://github.com/mudler/vllm.cpp/issues/1476) + gate. Owned by + [#1486](https://github.com/mudler/vllm.cpp/issues/1486). +- **`test_release_metadata` is red on every aarch64 host**: the fixture stages + the host `/bin/true` into an `x86_64`-named archive, so `agent-preflight` + cannot go green on the TT dev fleet. Found while running this row's preflight. + Owned by [#1487](https://github.com/mudler/vllm.cpp/issues/1487). +- **The TT `test_qwen3_paged_engine` golden is stale**: anchor drift + prompt[1] tok=10 (engine=14126, committed=62901), identical before and after + the #1476 fix (that test runs the default path). Needs the `VT_DUMP_IDS` + + near-tie-gap re-adjudication. Owned by + [#1488](https://github.com/mudler/vllm.cpp/issues/1488). + The operator must still rerun the 80-token no-hang gate and `test_qwen3_paged_engine` on a Blackhole P150. An implementer run is an input, not a gate result. @@ -282,8 +321,66 @@ input, not a gate result. ## Now `ACTIVE`. R1-R3b and the R2 on-device `cur_pos` / `update_idxs` advance are -implemented on this branch, env-gated by `VT_TT_HOST_FREE_DECODE`. A P150 -run of Qwen3-0.6B "Hello" at 80 tokens completed 79 replays with no hang -and 22/22 argmax vs the per-step-copy baseline. Next: operator rerun of -that 80-token gate and `test_qwen3_paged_engine` on card. A new batch -size after the first capture now throws instead of emitting wrong tokens. +implemented on this branch, env-gated by `VT_TT_HOST_FREE_DECODE`. + +The operator gate (2026-08-20, P150, `206afb63`) found +[#1476](https://github.com/mudler/vllm.cpp/issues/1476): captured replay went +degenerate at the first KV block boundary while host-free eager stayed +coherent — the recorded 22/22 argmax predated the final `cur_pos` integration +and did not reproduce on the landed tree. Root causes, both fixed in this +change: the RAC `page_table` was `[C,1]` where the tt-metal reader indexes the +full stick (`page_table_ptr[update_idx / block_size]`), so every write past +block 1 landed in a garbage physical block; and `WarmDecodePos` keyed its +skip on the process-global `GraphCapturesDone()`, which `Reset()` never +clears, so a post-boundary re-capture read `cur_pos` one position behind. + +Implementer verification on the P150 (2026-08-20, this change, full-answer +compares — never a `grep -m1` first-line artifact): + +- **The #1476 degeneration is gone.** Reverting either root cause in a + /tmp scratch clone (same TT build config) regenerates it: the G1 mutation + (page_table back to `[C,1]`, steady-state refresh removed) keeps 32 steps + of argmax agreement then reds at step 33 — the first step past + `block_size=32` — with a non-tie divergence (argmax 1555 gap 0.75 vs + eager 13) and the word salad ("straight line line line on road…"); + the G2 mutation (`GraphCapturesDone() > 0` early-return) reds the + recapture arm at step 11, the first step of the second capture cycle. + Both restores are sha256-verified byte-for-byte, rebuilt, and rerun + green with answers byte-identical (284B md5 `3b5a579d…`) to the + worktree gate runs. A 160-token captured run is coherent and the + 80-token captured answer is a strict byte-prefix of it (5 block + boundaries crossed). +- **Captured vs host-free eager is NOT byte-identical** — the prior + byte-identical claim was a first-line compare artifact. Full answers: + captured 284B md5 `3b5a579d82d58396fe4e344826946403`, eager 286B md5 + `f5ffdf6aa290e11fd187673c2f3c52bb`, first diff at byte 174, both arms + coherent. Adjudicated per-step (`VT_TT_DUMP_KV` top-2 dump; the top-2 + raw-logit gap is the logprob gap in nats, the `qwen3-neartie-gap.py` + bar): argmax identical for 45/80 steps with top-2 values agreeing to + ≤0.5 logits (≤4 bf16 ULP at the ~20-logit scale); the FIRST divergence, + decode step 46, is a swapped top-2 near-tie — captured `[11:19.75, + 311:19.50]` gap 0.25 nats vs eager `[311:19.625, 11:19.50]` gap 0.125 + nats (1-2 bf16 ULP), cross-arm deltas at the tied pair 0.125/0.25 — + inside the near-tie band this repo already tracks for Qwen3-0.6B on TT + ([#1488](https://github.com/mudler/vllm.cpp/issues/1488) owes the + teacher-forced golden re-adjudication). The 34 argmax differences after + step 46 are prefix divergence (each arm greedy-decodes its own prefix), + not numeric evidence. +- **The `cur_pos` re-seed regime (G2) is gate-covered.** + `VT_TT_RECAPTURE_EVERY=8` forces 9 captures / 71 replays (8 + destroy+re-capture cycles mid-generation); that arm is byte-identical to + the plain captured arm (same 284B md5) and carries the same single + step-46 near-tie vs eager. Restoring the old + `GraphCapturesDone() > 0` early-return in a scratch build reds this arm + at step 11 (non-tie divergence, incoherent text); restoring the fix + greens it. Without the arm the fixed-width 80-token gate never fires a + `Reset()` (the engine preallocates bt_cols=256), so exactly 2 re-seeds + run and the guarantee was undetected. + +`test_tenstorrent_backend` 23/23 + 831/831 green with and without an +ambient `VT_TT_HOST_FREE_DECODE` (its exit-time segfault is pre-existing, +[#1486](https://github.com/mudler/vllm.cpp/issues/1486)). + +Next: operator rerun of the 80-token captured-vs-eager gate and +`test_qwen3_paged_engine` on card; the paged-engine golden re-adjudication is +[#1488](https://github.com/mudler/vllm.cpp/issues/1488). diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index 2fc95892d..ebdcacca6 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -547,7 +547,7 @@ built on it rather than keeping the flattering one. | Gemma4 prefill-peer Finish-barrier cost (#1047 item 3) | **Attribution GREEN at T=2029, not a product ship number.** A (wait-only BEFORE) 1122.10 tok/s vs B (AFTER) 1094.24; +2.55% / 46.05 ms/req. Wait deletion is not authorized to land. | T=19 not run. Overlap/async retirement unmeasured. Detail: [benchmark-record](../.agents/benchmark-record.md) | | Tenstorrent Blackhole (`BACKEND-TENSTORRENT`) | **NOT APPLICABLE (speed).** Correctness: OPT-125m STRICT 6/6 e2e on real hardware. Qwen3-0.6B has a device-specific golden and short 4-token warm smoke (~0.28 tok/s), not a completed speed run | Full 16x16 Qwen3 gate, then device-resident tensors + `ttnn::sdpa_decode` before any performance comparison. [Spec](../.agents/specs/tenstorrent-backend.md) | | Mistral-7B-v0.3 on Tenstorrent (`BACKEND-TENSTORRENT-MISTRAL`) | **PENDING (speed).** First data point: 4.26 tok/s warm, batch 1, 32 tok, single run on a P150. Not a gate, not reproduced. No vLLM ratio exists or can (no TT backend). Correctness 16/16 | Reproduce idle with a same-binary A/B before quoting. [Record](../.agents/benchmark-record.md), [spec](../.agents/specs/tenstorrent-mistral.md) | -| Host-free decode graph (`BACKEND-TENSTORRENT-HOST-FREE-FORWARD`) | **PENDING (operator gate).** Implementer P150 run (2026-08-16): Qwen3-0.6B, 80 tok, 5.8x vs eager, 79 replays, 22/22 vs copy baseline. Env-gated. No vLLM ratio | Operator rerun, then full-engine `test_qwen3_paged_engine`. [Spec](../.agents/specs/tenstorrent-host-free-forward.md) | +| Host-free decode graph (`BACKEND-TENSTORRENT-HOST-FREE-FORWARD`) | **PENDING (golden re-adjudication).** Operator gate 2026-08-20: the #1476 block-boundary degeneration is gone; captured vs eager carries one step-46 near-tie flip (0.25/0.125 nats). Env-gated. No vLLM ratio | [#1488](https://github.com/mudler/vllm.cpp/issues/1488) golden re-adjudication, then a same-binary A/B. [Record](../.agents/benchmark-record.md), [spec](../.agents/specs/tenstorrent-host-free-forward.md) | | Prompt logprobs (`SAMPLE-PROMPT-LOGPROBS`, #223) | **NO number measured, claimed or owed.** Correctness-only, CPU. Upstream ships this path explicitly unoptimized (`gpu_model_runner.py:5622-5623`); a step where no request asks is unchanged | Floor if one is ever wanted: vLLM's own `prompt_logprobs=k`, same model and prompt | | `logprobs_mode` (`SAMPLE-LOGPROB-TOKEN-IDS`, #238) | **NO number measured, claimed or owed.** Correctness-only, CPU. One [n, vocab] device->host copy per step when a processed mode is engaged, nothing when not | Nothing to close: observation modes, not a path vLLM optimizes either | | `logprob_token_ids` scoring (#264) | **No number owed:** correctness-only, CPU-gated; inert unless set | Owed once the OpenAI field is wired | diff --git a/docs/STATUS.md b/docs/STATUS.md index aaaf5c3bf..89150a0fc 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -69,7 +69,7 @@ feature once. The [Usage index](USAGE.md) links each runnable workflow. | Metal | Partial runtime support on Apple Silicon | Only the operations and models listed in Features are covered | | Vulkan | Partial runtime support | The documented OPT and Qwen gates define the proved scope | | ROCm | Build and focused community test evidence | Model and oracle runtime gates remain open | -| Tenstorrent Blackhole | Active, partial runtime support | The full rerun and performance path remain open | +| Tenstorrent Blackhole | Active, partial runtime support | Host-free decode fidelity rerun passed (#1476 fixed); the paged-engine golden re-adjudication (#1488) and the performance path remain open | | Intel XPU | Hardware-blocked | No accepted runtime gate | The eight v0.0.2 server bundles are published; Windows v0.0.3-pre.1 remains diff --git a/src/vllm/model_executor/models/qwen3.cpp b/src/vllm/model_executor/models/qwen3.cpp index d86c4fcf8..085258738 100644 --- a/src/vllm/model_executor/models/qwen3.cpp +++ b/src/vllm/model_executor/models/qwen3.cpp @@ -89,6 +89,29 @@ bool TtDumpKv(const Dev& d) { std::getenv("VT_TT_DUMP_KV") != nullptr; } +// Near-tie adjudication for the [TT-DUMP-LOGITS] prints: the top-2 RAW-logit +// gap IS the top-2 logprob gap in nats (softmax is monotone), so an argmax +// flip between two arms whose own top-2 gaps are both tiny is a bf16 near-tie +// resolution difference, not a forward divergence — the same bar +// scripts/qwen3-neartie-gap.py applies (gap <= ~0.5 nats = structurally +// correct). Prints the pair as `top2=[id1:v1 id2:v2] gap2=g`. +void TtDumpTop2(const float* v, int64_t n, char* out, size_t out_n) { + int64_t i1 = 0, i2 = -1; + float v1 = v[0], v2 = -1.0e30f; + for (int64_t i = 1; i < n; ++i) { + if (v[i] > v1) { + i2 = i1; v2 = v1; i1 = i; v1 = v[i]; + } else if (v[i] > v2) { + i2 = i; v2 = v[i]; + } + } + if (i2 < 0) i2 = 0; // n==1 degenerate + std::snprintf(out, out_n, "top2=[%lld:%.6f %lld:%.6f] gap2=%.6g", + (long long)i1, static_cast(v1), + (long long)i2, static_cast(v2), + static_cast(v1 - v2)); +} + // Dense SwiGLU MLP (qwen3.py::Qwen3MLP=Qwen2MLP): merged gate_up_proj -> // SiluAndMul -> down_proj. `dh2` is the post-norm hidden [T,H] bf16. // @@ -494,8 +517,10 @@ std::vector Qwen3DenseModel::Forward( for (int64_t i = 1; i < n_out * config.vocab_size; ++i) if (logits[static_cast(i)] > logits[static_cast(argmax)]) argmax = static_cast(i); - fprintf(stderr, "[TT-DUMP-LOGITS] Forward eager argmax=%d first5=[%f,%f,%f,%f,%f]\n", - argmax, logits[0], logits[1], logits[2], logits[3], logits[4]); + char top2[96]; + TtDumpTop2(logits.data(), n_out * config.vocab_size, top2, sizeof(top2)); + fprintf(stderr, "[TT-DUMP-LOGITS] Forward eager argmax=%d first5=[%f,%f,%f,%f,%f] %s\n", + argmax, logits[0], logits[1], logits[2], logits[3], logits[4], top2); } return logits; } @@ -516,8 +541,10 @@ ForwardLogits Qwen3DenseModel::ForwardDevice( for (size_t i = 1; i < logits_dump.size(); ++i) if (logits_dump[i] > logits_dump[static_cast(argmax)]) argmax = static_cast(i); - fprintf(stderr, "[TT-DUMP-LOGITS] ForwardDevice eager argmax=%d first5=[%f,%f,%f,%f,%f]\n", - argmax, logits_dump[0], logits_dump[1], logits_dump[2], logits_dump[3], logits_dump[4]); + char top2[96]; + TtDumpTop2(logits_dump.data(), static_cast(logits_dump.size()), top2, sizeof(top2)); + fprintf(stderr, "[TT-DUMP-LOGITS] ForwardDevice eager argmax=%d first5=[%f,%f,%f,%f,%f] %s\n", + argmax, logits_dump[0], logits_dump[1], logits_dump[2], logits_dump[3], logits_dump[4], top2); } return WrapDeviceLogits(d, std::move(dlogits), n_out, config.vocab_size); } @@ -722,8 +749,10 @@ ForwardLogits Qwen3DenseDecodeGraph::Step( for (int64_t i = 1; i < vocab; ++i) if (logits_dump[static_cast(i)] > logits_dump[static_cast(argmax)]) argmax = static_cast(i); - fprintf(stderr, "[TT-DUMP-LOGITS] eager path argmax=%d first5=[%f,%f,%f,%f,%f]\n", - argmax, logits_dump[0], logits_dump[1], logits_dump[2], logits_dump[3], logits_dump[4]); + char top2[96]; + TtDumpTop2(logits_dump.data(), vocab, top2, sizeof(top2)); + fprintf(stderr, "[TT-DUMP-LOGITS] eager path argmax=%d first5=[%f,%f,%f,%f,%f] %s\n", + argmax, logits_dump[0], logits_dump[1], logits_dump[2], logits_dump[3], logits_dump[4], top2); } return WrapDeviceLogits(d, std::move(lg), B, vocab); } @@ -790,9 +819,14 @@ ForwardLogits Qwen3DenseDecodeGraph::Step( // R2: seed the on-device-advanced cur_pos BEFORE WarmRacIdx, so the RAC // path can alias update_idxs to it (eliminating the per-replay // update_idxs copy_to_device — the toxic ~38-replay hang class). + // replay_regime = this slot's graph is captured (the warm hooks run + // BEFORE the boundary Reset below): replay steps leave cur_pos to the + // captured plus_one; cold/warm/capture steps re-seed it, which is what + // makes a post-boundary RE-capture read the right position (#1476). if (!pam.seq_lens.empty()) { vt::tenstorrent::WarmDecodePos( - pam.seq_lens.data(), static_cast(pam.num_reqs)); + pam.seq_lens.data(), static_cast(pam.num_reqs), + /*replay_regime=*/s.graph.captured()); } vt::tenstorrent::WarmRacIdx( pam.slot_mapping.data(), pam.slot_mapping.data(), @@ -881,8 +915,10 @@ ForwardLogits Qwen3DenseDecodeGraph::Step( for (int64_t i = 1; i < vocab; ++i) if (logits_dump[static_cast(i)] > logits_dump[static_cast(argmax)]) argmax = static_cast(i); - fprintf(stderr, "[TT-DUMP-LOGITS] replay step argmax=%d first5=[%f,%f,%f,%f,%f]\n", - argmax, logits_dump[0], logits_dump[1], logits_dump[2], logits_dump[3], logits_dump[4]); + char top2[96]; + TtDumpTop2(logits_dump.data(), vocab, top2, sizeof(top2)); + fprintf(stderr, "[TT-DUMP-LOGITS] replay step argmax=%d first5=[%f,%f,%f,%f,%f] %s\n", + argmax, logits_dump[0], logits_dump[1], logits_dump[2], logits_dump[3], logits_dump[4], top2); } return ViewDeviceLogits(s.logits->ptr(), d.q.device, B, vocab); } @@ -1005,8 +1041,10 @@ ForwardLogits Qwen3DenseDecodeGraph::Step( for (int64_t i = 1; i < vocab; ++i) if (logits_dump[static_cast(i)] > logits_dump[static_cast(argmax)]) argmax = static_cast(i); - fprintf(stderr, "[TT-DUMP-LOGITS] capture step argmax=%d first5=[%f,%f,%f,%f,%f]\n", - argmax, logits_dump[0], logits_dump[1], logits_dump[2], logits_dump[3], logits_dump[4]); + char top2[96]; + TtDumpTop2(logits_dump.data(), vocab, top2, sizeof(top2)); + fprintf(stderr, "[TT-DUMP-LOGITS] capture step argmax=%d first5=[%f,%f,%f,%f,%f] %s\n", + argmax, logits_dump[0], logits_dump[1], logits_dump[2], logits_dump[3], logits_dump[4], top2); } return ViewDeviceLogits(s.logits->ptr(), d.q.device, B, vocab); } @@ -1026,8 +1064,10 @@ ForwardLogits Qwen3DenseDecodeGraph::Step( for (int64_t i = 1; i < vocab; ++i) if (logits_dump[static_cast(i)] > logits_dump[static_cast(argmax)]) argmax = static_cast(i); - fprintf(stderr, "[TT-DUMP-LOGITS] cold step argmax=%d first5=[%f,%f,%f,%f,%f]\n", - argmax, logits_dump[0], logits_dump[1], logits_dump[2], logits_dump[3], logits_dump[4]); + char top2[96]; + TtDumpTop2(logits_dump.data(), vocab, top2, sizeof(top2)); + fprintf(stderr, "[TT-DUMP-LOGITS] cold step argmax=%d first5=[%f,%f,%f,%f,%f] %s\n", + argmax, logits_dump[0], logits_dump[1], logits_dump[2], logits_dump[3], logits_dump[4], top2); } s.warm = true; // lg is [S,vocab]; hand ownership out but expose only the first B (real) rows. diff --git a/src/vt/tenstorrent/tenstorrent_device.h b/src/vt/tenstorrent/tenstorrent_device.h index 298a83dd0..5e4cdf2b3 100644 --- a/src/vt/tenstorrent/tenstorrent_device.h +++ b/src/vt/tenstorrent/tenstorrent_device.h @@ -116,14 +116,21 @@ inline void WarmPaMeta(const int32_t*, int64_t, int64_t, int64_t, int64_t, #endif // R2 (on-device state advance): seed the persistent cur_pos device tensor -// (= seq_lens - 1) for this step. Called on the capture/warm step (re-seed), -// NOT every replay step — the captured plus_one advances it on-device. -// Also warms the plus_one program (program cache) so CaptureDecodePosAdvance -// can run inside the trace without a "load new binaries during capture" fatal. +// (= seq_lens - 1) for this step. Called on every cold/warm/capture step +// (re-seed), never on a replay step — the captured plus_one advances it +// on-device. Also warms the plus_one program (program cache) so +// CaptureDecodePosAdvance can run inside the trace without a "load new +// binaries during capture" fatal. #ifdef VLLM_CPP_TENSTORRENT -void WarmDecodePos(const int32_t* seq_lens, int64_t num_reqs); +// replay_regime=true: this step replays a captured trace whose plus_one owns +// the cur_pos advance — do not touch the tensor. replay_regime=false (cold, +// warm, or capture step): re-seed cur_pos = seq_lens-1 for THIS step. The +// re-seed is what makes a RE-capture correct (#1476): Reset() releases the +// trace but the cold eager step that follows runs no plus_one, so without a +// re-seed the newly captured trace reads cur_pos one position behind. +void WarmDecodePos(const int32_t* seq_lens, int64_t num_reqs, bool replay_regime); #else -inline void WarmDecodePos(const int32_t*, int64_t) {} +inline void WarmDecodePos(const int32_t*, int64_t, bool) {} #endif // R2: capture ttnn::plus_one(cur_pos) at the END of the trace body (after all diff --git a/src/vt/tenstorrent/tenstorrent_ops.cpp b/src/vt/tenstorrent/tenstorrent_ops.cpp index ad6cd0353..6f9f09741 100644 --- a/src/vt/tenstorrent/tenstorrent_ops.cpp +++ b/src/vt/tenstorrent/tenstorrent_ops.cpp @@ -1846,8 +1846,14 @@ std::map& DecodePosCache() { namespace { struct RacIdxEntry { ttnn::Tensor update_idxs; // int32 [C] device (persistent, content refreshed) - ttnn::Tensor page_table; // int32 [C,1] device (persistent, content refreshed) - std::vector idx_host; // content identity for reuse check + ttnn::Tensor page_table; // int32 [C,pt_width] device (persistent, content refreshed) + std::vector pt_host; // last page-table content copied to device + int64_t pt_width = 0; // columns of the allocated page_table + // Retired page-table tensors from width growth, kept ALIVE deliberately: + // a freed device buffer can hand its address to a new allocation while a + // (doomed, never-replayed-again) trace still records it. Bounded by the + // number of block boundaries crossed (~context/block_size). + std::vector retired_pts; // Persistent height-sharded RAC input: logical [1,1,nkv,d], padded // [1,1,nkv_pad,d] (shard [nkv_pad,d] on one core). The in-region RAC // ttnn::copy's the rope output into it; paged_update_cache reads only the @@ -3536,9 +3542,11 @@ bool CopyDeviceDeviceIfCapture(void* dst, const void* src) { // (which the decode-graph framework runs BEFORE capture) also exercises // ttnn::empty+ttnn::copy, compiling those programs into the cache so the // subsequent capture doesn't hit "Cannot load new binaries during trace - // capture." - static const bool host_free = - std::getenv("VT_TT_HOST_FREE_DECODE") != nullptr; + // capture." Read LIVE, not cached in a static: the inertness-guard case in + // test_tenstorrent_backend unsets the env mid-process and must observe the + // decline, and a suite run under an ambient flag must not pin the armed + // behavior for cases that unset it. + const bool host_free = std::getenv("VT_TT_HOST_FREE_DECODE") != nullptr; if (!tt_capture_active() && !host_free) return false; static bool once = [&] { // Enable program cache once on the first host-free path use — ttnn trace @@ -3583,8 +3591,8 @@ bool CopyDeviceDeviceIfCapture(void* dst, const void* src) { // Reinterprets the buffer as a 2D [rows, cols] f32 tensor matching the // existing device shadow's numel (zeros is the only value the forward uses). bool MemsetDeviceIfCapture(void* p, int value) { - static const bool host_free = - std::getenv("VT_TT_HOST_FREE_DECODE") != nullptr; + // Live read for the same reason as CopyDeviceDeviceIfCapture above. + const bool host_free = std::getenv("VT_TT_HOST_FREE_DECODE") != nullptr; if (!tt_capture_active() && !host_free) return false; if (value != 0) return false; // only zero-fill is handled on-device // Need an existing shadow to know shape/dtype; or allocate from the slot. @@ -3750,13 +3758,19 @@ void WarmRacIdx(const void* /*slot_mapping_owner*/, const int64_t* slots, // exactly that position for the PA to see the current token's KV. std::vector ptv; std::vector idxv; - ptv.reserve(static_cast(num_slots)); + // The kernel maps update_idx -> page_table_ptr[update_idx / block_size] + // (reader_{,paged_fused_}update_cache: virtual_block_id indexes the + // page-table STICK), so the device page_table must carry the user's WHOLE + // block-table row, not just the current virtual block (#1476: the old + // [C,1] tensor made every write past the first block land in a garbage + // physical block the moment cur_pos crossed block_size). + ptv.reserve(static_cast(num_slots * block_table_cols)); idxv.reserve(static_cast(num_slots)); for (int64_t t = 0; t < num_slots; ++t) { const int64_t slot = slots[t]; if (slot < 0 || seq_lens == nullptr || block_table == nullptr) { // Padding slot: paged_update_cache skips when update_idx == -1. - ptv.push_back(0); + for (int64_t c = 0; c < block_table_cols; ++c) ptv.push_back(0); idxv.push_back(-1); continue; } @@ -3766,14 +3780,18 @@ void WarmRacIdx(const void* /*slot_mapping_owner*/, const int64_t* slots, // which must equal the slot_mapping from the scheduler. const int32_t cur_pos = seq_lens[t] - 1; idxv.push_back(cur_pos); - // page_table = the physical block for cur_pos's virtual block. - const int32_t vblk = cur_pos / static_cast(block_size); - const int32_t pblk = block_table[t * block_table_cols + vblk]; - ptv.push_back(pblk); - if (std::getenv("VT_TT_TRACE_DEBUG") != nullptr) + // Full row: every virtual block the kernel may resolve this step and + // later in this width regime (steady state refreshes on content change). + for (int64_t c = 0; c < block_table_cols; ++c) { + ptv.push_back(block_table[t * block_table_cols + c]); + } + if (std::getenv("VT_TT_TRACE_DEBUG") != nullptr) { + const int32_t vblk = cur_pos / static_cast(block_size); + const int32_t pblk = block_table[t * block_table_cols + vblk]; std::fprintf(stderr, "[TT-TRACE] WarmRacIdx user=%lld slot=%lld cur_pos=%d vblk=%d pblk=%d bt_cols=%lld (expect slot=%d)\n", (long long)t, (long long)slot, cur_pos, vblk, pblk, (long long)block_table_cols, pblk * static_cast(block_size) + cur_pos % static_cast(block_size)); + } } const auto key = std::make_pair(num_slots, block_size); std::lock_guard g(RacIdxMutex()); @@ -3782,15 +3800,18 @@ void WarmRacIdx(const void* /*slot_mapping_owner*/, const int64_t* slots, // refreshed in place each step (copy_to_device, outside capture). The // captured paged_update_cache replays against the stable address and reads // the fresh values device-side. - ttnn::Tensor pt_host = ttnn::Tensor::from_vector( - ptv, SpecOf(tt::tt_metal::Shape({static_cast(num_slots), 1u}), - ttnn::DataType::INT32, ttnn::Layout::ROW_MAJOR)); - ttnn::Tensor idx_host = ttnn::Tensor::from_vector( - idxv, SpecOf(tt::tt_metal::Shape({static_cast(num_slots)}), - ttnn::DataType::INT32, ttnn::Layout::ROW_MAJOR)); - if (!e.allocated) { + if (!e.allocated || block_table_cols != e.pt_width) { + // ANY width change (block boundary growth, or the shrink when the longest + // request of a multi-request batch finishes and block_table_num_cols drops) + // reallocates: the else-branch copy_to_device would TT_FATAL on a shape + // mismatch, and the driver resets + re-captures on any column-count change + // (`cols_changed` compares with `!=`), so the new address is what the next + // capture records. The retired tensor stays alive (see the field) — never + // free a buffer a recorded trace addresses. + if (e.allocated) e.retired_pts.push_back(std::move(e.page_table)); e.page_table = ttnn::Tensor::from_vector( - ptv, SpecOf(tt::tt_metal::Shape({static_cast(num_slots), 1u}), + ptv, SpecOf(tt::tt_metal::Shape({static_cast(num_slots), + static_cast(block_table_cols)}), ttnn::DataType::INT32, ttnn::Layout::ROW_MAJOR), &device); // R2: alias update_idxs to the on-device-advanced cur_pos (DecodePosCache) @@ -3819,19 +3840,25 @@ void WarmRacIdx(const void* /*slot_mapping_owner*/, const int64_t* slots, &device); } e.allocated = true; + e.pt_width = block_table_cols; } else { - // R2 steady state: skip the per-step RAC page_table + update_idxs copies. - // update_idxs is aliased to the on-device-advanced cur_pos (plus_one). - // The RAC page_table (where paged_update_cache WRITES) only matters at - // block boundaries for decode T=1; for now skip in steady state (the PA - // page_table in WarmPaMeta handles the sdpa_decode read with on-change - // refresh). Phase 2 full: on-change refresh for RAC page_table too. - if (!r2_steady) { + // Steady state within one width: update_idxs is aliased to the + // on-device-advanced cur_pos (plus_one on replay steps, WarmDecodePos + // re-seed on cold/capture steps) — never copied here. The RAC page_table + // refreshes ONLY when its content changed (a new block was mapped): + // zero copies inside a block, so the toxic every-step-interleaved-write + // class stays out of the steady state; the copy that does fire rides the + // same step as the boundary re-capture (#1476 — the "Phase 2 full" + // refresh the old comment owed but never implemented). + if (ptv != e.pt_host) { + ttnn::Tensor pt_host = ttnn::Tensor::from_vector( + ptv, SpecOf(tt::tt_metal::Shape({static_cast(num_slots), + static_cast(block_table_cols)}), + ttnn::DataType::INT32, ttnn::Layout::ROW_MAJOR)); ttnn::copy_to_device(pt_host, e.page_table); - ttnn::copy_to_device(idx_host, e.update_idxs); } } - e.idx_host = idxv; + e.pt_host = ptv; // Build the persistent sharded RAC input ONCE from the first available // paged-KV shadow's geometry (same nkv/d as the cache): logical // [1,1,nkv,d], padded [1,1,nkv_pad,d], HEIGHT_SHARDED L1, shard @@ -3989,15 +4016,13 @@ void WarmPaMeta(const int32_t* block_table, int64_t num_reqs, int64_t max_blocks // R2: seed the persistent cur_pos device tensor (= seq_lens - 1) and warm the // plus_one program (program cache) so CaptureDecodePosAdvance can run inside // the trace. Called on the capture/warm step (re-seed), NOT every replay. -void WarmDecodePos(const int32_t* seq_lens, int64_t num_reqs) { +void WarmDecodePos(const int32_t* seq_lens, int64_t num_reqs, bool replay_regime) { if (std::getenv("VT_TT_HOST_FREE_DECODE") == nullptr) return; if (num_reqs < 1 || seq_lens == nullptr) return; - // R2: only seed/warm on the capture step (GraphCapturesDone()==0). On replay - // steps, cur_pos is advanced on-device by the captured plus_one — re-seeding - // here would overwrite the advance and break correctness. - // A new num_reqs after the first capture was never seeded: refuse rather - // than return and let WarmPaMeta/WarmRacIdx allocate a frozen standalone. - if (GraphCapturesDone() > 0) { + // Replay regime: cur_pos advances on-device by the captured plus_one — + // re-seeding here would overwrite the advance and break correctness. + // A new num_reqs that was never seeded is refused rather than left frozen. + if (replay_regime) { std::lock_guard g(DecodePosMutex()); auto it = DecodePosCache().find(num_reqs); VT_CHECK(it != DecodePosCache().end() && it->second.allocated, @@ -4006,6 +4031,14 @@ void WarmDecodePos(const int32_t* seq_lens, int64_t num_reqs) { "cache entry; recapture does NOT clear this (#1105)."); return; } + // Cold/warm/capture step: (re-)seed cur_pos = seq_lens - 1 for THIS step. + // The regime flag comes from the driver (graph captured?), NOT from + // GraphCapturesDone(): Reset() releases the trace without clearing that + // process-global counter, and the cold eager step that follows a Reset + // runs no plus_one — so after a re-capture the on-device cur_pos is one + // position behind unless it is re-seeded here (#1476). This also keeps the + // FIRST capture correct (its capture step re-seeds, which is why the bug + // only surfaced at the first block boundary, where Reset+re-capture runs). MeshDevice& device = SharedMeshDevice(); std::vector cpos(static_cast(num_reqs)); for (int64_t r = 0; r < num_reqs; ++r) diff --git a/tests/vt/test_tenstorrent_backend.cpp b/tests/vt/test_tenstorrent_backend.cpp index c34c63505..6f76356f4 100644 --- a/tests/vt/test_tenstorrent_backend.cpp +++ b/tests/vt/test_tenstorrent_backend.cpp @@ -584,6 +584,12 @@ TEST_CASE("kTENSTORRENT kRopeNeox is BIT-EXACT vs a host F32 reference (small)") MESSAGE("SKIPPED: no Tenstorrent device on this box"); return; } + // This case asserts the HOST apply path (bit-exact). An ambient + // VT_TT_HOST_FREE_DECODE (e.g. a suite run under the host-free gate) flips + // PreferDeviceRope to the device BF16 path even at small T*H and reds the + // bit-exact checks — so the case owns its own default-path env, mirroring + // the inertness-guard case below. + ::unsetenv("VT_TT_HOST_FREE_DECODE"); REQUIRE(vt::OpRegistered(vt::OpId::kRopeNeox, DeviceType::kTENSTORRENT)); constexpr int64_t T = 4, Hq = 2, Hk = 1, Dh = 8, Rot = 8; From 4d08778b7cd5c3de208aee4b6dfeea454f1fb945 Mon Sep 17 00:00:00 2001 From: Luca Barbato Date: Fri, 21 Aug 2026 00:09:11 +0200 Subject: [PATCH 2/5] fix(BACKEND-TENSTORRENT-HOST-FREE-FORWARD): refresh the stale TT Qwen3-0.6B paged-engine golden pair after oracle re-adjudication (#1488) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The committed Tenstorrent anchor for qwen3_greedy_0_6b encoded tokens the current default decode path no longer produces: REQUIRE(anchor_ok) red at prompt[1] tok=10 (engine 14126, committed 6290 — the logged 62901 is a print artifact, #1508). The golden was captured 2026-08-10 under the then-current default path, so the tokens drifted when that path changed, not because the engine diverged. Re-derivation (2026-08-20, Blackhole P150, fix-applied tree 0c3788b9c): VT_DUMP_IDS captured the full [16,16] battery twice, eager (VLLM_CPP_CUDAGRAPH=0) and captured, byte-identical dumps (md5 b5307e3396d2c5121946c0b746ba5acc) — the #1476 captured/eager fidelity holds across the whole battery, not only the focused pair. The refreshed pair differs from the old golden on 53 cells across 7 prompts: six prompts diverge at exactly one token (p1 tok10, p5 tok10, p10 tok5, p11 tok4, p12 tok13, p15 tok12) and their greedy continuations then follow the new prefix, and prompt 7 is rewritten from tok0 — the new p7 row matches the vLLM greedy sequence exactly, where the old anchor matched neither vLLM nor CUDA. scripts/qwen3-neartie-gap-transformers.py (transformers 4.57.1, torch 2.10.0a0+cpu) teacher-forced the oracle on the new sequence: max gap 375 mnats, zero cells above the 500-mnat band, zero outside-top-K — every refreshed token is a legitimate bf16 near-tie resolution, so the anchor is exact-deterministic and the gap golden is re-derived, not widened. The refreshed pair is derived from the fix-applied tree's TT output and this change therefore stacks on row/BACKEND-TENSTORRENT-HOST-FREE-F1476. The earlier before/after-#1476 comparison recorded in #1488 established only that the p1-tok-10 engine token matched in both arms (exit codes were compared, not the battery); pre-fix TT behavior at the other refreshed cells was never recorded, so a main-based branch could not be gated. Closes #1488. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:zai-glm-5.3 [Maki] --- .../neartie_gap_mnats_tenstorrent.npy | Bin 1152 -> 1152 bytes .../qwen3_greedy_0_6b/our_ids_tenstorrent.npy | Bin 1152 -> 1152 bytes 2 files changed, 0 insertions(+), 0 deletions(-) diff --git a/tests/parity/goldens/qwen3_greedy_0_6b/neartie_gap_mnats_tenstorrent.npy b/tests/parity/goldens/qwen3_greedy_0_6b/neartie_gap_mnats_tenstorrent.npy index d744aa4a182e2ae04af47eb15c111e7df3059489..0db07bdfa634484bd4532c804388d50013971357 100644 GIT binary patch delta 63 zcmZqRY~Y+QWpWOS$ixX8lN*>VAWRR&UmF|uGftc!0O3Q0c^H39Y`hFp9l$IA0P1iT AEdT%j delta 56 zcmZqRY~Y+QWpe`K8b;nPj0|wFXLBGUGvnj}7J-TT6((-50I?>=GxkhQV3q)45dg2- B4{ZPd diff --git a/tests/parity/goldens/qwen3_greedy_0_6b/our_ids_tenstorrent.npy b/tests/parity/goldens/qwen3_greedy_0_6b/our_ids_tenstorrent.npy index 20301f04497be8014173cf5e4e48da887fa31516..e94fb91704e9b6a84e9faedde64b6893919a1138 100644 GIT binary patch delta 225 zcmZqRY~Y;mf=SPO;;VmR){G1c;y`?yn}NXqNEh=kGE8Ql{EyMe4kYZ$%fJx9!oc7R zWUmF{G9U(tHwh3DpIpZ@&F?lB149rK1H&O9h}O?a3=En;KFEkDb_NEJ?TYOT3;{r~ z`4S8aKu3VB0?Drd>g5HB?=%I{Kye4A$$rdolN*>#xDd)F?_=I1U@phNunnl=z9a*~ IWDy1i03awK`~Uy| delta 249 zcmZqRY~Y;mf@zY(#8>~s3>X<07Qbg;a0b$`$&3t>*%=rfvP}NR=w!#uz`$S*6mACM zJ3uVSgy4hZ%79`VKs-Z-fkBOjfx#Ne&-Mix$ve4@X)Z78e1C=xIL$NY5P(2Gs j{p31kGcI1BumjKvHIB*qm^TURi-(x94QSeZNd^W0-JB@Q From dc877df5dae12825e4a68c4d06588042efe33920 Mon Sep 17 00:00:00 2001 From: Luca Barbato Date: Thu, 20 Aug 2026 23:59:16 +0200 Subject: [PATCH 3/5] fix(BACKEND-TENSTORRENT-HOST-FREE-FORWARD): print truthful anchor-drift values in test_qwen3_paged_engine (#1508) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The pinned doctest MessageBuilder stream has no const char* overload, so every separately-bound const char* streamed through MESSAGE or REQUIRE_MESSAGE falls back to the implicit pointer-to-bool conversion and renders as 1. In this test that garbled the anchor-drift failure into "committed anchor=62901" for a golden holding 6290 (and 96251 for 9625): the trailing device-golden ternary printed its bool as a digit glued to the last numeric field, and label ("qwen3-0.6B") printed as 1 in every message. During #1488 the garbled value was misread as corruption of the loaded golden buffers until a byte-level npy verification cleared them; a 7-line harness against the pinned header reproduces the artifact. label, ids_name and gap_name are now std::string and the message ternaries are wrapped, so the drift line prints the model, the real committed anchor, and the device-appropriate remediation arm. Verified on card: the failing case now logs "qwen3-0.6B anchor drift prompt[0] tok=5 engine=15344 committed anchor=9625 — re-run qwen3-neartie-gap.py to refresh the gap golden" (that arm was itself invisible before — the two arms differed only by the same swallowed string). Closes #1508. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:zai-glm-5.3 [Maki] --- tests/parity/test_qwen3_paged_engine.cpp | 17 +++++++++++------ 1 file changed, 11 insertions(+), 6 deletions(-) diff --git a/tests/parity/test_qwen3_paged_engine.cpp b/tests/parity/test_qwen3_paged_engine.cpp index 3d11cd1e2..3df17d654 100644 --- a/tests/parity/test_qwen3_paged_engine.cpp +++ b/tests/parity/test_qwen3_paged_engine.cpp @@ -142,7 +142,7 @@ const int32_t* AsI32(const parity::NpyArray& a) { // vLLM's argmax in vLLM's own logits (strict where our token IS vLLM's argmax). // Reports the strict token-exact count and the max gap. void RunGate(const std::string& repo_dir, const std::string& golden_subdir, - const char* label) { + const std::string& label) { const std::string snap = FindSnapshot(repo_dir); if (snap.empty()) { MESSAGE(label << " checkpoint absent; skipping (dgx-only) — " << repo_dir @@ -262,10 +262,13 @@ void RunGate(const std::string& repo_dir, const std::string& golden_subdir, const int32_t* anchor_ids = od; // hard anchor for THIS device const int32_t* gap_ids = gapd; // vLLM teacher-forced gaps for THIS device parity::NpyArray o_dev, gap_dev; // keep device arrays alive for the loop - const char* ids_name = + // std::string, not const char*: doctest's MessageBuilder stream has no + // const char* overload and would render a bare pointer as bool "1" in the + // REQUIRE_MESSAGE below (#1508). + const std::string ids_name = metal ? "our_ids_metal.npy" : (rocm ? "our_ids_rocm.npy" : "our_ids_tenstorrent.npy"); - const char* gap_name = + const std::string gap_name = metal ? "neartie_gap_mnats_metal.npy" : (rocm ? "neartie_gap_mnats_rocm.npy" : "neartie_gap_mnats_tenstorrent.npy"); @@ -277,7 +280,8 @@ void RunGate(const std::string& repo_dir, const std::string& golden_subdir, // qwen3-neartie-gap.py then teacher-forces vLLM on that sequence. bootstrap_only = true; MESSAGE(label << ": BOOTSTRAP dump (device golden absent) for " - << (metal ? "Metal" : (rocm ? "ROCm" : "Tenstorrent")) << "..."); + << std::string(metal ? "Metal" : (rocm ? "ROCm" : "Tenstorrent")) + << "..."); } else { REQUIRE_MESSAGE(have_dev, label << ": device oracle golden absent (" << ids_name << " / " @@ -338,7 +342,7 @@ void RunGate(const std::string& repo_dir, const std::string& golden_subdir, << " engine=" << (first_div < 0 ? -1 : got[static_cast(first_div)]) << " committed anchor=" << (first_div < 0 ? -1 : anchor_ids[i * T + first_div]) - << (device_golden + << std::string(device_golden ? " — re-capture the device golden pair via " "qwen3-neartie-gap.py" : " — re-run qwen3-neartie-gap.py to refresh the gap " @@ -427,7 +431,8 @@ void RunGate(const std::string& repo_dir, const std::string& golden_subdir, << strict_exact << "/" << N << "; near-tie-band only: " << neartie_only << "/" << N << "; max gap " << (worst_gap / 1000.0) << " nats @ prompt[" << worst_i << "] tok=" << worst_j << "; " << fail << " forward-divergent" - << (device_golden ? "; anchor+gaps = device oracle-backed golden" : "") + << std::string(device_golden ? "; anchor+gaps = device oracle-backed golden" + : "") << ")"); REQUIRE(fail == 0); } From 4496116e8e4c6c5e60fca26980d255bae6bee8a4 Mon Sep 17 00:00:00 2001 From: Luca Barbato Date: Thu, 20 Aug 2026 23:59:17 +0200 Subject: [PATCH 4/5] docs(BACKEND-TENSTORRENT-HOST-FREE-FORWARD): record the #1488 golden re-adjudication outcome and index #1508 Moves the stale TT paged-engine golden bullet in the spec's Owed list to its resolved form with the re-derivation evidence (eager/captured dump identity, transformers teacher-forced max gap 375 mnats, zero band or top-K violations), corrects the #1488 record to say the before/after-#1476 comparison covered exit codes and the p1-tok-10 token only, and appends the #1508 index row for the doctest const char* print defect found and fixed in the same flow. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:zai-glm-5.3 [Maki] --- .agents/issue-index.md | 1 + .../specs/tenstorrent-host-free-forward.md | 32 ++++++++++++++++--- 2 files changed, 28 insertions(+), 5 deletions(-) diff --git a/.agents/issue-index.md b/.agents/issue-index.md index bab9e421e..e26bcd385 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -510,3 +510,4 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1487](https://github.com/mudler/vllm.cpp/issues/1487) | `BACKEND-TENSTORRENT-HOST-FREE-FORWARD` | **`test_release_metadata` is red on every aarch64 host, so `agent-preflight` cannot go green on the TT dev fleet.** `tests/scripts/test_release_metadata.py:55` stages the HOST arch — `shutil.copy2("/bin/true", ...)` — into the hardcoded `vllm.cpp-0.0.1-linux-x86_64-glibc-cpu.tar.gz` (`:143`); on x86_64 the ELF matches the manifest by accident, on aarch64 the validator correctly refuses: `ELF host architecture does not match manifest: ARM aarch64`. Measured at `206afb63` == `origin/main` with zero local commits and an uncommitted src-only diff; NOT the #1353 disk-full mode (902G free, no scratch-repo failures, the message is the mismatch itself). Repair direction: make the fixture self-consistent (stage an ELF matching the declared arch, or derive the declared arch from the staged ELF) so the validator is exercised for a manifest that lies. Listed under `## Owed` in [`tenstorrent-host-free-forward.md`](specs/tenstorrent-host-free-forward.md) | bug | | [#1488](https://github.com/mudler/vllm.cpp/issues/1488) | `BACKEND-TENSTORRENT-HOST-FREE-FORWARD` | **The TT `test_qwen3_paged_engine` golden is stale: `REQUIRE(anchor_ok)` fails with anchor drift prompt[1] tok=10 (engine=14126, committed anchor=62901).** Measured IDENTICALLY before and after the #1476 fix (same engine token both arms, ctest exit 8 both arms) — that test runs the DEFAULT path with no `VT_TT_HOST_FREE_DECODE`, which the fix does not touch, so the drift is not the captured-decode defect. The golden was committed `971d5506` (2026-08-09) and 13 TT ops commits landed Aug 9-18 with no on-card golden re-run. Owed: the `VT_DUMP_IDS` + `qwen3-neartie-gap.py` re-adjudication, then either a re-derived golden or an escalation to a real divergence. Split out of #1476 so the fidelity fix can close without orphaning the finding. Listed under `## Owed` in [`tenstorrent-host-free-forward.md`](specs/tenstorrent-host-free-forward.md) | bug | | [#1504](https://github.com/mudler/vllm.cpp/issues/1504) | `SPEC-DFLASH2` | **`test_audit_live_rows` is red on `main`: `KERNEL-DFLASH2-SELECTOR-EDGES` and `KERNEL-TOPK-PAIRS` remain `ACTIVE` after their implementation landed.** Measured at `origin/main` `cffe59b0` in a clean detached worktree: `scripts/audit-live-rows.py` reports `248 live rows; 2 abandoned ACTIVE` and `test_audit_live_rows` fails `test_shipped_record_has_no_abandoned_active_row` with exactly those two IDs. Both kernel-matrix rows (.agents/kernel-matrix.md:144-145) record their implementation as landed and gated on 2026-08-20 via SPEC-DFLASH2 W3 (PR #1497) yet still carry `ACTIVE`, so every preflight on current main inherits the red. Fourth instance of the #1469 class. NOT FIXED IN FLOW: the repair is a lifecycle move on `SPEC-DFLASH2`-owned rows and owes `docs/STATUS.md`, `docs/BENCHMARKS.md` and the moved spec's `## Now`; `KERNEL-TOPK-PAIRS` also carries the open CUDA NaN-tie disagreement #1489 that the move should reconcile or name as its open gate. Found while rebasing `BACKEND-TENSTORRENT-HOST-FREE-FORWARD` onto main; filed in flow | bug | +| [#1508](https://github.com/mudler/vllm.cpp/issues/1508) | `BACKEND-TENSTORRENT-HOST-FREE-FORWARD` | doctest `MessageBuilder` streams `const char*` as bool, so `test_qwen3_paged_engine` prints `committed anchor=62901` for a golden holding `6290` (and `96251` for `9625`) and a useless `1` label; proven with a 7-line harness against the pinned header, the garbled value was misread during #1488 as golden-buffer corruption before byte-level npy verification cleared the buffers; `label`/`ids_name`/`gap_name` now `std::string` and message ternaries wrapped, in the same flow | bug | diff --git a/.agents/specs/tenstorrent-host-free-forward.md b/.agents/specs/tenstorrent-host-free-forward.md index 36abbb8c8..6e616846d 100644 --- a/.agents/specs/tenstorrent-host-free-forward.md +++ b/.agents/specs/tenstorrent-host-free-forward.md @@ -308,11 +308,32 @@ not a process-global `GraphCapturesCounter`. Tracked on the host `/bin/true` into an `x86_64`-named archive, so `agent-preflight` cannot go green on the TT dev fleet. Found while running this row's preflight. Owned by [#1487](https://github.com/mudler/vllm.cpp/issues/1487). -- **The TT `test_qwen3_paged_engine` golden is stale**: anchor drift - prompt[1] tok=10 (engine=14126, committed=62901), identical before and after - the #1476 fix (that test runs the default path). Needs the `VT_DUMP_IDS` + - near-tie-gap re-adjudication. Owned by +- **The TT `test_qwen3_paged_engine` golden was stale** — RESOLVED 2026-08-20 + in this flow: the anchor drift prompt[1] tok=10 (engine=14126, committed + 6290 — the logged `62901` was a print artifact, #1508) came from the default + decode path having moved since the 2026-08-10 capture. The before/after-#1476 + comparison recorded in #1488 compared exit codes and only the p1-tok-10 + engine token, so it does not establish that pre-fix TT matches the refreshed + cells (prompt 5 tail, prompt 7); the golden is derived from this branch's TT + output and the refresh stacks here. Re-adjudicated with `VT_DUMP_IDS` (eager + `VLLM_CPP_CUDAGRAPH=0` and captured dumps byte-identical, md5 + `b5307e33…`) and `qwen3-neartie-gap-transformers.py` (transformers 4.57.1, + torch 2.10.0a0+cpu): 53 cells refreshed across 7 prompts — single + near-tie divergences at one token each (p1 tok10, p5 tok10, p10 tok5, + p11 tok4, p12 tok13, p15 tok12) whose greedy continuations then follow + the new prefix, plus prompt 7 rewritten from tok0; the new p7 row + matches the vLLM greedy sequence exactly, max gap 375 mnats, zero cells + above the 500-mnat band, zero outside-top-K. Fixed by [#1488](https://github.com/mudler/vllm.cpp/issues/1488). +- **doctest `MessageBuilder` streams `const char*` as bool**: every + separately-bound `const char*` in a `MESSAGE`/`REQUIRE_MESSAGE` renders as + `1`, so the anchor-drift message printed `committed anchor=62901` for a + golden holding `6290` (and `96251` for `9625`), and `label` printed as `1` + instead of `qwen3-0.6B`. Reproduced against the pinned header with a + 7-line harness; fixed by passing `std::string` in + `test_qwen3_paged_engine.cpp`. Found during the #1488 re-adjudication after + the garbled value had been misread as golden-buffer corruption. Fixed by + [#1508](https://github.com/mudler/vllm.cpp/issues/1508). The operator must still rerun the 80-token no-hang gate and `test_qwen3_paged_engine` on a Blackhole P150. An implementer run is an @@ -383,4 +404,5 @@ ambient `VT_TT_HOST_FREE_DECODE` (its exit-time segfault is pre-existing, Next: operator rerun of the 80-token captured-vs-eager gate and `test_qwen3_paged_engine` on card; the paged-engine golden re-adjudication is -[#1488](https://github.com/mudler/vllm.cpp/issues/1488). +done ([#1488](https://github.com/mudler/vllm.cpp/issues/1488) closed by the +stacked golden-refresh commit on this branch). From d871af49a3ad0dc8269e07712a193a96dbedc85d Mon Sep 17 00:00:00 2001 From: Luca Barbato Date: Fri, 21 Aug 2026 16:02:56 +0200 Subject: [PATCH 5/5] fix(BACKEND-TENSTORRENT-HOST-FREE-FORWARD): add the missing to test_qwen3_dflash2_gguf (#1595) Merging current main into this branch broke every rebuild on the clang-20 hosts: the SPEC-DFLASH2 GGUF drafter arm landed tests/vllm/models/test_qwen3_dflash2_gguf.cpp calling ::getpid() with no POSIX include, so the TU fails to compile and ninja stops before the TT gate targets can relink. The file is byte-identical to origin/main at e2a9e035d, so the defect is main's, inherited by any branch that merges main. The sibling tests include unguarded (test_kimi_linear_paged.cpp:44, test_loader_unaligned_offsets.cpp:46, test_ltx2_loader.cpp:26), so the one-line include is the established convention here. Closes #1595. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:zai-glm-5.3 [Maki] --- .agents/issue-index.md | 1 + tests/vllm/models/test_qwen3_dflash2_gguf.cpp | 1 + 2 files changed, 2 insertions(+) diff --git a/.agents/issue-index.md b/.agents/issue-index.md index ee2e08d6a..8804eb75a 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -532,3 +532,4 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1563](https://github.com/mudler/vllm.cpp/issues/1563) | `GATE-SQUASH-SEPARATOR` | **A markdown `---` horizontal rule anywhere in a pull request body silently voids the trailer block, and `check-commit-trailers.py` blames the trailers instead of the framing.** Found 2026-08-21 writing the body for PR [#1550](https://github.com/mudler/vllm.cpp/pull/1550) ([#1542](https://github.com/mudler/vllm.cpp/issues/1542)). `parsed_trailers()` shells out to git's trailer parser, and **git treats a line of exactly `---` as the start of the patch section**, so everything after the first one is not part of the message and a trailer block below it is invisible. Reproduced with no repository state: a body of `subject / prose / --- / more prose / FOLLOWING_AGENTS_PROTOCOL / the three trailers` reports `[trailers] Following-Agents-Protocol must appear exactly once` and `[attribution] AI-Assisted must appear exactly once`; `sed -i '/^---$/d'` on that same file reports `OK: commit trailer contract`, and the `---` is the only difference. **The MESSAGE is the defect, not only the behaviour**: `Following-Agents-Protocol` appears EXACTLY ONCE in the body while the checker says it must appear exactly once, so a reader counts occurrences, finds one, counts again and dumps bytes before thinking to test the parser's own framing. `_strict_errors` already computes `_paragraphs(body)[-1]` correctly as the three trailers verbatim, so the checker holds the information needed to say "the trailer paragraph is present but git could not parse it; a `---` line at line N ends the message". Worse, the neighbouring `FOLLOWING_AGENTS_PROTOCOL must appear exactly once as a separate paragraph before the trailer paragraph` check stays SILENT, so the two errors that fire both point away from the cause. **Beyond one confusing message**: the repository sets `squash_merge_commit_message = PR_BODY`, so the body IS the landed commit message, and a body carrying a `---` lands a commit whose trailers `git interpret-trailers` cannot see, on a branch that is never force-pushed. Same permanent-damage shape AGENTS.md records for the `---------` separator GitHub wrote under `COMMIT_MESSAGES`, arriving from the AUTHOR side rather than the forge side. `scripts/agent-pr-body.py --pr ` DOES catch it and caught it here before the merge; the exposure is a body never passed through that command, which AGENTS.md notes is not a gate and cannot be one because it reaches the network, while the CI guard reads the frozen `pull_request` payload and so does not re-read a body edited after the final push. NOT FIXED HERE: it changes a checker's semantics and its message, so under `## Changing the rules or a checker` it needs its own row, a red-before test and green-after evidence. Two candidate repairs, neither chosen: name the `---` line, or strip patch-section framing before parsing so a markdown rule is inert -- the second changes what the contract accepts and is the larger decision. Suggested minimum: `tests/scripts/test_check_commit_trailers.py` gains a case pinning the reproduction above | bug | | [#1454](https://github.com/mudler/vllm.cpp/issues/1454) | `SPEC-MTP-GGUF` | **`test_qwen3_5_gguf_mtp.cpp` reported `Status: SUCCESS!` with `assertions: 0` on every CI run, and its one arithmetic guarantee was a tautology.** Both cases opened `if (path == nullptr) return;` on `VLLM_MTP_GGUF_MODEL`, and a bare `return` from a doctest case is a PASS: re-derived on a clean Release build at `947e5f648`, unset, the file printed `test cases: 2 \| 2 passed \| 0 failed \| 0 skipped`, `assertions: 0`, `Status: SUCCESS!`, exit 0, and printed nothing else. The variable is set nowhere in `.github/workflows/`, so that was the state of every run. Second defect in the same file: the comment at `:52` stated `num_hidden_layers + depth == block_count` and the line under it asserted `CHECK(c.num_hidden_layers > 0)`, true of every valid model. MEASURED, not argued: mutating `src/vllm/model_executor/models/qwen3_5_gguf_weights.cpp:889` to `c.num_hidden_layers = block_count;` compiled clean and left the file at 2/2 cases, 0 assertions, `SUCCESS!`, exit 0. FIXED IN FLOW. The invariant is now pinned **HERMETICALLY** on KV-only synthetic GGUFs carrying no weight bytes, so CI checks it every run rather than never - 65/1 (the shipped Qwen3.8-27B pair), 25/1 (the Qwen3.5-2B reference this suite was developed against) and 28/3, the third arm separating `- nextn` from `- 1` - plus a head-less arm asserting the key is NOT published, which is the half `NumMtpLayers` cannot express because it answers 1 for an absent key. The two env-gated cases stay, now skipping with a `MESSAGE` naming the variable as `test_gguf_mmproj_reach.cpp` does, and the live one re-derives the invariant from the file's own `block_count` kv. Unset 4 cases / 18 assertions / `SUCCESS!` / rc 0; live on `Qwen3.8-27B-Q4_K_M.gguf` 4 / 38 / `SUCCESS!` / rc 0. Both mutants now red (9/18 and 5/18, exit 1), compiled clean, restored against a pre-taken sha256. **The production line is CORRECT and was not touched**: `block_count - nextn` landed `1a4db5c3c`, the `mtp_num_hidden_layers` republication `493327b4e`. Related but distinct: [#821](https://github.com/mudler/vllm.cpp/issues/821) W2 (`0adeb8b0e`) pins the same arithmetic for the 27B artifact on a committed manifest in `tests/vllm/models/test_qwen38_27b_gguf_manifest.cpp`, and that gate DOES catch both mutants - so the invariant was not globally unpinned, it was unpinned in this row's own file | bug | | [#1434](https://github.com/mudler/vllm.cpp/issues/1434) | `GATE-DOC-CHECKPOINT-STATES` | **`scripts/check-doc-checkpoint.py` could not see `PARTIAL`, so 118 state cells could move with no gate observing them.** `STATES` (`:56-66`) is the whole definition of what a lifecycle state IS for the gate that enforces AGENTS.md's `docs/STATUS.md` / `docs/BENCHMARKS.md` / spec `## Now` triple, and `row_states` drops any row it cannot match. `lifecycle_moves` and `moved_rows` then iterate the AFTER map, so leaving the matched set is silent by construction. Re-derived at `947e5f648` (the report measured `63d87805c`): `PARTIAL` **118** cells and `ANCHOR-BACKFILL` **73**, against `DONE` 77 and `BLOCKED` 9 — `PARTIAL` is the second most used state in the matrices and the gate was blind to it. Over the seven tables `ROW_TABLES` actually reads, the resolved population goes from **153 rows to 226**, a 47.7 % widening. Two of the transitions the report names behave differently from its description, measured with scratch commits at `947e5f648` on an unmodified checker: `READY -> PARTIAL` rc **0** and `PARTIAL -> READY` rc **0** are the real blind spots, while the report's suggested `PARTIAL -> ACTIVE` already reds — by accident, reporting **`added as ACTIVE`** for a row that has existed for months, because it is absent from the BEFORE map. FIXED IN FLOW for `PARTIAL` only. **`ANCHOR-BACKFILL` is deliberately excluded**: `.agents/feature-matrix.md:14-17` defines it as a property of the RECORD (*a legacy implemented row without exact code, test and real-spec anchors*), `docs/STATUS.md` carries no such term and would have nothing true to write on a `DONE <-> ANCHOR-BACKFILL` move, and `REQUIRED["lifecycle"]` cannot demand the spec's `## Now` alone — so admitting it would demand a public-document edit with nothing to say, which is the exact shape `check-doc-checkpoint.py:4-17` records as the reason the file was rewritten (16 of 20 red CI runs, six hardcoded escape hatches). One row's resolved state moves and the move is a REPAIR: `KV-BLOCK-POOL` says `` `PARTIAL` (not `DONE`) `` in its prose and the last-match heuristic believed the parenthesis, resolving `DONE`. No pinned counter moves — `check-gate-commands.py` has its own `GATED_STATES` and `RUNNABLE_BASELINE` is keyed on matrix rows, `UNOWNED_HIGH_WATER` is unmoved because this row names an owner, and no matrix row or public document changes — which was measured, not assumed, because this is the [#1376](https://github.com/mudler/vllm.cpp/issues/1376) ratchet shape. Remainder listed under `## Owed` in [doc-checkpoint-lifecycle-states.md](specs/doc-checkpoint-lifecycle-states.md): `ANCHOR-BACKFILL` moves, `.agents/sglang-matrix.md` never entering `ROW_TABLES`, a row that leaves the matched set entirely, and a new row added directly as `PARTIAL` | bug | +| [#1595](https://github.com/mudler/vllm.cpp/issues/1595) | `BACKEND-TENSTORRENT-HOST-FREE-FORWARD` | `tests/vllm/models/test_qwen3_dflash2_gguf.cpp:547` calls `::getpid()` with no `` in its include block, so the TU does not compile under clang-20 (`no member named 'getpid' in the global namespace`); landed on main in the SPEC-DFLASH2 GGUF arm and blocks every post-merge rebuild on a clang host, including the TT card gates; fixed in-flow with the one-line missing include, matching the sibling-test convention | bug | diff --git a/tests/vllm/models/test_qwen3_dflash2_gguf.cpp b/tests/vllm/models/test_qwen3_dflash2_gguf.cpp index de00dc735..8b9336303 100644 --- a/tests/vllm/models/test_qwen3_dflash2_gguf.cpp +++ b/tests/vllm/models/test_qwen3_dflash2_gguf.cpp @@ -59,6 +59,7 @@ #include #include #include +#include #include #include