diff --git a/.agents/issue-index.md b/.agents/issue-index.md index 6b66f8ea2..6a8340284 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -388,6 +388,7 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1198](https://github.com/mudler/vllm.cpp/issues/1198) | — | Three specs assert `model_loader.cpp` behaviour the loader no longer has, found while verifying line citations for conversion under [#1143](https://github.com/mudler/vllm.cpp/issues/1143): `safetensors-windowed-load.md:63,108` quotes a `shards.clear()` that does not exist (the shard vector is a `shared_ptr` released by the deferred-expert closure, `model_loader.cpp:1636-1643`), `gguf-dflash-draft.md:17-18` calls `LoadDflashDraft` "still typed on `std::vector`" when it takes a `SharedHeadSource`, its `A5`/`B2` rows plan around a GGUF refusal the loader says at `:906-910` is GONE, and `model-factory-registry.md:91` cites `IsDenseArch` which survives only in a comment saying the registry superseded it. Filed rather than repaired because a citation sweep can see the claim is false and not what the true statement is — that is the owning row's judgement. Owned under `## Owed` in [`citation-anchor-freshness.md`](specs/citation-anchor-freshness.md) | bug | | [#1199](https://github.com/mudler/vllm.cpp/issues/1199) | — | `scripts/check-symbol-anchors.py --upstream-root` resolved 354 `vllm/...::Symbol` citations against the parity pin `555967922`: **343 fresh, 0 stale**, which is the measurement that decides the symbol convention over line anchors, since the same pin advance broke every line anchor [#1139](https://github.com/mudler/vllm.cpp/issues/1139) examined. The 11 that did not resolve name a FILE absent at the pin: six `.agents/model-matrix.md` rows (`olmo.py`, `olmo2.py`, `ouro.py`, `persimmon.py`, `plamo2.py`, `fuyu.py` — `registry.py:658` routes `OlmoForCausalLM` to the `transformers` fallback and `registry.py:765` lists `FuyuForCausalLM` as removed at 0.25.0, so our rows claim a mirror source vLLM deleted), two malformed `vllm/tests/kernels/...` paths, `vllm/v1/worker/gpu/worker.py` for `gpu_worker.py`, and one fixture-text false positive. Filed rather than repaired because the model-matrix half is a claim about vLLM, not a path edit. Owned under `## Owed` in [`citation-anchor-freshness.md`](specs/citation-anchor-freshness.md) | bug | | [#1213](https://github.com/mudler/vllm.cpp/issues/1213) | `ENV-LEASE-RUNTIME-STAGING` | `AGENTS.md` stated that a leased worker "has no compiler, no downloader and no Python, so it cannot produce a runtime in place", and `.agents/environment.md` carried the matching clause twice for `dgx:gpu0`. All three negatives are false. `rc describe dgx:gpu0` states that a job runs as root in an Ubuntu 24.04 container carrying `git`, `curl`, `wget`, `ssh`, `gcc`, `g++`, `make`, `cmake`, `ninja`, `pkg-config`, `python3`, `pip` and `venv`, and it instructs the reader to install anything missing; the one limit it names is the absent CUDA toolkit. Two jobs then compiled inside a lease on 2026-08-18: `claude/mudler-ubuntu-box/qwen38-gate` apt-installed `cuda-nvcc-13-0` from the `ubuntu2404/sbsa` lane and built this tree 1791/1791 to `BUILD_RC=0` (`/mnt/nas_share/rc/qwen38-gate/out-main/cfg.log` records `nvcc` 13.0.88 and `CUDA feature cutlass-fp8: ENABLED for [121a]`), and `/mnt/nas_share/rc/mtp_test/build.sh` cloned `github.com/mudler/llama.cpp` from inside a job and left a 97 MB `libggml-cuda.so` on the share. **Why it matters:** "the lease cannot produce a runtime in place" is the stated basis for treating the pinned vLLM oracle as unreachable from a lease, and that oracle is the denominator for every speed-parity number the project owes, so the premise needs re-testing rather than inheriting. This claims nothing about a model run; #1185 owns that and stays open. FIXED IN FLOW: the `AGENTS.md` paragraph and both `.agents/environment.md` clauses now say what the measurement supports, and the four real limits (no preinstalled CUDA toolkit, global installs leak until the pod restarts, CIFS `/workspace` holds no symlink so build in `/tmp` and `cp -rL`, and `-j 4` because unconstrained parallelism OOM-reboots the box) plus the host-versus-container egress distinction ride with the correction. | record | +| [#1222](https://github.com/mudler/vllm.cpp/issues/1222) | `BACKEND-ROCM` | Qwen3.5 paged-engine gate reports success when a required greedy, anchor, or near-tie-gap artifact is absent | bug | | [#1225](https://github.com/mudler/vllm.cpp/issues/1225) | `SPEC-DSPARK-BLOCK-SIZE-GUARD` | A DSpark speculative length below the draft's block is accepted silently. `ResolveDspark` carries upstream's `k >= dspark_block_size` hard error (`include/vllm/config/speculative.h:179-185`, from `vllm/config/speculative.py:1003-1027` @ `555967922`) and both production call sites pass `std::nullopt` for `n_predict` and for `dspark_block_size` (`src/vllm/entrypoints/model_loader.cpp:881-883` and `:1675-1677`), so the floor reaches no user path and only `tests/vllm/config/test_speculative_dspark.cpp:99-107` drives it. Nothing in our draft path reads the checkpoint's block key — the block layout is sized by `k` alone (`include/vllm/v1/worker/gpu/spec_decode/dspark/speculator.h:56`) and no weight is block-shaped — so a short `k` raises no shape error and drafts a structurally wrong block while the tokens keep flowing. A literal port does NOT close it: `dspark_block_size` appears in no pinned file but `speculative.py`, and neither published Qwen3 draft sets it (`deepseek-ai/dspark_qwen3_4b_block7` and `RadixArk/Qwen3.8-27B-DSpark` @ `85ef153b` both carry `block_size: 7`, no `n_predict`), while upstream's `block_size` normalization at `:945-961` is Gemma4-only — so upstream accepts `k=6` on both sides of vllm#52197. Closing it for the lane we ship needs `block_size` as the floor fallback, one tracked divergence argued in the spec and the commit | bug | | [#1226](https://github.com/mudler/vllm.cpp/issues/1226) | `ENG-CUDAGRAPH-DEDUP` | The owed device A/B for #1162 RAN on `dgx:gpu0` 2026-08-18 (GB10, driver 580.173.02, nvcc 13.0.88, `rc` job f88d484b, gated commit `72de552c8`, whose four dedup sources are byte-identical to the merged `2a976eb9f`) and it SPLIT. CORRECTNESS PASSED and closes [#1184](https://github.com/mudler/vllm.cpp/issues/1184): 12/12 cells exit 0, zero `invalid device function` and zero `engine-fatal` where the pre-fix head `e4ce5571a` died after exactly one replay, ON replays as often as OFF (60=60, 33=33, 43=43), and `--output-token-ids` is IDENTICAL over 10/10 comparisons with the three OFF/OFF controls passing FIRST and three workloads hashing to three DIFFERENT values. THE BENEFIT IS REFUTED for the case the row was filed for: `N == M` in every `VT_CUDA_GRAPH_DEDUP=1` cell — 3 graphs to 3 execs on sizes [24 16 8], 2 to 2 on [16 8], 2 to 2 on [32 24] — with the registry count CLIMBING 1→1, 2→2, 3→3, so more than one capture reached it and the 1:1 is a measurement rather than the single-capture artefact of the first attempt. The cause is structural and was pre-registered before the run: `AppendKernelPayload` hashes (`func`, `gridDim.{x,y,z}`, `blockDim.{x,y,z}`, `sharedMemBytes`) at `src/vt/graph_dedup_runtime.h:121-128` and the memcpy payload hashes the copy extent, so the padded batch dimension is IN THE KEY, no candidate group forms and `cudaGraphExecUpdate` is never attempted. That contradicts the row's own premise, and SGLang keys the same fields (`cuda_graph_dedup_mixin.py:105-114`), so whatever folds upstream is not decode buckets either. NO throughput or memory number is recorded — clocks unpinned AND the ON arm allocated exactly as many executables as OFF. `VT_CUDA_GRAPH_DEDUP` therefore stays OFF and the default flip is NOT JUSTIFIED on this evidence. THE OPEN HYPOTHESIS this issue owns, deliberately not decided by the record that filed it: a COARSER key keeping the function addresses and topology but dropping the launch dimensions and memcpy extents would let two padded buckets form a candidate group at all, and the probe-before-fold design means a wrong grouping costs one wasted `cudaGraphExecUpdate` probe and a private executable rather than a wrong replay, so it is a cost question needing its own spec, a red-first discrimination test and a device probe-refusal rate. Also owed here: whether ANY other capture site or two models sharing the process-singleton registry can fold (only the Qwen3 dense decode driver was exercised), and a supporting `orin:gpu0` leg, BLOCKED because the Jetson 540.4.0 driver cannot run a CUDA 13 runtime (`cudaGetDeviceCount err=35`; a CUDA 12.x toolkit is the untried route). Honest gaps carried into the record: per-shape replay counts are unavailable, so B's ~30-per-shape is arithmetic over a printed TOTAL; the driver's "N captured size(s)" counts SLOTS not captures; the container's own cuBLASLt was never re-tested at CUDA 13.0. Full evidence [`benchmark-record.md`](benchmark-record.md) entry `ENG-CUDAGRAPH-DEDUP W4`; spec [`eng-cudagraph-dedup.md`](specs/eng-cudagraph-dedup.md) `## Outcome` | perf | | [#1229](https://github.com/mudler/vllm.cpp/issues/1229) | `ENG-CUDAGRAPH-DEDUP` | `origin/main` went RED at `fe24a3029` because `tests/scripts/test_check_gate_commands.py:670` pinned the literal string `Device byte-identity A/B (owed` inside `.agents/specs/eng-cudagraph-dedup.md`, and that commit legitimately rewrote the sentence from recording the device A/B as owed to recording that it RAN and what it found. The record edit is correct and the assertion described a state the record left, so every branch that merged `origin/main` inherited the failure. The protected property is stated by the assertion's own comment: the row is credited in `RUNNABLE_BASELINE` for `ctest -R test_graph_dedup` and `./scripts/agent-preflight.sh`, both CPU-tier, while gate item 6 of the spec contributes NO runnable command, so the credit covers the CPU tier and the record must not go silent about the leg it does not cover. That property is STATE-INDEPENDENT: `owed` was the leg's value on the day, not the property, and silence is the defect whether the leg was never run or the line was deleted after it ran. Asserting the NEW wording would only move the pin from `(owed` to `RAN 2026-08-18, PASS`, and the rerun this row already owes under [#1226](https://github.com/mudler/vllm.cpp/issues/1226) against a coarser signature key would red it again. FIXED by keying on structure: `gate_items`, `item_lead` and `gate_disposition` land in `scripts/check-gate-commands.py`, and the test now asserts that exactly one gate item names the device byte-identity A/B as its SUBJECT, that `runnable_commands` extracts nothing from it, and that its bold lead declares a disposition from a closed vocabulary. Lead-scoping is measured rather than assumed: a whole-item search credits item 5 with a disposition it does not declare, because its body contains `fail` in the phrase "proving the instrument can fail". Both historical wordings of item 6 are pinned as fixtures, which is what proves the new assertion is state-independent rather than re-pinned to today's sentence. The rule is NOT swept over every gate item and the reason is measured, not assumed: items 1 and 5 yield no command and declare no disposition, so a sweep is red on arrival, and widening the vocabulary until they pass would make it match ordinary English. That general rule is carried under `## Owed`. Also established and NOT a defect: `ERROR: these baseline rows left the gated population: ROW-THAT-IS-NOT-THERE` on the same run is the expected negative-path output of `test_check_mode_is_never_silently_swallowed_by_json`, confirmed by running that fixture alone. Spec [`fix-gate-commands-prose-pin.md`](specs/fix-gate-commands-prose-pin.md) | bug | diff --git a/.agents/parity-ledger.md b/.agents/parity-ledger.md index 68e733335..d30c8b9e1 100644 --- a/.agents/parity-ledger.md +++ b/.agents/parity-ledger.md @@ -940,3 +940,5 @@ Columns: | 2026-08-10 (**`FUSION-DENSE-MIGRATE` — the merged-GEMM allowlist's no-blocker half CLOSED**; issue [#299](https://github.com/mudler/vllm.cpp/issues/299), spec `specs/fusion-dense-migrate.md`; `row/FUSION-DENSE-MIGRATE`, base `origin/main` `c70f42b9` (rebased from the `e17e8f8a` pinned at worktree creation, via `688eea12`), isolated worktree `/home/mudler/_git/vllm.cpp-fusion-fold`; CPU-only, `benchmark_binding=false`) | Routes the gate/up of the five plain bf16 dense SwiGLU MLPs that bypassed the MUST-route merged-GEMM seam with no stated blocker — `commandr`, `glm4`, `minicpm`, `minicpm3`, `phi3` — through `layers::UnquantizedMlpGateUpMethod`, and deletes their five `pending FOLD-MIGRATE` entries from `scripts/merged-gemm-consistency-allowlist.txt` (11 → 6). The six survivors all name a blocker that needs the SHARED LAYER extended (f32 activations + rank-1 biases, ggml block-quant vs the explicitly UNQUANTIZED arm, up-front device staging vs `OwnedTensor`/`ResidentWeight` residency, a GeGLU arm on the SwiGLU-only grouped MoE op for `gemma4_moe`, an NVFP4-Marlin-resident arm for `laguna`'s raw-device-pointer w13) and were NOT touched: extending the seam is its own row. The `gemma4_moe` and `laguna` reasons were REWRITTEN in this branch: both previously read as a bare `pending fold`, which is the vague reason this row's own thesis says let the allowlist grow. Direct Unquantized arm, not `MakeMlpGateUpMethod` — no loader for the three `Qwen3DenseMlpWeights` models ever populates `*_fp4`. Shared-layer TUs (`linear.h`/`nvfp4.h`) and every `vt::` op UNTOUCHED. No env flag: the fold is unconditional and bit-exact, so a rollback arm would be dead code. | Upstream expresses all five as ONE `MergedColumnParallelLinear(gate_up_proj, output_sizes=[I]*2)` + `SiluAndMul()` at pin `555967922` (0.26.0.dev0): `commandr.py:91,102-108,116`; `glm4.py:46` → `llama.py:79,92-99,113`; `minicpm.py:193,204-211,219`; `minicpm3.py:186` (inherits `MiniCPMMLP`); `phi3.py:10` → `llama.py:79`. The fold moves us TOWARD that shape; no new upstream code was ported. | **BYTE-EXACT BY CONSTRUCTION, and that construction is what was gated** (a routing change, so the bar is op-sequence identity, never a tolerance). Each replaced body was the seam's own `{ResidentWeight; MatmulBT[2I,H]; SiluAndMul}` with `M` spelled `T`, and every call site passes a `DBuf{T,H}` so `x.shape[0] == T` identically. NEW byte-exact CPU case in `test_linear_method` for the DIRECTLY-constructed arm at BOTH the decode shape (M=1) and a prefill shape (M=4) vs the standalone sequence — RED-first PROVEN (mutating the seam's `SiluAndMul`→`GeluAndMul` fails 2 cases / 11 assertions; GREEN byte-for-byte on restore) → **6 cases / 76 assertions GREEN**. `check-fusion-consistency.py` RC=0 with 0 drift; two NEW mutation cases in `tests/scripts/test_check_fusion_consistency.py` make BOTH regressions RED (re-allowlisting a folded stem; reverting a fold) → **20/20**. MECHANICAL CORRECTION to issue #299: a fully-folded TU has NO hand-call left, so it leaves the detector's DENOMINATOR rather than entering its numerator — the checker reads `15 scanned / 6 routed / 11 allowlisted` → `10 / 6 / 6`, not the `11 routed / 15` the issue predicted; drift is 0 either way, which is what the gate asserts. Full CPU `ctest -j 6` on the final rebased base `60e71a0e`: **369/369, 0 failed**, 1288.75 s. RE-RUN on the MERGED head `4b99cefb` (after merging `origin/main` `5812b8b6` and adding the new coverage binary): **369/370, 1 failed**, 1441.85 s — the single failure is `test_engine_core_proc`, a STARVATION FLAKE and not a regression: serially on the same binary it is **10/10 cases, 93 assertions GREEN**, whereas under `-j 6` at load average ~170 (two other worktrees running their own suites) it failed in 0.06 s on `CHECK(abort_seen)` (`test_engine_core_proc.cpp:345`) having spun to 1089 assertions waiting for the abort. It is one of the four known starvation-prone binaries, loads none of the five folded TUs, and the other three passed in this run. `test_dense_gate_up_seam_forward` passed inside the suite (32.96 s), proving the new binary is registered and wired. An earlier run on the intermediate base read **367/369**, and both numbers are reported rather than the red one deleted: `test_async_llm` and `test_openai_conformance` failed while a SECOND worktree ran its own suite at load average 89-122, and re-run alone on a quiet box (load 12.8) passed in **0.04 s** and **20.76 s** (vs 605 s of `statuses == -1`, the client's 30 s read timeout, under load); `test_openai_conformance` builds a synthetic `Qwen3_5MoeForConditionalGeneration` in-process and loads none of the five folded TUs, so it could not be this row's, and the clean 369/369 settles it. The #274 ASan/UBSan five did not appear in either run — this gate is Release with no sanitizer. **OWED / NOT CLAIMED:** the five paged-engine SACRED gates (`test_{commandr,glm4,minicpm,minicpm3,phi3}_paged_engine`) are checkpoint-gated dgx-only and emit a loud SKIP on a CPU box, so empirical token-exact confirmation is PENDING for the next GPU holder — the same shape A1 carried for Granite/StableLM — and is tracked by [#337](https://github.com/mudler/vllm.cpp/issues/337) so the handle outlives #299. EXECUTED CPU COVERAGE OF THE FOLDED TUs (added on review finding F4, which proved that mutating `phi3`'s `I` to `I - 1` at the call site survived 176 CPU tests): NEW `tests/vllm/models/test_dense_gate_up_seam_forward.cpp` drives the REAL forward of four of the five folded TUs (`commandr`/`glm4`/`minicpm`/`phi3`) over synthetic in-memory weights with no checkpoint and no GPU — **4 cases / 1940 assertions GREEN** — pinning the gate/up split analytically (zeroing the UP half and zeroing `down_proj` both make the MLP contribute exactly zero, so they must be BYTE-IDENTICAL, which holds only at a split of exactly `I`) plus the half order and a vacuity guard. RED-first EXECUTED both ways: `I` → `I - 1` at all four call sites fails all four cases (`phi3` alone fails only its own, so each TU is genuinely reached), and a SHAPE-PRESERVING mutation of `vt::SiluAndMul`'s CPU kernel (`silu * up` → `silu + up`) fails all four on the split assertion itself while `test_linear_method` stays GREEN at 76/76 — a defect class the seam unit test structurally cannot see. Both restored byte-for-byte (md5 verified). This also retires the row's original "no e2e evidence is possible on a CPU box" framing as too strong: ORACLE evidence needs the GPU, SELF-CONSISTENCY evidence does not. `minicpm3` is the one folded TU that harness does not drive (MLA + load-time `kv_b_proj` absorption belong to the DeepSeek-V2 synthetic harness). Also REJECTED as evidence: an object-code A/B — compiling each of the five TUs from main's source and from the folded source with the identical production command and diffing the disassembly gives 3689-6888 differing instruction lines per TU, which is GCC re-allocating registers around a header-inlined method, not a numerical signal in either direction. STATED RATHER THAN HIDDEN: the OP sequence is identical but the ALLOCATION sequence is not — `DBuf` is pooled and returns its block on destruction (`dense_device_glue.h:99`), so the `[T,2I]` gate_up buffer is now released when `Apply` returns instead of at the end of the MLP block and the `[T,H]` output may reuse it. That cannot move a value (`vt::MatmulBT` writes every output element from a fresh f32 accumulator rather than accumulating into `out`) and peak pool usage is unchanged or lower; it is the same shape every earlier fold produced. ALSO REPOINTED: `glm4`/`phi3` sit on the OTHER (glue) allowlist reading `pending FUSION-DENSE-MIGRATE`; closing this row would leave both aimed at closed work, so their reasons now name [#314](https://github.com/mudler/vllm.cpp/issues/314), which owns the glue half — the glue fold itself is NOT done here. | | 2026-08-11 (`ENG-TRAILER-MERGE-ARTIFACTS`; implementation merge `157080c8`, issue [#406](https://github.com/mudler/vllm.cpp/issues/406)) | The trailer gate was rejecting CORRECT commits. `check-commit-trailers.py` reads through `git interpret-trailers --parse`, which treats only the FINAL paragraph as the block, and GitHub appends `Co-authored-by:` as a separate trailing paragraph on squash merge — so a complete block became invisible and 13 of the last 30 commits on `main` failed the check, unnoticed because those runs were cancelled (#274). Fixed by fusing consecutive trailing TRAILER-SHAPED paragraphs before parsing; no rule relaxed, and a prose paragraph still terminates the block. | `test_check_commit_trailers.py` 21/21 — the RED-BEFORE appended-co-author case plus four guards (doubled block, contradictory declarations, no-trailer merge message, prose after the block) green before and after; the five real `main` commits re-verified per commit and the verdict table reproduces. | Closes ONE of five shapes. `f64f2b71` (bot co-author) is a REAL violation the parse had hidden and now correctly fails; `87308dea`, `b8293c88` and `b580452d` stay red by design — merge-method artifacts, not checker defects. A first attempt that also collapsed identical duplicate trailers was REVERTED in full: it relaxed a uniqueness rule an existing test pins. | | 2026-08-12 (**Audio ENCODER TTFT — FA-2 TENSOR CORES for the hd-64 non-causal encoder attention; 5.50x encoder forward / 115.8x kernel, LANDS OPT-IN because it costs precision**; `CLAIM-MM-SPEED-AUDIO-ENC-FA2`, multimodal-speed.md §17, issue [#432](https://github.com/mudler/vllm.cpp/issues/432); worktree `/home/mudler/_git/vllm.cpp-mm-speed` branch `row/MM-SPEED-ENC-FA2` off `dc7a1392`) | New dense hd-64 non-split FA-2 instantiation (`flash_fwd_hdim64_bf16_sm80.cu`, upstream's own `run_mha_fwd_hdim64` via the plain batch entry — both firsts here) + `LaunchDenseFA2Bf16` (b=1, null `cu_seqlens_q`) + additive op `vt::AttentionDenseFa2`, routing the Whisper encoder self-attention to the kernel vLLM itself dispatches (`whisper.py:298-317` -> `flash_attn_varlen_func`). dgx GB10 sm_121a, all three mandatory banners CONFIRMED, `-Werror` 0-warn, GPU under `flock` on a CONTENDED box (3 other agents' jobs queued). **SPEED:** encoder attention 19,278 -> 166.5 us/layer (115.8x, nsys 32 inst both arms, SAME tool); encoder forward 731.7 -> 133.0 ms (5.50x, non-overlapping bands, rep0 dropped) => the ENCODER FORWARD goes from 15.90x to **2.89x** of vLLM's whole TTFT against the PIN `555967922` (TTFT median 46.02 ms, 45.60-46.41, production/graphed, 6 reps rep0 dropped, re-measured 2026-08-12). NOT a TTFT ratio: our projector, merge and prefill are unmeasured (F3). The originally published `~17.1x -> ~3.11x` used the carried-forward 0.25.0 denominator (42.8 ms), which the pin is 7.5% faster than - so the published figure was CONSERVATIVE. Oracle asserted BY COMMIT: the venv's `0.23.1rc1.dev1511+g555967922` version string is a setuptools_scm nearest-ancestor-tag artefact, not a mismatch; `soundfile==0.14.0` had to be installed into `~/venvs/vllm-oracle-next` before the pin could tokenize Voxtral at all (recorded against #375). **CORRECTNESS:** oracle teacher-force (fixture's own stack, asserted live as vLLM 0.25.0 + mistral_common 1.11.5) = 3 divergent, worst 0.1250 nats, over-band 0, **RESULT PASS** — inside the ratified band but the shipping kernel has 0 divergent at gap 0.0. Cause is a HYPOTHESIS, not grounded (corrected 2026-08-12 by the fresh review of PR #439): the stated cause - FA-2 converting P to bf16 before the PV MMA (`flash_fwd_kernel.h:347`) - was REFUTED by mutation M4, which forced exactly that rounding into the shipping scalar kernel, rebuilt clean, and got the SAME token md5 `89923566...`. Five candidate differences remain and none is isolated. **DISPOSITION: default UNCHANGED (byte-exact, `test_voxtral_e2e` 16/16, goldens md5 unchanged); FA-2 opt-in behind `VT_WHISPER_ENC_FA2=1`; ADOPTION = DEVELOPER DECISION.** Not a ceiling, RE-RANKED after M4: (1) attention is now only 5.33 ms of the 133 ms, so the §15.1-deferred device im2col kernel is the top encoder lever; (2) measure our ACTUAL TTFT (projector + merge + prefill); (3) isolate which of the five FA-2/scalar differences flips the tokens, one M4-style mutation at a time; (4) the FA-3-style f32-correction rescale, DEMOTED from #1 because M4 refuted its premise. Repairs landed on the branch: issue [#463](https://github.com/mudler/vllm.cpp/issues/463) for the zero-assertion SUCCESS skip (now `exit 77` + CTest `SKIP_RETURN_CODE`), `tests/vt/test_ops_attention_dense_fa2.cpp` for the op's missing tests, and a `causal` parameter on `LaunchDenseFA2Bf16` that THROWS instead of silently answering non-causally. No mm row advances to DONE. | +| 2026-08-12 (**M4 on gfx1100 — the pinned vLLM-ROCm oracle stood up on the only board that can host it + the ROCm device-golden lane in the Qwen3-0.6B near-tie gate; issue #41; `CLAIM-ROCM-M4`; worktree `/home/vikash/vllm.cpp-rocm-m4` branch `row/ROCM-M4-GATE`, base `8b00f79f`) | The M4 milestone for `BACKEND-ROCM`: build the PINNED vLLM oracle (`555967922`, runtime `0.23.1rc1.dev1511+g555967922` — the parity-pin block's identity, verified) as a ROCm wheel inside `rocm/vllm-dev:base` for `gfx1100` (`vllm-...rocm723-cp312` wheel, 2,696 files incl. `_rocm_C` ABI), commit the image (`vllm-rocm-oracle:555967922-gfx1100`, recipe in `.agents/specs/rocm-m4-oracle.md`), and capture the Qwen3-0.6B goldens FROM THAT ORACLE on this box. The gate test's existing Metal/Tenstorrent device-golden lane gains the `kROCM` arm (`our_ids_rocm.npy` + `neartie_gap_mnats_rocm.npy` + the `*_rocm` greedy/evidence pair); base dgx goldens byte-untouched. | **Oracle = upstream vLLM AT THE PIN ITSELF** — not a port: `scripts/qwen3-oracle-capture.py` + `scripts/qwen3-neartie-gap.py` run against the pinned wheel in the container; the gate lane mirrors the Metal M3b precedent (PR-set around row 632/633) file:line (`test_qwen3_paged_engine.cpp` device-golden branch). | **PASS.** Oracle K=10 per-prompt greedy: deterministic in every (prompt,pos) cell (0 multi-member). Our ROCm engine vs the same-box oracle: **16/16 prompts PASS** — STRICT token-exact 11/16 vs base greedy, 5/16 near-tie-band only, **max teacher-forced gap 0.125 nats** (prompt[3] tok1), 0 forward-divergent, 125/125 assertions; backend proof all 8 Qwen3-dense ops `selections>0 ∧ declines==0` on device 5 (kPagedAttention 7,168 selections; RoPE served via the fused-cache arm, `rope_alt` counted). The 28 token-divergent positions vs oracle greedy are all in-band; the p0 France/Italy flip is a literal 0.0000-nat tie in the oracle's OWN logits (vLLM's teacher-forced argmax on our prefix IS 15344 — the Metal M3b signature reproduced on ROCm). Hardware: 4x RX 7900 XTX gfx1100, ROCm 7.14 host, HIP 7.2.5 userland, `enforce_eager=True`, batch=1, GPU under `flock ~/gpu.lock`. Not a speed change (`benchmark_binding=false`). NEXT: Qwen3.5-0.8B (GDN) + 35B keep-quant gates under the same oracle. | +| 2026-08-12 (**M4 second model — the Qwen3.5-0.8B GDN oracle-backed gate, landed GREEN-shaped after review; the pre-fix RED capture is evidence, not goldens**; issue #41; `CLAIM-ROCM-M4`; branch `row/ROCM-M4-GDN-GATE` stacked on the dispatch fix) | NEW `tests/parity/goldens/qwen35_greedy_0_8b/` + NEW `tests/parity/test_qwen35_paged_engine.cpp` — the GDN sibling of the Qwen3-dense SACRED gate: 16-prompt battery through the full paged engine, anchor + <=0.5-nat near-tie band, GDN backend proof (15 ops, selections>0, declines==0). Goldens captured from the pinned vLLM-ROCm oracle on gfx1100 (K=10 deterministic in every cell) with a `manifest.json` (oracle identity + pinned HF revision `2fc06364` via `parity::Qwen35_08BSnapshot`). The gate FAILS SAFE: exit-77 (CTest Skipped) off-ROCm and when the checkpoint is absent (issue #463's pattern), never a false green. | Oracle = upstream vLLM at the parity pin itself; gate methodology = the Qwen3-dense [[near-tie-distributional-gate]]; the review sweep's three findings (runner-less manifest op -> `PendingRunnerOps()`; device fail-safe; green-shaped landing with the fix stacked below) are recorded in `.agents/specs/rocm-m4-oracle.md`. | **16/16 PASS with the fix stacked** (15/16 strict token-exact vs the oracle, 1/16 near-tie band, max gap 0.125 nats, 0 forward-divergent; 0.6B gate re-run green; backend proof kPagedAttention 1,536 / kGdnDecode 4,320 selections, 0 declines). The PRE-FIX capture — 13/16 forward-divergent, max gap 14.125 nats, first-token divergences on 6 prompts (0.375-1.062 nats, 5 of 6 over band) — is retained as evidence here; the committed goldens are the FIXED engine's oracle-re-derived pair, so the gate is green from the commit that introduces it and cannot hide a reintroduction (anchor drift = hard REQUIRE naming REGRESSION). `benchmark_binding=false`. | diff --git a/.agents/specs/rocm-gdn-08b-fix.md b/.agents/specs/rocm-gdn-08b-fix.md new file mode 100644 index 000000000..1078b7eac --- /dev/null +++ b/.agents/specs/rocm-gdn-08b-fix.md @@ -0,0 +1,98 @@ +# ROCm Qwen3.5-0.8B forward-divergence fix — the RED gate's target + +## The problem (oracle-measured) + +The M4 gate (`test_qwen35_paged_engine`, #559) FAILS 13/16 prompts against the +pinned vLLM-ROCm oracle on gfx1100, with **first-token divergences on 6 prompts** +(gaps 0.375–1.188 nats, over the 0.5-nat band) and max gap 14.125 nats deep in a +diverged continuation. This is a REAL forward divergence, not a bf16 near-tie. + +## The evidence (all measurements on gfx1100, GPU-locked) + +1. **The split is ROCm-kernel-specific, not model math.** For prompt + `The capital of France is` (p0): + - pinned oracle (vLLM `555967922`, enforce_eager): argmax **11751 " Paris"** (-1.9957), token 25 ":" at -2.8082 (#4) + - **our CPU backend: matches the oracle exactly** (` Paris.\nThe capital of France is`) + - our ROCm backend: argmax **25 ":"** (-1.6535), 11751 at -3.0910 (#9) + The candidate SETS are the same with similar spread; the rankings shift ~1.1 + nats in opposite directions — a systematic numerical offset in some block + output, not a distribution reshape. +2. **Backend proof is green**: all 15 dispatched GDN/full-attn ops run natively + on device 5 with 0 declines (kPagedAttention 1,536, kGdnDecode 4,320) — the + divergence is numerical, not dispatch. +3. **Not the ROCm paged-attention kernel**: `VT_ROCM_ATTN_CPU_REF=1` (the CPU-ref + attention path) produces the SAME divergent tokens. +4. **GDN ops were cleared in isolation**: the merged cross-device cases + (tests/vt/test_backend_cross_device.cpp) pass at real dims vs the CPU oracle. +5. **The 0.6B dense gate passes 16/16** with the same ROCm attention/GEMM/RMSNorm + kernels — so the diverging op is exercised by 0.8B's GDN hybrid path but not + (or differently) by dense 0.6B. +6. Prior characterization (issue #41, op-level): attention-block localized + (layer-3 attn block 20% gated / 4.4% raw mismatch); GDN ops + op 67 + op 66 + + paged-attention + GEMM all cleared at real dims; conclusion then was + "bf16-softmax amplification" — but (3) refutes the paged-attention softmax as + the sole source, so the amplification must be fed by an earlier drift. + +## The plan (W-plan) + +- **W1 — instrument.** Add `VT_DUMP_ACT`-style per-layer hidden-state dumps to + `qwen3_5.cpp` (the deepseek_v4.cpp precedent), plus a scratch host comparator + (CPU vs ROCm, per-layer max-abs/rel delta at the residual stream) in + `/tmp/vllm-diag` (never committed). +- **W2 — localize.** Find the first layer+op where the ROCm hidden state drifts + past ~1e-3 rel vs CPU on the p0 prefill. Suspects in order: the GDN + prefill/decode recurrence IN CONTEXT (state carry across the chunk boundary), + the fused preamble `kAttnQkNormRopeGate` at 0.8B geometry, the conv1d state + handoff, the GDN output gated RMSNorm. +- **W3 — fix.** The minimal kernel fix at the localized op (portable-math + corrections first; e.g. the reduction-order or accumulation-dtype issue the + in-isolation tests couldn't see). +- **W4 — gate GREEN.** `test_qwen35_paged_engine` 16/16 (strict + band), the + 0.6B gate re-run green, focused + full gates, fresh review, operator gate. + +## RESULT (W1–W4 complete — fix verified; the gate lands GREEN stacked on top) + +**Stack shape (post-review):** this branch carries ONLY the kernel fix + its +teeth + the debug instrumentation. The M4 0.8B gate with the GREEN (fixed-engine, +oracle-re-derived) goldens lands in the stacked sibling commit on +`row/ROCM-M4-GDN-GATE` — so every commit in the stack is individually green on +the gate host, and the gate never exists in a state whose goldens its own tree +cannot reproduce. + +**Root cause:** `AttnQkNormRopeGateKernelRocm` (src/vt/rocm/rocm_gdn_fused.hip) +dispatched its template on the SOURCE dtype (`qgate.dtype`) instead of the +OUTPUT dtype like the CUDA lane (`LaunchAttnPreambleOut` switches on +`q_out.dtype`). The 0.8B bf16 model runs bf16 QKV-projection output + f32 +q/k/gate outs (the f32-attention path — no FA-2 on ROCm), which the src-keyed +dispatch silently mis-launched as all-bf16: the kernel wrote bf16 bits through +the f32 out pointers. The cross-device test covered only f32-src, so the bug +class was invisible in isolation; the in-context logits comparison (W2) showed +`fa0_q` rms-rel 1.196 with `fa0_qkv`/`fa0_gate` clean, localizing it to the +preamble's q/k path. + +**Fix:** dispatch on the output dtype (bf16 src + f32 out is now a first-class +combo); VT_CHECK bf16-out requires bf16-src. + +**Teeth:** the cross-device AttnQkNormRopeGate case gains the bf16-src/f32-out +combo at the real 0.8B dims (Hq=8, Hkv=2, Dh=256, rot=64) — FAILS with the fix +reverted (mutation-proven), passes with it. + +**Gates (gfx1100, flock):** +- `test_qwen35_paged_engine` (the RED anchor): **16/16 PASS — 15/16 strict + token-exact vs the pinned oracle, 1/16 near-tie band, max gap 0.125 nats, + 0 forward-divergent** (was 3/16, 13 divergent, max 14.125 nats) +- `test_qwen3_paged_engine` (0.6B, regression): 16/16 PASS unchanged +- `test_backend_cross_device`: 19/19, 346 assertions +- e2e: `The capital of France is` → ` Paris.\nThe capital of France is`, + matching the pinned oracle AND our CPU backend token-for-token + +**Also fixed (pre-existing main bug, blocked the full HIP build):** +same-line double `CAPTURE` in test_qwen3_5_gdn_spec_routing.cpp:357 (doctest +redefinition that clang/hipcc rejects). + +## Boundaries + +- The RED gate stays RED until W4; no gate-weakening (no band-widening, no + goldens recapture against a "fixed" engine without the oracle re-derivation). +- ROCm-only files + the additive dump hook; no shared-path behavior change. +- Rollback: the dump hook is env-gated (off by default). diff --git a/.agents/specs/rocm-m4-oracle.md b/.agents/specs/rocm-m4-oracle.md new file mode 100644 index 000000000..d91dd1aaa --- /dev/null +++ b/.agents/specs/rocm-m4-oracle.md @@ -0,0 +1,253 @@ +# ROCm M4 — the pinned vLLM-ROCm oracle on gfx1100 and the ROCm near-tie gate lane + +## Goal + +Give `BACKEND-ROCM` (issue #41) the milestone that only this box can host: the M4 +correctness gate — our ROCm paged engine held against a **pinned vLLM-ROCm oracle +running on the same gfx1100 hardware**, not against the dgx (CUDA) capture. + +## The oracle + +- **Pin:** upstream vLLM commit `5559679229bc961848b121ccdeaa8fa5d79bec98` + (runtime identity `0.23.1rc1.dev1511+g555967922`, verified against the + parity-pin block in `.agents/upstream-sync.md`). +- **Base image:** `rocm/vllm-dev:base` (HIP 7.2.5 userland, host ROCm 7.14 + driver). `PYTORCH_ROCM_ARCH=gfx1100` covers the 7900 XTX. +- **Build:** `python3 setup.py bdist_wheel` inside the container after + `apt-get install binutils mold` (the base image ships no linker) and the + pyproject build-system requirements. Wheel: + `vllm-0.23.1rc1.dev1511+g555967922.rocm723-cp312-cp312-linux_x86_64.whl` + (2,696 files; `_C`/`_rocm_C`/`_moe_C` ABI extensions present). +- **Committed image:** `vllm-rocm-oracle:555967922-gfx1100` + (`/home/vikash/oracle/Dockerfile.oracle`), wheel installed with deps. +- **Determinism:** K=10 per-prompt greedy capture on the 16 gate prompts is + **deterministic in every cell** (0 multi-member (prompt,pos) cells) — a + well-posed strict gate on this board, matching the dgx finding for batch=1 + `enforce_eager=True`. + +## The gate lane + +`tests/parity/test_qwen3_paged_engine.cpp` already has device-aware goldens for +Metal and Tenstorrent (`our_ids_.npy` + `neartie_gap_mnats_.npy` under +the shared `qwen3_greedy_0_6b/` dir, with the dgx CUDA pair as base). This spec +adds `kROCM` to that lane, unchanged logic: + +- hard anchor REQUIRE (our deterministic tokens vs the committed ROCm anchor), +- near-tie band ≤ 500 milli-nats (vLLM teacher-forced on OUR exact prefix), +- backend proof (all 8 Qwen3-dense ops `selections>0 ∧ declines==0` on kROCM, + with the fused-RoPE alternative counted as in the existing lane). + +New committed goldens (captured on gfx1100, 4x RX 7900 XTX, ROCm 7.14, +`enforce_eager=True`, batch=1): + +| file | content | +|---|---| +| `our_ids_rocm.npy` | our engine's greedy tokens (16×16) | +| `neartie_gap_mnats_rocm.npy` | oracle teacher-forced gaps in milli-nats (16×16) | +| `greedy_ids_rocm.npy` | the ROCm oracle's own greedy (evidence; the base `greedy_ids.npy` stays the dgx capture) | +| `greedy_dist_rocm.npy` | K=10 run set (evidence) | + +Flow (the three commands, all GPU-locked): +1. oracle capture: `scripts/qwen3-oracle-capture.py --runs 10 --per-prompt` in the + committed container against `/models` (Qwen3-0.6B). +2. bootstrap dump: `VT_DUMP_IDS=1 ./build-hip/tests/test_qwen3_paged_engine` + → `our_ids_rocm.i32`. +3. teacher-forced gaps: `scripts/qwen3-neartie-gap.py --golden-dir ...` in the + container → `our_ids_rocm.npy` + `neartie_gap_mnats_rocm.npy`. + +## Result (Qwen3-0.6B, gfx1100) + +**16/16 prompts PASS** — STRICT token-exact 11/16 vs the base greedy, 5/16 via +the near-tie band, **max gap 0.125 nats** (prompt[3] tok=1), **0 +forward-divergent**; backend proof: kPagedAttention selections 7,168, 0 declines; +125/125 assertions. The 28 token-divergent positions vs the oracle's own greedy +are all within the band; the known p0 France/Italy flip is a literal 0.0000-nat +tie in the oracle's own logits. + +## Boundaries + +- No source/kernel change; the only edited file is the parity gate test. +- The dgx base goldens are byte-untouched; the ROCm additions are purely additive. +- Qwen3-4B (the strict deterministic dense) is not on this box (disk); the 0.6B + near-tie-robust gate is the M4 evidence for the lane. +- The oracle image and `/home/vikash/oracle/` scratch are machine-local, not + committed; this spec is the reproduction recipe. + +## Review rework (2026-08-13 sweep, localai-bot on #559) + +The first shape landed the 0.8B gate RED with the pre-fix goldens committed. +Review findings, all accepted and fixed in the current shape: +1. **CI-red manifest**: `qwen35_0_8b_greedy` had no runner — added to + `PendingRunnerOps()` in test_op_parity.cpp (the `qwen36_gguf_greedy` + precedent); the manifest itself stays (its oracle identity is the point). +2. **Fail-safe by device**: the gate now exits 77 (CTest Skipped) on any + non-ROCm device instead of comparing a foreign engine against ROCm-derived + goldens, and the checkpoint-absent path exits 77 too (issue #463's pattern) + instead of a silent `return` that printed SUCCESS with 0 assertions. +3. **Green-shaped landing**: the committed goldens are the FIXED engine's + sequence, oracle-re-derived (the kernel fix lands below this commit in the + stack); the RED capture (13/16, first-divergence gaps 0.375-1.062 nats — + 5 of 6 over band) stays as evidence here and in the parity ledger, not as + committed goldens that no code can pass. +4. **Anchor message honesty**: anchor drift now reads "REGRESSION SUSPECTED: + bisect the engine change first" — re-deriving goldens is the last step of a + justified re-capture, never the response to a failure. + +## Issue #1222: missing-artifact fail-safe repair + +### Scope + +The Qwen3.5-0.8B paged-engine gate promises that an unavailable prerequisite is +not a correctness pass. Repair the three required committed-artifact paths in +`tests/parity/test_qwen35_paged_engine.cpp`: `greedy_ids.npy`, `our_ids.npy`, +and `neartie_gap_mnats.npy`. Issue #1222 is an in-flow `BACKEND-ROCM` bug for PR +#559. Snapshot absence, non-ROCm execution, malformed present arrays, and the +numerical acceptance policy keep their existing behavior. + +### Design + +Put the three existence checks behind one prerequisite helper used by the real +gate and by a test-only golden-directory probe in the same executable. Missing +required files route through the existing `SkipGate` process exit 77. A +deterministic CMake subprocess test creates temporary golden directories, +removes each required file in turn, invokes the actual gate executable through +the probe, and rejects any child result other than 77. + +`VT_DUMP_IDS` remains the explicit bootstrap path when the greedy capture exists +but the anchor/gap pair is incomplete. It can run the model and write +`our_ids.i32`; after that bootstrap work it exits 77 because no correctness gate +ran. Present files continue into `LoadNpy` and the existing dtype/shape checks, +so malformed arrays remain hard failures rather than becoming skips. + +### Risks + +- A probe that reimplements the condition could pass while the real gate stays + unsafe. The probe therefore calls the same prerequisite helper in the same + process as the production test case. +- Treating any parse failure as unavailable would hide corrupt evidence. Only + filesystem absence is eligible for exit 77. +- Returning normally after bootstrap would preserve a second false-success + path; bootstrap completion must terminate through `SkipGate`. +- Artifact checks must happen before model construction only in probe mode; + the ordinary snapshot and device gates retain their ordering and behavior. + +### Tests + +1. Add the subprocess regression and capture semantic RED before the repair: + children missing greedy, anchor, and gap return 0 instead of 77. +2. Rebuild every affected target and require all three children to return 77. +3. Under `flock /home/vikash/gpu.lock` with + `LD_LIBRARY_PATH=/opt/rocm/lib:/opt/rocm/lib/llvm/lib`, run the real pinned + Qwen3.5 gate and require 16/16 prompts and 137 assertions, not exit 77. +4. In scratch, restore one repaired missing-artifact branch to a normal return, + freshly rebuild, and require the deterministic subprocess test to fail for + the corresponding child. Restore byte-for-byte, rebuild, and require green. +5. Run `PATH=/usr/bin:$PATH GIT_CONFIG_GLOBAL=/dev/null + scripts/agent-preflight.sh` and require literal `All gates green.` + +### Evidence + +- RED build: `flock /home/vikash/gpu.lock env + LD_LIBRARY_PATH=/opt/rocm/lib:/opt/rocm/lib/llvm/lib cmake --build build-hip + --target test_qwen35_paged_engine --parallel 4` exited 0. The new test then + ran with `ctest --test-dir build-hip --output-on-failure -R + '^test_qwen35_paged_engine_prerequisites$'`. CTest exited 8 because the + greedy, anchor, and gap children each exited 0 with zero assertions and + doctest `Status: SUCCESS!`. +- Focused GREEN: the same build command compiled the changed test object and + linked the HIP executable, then exited 0. The same CTest command exited 0. + Direct CMake-script execution exited 0 and printed child exit 77 for + `greedy_ids.npy`, `our_ids.npy`, and `neartie_gap_mnats.npy`. +- Real model gate: `flock /home/vikash/gpu.lock env + LD_LIBRARY_PATH=/opt/rocm/lib:/opt/rocm/lib/llvm/lib + ./build-hip/tests/test_qwen35_paged_engine` exited 0. It loaded snapshot + `2fc06364715b967f1860aea9cf38778875588b17`, ran on device type 5, passed + 16 of 16 prompts and 137 of 137 assertions, and reported zero declines. +- Mutation: the repaired greedy branch changed to a normal return. The source + hash changed from `7d7dfede853fe46746cbefec087130b17dc7e65b93376f28ff932ee32739bc8e` + to `27cbf5cba2bce704c2df737da9397f6d42afa92232f1c3bad0715ac745508116`. + A fresh rebuild exited 0 and changed the binary hash from + `a5fcd67f7d926efd90ce50732a6eb667fab045ccc0036248a9647f89cc1f3df4` + to `b758bcb7ab8af04d22c5f99dcefdc5c0d5db802f28d1543f79f755a8f33bed6d`. + CTest exited 8 because the greedy child exited 0. The other two children + still exited 77. Restoration recovered both original hashes, the same + tracked status, and CTest exit 0 after a fresh rebuild. +- Full gate: `PATH=/usr/bin:$PATH GIT_CONFIG_GLOBAL=/dev/null + scripts/agent-preflight.sh` exited 0 and printed literal `All gates green.` + +#### Reviewer repair loop for the complete probe + +- Complete-control RED: after the CMake regression created all three + placeholder artifacts, `flock /home/vikash/gpu.lock env + LD_LIBRARY_PATH=/opt/rocm/lib:/opt/rocm/lib/llvm/lib ctest --test-dir + build-hip --output-on-failure -R + '^test_qwen35_paged_engine_prerequisites$'` exited 8. The three + missing-artifact children stayed at exit 77. The complete child exited 0, + and doctest reported one passed case with zero assertions. +- Focused GREEN: a fresh `cmake --build build-hip --target + test_qwen35_paged_engine --parallel 4` under the same mutex and library path + exited 0. The same CTest command exited 0. Direct CMake-script execution + exited 0 and reported child exits 77, 77, 77, and 86. +- Real model gate: `flock /home/vikash/gpu.lock env + LD_LIBRARY_PATH=/opt/rocm/lib:/opt/rocm/lib/llvm/lib + ./build-hip/tests/test_qwen35_paged_engine` exited 0. The gate loaded snapshot + `2fc06364715b967f1860aea9cf38778875588b17`, ran on ROCm device type 5, + passed 16 of 16 prompts and 137 of 137 assertions, and reported zero + declines. +- Mutation: the complete-probe branch changed back to a normal return while + retaining a reference to the sentinel so the mutation compiled. The source + hash changed from + `7a2c3782b1546a8a90ba1599610e13068addeb631f7286e66dac7ce4ec1215f7` + to `b3de308eb7805d575248258eda60e2acbca8e2db9b8e25b8a16cfbacd88bb465`. + The fresh rebuild exited 0 and changed the binary hash from + `acf873bebec14804e9763bdee8c9c2e51dddd0b5949fb8cbff7091b8935bb931` + to `febc7bcc6e4cf81668651aff0c21bf3dc7be75c8593c7235dfc1f7da68e56947`. + CTest exited 8 because the complete child returned 0. The three + missing-artifact children still returned 77. Restoration recovered both + original hashes, and the fresh rebuild and CTest each exited 0. A prior + one-line mutation was rejected as evidence because `-Werror` stopped the + build on the unused sentinel before the regression ran. +- Full gate: `env PATH=/usr/bin:$PATH GIT_CONFIG_GLOBAL=/dev/null + LD_LIBRARY_PATH=/opt/rocm/lib:/opt/rocm/lib/llvm/lib + scripts/agent-preflight.sh` exited 0 and printed literal + `All gates green.` + +### Stop conditions + +Stop rather than weaken the gate if the deterministic test cannot exercise the +actual executable, if absent artifacts cannot be distinguished from malformed +present arrays, if the pinned local snapshot or ROCm gate cannot run under the +required mutex, or if any change outside this test, its registration, this +specification, and the issue-index append is required. + +### Outcome + +Issue #1222 was the two normal returns named by the finding. The greedy branch +returned directly, and one shared branch returned when either the anchor or gap +was absent. Both paths produced a zero-assertion doctest success. + +The gate now uses one existence-only helper. The helper exits 77 for each +missing required artifact and returns only when all three paths exist. Array +loading and dtype and shape checks remain after the helper, so a malformed +present array still takes the hard-failure path. `VT_DUMP_IDS` remains the one +bootstrap opt-in. Its model run and `our_ids.i32` write remain intact, but the +bootstrap path now exits 77 after the dump because it did not run correctness. + +The test uses an environment probe in the real executable instead of a second +artifact predicate. A separate validator was rejected because it could drift +from the gate. Exit 77 remains the existing CTest unavailable-gate contract. +No numerical threshold, snapshot pin, model setting, or device default changed. + +The reviewer found that the probe's complete-artifact path still returned +normally. That return bypassed snapshot resolution, model construction, and the +device check, so a test-only path could report production success. The +subprocess regression now includes one complete-artifact control. That control +requires the named test-only sentinel exit 86. Exit 86 is distinct from process +success 0 and the unavailable-gate skip 77, and it fits the portable process +exit range. + +After the shared prerequisite helper returns in probe mode, the executable now +prints that the test probe completed and that the production gate did not run. +It then exits 86. The missing greedy, anchor, and gap probes still exit 77. +Normal execution without the probe is unchanged and continues through snapshot +resolution, array validation, model construction, and the ROCm device gate. diff --git a/docs/FEATURES.md b/docs/FEATURES.md index 60e2c187a..0c1549a0e 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -266,7 +266,7 @@ both refuse, naming what is missing. | CPU (x86, Arm i8mm; A76 assembly correct/default, llama speed gate open, and the closed 20-core floor ran a SUPERSEDED fork denominator rather than the stock `b10451` pin, re-take owed #1003) | ✅ | ◐ | ☐ | ✅ | | Metal (Apple Silicon) | ✅ builds under Apple Clang with project warnings promoted to errors, the Qwen3.5 MoE loader included; its layout-refusal path uses the same messages and behavior on every platform (#1054) | ☐ | ☐ | ✅ | | Vulkan | ◐ | ☐ | ☐ | ✅ | -| ROCm | W0 verified on 5 gfx archs; dense and GDN run all-native. **M3: `ROCM_ATTN` registered and selected per attention group** (#1056/#1065, [spec](../.agents/specs/rocm-attn-backend.md)). CPU parity open (#269) | 44 registered ops including full GDN; ctest-green gfx1151/1103/1100/1201/1200 ([#41](https://github.com/mudler/vllm.cpp/issues/41)). APU managed allocation is unverified. [ROCM.md](ROCM.md) | ✅ | ✅ | +| ROCm | W0: 5 gfx archs; dense/GDN all-native; 0.8B dispatch fixed. **M4:** Qwen3-0.6B/3.5-0.8B 16/16 (#41). **M3:** `ROCM_ATTN` registered (#1056/#1065, [spec](../.agents/specs/rocm-attn-backend.md)). CPU parity open (#269) | 44 registered ops including full GDN; ctest-green gfx1151/1103/1100/1201/1200 ([#41](https://github.com/mudler/vllm.cpp/issues/41)). APU managed allocation is unverified. [ROCM.md](ROCM.md) | ✅ | ✅ | | XPU / TPU | ☐ | ✅ | ◐ | ☐ | | Tenstorrent Blackhole | ◐ `ACTIVE`, OPT-125m 6/6; Qwen3-0.6B wired; Mistral-7B-v0.3 16/16 on P150 ([spec](../.agents/specs/tenstorrent-mistral.md)). 16x16 rerun and residual-RMS owed ([spec](../.agents/specs/tenstorrent-backend.md)) | ✅ | ☐ | ☐ | | Tenstorrent host-free decode | ◐ env-gated `VT_TT_HOST_FREE_DECODE`; implementer P150 79-replay/5.8x. Default inert. New batch after capture refused. Engine golden owed | ☐ | ☐ | ☐ | diff --git a/docs/STATUS.md b/docs/STATUS.md index 5a22b0429..82ac0e8f9 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -560,18 +560,25 @@ recurrences + fused attn preamble; 27B prefill 21.5x, decode [campaign](../.agents/specs/vulkan-full-support.md)), ROCm (W0 community-green on 5 gfx archs; the APU unified-memory fix remains unverified; gfx1200 runs Gemma-3 and Qwen3 all-native, with Gemma-3 strict 48/48 against two vLLM-ROCm -oracles and Qwen3 in a measured near-tie regime; Qwen3.5-0.8B GDN runs all-native -but its CPU/ROCm divergence remains open; gfx1201 Gemma-4 FP8 MoE is -contributor-measured on 2x R9700 and CPU-link-verified our side; a `head_dim=128` -decode arm lands opt-in behind `VT_ATTN_DECODE_D128`, default OFF, which moves -gfx1200 per-token decode from 6.35x to 1.75x slower than the pinned vLLM oracle -on one shape, a directional figure that leaves the ROCm throughput axis PENDING; +oracles and Qwen3 in a measured near-tie regime; on gfx1100 the M4 gate now +runs against the **pinned vLLM-ROCm oracle built on the same box** +(`555967922` / `0.23.1rc1.dev1511+g555967922`): Qwen3-0.6B **16/16 PASS** +(11/16 strict token-exact, 5/16 near-tie band, max teacher-forced gap 0.125 +nats, 0 forward-divergent; oracle K=10 deterministic in every cell) with the +ROCm device-golden lane in `test_qwen3_paged_engine`; Qwen3.5-0.8B GDN runs +all-native and is gated **16/16 against the same-box pinned oracle** (15/16 +strict token-exact, max gap 0.125 nats) since the `AttnQkNormRopeGate` +output-dtype dispatch fix; gfx1201 Gemma-4 FP8 MoE is contributor-measured on +2x R9700 and CPU-link-verified our side; a `head_dim=128` decode arm lands +opt-in behind `VT_ATTN_DECODE_D128`, default OFF, which moves gfx1200 per-token +decode from 6.35x to 1.75x slower than the pinned vLLM oracle on one shape, a +directional figure that leaves the ROCm throughput axis PENDING; [guide](ROCM.md)), inference-time CPU weight offload (`ENG-WEIGHT-OFFLOAD` ACTIVE; the config surface landed W0a (the backend enum, both sub-configs, the validator's two errors and three warnings, and the dot-anchored segment match), all UNREACHABLE for now because nothing constructs an `OffloadConfig` yet, so -no engine behaviour changes. Still owed: vLLM's `cpu_offload_gb` UVA arm with dotted-segment -`cpu_offload_params` targeting, plus the layer-group `PrefetchOffloader` — a +no engine behaviour changes. Still owed: vLLM's `cpu_offload_gb` UVA arm with +dotted-segment `cpu_offload_params` targeting, plus the layer-group `PrefetchOffloader` — a pure mirror floor, and #149's dense half. Its memory and speed gates need a discrete-GPU rig, because on unified-memory GB10 offloading to "CPU" frees nothing, so those gates are blocked rather than pending diff --git a/scripts/env-doc-allowlist.txt b/scripts/env-doc-allowlist.txt index 1422e0210..fe66349a4 100644 --- a/scripts/env-doc-allowlist.txt +++ b/scripts/env-doc-allowlist.txt @@ -227,3 +227,5 @@ VT_W4A4_TRUE # selects a kernel or a batching strategy for the expert GEMMs and each # treats a name here as a known kernel-internal knob. A NEW production env # var must be documented in docs/ENVIRONMENT.md OR added here. +VT_DUMP_ACT_SUB +VT_DUMP_ATTN diff --git a/src/vllm/model_executor/models/qwen3_5.cpp b/src/vllm/model_executor/models/qwen3_5.cpp index 093f982b4..dabdbbfdc 100644 --- a/src/vllm/model_executor/models/qwen3_5.cpp +++ b/src/vllm/model_executor/models/qwen3_5.cpp @@ -5240,6 +5240,34 @@ DBuf FullAttnBlockPaged(Dev d, const FullAttnLayerWeights& w, const HfConfig& cf pa_args.max_seq_len = meta.max_seq_len; vt::PagedAttention(d.q, dattn.t(), qn3, k_cache, v_cache, dblk, dsl, dqsl, pa_args); + // VT_DUMP_ATTN (issue #41, 0.8B ROCm divergence spike W1/W2): dump the + // full-attn block's internals per full-attn-layer call index (0-based), as + // raw little-endian dumps under $VT_DUMP_ATTN/fa_{qkv,q,attn,gate}.bin. + // Inert when unset; the Downloads sync, so never set on a graph path. + static thread_local int64_t dump_fa_idx = -1; + dump_fa_idx++; + if (std::getenv("VT_DUMP_ATTN") != nullptr) { + auto DumpT = [&](const char* stage, const Tensor& t) { + int64_t n = 1; + for (int i = 0; i < t.rank; ++i) n *= t.shape[i]; + const size_t es = vt::SizeOf(t.dtype); + // contiguous check: innermost stride 1 and packed + std::vector raw(static_cast(n) * es); + DBuf tmp(d, t.dtype, {n}); + 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_ATTN")) + + "/fa" + std::to_string(dump_fa_idx) + "_" + + 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); } + }; + DumpT("qkv", qgate); + DumpT("q", qn3); + DumpT("attn", dattn.t()); + DumpT("gate", gatef.t()); + } + // Sigmoid output gate, folded into the o_proj activation quant on the true-W4A4 // path (§5) — see SigmoidGateOProjD. return SigmoidGateOProjD(d, Reshape(dattn.t(), {T, Hq * Dh}), @@ -7206,9 +7234,27 @@ void RunDenseLayerPaged(Dev d, const Qwen3_5DenseLayerWeights& layer, const int64_t H = cfg.hidden_size; const float eps = static_cast(cfg.rms_norm_eps); + // VT_DUMP_ACT layer index: the paged loop's per-layer counter (0-based), used + // by the sub-stage dump below. Declared per-process; only the dump reads it. + static thread_local int64_t dump_layer_idx = -1; + dump_layer_idx++; + const bool dump_sub = std::getenv("VT_DUMP_ACT_SUB") != nullptr; + auto DumpStage = [&](const char* stage, DBuf& buf) { + if (!dump_sub) return; + std::vector raw(static_cast(T) * static_cast(H) * + vt::SizeOf(buf.t().dtype)); + buf.Download(d, raw.data()); + const std::string path = std::string(std::getenv("VT_DUMP_ACT_SUB")) + + "/layer_" + std::to_string(dump_layer_idx) + "_" + + 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); } + }; + 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()); + DumpStage("post_input_norm", dhn); DBuf attn = [&] { if (layer.is_linear_attention) { @@ -7221,12 +7267,15 @@ void RunDenseLayerPaged(Dev d, const Qwen3_5DenseLayerWeights& layer, return FullAttnBlockPaged(d, layer.attn, cfg, dhn.t(), sdi, attn_meta, *attn_kv, T); }(); + DumpStage("block_out", attn); 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()); + DumpStage("post_attn_norm", dh2); hidden = DenseMlpBlock(d, layer.mlp, cfg, dh2.t(), T); + DumpStage("mlp_out", hidden); } // ── Qwen3.5/3.6 MTP head shared preamble (SPEC-MTP I5c). ──────────────────── @@ -8698,6 +8747,24 @@ static DBuf DenseForwardLayers(Dev d, const Tensor& hidden_in, // DFlash DF-AUX-TAPS: capture (hidden+res) at configured boundaries. Inert // (no-op) when aux_out is null — every non-DFlash caller. MaybeCaptureAuxTap(d, l, aux_layer_ids, aux_out, hidden.t(), res.t(), T, H); + // VT_DUMP_ACT (issue #41, ROCm 0.8B forward-divergence fix spike W1): dump + // the residual stream after each layer as raw little-endian bf16 to + // $VT_DUMP_ACT/layer_.bin (inert when unset; a debug hook, never the + // hot path — the Download forces a sync, so capture-graph paths must not + // set the env). + if (std::getenv("VT_DUMP_ACT") != nullptr) { + std::vector raw(static_cast(T) * static_cast(H) * + vt::SizeOf(hidden.t().dtype)); + hidden.Download(d, raw.data()); // Copy + Synchronize + const char* dir = std::getenv("VT_DUMP_ACT"); + const std::string path = + std::string(dir) + "/layer_" + std::to_string(l) + ".bin"; + std::FILE* f = std::fopen(path.c_str(), "wb"); + if (f != nullptr) { + std::fwrite(raw.data(), 1, raw.size(), f); + std::fclose(f); + } + } } // Final RMSNorm over the fused stream (res += hidden; norm), then lm_head. diff --git a/src/vt/rocm/rocm_gdn_fused.hip b/src/vt/rocm/rocm_gdn_fused.hip index 554912752..9e5c60f7d 100644 --- a/src/vt/rocm/rocm_gdn_fused.hip +++ b/src/vt/rocm/rocm_gdn_fused.hip @@ -250,16 +250,28 @@ void AttnQkNormRopeGateKernelRocm(Queue& q, Tensor& q_out, Tensor& k_out, Tensor kf.Ptr(), q_norm.Ptr(), k_norm.Ptr(), cos_sin.Ptr(), t, hq, hkv, dh, qgate.stride[0], kf.stride[0], ra.rotary_dim, na.eps, na.gemma); }; - // Validated combos (mirror the CUDA lane): (f32 src, f32 out); (bf16 src, - // bf16 out + bf16 gate); (bf16 src, bf16 out + f32 gate). - if (qgate.dtype == DType::kBF16) { + // Dispatch on the OUTPUT dtype like the CUDA lane (cuda_ops.cu + // LaunchAttnPreambleOut): the caller picks f32 out (token-exact path) or bf16 + // out (FA-2 combos) INDEPENDENTLY of the source dtype — the 0.8B bf16 model + // runs bf16 src + f32 q/k out + f32 gate on the f32-attention path, which the + // old src-keyed dispatch silently mis-launched as all-bf16 (writing bf16 bits + // through the f32 out pointer; issue #41 M4 W2 finding). Validated combos + // (mirror CUDA): (f32 src, f32 out); (bf16 src, f32 out + f32 gate); + // (bf16 src, bf16 out + bf16 gate); (bf16 src, bf16 out + f32 gate). + const bool src_bf16 = qgate.dtype == DType::kBF16; + if (q_out.dtype == DType::kF32) { + if (src_bf16) { + launch(__hip_bfloat16{}, float{}, float{}); + } else { + launch(float{}, float{}, float{}); + } + } else { + VT_CHECK(src_bf16, "rocm attn_qk_norm_rope_gate: bf16 out needs bf16 src"); if (gate_out.dtype == DType::kF32) { launch(__hip_bfloat16{}, __hip_bfloat16{}, float{}); } else { launch(__hip_bfloat16{}, __hip_bfloat16{}, __hip_bfloat16{}); } - } else { - launch(float{}, float{}, float{}); } Check(hipGetLastError(), "attn_qk_norm_rope_gate launch"); } diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 6bf3cf22c..b8c052162 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -2010,6 +2010,21 @@ target_compile_definitions(test_qwen3_paged_engine PRIVATE target_include_directories(test_qwen3_paged_engine PRIVATE ${CMAKE_SOURCE_DIR}/tests/parity ${CMAKE_SOURCE_DIR}/src) +# ROCm M4 — the GDN sibling of the Qwen3-dense gate: Qwen3.5-0.8B greedy decode +# vs the pinned vLLM-ROCm oracle (the ONLY capture of this model — no dgx +# capture exists; the base golden pair is ROCm-captured on gfx1100). +# Checkpoint-gated: skips loudly when the HF snapshot is absent. +vllm_cpp_add_test(test_qwen35_paged_engine parity/test_qwen35_paged_engine.cpp) +target_compile_definitions(test_qwen35_paged_engine PRIVATE + PARITY_GOLDENS_DIR="${CMAKE_SOURCE_DIR}/tests/parity/goldens") +target_include_directories(test_qwen35_paged_engine PRIVATE + ${CMAKE_SOURCE_DIR}/tests/parity ${CMAKE_SOURCE_DIR}/src) +add_test(NAME test_qwen35_paged_engine_prerequisites + COMMAND ${CMAKE_COMMAND} + "-DTEST_EXE=$" + "-DPROBE_ROOT=${CMAKE_CURRENT_BINARY_DIR}/qwen35_paged_engine_prerequisites" + -P "${CMAKE_CURRENT_SOURCE_DIR}/parity/test_qwen35_paged_engine_prerequisites.cmake") + # ROW-SERVE-ASYNC-DENSE-MIRROR — the classic-dense async-serving greedy gate: # Qwen3-0.6B/4B async batch-1 + concurrency reproduce the race-free SYNC engine # continuation token-for-token. RED on VT_ASYNC_DEVICE_MIRROR=0 (dense EmbedInto diff --git a/tests/parity/goldens/qwen35_greedy_0_8b/greedy_dist.npy b/tests/parity/goldens/qwen35_greedy_0_8b/greedy_dist.npy new file mode 100644 index 000000000..5282c5800 Binary files /dev/null and b/tests/parity/goldens/qwen35_greedy_0_8b/greedy_dist.npy differ diff --git a/tests/parity/goldens/qwen35_greedy_0_8b/greedy_ids.npy b/tests/parity/goldens/qwen35_greedy_0_8b/greedy_ids.npy new file mode 100644 index 000000000..b101b8145 Binary files /dev/null and b/tests/parity/goldens/qwen35_greedy_0_8b/greedy_ids.npy differ diff --git a/tests/parity/goldens/qwen35_greedy_0_8b/manifest.json b/tests/parity/goldens/qwen35_greedy_0_8b/manifest.json new file mode 100644 index 000000000..ad24cf746 --- /dev/null +++ b/tests/parity/goldens/qwen35_greedy_0_8b/manifest.json @@ -0,0 +1,23 @@ +{ + "op": "qwen35_0_8b_greedy", + "oracle": { + "engine": "pinned vLLM-ROCm oracle built from source on the gate host (issue #41 M4)", + "vllm_commit": "5559679229bc961848b121ccdeaa8fa5d79bec98", + "vllm_runtime_version": "0.23.1rc1.dev1511+g555967922", + "wheel": "vllm-0.23.1rc1.dev1511+g555967922.rocm723-cp312-cp312-linux_x86_64.whl", + "container_image": "vllm-rocm-oracle:555967922-gfx1100 (rocm/vllm-dev:base + the wheel)", + "hardware": "4x RX 7900 XTX (gfx1100), ROCm 7.14 host driver, HIP 7.2.5 userland", + "invocation": "scripts/qwen3-oracle-capture.py --model Qwen/Qwen3.5-0.8B --runs 10 --per-prompt --gpu-mem 0.8 (enforce_eager=True, batch=1); teacher-forced gaps via scripts/qwen3-neartie-gap.py", + "date": "2026-08-12", + "note": "There is no dgx/CUDA capture of this model: the 0.8B GDN checkpoint was only ever stood up on the gfx1100 box, which is also the only board hosting a vLLM-ROCm oracle. The base golden pair is ROCm-captured; K=10 per-prompt greedy is deterministic in every (prompt,pos) cell (greedy_dist.npy). GATE STATUS: GREEN -- 16/16 PASS (15/16 strict token-exact, 1/16 near-tie band, max gap 0.125 nats) with the AttnQkNormRopeGate output-dtype dispatch fix (row/ROCM-GDN-08B-FIX). The PRE-FIX capture showed 13/16 forward-divergent with first-token divergences on 6 prompts (first-divergence gaps 0.375-1.062 nats; 5 of 6 over the 0.5-nat band) -- recorded in the spec and parity ledger." + }, + "model": { + "repo": "Qwen/Qwen3.5-0.8B", + "revision": "2fc06364715b967f1860aea9cf38778875588b17", + "architectures": [ + "Qwen3_5ForConditionalGeneration" + ], + "note": "revision verified via the checkpoint's huggingface_hub download metadata (2026-08-12)" + }, + "prompts": "the standard 16-prompt battery shared with scripts/qwen3-oracle-capture.py / test_qwen3_paged_engine.cpp" +} \ No newline at end of file diff --git a/tests/parity/goldens/qwen35_greedy_0_8b/neartie_gap_mnats.npy b/tests/parity/goldens/qwen35_greedy_0_8b/neartie_gap_mnats.npy new file mode 100644 index 000000000..0fea365bf Binary files /dev/null and b/tests/parity/goldens/qwen35_greedy_0_8b/neartie_gap_mnats.npy differ diff --git a/tests/parity/goldens/qwen35_greedy_0_8b/our_ids.npy b/tests/parity/goldens/qwen35_greedy_0_8b/our_ids.npy new file mode 100644 index 000000000..aad44126e Binary files /dev/null and b/tests/parity/goldens/qwen35_greedy_0_8b/our_ids.npy differ diff --git a/tests/parity/goldens/qwen35_greedy_0_8b/p0_prompt.i32 b/tests/parity/goldens/qwen35_greedy_0_8b/p0_prompt.i32 new file mode 100644 index 000000000..191546743 Binary files /dev/null and b/tests/parity/goldens/qwen35_greedy_0_8b/p0_prompt.i32 differ diff --git a/tests/parity/goldens/qwen35_greedy_0_8b/p10_prompt.i32 b/tests/parity/goldens/qwen35_greedy_0_8b/p10_prompt.i32 new file mode 100644 index 000000000..454776427 Binary files /dev/null and b/tests/parity/goldens/qwen35_greedy_0_8b/p10_prompt.i32 differ diff --git a/tests/parity/goldens/qwen35_greedy_0_8b/p11_prompt.i32 b/tests/parity/goldens/qwen35_greedy_0_8b/p11_prompt.i32 new file mode 100644 index 000000000..623cdcfaf Binary files /dev/null and b/tests/parity/goldens/qwen35_greedy_0_8b/p11_prompt.i32 differ diff --git a/tests/parity/goldens/qwen35_greedy_0_8b/p12_prompt.i32 b/tests/parity/goldens/qwen35_greedy_0_8b/p12_prompt.i32 new file mode 100644 index 000000000..229d6e926 Binary files /dev/null and b/tests/parity/goldens/qwen35_greedy_0_8b/p12_prompt.i32 differ diff --git a/tests/parity/goldens/qwen35_greedy_0_8b/p13_prompt.i32 b/tests/parity/goldens/qwen35_greedy_0_8b/p13_prompt.i32 new file mode 100644 index 000000000..a2b1cb6c5 Binary files /dev/null and b/tests/parity/goldens/qwen35_greedy_0_8b/p13_prompt.i32 differ diff --git a/tests/parity/goldens/qwen35_greedy_0_8b/p14_prompt.i32 b/tests/parity/goldens/qwen35_greedy_0_8b/p14_prompt.i32 new file mode 100644 index 000000000..ea74bf755 Binary files /dev/null and b/tests/parity/goldens/qwen35_greedy_0_8b/p14_prompt.i32 differ diff --git a/tests/parity/goldens/qwen35_greedy_0_8b/p15_prompt.i32 b/tests/parity/goldens/qwen35_greedy_0_8b/p15_prompt.i32 new file mode 100644 index 000000000..c72735e10 Binary files /dev/null and b/tests/parity/goldens/qwen35_greedy_0_8b/p15_prompt.i32 differ diff --git a/tests/parity/goldens/qwen35_greedy_0_8b/p1_prompt.i32 b/tests/parity/goldens/qwen35_greedy_0_8b/p1_prompt.i32 new file mode 100644 index 000000000..1130fee1e Binary files /dev/null and b/tests/parity/goldens/qwen35_greedy_0_8b/p1_prompt.i32 differ diff --git a/tests/parity/goldens/qwen35_greedy_0_8b/p2_prompt.i32 b/tests/parity/goldens/qwen35_greedy_0_8b/p2_prompt.i32 new file mode 100644 index 000000000..083e5883d Binary files /dev/null and b/tests/parity/goldens/qwen35_greedy_0_8b/p2_prompt.i32 differ diff --git a/tests/parity/goldens/qwen35_greedy_0_8b/p3_prompt.i32 b/tests/parity/goldens/qwen35_greedy_0_8b/p3_prompt.i32 new file mode 100644 index 000000000..e69ad7489 Binary files /dev/null and b/tests/parity/goldens/qwen35_greedy_0_8b/p3_prompt.i32 differ diff --git a/tests/parity/goldens/qwen35_greedy_0_8b/p4_prompt.i32 b/tests/parity/goldens/qwen35_greedy_0_8b/p4_prompt.i32 new file mode 100644 index 000000000..a1138d117 Binary files /dev/null and b/tests/parity/goldens/qwen35_greedy_0_8b/p4_prompt.i32 differ diff --git a/tests/parity/goldens/qwen35_greedy_0_8b/p5_prompt.i32 b/tests/parity/goldens/qwen35_greedy_0_8b/p5_prompt.i32 new file mode 100644 index 000000000..92664cad2 Binary files /dev/null and b/tests/parity/goldens/qwen35_greedy_0_8b/p5_prompt.i32 differ diff --git a/tests/parity/goldens/qwen35_greedy_0_8b/p6_prompt.i32 b/tests/parity/goldens/qwen35_greedy_0_8b/p6_prompt.i32 new file mode 100644 index 000000000..3561c6f4d Binary files /dev/null and b/tests/parity/goldens/qwen35_greedy_0_8b/p6_prompt.i32 differ diff --git a/tests/parity/goldens/qwen35_greedy_0_8b/p7_prompt.i32 b/tests/parity/goldens/qwen35_greedy_0_8b/p7_prompt.i32 new file mode 100644 index 000000000..b8bc06310 Binary files /dev/null and b/tests/parity/goldens/qwen35_greedy_0_8b/p7_prompt.i32 differ diff --git a/tests/parity/goldens/qwen35_greedy_0_8b/p8_prompt.i32 b/tests/parity/goldens/qwen35_greedy_0_8b/p8_prompt.i32 new file mode 100644 index 000000000..db62992ee Binary files /dev/null and b/tests/parity/goldens/qwen35_greedy_0_8b/p8_prompt.i32 differ diff --git a/tests/parity/goldens/qwen35_greedy_0_8b/p9_prompt.i32 b/tests/parity/goldens/qwen35_greedy_0_8b/p9_prompt.i32 new file mode 100644 index 000000000..be6f068ae Binary files /dev/null and b/tests/parity/goldens/qwen35_greedy_0_8b/p9_prompt.i32 differ diff --git a/tests/parity/goldens/qwen3_greedy_0_6b/greedy_dist_rocm.npy b/tests/parity/goldens/qwen3_greedy_0_6b/greedy_dist_rocm.npy new file mode 100644 index 000000000..ed0bf47c4 Binary files /dev/null and b/tests/parity/goldens/qwen3_greedy_0_6b/greedy_dist_rocm.npy differ diff --git a/tests/parity/goldens/qwen3_greedy_0_6b/greedy_ids_rocm.npy b/tests/parity/goldens/qwen3_greedy_0_6b/greedy_ids_rocm.npy new file mode 100644 index 000000000..b8b49ca47 Binary files /dev/null and b/tests/parity/goldens/qwen3_greedy_0_6b/greedy_ids_rocm.npy differ diff --git a/tests/parity/goldens/qwen3_greedy_0_6b/neartie_gap_mnats_rocm.npy b/tests/parity/goldens/qwen3_greedy_0_6b/neartie_gap_mnats_rocm.npy new file mode 100644 index 000000000..a845e0408 Binary files /dev/null and b/tests/parity/goldens/qwen3_greedy_0_6b/neartie_gap_mnats_rocm.npy differ diff --git a/tests/parity/goldens/qwen3_greedy_0_6b/our_ids_rocm.npy b/tests/parity/goldens/qwen3_greedy_0_6b/our_ids_rocm.npy new file mode 100644 index 000000000..cff630492 Binary files /dev/null and b/tests/parity/goldens/qwen3_greedy_0_6b/our_ids_rocm.npy differ diff --git a/tests/parity/hf_snapshot.h b/tests/parity/hf_snapshot.h index f0e2ff311..ebb2cd398 100644 --- a/tests/parity/hf_snapshot.h +++ b/tests/parity/hf_snapshot.h @@ -378,6 +378,20 @@ inline std::string Qwen27DFlashDraftSnapshot() { return HfSnapshot("models--z-lab--Qwen3.6-27B-DFlash", kQwen27DFlashDraftRevision, "VT_DFLASH_DRAFT_SNAPSHOT");} +// The Qwen3.5-0.8B GDN gate model (issue #41, the ROCm M4 lane). This revision +// is what the committed `goldens/qwen35_greedy_0_8b` were captured against on +// gfx1100 2026-08-12 (verified via the checkpoint's huggingface_hub download +// metadata; `manifest.json` in the golden dir records the oracle identity). +inline constexpr const char* kQwen35_08BRevision = + "2fc06364715b967f1860aea9cf38778875588b17"; + +// The 0.8B GDN gate model, pinned to the goldens' revision. Absent cache => "" +// and the caller emits its loud SKIP (the intended behavior off the gate host). +inline std::string Qwen35_08BSnapshot() { + return HfSnapshot("models--Qwen--Qwen3.5-0.8B", kQwen35_08BRevision, + "VT_QWEN35_08B_SNAPSHOT"); +} + } // namespace parity #endif // VLLM_TESTS_PARITY_HF_SNAPSHOT_H_ diff --git a/tests/parity/test_op_parity.cpp b/tests/parity/test_op_parity.cpp index f9699f6be..1a316d6ee 100644 --- a/tests/parity/test_op_parity.cpp +++ b/tests/parity/test_op_parity.cpp @@ -1836,6 +1836,9 @@ const std::set& PendingRunnerOps() { // Full-engine GGUF greedy acceptance fixture consumed by // test_qwen36_gguf_engine.cpp. "qwen36_gguf_greedy", + // Full-engine paged greedy gate fixture consumed by + // test_qwen35_paged_engine.cpp (the ROCm M4 0.8B GDN gate, issue #41). + "qwen35_0_8b_greedy", }; return kPending; } diff --git a/tests/parity/test_qwen27n_fp8_tower_paged_engine.cpp b/tests/parity/test_qwen27n_fp8_tower_paged_engine.cpp index cf5cabcdc..6a619864d 100644 --- a/tests/parity/test_qwen27n_fp8_tower_paged_engine.cpp +++ b/tests/parity/test_qwen27n_fp8_tower_paged_engine.cpp @@ -92,7 +92,10 @@ namespace { // Qwen3.6-27B has 48 GDN (`linear_attention`) layers. One merged FP8 qkvz GEMM // per layer on the default arm; two split GEMMs per layer on the // VT_GDN_MERGED_QKVZ_FP8=0 rollback. -constexpr uint64_t kGdnLayers = 48; +// [[maybe_unused]]: the uses below live behind the CUDA-only fp8-tower arms; +// on a HIP-only build they compile out and the constant would be an +// -Wunused-const-variable error under clang/hipcc. +[[maybe_unused]] constexpr uint64_t kGdnLayers = 48; // Snapshot dir of the FP8-tower 27B, or "" to refuse. Pinned to the revision // its goldens were captured against; a cache holding some other revision of the diff --git a/tests/parity/test_qwen35_paged_engine.cpp b/tests/parity/test_qwen35_paged_engine.cpp new file mode 100644 index 000000000..542825d5c --- /dev/null +++ b/tests/parity/test_qwen35_paged_engine.cpp @@ -0,0 +1,391 @@ +// vllm.cpp original (ADDITIVE-MODEL W4 — THE SACRED correctness gate); no +// upstream mirror. +// +// THE PAGED-ENGINE Qwen3.5-0.8B (GDN) GREEDY CORRECTNESS GATE. The GDN sibling of +// test_qwen3_paged_engine.cpp: drives the standard prompt battery through the FULL +// PAGED LLMEngine stack (InputProcessor -> Scheduler -> GDN recurrence + paged +// attention on the full-attention layers + Sampler -> OutputProcessor) via +// LoadedEngine::FromModelDir, and checks greedy (temperature-0) decode against the +// pinned vLLM oracle (commit 555967922, runtime 0.23.1rc1.dev1511+g555967922). +// +// ORACLE PROVENANCE — DIFFERENT FROM THE QWEN3-DENSE GATE. There is no dgx/CUDA +// capture for this model: the 0.8B GDN checkpoint was only ever stood up on the +// gfx1100 box (issue #41 lane), which is also the only board that hosts a +// vLLM-ROCm oracle. The base golden pair here is therefore ROCm-captured: +// greedy_ids.npy [N,T] i32 the pinned ROCm oracle's per-prompt greedy +// (K=10 per-prompt runs DETERMINISTIC in every +// cell — see greedy_dist.npy). +// our_ids.npy [N,T] i32 OUR ROCm engine's greedy (the anchor). +// neartie_gap_mnats.npy [N,T] oracle teacher-forced gap (milli-nats) for OUR +// 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. +// +// PROVENANCE OF THE GREEN: the committed our_ids/near-tie pair is the +// FIXED engine's sequence, oracle-re-derived after the AttnQkNormRopeGate +// output-dtype dispatch fix (row/ROCM-GDN-08B-FIX). The PRE-FIX capture +// (13/16 forward-divergent, max first-divergence gap 1.062 nats) is recorded +// as evidence in .agents/specs/rocm-m4-oracle.md and the parity ledger — the +// gate landed green-shaped per review, with the RED capture kept as history +// rather than as committed goldens. +// +// METHODOLOGY — identical to the Qwen3-dense gate (see that file's header and +// [[near-tie-distributional-gate]]): STRICT token-exact is reported, but the +// PASS bar is the near-tie band: given OUR EXACT PREFIX, the oracle's OWN logits +// must place OUR token within kNearTieMnats of the oracle argmax. Strict where +// well-posed, near-tie-robust only where the oracle itself cannot separate the +// tokens. +// +// BACKEND PROOF — GDN op set (not the Qwen3-dense list): the GDN layers dispatch +// kCausalConv1dFwd/Update (prefill/decode conv), kGdnPrefill/kGdnDecode (the +// recurrence), kGdnPostConv, kRmsNormGated, kSigmoidGateBf16, +// kAttnQkNormRopeGate (the GDN preamble), and the full-attention layers dispatch +// kReshapeAndCache + kPagedAttention + kAttnQkNormRope; shared: kEmbedding, +// kMatmulBT, kRmsNorm, kSiluAndMul, kGreedyArgmax. All must show selections>0 +// and declines==0 on the running device. +// +// Checkpoint-gated: resolves the HF snapshot under ~/.cache/huggingface/hub/. +// Absent snapshot (CI/CPU) => loud SKIP. +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "hf_snapshot.h" +#include "npy.h" +#include "vllm/entrypoints/model_loader.h" +#include "vllm/sampling_params.h" +#include "vt/op_provider.h" // the "which backend actually ran" proof +#include "vt/ops.h" + +namespace fs = std::filesystem; + +namespace { + +// Near-tie acceptance threshold in milli-nats — same bar as the Qwen3-dense +// gate (0.5 nats); see that file for the derivation. +constexpr int32_t kNearTieMnats = 500; +// Stable test-only sentinel: distinct from process success and CTest skip 77. +constexpr int kPrerequisiteProbeCompleteExit = 86; + +// The standard prompt battery — MUST match scripts/qwen3-oracle-capture.py (and +// qwen3-neartie-gap.py) exactly; those scripts are model-agnostic and take +// --model/--out-dir/--golden-dir. +const std::vector& Prompts() { + static const std::vector p = { + "The capital of France is", + "Once upon a time,", + "In the beginning God created", + "The quick brown fox jumps over", + "def fibonacci(n):", + "Water boils at a temperature of", + "The theory of relativity was developed by", + "To be or not to be, that is", + "The largest planet in our solar system is", + "Machine learning is a subfield of", + "The mitochondria is the powerhouse of", + "Roses are red, violets are", + "The first president of the United States was", + "E equals m c", + "A journey of a thousand miles begins with", + "The chemical symbol for gold is", + }; + return p; +} + +vllm::SamplingParams Greedy(int max_tokens) { + vllm::SamplingParams sp; + sp.temperature = 0.0; + sp.max_tokens = max_tokens; + sp.PostInit(); + return sp; +} + +const int32_t* AsI32(const parity::NpyArray& a) { + return reinterpret_cast(a.data.data()); +} + +// A gate that cannot run must not report success: exit 77 makes CTest report +// **Skipped**, never a false green (issue #463; SKIP_RETURN_CODE is wired for +// every parity test in tests/CMakeLists.txt). +[[noreturn]] void SkipGate(const char* label, const std::string& why) { + std::fprintf(stderr, + "\n*** GATE NOT RUN — SKIPPED (exit 77), this is NOT a pass ***\n" + "*** %s: %s\n\n", + label, why.c_str()); + std::fflush(stderr); + std::exit(77); +} + +enum class GateArtifactState { kReady, kBootstrap }; + +GateArtifactState RequireGateArtifacts(const fs::path& gdir, const char* label, + bool dump) { + if (!fs::exists(gdir / "greedy_ids.npy")) { + SkipGate(label, + "required greedy_ids.npy is absent. Capture it with the pinned " + "ROCm oracle using qwen3-oracle-capture.py --per-prompt"); + } + + const bool have_anchor = fs::exists(gdir / "our_ids.npy"); + const bool have_gap = fs::exists(gdir / "neartie_gap_mnats.npy"); + if (have_anchor && have_gap) return GateArtifactState::kReady; + if (dump) return GateArtifactState::kBootstrap; + + std::string missing; + if (!have_anchor) missing = "our_ids.npy"; + if (!have_gap) { + if (!missing.empty()) missing += " and "; + missing += "neartie_gap_mnats.npy"; + } + SkipGate(label, "required " + missing + + " is absent. Run with VT_DUMP_IDS=1, then run " + "qwen3-neartie-gap.py"); +} + +void RunGate(const std::string& golden_subdir, const char* label) { + const char* probe_dir = std::getenv("VT_QWEN35_GATE_PREREQ_PROBE_DIR"); + const bool probe = probe_dir != nullptr; + const std::string snap = probe ? std::string() : parity::Qwen35_08BSnapshot(); + if (!probe && snap.empty()) { + SkipGate(label, "models--Qwen--Qwen3.5-0.8B snapshot at the pinned revision " + "2fc06364 not cached — this gate runs where the ROCm oracle " + "was captured (gfx1100)"); + } + const fs::path gdir = probe ? fs::path(probe_dir) + : fs::path(PARITY_GOLDENS_DIR) / golden_subdir; + const bool dump = !probe && std::getenv("VT_DUMP_IDS") != nullptr; + const GateArtifactState artifacts = + RequireGateArtifacts(gdir, label, dump); + if (probe) { + std::fprintf( + stderr, + "\n*** TEST PROBE COMPLETE (exit %d), production gate NOT RUN ***\n" + "*** %s: all three required artifact paths are present ***\n\n", + kPrerequisiteProbeCompleteExit, label); + std::fflush(stderr); + std::exit(kPrerequisiteProbeCompleteExit); + } + // BOOTSTRAP: with VT_DUMP_IDS set and no gap golden yet, generate + dump OUR + // token ids (our_ids.i32) so qwen3-neartie-gap.py can build the gap golden. + if (artifacts == GateArtifactState::kBootstrap) { + MESSAGE(label << ": BOOTSTRAP dump (gap golden absent) via FromModelDir(" << snap << ")..."); + std::unique_ptr le = + vllm::entrypoints::LoadedEngine::FromModelDir( + snap, vllm::entrypoints::EngineParams{}); + const parity::NpyArray gg = parity::LoadNpy((gdir / "greedy_ids.npy").string()); + const int64_t NN = gg.shape[0], TT = gg.shape[1]; + std::vector buf(static_cast(NN * TT), -1); + for (int64_t i = 0; i < NN; ++i) { + const vllm::RequestOutput out = le->engine().generate( + Prompts()[static_cast(i)], Greedy(static_cast(TT)), + "boot" + std::to_string(i)); + const std::vector& got = out.outputs[0].token_ids; + for (int64_t j = 0; j < TT && j < static_cast(got.size()); ++j) + buf[static_cast(i * TT + j)] = got[static_cast(j)]; + } + const std::string path = (gdir / "our_ids.i32").string(); + std::FILE* f = std::fopen(path.c_str(), "wb"); + if (f != nullptr) { std::fwrite(buf.data(), sizeof(int32_t), buf.size(), f); std::fclose(f); } + MESSAGE(label << " BOOTSTRAP dumped our token ids -> " << path); + SkipGate(label, "bootstrap does not run the correctness gate. Run " + "qwen3-neartie-gap.py, then rerun without VT_DUMP_IDS"); + } + + const parity::NpyArray g = parity::LoadNpy((gdir / "greedy_ids.npy").string()); + const parity::NpyArray o = parity::LoadNpy((gdir / "our_ids.npy").string()); + const parity::NpyArray gap = parity::LoadNpy((gdir / "neartie_gap_mnats.npy").string()); + REQUIRE(g.dtype == "(N) == Prompts().size()); + const int32_t* gd = AsI32(g); + const int32_t* od = AsI32(o); + const int32_t* gapd = AsI32(gap); + + std::vector our_dump; + if (dump) our_dump.assign(static_cast(N * T), -1); + + MESSAGE(label << ": loading via FromModelDir(" << snap << ")..."); + std::unique_ptr loaded = + vllm::entrypoints::LoadedEngine::FromModelDir( + 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. + const vt::DeviceType run_dev = loaded->runner().device().type; + const bool rocm = run_dev == vt::DeviceType::kROCM; + // 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) { + 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"); + } + + // The GDN op set this model dispatches — all must be proven on the running + // device (selections > 0, declines == 0; fan-out spike Risk 4). + const std::vector kGdnOps = { + vt::OpId::kEmbedding, vt::OpId::kMatmulBT, + vt::OpId::kRmsNorm, vt::OpId::kRmsNormGated, + vt::OpId::kCausalConv1dFwd, vt::OpId::kCausalConv1dUpdate, + vt::OpId::kGdnPrefill, vt::OpId::kGdnDecode, + vt::OpId::kGdnPostConv, vt::OpId::kSigmoidGateBf16, + vt::OpId::kAttnQkNormRopeGate, + vt::OpId::kReshapeAndCache, vt::OpId::kPagedAttention, + vt::OpId::kSiluAndMul, vt::OpId::kGreedyArgmax}; + if (rocm) { + 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"); + } + + const int32_t* anchor_ids = od; + const int32_t* gap_ids = gapd; + + int strict_exact = 0; + int neartie_only = 0; + int fail = 0; + int32_t worst_gap = 0; + int worst_i = -1, worst_j = -1; + for (int64_t i = 0; i < N; ++i) { + const vllm::RequestOutput out = loaded->engine().generate( + Prompts()[static_cast(i)], Greedy(static_cast(T)), + "gate" + std::to_string(i)); + REQUIRE(out.finished); + REQUIRE(out.outputs.size() == 1); + const std::vector& got = out.outputs[0].token_ids; + REQUIRE(static_cast(got.size()) == T); + if (dump) { + for (int64_t j = 0; j < T; ++j) + our_dump[static_cast(i * T + j)] = got[static_cast(j)]; + } + + // Anchor: the committed anchor is the exact deterministic sequence OUR ROCm + // engine produces. Drift is a hard REQUIRE — no cross-device latitude. + int first_div = -1; + for (int64_t j = 0; j < T; ++j) { + if (got[static_cast(j)] != anchor_ids[i * T + j]) { + first_div = static_cast(j); + break; + } + } + const bool anchor_ok = dump || first_div < 0; + // A drift is a REGRESSION until proven otherwise: this anchor is the fixed + // engine's oracle-verified sequence. Refreshing it is the LAST step of a + // justified re-capture, never the response to a failure. + REQUIRE_MESSAGE(anchor_ok, + label << " anchor drift prompt[" << i << "] tok=" << first_div + << " engine=" << (first_div < 0 ? -1 : got[static_cast(first_div)]) + << " committed anchor=" + << (first_div < 0 ? -1 : anchor_ids[i * T + first_div]) + << " — REGRESSION SUSPECTED: bisect the engine change first; " + "only re-derive goldens (VT_DUMP_IDS=1 + " + "qwen3-neartie-gap.py) after the drift is proven to be a " + "justified numerical change, never to silence the gate"); + + bool exact = true; + bool prompt_ok = true; + int first_bad = -1; + for (int64_t j = 0; j < T; ++j) { + if (got[static_cast(j)] != gd[i * T + j]) exact = false; + const int32_t mn = gap_ids[i * T + j]; + if (mn > worst_gap) { + worst_gap = mn; + worst_i = static_cast(i); + worst_j = static_cast(j); + } + if (mn > kNearTieMnats) { + prompt_ok = false; + if (first_bad < 0) first_bad = static_cast(j); + } + } + if (!prompt_ok) { + ++fail; + MESSAGE(label << " FORWARD DIVERGENCE prompt[" << i << "] tok=" << first_bad + << " our=" << got[static_cast(first_bad)] + << " vLLM_greedy=" << gd[i * T + first_bad] + << " gap=" << (gap_ids[i * T + first_bad] / 1000.0) << " nats (> " + << (kNearTieMnats / 1000.0) << ") \"" << out.outputs[0].text << "\""); + } else if (exact) { + ++strict_exact; + } else { + ++neartie_only; + } + CHECK(prompt_ok); + } + + // Backend proof: token equality alone does not prove which device ran. + if (rocm) { + vt::EnableOpProviderCallStats(false); + for (vt::OpId op : kGdnOps) { + const auto st = vt::GetOpProviderStats(op, run_dev); + CHECK_MESSAGE(st.selections > 0, + label << ": op " << static_cast(op) + << " was never dispatched on device type " + << static_cast(run_dev)); + CHECK_MESSAGE(st.declines == 0, + label << ": op " << static_cast(op) + << " DECLINED and fell back"); + } + MESSAGE(label << ": BACKEND PROOF — Qwen3.5 GDN ops on device type " + << static_cast(run_dev) << " with 0 declines (kPagedAttention " + "selections=" + << vt::GetOpProviderStats(vt::OpId::kPagedAttention, run_dev).selections + << ", kGdnDecode selections=" + << vt::GetOpProviderStats(vt::OpId::kGdnDecode, run_dev).selections + << ")"); + } + + if (dump) { + const std::string path = (gdir / "our_ids.i32").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); + std::fclose(f); + MESSAGE(label << " dumped our token ids -> " << path); + } + } + 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 + << "/" << N << "; max gap " << (worst_gap / 1000.0) << " nats @ prompt[" + << worst_i << "] tok=" << worst_j << "; " << fail + << " forward-divergent)"); + REQUIRE(fail == 0); +} + +} // namespace + +// Qwen3.5-0.8B (GDN hybrid: linear-attention recurrence + full-attention +// layers) — the first GDN-architecture gate, ROCm-oracle-backed (issue #41 M4). +TEST_CASE("qwen3.5-0.8B GDN paged-engine greedy near-tie correctness gate (ROCm, SACRED)") { + RunGate("qwen35_greedy_0_8b", "qwen3.5-0.8B"); +} diff --git a/tests/parity/test_qwen35_paged_engine_prerequisites.cmake b/tests/parity/test_qwen35_paged_engine_prerequisites.cmake new file mode 100644 index 000000000..209205478 --- /dev/null +++ b/tests/parity/test_qwen35_paged_engine_prerequisites.cmake @@ -0,0 +1,59 @@ +if(NOT DEFINED TEST_EXE OR NOT DEFINED PROBE_ROOT) + message(FATAL_ERROR "TEST_EXE and PROBE_ROOT are required") +endif() + +set(_required greedy_ids.npy our_ids.npy neartie_gap_mnats.npy) +# Stable test-only sentinel shared with test_qwen35_paged_engine.cpp. +set(_complete_probe_exit 86) +set(_failures "") +foreach(_missing IN LISTS _required) + set(_case_dir "${PROBE_ROOT}/${_missing}") + file(REMOVE_RECURSE "${_case_dir}") + file(MAKE_DIRECTORY "${_case_dir}") + foreach(_present IN LISTS _required) + if(NOT _present STREQUAL _missing) + file(WRITE "${_case_dir}/${_present}" "present") + endif() + endforeach() + + execute_process( + COMMAND "${CMAKE_COMMAND}" -E env + "VT_QWEN35_GATE_PREREQ_PROBE_DIR=${_case_dir}" + "${TEST_EXE}" + RESULT_VARIABLE _rc + OUTPUT_VARIABLE _stdout + ERROR_VARIABLE _stderr) + message(STATUS "missing ${_missing}: child exit ${_rc}") + if(NOT _rc EQUAL 77) + string(APPEND _failures + "\nmissing ${_missing}: expected exit 77, got ${_rc}" + "\nstdout:\n${_stdout}\nstderr:\n${_stderr}") + endif() +endforeach() + +set(_complete_dir "${PROBE_ROOT}/complete") +file(REMOVE_RECURSE "${_complete_dir}") +file(MAKE_DIRECTORY "${_complete_dir}") +foreach(_present IN LISTS _required) + file(WRITE "${_complete_dir}/${_present}" "present") +endforeach() + +execute_process( + COMMAND "${CMAKE_COMMAND}" -E env + "VT_QWEN35_GATE_PREREQ_PROBE_DIR=${_complete_dir}" + "${TEST_EXE}" + RESULT_VARIABLE _complete_rc + OUTPUT_VARIABLE _complete_stdout + ERROR_VARIABLE _complete_stderr) +message(STATUS + "complete artifacts: child exit ${_complete_rc} (probe sentinel ${_complete_probe_exit})") +if(NOT _complete_rc EQUAL _complete_probe_exit) + string(APPEND _failures + "\ncomplete artifacts: expected probe sentinel exit ${_complete_probe_exit}, got ${_complete_rc}" + "\nstdout:\n${_complete_stdout}\nstderr:\n${_complete_stderr}") +endif() + +file(REMOVE_RECURSE "${PROBE_ROOT}") +if(_failures) + message(FATAL_ERROR "Qwen3.5 gate prerequisite failures:${_failures}") +endif() diff --git a/tests/parity/test_qwen3_paged_engine.cpp b/tests/parity/test_qwen3_paged_engine.cpp index d0ded7df7..3d11cd1e2 100644 --- a/tests/parity/test_qwen3_paged_engine.cpp +++ b/tests/parity/test_qwen3_paged_engine.cpp @@ -227,8 +227,9 @@ void RunGate(const std::string& repo_dir, const std::string& golden_subdir, // not against CUDA's — identical anchor+band logic, device-appropriate goldens. const vt::DeviceType run_dev = loaded->runner().device().type; const bool metal = run_dev == vt::DeviceType::kMETAL; + const bool rocm = run_dev == vt::DeviceType::kROCM; const bool tenstorrent = run_dev == vt::DeviceType::kTENSTORRENT; - const bool device_golden = metal || tenstorrent; + const bool device_golden = metal || tenstorrent || rocm; // The forward + greedy ops Qwen3-dense dispatches on the DEFAULT // (VT_QWEN3_ROPE_CACHE) path. kRopeCosSinCache + kRopeFromCache are the M3b // additions (build the per-step cos|sin cache, then apply it); the rest are @@ -252,7 +253,7 @@ void RunGate(const std::string& repo_dir, const std::string& golden_subdir, } vt::EnableOpProviderCallStats(true); MESSAGE(label << ": running on device type " << static_cast(run_dev) - << " (2=METAL, 6=TENSTORRENT) — gated against this device's OWN " + << " (2=METAL, 5=ROCM, 6=TENSTORRENT) — gated against this device's OWN " "oracle-backed golden"); } @@ -261,9 +262,13 @@ void RunGate(const std::string& repo_dir, const std::string& golden_subdir, const int32_t* anchor_ids = od; // hard anchor for THIS device const int32_t* gap_ids = gapd; // vLLM teacher-forced gaps for THIS device parity::NpyArray o_dev, gap_dev; // keep device arrays alive for the loop - const char* ids_name = metal ? "our_ids_metal.npy" : "our_ids_tenstorrent.npy"; + const char* ids_name = + metal ? "our_ids_metal.npy" + : (rocm ? "our_ids_rocm.npy" : "our_ids_tenstorrent.npy"); const char* gap_name = - metal ? "neartie_gap_mnats_metal.npy" : "neartie_gap_mnats_tenstorrent.npy"; + metal ? "neartie_gap_mnats_metal.npy" + : (rocm ? "neartie_gap_mnats_rocm.npy" + : "neartie_gap_mnats_tenstorrent.npy"); bool bootstrap_only = false; if (device_golden) { const bool have_dev = fs::exists(gdir / ids_name) && fs::exists(gdir / gap_name); @@ -272,7 +277,7 @@ void RunGate(const std::string& repo_dir, const std::string& golden_subdir, // qwen3-neartie-gap.py then teacher-forces vLLM on that sequence. bootstrap_only = true; MESSAGE(label << ": BOOTSTRAP dump (device golden absent) for " - << (metal ? "Metal" : "Tenstorrent") << "..."); + << (metal ? "Metal" : (rocm ? "ROCm" : "Tenstorrent")) << "..."); } else { REQUIRE_MESSAGE(have_dev, label << ": device oracle golden absent (" << ids_name << " / " @@ -402,7 +407,8 @@ void RunGate(const std::string& repo_dir, const std::string& golden_subdir, if (dump) { const std::string dump_name = tenstorrent ? "our_ids_tenstorrent.i32" - : (metal ? "our_ids_metal.i32" : "our_ids.i32"); + : (metal ? "our_ids_metal.i32" + : (rocm ? "our_ids_rocm.i32" : "our_ids.i32")); const std::string path = (gdir / dump_name).string(); std::FILE* f = std::fopen(path.c_str(), "wb"); if (f != nullptr) { diff --git a/tests/vllm/models/test_qwen3_5_gdn_spec_routing.cpp b/tests/vllm/models/test_qwen3_5_gdn_spec_routing.cpp index e0fe74185..9d5cbe19e 100644 --- a/tests/vllm/models/test_qwen3_5_gdn_spec_routing.cpp +++ b/tests/vllm/models/test_qwen3_5_gdn_spec_routing.cpp @@ -363,7 +363,12 @@ void RunMixedRoutingCase(vt::DeviceType dev, const GdnDims& g, bool bit_exact) { INFO("dims := ", std::string(g.name)); CAPTURE(bad); CAPTURE(first); - if (bad != 0) { CAPTURE(mixed_out[first]); CAPTURE(ref[first]); } + if (bad != 0) { + // One CAPTURE per line: same-line pairs collide on DOCTEST_CAPTURE_ + // (clang/hipcc reject the redefinition; GCC happened not to). + CAPTURE(mixed_out[first]); + CAPTURE(ref[first]); + } CHECK(bad == 0); } else { float maxabs = 0.0f; diff --git a/tests/vt/test_backend_cross_device.cpp b/tests/vt/test_backend_cross_device.cpp index fc5d0e61b..35de30019 100644 --- a/tests/vt/test_backend_cross_device.cpp +++ b/tests/vt/test_backend_cross_device.cpp @@ -1885,6 +1885,67 @@ TEST_CASE("AttnQkNormRopeGate matches the CPU oracle within NMSE <= 5e-4") { } } + // The in-context production mix (issue #41 M4 W2): the 0.8B bf16 model feeds + // the preamble a BF16 projection output but wants F32 q/k/gate out (the f32 + // attention path — FA-2 is off on ROCm). The ROCm dispatcher once keyed on + // the SOURCE dtype and mis-launched all-bf16, writing bf16 bits through the + // f32 out pointers; this arm pins the (bf16 src -> f32 out) combo at the real + // 0.8B dims so the bug class cannot return silently. + { + const int64_t HQr = 8, HKVr = 2, DHr = 256, ROTr = 64; + const std::vector qg = RandomVec(static_cast(T * HQr * 2 * DHr), 991, -0.5f, 0.5f); + const std::vector kfv = RandomVec(static_cast(T * HKVr * DHr), 992, -0.5f, 0.5f); + const std::vector qnr = RandomVec(static_cast(DHr), 993, 0.2f, 1.0f); + const std::vector knr = RandomVec(static_cast(DHr), 994, 0.2f, 1.0f); + const std::vector csr = RandomVec(static_cast(T * ROTr), 995, -1.0f, 1.0f); + const std::vector qg_bf = Bf16Bits(qg), kf_bf = Bf16Bits(kfv); + vt::RmsNormArgs na3; na3.eps = 1e-6f; na3.gemma = true; + vt::RopeArgs ra3; ra3.rotary_dim = static_cast(ROTr); + // CPU reference: bf16 in (exact upcast inside the op) -> f32 out. + std::vector rq(static_cast(T * HQr * DHr)); + std::vector rk(static_cast(T * HKVr * DHr)); + std::vector rg(static_cast(T * HQr * DHr)); + { + vt::Backend& cpu = vt::GetBackend(DeviceType::kCPU); + Queue cq = cpu.CreateQueue(); + const Device cd{DeviceType::kCPU, 0}; + std::vector a = qg_bf, b = kf_bf; std::vector e = qnr, f = knr, g = csr; + Tensor tqg = Tensor::Contiguous(a.data(), DType::kBF16, cd, {T, HQr * 2 * DHr}); + Tensor tkf = Tensor::Contiguous(b.data(), DType::kBF16, cd, {T, HKVr * DHr}); + Tensor tqn = T1(e.data(), cd, DHr), tkn = T1(f.data(), cd, DHr); + Tensor tcs = T2(g.data(), cd, T, ROTr); + Tensor tqo = Tensor::Contiguous(rq.data(), DType::kF32, cd, {T, HQr, DHr}); + Tensor tko = Tensor::Contiguous(rk.data(), DType::kF32, cd, {T, HKVr, DHr}); + Tensor tgo = Tensor::Contiguous(rg.data(), DType::kF32, cd, {T, HQr, DHr}); + vt::AttnQkNormRopeGate(cq, tqo, tko, tgo, tqg, tkf, tqn, tkn, tcs, na3, ra3); + cpu.DestroyQueue(cq); + } + for (DeviceType dt : RegisteredDevices()) { + if (!OpAvailable(vt::OpId::kAttnQkNormRopeGate, dt)) continue; + CAPTURE(DeviceName(dt)); + vt::Backend& dev = vt::GetBackend(dt); + Queue q = dev.CreateQueue(); + const Device d{dt, 0}; + DevBufBytes dqg(dev, q, qg_bf.size() * 2), dkf(dev, q, kf_bf.size() * 2); + DevBuf dqn(dev, q, DHr), dkn(dev, q, DHr), dcs(dev, q, csr.size()); + DevBuf dqo(dev, q, rq.size()), dko(dev, q, rk.size()), dgo(dev, q, rg.size()); + dqg.Upload(qg_bf.data()); dkf.Upload(kf_bf.data()); + dqn.Upload(qnr); dkn.Upload(knr); dcs.Upload(csr); + Tensor tqg = Tensor::Contiguous(dqg.ptr(), DType::kBF16, d, {T, HQr * 2 * DHr}); + Tensor tkf = Tensor::Contiguous(dkf.ptr(), DType::kBF16, d, {T, HKVr * DHr}); + Tensor tqn = T1(dqn.ptr(), d, DHr), tkn = T1(dkn.ptr(), d, DHr); + Tensor tcs = T2(dcs.ptr(), d, T, ROTr); + Tensor tqo = Tensor::Contiguous(dqo.ptr(), DType::kF32, d, {T, HQr, DHr}); + Tensor tko = Tensor::Contiguous(dko.ptr(), DType::kF32, d, {T, HKVr, DHr}); + Tensor tgo = Tensor::Contiguous(dgo.ptr(), DType::kF32, d, {T, HQr, DHr}); + vt::AttnQkNormRopeGate(q, tqo, tko, tgo, tqg, tkf, tqn, tkn, tcs, na3, ra3); + CHECK(Nmse(rq, dqo.Download()) <= kNmseTol); + CHECK(Nmse(rk, dko.Download()) <= kNmseTol); + CHECK(Nmse(rg, dgo.Download()) <= kNmseTol); + dev.DestroyQueue(q); + } + } + const int64_t HQ = 3, HKV = 2, DH = 32, ROT = 16; const int64_t qg_outer = HQ * 2 * DH + 7, kf_outer = HKV * DH + 5; const std::vector qgate = RandomVec(static_cast(T * qg_outer), 881, -0.5f, 0.5f); diff --git a/tests/vt/test_ops_quant_repack.cpp b/tests/vt/test_ops_quant_repack.cpp index a1bb02b97..f3f80707f 100644 --- a/tests/vt/test_ops_quant_repack.cpp +++ b/tests/vt/test_ops_quant_repack.cpp @@ -201,8 +201,13 @@ TEST_CASE("G7 repacked GEMM matches plain at real model shapes / dtypes") { for (const Shape& s : shapes) { for (vt::DType odt : {vt::DType::kF32, vt::DType::kBF16}) { for (bool strided : {false, true}) { - CAPTURE(s.m); CAPTURE(s.n); CAPTURE(s.k); - CAPTURE(static_cast(odt)); CAPTURE(strided); + // One CAPTURE per line: same-line groups collide on + // DOCTEST_CAPTURE_ (clang/hipcc reject the redefinition). + CAPTURE(s.m); + CAPTURE(s.n); + CAPTURE(s.k); + CAPTURE(static_cast(odt)); + CAPTURE(strided); const int64_t nblocks = s.k / vt::cpu::kQK8_0; const std::vector plain = MakePlainQ8_0(s.n, nblocks, 0x11U); // Optionally over-allocate the activation row stride (a column slice of