diff --git a/.agents/backend-matrix.md b/.agents/backend-matrix.md index ad0b65d67..ef068a8c5 100644 --- a/.agents/backend-matrix.md +++ b/.agents/backend-matrix.md @@ -247,6 +247,7 @@ this repository. State remains `ACTIVE`; no lifecycle transition is claimed. | `BACKEND-TENSTORRENT-RESIDUAL-GOLDEN` | Child of `BACKEND-TENSTORRENT` — the owed op-level numerics evidence at the residual-RMS device boundary (`kDeviceResidualMinRows == 32`): device path does `ttnn::add`+`ttnn::rms_norm` in bf16; host/CPU path accumulates in f32. Bot-flagged on #289; never measured at the boundary. | vllm.cpp CPU oracle `RmsNormKernel` mirrors vLLM `fused_add_rms_norm` (add in model dtype, variance in f32); `src/vt/cpu/cpu_ops.cpp:371-398` | `src/vt/tenstorrent/tenstorrent_ops.cpp:1067-1117` (host/device split, `kDeviceResidualMinRows=32`) | [test_tenstorrent_backend.cpp](../tests/vt/test_tenstorrent_backend.cpp) `kRmsNorm residual: device vs CPU f32 oracle across the rows=32 boundary`: 22/22 cases on real Blackhole P150. **Measured 2026-08-11:** host path `rows<32` bit-identical to CPU (`max_abs=0`); device bf16 path `rows>=32` diverges by constant **0.0459 abs** (1.9–2.6× rel on near-zero outputs) — bf16 rounding signature, not accumulation. Decision pending the e2e golden tie-break | [tenstorrent-residual-golden.md](specs/tenstorrent-residual-golden.md) | `SPIKE` | `CLAIM-BACKEND-TENSTORRENT-RESIDUAL-GOLDEN` | | `BACKEND-TENSTORRENT-MISTRAL` | Child of `BACKEND-TENSTORRENT` — allowlist `MistralForCausalLM` (Mistral-7B-v0.3: GQA 32/8, head_dim 128, plain rope theta 1e6, untied lm_head, full attention) on the TT platform + device-aware SACRED gate. Mistral reuses the Qwen3-dense forward verbatim (qk-norm skipped); every op already registered. No new kernel. | vLLM `mistral.py::MistralForCausalLM(LlamaForCausalLM)` (already ported to the shared dense machinery); gate pattern mirrored from `test_qwen3_paged_engine.cpp:221-296` | `src/vllm/platforms/tenstorrent.cpp:52-54` (allowlist) + `tests/parity/test_mistral_paged_engine.cpp` (device-aware wiring + Backend Proof) | **Gate PASSED on Blackhole P150 (2026-08-12):** [test_mistral_paged_engine.cpp](../tests/parity/test_mistral_paged_engine.cpp) 16/16 prompts PASS (12/16 strict-exact, 4/16 near-tie, 0 forward-divergent), max gap **0.062 nats**, BACKEND PROOF 0 declines (kMatmul selections=256 = untied lm_head on device, kPagedAttention=8192). Goldens `our_ids_tenstorrent.npy` + `neartie_gap_mnats_tenstorrent.npy` (transformers alternative-oracle; POL-ORACLE deviation recorded, same as Qwen3-0.6B TT precedent). Exit SIGSEGV 139 is the known MeshDevice teardown crash, not a gate failure | [tenstorrent-mistral.md](specs/tenstorrent-mistral.md) | `ACTIVE` | `CLAIM-BACKEND-TENSTORRENT-MISTRAL` | | `BACKEND-TENSTORRENT-GDN` | Child of `BACKEND-TENSTORRENT` — the GDN linear-attention op chain as native TT kernels, the hard prerequisite for the Qwen3.5/3.8 family (#1715): `kGdnPrefill`, `kGdnDecode`, `kL2Norm`, `kRmsNormGated`, `kCausalConv1dFwd`/`kCausalConv1dUpdate`, `kGdnStateGather`/`kGdnStateScatter`. The P150 is discrete, so an op miss refuses by name — the ops must land before any `Qwen3_5*` arch registration. Correctness oracle is our own CPU f32 arm (residual-golden precedent); no vLLM mirror exists for TT | Substrate: pinned tt-metal `ttnn::transformer::chunk_gated_delta_rule` (FLA chunked GDN forward, on-core recurrent state, `initial_state`/`final_state`) behind a varlen+state-permute adapter for `kGdnPrefill`; decode = rank-1 update composed from ttnn matmul+eltwise with a device shadow keyed by host pointer (`PagedKvShadow` pattern); contracts at `src/vt/ops.cpp:1823-2500`, CPU reference `src/vt/cpu/cpu_ops.cpp:1537-1740` | `src/vt/tenstorrent/tenstorrent_ops.cpp` (kernels + `RegisterOp` block `:3165-3199`) + `tests/vt/test_tenstorrent_backend.cpp` (op-level cases vs CPU f32 oracle, T-sweep, indexed-`state_idx` form, prefill↔decode state round-trip) | Spec committed 2026-08-22; implementation waves W1 (prefill set) then W2 (decode+state-I/O set) owed; lands production-unreached until the wiring row registers `Qwen3_5*` (named in `## Owed` per Nothing-lands-dead) | [tenstorrent-gdn.md](specs/tenstorrent-gdn.md) | `ACTIVE` | `CLAIM-BACKEND-TENSTORRENT-GDN` | +| `BACKEND-TENSTORRENT-QWEN35` | Child of `BACKEND-TENSTORRENT` — the wiring row: `Qwen3_5ForConditionalGeneration` (dense text GDN hybrid) on the TT allow-list, the op delta the family forward refuses by name (`kGdnPostConv`, `kSigmoidGateBf16`, `kAttnQkNormRopeGate`, `kAttnQkNormRope` — pinned empirically by a W0 refusal sweep), and the first e2e gate. Makes the GDN row's ops production-reached (#1715 stays open until the family runs; GDN row lifecycle moves in the same change this lands) | Substrate: the GDN row's 8 kernels + the 27-op TT registry; e2e mirrors the ratified Mistral TT golden treatment (`VT_DUMP_IDS` bootstrap → `scripts/qwen3-neartie-gap-transformers.py` teacher-forced near-tie gaps → committed device-golden pair, POL-ORACLE deviation); op-level oracle is the CPU f32 arm. Capacity: `Qwen/Qwen3.5-0.8B` bf16 ≈1.6 GB fits (proven envelope Mistral-7B ≈14.5 GB); 27B bf16 ≈53.8 GB and GGUF k-quant arms (no TT kernels) refused by name, owed | `src/vllm/platforms/tenstorrent.cpp` (allow-list, lands last) + `src/vt/tenstorrent/tenstorrent_ops.cpp` (op delta) + `tests/parity/test_qwen35_paged_engine.cpp` (TT arm) | Spec committed 2026-08-23; W0 sweep, W1 op delta, W2 e2e, W3 GDN-reviewer leftovers (d2h counter completeness, `conv_transposed` fast-path check) owed | [tenstorrent-qwen35.md](specs/tenstorrent-qwen35.md) | `ACTIVE` | `CLAIM-BACKEND-TENSTORRENT-QWEN35` | | `BACKEND-TENSTORRENT-TRACE-RUNNER` | Child of `BACKEND-TENSTORRENT` — wire the landed graph-capture foundation (#354 / `59568772`) into a capturable forward region. Handoff §8/§9 "not done". **Decision (2026-08-13): NO-GO for pure T=1 decode capture — MEASURED, not assumed.** | CUDA is the ONLY backend with `SupportsGraphCapture()==true` (`cuda_backend.cu:184-240`); Metal/Vulkan `false`. Shared decode-graph framework `Qwen3DenseDecodeGraph` (qwen3.cpp:489, used by Qwen3/Mistral/Llama/InternLM2) gated on `support_static_graph_mode()` which TT does NOT override (base `false`) | `src/vt/tenstorrent/tenstorrent_backend.cpp:70-76` (capture surface, landed) + `tenstorrent_ops.cpp` `Trace*` (landed); NOT wired into a dense forward | **Measured on Blackhole P150 (2026-08-13):** Q1 no host-free region at T=1; Q2 all-device-at-T=1 costs 12.5→10.7 tok/s; Q2b capture attempt with both overrides flipped → ttnn **`TT_FATAL: Reads are not supported during trace capture`** (backtrace through `to_vector`), `0 replays`. The T=1 forward does device→host readbacks that ttnn trace prohibits; flipping the two thresholds is insufficient. Requires a host-free `ForwardLayers` redesign, not a threshold tweak. Next: prefill capture (separate row) must first audit its `to_vector` readbacks | [tenstorrent-trace-runner.md](specs/tenstorrent-trace-runner.md) | `SPIKE` (decision record complete) | `CLAIM-BACKEND-TENSTORRENT-TRACE-RUNNER` | | `BACKEND-TENSTORRENT-HOST-FREE-FORWARD` | Child of `BACKEND-TENSTORRENT` — make the per-decode-layer forward region host-free (zero `to_vector` readbacks) so mesh-trace capture can run. Decomposes into R1 RmsNorm+RoPE all-device, R2 QkvSplit+ReshapeAndCache device, R3 PA decode device-resident metadata, R4 capture wire+measure. Prerequisite revealed by the trace-runner spike (capture aborts on `to_vector`). | CUDA decode-graph capture contract (`cuda_backend.cu:184-197`): captured region is async, no host sync, no malloc, fixed ptrs. TT must match: no `to_vector` between Begin/EndCapture | `src/vt/tenstorrent/tenstorrent_ops.cpp` (RmsNorm:1067, PreferDeviceRope:1344, QkvSplit:1460, ReshapeAndCache:1527, PagedAttention:2009) | **R1-R3b MEASURED on P150** (env-gated `VT_TT_HOST_FREE_DECODE`, inert by default; 23/23 TT tests incl. a dedicated default-path inertness guard, M1-mutation-proved): R1 threshold flip landed; R2 device->device copy (`CopyDeviceDeviceIfCapture`, ttnn::copy+empty) landed; R3 program-cache warm (`enable_program_cache` + eager-warm) landed; R3b device zero-fill (`MemsetDeviceIfCapture`) landed. Capture enters the forward and reaches layer ops (CastBf16/RmsNorm fire). Remaining item-5 blocker = per-op enqueue_write; answer = persistent device tensors + before-replay populate (tt-metal vLLM plugin design). Full blocker map + architecture in [tenstorrent-host-free-r1.md](specs/tenstorrent-host-free-r1.md). **R5 (2026-08-21, #1604): host-free decode is the DEFAULT** — `HostFreeDecodeEnabled()` centralizes the polarity (exact `0` = pre-flip opt-out), both device golden pairs re-dumped and re-adjudicated under the new default (Qwen3 max 375 mnats, Mistral max 250 mnats, 0 outside top-K), both paged-engine gates 16/16 green (125/125 + 128/128 assertions), default leg 10.94-11.06 tok/s vs 5.34 opt-out (2.1x, Qwen3-0.6B b1). Capture declined by default: multi-request captured hangs (#1625); TT async scheduling stays off, no `SupportsAsyncSampledTokenReadback` override (#1627) | [tenstorrent-host-free-forward.md](specs/tenstorrent-host-free-forward.md) | `ACTIVE` (R1-R3b + R5 default flip gated on `row/BACKEND-TENSTORRENT-HOST-FREE-1604`; capture declined per #1625, TT async readback owed per #1627) | `CLAIM-BACKEND-TENSTORRENT-HOST-FREE-FORWARD` | ## Native competitor and performance gates diff --git a/.agents/claims/CLAIM-BACKEND-TENSTORRENT-QWEN35.md b/.agents/claims/CLAIM-BACKEND-TENSTORRENT-QWEN35.md new file mode 100644 index 000000000..c89bbe51f --- /dev/null +++ b/.agents/claims/CLAIM-BACKEND-TENSTORRENT-QWEN35.md @@ -0,0 +1,5 @@ +# CLAIM-BACKEND-TENSTORRENT-QWEN35 + +| Claim | Row IDs | Agent | Worktree / remote dir | Branch | Owned scope | State | Last update | +|---|---|---|---|---|---|---|---| +| `CLAIM-BACKEND-TENSTORRENT-QWEN35` | `BACKEND-TENSTORRENT-QWEN35` (`ACTIVE`, new) | Maki (zai-glm-5.3), helper role per `scripts/agent-role.py`, working from the committed spec `.agents/specs/tenstorrent-qwen35.md` | isolated worktree `/home/lu_zero/Sources/vllmcpp-tt-qwen35`; real Blackhole P150 (`thalia`, local board, `flock` mutex per GPU rule — not an `rc` fleet device) | `row/BACKEND-TENSTORRENT-QWEN35`, base `origin/main` `175733000`, issue [#1715](https://github.com/mudler/vllm.cpp/issues/1715) (tracks the wiring row per the GDN spec's `## Owed`) | Owns ONLY: the `Qwen3_5ForConditionalGeneration` TT allow-list entry (lands last), the W0 refusal sweep that pins the op-delta scope, the W1 op delta (`kGdnPostConv`, `kSigmoidGateBf16`, `kAttnQkNormRopeGate`, `kAttnQkNormRope`, plus whatever the sweep adds), the W2 e2e TT golden-pair gate for `tests/parity/test_qwen35_paged_engine.cpp` + the `docs/USAGE.md` weights entry (`Qwen/Qwen3.5-0.8B` @ `2fc0636471`, download authorized 2026-08-23), the W3 GDN-row reviewer leftovers (d2h counter completeness for `EnsureGdnCacheDevice` slow path + `CommitConvTransposed` fallback; the `conv_transposed` fast-path check), the spec, the backend-matrix row, and the two checker re-pins. EXCLUDES: MoE arches, GGUF k-quant arms (refused by name, owed), the VL tower, 27B-on-TT (no fitting arm), `src/vt/cpu/` (the oracle stays untouched), `src/vt/cuda/`, #1625/#1627 | `ACTIVE` | 2026-08-24 — W2b: pool-tenancy defect fixed (mutation-pinned), gemma `+1` defect fixed (mutation-pinned, op-level case added); ambient collapse healed — fixed engine matches the oracle on p0 and sits near-tie on 247/256 steps. Open: 2b (six out-of-band steps, leg-independent, intrinsic; layer bisect needs a GDN dump hook). Goldens in worktree describe the fixed engine. Nothing committed yet | diff --git a/.agents/specs/tenstorrent-qwen35.md b/.agents/specs/tenstorrent-qwen35.md new file mode 100644 index 000000000..58630c34a --- /dev/null +++ b/.agents/specs/tenstorrent-qwen35.md @@ -0,0 +1,935 @@ +# BACKEND-TENSTORRENT-QWEN35: the Qwen3.5 GDN family runs on Tenstorrent — allow-list, op delta, and the first e2e gate + +**Row:** `BACKEND-TENSTORRENT-QWEN35` (child of `BACKEND-TENSTORRENT`) +**Issue:** [#1715](https://github.com/mudler/vllm.cpp/issues/1715) (open; tracks the +wiring row per `tenstorrent-gdn.md` `## Owed` — "stays open until the family runs") +**Lifecycle:** `ACTIVE` (claimed 2026-08-23, helper) +**Prerequisite row:** `BACKEND-TENSTORRENT-GDN` (landed `175733000`; its eight ops are +the substrate this row makes reachable) +**Exemplar row:** `BACKEND-TENSTORRENT-MISTRAL` (`tenstorrent-mistral.md`) — the third +arch wired onto TT and the e2e recipe this row mirrors + +## Now + +`ACTIVE`. Spec committed spec-first 2026-08-23 with the matrix row and the two +checker re-pins. **W0 (refusal sweep), W1 (op delta), and W2a are done** — see +`## Evidence`. W2a flipped `SupportsCompressedConvState()` / +`SupportsCompressedGdnState()` (the production bf16 mamba-cache arms), pinned +both arms against the CUDA bf16-STORAGE emulation, added the arch allow-list +entry, and fixed the `ScatterRowsExact` L1 overflow that killed the first e2e +bootstrap (`SplitFactor` born-split shadows). Owed next: W2b — teacher-force +the dumped TT ids via `scripts/qwen3-neartie-gap-transformers.py`, commit the +TT golden pair, run the full 16/16 gate with BACKEND PROOF, and `docs/USAGE.md`; +then W3, the GDN-row reviewer leftovers. + +## Scope + +Make `Qwen3_5ForConditionalGeneration` (the dense text-only Qwen3.5 GDN hybrid) +run end to end on the Blackhole P150, and prove it: + +1. **Allow-list.** Add the arch to `TenstorrentPlatform::supports_model_architecture` + (`src/vllm/platforms/tenstorrent.cpp`). The P150 is discrete — an op miss refuses + by name, never falls back to CPU — so the allow-list entry lands LAST in the + implementation order, after the op delta is complete. +2. **Op delta (W1).** Complete exactly the op set the 0.8B forward refuses by name on + TT (W0 pins the list empirically). Named candidates from source inspection: + - `kGdnPostConv` — direct dispatch, no composite fallback + (`src/vt/ops.cpp:4270-4292`; model sites `qwen3_5.cpp:4146`, `:4600`, `:4980`). + - `kSigmoidGateBf16` — direct dispatch, no composite fallback + (`src/vt/ops.cpp:4142-4154`; model site `qwen3_5.cpp:2546`). + - `kAttnQkNormRopeGate` — the GDN full-attention preamble. The model checks + `OpRegistered` first (`qwen3_5.cpp:5192`, `:5209`, `:5317`); unregistered it + takes the unfused path (`AttnGateSplit` + `RmsNorm` + rope), and `AttnGateSplit` + is ALSO unregistered on TT — so W1 registers the fused op (the form the model + prefers and the ROCm gate's BACKEND PROOF names), not the unfused pair. + - `kAttnQkNormRope` — the plain (non-gated) preamble, same + registration-guarded shape; same decision. + The sweep is authoritative: if it refuses an op this list does not name, that op + is in scope too; if a candidate is not refused, it is not implemented. +3. **First e2e gate (W2).** The TT device-golden treatment for + `tests/parity/test_qwen35_paged_engine.cpp`, mirroring + `test_mistral_paged_engine.cpp`: `VT_DUMP_IDS=1` bootstrap on the P150 → + `our_ids_tenstorrent.i32`; transformers teacher-forced near-tie gaps via + `scripts/qwen3-neartie-gap-transformers.py` (the ratified POL-ORACLE deviation, + same tool and precedent as the Qwen3-0.6B and Mistral-7B TT goldens); commit the + TT golden pair; full gate 16/16 PASS with near-tie ≤500 milli-nats, strict + token-exact reported, and BACKEND PROOF (GDN op set selections > 0, declines == 0). +4. **Capacity decision (recorded, not deferred).** `Qwen/Qwen3.5-0.8B` bf16 + ≈1.6 GB fits the P150 trivially (proven envelope: Mistral-7B bf16 ≈14.5 GB). + Qwen3.8-27B bf16 ≈53.8 GB does NOT fit, and the GGUF k-quant arms have no TT + kernels, so the 27B-on-TT lane is refused by name and owed (see `## Owed`). + The first family gate is therefore the 0.8B bf16 arm — the cheapest checkpoint + that exercises the complete GDN chain end to end. +5. **Reviewer leftovers from the GDN row (W3).** (a) The d2h traffic counter misses + the `EnsureGdnCacheDevice` slow-path download and the `CommitConvTransposed` + untracked-buffer fallback; (b) `EnsureGdnCacheDevice`'s fast path does not check + `conv_transposed`, so a host pointer reused across roles would confuse + geometries (`qwen3_5.cpp` uses distinct buffers today, so this is a hardening + with a test, not a live bug). + +**Excludes:** the MoE arches (`Qwen3_5MoeForCausalLM`, +`Qwen3_5MoeForConditionalGeneration` — no TT MoE kernels, no fitting checkpoint); +GGUF k-quant arms (no TT kernels; refused by name — owed); the VL tower; 27B e2e; +`src/vt/cpu/` (the oracle stays untouched) and `src/vt/cuda/`; `#1625` capture and +`#1627` async readback (independent, as in the GDN row). + +## Upstream chain + +No vLLM Tenstorrent implementation exists (`vllm-omni` does not register TT either); +`tt-forge` is the registered secondary oracle for the hardware but the model-level +behavior reference is vLLM as everywhere. Two oracle lanes, both ratified: + +- **Op level:** our own CPU f32 arm (`src/vt/cpu/cpu_ops.cpp`) — the same + residual-golden precedent the GDN row used. +- **E2e level:** the pinned vLLM-ROCm greedy capture is the base golden + (`tests/parity/goldens/qwen35_greedy_0_8b/manifest.json`, vLLM `555967922`, + model revision `2fc06364715b967f1860aea9cf38778875588b17`), but the gate compares + TT prefixes, so the near-tie gaps are re-derived teacher-forced via + `scripts/qwen3-neartie-gap-transformers.py` — the exact deviation the Mistral row + recorded and the Qwen3-0.6B TT golden ratified before it. + +## Our baseline + +TT registers 27 ops today (`tenstorrent_ops.cpp:4434+`), including the GDN row's +eight. The Mistral row proved the full e2e recipe on this exact board: bootstrap → +transformers gap (max 0.0625 nats there) → committed golden pair → 16/16 with +BACKEND PROOF, tolerating the known exit-139 MeshDevice teardown (#1486) after the +green summary. + +## Port map + +| Op | Contract | CPU oracle | Model sites | TT realization | +|---|---|---|---|---| +| `kGdnPostConv` | `ops.cpp:4270-4292` (g/beta/a_log/dt_bias f32; a/b f32-or-bf16 views) | `cpu_ops.cpp:3800` | `qwen3_5.cpp:4146`, `:4600`, `:4980` | ttnn eltwise composition (sigmoid/exp/mul chains), the GDN row's elementwise pattern | +| `kSigmoidGateBf16` | `ops.cpp:4142-4154` (attn f32/bf16, gate f32, out bf16) | `cpu_ops.cpp` (grep `kSigmoidGateBf16`) | `qwen3_5.cpp:2546` | ttnn sigmoid + mul + cast | +| `kAttnQkNormRopeGate` | `ops.cpp` recipe `:1124-1136`; wrapper `AttnQkNormRopeGate` | CPU standalone op | `qwen3_5.cpp:5220`, `:5378` | fused: gemma-RMSNorm(q,k) + partial NeoX rope + gate split, one launch or the row's fused pattern | +| `kAttnQkNormRope` | FusedChain fast realization `ops.cpp` (`DispatchFusedFast`) | CPU standalone op | full-attn layers (guarded) | same shape minus the gate | + +The implementer pins exact CPU line anchors when writing each kernel; the contract +column above is the entry point, not the whole chain. + +## Design + +- **Fused over unfused for both preambles.** The model already prefers the fused op + when registered; the unfused path needs `AttnGateSplit`, which would be a fifth + kernel for no benefit. One fused kernel each mirrors the registration-guard + semantics and keeps the BACKEND PROOF list identical to the ROCm gate's. +- **bf16 storage, f32 intermediates** exactly as the GDN row: model-path buffers + stay the model dtype; an f32 intermediate names its reason inline. +- **Allow-list lands last.** The discrete device's by-name refusal is the safety + property; the arch string only reaches the allow-list when the sweep is empty. +- **Host-free decode stays default-on** (`VT_TT_HOST_FREE_DECODE` unset); the e2e + gate runs ambient plus the `=0` leg, as the GDN row did, so the flip is exercised + both ways at e2e depth, not just op depth. +- **Weights.** `Qwen/Qwen3.5-0.8B`, revision `2fc06364715b967f1860aea9cf38778875588b17` + (verified 2026-08-12 via download metadata, `qwen35_greedy_0_8b/manifest.json`), + bf16, ~1.6 GB. Download to the local HF cache authorized 2026-08-23. + `docs/USAGE.md` gains the TT arm entry (repo, revision, size) in the landing + change, with the refused arms named beside it. + +## Tests to port + +- Op-level doctest cases in `tests/vt/test_tenstorrent_backend.cpp`, one per new op, + `CompareVsOracle` against the CPU f32 arm, in the GDN row's case style (shape + sweep + dtype arms + NaN-hardened comparator). +- Registration-sweep case: every op the 0.8B forward dispatches has a TT + registration (the empty-refusal property, asserted by name list). +- The TT arm of `test_qwen35_paged_engine.cpp`: device-golden pair load, near-tie + methodology, BACKEND PROOF assertions — mirroring `test_mistral_paged_engine.cpp`. +- W3: counter/hardening tests per leftover. + +## Gates + +1. **Focused:** new op cases green on the P150, both ambient legs + (`VT_TT_HOST_FREE_DECODE` unset and `=0`). +2. **E2e:** `tests/test_qwen35_paged_engine` 16/16 prompts PASS on the P150, + near-tie ≤500 milli-nats, strict-exact reported, BACKEND PROOF selections > 0 + and declines == 0 for the GDN op set, both ambient legs. Exit 139 after the green summary is the known #1486 teardown and + counts green. +3. **Full TT suite green; CPU gate green; `scripts/agent-preflight.sh` all-green.** +4. **Mutation evidence** per asserted guarantee, re-run by the fresh reviewer. + +## Dependencies + +- The pinned tt-metal checkout and the GDN row's adapter/shadow substrate + (unchanged by this row). +- Real Blackhole P150 (`thalia`) under the file mutex; not an `rc` fleet device. +- The 0.8B checkpoint in the local HF cache (authorized). +- A python env with torch+transformers for `qwen3-neartie-gap-transformers.py` + (the Mistral row used torch 2.7.1+cpu / transformers 5.8.1; reuse or rebuild). +- Independent of #1625/#1627. + +## Work breakdown + +- **W0 — sweep (pins W1).** Scratch-wire the allow-list, run the 0.8B forward on + the P150, collect every by-name refusal. Recorded in `## Evidence`. Not committed + as product; the sweep's op list is the W1 scope contract. +- **W1 — op delta.** Implement each refused op as a TT kernel + oracle case + + mutations, the GDN row's method. +- **W2 — e2e.** Allow-list entry (one line, final commit of the wave), bootstrap, + transformers gap, committed TT golden pair, full gate, `docs/USAGE.md`. +- **W3 — leftovers.** d2h counter completeness; `conv_transposed` fast-path check; + tests for both. + +Each wave lands focused-green before the next; the full gate + fresh review close +the row. + +## Risks + +- **Cold JIT window.** 0.8B is smaller than Mistral-7B's cold run but larger than + Qwen3-0.6B's ~30 min; the e2e bootstrap needs an uninterrupted board window with + a warm cache on retries. +- **Near-tie deviation.** The gap golden is transformers-teacher-forced, not + vLLM-ROCm — ratified twice already; the spec records it a third time rather than + silently inheriting. +- **A real divergence** (outside the near-tie band) is an open gap: report, do not + tune the band. Stop condition. +- **Host-free decode interplay** with the new elementwise ops (capture stays + opt-out; ambient and `=0` legs both gate). + +## Stop conditions + +- Any need to edit `src/vt/cpu/cpu_ops.cpp` (the oracle) — out of scope, escalate. +- A non-near-tie e2e divergence that does not reproduce on the CPU arm. +- The P150 or the checkpoint download becoming unavailable mid-row (gate stays + PENDING with the named blocker; never waived silently). + +## Owed + +- **GGUF k-quant arms for the family on TT** — no TT kernels; refused by name. + A standing family requirement (AGENTS `## Shared seams`), not a choice. +- **Qwen3.8-27B on TT** — no arm fits the P150 (bf16 53.8 GB; quant kernels owed + above). Refused by name at load. +- **MoE arches / 2.4T lane** — TT MoE kernels do not exist. +- **GDN under capture** — unchanged from the GDN row's `## Owed` (#1625 first). +- **`kCausalConv1dUpdate` production bf16 arm** — capability refusal (W0 item 4): + TT already computes the bf16 cache via f32 shadows; enabling it needs the + `SupportsCompressedConvState()` flip in `src/vt/tenstorrent/tenstorrent_backend.cpp`, + outside the W1 delegated file set. Escalated to the operator 2026-08-23; blocks the + W2 e2e gate on the default (bf16-state) arm, not on `VT_GDN_STATE_BF16=0`. +- On landing, the GDN row's lifecycle moves (`ACTIVE` → `DONE` + `## Outcome`) in + the same change: its ops become production-reached. + +## Git integration + +One pull request for spec and implementation (row claim answer 2026-08-23, recorded +in `.agents/developer-preferences.md`). Base `origin/main` @ `175733000`. Branch +`row/BACKEND-TENSTORRENT-QWEN35`, worktree `/home/lu_zero/Sources/vllmcpp-tt-qwen35`. + +## Evidence + +All board runs on the P150 (`thalia`) inside the `${GPU_LOCK:-$HOME/gpu.lock}` file +mutex, `TT_METAL_HOME=/home/lu_zero/Sources/tt/tt-metal` (pinned tree), build +`ninja -C build tests/test_tenstorrent_backend`. Exit 139 after a green doctest +summary is the known #1486 teardown, not a gate failure. + +### W0 — refusal sweep (runs 1-8, `/tmp/w0_sweep_run{1..8}.log`) + +Scratch-wired the arch allow-list (reverted before commit; `src/vllm/platforms/ +tenstorrent.cpp` carries no Qwen3.5 entry in this change). Refusals, in sweep +order: + +1. `kGdnPostConv` — `no kernel for op GdnPostConv` (run 1). +2. `kAttnQkNormRopeGate` — unregistered on TT, reached 12x. +3. `kSigmoidGateBf16` — unregistered on TT, reached 12x. +4. `kCausalConv1dUpdate` — capability refusal on the PRODUCTION bf16 arm + (`conv_state` bf16 + TT `SupportsCompressedConvState()`=false, ops.cpp:1721). + The TT kernel already handles bf16 caches via f32 shadows; the fix is the + backend-flag flip in `tenstorrent_backend.cpp`, OUTSIDE this task's file set — + **escalated to the operator** (owed below). + +Kernel DEFECT surfaced (not a refusal): `ScatterRowsExact` — TILE rank-4 reshape +inflated physical volume (32 GiB OOM, runs 2-4); ROW_MAJOR reshape overflowed L1 +circular buffers (runs 5+7). Fixed in W1 (below). + +NOT refused, because never dispatched (fused inside the two compositions above): +`kAttnQkNormRope`, `kGdnGBeta`, `kGdnConvSplit`, `kAttnGateSplit`. Their W0 CPU +sweep-stubs are deleted; no TT kernels owed. + +Sweep completion: run 8 (f32 state arm, CPU stubs) generated 2 tokens end to end +(`first=220`). Sweep reach counts: GdnPostConv 36, GdnStateScatter 36, GdnDecode +18, AttnQkNormRopeGate 12, SigmoidGateBf16 12. + +### W1 — op delta (all focused logs under `/tmp/w1_*.log`) + +Three new TT kernels in `src/vt/tenstorrent/tenstorrent_ops.cpp`, each red-first +(Registration removed → `no kernel for op ...` refusal captured), then green +against the CPU f32 oracle, then negatively mutated and restored: + +- `SigmoidGateBf16Kernel` — red `/tmp/w1_red_sigmoidgate.log` (exit 1, `no kernel + for op SigmoidGateBf16 (id 64)`), green `/tmp/w1_green_sigmoidgate.log`: + 16/16 configs, `max_abs=0` (SFPU sigmoid rounds to the same bf16 values as + the CPU oracle in every tested element). +- `GdnPostConvKernel` — red `/tmp/w1_red_gdnpostconv.log` (id 71), green + `/tmp/w1_green_gdnpostconv.log`: 100/100. First green attempt used one-shot + `ttnn::softplus` and FAILED 5/100 at `g max_rel 4.4e-4` (SFPU poly ~1e-6 + ABSOLUTE fit error); the committed form composes `relu(x) + + log1p(exp(-|x|))` in f32 (`≤1.9e-7` rel, reproduces the threshold-20 branch). +- `AttnQkNormRopeGateKernel` — red `/tmp/w1_red_attnqknormropegate.log` (id 73, + op_provider.cpp:563), green `/tmp/w1_green_attnqknormropegate.log`: 6/6 configs + (T=1/3/65, Hq/Hkv GQA 4:2 and 32:8, rot 64/128, gemma on/off, f32/bf16 in), + q/k `max_abs ≤ 5.5e-4` (envelope 0.02), gate `max_abs=0` (exact passthrough). + A first implementation sliced+reshaped ONE `[T, Hq*2Dh]` upload on device; the + q/gate legs returned full-scale wrong data while the full-width k leg was + correct — slice+reshape on a fresh `from_vector` TILE upload is unsafe at this + pin. Committed form: the q|gate split happens in the host gather; every leg + uploads already-shaped (no device slice/reshape). +- `ScatterRowsExact` zero-copy fix (ROW_MAJOR convert + `ttnn::experimental::view` + `[slots,1,1,cols]` + indexed_fill dim=0, ~256 KB): verified on the REAL kernel + paths (`/tmp/w1_evidence_scatter_decode_roundtrip.log`): GdnStateGather/Scatter + `max_abs=0` (B1), GdnDecode-vs-oracle (B2) and the Prefill<->Decode final-state + round-trip (B3) both exit 0. `kGdnDecode`/`kGdnStateScatter` now register the + REAL kernels; the W0 CPU sweep-stubs are deleted from the tree. + +Negative mutations (each focused gate red, then restored byte-for-byte; +`/tmp/w1_mutations_m1_m2.log`, `/tmp/w1_mutations_m3_m4.log`): + +- M1 sigmoid dropped from SigmoidGateBf16 → 16/16 assertions failed + (`max_rel 1.5e4`), exit 1. +- M2 composed softplus → one-shot `ttnn::softplus` (the historical defect) → + 95/100, 5 failed at `g max_rel 4.4e-4`, Status FAILURE. +- M3 NeoX rope sign flip (`subtract`→`add`) → q/k failed in all 6 configs. +- M4 gate round-tripped through bf16 → gate failed the exact-passthrough check + on every f32-in config (`max_abs 0.00195` = one bf16 ULP). + +Full gate (no scratch stubs, allow-list reverted) — BOTH legs GREEN: +`/tmp/w1_fullgate_leg1.log` (ambient) and `/tmp/w1_fullgate_leg2.log` +(`VT_TT_HOST_FREE_DECODE=0`): 36/36 test cases, 2259/2259 assertions, +`Status: SUCCESS!` on each (exit 139 = #1486 teardown after the summary). + +### W2a — bf16 cache arms + allow-list + the L1 scatter fix (all focused logs +### under `/tmp/w2a_*.log`) + +**The capability flips** (`src/vt/tenstorrent/tenstorrent_backend.cpp`): +`SupportsCompressedConvState()` and `SupportsCompressedGdnState()` now return +true. The TT kernels compute through f32 shadows and honor bf16 STORAGE at the +boundary: every committed shadow value re-rounds through bf16 on device (RNE +typecast round-trip, zero PCIe), mirroring CUDA's "read/written in f32 +registers" (`cuda_backend.cu:119`, `cuda_gdn.cu`). This states as a capability +what `CheckConvCommon` asks; it un-refuses the W0 item-4 production arm. + +Red-first / mutation evidence for the flips and their oracle arms: + +- RED `/tmp/w2a_red_conv_bf16.log`: with the flip reverted, the bf16 conv_state + arm refuses by name (the ops.cpp compressed-state check). GREEN + `/tmp/w2a_green1_conv_bf16.log`, `/tmp/w2a_green2_conv_bf16.log`: the + bf16-state arm matches the CUDA bf16-STORAGE emulation — step 1 exact, + steps 2+ need the per-step store rounding (without it a full-mantissa x tap + diverges in step 2+; with it `max_abs ≈ 5.8e-4` full-mantissa arm and + bit-exact values where x is bf16-representable, the production activation + dtype). +- RED `/tmp/w2a_red_gdn_bf16.log`: gdn_decode refused for a bf16 state before + the flip ("state must be f32, or fp16/bf16 on a backend whose GDN kernels + support a compressed state", ops.cpp:1783). GREEN `/tmp/w2a_green_gdn_bf16.log`: + per-step outs match the f32 path with host-side bf16 round-trips bit-for-bit + after the commit-time rounding. +- Reviewer LOW-a/b repairs: std::string labels, out_bf16 arms, and the + kAttnQkNormRopeGate envelope tightened to `1e-3` f32-out with M5 redone — + `/tmp/w2a_green_lowb.log`, `/tmp/w2a_green_gate_tight.log` (q/k max_abs ≤ + 5.3e-4 against the tightened envelope), mutations `/tmp/w2a_mut_m1_conv.log` + (M1: drop the conv store-rounding → steps 2+ fail), `/tmp/w2a_mut_m2_gdn.log` + (M2: drop the GDN commit rounding → bit-exactness fails), M5 redo + `/tmp/w2a_mut_m5_gate_redo.log`; each red, each restored byte-for-byte. +- Allow-list: `Qwen3_5ForConditionalGeneration` added to + `TenstorrentPlatform::supports_model_architecture` LAST in implementation + order (the discrete device's by-name refusal is the safety property). + +**The L1 finding and the split-shadow fix.** The first e2e bootstrap +(`/tmp/w2a_e2e_bootstrap.log`) threw in `ttnn::prim::IndexedFillDeviceOperation`: +"Statically allocated dataflow buffers on core range [0-0 - 10-9] grow to +2208704 B which is beyond max L1 size of 1572864 B". Root cause: +`ScatterRowsExact` staged TWO FULL PAGES OF THE LAST DIM through indexed_fill's +generic interleaved path (`indexed_fill_program_factory.cpp`: page_size = +padded_last_dim × elem_size, data DFB num_entries = 2); its comment assumed +cols = 32768 (256 KB staging), but the Qwen3.5 GDN ssm_state row is +Hv·Dk·Dv = 16·128·128 = 262144 f32 elems → 2 × 1 MB > the 1.5 MB budget. + +The first repair attempt (view `[rows, cols]` metadata-only as `[rows*nb, cb]`) +is REFUTED and recorded here: `tt::tt_metal::view` launches no program even +when the last dim changes (`tensor_ops.cpp` recreates the MeshBuffer over the +same address), BUT device pages are the interleave unit +(`buffer.cpp Buffer::page_address`: bank_offset = aligned_page_size · +(page_index / num_banks)) — a last-dim-changing view re-maps every page index +to a different bank and scrambles flat order. Measured: one block landed, +three came back scrambled (`/tmp/w2a_diag.log` era). The landed design makes +the SHADOW BORN SPLIT instead: + +- `SplitFactor(cols, esz)` picks the smallest divisor F of cols with + 2·(cols/F)·esz ≤ 1 MiB of L1 staging (F=1 → byte-identical legacy form; + F=2 at cols=262144 → blk=131072, 512 KB pages, 1 MB DFB — inside budget). +- `EnsureGdnCacheDevice` uploads the SAME flat host bytes as + `[rows·F, cols/F]` (pure geometry at H2D time — no extra copy); logical dims + stay in the slot bookkeeping, so downloads and volume checks are unchanged. +- `ScatterRowsExact` takes `factor`, expands each live slot s into block ids + s·F+j, expands the NULL-compaction row list the same way, and views rank-4 + with the last dim UNCHANGED (the safe metadata-only case); still ONE launch. + Per-(slot, block) last-of-duplicates wins because entry order survives + inside every block group. +- `GdnDecodeKernel` gathers through a factor-adapted one-hot + (`UploadOneHot(..., F)`), reshapes S_new to `[B·F, blk]` (one exact + per-tile-CB regroup when F>1 — the only added device movement), applies the + bf16 storage rounding, then scatters. `GdnStateGatherKernel` expands its + index rows and `has_initial_state` mask to block-rows. +- The `CausalConv1dUpdate` scratch-row indexed_fill is NOT in this hazard + class — computed, not guessed: its tensors are TILE, where indexed_fill + page = tile size (4 KB), not row width; its `ttnn::reshape`s go through + reshape_tiled with per-tile CBs. Unchanged. + +Focused tests (`tests/vt/test_tenstorrent_backend.cpp`): a wide-row +gather/scatter arm at cols = 262144 (bit-exact vs CPU oracle, untouched slot +kept) and a wide-state decode arm with a NULL row (compaction feeding the +split launch, out AND cache within the standing envelope). + +Mutation evidence for the fix: + +- M-A (chunking removed — `SplitFactor` pinned to 1): the wide arms throw the + EXACT bootstrap throw ("grow to 2208704 B ... beyond max L1 size of + 1572864 B"), test case FAILED → `/tmp/w2a_red_l1_wide.log`. Restored → green + `/tmp/w2a_green_l1_wide.log` lineage. +- M-B (bid expansion drops the block offset — `s*F` instead of `s*F+j`): wrong + columns written, wide-arm cache compare fails → `/tmp/w2a_mut_bid_offset.log` + (1058/1059, Status FAILURE). Restored → green `/tmp/w2a_restore_check.log`. + +Focused suites green after the fix: gather/scatter + decode (116 assertions), +round-trip (71), conv-update + prefill — `/tmp/w2a_focused_green.log`, +`/tmp/w2a_focused_green2.log`, `/tmp/w2a_focused_green3.log`. + +Known upstream instability surfaced by the new width (recorded, not waived): +with SEVERAL consecutive wide-cache arms in one process, tt-metal's +RealtimeProfilerManager receiver thread can abort mid-test with host heap +corruption ("free(): invalid size", gdb lands in +`D2HSocket::pages_available`) — an upstream D2H-socket fragility in the same +class as the recorded #1486 teardown noise, reachable only because W2a makes +262144-wide shadows runnable at all. The landed focused arms use the minimal +shape that stays stable across reruns; if a future gate hits it, rerun and +record, do not widen the tolerance or skip the arm. + +### W2b — e2e bootstrap: the DevicePool tenancy defect (fix in flight) + +The first W2b bootstrap died at a readback guard: +`unexpected result size: got 2048 want 4096 out_shape=4x1024` from +`DownloadToHost` (`/tmp/w2b_diag_run1.log`, after runs 2-5 narrowed it — run 5 +is `/tmp/w2a_e2e_bootstrap_run5.log`). The slot-trace instrumentation +(`VT_TT_SLOT_TRACE=1`, `/tmp/w2b_diag_run5.log`) pinned it: + +- `RmsNormKernel -> EnsureDevice2D -> EnsureHost` reads the layer input + `[T,1024]` f32 and finds the slot still holding the PREVIOUS tenant's device + shadow (`[8,256]` bf16 — the full-attn `k_out` geometry of an earlier step). +- The buffer was registered ONCE and never freed or re-registered: it is a + `DevicePool` block. `DBuf` draws scratch from the pool; a pool HIT returns a + retained block WITHOUT calling `Backend::Alloc`, so nothing ever told the TT + backend that the block changed tenants. The slot kept the dead tensor's + shadow (`host_current=false`) and the next reader downloaded it. +- This is why the row (and Mistral/GDN before it) never saw it: their paths do + not recycle pooled scratch into read-first tensors on TT; Qwen3/Qwen3.5's + `DBuf` pool does. It is ALSO silent-corruption shaped: when the stale + geometry happens to match, the new tenant inherits the old bytes marked as + current. + +Fix (three pieces, in this change): + +1. `vt::Backend::OnScratchBlockAcquired(void*)` — new defaulted virtual; + `DevicePool::Get` calls it on every free-list hit. Default no-op: no other + backend keys residency by pointer. +2. `TenstorrentBackend` overrides it with `MarkHostWritten(p)` — the same state + a fresh Alloc registers. +3. `CommitDeviceLogical2D` now asserts the committed device tensor's volume + against `rows*cols`, so a producer bug dies at its producer instead of at a + later reader. + +Red evidence for the fix is the bootstrap failure itself (runs above). The +mutation (pool hit stops notifying) must rethrow the same mismatch — recorded +below once run. + +### W2b — e2e bootstrap GREEN, then the output-quality bisect + +With the tenancy fix in, the bootstrap completes: `Status: SUCCESS!`, 104/104 +assertions, ids dumped to `our_ids_tenstorrent.i32` (16x16), BACKEND PROOF +`kPagedAttention=1536 kGdnDecode=4320 kCausalConv1dUpdate=4320, 0 declines` +(`/tmp/w2b_fix_run1.log`; exit 139 is the #1486 teardown). The teacher-forced +golden pair was derived (`qwen3-neartie-gap-transformers.py`, torch +2.7.1+cpu / transformers 5.8.1 in `/home/lu_zero/Sources/tt/venv-qwen35gap`) +— and the DUMPED SEQUENCE IS DEGENERATE: every prompt collapses to +`,`/space (`11`/`220`) within a step or two. transformers teacher-forcing +AGREES with our tokens per step (`tf_argmax == our` throughout), so the gap +golden faithfully describes what our engine emitted; the engine's emission is +what is wrong. + +Bisect (all on one prompt, `vllm-cli --device auto --max-tokens 12`, +~4 min/run after the first JIT warm run): + +- **CPU arm clean**: `--device cpu` prints ` Paris.` — identical to the ROCm + golden. Model/metadata plumbing is correct; the divergence is TT-resident. + (First CLI attempt used a STALE pre-W2a binary from 12:47 and replayed the + already-fixed L1 throw — void run, rebuilt before any conclusion.) +- **bf16-state arms exonerated**: `VT_GDN_STATE_BF16=0` collapses identically. +- **Host-free decode is the carrier**: `VT_TT_HOST_FREE_DECODE=0` prints ` the + capital of the United States.` — coherent greedy text. Ambient (default ON) + collapses. + +This matches the spec's own risk line ("host-free decode interplay with the +new elementwise ops"). MECHANISM CORRECTED after reading #1625: on TT, +trace capture is OPT-IN (`VT_TT_DECODE_CAPTURE`; the multi-request hang made +`support_static_graph_mode()` decline by default), so the AMBIENT leg here is +HOST-FREE EAGER decode, not capture. The carrier statement stands +(`=0` — the legacy host-staged decode — is coherent while ambient collapses); +what breaks is the host-free eager chaining of the NEW ops' per-step inputs +and device-resident outputs (the dense ops each carry explicit host-free +plumbing — DecodeIdsCache, RacIdxCache warm hooks, rope warm-before-capture; +none exists for kGdnPostConv/kSigmoidGateBf16/kAttnQkNormRopeGate or the GDN +state path). The eager op-level suites (36/36) cannot see it; the dump-only +bootstrap had no correctness bar; only the `=0` leg runs the ops uncaptured +by host-free. Logs: +`/tmp/w2b_tt_ambient2.log` (collapse), `/tmp/w2b_tt_f32state.log` (collapse), +`/tmp/w2b_tt_nohostfree.log` (coherent). + +### W2b — the =0 gate red is a SEQUENCING error, and it sharpened the diagnosis + +The first golden pair was derived from the DEGENERATE bootstrap sequence +(before any output-sanity check) — a process mistake, recorded here so it does +not repeat: validate emitted text BEFORE deriving goldens. The pair was +quarantined to `/tmp/w2b_quarantined_goldens/` (never committed). + +The `=0` full-gate run against that pair then failed at the anchor REQUIRE — +correctly, since the committed anchor described the broken engine: +`prompt[0] tok=0 engine=279 committed anchor=11` +(`/tmp/w2b_gate_nohostfree.log`). Decoded: our `=0` leg emits `" the"` where +the ROCm oracle (and OUR CPU ARM, verbatim) emit `" Paris"`, and the broken +ambient leg emitted `","`. Three-way split at token 0: + +| leg | tok0 | character | +|---|---|---| +| ROCm oracle + our CPU arm | ` Paris` (11751) | reference | +| TT `VT_TT_HOST_FREE_DECODE=0` | ` the` (279) | fluent; diverges from both oracles | +| TT ambient (host-free eager) | `,` (11) | degenerate collapse | + +So there are TWO open defects, not one: + +1. **Ambient collapse** (host-free eager). The host-free branches are NOT + decode-only — `CopyDeviceDeviceIfCapture`, `MemsetDeviceIfCapture`, + `PreferDeviceRope`, and the forced-device residual+RMS all activate during + PREFILL too, and ambient's TOKEN 0 is already wrong. Next probe (board + freeing): `vllm-cli --max-tokens 1` ambient vs `=0` — pure-prefill + divergence convicts the prefill-active host-free branches without decode + in the picture; then `VT_TT_TRACE_DEBUG=1` for the tensor flow. +2. **`=0` numeric divergence** (" the" vs " Paris"). Fluent text but a real + first-token divergence from BOTH oracles. Either an honest near-tie (the + teacher-forced gap band decides) or a second milder numeric defect in the + eager path. The `=0` re-bootstrap WITH `VT_DUMP_IDS=1` + (`/tmp/w2b_bootstrap_nohostfree.log`) captures this leg's true sequence; + teacher-forcing it answers near-tie-vs-defect directly. + +### W2b — the =0 leg is near-tie-clean EXCEPT 4 steps; two distinct defects + +The `=0` bootstrap completed green (`/tmp/w2b_bootstrap_nohostfree.log`, +Status SUCCESS, ids dumped) and its sequence is HEALTHY: 127/256 tokens +identical to the ROCm oracle, fluent throughout (p4 emits correct fibonacci +code; p15 tracks the oracle's structure). Teacher-forced gaps +(`qwen3-neartie-gap-transformers.py`, committed pair +`our_ids_tenstorrent.npy` + `neartie_gap_mnats_tenstorrent.npy`): + +- **252/256 steps inside the 500 mnats band**, most at exactly 0. +- p0 tok0 IS a genuine near-tie: our `" the"` vs transformers' `" Paris"` at + **375 mnats** — inside the band; not a defect. +- FOUR steps over: p1 tok0 (**6813**), p1 tok8 (4688), p10 tok14 (1000), + p15 tok13 (625). Real divergences, localized — NOT a global math error. + +So the row owes TWO fixes, in this order: + +- **2a (primary, blocks the default path): ambient host-free collapse.** + Prefill-active suspicion recorded above; trace probe queued. +- **2b (blocks the =0 leg): 4 out-of-band steps.** Localized numeric + divergence under the eager path; bisect by prompt/layer after 2a, since the + full gate needs both legs anyway. + +### W2b — DEFECT 2A ROOT-CAUSED AND FIXED: the gemma `+1` dropped on the device arm + +The bisection chain, each step board-pinned: + +1. **Prefill, not decode.** `vllm-cli --max-tokens 1` (pure prefill, zero + decode steps): ambient `,` vs `=0` ` the` (`/tmp/w2b_prefill_{AMB,NOHF}.log`). + The host-free branches active during PREFILL are the suspects. +2. **Kill-switch bisect** (scratch `VT_TT_HF_DISABLE`, since removed): with + ALL FOUR prefill-active branches declined — device rope + (`PreferDeviceRope`), forced-device residual+RMS, d2d `Copy` hook, device + `Memset` — ambient emits ` the` (`/tmp/w2b_hfdisable.log`). Per-branch + split (`/tmp/w2b_split_*.log`): **only `residual` convicts** (armed → `, +`; declined → ` the`); rope/copy/memset are innocent. +3. **Root cause:** `RmsNormKernel`'s device arm hands ttnn::rms_norm the RAW + affine and has ALWAYS been gemma-host-only — the file said so ("Gemma + style (w+1) is host-only for now — Qwen3 does not set gemma=true"). + Qwen3-dense never sets gemma; **Qwen3.5 sets gemma=true at 21 RmsNorm + sites**, so under host-free every forced-device norm silently dropped the + `+1`. Under `=0` the `args.gemma` clause routes to the host arm, which is + why that leg was merely near-tie-flavored instead of collapsed. + +Fix: when `args.gemma`, bake `w+1` host-side in f32 (the oracle's order, same +treatment as the fused preamble's `weff`) and upload as an F32 gamma — a +TRANSIENT upload that deliberately bypasses the `EnsureAffine1D` slot cache, +because the cached form is the RAW weight and a non-gemma consumer of the +same buffer must never read the baked one. + +Evidence ladder: + +- GREEN e2e: ambient CLI now emits ` Paris.` — the ROCm oracle's AND the CPU + arm's exact continuation (`/tmp/w2b_gemma_fix.log`). The fix even improves + on the pre-fix `=0` leg's `" the"` near-tie. +- Op-level pin added: + `kTENSTORRENT kRmsNorm gemma matches a host F32 reference (w+1)` in + `tests/vt/test_tenstorrent_backend.cpp` (small weights make the +1 dominate). +- Mutation RED: dropping the `+1` fails that case + (`/tmp/w2b_gemma_red.log`, CHECK max_abs_diff < 0.5f NOT correct); restored + → green `/tmp/w2b_gemma_green2.log`. + +### W2b — post-fix state: collapse healed, 2b isolated as leg-independent + +The post-fix ambient bootstrap ran against the stale pre-fix goldens still on +disk, so it took the COMPARE path: 3 `prompt_ok` reds vs an obsolete anchor +and a correct final `REQUIRE(fail==0)` failure (`/tmp/w2b_bootstrap_fixed.log`) +— expected, not a regression. Fresh ids were dumped anyway and re-derived into +a new committed-to-worktree pair describing the FIXED engine. + +Fresh ambient sequence: 247/256 steps inside the band; out-of-band steps are +p1 tok0 (**6813**), p1 tok10 (3125), p1 tok4 (1250), p1 tok5 (1188), +p10 tok14 (1000), p7 tok11 (812). Prompt 1 carries four of six. + +**2b is INTRINSIC, not cross-prompt contamination**: standalone CLI run of +prompt 1 reproduces the identical divergent continuation +(` when the world was a place of wonder...` vs oracle +` in a world where everything was made of atoms...`, +`/tmp/w2b_p1_solo.log`). The engine is self-consistent (solo == batch), +fluent, near-tie on >96% of steps — but computes specific contexts +differently from BOTH oracles by up to 6.8 nats. + +W2c queue (next sessions): + +1. **2b bisect by layer** on prompt 1: `VT_DUMP_ATTN` covers full-attn layers + only; GDN layers need an equivalent per-layer dump hook before the + first-drift layer can be named. Then op-level replay of the drifting + layer's real inputs vs the CPU arm. +2. Full gates both legs (blocked by 2b only). +3. Fresh review + landing of the whole wave (pool tenancy fix, gemma fix, + diagnostics, golden pair, USAGE entry). + +### W2c — 2b localized to the qkvz projection output + +Instrumentation added (env-gated debug hooks): `DumpGdnStage` inside +`GdnBlockPaged` (dumps `mixed`, `conv`, `postconv_q/v`, `core`, `gated` per +invocation under `$VT_DUMP_ACT`) and a pre-layer dual-stream snapshot +(`layer_-1_{hidden,res}`) beside the existing per-layer `VT_DUMP_ACT` loop. + +Bisect chain on prompt 1, prefill-only, ambient vs CPU: + +1. Residual stream after EVERY layer diverges from layer 0 onward + (`layer_0.bin`: max_abs 2.12, mean 0.11) — born in the first layer, not + accumulated. +2. Pre-layer inputs (`layer_-1_*`): BIT-IDENTICAL between arms. The defect + lives inside layer 0's mixer (a linear-attention/GDN layer). +3. `VT_DUMP_ACT_SUB` stage probes for layer 0: `post_input_norm` max 0.0098 + corr=1.000000 (the gemma-baked device residual+RMS arm is numerically + sound); `post_attn_norm` max 3.30 corr 0.962. The divergence enters in + the mixer. +4. First mixer checkpoint — the qkvz PROJECTION OUTPUT (`mixed`, + [5,6144]): max_abs 2.1, corr 0.9986, with ~92/30720 elements grossly + wrong including SIGN FLIPS (cpu +1.14 vs tt -0.96), scattered across + columns and rows. Not a layout block, not accumulation noise. +5. `VT_GDN_IN_BF16=0` (f32 activations both arms): IDENTICAL divergence — + the bf16 activation path is exonerated; so is the dtype policy. +6. Standalone-vs-batch prompt 1 identical ⇒ not cross-prompt state leak. + +**Suspect: the qkvz projection GEMM on TT** (kMatmul / merged-qkvz leaf) +producing scattered grossly-wrong elements at [T,1024]x[1024,6144] with +real layer-0 weights and verified-clean inputs. Next session: op-level +replay of exactly that shape with the dumped `post_input_norm` as input +and the real weight column, compared element-wise against the CPU result; +then read the winning kernel's accumulation/padding path for the defect. + +### W2c — replay verdict: the op is clean; the CAPTURES under host-free are not + +Op-level replays with the REAL captured bytes, isolated: + +- kMatmulBT bf16 x bf16 -> bf16, [5,1024]x[8192,1024] (qkvz): worst err + 0.163 over ±20 values — envelope. CLEAN. +- kMatmulBT bf16 x bf16 -> F32, synthetic: CLEAN (0.047). +- kMatmulBT with the REAL h0 + REAL w_ba bytes, split-arm signature + ([5,1024]x[16,1024] -> F32): CLEAN (0.045 / 0.057). +- Resident weights captured from BOTH arms byte-identical and clean + (`w_ba.bin`: no NaNs, ±0.18 range). + +And the decisive behavioral test: **`VT_POOL_BYPASS=1` produces output +IDENTICAL to the pooled run** — pool-block state cannot be carrying the +divergence, and (critically) the engine's TEXT IS FLUENT even where the +mm-captures claimed `b/a` outputs were 1e38/NaN. If `b/a` were truly +garbage in the live dataflow, exp()/sigmoid would explode and the text +would collapse. It does not. + +**Therefore: the stage-dump capture path itself is UNRELIABLE under +host-free decode.** The DBuf-tmp + Backend::Copy download pattern used by +every probe hook (DumpStage, DumpGdnStage, the mm/qkvz captures) can serve +bytes that do not match what the live chain consumes — most plausibly via +the CopyDeviceDeviceIfCapture clone interaction (clone enqueued on the +device while the readback path resolves different residency state). This +invalidates the specific "projection GEMM garbage" and "ba garbage" +findings above as statements about live data, and with them the +layer-localization derived from those dumps. + +What SURVIVES (behavioral ground truth, independent of captures): + +- Post-gemma-fix engine: fluent, deterministic, 247/256 steps inside the + near-tie band; six out-of-band steps concentrated on prompt 1 (worst + 6.8 nats at tok 0); solo == batch; f32-input leg reproduces it; + VT_POOL_BYPASS neutral. + +### W2c — slice-view cache poisoning FIXED; defect narrowed to kCausalConv1dFwd + +**Root cause #3 found and fixed (mutation-pinned):** `EnsureDevice2D` keyed its +staging cache on (base slot, dims) — an INTERIOR slice view +(`packed_weight.Slice(0, Hv, 2*Hv)` fed to the BA matmul) resolved to the base +slot and consumed ANOTHER slice's staged weights. Engine proof: TT's `a` +projection output equaled CPU's `b` output (corr 0.99998) while sharing the +`b` input bit-for-bit. Fix: hits and stores require `t.data == slot->host`; +interior views upload as unregistered transients; interior views of a +device-current base refuse loudly rather than serve stale bytes. Same guard +applied to `EnsureAffine1D`. Permanent regression case +("kMatmulBT slice views do not consume the base staging") runs the engine's +exact b-then-a sequence twice against distinct-half weights; mutation +(neuter the base-pointer check) goes RED at worst 49.5. + +Post-fix engine state: layer-0 mixer divergence SHRANK (post_attn_norm max +3.28 → 2.25) but persists. Trusted (=0-leg) stage dumps now name the next +suspect precisely: + +- `post_input_norm`: bit-identical ✓ (trusted) +- **`conv` output: max 1.99 corr 0.99908 — DIVERGES** (both arms' own + dedicated-allocation dumps; trustworthy) +- postconv/core/gated inherit it. + +The two arms even take different conv sub-paths (TT: indexed gather + +`kGdnStateGather`; CPU: manual gather), so the candidate defects are the +indexed-gather state contents (stale slot rows leaking past `has_initial`) +or `kCausalConv1dFwd` itself under real activations. Scattered O(1) +outliers across all tokens and q/k/v segments — not precision noise. + +Next session: capture `mixed`'s BASE (the packed projection output), conv +weight, and gathered state via dedicated whole-tensor dumps; replay +`kCausalConv1dFwd` (both sub-paths) element-wise against the CPU result. +Note: view-shaped dumps (numel ≠ base volume) are UNRELIABLE — always dump +whole allocations. + +Also recorded: after this fix the full-bootstrap teacher-forced gaps show +18 out-of-band steps (was 6) with top-K misses at tok0/tok1 of five prompts +— the sequences shifted because g/beta changed; the gate re-judges against +transformers each time. The remaining numeric defect(s) above are why. + +### W2c — trusted (=0-leg) localization results and the OPEN IDENTITY ANOMALY + +All captures below are on the `=0` leg where no host-free hooks fire; +downloads are plain EnsureHostBytes+memcpy. Prompt 1, prefill-only, +CPU arm vs TT arm: + +1. Residual stream diverges FROM LAYER 0 (max 2.11, corr 0.94) — confirmed + under the trusted path; the layer-0 finding was real. +2. Layer 0 `post_input_norm`: **BIT-IDENTICAL** between arms (max 0.0000). + Under `=0` the gemma clause routes this norm to the host arm on both + arms, so both run identical host math over identical inputs. ✓ +3. Layer 0 `post_attn_norm`: max 3.28 corr 0.965 ⇒ the divergence is born + INSIDE the layer-0 GDN mixer. Same conclusion as the ambient runs. +4. Mixer stages (real pass): `mixed` max 4.21 corr 0.82; gated max 3.98 + corr 0.66. Consistent chain-level divergence. + +Then the anomaly that consumed the session — and it is PRECISELY bounded: + +- Capturing `h` at the top of `ProjectGdnQkvz` yields bytes B that differ + from `post_input_norm` (A) by max 4.06 / corr 0.85 — IDENTICALLY ON BOTH + ARMS (cpu-vs-tt of the captures: 0.0). +- Both hooks print the SAME data pointer for the same layer. +- A RECHECK download after the mixer returns A again; a DIRECT-pattern + download beside DumpStage returns A. Only the ProjectGdnQkvz-time read + sees B. +- VT_POOL_BYPASS output ≡ pooled output; VT_TT_SLOT_TRACE around the + pointer shows exactly ONE register and no other events. +- Isolated replays of kMatmulBT on the captured bytes are envelope-clean; + if B were the true GEMM input on TT while CPU consumed A, the scattered + mixed/column differences (~0.3% gross errors incl. sign flips) follow + naturally from a ~0.09-magnitude input perturbation through K=1024. + +So EITHER something transiently swaps the mixer-input bytes during the +mixer and restores them (writer unidentified; slot trace shows none), OR +B is the true input and A the stale one — but then the arms' GEMMs consume +identical B and `mixed` should match, which it does not. Both horns are +contradictory; the next session must break the tie with an in-process +arbitration: hash the buffer at THREE points (post-norm, pre-GEMM, +post-mixer) inside ONE process on ONE arm and also feed the pre-GEMM bytes +through the op immediately, comparing against the committed result — that +answers "is the GEMM consuming what I captured" without any cross-process +assumption. + +Withdrawn claims stay withdrawn (projection-GEMM garbage etc.). What +stands: layer-0 mixer divergence, six band violations, all behavioral +ground truth. + +The pool-tenancy fix is ORTHOGONAL to both and already proven necessary: +without it no run reaches a summary (runs 2-5 crashed). Its mutation cell +(same bootstrap config, guard neutered) is in flight — expected to rethrow +the stale-shadow mismatch, closing the red/green pair with +`/tmp/w2b_fix_run1.log`. + +**Mutation cell CLOSED**: guard neutered + same bootstrap config rethrows the +EXACT original signature (`unexpected result size: got 2048 want 4096 +out_shape=4x1024 ctx=EnsureHost dev[8x256 dt=1]`, Status FAILURE — +`/tmp/w2b_mut_pool_gate2.log`); restored byte-for-byte → green lineage +`/tmp/w2b_fix_run1.log`. The fix's guarantee ("a pooled block never inherits +its previous tenant's residency") is mutation-pinned. + +### W2c — device-readback verification and the measurement reset + +New seam: DebugDeviceReadbackF32 (ops-layer, declared in tenstorrent_device.h) +— EnsureDevice2D + to_vector, so probes can compare the DEVICE-STAGED bytes +against the host master. First result: the TT-resident merged in_proj_qkvz +staging is BIT-PERFECT vs host (8,388,608 elements, zero diff). Weight staging +is exonerated. + +Also established, and now load-bearing for every future probe: the loader's +merged in_proj_qkvz = concat[in_proj_qkv; in_proj_z] ([8192,1024]) exists on +both arms and host masters are byte-identical across arms. + +MEASUREMENT RESET declared: several cross-run divergence numbers in the +earlier W2c notes mixed dtypes (bf16 bytes read as f32 in analysis scripts) +and paired matmul calls across arms that take DIFFERENT projection arms +(merged-vs-split), producing meaningless comparisons — including one that +motivated the earlier projection-GEMM-garbage claim. The trusted-facts list +shrinks to: + +- Behavioral: fluent text, ~157/256 oracle-token matches post-BA-fix, + teacher-forced band violations that shift with each numeric change + (deterministic per build). +- Layer-level (=0, whole-allocation dumps): residual divergence starts at + layer 0's mixer; input norm bit-identical. +- The BA slice-view poisoning was REAL (fixed, mutation-pinned). + +Next session MUST start from ONE dtype-explicit, dual-read-verified dump +utility (whole allocations only, header-recorded dtype+shape) and redo the +mixed -> conv -> postconv -> core -> gated localization through it. No numeric +claim made through the old ad-hoc captures survives. + +### W2c — TrustDump harness results: gated matches, pc_q uncorrelated; contradiction open + +Built the trusted measurement utility (VT_DUMP_TRUST): whole allocations, +typed header ('TDMP': dtype/rank/dims/numel/verified), DUAL-READ verified +(two independent Synchronize+Copy passes must agree byte-for-byte or no +payload is written). Pool double-hand hypothesis tested separately: a +200k-cycle Get/Put hammer over the real pool found zero collisions — the +pool is exonerated. + +Trusted (=0, p0, whole-allocation) stage matrix: + +- post_input_norm: BIT-IDENTICAL +- conv: max 0.035, corr 0.99998 — MATCHES +- gated: max 0.039, corr 0.99972 — MATCHES (envelope) +- pc_q (q after l2norm inside GdnPostConv): max 1.09, corr +0.028 — + UNCORRELATED, yet properly l2-normalized per head on BOTH arms (norms + ~1.0, std 0.088 both). Not a layout permutation: mutual-nearest-neighbor + bijectivity 1/80; zero rows have any sub-2.0-L1 partner. + +Physical contradiction: gated (which consumes prefill output over q/k/v/g/ +beta) matches within envelope while its upstream ql2 reads uncorrelated. +Either the pc_q dump reads the wrong allocation consistently (verified +stable-wrong), or layer-0's gated match is coincidental at envelope scale. +Next session instruments INSIDE kGdnPostConvKernel (dump dconv-in, ql2-out, +at the commit site in ops.cpp — same TU, same tensor objects) and dumps +g/beta (dg/dbeta) which were never captured. Also add TrustDump to the +merged-arm packed output and to every remaining stage so the whole chain +is covered by the verified instrument. + +### W2c — RESOLVED: the pc_q "contradiction" was layer misalignment in analysis + +Root cause of every anomalous reading this round: cross-arm comparisons took +`sorted()[0]` per stage name, which paired CPU's layer-0 file against TT's +layer-11 file (the arms emit different site mixes, so global counters drift). +Emission-aligned comparison across all 18 layers: + +- pc_q: corr >= 0.9991 on every layer — MATCHES +- gated: corr >= 0.9983 on every layer — MATCHES +- conv: corr >= 0.9996 on every layer — MATCHES + +The layer-0 mixer chain is clean end-to-end on the =0 leg within the bf16 +envelope. There is no open mixer defect. Supporting verifications made along +the way: kernel-commit-site q_out == model-side dump bit-exact (no buffer +recycling); split-path BA gate inputs device==host==real values (=0 leg); +k_beta/k_g oddities were f32-payload files decoded bf16-style by analysis +scripts (writer stores raw device bytes; dtype tag 3 vs 2 must be honored). + +Instrument caveats recorded: + +- TrustDump under AMBIENT (VT_TT_HOST_FREE_DECODE unset) READ STALE HOST + BYTES during this row's interim builds (corrected per review finding F2: + Backend::Copy already refreshes the source via EnsureHostBytes since + 7faa9c6ba, so the shipped Copy path is sound; the stale readings came from + the pre-fix tree where ENSUREDEVICE2D — not Copy — was the broken reader, + and from analysis scripts decoding f32 payloads bf16-style). TrustDump now + states an explicit entry refresh as defense-in-depth; the necessity claim + in earlier W2c notes is withdrawn. +- VT_GLUE_FUSE=0 is not runnable on Tenstorrent: `GdnConvSplit` has no native + kernel (op_provider.cpp refuses the CPU reference tier). The fused chain is + load-bearing; A/B tests must vary other levers. + +Where the remaining gap can live, given the mixer chain is exonerated: +decode-phase packed path and state carry across steps, the ambient host- +staleness family itself, or logits/sampling. Next session opens there, from +emission-aligned dumps only. + +### W2c — ROOT CAUSE CLOSED: EnsureDevice2D consumed stale host bytes under host-free decode + +Behavioral split, same build, same prompt ("The capital of France is"): +ambient (default) emitted `!!!(1, 2, 3, 4, 5`; VT_TT_HOST_FREE_DECODE=0 +emitted `Paris.`. Emission-aligned trusted dumps localized the zeroing: +embedding output IDENTICAL on both legs; the residual stream entering +layer 0's mixer ALL-ZERO on ambient only — every downstream stage (conv, +pc_q/pc_v across all 288 captures, gated) zero or decorrelated as a +mechanical consequence. + +Mechanism: EnsureDevice2D builds its upload buffer from HOST bytes +(LoadElemF32 loop) without checking slot residency, then marks the slot +host_current=true. Under host-free decode a producer commits device-only +(host_current=false), so the next consumer staging through here uploaded +pool-fresh zeros AND poisoned the residency record. + +Fix: EnsureHostBytes(t.data) before the read loop. TrustDump gets the same +one-line refresh (its Copy resolves t.data to whichever memory the address +maps to; ambient dual-read verification could not see stale host bytes). + +Evidence: ambient WITHOUT any instrumentation now completes coherently +("Water boils at" -> "100°C. If a 100"; instrumented ambient run reached +"Paris!"). The =0 leg is unchanged. Heisenbug note for the record: the +interim dump instrumentation masked the defect because EnsureHostBytes' +refresh side effect heals exactly the state the defect corrupts — several +"fixed it by adding a dump" observations during this session were that +masking, not progress. + +### W2c — SACRED GATE GREEN on ambient after the residency fix; golden pair re-derived + +Post-fix parity quantification on the DEFAULT (ambient) configuration, +16 prompts x 16 tokens against the pinned ROCm oracle greedy_ids: + +- exact token cells: 135/256 (stale anchor) -> 213/256 (fixed engine) +- fully-exact prompts: 5/16 -> 10/16 +- re-derived TT golden pair via the sanctioned procedure (VT_DUMP_IDS=1 + bootstrap dump, then qwen3-neartie-gap-transformers.py secondary-oracle + teacher-forcing): max gap 0.375 nats — every divergence inside the + 0.5-nat near-tie band. +- Full gate verdict: 16/16 prompts PASS (10 strict-exact, 6 near-tie), + 0 forward-divergent, doctest 146/146 SUCCESS. + +Justification for re-derivation per the gate's own drift rule: the fixed +engine's tokens match the pinned oracle EXACTLY on cells where the stale +anchor diverged (prompt[1] tok0 our==oracle==303 vs anchor 948; prompt[12] +tok0 our==oracle==9565). The stale pair encoded the corrupt-ambient zeros. + +Mutation proof for the fix: disabling the EnsureDevice2D refresh in a +scratch build regressed ambient to the exact pre-fix garbage ("!!!(1, 2, +3, "), restored byte-for-byte afterward. + +Owed (recorded, not blocking): the test binary SEGFAULTS during teardown +after printing its verdict (ttnn::Tensor deallocate -> GraphTracker:: +is_enabled, device-destruction order). Verdict unaffected; file the issue +and fix the teardown ordering separately. diff --git a/include/vllm/model_executor/models/device_pool.h b/include/vllm/model_executor/models/device_pool.h index 84e089950..68aa34649 100644 --- a/include/vllm/model_executor/models/device_pool.h +++ b/include/vllm/model_executor/models/device_pool.h @@ -125,75 +125,89 @@ class DevicePool { // storm this pool exists to remove, so it is never a timing configuration. if (Bypass()) return b.Alloc(bytes); const size_t key = ClassOf(bytes); + void* hit = nullptr; { std::lock_guard lk(mu_); ClassState& cs = classes_[key]; ++cs.live; if (cs.live - cs.base > cs.peak) cs.peak = cs.live - cs.base; if (!cs.free.empty()) { - void* p = cs.free.back(); + hit = cs.free.back(); cs.free.pop_back(); retained_ -= key; - block_class_[p] = key; + block_class_[hit] = key; ++hits_; - return p; - } - // BEST FIT OVER THE RETAINED POOL (#1922). A block held free in a LARGER - // class already satisfies this request, and refusing to lend it is what - // made retention a function of how many distinct shapes the traffic has - // shown rather than of how much one step concurrently needs. Measured on - // this tree: twelve sequential requests through `LoadedEngine::generate`, - // the LARGEST one first so every later request demanded strictly less, - // still grew the process heap from 1.71 MiB to 4.24 MiB — every buffer - // the later requests needed had already been allocated and returned, and - // none of it could be reused because a freed block could only ever serve - // its own class. - // - // MIRROR. torch's caching allocator, which is the allocator vLLM's - // activations come out of, searches its cached pool for the SMALLEST - // block at least as large as the request before it asks the driver - // (`c10/cuda/CUDACachingAllocator.cpp::get_free_block`). This is that - // search, over the class ladder instead of over a sorted block set, and - // it is the one structural difference that made a bounded upstream - // working set unbounded here. - // - // The borrow is BOUNDED at kBorrowMaxRatio, so a caller never holds more - // than twice the bytes it asked for while it holds a borrowed block. The - // line that DELIVERS that bound is the `kBorrowMaxSteps` budget, not the - // `probe > limit` test beside it — see `kBorrowMaxRatio`, which is the - // same bound written from the other end and is `static_assert`ed against - // this budget. The `limit` test is kept because it is what makes the - // guarantee hold for a request below `2^kClassBits` bytes, where the - // ladder keys exactly and 16 rungs is more than one octave. - // - // Upstream bounds the same waste with `kMaxSplitSize` plus its - // small/large pool split; we cannot split a driver allocation, so this - // pair is the whole bound. - // - // The block keeps its OWN class: `block_class_` records what the driver - // actually allocated, and `Put` returns it there. A borrow is therefore a - // loan and never a demotion — the large class gets its block back and can - // still serve a large request — which is what keeps the borrow from - // starving the class it came from. - if (BorrowEnabled()) { - size_t probe = key; - const size_t limit = (key > std::numeric_limits::max() / kBorrowMaxRatio) - ? key - : key * kBorrowMaxRatio; - for (int step = 0; step < kBorrowMaxSteps; ++step) { - probe = NextClassAbove(probe); - if (probe == 0 || probe > limit) break; - auto it = classes_.find(probe); - if (it == classes_.end() || it->second.free.empty()) continue; - void* p = it->second.free.back(); - it->second.free.pop_back(); - retained_ -= probe; - block_class_[p] = probe; - ++hits_; - return p; + } else { + // BEST FIT OVER THE RETAINED POOL (#1922). A block held free in a LARGER + // class already satisfies this request, and refusing to lend it is what + // made retention a function of how many distinct shapes the traffic has + // shown rather than of how much one step concurrently needs. Measured on + // this tree: twelve sequential requests through `LoadedEngine::generate`, + // the LARGEST one first so every later request demanded strictly less, + // still grew the process heap from 1.71 MiB to 4.24 MiB — every buffer + // the later requests needed had already been allocated and returned, and + // none of it could be reused because a freed block could only ever serve + // its own class. + // + // MIRROR. torch's caching allocator, which is the allocator vLLM's + // activations come out of, searches its cached pool for the SMALLEST + // block at least as large as the request before it asks the driver + // (`c10/cuda/CUDACachingAllocator.cpp::get_free_block`). This is that + // search, over the class ladder instead of over a sorted block set, and + // it is the one structural difference that made a bounded upstream + // working set unbounded here. + // + // The borrow is BOUNDED at kBorrowMaxRatio, so a caller never holds more + // than twice the bytes it asked for while it holds a borrowed block. The + // line that DELIVERS that bound is the `kBorrowMaxSteps` budget, not the + // `probe > limit` test beside it — see `kBorrowMaxRatio`, which is the + // same bound written from the other end and is `static_assert`ed against + // this budget. The `limit` test is kept because it is what makes the + // guarantee hold for a request below `2^kClassBits` bytes, where the + // ladder keys exactly and 16 rungs is more than one octave. + // + // Upstream bounds the same waste with `kMaxSplitSize` plus its + // small/large pool split; we cannot split a driver allocation, so this + // pair is the whole bound. + // + // The block keeps its OWN class: `block_class_` records what the driver + // actually allocated, and `Put` returns it there. A borrow is therefore a + // loan and never a demotion — the large class gets its block back and can + // still serve a large request — which is what keeps the borrow from + // starving the class it came from. + if (BorrowEnabled()) { + size_t probe = key; + const size_t limit = + (key > std::numeric_limits::max() / kBorrowMaxRatio) + ? key + : key * kBorrowMaxRatio; + for (int step = 0; step < kBorrowMaxSteps; ++step) { + probe = NextClassAbove(probe); + if (probe == 0 || probe > limit) break; + auto it = classes_.find(probe); + if (it == classes_.end() || it->second.free.empty()) continue; + hit = it->second.free.back(); + it->second.free.pop_back(); + retained_ -= probe; + block_class_[hit] = probe; + ++hits_; + break; + } } + if (hit == nullptr) ++misses_; } - ++misses_; + } + if (hit != nullptr) { + // The block changes tenants here: whatever the previous DBuf left at + // `hit` is dead. Backends that keep per-pointer residency must drop it — + // a fresh Alloc would have registered the block anew, and a pool HIT + // bypasses Alloc entirely (Tenstorrent keys its f32 device shadows on + // the host pointer; a stale shadow was downloaded into an unrelated + // tensor, #1715). Applies to BORROWED blocks identically: a loan from a + // larger class had its own previous tenant. Outside the pool mutex: the + // notification takes the backend's own locks. + b.OnScratchBlockAcquired(hit); + return hit; } void* fresh = nullptr; try { diff --git a/include/vt/backend.h b/include/vt/backend.h index 2dd019422..cd96768d3 100644 --- a/include/vt/backend.h +++ b/include/vt/backend.h @@ -26,6 +26,15 @@ class Backend { // Returns memory aligned to at least 64 bytes; StepArena depends on this. virtual void* Alloc(size_t bytes) = 0; virtual void Free(void* p) = 0; + // The DevicePool hands a RETAINED block back out (DevicePool::Get free-list + // hit). The block's contents are undefined and its PREVIOUS tenant's device + // residency — any state the backend keyed on the pointer — must be dropped + // before the new tenant is used. Default no-op suits backends that keep no + // per-pointer residency; Tenstorrent overrides, because its f32 device + // shadows are keyed by the host pointer (RegisterHostBuffer) and a stale + // shadow would otherwise be downloaded into (or committed over) an + // unrelated tensor. + virtual void OnScratchBlockAcquired(void* p) { (void)p; } virtual void Memset(Queue& q, void* p, int value, size_t bytes) = 0; // Same-device or host<->device transfer; on CPU this is memcpy. virtual void Copy(Queue& q, void* dst, const void* src, size_t bytes) = 0; diff --git a/scripts/check-agent-record.py b/scripts/check-agent-record.py index 136d68ce5..92de8bd76 100644 --- a/scripts/check-agent-record.py +++ b/scripts/check-agent-record.py @@ -403,7 +403,7 @@ # linear-attention op chain as native TT kernels — the hard prerequisite # for every Qwen3.5/3.8 arch on Tenstorrent. ACTIVE, spec-first; no # implementation yet. - "BACKEND": (AGENTS / "backend-matrix.md", 86), + "BACKEND": (AGENTS / "backend-matrix.md", 87), } ENGINE_MATRIX = AGENTS / "engine-matrix.md" diff --git a/scripts/check-gate-commands.py b/scripts/check-gate-commands.py index ca8a442a0..c8a472e5f 100755 --- a/scripts/check-gate-commands.py +++ b/scripts/check-gate-commands.py @@ -518,6 +518,12 @@ def audit() -> list[dict]: # decode goldens the ten kernel arms were transcribed from). Neither is a # `git diff`, so the row earns the entry rather than being carried by one. "KERNEL-LTX2-VAE", + # 2026-08-23: +BACKEND-TENSTORRENT-QWEN35 enters the runnable population + # when its spec lands (#1715 wiring row). Same inherited-credit shape as + # BACKEND-TENSTORRENT-GDN above: the credit is the pre-existing full TT + # suite, CPU gate, and agent-preflight.sh; the row's own sweep/e2e gates + # are owed (spec-first). + "BACKEND-TENSTORRENT-QWEN35", }) diff --git a/scripts/env-doc-allowlist.txt b/scripts/env-doc-allowlist.txt index 8f6c9a827..93409e181 100644 --- a/scripts/env-doc-allowlist.txt +++ b/scripts/env-doc-allowlist.txt @@ -234,3 +234,6 @@ VT_W4A4_TRUE # var must be documented in docs/ENVIRONMENT.md OR added here. VT_DUMP_ACT_SUB VT_DUMP_ATTN +VT_TT_SLOT_TRACE +VT_DUMP_QKVZ +VT_DUMP_TRUST diff --git a/src/vllm/model_executor/models/qwen3_5.cpp b/src/vllm/model_executor/models/qwen3_5.cpp index c0cb7b9f0..18b5dacfb 100644 --- a/src/vllm/model_executor/models/qwen3_5.cpp +++ b/src/vllm/model_executor/models/qwen3_5.cpp @@ -20,6 +20,7 @@ #include "vllm/model_executor/models/kv_cache_route.h" // KV-FP8 W3 store/read route #include "vllm/model_executor/models/dense_fp8_block_gemm.h" // MODEL-FP8-BLOCK-LINEAR (#1189 M4) #include "vllm/model_executor/models/device_pool.h" // DevicePool/Pool/AuxPool/ActivePool (shared) +#include "vt/tenstorrent/tenstorrent_device.h" // DebugDeviceReadbackF32 (TT-only debug seam) #include "vllm/model_executor/models/qwen3_5_dense.h" #include "vllm/model_executor/models/qwen3_5_internal.h" @@ -3639,6 +3640,42 @@ DBuf MatmulBTRawD(Dev d, const Tensor& x, const Tensor& weight, VT_CHECK(weight.shape[1] == x.shape[1], "qwen3_5 merged GDN proj: input/weight K mismatch"); DBuf out(d, out_dtype, {x.shape[0], weight.shape[0]}); + if (const char* qdir = std::getenv("VT_DUMP_QKVZ")) { + static std::atomic mmseq{0}; + const int call = mmseq.fetch_add(1, std::memory_order_relaxed); + auto cap = [&](const char* tag, const void* bytes, size_t n) { + std::FILE* f = std::fopen( + (std::string(qdir) + "/mm" + std::to_string(call) + "_" + tag + ".bin") + .c_str(), "wb"); + if (f) { std::fwrite(bytes, 1, n, f); std::fclose(f); } + }; + // In-process arbitration: hash-and-capture x BEFORE, run the REAL GEMM, + // capture out, then run a SHADOW GEMM from the same tensor and capture + // its output, then re-read x. Answers, without cross-process + // assumptions: did the real GEMM consume these bytes, is consumption + // deterministic, and does x change across the op? + std::vector xpre(static_cast(x.Numel()) * vt::SizeOf(x.dtype)); + d.b.Synchronize(d.q); + d.b.Copy(d.q, xpre.data(), x.data, xpre.size()); + vt::MatmulBT(d.q, out.t(), x, weight); + std::vector ore(static_cast(out.t().Numel()) * + vt::SizeOf(out_dtype)); + d.b.Synchronize(d.q); + d.b.Copy(d.q, ore.data(), out.t().data, ore.size()); + std::vector xpost(static_cast(x.Numel()) * vt::SizeOf(x.dtype)); + d.b.Copy(d.q, xpost.data(), x.data, xpost.size()); + DBuf shadow(d, out_dtype, {x.shape[0], weight.shape[0]}); + vt::MatmulBT(d.q, shadow.t(), x, weight); + std::vector osh(static_cast(shadow.t().Numel()) * + vt::SizeOf(out_dtype)); + d.b.Synchronize(d.q); + d.b.Copy(d.q, osh.data(), shadow.t().data, osh.size()); + cap("xpre", xpre.data(), xpre.size()); + cap("xpost", xpost.data(), xpost.size()); + cap("out_real", ore.data(), ore.size()); + cap("out_shadow", osh.data(), osh.size()); + return out; + } vt::MatmulBT(d.q, out.t(), x, weight); return out; } @@ -3661,6 +3698,19 @@ GdnBaOutput ProjectGdnBA(Dev d, const GdnLayerWeights& weights, weights.in_proj_ba.shape[1] == hidden.shape[1], "qwen3_5 merged GDN BA: invalid packed owner"); Tensor packed_weight = ResidentWeight(d, weights.in_proj_ba); + if (const char* qdir = std::getenv("VT_DUMP_QKVZ")) { + static std::atomic baseq{0}; + const int call = baseq.fetch_add(1, std::memory_order_relaxed); + if (call == 0) { + std::vector raw(static_cast(packed_weight.Numel()) * + vt::SizeOf(packed_weight.dtype)); + DBuf tmp(d, packed_weight.dtype, {packed_weight.Numel()}, packed_weight.data); + d.b.Copy(d.q, tmp.ptr(), packed_weight.data, raw.size()); + tmp.Download(d, raw.data()); + std::FILE* f = std::fopen((std::string(qdir) + "/w_ba.bin").c_str(), "wb"); + if (f) { std::fwrite(raw.data(), 1, raw.size(), f); std::fclose(f); } + } + } if (MergedGdnBaEnabled(d)) { out.packed_owner.emplace( MatmulBTRawD(d, hidden, packed_weight, @@ -3668,6 +3718,17 @@ GdnBaOutput ProjectGdnBA(Dev d, const GdnLayerWeights& weights, Tensor packed = out.packed_owner->t(); out.b = packed.Slice(1, 0, value_heads); out.a = packed.Slice(1, value_heads, 2 * value_heads); + if (const char* td = std::getenv("VT_DUMP_TRUST")) { + static std::atomic ba_seq{0}; + if (ba_seq.fetch_add(1, std::memory_order_relaxed) == 0) { + // Device truth of the packed matmul result and of the interior + // a-window, captured through the same verified instrument the + // kernel-side probes use. + vt::tenstorrent::TrustDump(d.q, td, "ba_packed_dev", packed); + vt::tenstorrent::TrustDump(d.q, td, "ba_a_win_dev", out.a); + vt::tenstorrent::TrustDump(d.q, td, "ba_b_win_dev", out.b); + } + } return out; } @@ -3683,6 +3744,16 @@ GdnBaOutput ProjectGdnBA(Dev d, const GdnLayerWeights& weights, } out.b = out.b_owner->t(); out.a = out.a_owner->t(); + if (const char* td = std::getenv("VT_DUMP_TRUST")) { + static std::atomic bas_seq{0}; + if (bas_seq.fetch_add(1, std::memory_order_relaxed) == 0) { + vt::tenstorrent::TrustDump(d.q, td, "sba_a_dev", out.a); + vt::tenstorrent::TrustDump(d.q, td, "sba_b_dev", out.b); + vt::tenstorrent::TrustDump(d.q, td, "sba_h_dev", hidden); + Tensor sba_wa = ResidentWeight(d, weights.in_proj_a); + vt::tenstorrent::TrustDump(d.q, td, "sba_wa_dev", sba_wa); + } + } return out; } @@ -3922,6 +3993,48 @@ GdnQkvzOutput ProjectGdnQkvz(Dev d, const GdnLayerWeights& w, const Tensor& h, w.in_proj_qkvz.shape[1] == h.shape[1], "qwen3_5 merged GDN qkvz: invalid packed owner"); Tensor packed_weight = ResidentWeight(d, w.in_proj_qkvz); + if (const char* qdir = std::getenv("VT_DUMP_QKVZ")) { + // Per-invocation replay capture (the engine WARMS UP with a dummy + // forward, so the FIRST call is not the real step): h per invocation, + // the resident merged weight once (device-independent). + static std::atomic qseq{0}; + const int call = qseq.fetch_add(1, std::memory_order_relaxed); + const std::string dir = qdir; + if (call == 0) { + const Tensor& wt = packed_weight; + // HOST master bytes... + std::vector raw(static_cast(wt.Numel()) * vt::SizeOf(wt.dtype)); + DBuf tmp(d, wt.dtype, {wt.Numel()}, wt.data); + d.b.Copy(d.q, tmp.ptr(), wt.data, raw.size()); + tmp.Download(d, raw.data()); + std::FILE* f = std::fopen((dir + "/w_host.bin").c_str(), "wb"); + if (f) { std::fwrite(raw.data(), 1, raw.size(), f); std::fclose(f); } + // ...and the DEVICE-STAGED copy the GEMM will actually consume, + // read back through ttnn via the ops seam. + { + std::vector vec = + vt::tenstorrent::DebugDeviceReadbackF32(d.q, wt); + std::FILE* f2 = std::fopen((dir + "/w_device.bin").c_str(), "wb"); + if (f2) { + std::fwrite(vec.data(), 4, vec.size(), f2); + std::fclose(f2); + } + } + } + { + if (call == 0) + std::fprintf(stderr, "[TT-DUMP] qkvz-h0 ptr=%p rows=%lld\n", + static_cast(h.data), (long long)h.shape[0]); + // NO DBuf here: a pool-backed tmp aliases live blocks (see the + // capture-reliability finding); read straight into a host vector. + std::vector raw(static_cast(h.Numel()) * vt::SizeOf(h.dtype)); + d.b.Synchronize(d.q); + d.b.Copy(d.q, raw.data(), h.data, raw.size()); + std::FILE* f = std::fopen( + (dir + "/h" + std::to_string(call) + ".bin").c_str(), "wb"); + if (f) { std::fwrite(raw.data(), 1, raw.size(), f); std::fclose(f); } + } + } if (detail::ShouldUseMergedGdnQkvz(detail::GdnMergedQkvzEligibility{ MergedGdnQkvzEnabled(d), vllm::platforms::GetPlatform(d.q.device.type).needs_weight_staging(), @@ -3931,6 +4044,8 @@ GdnQkvzOutput ProjectGdnQkvz(Dev d, const GdnLayerWeights& w, const Tensor& h, out.mixed = packed.Slice(1, 0, conv_dim); out.z = packed.Slice(1, conv_dim, conv_dim + value_dim); return out; + if (const char* td = std::getenv("VT_DUMP_TRUST")) + vt::tenstorrent::TrustDump(d.q, td, "packed", packed); } Tensor qkv_weight = packed_weight.Slice(0, 0, conv_dim); Tensor z_weight = packed_weight.Slice(0, conv_dim, conv_dim + value_dim); @@ -3939,6 +4054,10 @@ GdnQkvzOutput ProjectGdnQkvz(Dev d, const GdnLayerWeights& w, const Tensor& h, out.mixed = out.mixed_owner->t(); out.z = out.z_owner->t(); return out; + if (const char* td = std::getenv("VT_DUMP_TRUST")) { + vt::tenstorrent::TrustDump(d.q, td, "mixed", out.mixed); + vt::tenstorrent::TrustDump(d.q, td, "z", out.z); + } } // PERF-FP8-ALPHA-FOLD / #417 — the output dtype of the merged fp8 in_proj leaf. // vLLM's ModelOpt fp8 linear emits the model dtype (bf16); ours hardcoded f32, @@ -4732,6 +4851,24 @@ DBuf GdnBlockPagedMixedSpec(Dev d, const GdnLayerWeights& w, const HfConfig& cfg : MatmulBf16D(d, gated_bf16.t(), w.out_proj); // [T,H] } +// VT_DUMP_ACT stage probe (GDN): dump named intermediates per invocation so a +// layer-level divergence can be pinned to one kernel. Debug-only; Download +// syncs, never set on capture paths. +void DumpGdnStage(Dev d, const char* stage, const Tensor& t) { + if (std::getenv("VT_DUMP_ACT") == nullptr) return; + static std::atomic gdn_seq{0}; + const int call = gdn_seq.fetch_add(1, std::memory_order_relaxed); + const int64_t n = t.Numel(); + std::vector raw(static_cast(n) * vt::SizeOf(t.dtype)); + DBuf tmp(d, t.dtype, {n}, t.data); + d.b.Copy(d.q, tmp.ptr(), t.data, raw.size()); + tmp.Download(d, raw.data()); + const std::string path = std::string(std::getenv("VT_DUMP_ACT")) + "/gdn" + + std::to_string(call) + "_" + stage + ".bin"; + std::FILE* f = std::fopen(path.c_str(), "wb"); + if (f != nullptr) { std::fwrite(raw.data(), 1, raw.size(), f); std::fclose(f); } +} + DBuf GdnBlockPaged(Dev d, const GdnLayerWeights& w, const HfConfig& cfg, const Tensor& h, const StepDevInputs& sdi, const GDNAttentionMetadata& meta, @@ -4913,6 +5050,8 @@ DBuf GdnBlockPaged(Dev d, const GdnLayerWeights& w, const HfConfig& cfg, dcs.t(), sdi.gdn_non_spec_qsl.t(), sdi.gdn_has_initial.t(), conv_args); + if (const char* td = std::getenv("VT_DUMP_TRUST")) vt::tenstorrent::TrustDump(d.q, td, "conv", dconv.t()); + DumpGdnStage(d, "conv", dconv.t()); Tensor conv_cache = state.conv_state; vt::GdnStateScatter(d.q, conv_cache, dcs.t(), sdi.gdn_state_idx.t()); @@ -4928,6 +5067,8 @@ DBuf GdnBlockPaged(Dev d, const GdnLayerWeights& w, const HfConfig& cfg, dcs.t(), dqsl.t(), dhis.t(), conv_args); ScatterStateF32(d, state.conv_state, dcs, sidx, conv_row_elems); + if (const char* td = std::getenv("VT_DUMP_TRUST")) vt::tenstorrent::TrustDump(d.q, td, "conv", dconv.t()); + DumpGdnStage(d, "conv2", dconv.t()); } } else { // Pure decode: single-token conv step per sequence, IN PLACE on the persistent @@ -4998,6 +5139,13 @@ DBuf GdnBlockPaged(Dev d, const GdnLayerWeights& w, const HfConfig& cfg, vt::GdnPostConv(d.q, dql2.t(), dkl2.t(), vf.t(), dg.t(), dbeta.t(), dconv.t(), araw, braw, a_log_dev, dt_bias_dev, vt::L2NormArgs{1e-6F}); + DumpGdnStage(d, "mixed", mixed); + DumpGdnStage(d, "postconv_q", dql2.t()); + DumpGdnStage(d, "postconv_v", vf.t()); + if (const char* td = std::getenv("VT_DUMP_TRUST")) { + vt::tenstorrent::TrustDump(d.q, td, "pc_q", dql2.t()); + vt::tenstorrent::TrustDump(d.q, td, "pc_v", vf.t()); + } } else { DBuf qf(d, actdt, {T, Hk, Dk}); DBuf kf(d, actdt, {T, Hk, Dk}); @@ -5103,6 +5251,8 @@ DBuf GdnBlockPaged(Dev d, const GdnLayerWeights& w, const HfConfig& cfg, vt::GdnPrefill(d.q, o_pre, q_pre, k_pre, v_pre, g_pre, b_pre, dss.t(), sdi.gdn_prefill_qsl.t(), gdn_args); Tensor ssm_cache = state.ssm_state; + DumpGdnStage(d, "core", o_pre); + if (const char* td = std::getenv("VT_DUMP_TRUST")) vt::tenstorrent::TrustDump(d.q, td, "core", o_pre); vt::GdnStateScatter(d.q, ssm_cache, dss.t(), sdi.gdn_prefill_state_idx.t()); } else { @@ -5110,6 +5260,7 @@ DBuf GdnBlockPaged(Dev d, const GdnLayerWeights& w, const HfConfig& cfg, vt::GdnPrefill(d.q, o_pre, q_pre, k_pre, v_pre, g_pre, b_pre, dss.t(), dpqsl.t(), gdn_args); ScatterStateF32(d, state.ssm_state, dss, pidx, ssm_row_elems); + if (const char* td = std::getenv("VT_DUMP_TRUST")) vt::tenstorrent::TrustDump(d.q, td, "core", o_pre); } } } // end non-spec recurrence @@ -5169,6 +5320,8 @@ DBuf GdnBlockPaged(Dev d, const GdnLayerWeights& w, const HfConfig& cfg, : Reshape(gated_bf16.t(), {T * Hv, Dv}); vt::RmsNormGated(d.q, gated2, core2, z2, dnw, vt::RmsNormGatedArgs{eps, sigmoid_gate}); + DumpGdnStage(d, "gated", gated_bf16.t()); + if (const char* td = std::getenv("VT_DUMP_TRUST")) vt::tenstorrent::TrustDump(d.q, td, "gated", gated_bf16.t()); } else { DBuf dgated(d, DType::kF32, {T * Hv, Dv}); Tensor gated_f32 = z_strided ? Reshape(dgated.t(), {T, Hv, Dv}) : dgated.t(); @@ -7637,7 +7790,23 @@ void RunDenseLayerPaged(Dev d, const Qwen3_5DenseLayerWeights& layer, Tensor dw_in = ResidentWeight(d, layer.input_layernorm, {H}); DBuf dhn(d, DType::kBF16, {T, H}); vt::RmsNorm(d.q, dhn.t(), hidden.t(), dw_in, vt::RmsNormArgs{eps, true}, &res.t()); + if (std::getenv("VT_DUMP_ACT_SUB") != nullptr) + std::fprintf(stderr, "[TT-DUMP] post_input_norm ptr=%p rows=%lld\n", + static_cast(dhn.t().data), (long long)T); DumpStage("post_input_norm", dhn); + if (std::getenv("VT_DUMP_ACT_SUB") != nullptr) { + // Same-instant second read via the DIRECT pattern (no DBuf/pool): if + // these two disagree, the download patterns themselves diverge. + std::vector rawD(static_cast(T) * static_cast(H) * + vt::SizeOf(dhn.t().dtype)); + d.b.Synchronize(d.q); + d.b.Copy(d.q, rawD.data(), dhn.t().data, rawD.size()); + const std::string pd = std::string(std::getenv("VT_DUMP_ACT_SUB")) + + "/layer_" + std::to_string(dump_layer_idx) + + "_post_input_norm_DIRECT.bin"; + std::FILE* fd = std::fopen(pd.c_str(), "wb"); + if (fd != nullptr) { std::fwrite(rawD.data(), 1, rawD.size(), fd); std::fclose(fd); } + } DBuf attn = [&] { if (layer.is_linear_attention) { @@ -7652,6 +7821,21 @@ void RunDenseLayerPaged(Dev d, const Qwen3_5DenseLayerWeights& layer, }(); DumpStage("block_out", attn); + if (std::getenv("VT_DUMP_ACT_SUB") != nullptr) { + // Triple-read probe: re-download the INPUT-NORM buffer after the mixer + // ran. If it differs from the pre-mixer dump, something between them + // writes it; if equal, the earlier disagreement is a download artifact. + std::vector raw2(static_cast(T) * static_cast(H) * + vt::SizeOf(dhn.t().dtype)); + DBuf tmp2(d, dhn.t().dtype, {T * H}, dhn.t().data); + d.b.Copy(d.q, tmp2.ptr(), dhn.t().data, raw2.size()); + tmp2.Download(d, raw2.data()); + const std::string p2 = std::string(std::getenv("VT_DUMP_ACT_SUB")) + + "/layer_" + std::to_string(dump_layer_idx) + + "_post_input_norm_RECHECK.bin"; + std::FILE* f2 = std::fopen(p2.c_str(), "wb"); + if (f2 != nullptr) { std::fwrite(raw2.data(), 1, raw2.size(), f2); std::fclose(f2); } + } Tensor dw_post = ResidentWeight(d, layer.post_attention_layernorm, {H}); DBuf dh2(d, DType::kBF16, {T, H}); vt::RmsNorm(d.q, dh2.t(), attn.t(), dw_post, vt::RmsNormArgs{eps, true}, &res.t()); @@ -8239,7 +8423,21 @@ static DBuf ForwardLayers(Dev d, const Tensor& hidden_in, } int64_t fa_idx = 0, gdn_idx = 0; - for (int64_t l = 0; l < config.num_hidden_layers; ++l) { + if (std::getenv("VT_DUMP_ACT") != nullptr) { + for (int which = 0; which < 2; ++which) { + const Tensor& t = which == 0 ? hidden.t() : res.t(); + std::vector raw(static_cast(T) * static_cast(H) * + vt::SizeOf(t.dtype)); + DBuf tmp(d, t.dtype, {T * H}, t.data); + d.b.Copy(d.q, tmp.ptr(), t.data, raw.size()); + tmp.Download(d, raw.data()); + const std::string path = std::string(std::getenv("VT_DUMP_ACT")) + "/layer_-1_" + + (which == 0 ? "hidden" : "res") + ".bin"; + std::FILE* f = std::fopen(path.c_str(), "wb"); + if (f != nullptr) { std::fwrite(raw.data(), 1, raw.size(), f); std::fclose(f); } + } + } +for (int64_t l = 0; l < config.num_hidden_layers; ++l) { const Qwen3_5MoeLayerWeights& layer = weights.layers[static_cast(l)]; const PagedKvCache* kv = layer.is_linear_attention ? nullptr : &attn_kv[static_cast(fa_idx++)]; diff --git a/src/vllm/platforms/tenstorrent.cpp b/src/vllm/platforms/tenstorrent.cpp index 95403d900..28156fcc6 100644 --- a/src/vllm/platforms/tenstorrent.cpp +++ b/src/vllm/platforms/tenstorrent.cpp @@ -50,12 +50,16 @@ class TenstorrentPlatform final : public Platform { // was the first bring-up; Qwen3-dense is the second (same OPT→Qwen3 sequence // Metal used for M3a/M3b). Mistral-7B-v0.3 is the third: it reuses the // Qwen3-dense forward verbatim (qk-norm skipped, plain rope, untied lm_head), - // so every op is already registered — no new kernel. Anything else falls back - // to CPU via SelectQueue. + // so every op was already registered — no new kernel. Qwen3.5 (GDN hybrid) + // is the fourth: its op delta (kGdnPostConv, kSigmoidGateBf16, + // kAttnQkNormRopeGate, the GDN decode set) landed in the GDN/Qwen35 rows, + // with the bf16 mamba-cache arms enabled by the backend's compressed-state + // capabilities. Anything else falls back to CPU via SelectQueue. bool supports_model_architecture(std::string_view architecture) const override { return architecture == "OPTForCausalLM" || architecture == "Qwen3ForCausalLM" || - architecture == "MistralForCausalLM"; + architecture == "MistralForCausalLM" || + architecture == "Qwen3_5ForConditionalGeneration"; } // kPagedAttention + kReshapeAndCache are registered against the NHD diff --git a/src/vt/tenstorrent/tenstorrent_backend.cpp b/src/vt/tenstorrent/tenstorrent_backend.cpp index 71e844a40..9b84b6c60 100644 --- a/src/vt/tenstorrent/tenstorrent_backend.cpp +++ b/src/vt/tenstorrent/tenstorrent_backend.cpp @@ -49,6 +49,11 @@ class TenstorrentBackend final : public Backend { UnregisterHostBuffer(p); std::free(p); } + // DevicePool::Get hands a retained block to a NEW tensor without passing + // through Alloc, so the slot at that address still describes the previous + // tenant (device shadow committed, host stale). Drop the residency — same + // state a fresh Alloc registers: host current, no device copy (#1715). + void OnScratchBlockAcquired(void* p) override { MarkHostWritten(p); } void Memset(Queue&, void* p, int value, size_t bytes) override { // HOST-FREE-FORWARD R3: on-device zero-fill when capturing. if (MemsetDeviceIfCapture(p, value)) return; @@ -73,6 +78,26 @@ class TenstorrentBackend final : public Backend { // ttnn mesh-trace capture — see Trace* in tenstorrent_device.h / ops.cpp. bool SupportsGraphCapture() const override { return true; } + + // Production mamba_cache_dtype is bf16 (qwen3_5_common.cpp conv_dtype + // default), so the GDN decode conv-update addresses a bf16 conv_state. The + // TT conv kernel computes through its f32 transposed shadow and honors that + // STORAGE semantics at the host boundary — LoadElemF32 widens on the way in, + // StoreElemF32 narrows on the way out, and (the bf16-state arm's test in + // tests/vt/test_tenstorrent_backend.cpp) the shadow itself re-rounds + // through bf16 on every commit, mirroring CUDA's "read/written in f32 + // registers" (cuda_backend.cu:119, cuda_gdn.cu's conv kernels). This states + // as a CAPABILITY what the shared CheckConvCommon gate asks, the same + // device-agnostic query CUDA/ROCm/Vulkan answer. + bool SupportsCompressedConvState() const override { return true; } + + // The GDN decode recurrence passes the persistent ssm_state (production + // mamba_ssm_dtype = bf16) straight to kGdnDecode. The TT kernel computes in + // its f32 shadow but honors bf16 STORAGE semantics — the committed shadow + // re-rounds through bf16 every step (mirroring CUDA's "read/written in f32 + // registers", cuda_backend.cu:123 / rocm_backend.hip:353), pinned by the + // bf16-state arm of the kGdnDecode oracle test. + bool SupportsCompressedGdnState() const override { return true; } void BeginCapture(Queue&) override { TraceBeginCapture(); } void EndCapture(Queue&) override { TraceEndCapture(); } void Replay(Queue&) override { TraceReplay(); } diff --git a/src/vt/tenstorrent/tenstorrent_device.h b/src/vt/tenstorrent/tenstorrent_device.h index 7c7cab283..40955b893 100644 --- a/src/vt/tenstorrent/tenstorrent_device.h +++ b/src/vt/tenstorrent/tenstorrent_device.h @@ -4,6 +4,7 @@ #pragma once #include +#include #include #include #include @@ -16,6 +17,13 @@ namespace tt::tt_metal::distributed { class MeshDevice; } // namespace tt::tt_metal::distributed +// Forward declarations: this header stays light (the backend TU includes +// only vt/backend.h); the definition links against vt/ops.h types. +// MSVC C4099 (W-X): the real definitions are structs (vt/tensor.h:15, +// vt/device.h:107); forward-declaring them as class breaks /WX builds on +// windows the moment a TU includes vt/ops.h before this header. +namespace vt { struct Tensor; struct Queue; } + namespace vt::tenstorrent { using MeshDevice = tt::tt_metal::distributed::MeshDevice; @@ -50,6 +58,24 @@ MeshDevice& SharedMeshDevice(); // cached ttnn::Tensor that still owns device pages for that host pointer. // No-ops until the ops registrar has loaded (static init order: backend may // Free before ops if a test tears down early — Unregister is tolerant). +// Defined in tenstorrent_ops.cpp (needs ttnn). Declared here so model TUs +// can ask for a device-side readback of a staged tensor without linking +// ttnn themselves. `Queue`/`Tensor` are vt types (vt/ops.h is in scope +// wherever this header is included after it). +std::vector DebugDeviceReadbackF32(Queue& q, const Tensor& t); +// TRUSTED dump: whole tensor, typed header, dual-read verified (see ops.cpp). +void TrustDump(Queue& q, const char* dir, const char* name, const Tensor& t); +// Review finding F1 (#1715): these two are defined only in the TT-gated ops +// TU, but model TUs call them under a RUNTIME env check. Follow the +// WarmPagedKvShadow pattern: inline no-ops when the backend is not linked, +// so a default (non-tt-metal) configure of examples still links. +#ifdef VLLM_CPP_TENSTORRENT +#else +inline std::vector DebugDeviceReadbackF32(Queue&, const Tensor&) { + return {}; +} +inline void TrustDump(Queue&, const char*, const char*, const Tensor&) {} +#endif void RegisterHostBuffer(void* host, size_t bytes); void UnregisterHostBuffer(void* host); // Host bytes at `host` (or any interior pointer into that allocation) were diff --git a/src/vt/tenstorrent/tenstorrent_ops.cpp b/src/vt/tenstorrent/tenstorrent_ops.cpp index 6383b41bf..12cff379a 100644 --- a/src/vt/tenstorrent/tenstorrent_ops.cpp +++ b/src/vt/tenstorrent/tenstorrent_ops.cpp @@ -32,8 +32,11 @@ #include #include #include +#include #include +#include #include +#include #include #include #include @@ -52,6 +55,7 @@ #include #include #include +#include #include #include #include @@ -75,6 +79,7 @@ #include #include #include +#include #include // Forward declare clone (header not in installed includes) @@ -326,12 +331,44 @@ bool IsFloatDType(DType d) { return d == DType::kF32 || d == DType::kBF16 || d == DType::kF16; } -void DownloadToHost(ttnn::Tensor& dev, Tensor& out) { +namespace { +std::string DevShapeStr(const ttnn::Tensor& t) { + const auto s = t.logical_shape(); + std::string r; + for (uint32_t i = 0; i < s.rank(); ++i) { + if (i != 0) r += 'x'; + r += std::to_string(s[i]); + } + r += " dt=" + std::to_string(static_cast(t.dtype())) + + " lay=" + std::to_string(static_cast(t.layout())); + return r; +} +} // namespace + +void DownloadToHost(ttnn::Tensor& dev, Tensor& out, const char* ctx) { if (std::getenv("VT_TT_TRACE_DEBUG") != nullptr && tt_capture_active()) std::fprintf(stderr, "[TT-TRACE] to_vector readback DURING CAPTURE\n"); std::vector result = dev.to_vector(); + if (result.size() != static_cast(out.Numel())) { + void* bt[8]; + const int nbt = ::backtrace(bt, 8); + char** sym = ::backtrace_symbols(bt, nbt); + std::fprintf(stderr, "[TT-SLOT] MISMATCH self=%p frames=%d\n", + reinterpret_cast(&DownloadToHost), nbt); + for (int i = 0; sym != nullptr && i < nbt; ++i) + std::fprintf(stderr, "[TT-SLOT] bt[%d]=%p %s\n", i, bt[i], sym[i]); + std::fflush(stderr); + std::free(sym); + } VT_CHECK(static_cast(result.size()) == out.Numel(), - "tenstorrent: unexpected result size"); + std::string("tenstorrent: unexpected result size: got ") + + std::to_string(result.size()) + " want " + + std::to_string(out.Numel()) + " out_shape=" + + std::to_string(out.shape[0]) + "x" + std::to_string(out.shape[1]) + + "x" + std::to_string(out.shape[2]) + "x" + std::to_string(out.shape[3]) + + " rank=" + std::to_string(out.rank) + + " ctx=" + std::string(ctx) + + " dev[" + DevShapeStr(dev) + "]"); for (int64_t i = 0; i < out.Numel(); ++i) StoreElemF32(out, i, result[static_cast(i)]); } @@ -343,6 +380,14 @@ void EnsureHost(Tensor& t) { if (s == nullptr || s->host_current) return; VT_CHECK(s->device_current && s->device.has_value(), "tenstorrent: EnsureHost with no current device or host copy"); + if (std::getenv("VT_TT_SLOT_TRACE") != nullptr) + std::fprintf(stderr, + "[TT-SLOT] ensure t=%p numel=%" PRId64 + " slot=%p bytes=%zu dev=%ux%u dt=%d ht=%d dc=%d ct=%d\n", + static_cast(t.data), t.Numel(), + static_cast(s->host), s->bytes, s->dev_rows, + s->dev_cols, static_cast(s->device->dtype()), + s->host_current, s->device_current, s->conv_transposed); if (s->conv_transposed) { std::vector v = s->device->to_vector(); for (int64_t i = 0; i < t.Numel(); ++i) { @@ -355,7 +400,7 @@ void EnsureHost(Tensor& t) { s->host_current = true; return; } - DownloadToHost(*s->device, t); + DownloadToHost(*s->device, t, "EnsureHost"); s->host_current = true; } @@ -391,7 +436,23 @@ ttnn::Tensor EnsureDevice2D(const Tensor& t, MeshDevice& device) { "tenstorrent: EnsureDevice2D expects contiguous rank-2"); const uint32_t rows = static_cast(t.shape[0]); const uint32_t cols = static_cast(t.shape[1]); + // The per-slot cache describes the BASE allocation. An interior view + // (e.g. packed_weight.Slice(0, Hv, 2*Hv) fed to a matmul) resolves to the + // base slot but must NEVER hit nor store against it: the cached tensor is + // the BASE's staging, and returning it for a differently-offset view makes + // the consumer read another slice's bytes (Qwen3.5 BA: TT computed the `a` + // projection with the `b` weight rows — BACKEND-TENSTORRENT-QWEN35 W2c). + bool tracked_base = false; + bool base_needs_host_refresh = false; { + std::lock_guard g(SlotMutex()); + BufferSlot* s = FindSlot(t.data); + tracked_base = (s != nullptr && t.data == s->host); + if (!tracked_base && s != nullptr) + base_needs_host_refresh = s->device_current && !s->host_current; + } + if (base_needs_host_refresh) EnsureHostBytes(t.data); + if (tracked_base) { std::lock_guard g(SlotMutex()); BufferSlot* s = FindSlot(t.data); if (s != nullptr && s->device_current && s->device.has_value()) { @@ -411,10 +472,28 @@ ttnn::Tensor EnsureDevice2D(const Tensor& t, MeshDevice& device) { } } } - // Need host truth to upload (may download first if only device was current - // under a different shape — rare). - EnsureHost(t); - const auto host = ToHostF32(t); + // Host truth for the WINDOW: read the view's own bytes directly. For an + // interior view this is valid because uploads keep the base's host master + // current (store above sets host_current=true); if some future path makes + // the base device-only, refuse loudly rather than serve stale bytes. + { + std::lock_guard g(SlotMutex()); + BufferSlot* s = FindSlot(t.data); + VT_CHECK(!(s != nullptr && !tracked_base && !s->host_current), + "tenstorrent: EnsureDevice2D interior view of a device-current " + "base is unsupported (would read stale bytes); stage via the " + "base tensor"); + } + // HOST-FREE-FORWARD defect: this loop consumed HOST bytes without checking + // slot residency. Under host-free decode a producer commits device-only + // (host_current=false), so the next consumer staging through here uploaded + // pool-fresh zeros and then marked the slot host_current=true — corrupting + // both the value and the record. Refresh the host bytes from the device + // shadow before reading them. + EnsureHostBytes(t.data); + std::vector host(static_cast(rows) * static_cast(cols)); + for (int64_t i = 0; i < t.Numel(); ++i) + host[static_cast(i)] = LoadElemF32(t, i); ttnn::Tensor dev = UploadRows(host.data(), rows, cols, device); if (HostFreeDecodeEnabled()) { // Prime the persistent-zero cache for this spec during the eager warmup @@ -424,7 +503,7 @@ ttnn::Tensor EnsureDevice2D(const Tensor& t, MeshDevice& device) { } std::lock_guard g(SlotMutex()); BufferSlot* s = FindSlot(t.data); - if (s != nullptr) { + if (s != nullptr && t.data == s->host) { s->device = dev; s->dev_rows = rows; s->dev_cols = cols; @@ -434,6 +513,10 @@ ttnn::Tensor EnsureDevice2D(const Tensor& t, MeshDevice& device) { return dev; } +// DEBUG (BACKEND-TENSTORRENT-QWEN35 W2c): ensure the tensor is staged on +// device exactly as a consuming kernel would, then read the DEVICE copy back. +// Comparing this against the host master exposes staging corruption that a +// host-side dump cannot see. // True when `t` already has a device-resident shadow matching [rows, cols] // (exact shape). Used to pick device vs host kernels without forcing upload. bool DeviceShadowExact(const Tensor& t, uint32_t rows, uint32_t cols) { @@ -935,11 +1018,30 @@ void CommitDeviceLogical2D(Tensor& out, ttnn::Tensor dev, uint32_t rows, uint32_ VT_CHECK(out.IsContiguous(), "tenstorrent: CommitDeviceLogical2D expects contiguous out"); VT_CHECK(out.Numel() == static_cast(rows) * static_cast(cols), "tenstorrent: CommitDeviceLogical2D numel mismatch"); + { + int64_t vol = 1; + const auto ds = dev.logical_shape(); + for (uint32_t i = 0; i < ds.rank(); ++i) vol *= ds[i]; + VT_CHECK(vol == static_cast(rows) * static_cast(cols), + std::string("tenstorrent: CommitDeviceLogical2D device volume ") + + std::to_string(vol) + " != rows*cols " + + std::to_string(rows * cols)); + } std::lock_guard g(SlotMutex()); BufferSlot* s = FindSlot(out.data); + if (std::getenv("VT_TT_SLOT_TRACE") != nullptr) { + const auto ds = dev.logical_shape(); + int64_t vol = 1; + for (uint32_t i = 0; i < ds.rank(); ++i) vol *= ds[i]; + std::fprintf(stderr, + "[TT-SLOT] commit out=%p rows=%u cols=%u tracked=%d vol=%" PRId64 + " dt=%d\n", + static_cast(out.data), rows, cols, + s != nullptr ? 1 : 0, vol, static_cast(dev.dtype())); + } if (s == nullptr) { // Untracked buffer (e.g. stack/test scratch): fall back to host write. - DownloadToHost(dev, out); + DownloadToHost(dev, out, "CommitDeviceLogical2D(untracked)"); return; } s->device = std::move(dev); @@ -1214,8 +1316,11 @@ ttnn::Tensor EnsureAffine1D(const Tensor& t, uint32_t d, MeshDevice& device) { { std::lock_guard g(SlotMutex()); BufferSlot* s = FindSlot(t.data); - if (s != nullptr && s->device_current && s->device.has_value() && s->dev_rows == 1 && - s->dev_cols == d) { + // Hit only for the BASE pointer (same interior-view hazard as + // EnsureDevice2D — a differently-offset rank-1 view must not consume the + // base's staged affine). + if (s != nullptr && t.data == s->host && s->device_current && + s->device.has_value() && s->dev_rows == 1 && s->dev_cols == d) { return *s->device; } } @@ -1229,7 +1334,7 @@ ttnn::Tensor EnsureAffine1D(const Tensor& t, uint32_t d, MeshDevice& device) { &device); std::lock_guard g(SlotMutex()); BufferSlot* s = FindSlot(t.data); - if (s != nullptr) { + if (s != nullptr && t.data == s->host) { s->device = dev; s->dev_rows = 1; s->dev_cols = d; @@ -1356,7 +1461,27 @@ void RmsNormKernel(Queue&, Tensor& out, const Tensor& x, const Tensor& weight, // rms) when rows are large enough that launches amortize. MeshDevice& device = SharedMeshDevice(); ttnn::Tensor dev_x = EnsureDevice2D(x, device); - ttnn::Tensor dev_w = EnsureAffine1D(weight, d, device); + // Gemma (w+1) is baked host-side in the oracle's f32 order — the same + // treatment as the fused preamble's weff — because ttnn::rms_norm applies + // gamma raw. The baked upload is TRANSIENT and deliberately bypasses the + // EnsureAffine1D slot cache: the cached form is the RAW weight, and a + // non-gemma consumer of the same buffer must never read the +1 version. + // Qwen3 never sets gemma; Qwen3.5 sets it on every norm (BACKEND-TENSTORRENT- + // QWEN35 W2b: dropping the +1 here collapsed ambient prefill to `,`). + ttnn::Tensor dev_w; + if (args.gemma) { + EnsureHost(weight); + std::vector gw(static_cast(d)); + for (uint32_t i = 0; i < d; ++i) + gw[static_cast(i)] = LoadElemF32(weight, static_cast(i)) + 1.0f; + dev_w = ttnn::Tensor::from_vector( + std::move(gw), + SpecOf(tt::tt_metal::Shape({1, d}), ttnn::DataType::FLOAT32, + ttnn::Layout::TILE), + &device); + } else { + dev_w = EnsureAffine1D(weight, d, device); + } ttnn::Tensor to_norm = dev_x; if (residual != nullptr) { ttnn::Tensor dev_r = EnsureDevice2D(*residual, device); @@ -1475,6 +1600,402 @@ void CastF32Kernel(Queue&, Tensor& out, const Tensor& in) { CommitHost(out); } +// kSigmoidGateBf16: out[i] = F32ToBF16(attn[i] * sigmoid(gate[i])) — the +// full-attention o_proj gate (cpu_ops.cpp SigmoidGateBf16Kernel; wrapper +// ops.cpp:4136-4153). The gate must NOT be rounded before the sigmoid +// (ops.cpp:4140: "sigmoid input must not be rounded"), so both operands ride +// FLOAT32 tile shadows and the product runs in f32; the single RNE round to +// bf16 happens in ttnn::typecast (the device fp32->fp16b cast: +0x7FFF+lsb +// then mask, ckernel_sfpu_typecast.h:246-262 — bit-identical to F32ToBF16; +// ttnn::to_dtype is HOST-ONLY at this pin, tensor_ops.cpp:533). The +// only TT-vs-CPU delta is the SFPU f32 sigmoid (accurate exp + +// reciprocal_iter<2> vs std::exp) — a few f32 ULP that can flip at most one +// bf16 rounding; the doctest envelope is one bf16 ULP. +ttnn::Tensor UploadTensor(std::vector host, const ttnn::Shape& shape, + ttnn::DataType dtype, ttnn::Layout layout, + MeshDevice& device); // defined below (GDN uploads) +ttnn::Tensor DeviceRows(const Tensor& t, uint32_t rows, uint32_t cols, + MeshDevice& device); // defined below (shadow ensure) +void SigmoidGateBf16Kernel(Queue&, Tensor& out, const Tensor& attn, + const Tensor& gate) { + TT_OP_TRACE("SigmoidGateBf16"); + VT_CHECK(out.dtype == DType::kBF16, + "tenstorrent kSigmoidGateBf16: out must be bf16"); + VT_CHECK((attn.dtype == DType::kF32 || attn.dtype == DType::kBF16) && + gate.dtype == DType::kF32, + "tenstorrent kSigmoidGateBf16: attn must be f32/bf16, gate f32"); + VT_CHECK(out.Numel() == attn.Numel() && out.Numel() == gate.Numel(), + "tenstorrent kSigmoidGateBf16: out/attn/gate same element count"); + VT_CHECK(out.IsContiguous() && attn.IsContiguous() && gate.IsContiguous(), + "tenstorrent kSigmoidGateBf16: contiguous required"); + const uint32_t n = static_cast(out.Numel()); + MeshDevice& device = SharedMeshDevice(); + // UploadTensor FLOAT32 (not the bf16 EnsureDevice2D shadow): bf16 attn + // upcasts exactly, and f32 attn keeps its full mantissa — a bf16 shadow + // would pre-round the attn operand and widen the envelope (the f32 arm of + // the doctest exists to catch exactly that). + ttnn::Tensor dev_attn = + UploadTensor(ToHostF32(attn), ttnn::Shape({1, n}), ttnn::DataType::FLOAT32, + ttnn::Layout::TILE, device); + ttnn::Tensor dev_gate = + UploadTensor(ToHostF32(gate), ttnn::Shape({1, n}), ttnn::DataType::FLOAT32, + ttnn::Layout::TILE, device); + ttnn::Tensor sig = ttnn::sigmoid(dev_gate); + ttnn::Tensor prod = ttnn::multiply(dev_attn, sig); + ttnn::Tensor dev_y = ttnn::typecast(prod, ttnn::DataType::BFLOAT16); + CommitDeviceLogical2D(out, std::move(dev_y), 1, n); +} + +// kGdnPostConv: fused GDN post-conv prep = GdnConvSplit + per-head L2Norm(q/k) +// + GdnGBeta in one pass (cpu_ops.cpp GdnPostConvKernel:3472-3516; wrapper +// ops.cpp:4255-4287; the fla fused_recurrent_gated_delta_rule prefill preamble). +// q/k: conv slices [0,key_dim)/[key_dim,2*key_dim) reshape to [T*Hk, Dk] rows +// and run the kL2Norm device math (square -> sum -> +eps -> rsqrt -> mul) in +// bf16 tiles — the row's calibrated l2norm envelope. v: plain slice copy. +// g/beta: FLOAT32 tiles end to end (softplus threshold-20, exp(a_log) and +// sigmoid are f32 in the oracle); araw/braw are row-strided views, gathered +// on host exactly like the kRmsNormGated gate view. +void GdnPostConvKernel(Queue& q, Tensor& q_out, Tensor& k_out, Tensor& v_out, + Tensor& g_out, Tensor& beta_out, const Tensor& conv, + const Tensor& araw, const Tensor& braw, + const Tensor& a_log, const Tensor& dt_bias, + const L2NormArgs& args) { + TT_OP_TRACE("GdnPostConv"); + VT_CHECK((q_out.dtype == DType::kF32 || q_out.dtype == DType::kBF16) && + k_out.dtype == q_out.dtype && v_out.dtype == q_out.dtype, + "tenstorrent kGdnPostConv: q_out/k_out/v_out must be f32 or bf16, same dtype"); + VT_CHECK(conv.dtype == DType::kF32 || conv.dtype == DType::kBF16, + "tenstorrent kGdnPostConv: conv must be f32 or bf16"); + VT_CHECK(g_out.dtype == DType::kF32 && beta_out.dtype == DType::kF32 && + (araw.dtype == DType::kF32 || araw.dtype == DType::kBF16) && + braw.dtype == araw.dtype && a_log.dtype == DType::kF32 && + dt_bias.dtype == DType::kF32, + "tenstorrent kGdnPostConv: g/beta/a_log/dt_bias f32; a/b share f32 or bf16"); + VT_CHECK(q_out.IsContiguous() && k_out.IsContiguous() && v_out.IsContiguous() && + g_out.IsContiguous() && beta_out.IsContiguous() && + conv.IsContiguous() && araw.stride[1] == 1 && braw.stride[1] == 1 && + a_log.IsContiguous() && dt_bias.IsContiguous(), + "tenstorrent kGdnPostConv: contiguous required (a/b row views excepted)"); + const uint32_t t = static_cast(conv.shape[0]); + const uint32_t hk = static_cast(q_out.shape[1]); + const uint32_t dk = static_cast(q_out.shape[2]); + const uint32_t hv = static_cast(v_out.shape[1]); + const uint32_t dv = static_cast(v_out.shape[2]); + const uint32_t key_dim = hk * dk, value_dim = hv * dv; + const uint32_t conv_dim = 2 * key_dim + value_dim; + VT_CHECK(conv.shape[0] == q_out.shape[0] && conv.shape[1] == conv_dim && + g_out.shape[0] == t && g_out.shape[1] == hv, + "tenstorrent kGdnPostConv: shape mismatch"); + + MeshDevice& device = SharedMeshDevice(); + ttnn::Tensor dev_conv = DeviceRows(conv, t, conv_dim, device); + ttnn::Tensor q2 = ttnn::slice(dev_conv, ttsl::SmallVector{0, 0}, + ttsl::SmallVector{t, key_dim}, + ttsl::SmallVector{1, 1}); + ttnn::Tensor k2 = ttnn::slice(dev_conv, ttsl::SmallVector{0, key_dim}, + ttsl::SmallVector{t, 2 * key_dim}, + ttsl::SmallVector{1, 1}); + ttnn::Tensor v2 = ttnn::slice(dev_conv, ttsl::SmallVector{0, 2 * key_dim}, + ttsl::SmallVector{t, conv_dim}, + ttsl::SmallVector{1, 1}); + // Heads are laid contiguously along the sliced cols, so [t, hk*dk] -> + // [t*hk, dk] is a pure logical re-view; each row is one head's Dk vector. + auto l2 = [&](const ttnn::Tensor& cols) { + ttnn::Tensor rows = + ttnn::reshape(cols, ttnn::Shape({t * hk, dk})); + ttnn::Tensor sq = ttnn::multiply(rows, rows); + ttnn::Tensor s = ttnn::sum(sq, ttsl::SmallVector{1}, true); + ttnn::Tensor denom = ttnn::add(s, args.eps); + ttnn::Tensor inv = ttnn::rsqrt(denom); + return ttnn::multiply(rows, inv); + }; + CommitDeviceLogical2D(q_out, l2(q2), t * hk, dk); + CommitDeviceLogical2D(k_out, l2(k2), t * hk, dk); + CommitDeviceLogical2D(v_out, ttnn::reshape(v2, ttnn::Shape({t * hv, dv})), t * hv, dv); + + // g/beta in f32 (the row's "f32 intermediates" doctrine: softplus(x) with + // threshold 20, exp(a_log) and sigmoid must not round their inputs). + EnsureHost(araw); + EnsureHost(braw); + EnsureHost(a_log); + EnsureHost(dt_bias); + if (const char* td = std::getenv("VT_DUMP_TRUST")) { + static std::atomic gb_seq{0}; + if (gb_seq.fetch_add(1, std::memory_order_relaxed) == 0) { + // Raw HOST bytes of the a/b windows exactly as LoadElemF32 sees them, + // after EnsureHost. Compare against ba_a_win_dev/ba_b_win_dev. + std::FILE* fp = std::fopen( + (std::string(td) + "/0_ab_host_raw.f32").c_str(), "wb"); + if (fp) { + const float* ap = static_cast( + static_cast(araw.data)); + const float* bp = static_cast( + static_cast(braw.data)); + for (uint32_t i = 0; i < t; ++i) { + std::fwrite(ap + i * araw.stride[0], 4, hv, fp); + std::fwrite(bp + i * braw.stride[0], 4, hv, fp); + } + std::fclose(fp); + } + } + } + std::vector a(static_cast(t) * hv), b(a.size()), + al(hv), dt(hv); + for (uint32_t i = 0; i < t; ++i) + for (uint32_t h = 0; h < hv; ++h) { + const int64_t aidx = i * araw.stride[0] + h; + const int64_t bidx = i * braw.stride[0] + h; + a[static_cast(i) * hv + h] = LoadElemF32(araw, aidx); + b[static_cast(i) * hv + h] = LoadElemF32(braw, bidx); + } + for (uint32_t h = 0; h < hv; ++h) { + al[h] = LoadElemF32(a_log, h); + dt[h] = LoadElemF32(dt_bias, h); + } + ttnn::Tensor dev_a = + UploadTensor(std::move(a), ttnn::Shape({t, hv}), ttnn::DataType::FLOAT32, + ttnn::Layout::TILE, device); + ttnn::Tensor dev_b = + UploadTensor(std::move(b), ttnn::Shape({t, hv}), ttnn::DataType::FLOAT32, + ttnn::Layout::TILE, device); + ttnn::Tensor dev_al = UploadTensor(std::move(al), ttnn::Shape({1, hv}), + ttnn::DataType::FLOAT32, ttnn::Layout::TILE, + device); + ttnn::Tensor dev_dt = UploadTensor(std::move(dt), ttnn::Shape({1, hv}), + ttnn::DataType::FLOAT32, ttnn::Layout::TILE, + device); + // x = araw + dt_bias (row-vector broadcast); sp = relu(x) + + // log1p(exp(-|x|)) — the oracle's log1p(exp(x)) written in a form that is + // stable for every x (the exp argument is always <= 0). For x > 20 the + // log1p term sits below x's half-ULP, so the sum rounds to x EXACTLY and + // reproduces the oracle's threshold branch. Used instead of the one-shot + // SFPU softplus poly, whose ~1e-6 ABSOLUTE fit error becomes ~1e-3 + // RELATIVE on small softplus values (measured on the P150, first green + // run: g max_rel 4.4e-4 over |g|~0.04). g = -exp(a_log)*sp. + ttnn::Tensor x = ttnn::add(dev_a, dev_dt); + ttnn::Tensor sp = ttnn::add( + ttnn::relu(x), ttnn::log1p(ttnn::exp(ttnn::neg(ttnn::abs(x))))); + ttnn::Tensor neg_ea = ttnn::multiply(ttnn::exp(dev_al), -1.0f); + ttnn::Tensor g = ttnn::multiply(neg_ea, sp); + ttnn::Tensor beta = ttnn::sigmoid(dev_b); + CommitDeviceLogical2D(g_out, std::move(g), t, hv); + CommitDeviceLogical2D(beta_out, std::move(beta), t, hv); + if (const char* td = std::getenv("VT_DUMP_TRUST")) { + // Kernel-side trusted captures at the COMMIT SITE: same tensor objects, + // same TU — removes every cross-function identity assumption. + static std::atomic gpc_seq{0}; + const int call = gpc_seq.fetch_add(1, std::memory_order_relaxed); + if (call == 0) { + TrustDump(q, td, "k_conv_in", conv); + TrustDump(q, td, "k_q", q_out); + TrustDump(q, td, "k_k", k_out); + TrustDump(q, td, "k_v", v_out); + TrustDump(q, td, "k_g", g_out); + TrustDump(q, td, "k_beta", beta_out); + // Intermediates of the g chain (device truth via to_vector). + auto tdv = [&](const char* n, const ttnn::Tensor& t) { + auto v = t.to_vector(); + std::FILE* f = std::fopen( + (std::string(td) + "/" + std::to_string(call) + "_" + n + ".f32") + .c_str(), "wb"); + if (f) { std::fwrite(v.data(), 4, v.size(), f); std::fclose(f); } + }; + tdv("k_dev_a", dev_a); + tdv("k_dev_b", dev_b); + tdv("k_dev_al", dev_al); + tdv("k_dev_dt", dev_dt); + tdv("k_x", x); + tdv("k_sp", sp); + tdv("k_neg_ea", neg_ea); + } + } +} + +// kAttnQkNormRopeGate: fused full-attention preamble = split q|gate + +// per-head gemma qk-RMSNorm + partial NeoX RoPE-from-cos_sin + exact gate +// passthrough, in ONE launch (cpu_ops.cpp AttnQkNormRopeGateKernel:1216-1270; +// wrapper ops.cpp:1638-1689; production call qwen3_5.cpp:5206-5224, gemma=true, +// rot < Dh). qgate/kf arrive as merged-QKV strided views (rows +// inner-contiguous); the split into per-head [T*H, Dh] rows happens in the +// HOST gather, and each leg is uploaded already in its final logical shape +// (no device slice/reshape: measured on the P150, first green attempt, a +// column slice + ttnn::reshape chain returned wrong data for the qgate legs +// while the full-width kf leg stayed correct — only a fresh contiguous +// upload is a safe reshape input at this pin). Norm+rope ride FLOAT32 tiles +// end to end (the SigmoidGateBf16 doctrine): rsqrt, the weight mix and the +// cos/sin rotation never round before the single output typecast, so an f32 +// out sits within reduction-order ULPs of the scalar-f32 oracle and a bf16 +// out is its exact RNE round. The gate leg is a plain device copy — no +// arithmetic touches it, so the f32 passthrough is bit-exact (the sigmoid +// input must not be rounded, ops.cpp:1660-1662/4140). The gemma weight w+1 +// is computed host-side in f32, the same add GemmaNormElem does +// (cpu_ops.cpp:1204-1208). +void AttnQkNormRopeGateKernel(Queue&, Tensor& q_out, Tensor& k_out, + Tensor& gate_out, const Tensor& qgate, + const Tensor& kf, const Tensor& q_norm, + const Tensor& k_norm, const Tensor& cos_sin, + const RmsNormArgs& na, const RopeArgs& ra) { + TT_OP_TRACE("AttnQkNormRopeGate"); + VT_CHECK((q_out.dtype == DType::kF32 || q_out.dtype == DType::kBF16) && + k_out.dtype == q_out.dtype && + (gate_out.dtype == q_out.dtype || + (q_out.dtype == DType::kBF16 && gate_out.dtype == DType::kF32)), + "tenstorrent kAttnQkNormRopeGate: q/k/gate out f32 or bf16 " + "(gate f32 allowed with bf16 q/k)"); + VT_CHECK(IsFloatDType(qgate.dtype) && kf.dtype == qgate.dtype, + "tenstorrent kAttnQkNormRopeGate: qgate/kf float, same dtype"); + VT_CHECK(q_out.IsContiguous() && k_out.IsContiguous() && + gate_out.IsContiguous() && qgate.stride[1] == 1 && + qgate.stride[0] >= qgate.shape[1] && kf.stride[1] == 1 && + kf.stride[0] >= kf.shape[1] && q_norm.IsContiguous() && + k_norm.IsContiguous() && cos_sin.IsContiguous(), + "tenstorrent kAttnQkNormRopeGate: contiguous required " + "(qgate/kf row views excepted)"); + const int64_t t = q_out.shape[0], hq = q_out.shape[1], dh = q_out.shape[2]; + const int64_t hkv = k_out.shape[1]; + const int64_t rot = ra.rotary_dim, half = rot / 2; + const int64_t qrow = qgate.shape[1], krow = kf.shape[1]; + VT_CHECK(qrow == hq * 2 * dh && krow == hkv * dh, + "tenstorrent kAttnQkNormRopeGate: qgate [T, Hq*2*Dh], kf [T, Hkv*Dh]"); + + MeshDevice& device = SharedMeshDevice(); + // Host-gather the strided merged-QKV rows straight into per-head [T*H, Dh] + // legs (the q|gate split is the host half of this fused op). The FLOAT32 + // upload also upcasts a bf16 input exactly (LoadElemF32), so every later + // leg is dtype-uniform. + EnsureHost(qgate); + EnsureHost(kf); + std::vector qh(static_cast(t * hq * dh)), + gh(qh.size()), kh(static_cast(t * hkv * dh)); + for (int64_t i = 0; i < t; ++i) { + for (int64_t h = 0; h < hq; ++h) { + const int64_t base = i * qgate.stride[0] + h * 2 * dh; + const size_t dst = static_cast((i * hq + h) * dh); + for (int64_t j = 0; j < dh; ++j) { + qh[dst + j] = LoadElemF32(qgate, base + j); + gh[dst + j] = LoadElemF32(qgate, base + dh + j); + } + } + for (int64_t h = 0; h < hkv; ++h) { + const int64_t base = i * kf.stride[0] + h * dh; + const size_t dst = static_cast((i * hkv + h) * dh); + for (int64_t j = 0; j < dh; ++j) + kh[dst + j] = LoadElemF32(kf, base + j); + } + } + ttnn::Tensor dev_q = + UploadTensor(std::move(qh), + ttnn::Shape({static_cast(t * hq), + static_cast(dh)}), + ttnn::DataType::FLOAT32, ttnn::Layout::TILE, device); + ttnn::Tensor dev_gate = + UploadTensor(std::move(gh), + ttnn::Shape({static_cast(t * hq), + static_cast(dh)}), + ttnn::DataType::FLOAT32, ttnn::Layout::TILE, device); + ttnn::Tensor dev_k = + UploadTensor(std::move(kh), + ttnn::Shape({static_cast(t * hkv), + static_cast(dh)}), + ttnn::DataType::FLOAT32, ttnn::Layout::TILE, device); + // Gemma effective weight (host f32 add, oracle-identical) and per-head-row + // cos|sin expansion: row (tok, head) copies token tok's cos_sin halves. + EnsureHost(q_norm); + EnsureHost(k_norm); + auto weff = [&](const Tensor& w) { + std::vector v(static_cast(dh)); + for (int64_t j = 0; j < dh; ++j) { + const float wj = LoadElemF32(w, j); + v[static_cast(j)] = na.gemma ? wj + 1.0f : wj; + } + return v; + }; + EnsureHost(cos_sin); + auto rope_cs = [&](int64_t heads, std::vector& cexp, + std::vector& sexp) { + cexp.resize(static_cast(t * heads * half)); + sexp.resize(cexp.size()); + for (int64_t i = 0; i < t; ++i) + for (int64_t h = 0; h < heads; ++h) { + const size_t dst = static_cast((i * heads + h) * half); + for (int64_t j = 0; j < half; ++j) { + cexp[dst + j] = LoadElemF32(cos_sin, i * rot + j); + sexp[dst + j] = LoadElemF32(cos_sin, i * rot + half + j); + } + } + }; + // normed = x * rsqrt(mean(x^2)+eps) * (w+gemma); NeoX half-split rotation of + // the leading rot cols; tail [rot, Dh) passes through NORMED (not rotated). + auto norm_rope = [&](ttnn::Tensor x, std::vector w, + std::vector cexp, std::vector sexp, + int64_t nrows) { + const uint32_t nr = static_cast(nrows); + const uint32_t du = static_cast(dh); + const uint32_t halfu = static_cast(half); + const uint32_t rotu = static_cast(rot); + ttnn::Tensor dev_w = + UploadTensor(std::move(w), ttnn::Shape({1, du}), + ttnn::DataType::FLOAT32, ttnn::Layout::TILE, device); + ttnn::Tensor dev_cos = + UploadTensor(std::move(cexp), ttnn::Shape({nr, halfu}), + ttnn::DataType::FLOAT32, ttnn::Layout::TILE, device); + ttnn::Tensor dev_sin = + UploadTensor(std::move(sexp), ttnn::Shape({nr, halfu}), + ttnn::DataType::FLOAT32, ttnn::Layout::TILE, device); + ttnn::Tensor sq = ttnn::multiply(x, x); + ttnn::Tensor s = ttnn::sum(sq, ttsl::SmallVector{1}, true); + ttnn::Tensor denom = + ttnn::add(ttnn::multiply(s, 1.0f / static_cast(dh)), na.eps); + ttnn::Tensor inv = ttnn::rsqrt(denom); + ttnn::Tensor normed = ttnn::multiply(ttnn::multiply(x, inv), dev_w); + ttnn::Tensor x1 = ttnn::slice(normed, ttsl::SmallVector{0, 0}, + ttsl::SmallVector{nr, halfu}, + ttsl::SmallVector{1, 1}); + ttnn::Tensor x2 = ttnn::slice(normed, + ttsl::SmallVector{0, halfu}, + ttsl::SmallVector{nr, rotu}, + ttsl::SmallVector{1, 1}); + ttnn::Tensor o1 = ttnn::subtract(ttnn::multiply(x1, dev_cos), + ttnn::multiply(x2, dev_sin)); + ttnn::Tensor o2 = ttnn::add(ttnn::multiply(x1, dev_sin), + ttnn::multiply(x2, dev_cos)); + ttnn::Tensor out = + ttnn::concat(std::vector{o1, o2}, /*dim=*/1); + if (rot < dh) { + ttnn::Tensor tail = + ttnn::slice(normed, ttsl::SmallVector{0, rotu}, + ttsl::SmallVector{nr, du}, + ttsl::SmallVector{1, 1}); + out = ttnn::concat(std::vector{out, tail}, /*dim=*/1); + } + return out; + }; + // Gate leg FIRST (plain device copy of its own upload — no math, so the + // passthrough stays bit-exact), then the q/k norm+rope legs. + if (gate_out.dtype == DType::kBF16) + dev_gate = ttnn::typecast(dev_gate, ttnn::DataType::BFLOAT16); + CommitDeviceLogical2D(gate_out, std::move(dev_gate), + static_cast(t * hq), + static_cast(dh)); + std::vector qc, qs, kc, ks; + rope_cs(hq, qc, qs); + ttnn::Tensor q_dev = norm_rope(std::move(dev_q), weff(q_norm), + std::move(qc), std::move(qs), t * hq); + if (q_out.dtype == DType::kBF16) + q_dev = ttnn::typecast(q_dev, ttnn::DataType::BFLOAT16); + CommitDeviceLogical2D(q_out, std::move(q_dev), static_cast(t * hq), + static_cast(dh)); + rope_cs(hkv, kc, ks); + ttnn::Tensor k_dev = + norm_rope(std::move(dev_k), weff(k_norm), std::move(kc), std::move(ks), + t * hkv); + if (k_out.dtype == DType::kBF16) + k_dev = ttnn::typecast(k_dev, ttnn::DataType::BFLOAT16); + CommitDeviceLogical2D(k_out, std::move(k_dev), + static_cast(t * hkv), + static_cast(dh)); +} + // Llama-3 frequency rescale (cpu_ops Llama3ScaleFreq); no-op when scaling_factor // is unset. Kept so Qwen3 / Llama rope paths share one host implementation. inline double Llama3ScaleFreq(double freq, const RopeArgs& a) { @@ -1673,6 +2194,7 @@ void RopeApplyHost(Tensor& qs, Tensor* ks, const float* cos_t, const float* sin_ // Prefer device apply only when T*H amortizes the slice/mul/concat launches. // Short Qwen3 decode (T=1,H=16) is host-faster even when Q is already on device // (measured regression when always-device-for-resident was forced). + inline bool PreferDeviceRope(int64_t tokens, int64_t heads) { // HOST-FREE-FORWARD R1: force device RoPE at T=1 for capture (see RmsNorm note). if (HostFreeDecodeEnabled()) return true; @@ -3598,6 +4120,32 @@ std::atomic& GdnDecodeSteps() { // through exact 0/1 one-hot matmuls (0*x + v == v in f32), so the full cache // never moves either. +// The shadow-row split factor for a wide cache: indexed_fill's generic +// interleaved path stages TWO full pages of the LAST dim through its +// dataflow buffer (indexed_fill_program_factory.cpp: page_size = +// padded_shape[-1] * elem_size, data DFB num_entries = 2), so a row wider +// than kStageBytes/(2*esz) elems overflows L1 — the Qwen3.5 GDN ssm_state +// row Hv*Dk*Dv = 16*128*128 = 262144 f32 elems stages 2 x 1 MB and threw +// exactly that ("grow to 2208704 B beyond max L1 size of 1572864 B", first +// W2a e2e bootstrap, /tmp/w2a_e2e_bootstrap.log). The fix is NOT a +// last-dim-changing view: tt::tt_metal::view moves no bytes, but device +// pages are the INTERLEAVE UNIT (buffer.cpp Buffer::page_address: +// bank_offset = aligned_page_size * (page_index / num_banks)) — reinterpreting +// [rows, cols] as [rows*nb, cb] changes every page index's bank mapping and +// scrambles the flat order (measured in W2a: one block landed correctly, +// three came back scrambled). Instead the SHADOW IS BORN SPLIT: +// EnsureGdnCacheDevice uploads the SAME flat host bytes as +// [rows*F, cols/F], so every later op sees narrow rows natively and the +// fill stays ONE launch inside the budget. F is the smallest divisor of +// cols with 2*(cols/F)*esz <= 1 MiB; F == 1 (cols at or under the budget — +// every small-cache model) keeps the W1-verified form byte-identical. +int64_t SplitFactor(int64_t cols, int64_t esz) { + constexpr int64_t kStageBytes = 1 << 20; + int64_t f = std::max(1, (2 * cols * esz + kStageBytes - 1) / kStageBytes); + while (cols % f != 0) ++f; + return f; +} + // Ensure the f32 device shadow of a GDN state/conv CACHE tensor, viewed as // [rows, cols]. Uploads (counted) only when the shadow is missing, stale // (host wrote), or the wrong shape/dtype. The device tensor is FLOAT32 TILE @@ -3606,6 +4154,12 @@ std::atomic& GdnDecodeSteps() { ttnn::Tensor EnsureGdnCacheDevice(const Tensor& t, int64_t rows, int64_t cols, MeshDevice& device, ttnn::Layout layout = ttnn::Layout::TILE) { + // The shadow is stored SPLIT: [rows*F, cols/F] with F = SplitFactor(cols) + // (see the helper above — wide rows must not become indexed_fill pages). + // dev_rows/dev_cols stay LOGICAL; the flat bytes are identical either way, + // so every download and every volume check is unaffected. + const uint32_t sf = static_cast(SplitFactor(cols, /*esz=*/4)); + const uint32_t ublk = static_cast(cols / sf); { std::lock_guard g(SlotMutex()); BufferSlot* s = FindSlot(t.data); @@ -3615,9 +4169,12 @@ ttnn::Tensor EnsureGdnCacheDevice(const Tensor& t, int64_t rows, int64_t cols, static_cast(s->dev_rows) * s->dev_cols == static_cast(rows) * cols) { if (s->dev_rows == rows && s->dev_cols == cols) return *s->device; + // Same buffer re-served at new logical dims: reshape to the NEW dims' + // split geometry (equal volume — a pure data-movement program whose + // circular buffers are per-tile). ttnn::Tensor reshaped = - ttnn::reshape(*s->device, ttnn::Shape({static_cast(rows), - static_cast(cols)})); + ttnn::reshape(*s->device, ttnn::Shape({static_cast(rows) * sf, + ublk})); s->device = reshaped; s->dev_rows = static_cast(rows); s->dev_cols = static_cast(cols); @@ -3630,10 +4187,10 @@ ttnn::Tensor EnsureGdnCacheDevice(const Tensor& t, int64_t rows, int64_t cols, for (int64_t i = 0; i < n; ++i) host[static_cast(i)] = LoadElemF32(t, i); GdnStateH2dBytes().fetch_add(static_cast(n) * sizeof(float), std::memory_order_relaxed); - ttnn::Tensor dev = UploadTensor( - std::move(host), - ttnn::Shape({static_cast(rows), static_cast(cols)}), - ttnn::DataType::FLOAT32, layout, device); + ttnn::Tensor dev = UploadTensor(std::move(host), + ttnn::Shape({static_cast(rows) * sf, + ublk}), + ttnn::DataType::FLOAT32, layout, device); std::lock_guard g(SlotMutex()); BufferSlot* s = FindSlot(t.data); if (s != nullptr) { @@ -3702,15 +4259,22 @@ ttnn::Tensor GatherRowsExact(const ttnn::Tensor& cache2d, return lay == ttnn::Layout::TILE ? out : ttnn::to_layout(out, lay); } -// EXACT scatter of rows2d [rows, cols] into cache2d [slots, cols] via -// indexed_fill (torch.index_copy_ semantics): pure data movement, unnamed -// slots keep their bytes, and a slot named by SEVERAL rows keeps the LAST -// row — the CPU oracle's sequential loop order. idx<0 = NULL row: writes -// nothing (its row is compacted away before the fill). +// EXACT scatter of rows2d [rows*factor, cols/factor] into cache2d +// [slots*factor, cols/factor] (both shadows are born split by SplitFactor; +// the LOGICAL geometry both kernels serve is [rows, cols]) via indexed_fill +// (torch.index_copy_ semantics): pure data movement, unnamed slots keep +// their bytes, and a slot named by SEVERAL rows keeps the LAST row — the +// CPU oracle's sequential loop order. idx<0 = NULL row: writes nothing (its +// row is compacted away before the fill). Each live slot s expands into its +// `factor` block ids s*F+j, so source block-row (e, j) lands in slot +// idx[e]'s j-th block; relative entry order is preserved inside every block +// group, so per-(slot, block) last-of-duplicates wins exactly like the +// whole-row form. ttnn::Tensor ScatterRowsExact(const ttnn::Tensor& cache2d, const std::vector& idx, const ttnn::Tensor& rows2d, int64_t slots, - int64_t cols, MeshDevice& device) { + int64_t cols, int64_t factor, MeshDevice& device) { + const uint32_t blk = static_cast(cols / factor); std::vector bid; bid.reserve(idx.size()); for (int32_t ix : idx) @@ -3721,25 +4285,56 @@ ttnn::Tensor ScatterRowsExact(const ttnn::Tensor& cache2d, src = ttnn::to_layout(rows2d, cache2d.layout()); // pure copy, f32 exact if (bid.size() != idx.size()) { // Compact the live rows through the same exact gather: row_of[e] is the - // ORIGINAL row of the e-th live entry. + // ORIGINAL row of the e-th live entry — expanded to its `factor` + // block-rows, which sit contiguous and in order. std::vector row_of; row_of.reserve(bid.size()); for (int64_t r = 0; r < static_cast(idx.size()); ++r) if (idx[static_cast(r)] >= 0) row_of.push_back(static_cast(r)); - src = GatherRowsExact(src, row_of, cols, device); - } - const uint32_t un = static_cast(bid.size()); + std::vector row_of_exp; + row_of_exp.reserve(row_of.size() * static_cast(factor)); + for (int32_t r : row_of) + for (int64_t j = 0; j < factor; ++j) + row_of_exp.push_back(static_cast(r * factor + j)); + src = GatherRowsExact(src, row_of_exp, blk, device); + } + // Expand each live slot into its `factor` block ids before the upload. + std::vector blocks; + blocks.reserve(bid.size() * static_cast(factor)); + for (uint32_t s : bid) + for (int64_t j = 0; j < factor; ++j) + blocks.push_back(s * static_cast(factor) + static_cast(j)); + const uint32_t unb = static_cast(blocks.size()); ttnn::Tensor bid_dev = UploadIdxU32( - std::move(bid), ttnn::Shape({un}), ttnn::Layout::ROW_MAJOR, device); - // indexed_fill wants rank-4 on dim 0 (native TILE geometry). + std::move(blocks), ttnn::Shape({unb}), ttnn::Layout::ROW_MAJOR, device); + // indexed_fill wants rank-4 on dim 0, and the rank change must stay a + // ZERO-COPY view, never ttnn::reshape: a TILE rank-4 reshape of + // [rows, cols] pads the trailing 1-dims to the 32-wide tile (physical + // x1024 — the Qwen3.5 GDN ssm cache asked for 32 GiB and OOM'd the P150), + // and a ROW_MAJOR reshape launches a data-movement program whose circular + // buffers scale with the tensor (4.3 MB > 1.5 MB L1 for the 4 GiB cache) — + // both found by the W0 sweep (BACKEND-TENSTORRENT-QWEN35, runs 2-5). So: + // convert to ROW_MAJOR once (to_layout — a pure copy, bytes unchanged), + // then view rank-4. The views keep the LAST dim (the block width), which + // is the metadata-only case that preserves flat order; a last-dim-changing + // view would scramble the bank interleave (see SplitFactor above). The + // staging per launch is 2 x blk x elem_size, inside L1 by construction. + // Semantics are unchanged: pure data movement, unnamed slots keep their + // bytes, last-of-duplicates wins (the oracle's loop order). + const ttnn::Layout lay = cache2d.layout(); + ttnn::Tensor cache_rm = + lay == ttnn::Layout::ROW_MAJOR ? cache2d : ttnn::to_layout(cache2d, ttnn::Layout::ROW_MAJOR); + ttnn::Tensor src_rm = + lay == ttnn::Layout::ROW_MAJOR ? src : ttnn::to_layout(src, ttnn::Layout::ROW_MAJOR); ttnn::Tensor out4 = ttnn::indexed_fill( bid_dev, - ttnn::reshape(cache2d, ttnn::Shape({static_cast(slots), - static_cast(cols), 1, 1})), - ttnn::reshape(src, ttnn::Shape({un, static_cast(cols), 1, 1})), + ttnn::experimental::view( + cache_rm, ttnn::Shape({static_cast(slots * factor), 1, 1, blk})), + ttnn::experimental::view(src_rm, ttnn::Shape({unb, 1, 1, blk})), std::nullopt, /*dim=*/0); - return ttnn::reshape(out4, ttnn::Shape({static_cast(slots), - static_cast(cols)})); + ttnn::Tensor out2 = ttnn::experimental::view( + out4, ttnn::Shape({static_cast(slots * factor), blk})); + return lay == ttnn::Layout::ROW_MAJOR ? out2 : ttnn::to_layout(out2, lay); } // The [rows, slots] one-hot f32 matrix for `idx` (idx<0 = NULL row → zero @@ -3749,18 +4344,28 @@ ttnn::Tensor ScatterRowsExact(const ttnn::Tensor& cache2d, // the wide state rows); its matmul rounds the gathered state to tf32, which // sits inside the decode step's own tf32 envelope. Exact paths use // GatherRowsExact/ScatterRowsExact above. +// +// `factor` adapts the matrix to the SPLIT shadow geometry +// (EnsureGdnCacheDevice): the result is [rows*factor, slots*factor] with +// gmat[(b*F+j)][(s*F+j')] = 1 iff s == idx[b] and j == j', so each gathered +// block-row lands under its own batch row. Same values, same tf32 envelope. ttnn::Tensor UploadOneHot(const std::vector& idx, int64_t slots, - MeshDevice& device) { + MeshDevice& device, int64_t factor = 1) { const int64_t rows = static_cast(idx.size()); - std::vector gmat(static_cast(rows) * slots, 0.0f); + const int64_t scols = slots * factor; + std::vector gmat(static_cast(rows * factor) * + static_cast(scols), + 0.0f); for (int64_t r = 0; r < rows; ++r) { const int32_t ix = idx[static_cast(r)]; if (ix < 0) continue; // NULL row: gathers zeros - gmat[static_cast(r) * slots + ix] = 1.0f; + for (int64_t j = 0; j < factor; ++j) + gmat[static_cast(r * factor + j) * static_cast(scols) + + static_cast(ix * factor + j)] = 1.0f; } return UploadTensor(std::move(gmat), - ttnn::Shape({static_cast(rows), - static_cast(slots)}), + ttnn::Shape({static_cast(rows * factor), + static_cast(scols)}), ttnn::DataType::FLOAT32, ttnn::Layout::TILE, device); } @@ -4086,7 +4691,21 @@ void CausalConv1dUpdateKernel(Queue&, Tensor& out, const Tensor& x, const Tensor ttnn::DataType::FLOAT32, ttnn::Layout::TILE, device); }; rolled = ttnn::add(ttnn::multiply(rolled, up1(std::move(mv))), - ttnn::multiply(T, up1(std::move(mk)))); + ttnn::multiply(T, up1(std::move(mk)))); + } + // bf16 STORAGE semantics (SupportsCompressedConvState, cuda_backend.cu:119): + // a bf16 conv_state is "read/written in f32 registers" — every value that + // enters the cache rounds through bf16 at the store boundary, so the next + // step's window sees the STORED bits, not the unrounded f32 tap. The device + // shadow is f32, so honor that boundary here: round the committed shadow + // through bf16 on device (RNE, zero PCIe — the shadow stays resident). With + // bf16-representable inputs (the production activation dtype) this is a + // no-op; with f32-mantissa taps it is what keeps TT on the CUDA contract — + // pinned by the bf16-state arm in tests/vt/test_tenstorrent_backend.cpp + // (steps 1+: out max_rel 2-6.5% without this round). + if (conv_state.dtype == DType::kBF16) { + rolled = ttnn::typecast(ttnn::typecast(rolled, ttnn::DataType::BFLOAT16), + ttnn::DataType::FLOAT32); } CommitConvTransposed(conv_state, std::move(rolled), uslots, uc, usl); } @@ -4186,12 +4805,16 @@ void GdnDecodeKernel(Queue&, Tensor& out, const Tensor& q_in, const Tensor& k, ttnn::Tensor dev_b = UploadTensor(std::move(be), ttnn::Shape({bh, 1}), ttnn::DataType::FLOAT32, ttnn::Layout::TILE, device); + // The shadow is born split ([slots*F, blk], F = SplitFactor of the state + // row) — see EnsureGdnCacheDevice / SplitFactor. + const uint32_t sf = + static_cast(SplitFactor(hv * dv * dk, /*esz=*/4)); ttnn::Tensor cache2d = EnsureGdnCacheDevice(state, slots, hv * dv * dk, device); std::optional oh; ttnn::Tensor S; // [B*Hv, Dv, Dk] — the batch's state rows, on device if (state_idx != nullptr) { - oh = UploadOneHot(idxv, slots, device); + oh = UploadOneHot(idxv, slots, device, sf); S = ttnn::reshape(ttnn::matmul(*oh, cache2d), ttnn::Shape({bh, udv, udk})); } else { @@ -4283,12 +4906,30 @@ void GdnDecodeKernel(Queue&, Tensor& out, const Tensor& q_in, const Tensor& k, ttnn::reshape(o, ttnn::Shape({ub, uhv, udv})), static_cast(batch * hv), udv); - ttnn::Tensor rows2d = ttnn::reshape(S_new, ttnn::Shape({ub, uhv * udv * udk})); + // rows2d in the SPLIT shadow geometry ([B*F, blk]) — same flat bytes as + // [B, Hv*Dv*Dk]; the regroup across the head boundary is one exact + // data-movement program with per-tile circular buffers (only when F > 1). + ttnn::Tensor rows2d = ttnn::reshape( + S_new, ttnn::Shape({static_cast(ub * sf), + static_cast((uhv * udv * udk) / sf)})); + // bf16 STORAGE semantics (SupportsCompressedGdnState, cuda_backend.cu): a + // bf16 ssm_state is "read/written in f32 registers" — the state a step + // STORES rounds through bf16, and the next step reads those bits. The + // device shadow is f32, so round it in place on commit (RNE typecast + // round-trip, zero PCIe). With a persistent f32 shadow instead, a + // full-mantissa state value would re-enter unrounded and drift from the + // CUDA contract within a few steps — pinned by the bf16-state arm of the + // kGdnDecode oracle test (per-step outs must match the f32 path with + // host-side bf16 round-trips bit-for-bit). + if (state.dtype == DType::kBF16) { + rows2d = ttnn::typecast(ttnn::typecast(rows2d, ttnn::DataType::BFLOAT16), + ttnn::DataType::FLOAT32); + } if (oh.has_value()) { // Exact scatter (last-writer-wins on duplicate slots, NULL rows write // nothing) — the state bytes never round through tf32 on the way back. - ttnn::Tensor newc = - ScatterRowsExact(cache2d, idxv, rows2d, slots, hv * dv * dk, device); + ttnn::Tensor newc = ScatterRowsExact(cache2d, idxv, rows2d, slots, + hv * dv * dk, sf, device); CommitDeviceLogical2D(state, std::move(newc), static_cast(slots), static_cast(hv * dv * dk)); } else { @@ -4348,13 +4989,32 @@ void GdnStateGatherKernel(Queue&, Tensor& working, const Tensor& cache, MeshDevice& device = SharedMeshDevice(); ttnn::Tensor cache2d = EnsureGdnCacheDevice(cache, slots, cache_row, device); - ttnn::Tensor gathered = - GatherRowsExact(cache2d, idxv, cache_row, device); // exact, no NULLs + // Split-shadow geometry: expand the (NULL-free) row indices to their + // block-rows — contiguous, in order — and the has_initial_state mask the + // same way, so every gathered block-row keeps its own keep factor. + const uint32_t sgf = + static_cast(SplitFactor(cache_row, /*esz=*/4)); + std::vector idx_blk; + const std::vector* idxp = &idxv; + if (sgf > 1) { + idx_blk.reserve(idxv.size() * static_cast(sgf)); + for (int32_t ix : idxv) + for (uint32_t j = 0; j < sgf; ++j) + idx_blk.push_back(ix * static_cast(sgf) + + static_cast(j)); // BLOCK-row id, not slot id + idxp = &idx_blk; + std::vector keep_exp; + keep_exp.reserve(keepv.size() * static_cast(sgf)); + for (float kv : keepv) keep_exp.insert(keep_exp.end(), sgf, kv); + keepv.swap(keep_exp); + } + ttnn::Tensor gathered = GatherRowsExact(cache2d, *idxp, + cache_row / sgf, device); // exact, no NULLs if (has_his) { gathered = ttnn::multiply( gathered, UploadTensor(std::move(keepv), - ttnn::Shape({static_cast(rows), 1}), + ttnn::Shape({static_cast(rows) * sgf, 1}), ttnn::DataType::FLOAT32, ttnn::Layout::TILE, device)); } CommitDeviceLogical2D(working, std::move(gathered), static_cast(rows), @@ -4397,9 +5057,13 @@ void GdnStateScatterKernel(Queue&, Tensor& cache, const Tensor& working, } MeshDevice& device = SharedMeshDevice(); + const uint32_t ssf = + static_cast(SplitFactor(cache_row, /*esz=*/4)); ttnn::Tensor cache2d = EnsureGdnCacheDevice(cache, slots, cache_row, device); // working rows: reuse a resident f32 shadow when its element count and // dims match, else upload (and leave it unregistered — transient rows). + // The upload carries the SPLIT shadow geometry ([rows*F, blk]); the flat + // host bytes are unchanged. bool rows_resident = false; ttnn::Tensor rows2d; { @@ -4419,11 +5083,12 @@ void GdnStateScatterKernel(Queue&, Tensor& cache, const Tensor& working, wh[static_cast(i)] = LoadElemF32(working, i); rows2d = UploadTensor( std::move(wh), - ttnn::Shape({static_cast(rows), static_cast(work_row)}), + ttnn::Shape({static_cast(rows) * ssf, + static_cast(work_row / ssf)}), ttnn::DataType::FLOAT32, ttnn::Layout::TILE, device); } - ttnn::Tensor newc = - ScatterRowsExact(cache2d, idxv, rows2d, slots, cache_row, device); + ttnn::Tensor newc = ScatterRowsExact(cache2d, idxv, rows2d, slots, + cache_row, ssf, device); CommitDeviceLogical2D(cache, std::move(newc), static_cast(slots), static_cast(cache_row)); } @@ -4482,7 +5147,15 @@ struct Registrar { RegisterOp(OpId::kGdnStateGather, DeviceType::kTENSTORRENT, reinterpret_cast(static_cast(&GdnStateGatherKernel))); RegisterOp(OpId::kGdnStateScatter, DeviceType::kTENSTORRENT, - reinterpret_cast(static_cast(&GdnStateScatterKernel))); + reinterpret_cast( + static_cast(&GdnStateScatterKernel))); + RegisterOp(OpId::kGdnPostConv, DeviceType::kTENSTORRENT, + reinterpret_cast(static_cast(&GdnPostConvKernel))); + RegisterOp(OpId::kSigmoidGateBf16, DeviceType::kTENSTORRENT, + reinterpret_cast(static_cast(&SigmoidGateBf16Kernel))); + RegisterOp(OpId::kAttnQkNormRopeGate, DeviceType::kTENSTORRENT, + reinterpret_cast( + static_cast(&AttnQkNormRopeGateKernel))); } } registrar; @@ -4744,6 +5417,8 @@ void EmbedDeviceIdsInto(void* out_host, int64_t rows, int64_t cols, void RegisterHostBuffer(void* host, size_t bytes) { if (host == nullptr) return; std::lock_guard g(SlotMutex()); + if (std::getenv("VT_TT_SLOT_TRACE") != nullptr) + std::fprintf(stderr, "[TT-SLOT] register %p bytes=%zu\n", host, bytes); BufferSlot s; s.host = host; s.bytes = bytes; @@ -4756,6 +5431,8 @@ void UnregisterHostBuffer(void* host) { if (host == nullptr) return; { std::lock_guard g(SlotMutex()); + if (std::getenv("VT_TT_SLOT_TRACE") != nullptr) + std::fprintf(stderr, "[TT-SLOT] unregister %p\n", host); Slots().erase(reinterpret_cast(host)); } DropPagedKvShadow(host); @@ -4778,6 +5455,64 @@ void MarkHostWritten(void* host) { DropEmbedTableShadow(host); } +// TRUSTED dump (BACKEND-TENSTORRENT-QWEN35 W2c measurement reset): whole +// tensors only, typed header, DUAL-READ verified. Two independent +// Synchronize+Copy passes must agree byte-for-byte or the file records +// verified=0 (a mismatch means residency state changed under us and the +// payload must not be trusted). File layout: magic 'TDMP', dtype u32, +// rank u32, dims[4] u32, numel u64, verified u8, payload. +void TrustDump(Queue& q, const char* dir, const char* name, const Tensor& t) { + static std::atomic seq{0}; + // Defense-in-depth refresh (review finding F2, #1715): Backend::Copy below + // already runs EnsureHostBytes on the source before its memcpy, so this + // entry call is redundant today; it is kept explicit because TrustDump's + // contract ("the payload reflects current truth") must not depend on a + // copy-path implementation detail elsewhere. + EnsureHostBytes(t.data); + const size_t bytes = static_cast(t.Numel()) * vt::SizeOf(t.dtype); + std::vector a(bytes), b(bytes); + Backend& tb = GetBackend(DeviceType::kTENSTORRENT); + tb.Synchronize(q); + tb.Copy(q, a.data(), t.data, bytes); + tb.Synchronize(q); + tb.Copy(q, b.data(), t.data, bytes); + tb.Synchronize(q); + const bool verified = (a == b); + uint32_t dims[4] = {0, 0, 0, 0}; + for (int i = 0; i < t.rank && i < 4; ++i) + dims[i] = static_cast(t.shape[i]); + char hdr[48]; + const uint32_t magic = 0x504D4454; // 'TDMP' + const uint32_t dt = static_cast(t.dtype); + const uint32_t rk = static_cast(t.rank); + const uint64_t numel = static_cast(t.Numel()); + const uint8_t ver = verified ? 1 : 0; + std::memcpy(hdr, &magic, 4); + std::memcpy(hdr + 4, &dt, 4); + std::memcpy(hdr + 8, &rk, 4); + std::memcpy(hdr + 12, dims, 16); + std::memcpy(hdr + 28, &numel, 8); + std::memcpy(hdr + 36, &ver, 1); + const uint64_t sid = seq.fetch_add(1, std::memory_order_relaxed); + std::FILE* f = std::fopen( + (std::string(dir) + "/" + std::to_string(sid) + "_" + name + ".tdmp").c_str(), + "wb"); + if (f != nullptr) { + std::fwrite(hdr, 1, 40, f); + if (verified) std::fwrite(a.data(), 1, bytes, f); + std::fclose(f); + } +} + + +std::vector DebugDeviceReadbackF32(Queue& q, const Tensor& t) { + (void)q; + MeshDevice& device = SharedMeshDevice(); + ttnn::Tensor dev = EnsureDevice2D(t, device); + std::vector v = dev.to_vector(); + return v; +} + void EnsureHostBytes(void* host) { if (host == nullptr) return; if (std::getenv("VT_TT_TRACE_DEBUG") != nullptr && tt_capture_active()) @@ -4867,6 +5602,7 @@ bool CopyDeviceDeviceIfCapture(void* dst, const void* src) { }(); (void)once; ttnn::Tensor src_dev; + size_t copy_bytes = 0; { std::lock_guard g(SlotMutex()); BufferSlot* s = FindSlot(const_cast(src)); @@ -4875,9 +5611,21 @@ bool CopyDeviceDeviceIfCapture(void* dst, const void* src) { if (d == nullptr) return false; if (s->bytes != d->bytes) return false; src_dev = *s->device; + copy_bytes = d->bytes; + } + if (std::getenv("VT_TT_TRACE_DEBUG") != nullptr) { + void* bt[8]; + const int nbt = ::backtrace(bt, 8); + char** sym = ::backtrace_symbols(bt, nbt); + std::fprintf(stderr, + "[TT-TRACE] device->device copy (capture-safe) dst=%p src=%p " + "bytes=%zu frames=%d\n", + dst, src, copy_bytes, nbt); + for (int i = 1; sym != nullptr && i < nbt; ++i) + std::fprintf(stderr, "[TT-TRACE] bt[%d]=%p %s\n", i, bt[i], sym[i]); + std::fflush(stderr); + std::free(sym); } - if (std::getenv("VT_TT_TRACE_DEBUG") != nullptr) - std::fprintf(stderr, "[TT-TRACE] device->device copy (capture-safe)\n"); MeshDevice& device = SharedMeshDevice(); // Allocate a destination device tensor matching src's shape/dtype/layout, // then copy. No host readback. diff --git a/tests/parity/goldens/qwen35_greedy_0_8b/neartie_gap_mnats_tenstorrent.npy b/tests/parity/goldens/qwen35_greedy_0_8b/neartie_gap_mnats_tenstorrent.npy new file mode 100644 index 000000000..3bfc31b74 Binary files /dev/null and b/tests/parity/goldens/qwen35_greedy_0_8b/neartie_gap_mnats_tenstorrent.npy differ diff --git a/tests/parity/goldens/qwen35_greedy_0_8b/our_ids_tenstorrent.npy b/tests/parity/goldens/qwen35_greedy_0_8b/our_ids_tenstorrent.npy new file mode 100644 index 000000000..8356bd10d Binary files /dev/null and b/tests/parity/goldens/qwen35_greedy_0_8b/our_ids_tenstorrent.npy differ diff --git a/tests/parity/test_qwen35_paged_engine.cpp b/tests/parity/test_qwen35_paged_engine.cpp index 542825d5c..bb6da61e5 100644 --- a/tests/parity/test_qwen35_paged_engine.cpp +++ b/tests/parity/test_qwen35_paged_engine.cpp @@ -20,8 +20,11 @@ // token given OUR prefix. // On kROCM these base files are the gate; on ANY other device the gate SKIPS // (exit 77) rather than compare another engine's tokens against ROCm-derived -// goldens — fail-safe by device, not luck. When another device gains this -// model it gets the Qwen3-dense gate's device-golden pair treatment. +// goldens — fail-safe by device, not luck — with ONE ratified exception: the +// Tenstorrent lane, which gets the Mistral gate's device-golden pair treatment +// (our_ids_tenstorrent.npy + neartie_gap_mnats_tenstorrent.npy; gap teacher- +// forced via scripts/qwen3-neartie-gap-transformers.py — vLLM has no TT +// backend, so `transformers` is the secondary oracle per AGENTS.md's registry). // // PROVENANCE OF THE GREEN: the committed our_ids/near-tie pair is the // FIXED engine's sequence, oracle-re-derived after the AttnQkNormRopeGate @@ -230,17 +233,18 @@ void RunGate(const std::string& golden_subdir, const char* label) { snap, vllm::entrypoints::EngineParams{}); // The base golden pair for this model is ROCm-captured (see the file header). - // The Metal/Tenstorrent device lanes are kept for when those devices run this - // model; their pairs do not exist yet, so those devices skip loudly. + // The Tenstorrent device lane carries its OWN oracle-backed golden pair + // (the Mistral gate's treatment); every other device still skips loudly. const vt::DeviceType run_dev = loaded->runner().device().type; const bool rocm = run_dev == vt::DeviceType::kROCM; + const bool tenstorrent = run_dev == vt::DeviceType::kTENSTORRENT; + const bool device_golden = tenstorrent; // FAIL SAFE BY DEVICE: this model's only oracle-backed goldens are // ROCm-captured (the pinned vLLM-ROCm oracle on gfx1100 — no dgx/CUDA capture - // exists). On ANY other device the anchor would be compared against another - // engine's sequence, so skip loudly rather than misattribute a drift. When a - // Metal/CUDA/Tenstorrent capture lands it gets the device-golden pair - // treatment this file's header describes. - if (!rocm) { + // exists). On any device without its OWN pair the anchor would be compared + // against another engine's sequence, so skip loudly rather than misattribute + // a drift. Tenstorrent is the one lane with a committed device pair. + if (!rocm && !device_golden) { SkipGate(label, "goldens are ROCm-oracle-captured; this run is on device " "type " + std::to_string(static_cast(run_dev)) + " — capture that device's pair first"); @@ -257,14 +261,55 @@ void RunGate(const std::string& golden_subdir, const char* label) { vt::OpId::kAttnQkNormRopeGate, vt::OpId::kReshapeAndCache, vt::OpId::kPagedAttention, vt::OpId::kSiluAndMul, vt::OpId::kGreedyArgmax}; - if (rocm) { + if (rocm || device_golden) { for (vt::OpId op : kGdnOps) { CHECK(vt::OpRegistered(op, run_dev)); vt::ResetOpProviderStats(op, run_dev); } vt::EnableOpProviderCallStats(true); MESSAGE(label << ": running on device type " << static_cast(run_dev) - << " (5=ROCM) — gated against the ROCm oracle-backed golden pair"); + << " (5=ROCM, 6=TENSTORRENT) — gated against " + << (device_golden ? "this device's OWN oracle-backed golden" + : "the ROCm oracle-backed golden pair")); + } + + // Device-appropriate anchor + teacher-forced gap goldens. Base = the ROCm + // pair. Tenstorrent carries its own pair (captured via VT_DUMP_IDS=1 on the + // P150, then qwen3-neartie-gap-transformers.py teacher-forces `transformers` + // on that sequence — NOT vLLM, which has no Tenstorrent backend; same + // secondary-oracle lane and precedent as the Qwen3-0.6B and Mistral-7B TT + // goldens). + const char* ids_name = tenstorrent ? "our_ids_tenstorrent.npy" : "our_ids.npy"; + const char* gap_name = + tenstorrent ? "neartie_gap_mnats_tenstorrent.npy" : "neartie_gap_mnats.npy"; + parity::NpyArray o_dev, gap_dev; // keep the device arrays alive for the loop + bool bootstrap_only = false; + if (device_golden) { + const bool have_dev = fs::exists(gdir / ids_name) && fs::exists(gdir / gap_name); + if (!have_dev && dump) { + // Bootstrap dump path: generate tokens, write raw i32, skip the gate. + bootstrap_only = true; + MESSAGE(label << ": BOOTSTRAP dump (device golden absent) for Tenstorrent..."); + } else { + REQUIRE_MESSAGE(have_dev, + label << ": device oracle golden absent (" << ids_name << " / " + << gap_name + << ") — capture the sequence with VT_DUMP_IDS=1, then " + "teacher-force it: qwen3-neartie-gap-transformers.py " + "-> device golden pair"); + o_dev = parity::LoadNpy((gdir / ids_name).string()); + gap_dev = parity::LoadNpy((gdir / gap_name).string()); + REQUIRE(o_dev.dtype == "(i * T + j)] = got[static_cast(j)]; } + if (bootstrap_only) continue; // dump-only path; no anchor/gap yet // Anchor: the committed anchor is the exact deterministic sequence OUR ROCm // engine produces. Drift is a hard REQUIRE — no cross-device latitude. @@ -343,8 +389,10 @@ void RunGate(const std::string& golden_subdir, const char* label) { } // Backend proof: token equality alone does not prove which device ran. - if (rocm) { - vt::EnableOpProviderCallStats(false); + // The bootstrap dump path does not exercise the full op set to a comparison, + // so its stats prove reachability only (still selections > 0, declines == 0). + if (device_golden || rocm) vt::EnableOpProviderCallStats(false); + if (device_golden || rocm) { for (vt::OpId op : kGdnOps) { const auto st = vt::GetOpProviderStats(op, run_dev); CHECK_MESSAGE(st.selections > 0, @@ -361,11 +409,15 @@ void RunGate(const std::string& golden_subdir, const char* label) { << vt::GetOpProviderStats(vt::OpId::kPagedAttention, run_dev).selections << ", kGdnDecode selections=" << vt::GetOpProviderStats(vt::OpId::kGdnDecode, run_dev).selections + << ", kCausalConv1dUpdate selections=" + << vt::GetOpProviderStats(vt::OpId::kCausalConv1dUpdate, run_dev).selections << ")"); } if (dump) { - const std::string path = (gdir / "our_ids.i32").string(); + const std::string dump_name = + tenstorrent ? "our_ids_tenstorrent.i32" : "our_ids.i32"; + const std::string path = (gdir / dump_name).string(); std::FILE* f = std::fopen(path.c_str(), "wb"); if (f != nullptr) { std::fwrite(our_dump.data(), sizeof(int32_t), our_dump.size(), f); @@ -373,6 +425,16 @@ void RunGate(const std::string& golden_subdir, const char* label) { MESSAGE(label << " dumped our token ids -> " << path); } } + if (bootstrap_only) { + // BOOTSTRAP complete — no correctness bar was applied this run (every + // prompt hit the `continue` above). The summary below would print + // "0/16 prompts PASS ... 0 forward-divergent" over a vacuous REQUIRE — + // the Mistral gate returns here for the same reason. + MESSAGE(label << ": BOOTSTRAP complete -- ids dumped, NO correctness bar " + "was applied this run. Teacher-force the dumped sequence, " + "commit the golden pair, then re-run without VT_DUMP_IDS."); + return; + } MESSAGE(label << " correctness gate: " << (strict_exact + neartie_only) << "/" << N << " prompts PASS (STRICT token-exact vs oracle per-prompt greedy: " << strict_exact << "/" << N << "; near-tie-band only: " << neartie_only diff --git a/tests/scripts/test_agent_record.py b/tests/scripts/test_agent_record.py index 621f3a78e..33692d1cb 100644 --- a/tests/scripts/test_agent_record.py +++ b/tests/scripts/test_agent_record.py @@ -2192,3 +2192,40 @@ def test_the_kernel_pin_is_load_bearing_for_this_row(self) -> None: any("KERNEL rows" in e for e in errors), "removing the row must break the pin it was bumped for", ) + + +class Qwen35GdnBackendRowBacksTheRatchet(unittest.TestCase): + """The BACKEND ratchet bump 86 -> 87 is backed by the Qwen3.5 GDN row + (#1715). + + Same shape and reason as the BACKEND-TENSTORRENT-RESIDUAL-GOLDEN class + above: the count is re-pinned by hand, so a bump with no row behind it is + indistinguishable from a bump for a row that really landed. The first test + ties THIS value of the pin to a real matrix line; the second proves the pin + BINDS against the shipped matrix file through the checker's own entry point, + which is what makes the pair semantic evidence rather than a restatement. + """ + + ROW = "BACKEND-TENSTORRENT-QWEN35" + + def test_the_qwen35_gdn_row_exists_in_the_backend_matrix(self) -> None: + text = (ROOT / ".agents/backend-matrix.md").read_text(encoding="utf-8") + matching = [ + line for line in text.splitlines() if line.startswith(f"| `{self.ROW}` |") + ] + self.assertEqual(len(matching), 1, f"{self.ROW} must appear exactly once") + + def test_the_backend_checker_accepts_the_matrix_and_binds(self) -> None: + clean: list[str] = [] + agent_record.check_matrices(clean) + self.assertEqual([e for e in clean if "backend rows" in e.lower()], []) + path, count = agent_record.MATRICES["BACKEND"] + errors: list[str] = [] + with mock.patch.dict( + agent_record.MATRICES, {"BACKEND": (path, count - 1)} + ): + agent_record.check_matrices(errors) + self.assertTrue( + any("backend rows" in e.lower() for e in errors), + f"the BACKEND pin must bind; got {errors}", + ) diff --git a/tests/scripts/test_check_gate_commands.py b/tests/scripts/test_check_gate_commands.py index 0740e9d3a..694480adb 100644 --- a/tests/scripts/test_check_gate_commands.py +++ b/tests/scripts/test_check_gate_commands.py @@ -1006,3 +1006,44 @@ def test_bpe_quadratic_merge_left_the_gated_population_cleanly(self): if __name__ == "__main__": unittest.main() + + +class Qwen35GdnRunnablePopulationTests(unittest.TestCase): + """The BACKEND-TENSTORRENT-QWEN35 entry is the whole of what #1715's wiring + row changed in this checker, so the credit has to be checkable rather than + plausible: the row must be IN the pinned population, its audit verdict must + be earned by the record (not merely asserted), and removing the entry must + break the exact pin -- which is what proves the row was pinned because it + ENTERED the runnable population and not to quiet a gate.""" + + ROW = "BACKEND-TENSTORRENT-QWEN35" + + def test_the_row_is_pinned_and_its_verdict_is_runnable(self): + self.assertIn(self.ROW, gates.RUNNABLE_BASELINE) + verdicts = {r["id"]: r["verdict"] for r in gates.audit()} + self.assertEqual(verdicts.get(self.ROW), "runnable") + + def test_dropping_the_qwen35_gdn_entry_breaks_the_pin(self): + # MUTATION, in the direction this pin actually moved (#1715): remove the + # entry and the exact-population equality inside the checker has to go + # red, because the audit still reports the row as runnable from its spec. + reduced = set(gates.RUNNABLE_BASELINE) - {self.ROW} + self.assertNotEqual(reduced, set(gates.RUNNABLE_BASELINE)) + runnable = {r["id"] for r in gates.audit() if r["verdict"] == "runnable"} + self.assertNotEqual(runnable, reduced) + self.assertEqual(runnable - reduced, {self.ROW}) + + def test_the_qwen35_gdn_row_is_credited_for_real_commands(self): + # A record row can only be credited for the checker that reads the + # record, so the load-bearing assertion is the last one: the `## Gates` + # section of the row's own spec must name commands that genuinely fail + # on a broken tree -- the parity binary and the preflight chain -- not + # a command that exits 0 in any tree (`git diff`). + self.assertIn(self.ROW, gates.RUNNABLE_BASELINE) + spec = ( + gates.ROOT / ".agents/specs/tenstorrent-qwen35.md" + ).read_text(encoding="utf-8") + gate_section = spec.split("\n## Gates\n", 1)[1].split("\n## ", 1)[0] + self.assertIn("tests/test_qwen35_paged_engine", gate_section) + self.assertIn("scripts/agent-preflight.sh", gate_section) + self.assertNotIn("git diff", gate_section) diff --git a/tests/vt/test_tenstorrent_backend.cpp b/tests/vt/test_tenstorrent_backend.cpp index a84994a31..72c29ff72 100644 --- a/tests/vt/test_tenstorrent_backend.cpp +++ b/tests/vt/test_tenstorrent_backend.cpp @@ -28,6 +28,7 @@ #include "vllm/platforms/interface.h" #include "vt/backend.h" +#include "vt/dtype.h" #include "vt/ops.h" using vt::Backend; @@ -534,6 +535,73 @@ TEST_CASE("kTENSTORRENT kRmsNorm matches a host F32 reference (weight, no residu CHECK(max_abs_diff < 0.5f); } +// Gemma style (w+1) — the Qwen3.5 norm (every RmsNorm site passes gemma=true). +// The device arm must bake +1 into the gamma it hands ttnn::rms_norm; dropping +// it collapsed ambient host-free prefill to `,` (BACKEND-TENSTORRENT-QWEN35 +// W2b). Reference computes in f32 with the SAME wj = w+1 order as the kernel's +// host arm. +TEST_CASE("kTENSTORRENT kRmsNorm gemma matches a host F32 reference (w+1)") { + if (!TenstorrentPresent()) { + MESSAGE("SKIPPED: no Tenstorrent device on this box"); + return; + } + REQUIRE(vt::OpRegistered(vt::OpId::kRmsNorm, DeviceType::kTENSTORRENT)); + + constexpr int64_t Rows = 32, D = 32; + constexpr float Eps = 1e-6f; + Backend& backend = vt::GetBackend(DeviceType::kTENSTORRENT); + auto rms_norm = reinterpret_cast( + vt::GetOp(vt::OpId::kRmsNorm, DeviceType::kTENSTORRENT)); + + std::vector host_x(Rows * D), host_w(D), host_out(Rows * D, 0.0f); + for (size_t i = 0; i < host_x.size(); ++i) + host_x[i] = (static_cast(i % 17) - 8.0f) * 0.15f; + for (int64_t j = 0; j < D; ++j) + // Small weights make the +1 dominant: a dropped bake fails by ~1.0 per + // element, far outside any storage envelope. + host_w[static_cast(j)] = 0.05f * static_cast(j % 3); + + void* mem_x = backend.Alloc(host_x.size() * sizeof(float)); + void* mem_w = backend.Alloc(host_w.size() * sizeof(float)); + void* mem_out = backend.Alloc(host_out.size() * sizeof(float)); + Queue q = backend.CreateQueue(); + backend.Copy(q, mem_x, host_x.data(), host_x.size() * sizeof(float)); + backend.Copy(q, mem_w, host_w.data(), host_w.size() * sizeof(float)); + + Tensor x = + Tensor::Contiguous(mem_x, vt::DType::kF32, Device{DeviceType::kTENSTORRENT, 0}, {Rows, D}); + Tensor w = + Tensor::Contiguous(mem_w, vt::DType::kF32, Device{DeviceType::kTENSTORRENT, 0}, {D}); + Tensor out = + Tensor::Contiguous(mem_out, vt::DType::kF32, Device{DeviceType::kTENSTORRENT, 0}, {Rows, D}); + + vt::RmsNormArgs args; + args.eps = Eps; + args.gemma = true; + rms_norm(q, out, x, w, args, /*residual=*/nullptr); + + backend.Copy(q, host_out.data(), mem_out, host_out.size() * sizeof(float)); + backend.Free(mem_x); + backend.Free(mem_w); + backend.Free(mem_out); + + float max_abs_diff = 0.0f; + for (int64_t r = 0; r < Rows; ++r) { + float sumsq = 0.0f; + for (int64_t j = 0; j < D; ++j) { + const float v = host_x[r * D + j]; + sumsq += v * v; + } + const float inv = 1.0f / std::sqrt(sumsq / static_cast(D) + Eps); + for (int64_t j = 0; j < D; ++j) { + const float ref = + host_x[r * D + j] * inv * (host_w[static_cast(j)] + 1.0f); + max_abs_diff = std::max(max_abs_diff, std::fabs(host_out[r * D + j] - ref)); + } + } + CHECK(max_abs_diff < 0.5f); +} + // Qwen3-dense MLP SwiGLU half. Device path (slice + silu + mul) via BF16 tiles. TEST_CASE("kTENSTORRENT kSiluAndMul matches host F32 within BF16 envelope") { if (!TenstorrentPresent()) { @@ -2338,6 +2406,128 @@ TEST_CASE("kTENSTORRENT kCausalConv1dUpdate matches the CPU f32 oracle (read-old tt.Free(mo); tt.Free(mx); } + // --- bf16 conv_state (production mamba_cache_dtype default): the + // SupportsCompressedConvState arm. The oracle is the CUDA bf16-STORAGE + // emulation — widen bf16->f32, compute the step in f32, round the state + // back to bf16 EVERY step ("read/written in f32 registers", cuda_gdn.cu) — + // not a persistent f32 recurrence. x carries a FULL f32 mantissa in the + // first arm so a shadow that skips the per-step store-rounding diverges in + // step 2+ (the rounded tap feeds the next MAC); the second arm feeds + // bf16-representable x (the production activation dtype), where storage + // rounding is a no-op and both readings agree bit-for-bit. + { + Backend& tt = *vt::TryGetBackend(DeviceType::kTENSTORRENT); + Backend& cpu = *vt::TryGetBackend(DeviceType::kCPU); + const int64_t slots = 4, B = 2, state_len = K - 1; + const std::vector idx{3, 1}; + for (int x_full_mantissa : {1, 0}) { + uint32_t s = 66000u + static_cast(x_full_mantissa); + std::vector w(static_cast(C * K)), bias(static_cast(C)); + std::vector st_f(static_cast(slots * C * state_len)); + for (float& v : w) v = 0.4f * GdnLcg(s); + for (float& v : bias) v = 0.1f * GdnLcg(s); + for (float& v : st_f) v = GdnLcg(s); + // The cache holds bf16 bits (production storage); the initial values + // round ONCE so both sides start from identical storage. + std::vector st_bits(st_f.size()); + for (size_t i = 0; i < st_f.size(); ++i) st_bits[i] = vt::F32ToBF16(st_f[i]); + + // Emulation: per step, widen the bf16 cache, run the CPU f32 oracle + // in place, round the cache back to bf16 (the CUDA store boundary). + std::vector emu = st_bits; + auto emu_step = [&](const std::vector& x, std::vector& out) { + std::vector st(emu.size()); + for (size_t i = 0; i < emu.size(); ++i) st[i] = vt::BF16ToF32(emu[i]); + void* mx = cpu.Alloc(x.size() * sizeof(float)); + void* mw = cpu.Alloc(w.size() * sizeof(float)); + void* mb = cpu.Alloc(bias.size() * sizeof(float)); + void* ms = cpu.Alloc(st.size() * sizeof(float)); + void* mo = cpu.Alloc(out.size() * sizeof(float)); + void* mi = cpu.Alloc(idx.size() * sizeof(int32_t)); + Queue q = cpu.CreateQueue(); + cpu.Copy(q, mx, x.data(), x.size() * sizeof(float)); + cpu.Copy(q, mw, w.data(), w.size() * sizeof(float)); + cpu.Copy(q, mb, bias.data(), bias.size() * sizeof(float)); + cpu.Copy(q, ms, st.data(), st.size() * sizeof(float)); + cpu.Copy(q, mi, idx.data(), idx.size() * sizeof(int32_t)); + Tensor tx = Tensor::Contiguous(mx, vt::DType::kF32, Device{DeviceType::kCPU, 0}, {B, C}); + Tensor tw = Tensor::Contiguous(mw, vt::DType::kF32, Device{DeviceType::kCPU, 0}, {C, K}); + Tensor tb = Tensor::Contiguous(mb, vt::DType::kF32, Device{DeviceType::kCPU, 0}, {C}); + Tensor ts = Tensor::Contiguous(ms, vt::DType::kF32, Device{DeviceType::kCPU, 0}, + {slots, C, state_len}); + Tensor ti = Tensor::Contiguous(mi, vt::DType::kI32, Device{DeviceType::kCPU, 0}, + {static_cast(idx.size())}); + Tensor to = Tensor::Contiguous(mo, vt::DType::kF32, Device{DeviceType::kCPU, 0}, {B, C}); + vt::CausalConv1dArgs a; + a.silu_activation = true; + vt::CausalConv1dUpdate(q, to, tx, tw, &tb, ts, a, &ti); + cpu.Copy(q, out.data(), mo, out.size() * sizeof(float)); + cpu.Copy(q, st.data(), ms, st.size() * sizeof(float)); + for (size_t i = 0; i < emu.size(); ++i) emu[i] = vt::F32ToBF16(st[i]); + for (void* m : {mx, mw, mb, ms, mo, mi}) cpu.Free(m); + }; + + // TT side: one persistent bf16 cache buffer, chained steps (the + // transposed-shadow fast path across steps is exactly what must honor + // the storage rounding). + void* mw = tt.Alloc(w.size() * sizeof(float)); + void* mb = tt.Alloc(bias.size() * sizeof(float)); + void* ms = tt.Alloc(st_bits.size() * sizeof(uint16_t)); + void* mo = tt.Alloc(static_cast(B * C) * sizeof(float)); + void* mx = tt.Alloc(static_cast(B * C) * sizeof(float)); + void* mi = tt.Alloc(idx.size() * sizeof(int32_t)); + Queue q = tt.CreateQueue(); + tt.Copy(q, mw, w.data(), w.size() * sizeof(float)); + tt.Copy(q, mb, bias.data(), bias.size() * sizeof(float)); + tt.Copy(q, ms, st_bits.data(), st_bits.size() * sizeof(uint16_t)); // the ONE upload + tt.Copy(q, mi, idx.data(), idx.size() * sizeof(int32_t)); + for (int step_i = 0; step_i < 4; ++step_i) { + std::vector x(static_cast(B * C)); + for (float& v : x) + v = x_full_mantissa ? (2.0f * GdnLcg(s) + 0.03125f) // full f32 mantissa + : vt::BF16ToF32(vt::F32ToBF16(2.0f * GdnLcg(s))); + std::vector out_emu(static_cast(B * C), 0.0f); + emu_step(x, out_emu); + tt.Copy(q, mx, x.data(), x.size() * sizeof(float)); + Tensor tx = Tensor::Contiguous(mx, vt::DType::kF32, + Device{DeviceType::kTENSTORRENT, 0}, {B, C}); + Tensor tw = Tensor::Contiguous(mw, vt::DType::kF32, + Device{DeviceType::kTENSTORRENT, 0}, {C, K}); + Tensor tb = Tensor::Contiguous(mb, vt::DType::kF32, + Device{DeviceType::kTENSTORRENT, 0}, {C}); + Tensor ts = Tensor::Contiguous(ms, vt::DType::kBF16, + Device{DeviceType::kTENSTORRENT, 0}, + {slots, C, state_len}); + Tensor ti = Tensor::Contiguous(mi, vt::DType::kI32, + Device{DeviceType::kTENSTORRENT, 0}, + {static_cast(idx.size())}); + Tensor to = Tensor::Contiguous(mo, vt::DType::kF32, + Device{DeviceType::kTENSTORRENT, 0}, {B, C}); + vt::CausalConv1dArgs a; + a.silu_activation = true; + vt::CausalConv1dUpdate(q, to, tx, tw, &tb, ts, a, &ti); + std::vector out_tt(static_cast(B * C), 0.0f); + tt.Copy(q, out_tt.data(), mo, out_tt.size() * sizeof(float)); + GdnDiffStats d = CompareVsOracle(out_tt, out_emu, 1e-4f, 1e-5f); + MESSAGE("kCausalConv1dUpdate bf16-state x_full=", x_full_mantissa, + " step=", step_i, ": out max_abs=", d.max_abs, + " max_rel=", d.max_rel); + CHECK(std::isfinite(d.max_abs)); + CHECK(d.within); + } + // Storage truth: the downloaded bf16 cache must match the emulation's + // bf16 cache BIT-FOR-BIT (both round the same values through RNE). + std::vector got_bits(st_bits.size(), 0); + tt.Copy(q, got_bits.data(), ms, got_bits.size() * sizeof(uint16_t)); + size_t mism = 0; + for (size_t i = 0; i < emu.size(); ++i) + if (got_bits[i] != emu[i]) ++mism; + MESSAGE("kCausalConv1dUpdate bf16-state x_full=", x_full_mantissa, + ": cache bit mismatches=", mism, "/", emu.size()); + CHECK(mism == 0); + for (void* m : {mw, mb, ms, mo, mx, mi}) tt.Free(m); + } + } } TEST_CASE("kTENSTORRENT kGdnDecode matches the CPU f32 oracle (rank-1 step, both state_idx forms, NULL slot)") { @@ -2507,6 +2697,39 @@ TEST_CASE("kTENSTORRENT kGdnDecode matches the CPU f32 oracle (rank-1 step, both CHECK(out_tt[static_cast(1 * Hv * Dv + e)] == 0.0f); } + // --- W2a: the PRODUCTION state-row width Hv*Dk*Dv = 16*128*128 = 262144 + // through the exact indexed scatter, WITH a NULL row (so the live-row + // compaction feeds the column-chunked launch). Without ScatterRowsExact's + // chunking this throws "dataflow buffers ... beyond max L1 size of + // 1572864 B"; with it, out AND cache must match the oracle within the same + // envelope as every other arm. + { + const int64_t B = 2, Hk = 16, Hv = 16, slots = 3; + std::vector q, k, v, g, beta; + gen(77000u, B, Hk, Hv, q, k, v, g, beta); + std::vector cache(static_cast(slots * Hv * Dv * Dk)); + { + uint32_t s = 78000u; + for (float& x : cache) x = 0.05f * GdnLcg(s); + } + const std::vector idx{2, -1}; + std::vector ca_cpu = cache, ca_tt = cache; + std::vector out_cpu(static_cast(B * Hv * Dv), 0.0f), + out_tt(static_cast(B * Hv * Dv), 0.0f); + step(cpu, DeviceType::kCPU, B, Hk, Hv, q, k, v, g, beta, ca_cpu, &idx, out_cpu); + step(*vt::TryGetBackend(DeviceType::kTENSTORRENT), DeviceType::kTENSTORRENT, B, Hk, + Hv, q, k, v, g, beta, ca_tt, &idx, out_tt); + const float tol = 0.02f; + GdnDiffStats d = CompareVsOracle(out_tt, out_cpu, 0.0f, tol); + GdnDiffStats ds = CompareVsOracle(ca_tt, ca_cpu, 0.0f, tol); + MESSAGE("kGdnDecode wide-state (row 262144) + NULL: out max_abs=", d.max_abs, + " cache max_abs=", ds.max_abs, " tol=", tol); + CHECK(std::isfinite(d.max_abs)); + CHECK(d.within); + CHECK(std::isfinite(ds.max_abs)); + CHECK(ds.within); + } + // --- Chained steps on the SAME state buffer + traffic counters: the state // must stay device-resident across steps (one upload, zero downloads). { @@ -2577,6 +2800,118 @@ TEST_CASE("kTENSTORRENT kGdnDecode matches the CPU f32 oracle (rank-1 step, both CHECK(tr.state_h2d_bytes == want_up); // exactly ONE upload across 3 steps CHECK(tr.state_d2h_bytes == 0); } + + // --- bf16 ssm_state (production mamba_ssm_dtype default = conv dtype): + // the SupportsCompressedGdnState arm. The contract is CUDA bf16 STORAGE + // semantics — each step loads the bf16 state into f32 registers, computes, + // and STORES back to bf16 ("read/written in f32 registers", + // cuda_backend.cu) — not a persistent f32 recurrence. The reference here is + // the TT f32-state path EMULATING that boundary: after every step the f32 + // state is rounded to bf16 bits and widened back before the next step. The + // bf16-state path must match it — the shadow may keep f32 tiles internally, + // but the values that RE-ENTER the next step must be the stored bf16 ones. + { + const int64_t B = 2, Hk = 2, Hv = 2, slots = 3; + const std::vector idx{2, 0}; + vt::GdnArgs args; + args.scale = 1.0f / std::sqrt(static_cast(Dk)); + Backend& tt = *vt::TryGetBackend(DeviceType::kTENSTORRENT); + const size_t st_n = static_cast(slots * Hv * Dv * Dk); + std::vector st0(st_n); + { + uint32_t s = 77000u; + for (float& x : st0) x = 0.05f * GdnLcg(s); // full f32 mantissa + } + std::vector bits0(st_n); + for (size_t i = 0; i < st_n; ++i) bits0[i] = vt::F32ToBF16(st0[i]); + + // Emulation: f32 buffer, per-step round-trip through bf16 bits. + std::vector emu_st(st_n); + for (size_t i = 0; i < st_n; ++i) emu_st[i] = vt::BF16ToF32(bits0[i]); + std::vector> out_emu(3); // emulation out per step + void* mq = tt.Alloc(static_cast(B * Hk * Dk) * sizeof(float)); + void* mk = tt.Alloc(static_cast(B * Hk * Dk) * sizeof(float)); + void* mv = tt.Alloc(static_cast(B * Hv * Dv) * sizeof(float)); + void* mg = tt.Alloc(static_cast(B * Hv) * sizeof(float)); + void* mb = tt.Alloc(static_cast(B * Hv) * sizeof(float)); + void* mo = tt.Alloc(static_cast(B * Hv * Dv) * sizeof(float)); + void* mi = tt.Alloc(idx.size() * sizeof(int32_t)); + void* ms_f32 = tt.Alloc(st_n * sizeof(float)); + void* ms_bf16 = tt.Alloc(st_n * sizeof(uint16_t)); + Queue qq = tt.CreateQueue(); + tt.Copy(qq, mi, idx.data(), idx.size() * sizeof(int32_t)); + tt.Copy(qq, ms_bf16, bits0.data(), st_n * sizeof(uint16_t)); // the ONE upload + Tensor tidx = Tensor::Contiguous(mi, vt::DType::kI32, + Device{DeviceType::kTENSTORRENT, 0}, + {static_cast(idx.size())}); + auto step_tensors = [&](void* ms, vt::DType sdt, Tensor& tq, Tensor& tk, + Tensor& tv, Tensor& tg, Tensor& tb, Tensor& ts, + Tensor& to) { + tq = Tensor::Contiguous(mq, vt::DType::kF32, Device{DeviceType::kTENSTORRENT, 0}, + {B, Hk, Dk}); + tk = Tensor::Contiguous(mk, vt::DType::kF32, Device{DeviceType::kTENSTORRENT, 0}, + {B, Hk, Dk}); + tv = Tensor::Contiguous(mv, vt::DType::kF32, Device{DeviceType::kTENSTORRENT, 0}, + {B, Hv, Dv}); + tg = Tensor::Contiguous(mg, vt::DType::kF32, Device{DeviceType::kTENSTORRENT, 0}, + {B, Hv}); + tb = Tensor::Contiguous(mb, vt::DType::kF32, Device{DeviceType::kTENSTORRENT, 0}, + {B, Hv}); + ts = Tensor::Contiguous(ms, sdt, Device{DeviceType::kTENSTORRENT, 0}, + {slots, Hv, Dv, Dk}); + to = Tensor::Contiguous(mo, vt::DType::kF32, Device{DeviceType::kTENSTORRENT, 0}, + {B, Hv, Dv}); + }; + for (int step_i = 0; step_i < 3; ++step_i) { + std::vector q, k, v, g, beta; + gen(78000u + static_cast(step_i * 17), B, Hk, Hv, q, k, v, g, beta); + tt.Copy(qq, mq, q.data(), q.size() * sizeof(float)); + tt.Copy(qq, mk, k.data(), k.size() * sizeof(float)); + tt.Copy(qq, mv, v.data(), v.size() * sizeof(float)); + tt.Copy(qq, mg, g.data(), g.size() * sizeof(float)); + tt.Copy(qq, mb, beta.data(), beta.size() * sizeof(float)); + // Emulation step (f32 buffer + host-side bf16 round of the state). + tt.Copy(qq, ms_f32, emu_st.data(), st_n * sizeof(float)); + { + Tensor tq, tk, tv, tg, tb, ts, to; + step_tensors(ms_f32, vt::DType::kF32, tq, tk, tv, tg, tb, ts, to); + vt::GdnDecode(qq, to, tq, tk, tv, tg, tb, ts, args, &tidx); + out_emu[static_cast(step_i)].assign( + static_cast(B * Hv * Dv), 0.0f); + tt.Copy(qq, out_emu[static_cast(step_i)].data(), mo, + out_emu[static_cast(step_i)].size() * sizeof(float)); + tt.Copy(qq, emu_st.data(), ms_f32, st_n * sizeof(float)); + for (size_t i = 0; i < st_n; ++i) + emu_st[i] = vt::BF16ToF32(vt::F32ToBF16(emu_st[i])); // the store boundary + } + // bf16-state step on the PERSISTENT buffer (shadow fast path). + std::vector out_tt(static_cast(B * Hv * Dv), 0.0f); + { + Tensor tq, tk, tv, tg, tb, ts, to; + step_tensors(ms_bf16, vt::DType::kBF16, tq, tk, tv, tg, tb, ts, to); + vt::GdnDecode(qq, to, tq, tk, tv, tg, tb, ts, args, &tidx); + tt.Copy(qq, out_tt.data(), mo, out_tt.size() * sizeof(float)); + } + // Both paths fed IDENTICAL f32 inputs and identical state VALUES (the + // bf16 path's shadow rounds exactly where the emulation rounds), so the + // deterministic device kernels must produce bit-identical outs. + GdnDiffStats d = + CompareVsOracle(out_tt, out_emu[static_cast(step_i)], 0.0f, 0.0f); + MESSAGE("kGdnDecode bf16-state step=", step_i, + ": out max_abs=", d.max_abs, " (0 = storage semantics honored)"); + CHECK(d.max_abs == 0.0f); + } + // Final stored truth: the bf16 buffer's bits must equal the emulation's + // rounded bits exactly. + std::vector got_bits(st_n, 0); + tt.Copy(qq, got_bits.data(), ms_bf16, st_n * sizeof(uint16_t)); + size_t mism = 0; + for (size_t i = 0; i < st_n; ++i) + if (got_bits[i] != vt::F32ToBF16(emu_st[i])) ++mism; + MESSAGE("kGdnDecode bf16-state: cache bit mismatches=", mism, "/", st_n); + CHECK(mism == 0); + for (void* m : {mq, mk, mv, mg, mb, mo, mi, ms_f32, ms_bf16}) tt.Free(m); + } } TEST_CASE("kTENSTORRENT kGdnPrefill<->kGdnDecode round-trip (final states agree, both arms, traffic)") { @@ -2950,6 +3285,16 @@ TEST_CASE("kTENSTORRENT kGdnStateGather/Scatter match the CPU f32 oracle (indexe // Rank-2 cache. gather_scatter({5, Dk}, {3, Dk}, {2, 2, 4}, nullptr, "rank-2"); + // W2a: a row WIDER than indexed_fill's L1 staging budget. The Qwen3.5 GDN + // ssm_state row is Hv*Dk*Dv = 262144 elems; without the split-shadow + // geometry (SplitFactor) the generic interleaved indexed_fill stages + // 2 x 1 MB of dataflow buffer and throws "beyond max L1 size of 1572864 B". + // One arm pins all three properties at that width: bit-exact rows, + // per-(slot, block) last-of-duplicates wins (idx repeats slot 1), and the + // unnamed slot 2 keeps its bytes. + gather_scatter({3, 262144}, {2, 262144}, {1, 0}, nullptr, + "ssm wide-row (262144 > L1 page budget)"); + // --- Untouched rows: scatter must not write rows no index names. { const int64_t slots = 5, rows = 2; @@ -3163,3 +3508,523 @@ TEST_CASE("kTENSTORRENT GDN edge shapes (all-empty prefill early return, empty d MESSAGE("edge empty conv/gather: no-ops held"); } } + +// ==== BACKEND-TENSTORRENT-QWEN35 W1: the Qwen3.5 op delta vs the CPU f32 oracle +// Same doctrine as the GDN block above: identical inputs, both arms run the +// SAME public vt:: op on their own backend, outputs compared. A missing TT +// kernel refuses by name — this suite's red state before the kernel lands. + +TEST_CASE("kTENSTORRENT kSigmoidGateBf16 matches the CPU f32 oracle (attn dtype arms)") { + if (!TenstorrentPresent()) { + MESSAGE("SKIPPED: no Tenstorrent device on this box"); + return; + } + Backend& cpu = *vt::TryGetBackend(DeviceType::kCPU); + // Production shape is [T, K] with K = Hq*Dh (qwen3_5.cpp:2546). T=1 is the + // decode step; 65 crosses the 64-row tile boundary. Gate values span +-8 so + // the sigmoid input precision is load-bearing: at gate ~ -5 the sigmoid is + // ~7e-3 and a bf16 pre-round of the gate moves the product by >2% relative — + // exactly the arm that proves the gate stays f32 (ops.cpp:4140). + for (int64_t T : {int64_t{1}, int64_t{3}, int64_t{64}, int64_t{65}}) { + for (int64_t K : {int64_t{128}, int64_t{384}}) { + for (const vt::DType adt : {vt::DType::kF32, vt::DType::kBF16}) { + const int64_t n = T * K; + std::vector attn(static_cast(n)), gate(static_cast(n)); + { + uint32_t s = 4001u + static_cast(T * 7 + K + (adt == vt::DType::kF32 ? 0 : 1)); + for (int64_t i = 0; i < n; ++i) { + attn[static_cast(i)] = 4.0f * GdnLcg(s) + 0.125f; // full f32 mantissa + gate[static_cast(i)] = 16.0f * GdnLcg(s); // spans +-8 + } + } + auto run = [&](Backend& b, DeviceType dt, std::vector& out_bf) { + out_bf.assign(static_cast(n), 0); + const size_t a_bytes = adt == vt::DType::kF32 ? sizeof(float) : sizeof(uint16_t); + void* ma = b.Alloc(n * a_bytes); + void* mg = b.Alloc(n * sizeof(float)); + void* mo = b.Alloc(n * sizeof(uint16_t)); + Queue q = b.CreateQueue(); + if (adt == vt::DType::kF32) { + b.Copy(q, ma, attn.data(), attn.size() * sizeof(float)); + } else { // bf16 arm holds bf16-representable values (upcast exact) + std::vector bits(static_cast(n)); + for (int64_t i = 0; i < n; ++i) + bits[static_cast(i)] = vt::F32ToBF16(attn[static_cast(i)]); + b.Copy(q, ma, bits.data(), bits.size() * sizeof(uint16_t)); + } + b.Copy(q, mg, gate.data(), gate.size() * sizeof(float)); + Tensor ta = Tensor::Contiguous(ma, adt, Device{dt, 0}, {T, K}); + Tensor tg = Tensor::Contiguous(mg, vt::DType::kF32, Device{dt, 0}, {T, K}); + Tensor to = Tensor::Contiguous(mo, vt::DType::kBF16, Device{dt, 0}, {T, K}); + vt::SigmoidGateBf16(q, to, ta, tg); + b.Copy(q, out_bf.data(), mo, out_bf.size() * sizeof(uint16_t)); + b.Free(ma); + b.Free(mg); + b.Free(mo); + }; + std::vector out_cpu, out_tt; + run(cpu, DeviceType::kCPU, out_cpu); + run(*vt::TryGetBackend(DeviceType::kTENSTORRENT), DeviceType::kTENSTORRENT, out_tt); + std::vector ref(out_cpu.size()), got(out_tt.size()); + for (size_t i = 0; i < out_cpu.size(); ++i) { + ref[i] = vt::BF16ToF32(out_cpu[i]); + got[i] = vt::BF16ToF32(out_tt[i]); + } + // Both arms round once to bf16 (RNE); the only TT-vs-CPU delta is the + // SFPU f32 sigmoid (accurate exp + reciprocal_iter<2) vs std::exp — + // a few f32 ULP, which can flip at most one bf16 rounding. One bf16 + // ULP is <= 2^-7 relative on any binade boundary, so that is the + // envelope; abs_floor guards the exact-zero ref. + const float rel = 1.0f / 128.0f, abs_floor = 1e-6f; + GdnDiffStats d = CompareVsOracle(got, ref, rel, abs_floor); + // doctest quirk (this build): MESSAGE streams a `const char*` VARIABLE + // as its pointer→bool ("1"); only literals and std::string print as + // text (same fix as kAttnQkNormRopeGate below). + const std::string adt_name = adt == vt::DType::kF32 ? "f32" : "bf16"; + MESSAGE("kSigmoidGateBf16 T=", T, " K=", K, + " attn=", adt_name, + ": max_abs=", d.max_abs, " max_rel=", d.max_rel, " within=", d.within); + CHECK(d.within); + } + } + } +} + +TEST_CASE("kTENSTORRENT kGdnPostConv matches the CPU f32 oracle (fused split + l2norm + g/beta)") { + if (!TenstorrentPresent()) { + MESSAGE("SKIPPED: no Tenstorrent device on this box"); + return; + } + Backend& cpu = *vt::TryGetBackend(DeviceType::kCPU); + const vt::L2NormArgs args{1e-6f}; + // Production geometry (0.8B GDN layers): Hk=Hv=32 heads, Dk=Dv=128, so + // conv cols = 2*Hk*Dk + Hv*Dv (qwen3_5.cpp:4146). Sweep smaller heads plus + // the production shape; one config exercises Hk != Hv (the contract admits + // it even though the model does not use it). + struct Cfg { int64_t T, Hk, Dk, Hv, Dv; }; + const Cfg cfgs[] = {{1, 2, 128, 2, 128}, {3, 8, 128, 8, 128}, + {65, 2, 128, 2, 128}, {64, 32, 128, 32, 128}, + {3, 2, 128, 4, 128}}; + for (const Cfg& c : cfgs) { + // out_bf16: the PRODUCTION arm — VT_GDN_BF16 defaults the matmul-input + // activations q/k/v to bf16 (qwen3_5.cpp GdnActDType), while g/beta stay + // f32. Covers the kernel's per-out typecast on commit. + for (int out_bf16 : {0, 1}) { + for (int conv_bf16 : {0, 1}) { + for (int ab_bf16 : {0, 1}) { + const int64_t key_dim = c.Hk * c.Dk, value_dim = c.Hv * c.Dv; + const int64_t conv_dim = 2 * key_dim + value_dim; + const int64_t n = c.T * conv_dim; + const int64_t a_stride = c.Hv + 8; // padded row view (production form) + // f32 values; the bf16 arms round these to bf16 bits ONCE so both + // backends read identical bf16-representable inputs. + std::vector conv_f(n), araw_f(c.T * c.Hv), braw_f(c.T * c.Hv), + a_log(static_cast(c.Hv)), dt_bias(static_cast(c.Hv)); + { + uint32_t s = 9001u + static_cast(c.T * 11 + c.Hk * 5 + c.Hv + + conv_bf16 * 2 + ab_bf16 + + out_bf16 * 41); + for (float& v : conv_f) v = 2.0f * GdnLcg(s) + 0.0625f; + // araw spans +-25 so x = araw+dt_bias crosses the softplus + // threshold-20 branch; braw spans +-8 (sigmoid sensitivity). + for (float& v : araw_f) v = 50.0f * GdnLcg(s); + for (float& v : braw_f) v = 16.0f * GdnLcg(s); + for (float& v : a_log) v = -4.0f * (GdnLcg(s) + 0.5f); // exp in (0,1] + for (float& v : dt_bias) v = 4.0f * GdnLcg(s); + } + const vt::DType cdt = conv_bf16 ? vt::DType::kBF16 : vt::DType::kF32; + const vt::DType adt = ab_bf16 ? vt::DType::kBF16 : vt::DType::kF32; + const vt::DType odt = out_bf16 ? vt::DType::kBF16 : vt::DType::kF32; + const size_t cb = cdt == vt::DType::kF32 ? sizeof(float) : sizeof(uint16_t); + const size_t ab = adt == vt::DType::kF32 ? sizeof(float) : sizeof(uint16_t); + // bf16-bit staging helpers (round once, reuse for both backends). + auto to_bits = [&](const std::vector& v) { + std::vector bits(v.size()); + for (size_t i = 0; i < v.size(); ++i) bits[i] = vt::F32ToBF16(v[i]); + return bits; + }; + std::vector conv_bits = conv_bf16 ? to_bits(conv_f) : std::vector{}; + // Padded a/b backing buffers, garbage in the pad cols (production form); + // copied whole so both backends see identical bytes. + auto pad_ab = [&](const std::vector& live) { + std::vector bits(static_cast(c.T * a_stride), 0xABCDu); + for (int64_t t = 0; t < c.T; ++t) + for (int64_t h = 0; h < c.Hv; ++h) + bits[static_cast(t * a_stride + h)] = vt::F32ToBF16(live[static_cast(t * c.Hv + h)]); + return bits; + }; + std::vector a_bits, b_bits; + std::vector a_pad_f, b_pad_f; // f32-arm padded buffers + if (ab_bf16) { + a_bits = pad_ab(araw_f); + b_bits = pad_ab(braw_f); + } else { + a_pad_f.assign(static_cast(c.T * a_stride), std::numeric_limits::quiet_NaN()); + b_pad_f.assign(static_cast(c.T * a_stride), std::numeric_limits::quiet_NaN()); + for (int64_t t = 0; t < c.T; ++t) + for (int64_t h = 0; h < c.Hv; ++h) { + a_pad_f[static_cast(t * a_stride + h)] = araw_f[static_cast(t * c.Hv + h)]; + b_pad_f[static_cast(t * a_stride + h)] = braw_f[static_cast(t * c.Hv + h)]; + } + } + + std::vector q_cpu(static_cast(c.T * key_dim)), + k_cpu(q_cpu.size()), v_cpu(static_cast(c.T * value_dim)), + g_cpu(static_cast(c.T * c.Hv)), beta_cpu(g_cpu.size()); + std::vector q_tt(q_cpu.size()), k_tt(k_cpu.size()), v_tt(v_cpu.size()), + g_tt(g_cpu.size()), beta_tt(beta_cpu.size()); + auto run = [&](Backend& b, DeviceType dt, std::vector& qo, + std::vector& ko, std::vector& vo, + std::vector& go, std::vector& bo) { + void* mc = b.Alloc(n * cb); + // padded a/b backing (garbage in the pad rows' tail columns) + void* mab = b.Alloc(static_cast(c.T * a_stride) * ab); + void* mbb = b.Alloc(static_cast(c.T * a_stride) * ab); + void* mal = b.Alloc(a_log.size() * sizeof(float)); + void* md = b.Alloc(dt_bias.size() * sizeof(float)); + void* mq = b.Alloc(qo.size() * (out_bf16 ? sizeof(uint16_t) : sizeof(float))); + void* mk = b.Alloc(ko.size() * (out_bf16 ? sizeof(uint16_t) : sizeof(float))); + void* mv = b.Alloc(vo.size() * (out_bf16 ? sizeof(uint16_t) : sizeof(float))); + void* mg = b.Alloc(go.size() * sizeof(float)); + void* mb = b.Alloc(bo.size() * sizeof(float)); + Queue q = b.CreateQueue(); + if (conv_bf16) b.Copy(q, mc, conv_bits.data(), conv_bits.size() * sizeof(uint16_t)); + else b.Copy(q, mc, conv_f.data(), conv_f.size() * sizeof(float)); + if (ab_bf16) { + b.Copy(q, mab, a_bits.data(), a_bits.size() * sizeof(uint16_t)); + b.Copy(q, mbb, b_bits.data(), b_bits.size() * sizeof(uint16_t)); + } else { + b.Copy(q, mab, a_pad_f.data(), a_pad_f.size() * sizeof(float)); + b.Copy(q, mbb, b_pad_f.data(), b_pad_f.size() * sizeof(float)); + } + b.Copy(q, mal, a_log.data(), a_log.size() * sizeof(float)); + b.Copy(q, md, dt_bias.data(), dt_bias.size() * sizeof(float)); + Tensor tc = Tensor::Contiguous(mc, cdt, Device{dt, 0}, {c.T, conv_dim}); + auto view2 = [&](void* m) { // [T, Hv] inner-contiguous row view + Tensor t{}; + t.data = m; + t.dtype = adt; + t.device = Device{dt, 0}; + t.rank = 2; + t.shape[0] = c.T; + t.shape[1] = c.Hv; + t.stride[0] = a_stride; + t.stride[1] = 1; + return t; + }; + Tensor ta = view2(mab), tb2 = view2(mbb); + Tensor tal = Tensor::Contiguous(mal, vt::DType::kF32, Device{dt, 0}, {c.Hv}); + Tensor td = Tensor::Contiguous(md, vt::DType::kF32, Device{dt, 0}, {c.Hv}); + Tensor tq = Tensor::Contiguous(mq, odt, Device{dt, 0}, {c.T, c.Hk, c.Dk}); + Tensor tk = Tensor::Contiguous(mk, odt, Device{dt, 0}, {c.T, c.Hk, c.Dk}); + Tensor tv = Tensor::Contiguous(mv, odt, Device{dt, 0}, {c.T, c.Hv, c.Dv}); + Tensor tg = Tensor::Contiguous(mg, vt::DType::kF32, Device{dt, 0}, {c.T, c.Hv}); + Tensor tbe = Tensor::Contiguous(mb, vt::DType::kF32, Device{dt, 0}, {c.T, c.Hv}); + vt::GdnPostConv(q, tq, tk, tv, tg, tbe, tc, ta, tb2, tal, td, args); + // bf16 outs download as bits and widen once for the comparator. + auto read_out = [&](void* m, std::vector& o) { + if (!out_bf16) { + b.Copy(q, o.data(), m, o.size() * sizeof(float)); + return; + } + std::vector bits(o.size()); + b.Copy(q, bits.data(), m, bits.size() * sizeof(uint16_t)); + for (size_t i = 0; i < o.size(); ++i) + o[i] = vt::BF16ToF32(bits[i]); + }; + read_out(mq, qo); + read_out(mk, ko); + read_out(mv, vo); + b.Copy(q, go.data(), mg, go.size() * sizeof(float)); + b.Copy(q, bo.data(), mb, bo.size() * sizeof(float)); + for (void* m : {mc, mab, mbb, mal, md, mq, mk, mv, mg, mb}) b.Free(m); + }; + run(cpu, DeviceType::kCPU, q_cpu, k_cpu, v_cpu, g_cpu, beta_cpu); + run(*vt::TryGetBackend(DeviceType::kTENSTORRENT), DeviceType::kTENSTORRENT, q_tt, + k_tt, v_tt, g_tt, beta_tt); + // q/k: bf16 tile l2norm — the kL2Norm envelope (0.02 abs, the row's + // calibrated bound for unit-vector-scale outputs). bf16 outs add one + // store round on each side (a compute ULP can flip it: <= 2^-7 rel). + const float o_rel = out_bf16 ? 1.0f / 128.0f : 0.0f; + GdnDiffStats dq = CompareVsOracle(q_tt, q_cpu, o_rel, 0.02f); + GdnDiffStats dk = CompareVsOracle(k_tt, k_cpu, o_rel, 0.02f); + // v: slice copy through the bf16 tile — exact for bf16 conv inputs, + // one bf16 round for f32 conv inputs (<= 2^-7 rel on any binade), + // plus the store round when the out itself is bf16. + const float v_rel = (conv_bf16 ? 0.0f : 1.0f / 128.0f) + o_rel; + GdnDiffStats dv = CompareVsOracle(v_tt, v_cpu, v_rel, 1e-6f); + // g/beta: f32 tiles end to end (softplus threshold-20, exp(a_log), + // sigmoid) — SFPU f32 vs std::exp/log1p is ULP-level. + GdnDiffStats dg = CompareVsOracle(g_tt, g_cpu, 1e-4f, 1e-6f); + GdnDiffStats db = CompareVsOracle(beta_tt, beta_cpu, 1e-4f, 1e-6f); + // doctest quirk (this build): MESSAGE streams a `const char*` VARIABLE + // as its pointer→bool ("1"); only literals and std::string print as + // text (same fix as kAttnQkNormRopeGate). + const std::string conv_name = conv_bf16 ? "bf16" : "f32"; + const std::string ab_name = ab_bf16 ? "bf16" : "f32"; + const std::string out_name = out_bf16 ? "bf16" : "f32"; + MESSAGE("kGdnPostConv T=", c.T, " Hk=", c.Hk, " Hv=", c.Hv, + " conv=", conv_name, " ab=", ab_name, " out=", out_name, + ": q[max_abs=", dq.max_abs, "] k[max_abs=", dk.max_abs, + "] v[max_abs=", dv.max_abs, "] g[max_abs=", dg.max_abs, + " max_rel=", dg.max_rel, "] beta[max_abs=", db.max_abs, + " max_rel=", db.max_rel, "]"); + CHECK(dq.within); + CHECK(dk.within); + CHECK(dv.within); + CHECK(dg.within); + CHECK(db.within); + } + } + } + } +} + +TEST_CASE("kTENSTORRENT kAttnQkNormRopeGate matches the CPU f32 oracle (fused preamble, GQA + partial rope)") { + if (!TenstorrentPresent()) { + MESSAGE("SKIPPED: no Tenstorrent device on this box"); + return; + } + Backend& cpu = *vt::TryGetBackend(DeviceType::kCPU); + // Production call: q/k/gate f32 outs, gemma=true (qwen3_5.cpp:5206-5224), + // qgate rows [q(Dh)|gate(Dh)] per head with the merged-QKV row stride, kf + // rows likewise; rot < Dh exercises the partial-rope tail (normed, not + // rotated). Hq > Hkv covers GQA. + struct Cfg { int64_t T, Hq, Hkv, Dh, rot; bool gemma; int in_bf16; int out_bf16; }; + const Cfg cfgs[] = { + {1, 4, 2, 128, 64, true, 0, 0}, {3, 4, 2, 128, 64, true, 1, 0}, + {3, 4, 2, 128, 128, true, 0, 0}, {65, 4, 2, 128, 64, true, 1, 0}, + {3, 32, 8, 128, 64, true, 1, 0}, {3, 4, 2, 128, 64, false, 0, 0}, + // bf16-out arms — the kernel's per-out typecast branches + // (tenstorrent_ops.cpp AttnQkNormRopeGateKernel commit legs). bf16 in + + // bf16 out: the gate passthrough stays exact (rounding a bf16 value is + // identity). f32 in + bf16 out: every leg rounds once at the store. + {1, 4, 2, 128, 64, true, 1, 1}, {3, 32, 8, 128, 64, true, 0, 1}, + }; + for (const Cfg& c : cfgs) { + const vt::RmsNormArgs na{1e-6f, c.gemma}; + const vt::RopeArgs ra{10000.0, static_cast(c.rot)}; // base unused: cos_sin is data + const int64_t qrow = c.Hq * 2 * c.Dh, krow = c.Hkv * c.Dh; + const int64_t qgate_pad = 16, kf_pad = 16; + std::vector qgate_f(static_cast(c.T * (qrow + qgate_pad))); + std::vector kf_f(static_cast(c.T * (krow + kf_pad))); + std::vector qw(static_cast(c.Dh)), kw(static_cast(c.Dh)); + std::vector cs(static_cast(c.T * c.rot)); + { + uint32_t s = 11001u + static_cast(c.T * 13 + c.Hq * 7 + c.rot + c.gemma * 3 + c.in_bf16 + c.out_bf16 * 43); + for (float& v : qgate_f) v = 2.0f * GdnLcg(s); + for (float& v : kf_f) v = 2.0f * GdnLcg(s); + for (float& v : qw) v = 0.2f * GdnLcg(s); // gemma adds 1 + for (float& v : kw) v = 0.2f * GdnLcg(s); + for (float& v : cs) v = 2.0f * GdnLcg(s); // cos|sin as data in [-1,1) + } + const vt::DType idt = c.in_bf16 ? vt::DType::kBF16 : vt::DType::kF32; + const size_t ib = idt == vt::DType::kF32 ? sizeof(float) : sizeof(uint16_t); + std::vector qgate_bits, kf_bits; + if (c.in_bf16) { + qgate_bits.resize(qgate_f.size()); + kf_bits.resize(kf_f.size()); + for (size_t i = 0; i < qgate_f.size(); ++i) qgate_bits[i] = vt::F32ToBF16(qgate_f[i]); + for (size_t i = 0; i < kf_f.size(); ++i) kf_bits[i] = vt::F32ToBF16(kf_f[i]); + } + const int64_t qn = c.T * c.Hq * c.Dh, kn = c.T * c.Hkv * c.Dh; + std::vector q_cpu(static_cast(qn)), k_cpu(static_cast(kn)), + g_cpu(static_cast(qn)); + std::vector q_tt(q_cpu.size()), k_tt(k_cpu.size()), g_tt(g_cpu.size()); + auto run = [&](Backend& b, DeviceType dt, std::vector& qo, + std::vector& ko, std::vector& go) { + void* mqg = b.Alloc(qgate_f.size() * ib); + void* mkf = b.Alloc(kf_f.size() * ib); + void* mqw = b.Alloc(qw.size() * sizeof(float)); + void* mkw = b.Alloc(kw.size() * sizeof(float)); + void* mcs = b.Alloc(cs.size() * sizeof(float)); + void* mq = b.Alloc(qo.size() * (c.out_bf16 ? sizeof(uint16_t) : sizeof(float))); + void* mk = b.Alloc(ko.size() * (c.out_bf16 ? sizeof(uint16_t) : sizeof(float))); + void* mg = b.Alloc(go.size() * (c.out_bf16 ? sizeof(uint16_t) : sizeof(float))); + Queue q = b.CreateQueue(); + if (c.in_bf16) { + b.Copy(q, mqg, qgate_bits.data(), qgate_bits.size() * sizeof(uint16_t)); + b.Copy(q, mkf, kf_bits.data(), kf_bits.size() * sizeof(uint16_t)); + } else { + b.Copy(q, mqg, qgate_f.data(), qgate_f.size() * sizeof(float)); + b.Copy(q, mkf, kf_f.data(), kf_f.size() * sizeof(float)); + } + b.Copy(q, mqw, qw.data(), qw.size() * sizeof(float)); + b.Copy(q, mkw, kw.data(), kw.size() * sizeof(float)); + b.Copy(q, mcs, cs.data(), cs.size() * sizeof(float)); + auto strided2 = [&](void* m, int64_t rows, int64_t cols, int64_t stride) { + Tensor t{}; + t.data = m; + t.dtype = idt; + t.device = Device{dt, 0}; + t.rank = 2; + t.shape[0] = rows; + t.shape[1] = cols; + t.stride[0] = stride; + t.stride[1] = 1; + return t; + }; + Tensor tqg = strided2(mqg, c.T, qrow, qrow + qgate_pad); + Tensor tkf = strided2(mkf, c.T, krow, krow + kf_pad); + Tensor tqw = Tensor::Contiguous(mqw, vt::DType::kF32, Device{dt, 0}, {c.Dh}); + Tensor tkw = Tensor::Contiguous(mkw, vt::DType::kF32, Device{dt, 0}, {c.Dh}); + Tensor tcs = Tensor::Contiguous(mcs, vt::DType::kF32, Device{dt, 0}, {c.T, c.rot}); + const vt::DType odt = c.out_bf16 ? vt::DType::kBF16 : vt::DType::kF32; + Tensor tq = Tensor::Contiguous(mq, odt, Device{dt, 0}, {c.T, c.Hq, c.Dh}); + Tensor tk = Tensor::Contiguous(mk, odt, Device{dt, 0}, {c.T, c.Hkv, c.Dh}); + Tensor tg = Tensor::Contiguous(mg, odt, Device{dt, 0}, {c.T, c.Hq, c.Dh}); + vt::AttnQkNormRopeGate(q, tq, tk, tg, tqg, tkf, tqw, tkw, tcs, na, ra); + // bf16 outs download as bits and widen once for the comparator. + auto read_out = [&](void* m, std::vector& o) { + if (!c.out_bf16) { + b.Copy(q, o.data(), m, o.size() * sizeof(float)); + return; + } + std::vector bits(o.size()); + b.Copy(q, bits.data(), m, bits.size() * sizeof(uint16_t)); + for (size_t i = 0; i < o.size(); ++i) o[i] = vt::BF16ToF32(bits[i]); + }; + read_out(mq, qo); + read_out(mk, ko); + read_out(mg, go); + for (void* m : {mqg, mkf, mqw, mkw, mcs, mq, mk, mg}) b.Free(m); + }; + run(cpu, DeviceType::kCPU, q_cpu, k_cpu, g_cpu); + run(*vt::TryGetBackend(DeviceType::kTENSTORRENT), DeviceType::kTENSTORRENT, q_tt, k_tt, + g_tt); + // q/k: f32 tile norm + rope (inside the row's RopeApplyDeviceNeox + // envelope); outputs are O(1) (unit-RMS rows, rotation preserves + // magnitude). bf16 outs add one store round per side (<= 2^-7 rel). + // f32 outs carry no store round, so the abs floor is pinned at 1e-3: + // measured green headroom on P150 is <= 5.3e-4, while an unconditional + // pre-norm bf16 typecast of q (the ops.cpp:4140 branch this guards) + // lands at 5.4e-3 — 1e-3 separates both sides. + const float o_rel = c.out_bf16 ? 1.0f / 128.0f : 0.0f; + const float o_abs = c.out_bf16 ? 0.02f : 1e-3f; + GdnDiffStats dq = CompareVsOracle(q_tt, q_cpu, o_rel, o_abs); + GdnDiffStats dk = CompareVsOracle(k_tt, k_cpu, o_rel, o_abs); + // gate: EXACT passthrough when it stays f32 or when the inputs are + // already bf16 (rounding identity); one store round when f32 inputs meet + // a bf16 out (ops.cpp:1660-1662, 4140 — the sigmoid input must not be + // rounded BEFORE the gate; the out store is the only legal round). + const float g_rel = (c.out_bf16 && !c.in_bf16) ? 1.0f / 128.0f : 0.0f; + GdnDiffStats dg = CompareVsOracle(g_tt, g_cpu, g_rel, 0.0f); + // doctest quirk (this build): MESSAGE streams a `const char*` VARIABLE as + // its pointer→bool ("1"); only literals and std::string print as text. + const std::string in = c.in_bf16 ? "bf16" : "f32"; + const std::string out = c.out_bf16 ? "bf16" : "f32"; + MESSAGE("kAttnQkNormRopeGate T=", c.T, " Hq=", c.Hq, " Hkv=", c.Hkv, + " rot=", c.rot, " gemma=", c.gemma, " in=", in, " out=", out, + ": q[max_abs=", dq.max_abs, "] k[max_abs=", dk.max_abs, + "] gate[max_abs=", dg.max_abs, "]"); + CHECK(dq.within); + CHECK(dk.within); + CHECK(dg.within); + } +} + + +// SCRATCH (2b replay): real captured bytes through kMatmulBT. Reads +// /tmp/w2b_qkvz_cpu/{h0.bin,w_packed.bin}; skips when absent. REMOVE before +// landing. + +// Interior slice views of ONE tracked allocation must not share the base's +// staged device tensor. EnsureDevice2D keyed its hit on (base slot, dims) +// only, so ProjectGdnBA's `a` projection consumed the `b` weight rows and +// every Qwen3.5 GDN layer diverged from layer 0 (BACKEND-TENSTORRENT-QWEN35 +// W2c: TT-a == CPU-b, corr -0.23). The fix requires t.data == slot host base +// for hits AND stores; this case pins that with the engine's exact sequence: +// one packed [2N,K] resident, then b-view then a-view matmuls. +TEST_CASE("kTENSTORRENT kMatmulBT slice views do not consume the base staging") { + if (!TenstorrentPresent()) { + MESSAGE("SKIPPED: no Tenstorrent device on this box"); + return; + } + REQUIRE(vt::OpRegistered(vt::OpId::kMatmulBT, DeviceType::kTENSTORRENT)); + constexpr int64_t M = 5, K = 64, N = 16; + Backend& backend = vt::GetBackend(DeviceType::kTENSTORRENT); + std::vector hb(M * K), wb(2 * N * K); + // Activations O(1) and halves far apart (+1 vs -2 weights): consuming the + // wrong slice's staging moves outputs by ~3x the row sum — far past the + // threshold. Tiny activations would swallow the poisoning below it. + for (size_t i = 0; i < hb.size(); ++i) hb[i] = static_cast(0x3E80 + (i % 9)); + for (size_t i = 0; i < wb.size(); ++i) + wb[i] = static_cast(i < wb.size() / 2 ? 0x3F80 : 0xC000); + void* ma = backend.Alloc(M * K * 2); + void* mb = backend.Alloc(2 * N * K * 2); + void* mo = backend.Alloc(M * N * 4); + Queue q = backend.CreateQueue(); + backend.Copy(q, ma, hb.data(), M * K * 2); + backend.Copy(q, mb, wb.data(), 2 * N * K * 2); + Tensor a = Tensor::Contiguous(ma, vt::DType::kBF16, Device{DeviceType::kTENSTORRENT, 0}, {M, K}); + Tensor packed = Tensor::Contiguous(mb, vt::DType::kBF16, Device{DeviceType::kTENSTORRENT, 0}, {2 * N, K}); + auto mm = reinterpret_cast(vt::GetOp(vt::OpId::kMatmulBT, DeviceType::kTENSTORRENT)); + auto widen = [](uint16_t u) { + uint32_t bits = static_cast(u) << 16; + float f; std::memcpy(&f, &bits, 4); return f; + }; + // Run TWICE so the second view also meets a warm (b-staged) cache. + for (int rep = 0; rep < 2; ++rep) { + for (int half = 0; half < 2; ++half) { + Tensor wview = packed.Slice(0, half * N, (half + 1) * N); + Tensor o = Tensor::Contiguous(mo, vt::DType::kF32, Device{DeviceType::kTENSTORRENT, 0}, {M, N}); + mm(q, o, a, wview); + std::vector oh(M * N); + backend.Copy(q, oh.data(), mo, M * N * 4); + double worst = 0; + for (int64_t i = 0; i < M; ++i) + for (int64_t j = 0; j < N; ++j) { + double acc = 0; + for (int64_t k = 0; k < K; ++k) + acc += static_cast(widen(hb[i * K + k])) * + widen(wb[(half * N + j) * K + k]); + worst = std::max(worst, std::fabs(acc - oh[i * N + j])); + } + CHECK_MESSAGE(worst < 1.0, + "half " << half << " rep " << rep << " worst " << worst + << " — a view consumed another slice's staging"); + } + } + backend.Free(ma); backend.Free(mb); backend.Free(mo); +} + +// bf16 x bf16 -> F32 out through kMatmulBT (the ProjectGdnBA split-arm +// signature): output dtype differs from input dtype. +TEST_CASE("kTENSTORRENT kMatmulBT bf16 inputs to F32 output") { + if (!TenstorrentPresent()) { + MESSAGE("SKIPPED: no Tenstorrent device on this box"); + return; + } + REQUIRE(vt::OpRegistered(vt::OpId::kMatmulBT, DeviceType::kTENSTORRENT)); + constexpr int64_t M = 5, K = 1024, N = 16; + Backend& backend = vt::GetBackend(DeviceType::kTENSTORRENT); + std::vector hb(M * K), wb(N * K); + for (size_t i = 0; i < hb.size(); ++i) hb[i] = static_cast(0x3800 + (i % 13)); + for (size_t i = 0; i < wb.size(); ++i) wb[i] = static_cast(0x3c00 + (i % 7)); + void* ma = backend.Alloc(M * K * 2); void* mb = backend.Alloc(N * K * 2); + void* mo = backend.Alloc(M * N * 4); + Queue q = backend.CreateQueue(); + backend.Copy(q, ma, hb.data(), M * K * 2); + backend.Copy(q, mb, wb.data(), N * K * 2); + Tensor a = Tensor::Contiguous(ma, vt::DType::kBF16, Device{DeviceType::kTENSTORRENT, 0}, {M, K}); + Tensor b = Tensor::Contiguous(mb, vt::DType::kBF16, Device{DeviceType::kTENSTORRENT, 0}, {N, K}); + Tensor o = Tensor::Contiguous(mo, vt::DType::kF32, Device{DeviceType::kTENSTORRENT, 0}, {M, N}); + auto mm = reinterpret_cast(vt::GetOp(vt::OpId::kMatmulBT, DeviceType::kTENSTORRENT)); + mm(q, o, a, b); + std::vector oh(M * N); + backend.Copy(q, oh.data(), mo, M * N * 4); + auto widen = [](uint16_t u) { + uint32_t bits = static_cast(u) << 16; + float f; std::memcpy(&f, &bits, 4); return f; + }; + double worst = 0; + for (int64_t i = 0; i < M; ++i) + for (int64_t j = 0; j < N; ++j) { + double acc = 0; + for (int64_t k = 0; k < K; ++k) + acc += static_cast(widen(hb[i * K + k])) * widen(wb[j * K + k]); + worst = std::max(worst, std::fabs(acc - oh[i * N + j])); + } + CHECK(worst < 1.0); + backend.Free(ma); backend.Free(mb); backend.Free(mo); +}