From 45633aa315b728714c6edef53b845342b14925c6 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sun, 13 Sep 2026 05:56:49 +0000 Subject: [PATCH 01/10] spec(MODEL-MM-QWEN4-EXP): chunked H2D through a bounded pinned bounce ring The 67.56 GiB Qwen3.8-Flash-Next UD-IQ1_S still produces no token on gfx1151 after the two host-residency fixes landed, and the compute thread is still uninterruptible in svm_range_set_attr for about three quarters of its wchan samples. Releasing the spent source pages was necessary and is not sufficient, so the remaining suspect is the SHAPE of the transfer: RocmBackend::Copy hands the ROCr runtime one multi-GiB pageable, file-backed, CIFS-backed range and the runtime has to make all of it resident to DMA out of it. This spec takes llama.cpp's ring at the recorded pin -- four pinned buffers of 64 MiB with an event each, llama-model-loader.cpp:1440, :1449, :1496-1516 and :1591-1642 -- and states the one place it differs from ours, which is that the oracle takes that ring only when it is NOT using mmap. What we port is the shape, not the call, and the spec says so rather than claiming a parity it does not have. The reach is enumerated before the code because this is a shared seam and this row has already shipped a seam change whose blast radius nobody wrote down. One override of Backend::Copy is touched, no other backend is, and the staged path needs five terms to hold at once, so a 4 KiB norm keeps the existing single call byte for byte. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code] --- .agents/specs/rocm-chunked-pinned-h2d.md | 330 +++++++++++++++++++++++ 1 file changed, 330 insertions(+) create mode 100644 .agents/specs/rocm-chunked-pinned-h2d.md diff --git a/.agents/specs/rocm-chunked-pinned-h2d.md b/.agents/specs/rocm-chunked-pinned-h2d.md new file mode 100644 index 000000000..9036acc5b --- /dev/null +++ b/.agents/specs/rocm-chunked-pinned-h2d.md @@ -0,0 +1,330 @@ +# ROCm chunked H2D through a bounded pinned bounce ring + +Row: `MODEL-MM-QWEN4-EXP` +Issue: `ISSUE-LOCAL-01M2BZ5QK4XRETK48CXKSHKRDW` (row-owned, deferred by +`.agents/specs/rocm-host-residency-after-upload.md` § "Deferred, with an issue", +and made REQUIRED by that spec's §6a measurement) + +## 1. The defect + +`RocmBackend::Copy` is one call: + +```cpp +void Copy(Queue& q, void* dst, const void* src, size_t bytes) override { + Check(hipMemcpyAsync(dst, src, bytes, hipMemcpyDefault, AsStream(q)), "hipMemcpyAsync"); +} +``` + +(`src/vt/rocm/rocm_backend.hip:278-280` at `98e2cd7da`.) + +On `strix:gpu0` (gfx1151) that call is handed a **multi-GiB, pageable, +file-backed, CIFS-backed** source: a `PROT_READ MAP_PRIVATE` view of a GGUF +shard on `//192.168.68.102/Data`. To DMA out of a pageable range the ROCr +runtime has to make that range resident and describe it to the KFD, and +`svm_range_set_attr` is where that work lands. + +`.agents/specs/rocm-host-residency-after-upload.md` §6a measured the +consequence. The 67.56 GiB `Qwen3.8-Flash-Next UD-IQ1_S` loads with zero op +refusals, stages 29.69 GiB onto the board, and then never produces a token: two +1200 s runs, killed at the deadline, with the compute thread uninterruptible in +`svm_range_set_attr` for 153 of 196 and 148 of 198 `wchan` samples. Releasing +the spent source pages after upload (that spec's fix 1) and stopping the +prefault (fix 2) both landed and neither moved it: host `RssFile` during the +forward is still 20.2-20.4 GiB. That spec's §7 fourth risk names this outcome +and names this change as the consequence. + +So the remaining suspect is the SHAPE of the transfer, not the residency of +what it reads. + +## 2. Oracle + +`llama-cpp` pin `10bf611e5` (b10451), the recorded pin in +`.agents/oracles/llama-cpp.md`. The upload ring is in +`src/llama-model-loader.cpp`: + +- `:1440` — `constexpr size_t n_buffers = 4;` +- `:1449` — `const size_t buffer_size = alignment != 1 ? 64 * 1024 * 1024 + 2 * alignment : 1 * 1024 * 1024;` + under the comment "Buffer size: balance between memory usage and I/O + efficiency / 64MB works well for NVMe drives". +- `:1496-1516` — allocates `n_buffers` buffers from + `ggml_backend_dev_host_buffer_type(dev)` (pinned host memory) and one + `ggml_backend_event_t` per buffer. +- `:1591-1642` — the loop. For each chunk: `ggml_backend_event_synchronize` on + the buffer about to be reused, fill the buffer, `ggml_backend_tensor_set_async` + from the pinned buffer to the device, `ggml_backend_event_record`, advance + `buffer_idx` modulo `n_buffers`. +- `:1664` — the buffers are freed once, after the whole load. + +Peak PINNED host residency is therefore `n_buffers * buffer_size` for any model +size. Ours is O(the largest single weight) of pageable range that the driver has +to pin and describe, on every weight. + +**ONE HONEST DIFFERENCE, STATED RATHER THAN GLOSSED.** The oracle takes this +ring only when it is NOT using mmap (`:1456-1458`, `if (use_mmap || +check_tensors) return nullptr;`); with mmap it calls `ggml_backend_tensor_set` +straight from the mapped pointer, which is the same single-shot pageable copy we +do today. So this is not a case where the oracle refuses the shape we ship. What +the oracle supplies is the SHAPE — ring size, buffer size, event discipline, and +the fact that a bounded pinned ring is the sanctioned way to feed a device out of +host storage the device cannot address. We apply it to a source the oracle +reaches by `read()` and we reach by `memcpy` out of a mapping. That adaptation is +the design, and it is why this is its own spec rather than a port. + +## 3. Design + +A ring of `N` fixed pinned buffers inside `RocmBackend`, lazily created on the +first copy that qualifies. + +``` +for (off = 0; off < bytes; off += n) { + n = min(chunk, bytes - off); + i = next slot + if (slot[i] has an in-flight chunk) hipEventSynchronize(slot[i].event); + memcpy(slot[i].host, (const char*)src + off, n); + hipMemcpyAsync((char*)dst + off, slot[i].host, n, hipMemcpyHostToDevice, stream); + hipEventRecord(slot[i].event, stream); +} +``` + +The call still returns with work in flight on the stream, so `Backend::Copy`'s +asynchronous contract is unchanged. What changes is that the SOURCE is fully +consumed by the time `Copy` returns, which is strictly stronger than today. + +### 3a. N and the chunk size, justified against the oracle + +- **`N = 4`.** The oracle's `n_buffers` at `llama-model-loader.cpp:1440`, taken + rather than re-derived. Four slots keep one buffer filling while up to three + drain, which is what a ring buys over a double buffer, and four is the number + a runtime that loads models for a living ships. +- **`chunk = 64 MiB`.** The oracle's `buffer_size` at `:1449`. We drop its + `+ 2 * alignment` term, deliberately: that term exists to absorb the + `read_alignment()` padding of an `O_DIRECT`-style file read (`:1604-1631` + computes `aligned_offset`, `read_start`, `read_end` and trims the padding back + off). Our source is an in-memory mapping and our chunk boundaries are exact, so + there is no padding to absorb and a padded buffer would only waste pinned + bytes. +- **Total pinned residency: `4 * 64 MiB = 256 MiB`, for any model size.** That + is the oracle's bound, unchanged. +- **Threshold = one chunk (64 MiB).** Below one chunk the ring degenerates to + `memcpy` + `hipMemcpyAsync` with no overlap, no second slot ever used, and no + residency benefit that the driver's own internal staging does not already + give — it is strictly one extra copy of the bytes. A 4 KiB norm weight must + not pay a bounce, and at this threshold it does not: it takes the existing + single call, byte for byte. This boundary is not invented — it is where the + oracle's own `while` loop stops having more than one iteration. + +`VT_ROCM_PINNED_H2D_MIB` overrides the chunk in MiB and `0` disables the path +entirely, restoring `98e2cd7da` behaviour in the same binary. The A/B has to be +available in one binary because the model gate is a 20-minute run on a leased +box and rebuilding between arms is how a build difference gets to masquerade as +the effect. + +### 3b. THE REACH, ENUMERATED BEFORE THE CODE + +`Backend::Copy` is a shared seam. This change is confined to ONE override of it. +Files touched in `src/` and `include/`: +`src/vt/rocm/rocm_backend.hip` and the new HIP-free +`include/vt/rocm/rocm_pinned_h2d.h`. No other backend's `Copy`, no model, no +loader, no layer. + +**Which backends change: ROCm, and only ROCm.** `cuda_backend.cu`, +`cpu_backend.cpp`, `metal_backend.mm`, `vulkan_backend.cpp`, the XPU backend and +the Tenstorrent backend are not edited and their `Copy` is byte-identical. CUDA +is deliberately excluded: `dgx`, `thor` and `orin` produce tokens today, so +nothing there is owed this, and widening to CUDA would put a new host +`cudaEventSynchronize` on a path four measured campaigns depend on. A CUDA arm +is a later row if a CUDA measurement ever asks for one. + +**Which ROCm traffic changes.** The staged path is taken only when ALL FIVE +hold, and anything else takes the existing single call unchanged: + +1. `bytes >= chunk` (64 MiB by default). +2. `dst` resolves to DEVICE memory. D2H and H2H never bounce, so the sampler's + pinned-host download and every readback are untouched. +3. `src` is UNREGISTERED host memory. D2D never bounces. A source that is + already pinned (`hipMemoryTypeHost`) never bounces, because it is already + DMA-able and the ring exists only to create that property. A MANAGED source + (`hipMemoryTypeManaged` / `Unified`) never bounces either, because it is + already device-addressable — which is exactly the `VT_ROCM_MANAGED_ALLOC=1` + configuration, so that knob's behaviour is unchanged. +4. The stream is NOT capturing a graph. `hipEventSynchronize` inside a capture + region aborts the capture. The capture contract documented at + `rocm_backend.hip:334-348` already forbids "host<->device blocking copies" + inside the region, so a qualifying copy in there is already a contract + violation — but it would previously have been a silent one and would now be a + loud one, and changing WHICH failure a misuse produces is still a change. The + guard keeps capture byte-identical. +5. The ring resolves (`VT_ROCM_PINNED_H2D_MIB != 0` and the pinned allocation + succeeded). A failed `hipHostMalloc` falls back to the existing single call + rather than failing the load: 256 MiB of pinned memory is not worth refusing + a model over. + +**So what actually takes it.** Every H2D weight upload of 64 MiB or more on an +AMD board. That is a WIDER set than the release in +`rocm-host-residency-after-upload.md` §4a, and saying so is the point of this +paragraph: that release needed `mmap_fd >= 0`, so it fired only on the GGUF +keep-quant borrow. This predicate asks only "unregistered host source, big +enough", so it ALSO fires on safetensors borrows and on any heap buffer a loader +hands to `Copy`. On ROCm the families that reach it are therefore §4a's five +(Qwen4-Exp, GLM-MoE-DSA, GLM5-Next, Muse-Glimmer, the Qwen3.5 DFlash draft +head's rebound embedding table) PLUS every safetensors model whose weights clear +64 MiB — gemma, phi, minicpm, olmo2, stablelm, commandr, deepseek_v2, dots3, +nemotron_h, qwen3_vl and the rest — PLUS the EXL3 device loader's trellis +uploads, which §4a excluded for want of `mmap_fd`. The behaviour those models +see is identical bytes, bounded pinned host residency, and one extra host +`memcpy` per 64 MiB. + +**What is NOT covered by a test, and is recorded rather than claimed.** The only +family this change is MEASURED on is Qwen4-Exp, on one board. The device case in +§5 enters `RocmBackend::Copy` directly with a large pageable source, which is the +seam every one of those families reaches, so the CALL SITE is gated; no family +but Qwen4-Exp gets an end-to-end run here, for the same reason +`rocm-host-residency-after-upload.md` §5 gives — there is no harness that can +drive a second family's production entry point on a device. +`ISSUE-LOCAL-01M2CKN5516AKE7W2JVDV86Z8X` already owns that gap and this change +does not narrow it. + +### 3c. Where the decision lives + +The PURE parts — the chunk plan, the ring-reuse order, and the predicate that +decides whether to stage — go in `include/vt/rocm/rocm_pinned_h2d.h`, free of +HIP headers, and are table-tested in the ordinary CPU build. That mirrors +`include/vt/rocm/rocm_arch.h` and `ResolveMemoryPolicy` exactly, and for the +reason that header states about `CapabilityFromGcnArch`: the piece a wrong +answer breaks silently is the piece that must be gated on a runner with no AMD +GPU. `rocm_backend.hip` reads two `hipPointerGetAttributes`, one +`hipStreamIsCapturing`, and calls it. + +### 3d. The ring is deliberately never freed + +The pinned slots and events are allocated on first use and leaked at process +exit. `RocmBackend` instances live in a function-local `static +std::vector>` in the registrar, so a destructor +would run during static destruction, and calling `hipHostFree` / `hipEventDestroy` +after the HIP runtime has begun tearing down is a hazard this file does not +have today (`exec_`, the `hipGraphExec_t`, is likewise never destroyed). 256 MiB +returned to the OS at `exit()` buys nothing and a teardown-order crash costs a +measurement. Stated here because "it leaks" must be a decision on the record, +not something a reader discovers. + +## 4. Tests — red first + +Commit order is the red: the test commit lands before the implementation commit, +so the failing run is reproducible by building the tree at the test commit. + +**`tests/vt/test_rocm_pinned_h2d.cpp`** — new, UNCONDITIONAL (no HIP needed), +registered beside `test_rocm_arch`. Drives the pure header with fakes: + +1. **The plan.** `bytes = 200 MiB`, `chunk = 64 MiB` gives 4 chunks of + 64/64/64/8 MiB, contiguous offsets summing to `bytes`, and no chunk larger + than `chunk`. Mutating the chunk size to the whole buffer produces 1 chunk of + 200 MiB and fails this. +2. **The ring waits before it reuses.** With `N = 4` and 9 chunks, slot 0 is + waited on before chunks 4 and 8 and NOT before chunk 0. The fake records the + wait/stage/enqueue order; deleting the wait fails it. +3. **The bytes survive.** The fake stage/enqueue pair reassembles the output and + it is `memcmp`-identical to the input, over a size that is NOT a multiple of + the chunk. +4. **The predicate truth table.** Staged only for {unregistered host src, device + dst, `bytes >= chunk`, not capturing, `chunk != 0`}. Every other row of the + table is direct. Deleting the size term, the src-kind term, the dst term or + the capture term each flips a row. + +**`tests/vt/test_backend_cross_device.cpp`** — one new case, which is the +REACHABILITY conviction. It enters `RocmBackend::Copy` through +`vt::Backend&`, the production seam, with a 200 MiB pageable `std::vector` +source and a device destination, and asserts: + +- the downloaded bytes are `memcmp`-identical to the source (bit-exactness is + this file's declared bar for a pure copy path); +- `PinnedH2DSnapshot()` shows the copy took the ring — `staged_copies` grew by + one and `chunks` by four — and that `max_chunk_bytes <= 64 MiB`, which is the + bounded-host-residency assertion in the form this harness can actually make; +- a SMALL copy on the same backend grows `direct_copies` and not + `staged_copies`, so the threshold is gated in the same case. + +Deleting the staged branch from `rocm_backend.hip` leaves `staged_copies` at 0 +and fails this case. That is the mutation that proves the wiring, and it is the +one `AGENTS.md` "Nothing lands dead" asks for. + +This case measures nothing on a build with no ROCm device, exactly like every +other case in that file, and the gate line below therefore names the board it +was run on. + +## 5. Gates + +Every selector names its binary and prints its case and assertion counts. This +row has already shipped two selectors that matched nothing and reported +`Status: SUCCESS!`; a selector that matches nothing is indistinguishable from +one that passes. + +| binary | selector | cases | assertions | +|---|---|---|---| +| `test_rocm_pinned_h2d` | (none — whole binary) | RECORDED AT §7 | RECORDED AT §7 | +| `test_backend_cross_device` | `-tc=*pinned bounce*` | RECORDED AT §7 | RECORDED AT §7 | +| `test_backend_cross_device` | (none — whole binary) | RECORDED AT §7 | RECORDED AT §7 | +| `test_backend_cross_device` | `-tc=*DSA*` | RECORDED AT §7 | RECORDED AT §7 | + +Baseline to hold at `98e2cd7da`: `test_backend_cross_device` whole binary +**60 cases / 84833 assertions**, `-tc=*DSA*` **2 / 273**. + +```sh +python3 scripts/check-agent-record.py +python3 scripts/check-commit-style.py --range origin/main..HEAD +python3 scripts/check-commit-trailers.py --range origin/main..HEAD +python3 scripts/check-pr-size.py --base origin/main --head HEAD \ + --branch row/MODEL-MM-QWEN4-EXP-ROCM-CHUNKED-H2D +``` + +**The model gate is the deliverable.** `/workspace/ckpt/qwen4exp-flash-next-iq1s` +shard 1, `examples/vllm-server` with `--device auto`, `VT_ROCM_MANAGED_ALLOC` +unset, on `strix:gpu0` inside an `rc` lease. Does it produce a token? If it does, +that is this row's G3 and the numbers follow — TTFT, prefill and decode tok/s, +load seconds, peak host `VmHWM` / `RssFile` / `RssAnon`, device memory from +`/sys/class/drm/card*/device/mem_info_vram_used` (`rocm-smi` is not on `PATH` in +the leased container), the exact build and run recipe, revisions, artifact sizes, +environment and contention. Generation is repeated at least three times and +reported as a spread, because gfx1151 fails about two of five identical greedy +runs with an illegal GPU memory access +(`ISSUE-LOCAL-01M2BY2M2ATNVR3XQKV2DB1BJD`) and one green there is unreplicated. +If it does not produce a token, NO number is recorded and the `wchan` +distribution is reported instead, including whether it moved off +`svm_range_set_attr`. + +The artifact's compiled feature set is asserted before it is timed: `ldd` for +`libamdhip64.so.7` and the HIP arch it was built for. + +## 6. Risks + +- **The stall survives this too.** Then the trigger is neither the residency of + the source nor the shape of the transfer, and the next hypothesis is the + allocation side — 29.69 GiB of `hipMalloc` on a 33.27 GB board behind a 96 GiB + carve, or the CIFS mount, which §6a already named as an unseparated confound. + A negative result with a `wchan` distribution is the reportable outcome, not a + failure of the change. +- **An extra `memcpy` per 64 MiB slows the load.** Load seconds are recorded on + both arms of `VT_ROCM_PINNED_H2D_MIB` in one binary, so the cost is measured + rather than argued. +- **256 MiB of pinned memory on a 31 GiB host.** It is allocated once, it is the + oracle's own bound, and it replaces an unbounded pinned range the driver was + creating per copy. +- **A qualifying copy inside a graph capture.** Guarded by term 4, and the guard + is in the truth table. +- **`hipPointerGetAttributes` leaves a sticky error for an unregistered host + pointer.** It returns `hipErrorInvalidValue` for one, which is the very answer + we want, and the last error is cleared with `hipGetLastError()` immediately so + it cannot poison the next `Check`. + +## 7. Outcome + +RECORDED ON LANDING. + +## 8. Stop conditions + +- The ring cannot be placed without changing `Backend::Copy`'s asynchronous + contract: STOP and return `NEEDS_DECISION`. +- The change would have to widen past `RocmBackend::Copy` to be effective: STOP + and return `NEEDS_DECISION` rather than editing a second backend. +- `strix:gpu0` is unreachable or the controller is down: the model gate is + reported UNVERIFIED. It is never replaced by an `ssh` plus a file mutex the + fleet cannot see. From 054546bfeeee997dcd719e6b71dfe32bd3a10f53 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sun, 13 Sep 2026 06:02:32 +0000 Subject: [PATCH 02/10] test(MODEL-MM-QWEN4-EXP): gate the pinned H2D ring's plan, reuse order and predicate The decision behind the bounce ring goes in a HIP-free header and is table-tested on a runner with no AMD GPU, which is the same split include/vt/rocm/rocm_arch.h already makes for ResolveMemoryPolicy and for the same reason: a ring that silently never engages is indistinguishable from one that works. The bytes are identical either way, the model loads either way, and only the host residency differs, which no gate in this tree reads by accident. So the chunk arithmetic, the cross-call slot reuse and the five-term predicate are gated here, and the half that cannot be gated here -- that RocmBackend::Copy actually calls any of it -- gets a device case in test_backend_cross_device that asserts the instrument beside the bytes. That case is RED at this commit: nothing calls the ring yet, so staged_copies stays at zero. Eight mutations were run against the pure half and each convicted, with the binary proven changed by md5 on every arm: the chunk size widened to the whole buffer, the slot wait deleted, each of the four predicate terms deleted, the cross-call in-flight state reset, and the knob parse made to answer 0 on a typo instead of the default. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code] --- include/vt/rocm/rocm_pinned_h2d.h | 218 ++++++++++++++++++ tests/CMakeLists.txt | 7 + tests/vt/test_backend_cross_device.cpp | 120 ++++++++++ tests/vt/test_rocm_pinned_h2d.cpp | 301 +++++++++++++++++++++++++ 4 files changed, 646 insertions(+) create mode 100644 include/vt/rocm/rocm_pinned_h2d.h create mode 100644 tests/vt/test_rocm_pinned_h2d.cpp diff --git a/include/vt/rocm/rocm_pinned_h2d.h b/include/vt/rocm/rocm_pinned_h2d.h new file mode 100644 index 000000000..e029051d6 --- /dev/null +++ b/include/vt/rocm/rocm_pinned_h2d.h @@ -0,0 +1,218 @@ +// The DECISION and the LOOP behind RocmBackend::Copy's bounded pinned bounce +// ring, deliberately free of HIP headers. +// +// Same split, and the same reason, as include/vt/rocm/rocm_arch.h and +// `ResolveMemoryPolicy`: the piece a wrong answer breaks SILENTLY is the piece +// that has to be gated on an ordinary CPU runner with no AMD GPU. A ring that +// quietly never engages looks exactly like one that works — the bytes are still +// correct, the model still loads, and only the host residency differs — so the +// chunk arithmetic, the reuse order and the predicate are compiled and +// table-tested in tests/vt/test_rocm_pinned_h2d.cpp on a machine with no HIP at +// all. src/vt/rocm/rocm_backend.hip reads two pointer attributes and one +// capture status, and calls into here. +// +// Oracle: llama-cpp pin 10bf611e5 (b10451), src/llama-model-loader.cpp:1440 +// (`n_buffers = 4`), :1449 (64 MiB per buffer), :1591-1642 (the loop: wait on +// the slot's event, fill the pinned buffer, async-upload, record the event, +// advance modulo n_buffers). See .agents/specs/rocm-chunked-pinned-h2d.md §2 +// for the one place the oracle's ring and ours are NOT the same thing. +#pragma once + +#include +#include +#include +#include +#include + +namespace vt::rocm { + +// Oracle: llama-model-loader.cpp:1440. Four slots keep one buffer filling while +// up to three drain; taken rather than re-derived. +inline constexpr size_t kPinnedH2DBuffers = 4; + +// Oracle: llama-model-loader.cpp:1449, "64MB works well for NVMe drives". The +// oracle's `+ 2 * alignment` term is deliberately dropped: it absorbs the +// read-alignment padding of a file read (:1604-1631 computes it and then trims +// it back off), and our source is an in-memory mapping whose chunk boundaries +// are exact. +inline constexpr size_t kPinnedH2DChunkBytesDefault = static_cast(64) << 20; + +// Total pinned host residency, for a model of ANY size. This is the number the +// whole change exists to bound, so it is written down rather than left to be +// multiplied out by a reader. +inline constexpr size_t kPinnedH2DRingBytes = + kPinnedH2DBuffers * kPinnedH2DChunkBytesDefault; + +// What hipPointerGetAttributes said about a pointer, reduced to the four +// answers this decision needs. `kUnregisteredHost` is ordinary host storage the +// HIP runtime knows nothing about — a std::vector, a heap block, or the +// PROT_READ MAP_PRIVATE view of a GGUF shard that this change exists for. +// `kOther` covers managed and unified pointers, which are ALREADY device +// addressable and must never be bounced. +enum class PtrKind { kUnregisteredHost, kPinnedHost, kDevice, kOther }; + +// Everything RocmBackend::Copy knows at the moment it has to choose a path. +struct StagedH2DInputs { + PtrKind src = PtrKind::kOther; + PtrKind dst = PtrKind::kOther; + size_t bytes = 0; + size_t chunk_bytes = 0; // 0 disables the ring entirely + bool stream_capturing = false; + bool ring_available = false; +}; + +// FIVE terms, all required. Spelled as one expression so the truth table in +// tests/vt/test_rocm_pinned_h2d.cpp can flip exactly one at a time. +// +// * `chunk_bytes != 0` — VT_ROCM_PINNED_H2D_MIB=0 restores the pre-change +// single call in the same binary, which is what makes +// the load-time A/B one build instead of two. +// * `ring_available` — a failed hipHostMalloc falls back rather than +// refusing a model over 256 MiB of pinned memory. +// * `!stream_capturing` — hipEventSynchronize aborts a graph capture. The +// capture contract already forbids blocking copies in +// the region, so this guard changes no legal program; +// it keeps an ILLEGAL one failing the way it used to. +// * `dst == kDevice` — D2H and H2H never bounce, so every readback and the +// sampler's download are untouched. +// * `src == kUnregisteredHost` — D2D never bounces; an already-pinned source is +// already DMA-able; a managed source is already device +// addressable, so VT_ROCM_MANAGED_ALLOC=1 is unchanged. +// * `bytes >= chunk_bytes` — below one chunk the ring is strictly one extra +// copy of the bytes with no overlap and no second slot +// ever used. A 4 KiB norm weight must not pay a bounce. +constexpr bool ShouldStageH2D(const StagedH2DInputs& in) { + return in.chunk_bytes != 0 && in.ring_available && !in.stream_capturing && + in.dst == PtrKind::kDevice && in.src == PtrKind::kUnregisteredHost && + in.bytes >= in.chunk_bytes; +} + +// VT_ROCM_PINNED_H2D_MIB, in MiB. Absent or empty takes the default; an +// explicit `0` disables the ring; anything that does not parse as a +// non-negative integer takes the DEFAULT rather than 0, because a typo must not +// silently turn the feature off — a disabled ring and a working one differ only +// in host residency, which no gate in this tree reads by accident. +inline size_t ParsePinnedH2DChunkBytes(std::string_view v) { + if (v.empty()) return kPinnedH2DChunkBytesDefault; + size_t mib = 0; + for (char c : v) { + if (c < '0' || c > '9') return kPinnedH2DChunkBytesDefault; + mib = mib * 10 + static_cast(c - '0'); + if (mib > (static_cast(1) << 20)) return kPinnedH2DChunkBytesDefault; + } + return mib << 20; +} + +inline size_t PinnedH2DChunkBytes() { + static const size_t bytes = [] { + const char* e = std::getenv("VT_ROCM_PINNED_H2D_MIB"); + return ParsePinnedH2DChunkBytes(e != nullptr ? std::string_view(e) : std::string_view()); + }(); + return bytes; +} + +// Ring state, owned by the backend and LIVING ACROSS CALLS. That is the part +// worth testing: a second Copy must wait on the previous Copy's still-in-flight +// chunk before it refills that slot, and a ring that only tracked reuse within +// one call would overwrite bytes a DMA was still reading. +struct StagedH2DRing { + size_t n_buffers = kPinnedH2DBuffers; + size_t next_slot = 0; + bool in_flight[kPinnedH2DBuffers] = {}; +}; + +// The loop, generic over the three device actions so it is drivable with fakes. +// Returns the number of chunks enqueued (0 when the ring is disabled). +// +// wait(slot) — block until that slot's last upload is done +// stage(slot, offset, bytes) — copy source bytes into the slot's buffer +// enqueue(slot, offset, bytes) — async upload from the slot, then record +// +// The call returns with the last chunks still in flight on the stream, so +// Backend::Copy's asynchronous contract is unchanged. What IS stronger than +// before is that the SOURCE has been fully consumed by the time it returns. +template +inline size_t RunStagedH2D(StagedH2DRing& ring, size_t bytes, size_t chunk_bytes, + Wait wait, Stage stage, Enqueue enqueue) { + if (chunk_bytes == 0 || ring.n_buffers == 0 || ring.n_buffers > kPinnedH2DBuffers) { + return 0; + } + size_t chunks = 0; + for (size_t off = 0; off < bytes; off += chunk_bytes) { + const size_t n = (bytes - off) < chunk_bytes ? (bytes - off) : chunk_bytes; + const size_t slot = ring.next_slot; + if (ring.in_flight[slot]) { + wait(slot); + ring.in_flight[slot] = false; + } + stage(slot, off, n); + enqueue(slot, off, n); + ring.in_flight[slot] = true; + ring.next_slot = (slot + 1) % ring.n_buffers; + ++chunks; + } + return chunks; +} + +// The instrument. A bounded-residency guarantee changes no byte, so a +// byte-equality case cannot see whether the ring engaged at all — the same +// reason NoteGgufPrefaultedSpan exists for the prefault. These counters are what +// tests/vt/test_backend_cross_device.cpp asserts on to prove the production call +// site in rocm_backend.hip is reached. +struct PinnedH2DStats { + uint64_t staged_copies = 0; // Copy() calls that took the ring + uint64_t direct_copies = 0; // Copy() calls that stayed on the single call + uint64_t chunks = 0; // pinned chunks enqueued, in total + size_t max_chunk_bytes = 0; // largest single pinned -> device transfer + size_t ring_bytes = 0; // pinned bytes actually allocated, 0 until used +}; + +namespace detail { +struct PinnedH2DCounters { + std::atomic staged_copies{0}; + std::atomic direct_copies{0}; + std::atomic chunks{0}; + std::atomic max_chunk_bytes{0}; + std::atomic ring_bytes{0}; +}; +// Relaxed atomics, not a plain counter: direct_copies is incremented on EVERY +// Backend::Copy, which is a hot path (about 1,361 ResidentWeight calls per +// forward step on this row's checkpoint), and Copy is reachable from more than +// one thread. A relaxed fetch_add is nanoseconds beside a hipMemcpyAsync. +inline PinnedH2DCounters& Counters() { + static PinnedH2DCounters c; + return c; +} +} // namespace detail + +inline void NotePinnedH2DStaged(size_t chunks, size_t max_chunk_bytes) { + auto& c = detail::Counters(); + c.staged_copies.fetch_add(1, std::memory_order_relaxed); + c.chunks.fetch_add(chunks, std::memory_order_relaxed); + size_t prev = c.max_chunk_bytes.load(std::memory_order_relaxed); + while (max_chunk_bytes > prev && + !c.max_chunk_bytes.compare_exchange_weak(prev, max_chunk_bytes, + std::memory_order_relaxed)) { + } +} + +inline void NotePinnedH2DDirect() { + detail::Counters().direct_copies.fetch_add(1, std::memory_order_relaxed); +} + +inline void NotePinnedH2DRingBytes(size_t bytes) { + detail::Counters().ring_bytes.store(bytes, std::memory_order_relaxed); +} + +inline PinnedH2DStats PinnedH2DSnapshot() { + auto& c = detail::Counters(); + PinnedH2DStats s; + s.staged_copies = c.staged_copies.load(std::memory_order_relaxed); + s.direct_copies = c.direct_copies.load(std::memory_order_relaxed); + s.chunks = c.chunks.load(std::memory_order_relaxed); + s.max_chunk_bytes = c.max_chunk_bytes.load(std::memory_order_relaxed); + s.ring_bytes = c.ring_bytes.load(std::memory_order_relaxed); + return s; +} + +} // namespace vt::rocm diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 5d22ed30d..97ef3b12d 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -2685,6 +2685,13 @@ set_tests_properties(test_backend_cross_device_vt_attn_decode_d128 PROPERTIES # of that skeleton carrying a DECISION is gated on a CPU-only CI runner with no # AMD GPU. The rest of the ROCm skeleton is under VLLM_CPP_HIP below. vllm_cpp_add_test(test_rocm_arch vt/test_rocm_arch.cpp) +# The pure half of the ROCm pinned H2D bounce ring (chunk plan, cross-call slot +# reuse, the five-term predicate). UNCONDITIONAL for exactly the reason +# test_rocm_arch is: include/vt/rocm/rocm_pinned_h2d.h carries no HIP header, and +# a ring that silently never engages produces identical bytes -- only the host +# residency differs -- so the decision has to be gated on a runner with no AMD +# GPU. The reachability half is in test_backend_cross_device. +vllm_cpp_add_test(test_rocm_pinned_h2d vt/test_rocm_pinned_h2d.cpp) vllm_cpp_add_test(test_rocm_f16_contract vt/test_rocm_f16_contract.cpp) # The grow-only per-stream device scratch pool (#2712). UNCONDITIONAL for the # same reason test_rocm_arch is: src/vt/grow_only_stream_scratch.h is free of diff --git a/tests/vt/test_backend_cross_device.cpp b/tests/vt/test_backend_cross_device.cpp index 22240aae3..d0d152290 100644 --- a/tests/vt/test_backend_cross_device.cpp +++ b/tests/vt/test_backend_cross_device.cpp @@ -46,6 +46,7 @@ #include "vt/quant.h" #include "vt/recipes.h" #include "vt/rocm/rocm_arch.h" +#include "vt/rocm/rocm_pinned_h2d.h" #include "vt/rocm/rocm_runtime.h" namespace { @@ -229,6 +230,125 @@ TEST_CASE("device Copy/Memset are BIT-EXACT against the host bytes") { } } +// --- The bounded pinned bounce ring (.agents/specs/rocm-chunked-pinned-h2d.md) +// +// THE REACHABILITY CONVICTION for that change, and the only one there can be. +// The ring changes no byte: the model still loads, the tokens are still the +// same, and only the HOST RESIDENCY of the transfer differs, which is a property +// no byte-equality case in this file can see. So the instrument is asserted +// beside the bytes, and the mutation that fails this case is deleting the staged +// branch from RocmBackend::Copy in src/vt/rocm/rocm_backend.hip. +// +// It enters through `vt::Backend&`, the production seam every model loader +// reaches — dense_attn::ResidentWeight, the qwen3_5.cpp twin, the MoE expert +// towers, the EXL3 device loader — rather than through any test hook. +// +// TWO ARMS, and neither of them is a silent skip. On a board that takes the +// plain hipMalloc branch the destination is device memory and the copy MUST take +// the ring. On a board that takes the hipMallocManaged branch (gfx1103, or +// VT_ROCM_MANAGED_ALLOC=1) the destination is already device-addressable and the +// copy MUST NOT take it — the ring exists only to create a property managed +// memory already has. Each arm asserts its own direction and says which it ran. +TEST_CASE("a large pageable H2D takes the pinned bounce ring, and a small one does not") { + // One chunk short of three, so the last chunk is ragged and the count is not + // a number the loop could reach by accident. + const size_t kChunk = vt::rocm::kPinnedH2DChunkBytesDefault; + const size_t kBig = kChunk * 2 + (kChunk / 2); + const size_t kSmall = 4096; // a norm weight: it must not pay a bounce + + bool ran_on_a_board = false; + for (DeviceType dt : RegisteredDevices()) { + if (dt != DeviceType::kROCM) continue; + ran_on_a_board = true; + CAPTURE(DeviceTag(dt)); + vt::Backend& dev = vt::GetBackend(dt); + Queue q = dev.CreateQueue(); + + // A PAGEABLE host source: an ordinary std::vector the HIP runtime knows + // nothing about. That is what a GGUF mmap view is, as far as this decision + // is concerned, and it is the only source kind the ring accepts. + std::vector src(kBig); + { + uint32_t s = 20260913u; + for (size_t i = 0; i < kBig; ++i) { + s = s * 1664525u + 1013904223u; + src[i] = static_cast(s >> 24); + } + } + + const bool managed = vt::rocm::ManagedAllocActive(0); + const vt::rocm::PinnedH2DStats before = vt::rocm::PinnedH2DSnapshot(); + + void* p = dev.Alloc(kBig); + dev.Copy(q, p, src.data(), kBig); + dev.Synchronize(q); + + const vt::rocm::PinnedH2DStats after_big = vt::rocm::PinnedH2DSnapshot(); + + // THE BYTES FIRST. The bar for a pure copy path in this file is + // bit-exactness, and a ring that reassembled the buffer wrongly would be a + // far worse defect than the residency it fixes. + std::vector back(kBig, 0); + dev.Copy(q, back.data(), p, kBig); + dev.Synchronize(q); + CHECK(std::memcmp(src.data(), back.data(), kBig) == 0); + + // A SMALL copy on the same backend, in the same case, so the threshold is + // gated beside the path it guards. + void* small = dev.Alloc(kSmall); + dev.Copy(q, small, src.data(), kSmall); + dev.Synchronize(q); + const vt::rocm::PinnedH2DStats after_small = vt::rocm::PinnedH2DSnapshot(); + + // ONE std::string, not a `const char*` chained with `<<`: doctest prints the + // literal as `1` (see the DeviceTag comment at the top of this file), which + // would turn the one line that says WHICH arm ran into noise. + const std::string note = + std::string("pinned H2D: managed_alloc=") + std::to_string(managed ? 1 : 0) + + " staged=" + std::to_string(after_small.staged_copies - before.staged_copies) + + " direct=" + std::to_string(after_small.direct_copies - before.direct_copies) + + " chunks=" + std::to_string(after_big.chunks - before.chunks) + + " max_chunk=" + std::to_string(after_big.max_chunk_bytes) + + " ring_bytes=" + std::to_string(after_big.ring_bytes); + MESSAGE(note); + + if (managed) { + // The managed arm: device memory is host-addressable here, so the ring is + // deliberately not engaged and every copy is direct. + CHECK(after_small.staged_copies == before.staged_copies); + CHECK(after_small.direct_copies > before.direct_copies); + } else { + // The staging arm. This is the assertion the whole change is for. + CHECK(after_big.staged_copies == before.staged_copies + 1); + CHECK(after_big.chunks == before.chunks + 3); + // BOUNDED HOST RESIDENCY, in the form this harness can actually make: no + // single pinned-to-device transfer ever exceeds one chunk, whatever the + // size of the copy. Mutating the chunk size to the whole buffer makes this + // one kBig-sized transfer and fails here. + CHECK(after_big.max_chunk_bytes <= kChunk); + // Read against the RESOLVED chunk, not the compile-time constant, so the + // case still measures something when VT_ROCM_PINNED_H2D_MIB is set for an + // A/B. A ring that was never allocated reports 0 and fails either way. + CHECK(after_big.ring_bytes == + vt::rocm::kPinnedH2DBuffers * vt::rocm::PinnedH2DChunkBytes()); + CHECK(after_big.ring_bytes > 0); + // The D2H readback and the 4 KiB upload both stayed on the single call. + CHECK(after_small.staged_copies == after_big.staged_copies); + CHECK(after_small.direct_copies >= after_big.direct_copies + 2); + } + + dev.Free(small); + dev.Free(p); + dev.DestroyQueue(q); + } + // NOT an assertion: a CPU-only or CUDA-only build registers no ROCm backend + // and this case measures nothing, exactly like every other case in this file. + // It is printed so a green run says which of the two it was. + const std::string ran = std::string("pinned H2D case ran on a ROCm board: ") + + std::to_string(ran_on_a_board ? 1 : 0); + MESSAGE(ran); +} + // The bf16<->f32 casts are a pure ELEMENTWISE CODEC: no reduction, no // reassociation, one rounding on store. So the bar here is BIT-EXACTNESS against // the CPU reference, not NMSE — CastF32 (bf16 -> f32) is an exact widening, and diff --git a/tests/vt/test_rocm_pinned_h2d.cpp b/tests/vt/test_rocm_pinned_h2d.cpp new file mode 100644 index 000000000..14245df86 --- /dev/null +++ b/tests/vt/test_rocm_pinned_h2d.cpp @@ -0,0 +1,301 @@ +// The pure half of RocmBackend::Copy's bounded pinned bounce ring: +// .agents/specs/rocm-chunked-pinned-h2d.md §4. +// +// UNCONDITIONAL, no HIP required, for the reason test_rocm_arch.cpp is +// unconditional: include/vt/rocm/rocm_pinned_h2d.h is deliberately free of HIP +// headers, and it holds the only parts of this change that a wrong answer breaks +// SILENTLY. A ring that never engages produces identical bytes and loads the +// same model; only the host residency differs, and no gate in this tree reads +// that by accident. So the chunk arithmetic, the cross-call reuse order and the +// five-term predicate are gated here, on a runner with no AMD GPU. +// +// The OTHER half — that rocm_backend.hip actually calls this — cannot be gated +// here and is not pretended to be. It is gated by the device case +// "large pageable H2D takes the pinned bounce ring" in +// tests/vt/test_backend_cross_device.cpp, whose mutation is deleting the staged +// branch from the backend. +#include + +#include +#include +#include +#include + +#include "vt/rocm/rocm_pinned_h2d.h" + +namespace { + +using vt::rocm::kPinnedH2DBuffers; +using vt::rocm::kPinnedH2DChunkBytesDefault; +using vt::rocm::ParsePinnedH2DChunkBytes; +using vt::rocm::PtrKind; +using vt::rocm::RunStagedH2D; +using vt::rocm::ShouldStageH2D; +using vt::rocm::StagedH2DInputs; +using vt::rocm::StagedH2DRing; + +// A fake device: N pinned slots, an ordered event log, and a reassembled +// destination. Everything the real .hip does, minus HIP. +struct FakeRing { + explicit FakeRing(size_t n_buffers, size_t chunk_bytes) + : slots(n_buffers, std::vector(chunk_bytes, 0)) { + ring.n_buffers = n_buffers; + } + + StagedH2DRing ring; + std::vector> slots; + std::vector log; + // What the "device" ended up holding, and how big each upload was. + std::vector dst; + std::vector chunk_sizes; + std::vector chunk_offsets; + + size_t Run(const std::vector& src, size_t chunk_bytes) { + return RunStagedH2D( + ring, src.size(), chunk_bytes, + [&](size_t slot) { log.push_back("wait " + std::to_string(slot)); }, + [&](size_t slot, size_t off, size_t n) { + log.push_back("stage " + std::to_string(slot)); + std::memcpy(slots[slot].data(), src.data() + off, n); + }, + [&](size_t slot, size_t off, size_t n) { + log.push_back("enqueue " + std::to_string(slot)); + if (dst.size() < off + n) dst.resize(off + n, 0); + std::memcpy(dst.data() + off, slots[slot].data(), n); + chunk_sizes.push_back(n); + chunk_offsets.push_back(off); + }); + } +}; + +std::vector Pattern(size_t n, uint32_t seed) { + std::vector v(n); + uint32_t s = seed; + for (size_t i = 0; i < n; ++i) { + s = s * 1664525u + 1013904223u; + v[i] = static_cast(s >> 24); + } + return v; +} + +} // namespace + +// --------------------------------------------------------------------------- +// 1. The plan: the chunk count, the sizes, and the bound. +// --------------------------------------------------------------------------- +TEST_CASE("the pinned ring splits a large copy into bounded chunks") { + // 200 MiB against a 64 MiB chunk: 64 + 64 + 64 + 8. + constexpr size_t kChunk = static_cast(64) << 20; + constexpr size_t kBytes = static_cast(200) << 20; + + StagedH2DRing ring; + std::vector sizes; + std::vector offsets; + const size_t chunks = RunStagedH2D( + ring, kBytes, kChunk, [](size_t) {}, [](size_t, size_t, size_t) {}, + [&](size_t, size_t off, size_t n) { + offsets.push_back(off); + sizes.push_back(n); + }); + + CHECK(chunks == 4); + REQUIRE(sizes.size() == 4); + CHECK(sizes[0] == kChunk); + CHECK(sizes[1] == kChunk); + CHECK(sizes[2] == kChunk); + CHECK(sizes[3] == (static_cast(8) << 20)); + + // THE BOUND, which is the whole point: no single upload reads more than one + // chunk out of the pageable source, whatever the copy's total size. Mutating + // the chunk size to the whole buffer makes this one 200 MiB upload. + size_t total = 0; + size_t expect_off = 0; + for (size_t i = 0; i < sizes.size(); ++i) { + CHECK(sizes[i] <= kChunk); + CHECK(offsets[i] == expect_off); + expect_off += sizes[i]; + total += sizes[i]; + } + CHECK(total == kBytes); +} + +TEST_CASE("the pinned ring leaves a copy smaller than one chunk as one chunk") { + StagedH2DRing ring; + size_t chunks = 0; + size_t only_size = 0; + chunks = RunStagedH2D( + ring, 1000, 4096, [](size_t) {}, [](size_t, size_t, size_t) {}, + [&](size_t, size_t, size_t n) { only_size = n; }); + CHECK(chunks == 1); + CHECK(only_size == 1000); +} + +TEST_CASE("the pinned ring enqueues nothing when the chunk size is zero") { + // VT_ROCM_PINNED_H2D_MIB=0 restores the pre-change single call in the SAME + // binary, which is what makes the load-time A/B one build instead of two. + StagedH2DRing ring; + size_t enqueued = 0; + const size_t chunks = RunStagedH2D( + ring, 1u << 30, 0, [](size_t) {}, [](size_t, size_t, size_t) {}, + [&](size_t, size_t, size_t) { ++enqueued; }); + CHECK(chunks == 0); + CHECK(enqueued == 0); +} + +// --------------------------------------------------------------------------- +// 2. The reuse order — WITHIN a call and, the part that matters, ACROSS calls. +// --------------------------------------------------------------------------- +TEST_CASE("the pinned ring waits on a slot before it refills it, not before first use") { + FakeRing fake(kPinnedH2DBuffers, 16); + const std::vector src = Pattern(16 * 9, 7); // 9 chunks over 4 slots + const size_t chunks = fake.Run(src, 16); + CHECK(chunks == 9); + + // The first four chunks touch four distinct slots and wait on none of them. + for (size_t i = 0; i < 4; ++i) { + CHECK(fake.log[i * 2] == "stage " + std::to_string(i)); + CHECK(fake.log[i * 2 + 1] == "enqueue " + std::to_string(i)); + } + // Chunk 4 is the first reuse of slot 0, and it MUST wait first. Deleting the + // wait removes exactly this line and overwrites bytes a DMA is still reading. + CHECK(fake.log[8] == "wait 0"); + CHECK(fake.log[9] == "stage 0"); + // Chunk 8 is the second reuse of slot 0. + CHECK(fake.log[8 + 4 * 3] == "wait 0"); + + size_t waits = 0; + for (const auto& line : fake.log) { + if (line.rfind("wait ", 0) == 0) ++waits; + } + CHECK(waits == 5); // chunks 4..8 each reuse a slot; chunks 0..3 do not +} + +TEST_CASE("the pinned ring waits across calls, not only within one") { + // THIS is the case a ring that tracked reuse only within one Copy would fail. + // Two 2-chunk copies over a 4-slot ring: the second call starts at slot 2, so + // it waits on nothing; a THIRD 2-chunk copy comes back round to slot 0, which + // the FIRST call left in flight, and must wait. + FakeRing fake(kPinnedH2DBuffers, 16); + const std::vector src = Pattern(32, 11); + CHECK(fake.Run(src, 16) == 2); + CHECK(fake.Run(src, 16) == 2); + size_t waits_after_two = 0; + for (const auto& line : fake.log) { + if (line.rfind("wait ", 0) == 0) ++waits_after_two; + } + CHECK(waits_after_two == 0); + + fake.log.clear(); + CHECK(fake.Run(src, 16) == 2); + REQUIRE(fake.log.size() >= 2); + CHECK(fake.log[0] == "wait 0"); + CHECK(fake.log[3] == "wait 1"); +} + +// --------------------------------------------------------------------------- +// 3. The bytes survive the bounce, at a size that is not a chunk multiple. +// --------------------------------------------------------------------------- +TEST_CASE("the pinned ring reassembles the source byte for byte") { + constexpr size_t kChunk = 4096; + constexpr size_t kBytes = kChunk * 9 + 137; // deliberately ragged + FakeRing fake(kPinnedH2DBuffers, kChunk); + const std::vector src = Pattern(kBytes, 4242); + const size_t chunks = fake.Run(src, kChunk); + CHECK(chunks == 10); + REQUIRE(fake.dst.size() == kBytes); + CHECK(std::memcmp(fake.dst.data(), src.data(), kBytes) == 0); + CHECK(fake.chunk_sizes.back() == 137); + for (size_t n : fake.chunk_sizes) CHECK(n <= kChunk); +} + +// --------------------------------------------------------------------------- +// 4. The predicate: the five terms, one row per term. +// --------------------------------------------------------------------------- +TEST_CASE("the staged path needs all five terms, and refuses each missing one") { + StagedH2DInputs ok; + ok.src = PtrKind::kUnregisteredHost; + ok.dst = PtrKind::kDevice; + ok.bytes = kPinnedH2DChunkBytesDefault; + ok.chunk_bytes = kPinnedH2DChunkBytesDefault; + ok.stream_capturing = false; + ok.ring_available = true; + CHECK(ShouldStageH2D(ok)); + + SUBCASE("a copy smaller than one chunk stays on the single call") { + StagedH2DInputs in = ok; + in.bytes = kPinnedH2DChunkBytesDefault - 1; + CHECK_FALSE(ShouldStageH2D(in)); + // A 4 KiB norm weight is the case this term exists for. + in.bytes = 4096; + CHECK_FALSE(ShouldStageH2D(in)); + } + SUBCASE("a device source is never bounced (D2D)") { + StagedH2DInputs in = ok; + in.src = PtrKind::kDevice; + CHECK_FALSE(ShouldStageH2D(in)); + } + SUBCASE("an already-pinned source is never bounced") { + StagedH2DInputs in = ok; + in.src = PtrKind::kPinnedHost; + CHECK_FALSE(ShouldStageH2D(in)); + } + SUBCASE("a managed or unified source is never bounced") { + // VT_ROCM_MANAGED_ALLOC=1 is exactly this configuration, and it must be + // byte-identical to before: managed memory is already device addressable. + StagedH2DInputs in = ok; + in.src = PtrKind::kOther; + CHECK_FALSE(ShouldStageH2D(in)); + } + SUBCASE("a host destination is never bounced (D2H and H2H)") { + StagedH2DInputs in = ok; + in.dst = PtrKind::kPinnedHost; + CHECK_FALSE(ShouldStageH2D(in)); + in.dst = PtrKind::kUnregisteredHost; + CHECK_FALSE(ShouldStageH2D(in)); + in.dst = PtrKind::kOther; + CHECK_FALSE(ShouldStageH2D(in)); + } + SUBCASE("a capturing stream is never bounced") { + StagedH2DInputs in = ok; + in.stream_capturing = true; + CHECK_FALSE(ShouldStageH2D(in)); + } + SUBCASE("a ring that could not be allocated falls back instead of failing") { + StagedH2DInputs in = ok; + in.ring_available = false; + CHECK_FALSE(ShouldStageH2D(in)); + } + SUBCASE("chunk size zero disables the path") { + StagedH2DInputs in = ok; + in.chunk_bytes = 0; + CHECK_FALSE(ShouldStageH2D(in)); + } +} + +// --------------------------------------------------------------------------- +// 5. The knob parse. A typo must not silently disable a residency bound. +// --------------------------------------------------------------------------- +TEST_CASE("VT_ROCM_PINNED_H2D_MIB parses to bytes, and a typo takes the default") { + CHECK(ParsePinnedH2DChunkBytes("") == kPinnedH2DChunkBytesDefault); + CHECK(ParsePinnedH2DChunkBytes("64") == kPinnedH2DChunkBytesDefault); + CHECK(ParsePinnedH2DChunkBytes("0") == 0); + CHECK(ParsePinnedH2DChunkBytes("8") == (static_cast(8) << 20)); + CHECK(ParsePinnedH2DChunkBytes("128") == (static_cast(128) << 20)); + // Not a number: the DEFAULT, never 0. + CHECK(ParsePinnedH2DChunkBytes("off") == kPinnedH2DChunkBytesDefault); + CHECK(ParsePinnedH2DChunkBytes("-1") == kPinnedH2DChunkBytesDefault); + CHECK(ParsePinnedH2DChunkBytes("64MiB") == kPinnedH2DChunkBytesDefault); + // Absurdly large: the default, so nobody pins a terabyte by fat-fingering. + CHECK(ParsePinnedH2DChunkBytes("99999999") == kPinnedH2DChunkBytesDefault); +} + +// --------------------------------------------------------------------------- +// 6. The numbers are the ORACLE'S numbers, not invented ones. +// --------------------------------------------------------------------------- +TEST_CASE("the ring size and chunk size are llama.cpp's, and bound the residency") { + // llama-cpp 10bf611e5 (b10451), src/llama-model-loader.cpp:1440 and :1449. + CHECK(kPinnedH2DBuffers == 4); + CHECK(kPinnedH2DChunkBytesDefault == (static_cast(64) << 20)); + // The number the whole change exists to bound, for a model of ANY size. + CHECK(vt::rocm::kPinnedH2DRingBytes == (static_cast(256) << 20)); +} From 868db7d1b0340e70bbb6fb096928becf092f59d4 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sun, 13 Sep 2026 06:04:47 +0000 Subject: [PATCH 03/10] feat(MODEL-MM-QWEN4-EXP): stage a large pageable H2D through a bounded pinned ring on ROCm RocmBackend::Copy was one hipMemcpyAsync handed a multi-GiB pageable, file-backed, CIFS-backed source. To DMA out of a pageable range the ROCr runtime has to make all of it resident and describe it to the KFD, and that is where gfx1151 sits: 153 of 196 and 148 of 198 wchan samples in svm_range_set_attr across two 1200 s runs that never produced a token, with the two host-residency fixes already landed and the source pages already released. A copy of 64 MiB or more, from unregistered host storage to device memory, on a stream that is not capturing, now goes through four pinned 64 MiB buffers with an event each: wait on the slot, fill it, upload from it, record, advance modulo four. Host residency for the transfer is 256 MiB for a model of any size, which is llama.cpp's own bound at the recorded pin. Everything else takes the previous single call byte for byte -- every copy under one chunk, every readback, every D2D, every already-pinned or managed source, and anything inside a graph capture. VT_ROCM_PINNED_H2D_MIB=0 restores the old path in the same binary, so the load-time A/B is one build rather than two. The pointer classifier reads TWO shapes of "the runtime does not know this pointer", an error return and a success with null host and device pointers, because HIP has answered it both ways across versions and reading only one of them would leave the ring silently disengaged with identical bytes and an unchanged residency -- which is the failure this row has already shipped twice in the shape of a selector that matched nothing. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code] --- src/vt/rocm/rocm_backend.hip | 154 +++++++++++++++++++++++++++++++++++ 1 file changed, 154 insertions(+) diff --git a/src/vt/rocm/rocm_backend.hip b/src/vt/rocm/rocm_backend.hip index bea2c21e7..8b8704b82 100644 --- a/src/vt/rocm/rocm_backend.hip +++ b/src/vt/rocm/rocm_backend.hip @@ -38,7 +38,9 @@ #include #include +#include #include +#include #include #include #include @@ -50,6 +52,7 @@ #include "vt/graph_dedup.h" #include "vt/graph_dedup_runtime.h" #include "vt/rocm/rocm_arch.h" +#include "vt/rocm/rocm_pinned_h2d.h" namespace vt::rocm { namespace { @@ -62,6 +65,55 @@ void Check(hipError_t err, const char* what) { hipStream_t AsStream(const Queue& q) { return static_cast(q.handle); } +// --- the bounded pinned bounce ring's two runtime probes ------------------- +// .agents/specs/rocm-chunked-pinned-h2d.md. Everything with a DECISION in it +// lives in include/vt/rocm/rocm_pinned_h2d.h and is table-tested without HIP; +// these two read the runtime and answer in that header's vocabulary. + +// Which of the four kinds a pointer is. The UNREGISTERED answer is the one this +// change exists for -- a PROT_READ MAP_PRIVATE view of a GGUF shard, or any +// other ordinary host storage the HIP runtime knows nothing about. +// +// TWO SHAPES OF "the runtime does not know this pointer", because HIP has +// answered it both ways across versions and reading only one of them would make +// the ring silently never engage: an error return (historically +// hipErrorInvalidValue), and a success whose `type` is neither Host nor Device +// and whose recorded host and device pointers are both null (CUDA's +// cudaMemoryTypeUnregistered shape). Only `hipMemoryTypeHost` and +// `hipMemoryTypeDevice` are named, because those two enumerators exist in every +// HIP version this tree builds against; a managed or unified pointer falls out +// as kOther, which is what it must be -- it is already device addressable and +// must never be bounced. +// +// The last error is cleared immediately. hipPointerGetAttributes failing IS the +// answer we want here, and leaving it sticky would poison the next Check(). +vt::rocm::PtrKind ClassifyPointer(const void* p) { + hipPointerAttribute_t attr{}; + const hipError_t err = hipPointerGetAttributes(&attr, const_cast(p)); + if (err != hipSuccess) { + (void)hipGetLastError(); + return vt::rocm::PtrKind::kUnregisteredHost; + } + if (attr.type == hipMemoryTypeDevice) return vt::rocm::PtrKind::kDevice; + if (attr.type == hipMemoryTypeHost) return vt::rocm::PtrKind::kPinnedHost; + if (attr.devicePointer == nullptr && attr.hostPointer == nullptr) { + return vt::rocm::PtrKind::kUnregisteredHost; + } + return vt::rocm::PtrKind::kOther; +} + +// FAIL CLOSED: a probe that does not answer is treated as "capturing", so the +// staged path is refused rather than risking a hipEventSynchronize inside a +// capture region, which aborts the capture. +bool StreamIsCapturing(hipStream_t stream) { + hipStreamCaptureStatus status = hipStreamCaptureStatusNone; + if (hipStreamIsCapturing(stream, &status) != hipSuccess) { + (void)hipGetLastError(); + return true; + } + return status != hipStreamCaptureStatusNone; +} + // One probe gathers the residency attributes and the capability together. // Mirrors vt/cuda/cuda_device_caps.h in shape but is NOT cached (every consumer // is init-time or test-time, never per-step), and stays file-local: the public @@ -275,7 +327,18 @@ class RocmBackend final : public Backend { } // hipMemcpyDefault infers direction from the pointer values, so one entry // point covers h2d, d2h and d2d — and pageable host pointers on an APU. + // + // A LARGE H2D OUT OF PAGEABLE HOST STORAGE TAKES A BOUNDED PINNED RING + // INSTEAD (.agents/specs/rocm-chunked-pinned-h2d.md). Handing this one call a + // multi-GiB pageable, file-backed range makes the ROCr runtime pin and + // describe all of it to the KFD, which is where gfx1151 spends three quarters + // of its wchan samples in svm_range_set_attr while a 67.56 GiB checkpoint + // never produces a token. Everything else — every copy under one chunk, every + // readback, every D2D, every already-pinned or managed source, and anything + // inside a graph capture — takes the single call below, byte for byte. void Copy(Queue& q, void* dst, const void* src, size_t bytes) override { + if (StagedCopy(q, dst, src, bytes)) return; + vt::rocm::NotePinnedH2DDirect(); Check(hipMemcpyAsync(dst, src, bytes, hipMemcpyDefault, AsStream(q)), "hipMemcpyAsync"); } Queue CreateQueue() override { @@ -512,6 +575,88 @@ class RocmBackend final : public Backend { } private: + // The bounded pinned bounce ring. Oracle: llama-cpp 10bf611e5 (b10451), + // src/llama-model-loader.cpp:1440 (n_buffers = 4), :1449 (64 MiB), :1591-1642 + // (wait on the slot's event, fill it, upload from it, record, advance). + // + // Returns true when it handled the copy. The four cheap terms are checked + // before any HIP probe, deliberately: about 1,361 ResidentWeight calls per + // forward step reach Copy on this row's checkpoint, and a 4 KiB norm weight + // must not pay two hipPointerGetAttributes calls to learn that it is small. + // Those same two terms are then part of ShouldStageH2D, which stays the single + // authority on the decision and is the thing the truth table gates. + bool StagedCopy(Queue& q, void* dst, const void* src, size_t bytes) { + const size_t chunk = vt::rocm::PinnedH2DChunkBytes(); + if (chunk == 0 || bytes < chunk) return false; + hipStream_t stream = AsStream(q); + if (StreamIsCapturing(stream)) return false; + + std::lock_guard lock(h2d_mutex_); + vt::rocm::StagedH2DInputs in; + in.src = ClassifyPointer(src); + in.dst = ClassifyPointer(dst); + in.bytes = bytes; + in.chunk_bytes = chunk; + in.stream_capturing = false; // probed above, outside the lock + in.ring_available = EnsureRing(chunk); + if (!vt::rocm::ShouldStageH2D(in)) return false; + + size_t max_chunk = 0; + const size_t chunks = vt::rocm::RunStagedH2D( + h2d_ring_, bytes, chunk, + [&](size_t slot) { + Check(hipEventSynchronize(h2d_events_[slot]), "hipEventSynchronize"); + }, + [&](size_t slot, size_t off, size_t n) { + std::memcpy(h2d_slots_[slot], static_cast(src) + off, n); + }, + [&](size_t slot, size_t off, size_t n) { + Check(hipMemcpyAsync(static_cast(dst) + off, h2d_slots_[slot], n, + hipMemcpyHostToDevice, stream), + "hipMemcpyAsync"); + Check(hipEventRecord(h2d_events_[slot], stream), "hipEventRecord"); + if (n > max_chunk) max_chunk = n; + }); + if (chunks == 0) return false; + vt::rocm::NotePinnedH2DStaged(chunks, max_chunk); + return true; + } + + // Allocated on first use and NEVER FREED, deliberately. RocmBackend instances + // live in a function-local static vector in the registrar below, so a + // destructor would run during static destruction, and calling hipHostFree or + // hipEventDestroy after the HIP runtime has begun tearing down is a hazard + // this file does not have today (exec_, the hipGraphExec_t, is likewise never + // destroyed). 256 MiB returned to the OS at exit() buys nothing and a + // teardown-order crash costs a measurement. + // + // A FAILED ALLOCATION FALLS BACK rather than refusing the model: the ring is a + // residency bound, not a correctness requirement, and no checkpoint should + // fail to load because 256 MiB of pinned memory was unavailable. Called with + // h2d_mutex_ held. + bool EnsureRing(size_t chunk) { + if (h2d_state_ != 0) return h2d_state_ == 1; + h2d_state_ = 2; // failed, unless this function gets all the way through + if (hipSetDevice(device_) != hipSuccess) { + (void)hipGetLastError(); + return false; + } + for (size_t i = 0; i < vt::rocm::kPinnedH2DBuffers; ++i) { + if (hipHostMalloc(&h2d_slots_[i], chunk, hipHostMallocDefault) != hipSuccess) { + (void)hipGetLastError(); + return false; + } + if (hipEventCreateWithFlags(&h2d_events_[i], hipEventDisableTiming) != hipSuccess) { + (void)hipGetLastError(); + return false; + } + } + h2d_ring_.n_buffers = vt::rocm::kPinnedH2DBuffers; + h2d_state_ = 1; + vt::rocm::NotePinnedH2DRingBytes(vt::rocm::kPinnedH2DBuffers * chunk); + return true; + } + int device_ = 0; bool unified_memory_ = false; bool managed_alloc_ = false; @@ -524,6 +669,15 @@ class RocmBackend final : public Backend { // ENG-CUDAGRAPH-DEDUP (#1162): built on the first capture, and only when // VT_CUDA_GRAPH_DEDUP asked for it, so an unset environment allocates nothing. std::unique_ptr dedup_; + // 0 = not tried, 1 = ready, 2 = unavailable (tried once and failed). + int h2d_state_ = 0; + // Serialises the SHARED ring, not the copy: two threads staging at once would + // otherwise refill a slot another thread's DMA is still reading. A mutex on a + // path that only >= 64 MiB copies reach costs nothing measurable. + std::mutex h2d_mutex_; + vt::rocm::StagedH2DRing h2d_ring_; + void* h2d_slots_[vt::rocm::kPinnedH2DBuffers] = {}; + hipEvent_t h2d_events_[vt::rocm::kPinnedH2DBuffers] = {}; }; // Registers every visible AMD GPU at its own Device{kROCM, i} slot, mirroring From 556475c10af7e4829164c839cbc2b60560c0771f Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sun, 13 Sep 2026 06:48:34 +0000 Subject: [PATCH 04/10] record(MODEL-MM-QWEN4-EXP): the ring produces the row's first token on gfx1151, and the same binary without it does not MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The gate counts, the red, the mutations and the model gate are all measured now, so the spec stops saying RECORDED AT §7 and says what was read. The model gate is met. On strix:gpu0 the 67.56 GiB Qwen3.8-Flash-Next UD-IQ1_S produced three generations of 32 tokens, exit 0, no deadline kill, and the compute thread spent none of its 122 wchan samples in svm_range_set_attr. The same binary minutes later with VT_ROCM_PINNED_H2D_MIB=0 reproduced the recorded wedge to within noise -- 135 of 197 samples in svm_range_set_attr, device memory stuck at the same 29.69 GiB plateau, killed at the 1200 s deadline with no token. One binary, one boot, the knob read at runtime, so the ring is the cause rather than a coincidence of the day. Three things this deliberately does not claim. There is no throughput comparison, because no oracle ran this workload. There is no load-time verdict, because the two arms read the checkpoint at different page-cache temperatures. And the warm decode pair is TWO samples, not three, because --repeat 3 gives one cold generation that carries the whole 72 GiB device staging; the number is recorded as thin instead of being presented as a spread it is not. VT_ROCM_PINNED_H2D_MIB is a new config key, so ENVIRONMENT.md and USAGE.md carry it, including the fact that turning it off is what makes this artifact stop producing tokens. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code] --- .../ISSUE-LOCAL-01M2BZ5QK4XRETK48CXKSHKRDW.md | 8 +- .agents/specs/rocm-chunked-pinned-h2d.md | 159 ++++++++++++++++-- docs/ENVIRONMENT.md | 1 + docs/USAGE.md | 11 ++ 4 files changed, 164 insertions(+), 15 deletions(-) diff --git a/.agents/issues/MODEL-MM-QWEN4-EXP/ISSUE-LOCAL-01M2BZ5QK4XRETK48CXKSHKRDW.md b/.agents/issues/MODEL-MM-QWEN4-EXP/ISSUE-LOCAL-01M2BZ5QK4XRETK48CXKSHKRDW.md index 8893e41c1..0f76e4829 100644 --- a/.agents/issues/MODEL-MM-QWEN4-EXP/ISSUE-LOCAL-01M2BZ5QK4XRETK48CXKSHKRDW.md +++ b/.agents/issues/MODEL-MM-QWEN4-EXP/ISSUE-LOCAL-01M2BZ5QK4XRETK48CXKSHKRDW.md @@ -1,14 +1,14 @@ ID: ISSUE-LOCAL-01M2BZ5QK4XRETK48CXKSHKRDW Title: OWED: chunked H2D through a pinned bounce buffer, llama.cpp's 4 x 64 MiB shape Row: MODEL-MM-QWEN4-EXP -State: OPEN +State: CLOSED Kind: task GitHub: - Mirror: PENDING Availability: FULL Created: 2026-09-12 -Updated: 2026-09-12 -Closed: - +Updated: 2026-09-13 +Closed: 2026-09-13 ## Problem @@ -16,4 +16,4 @@ ResidentWeight stages every weight with one hipMemcpyAsync straight out of a PAG ## Resolution -- +Built and measured 2026-09-13. RocmBackend::Copy now stages a host-to-device copy of 64 MiB or more, out of unregistered host storage, into device memory, on a non-capturing stream, through four pinned 64 MiB buffers with an event each -- llama.cpp's shape at 10bf611e5, src/llama-model-loader.cpp:1440 and :1449. Spec .agents/specs/rocm-chunked-pinned-h2d.md; decision and loop in the HIP-free include/vt/rocm/rocm_pinned_h2d.h, table-tested in tests/vt/test_rocm_pinned_h2d.cpp (9 cases / 84 assertions) and reached through Backend::Copy in tests/vt/test_backend_cross_device.cpp (61 cases / 84841 assertions on strix:gpu0, rc job e8bf3b66). It is the change that makes the row's model gate pass: on strix:gpu0 the 67.56 GiB Qwen3.8-Flash-Next UD-IQ1_S produced three generations of 32 tokens with the ring on, and the SAME BINARY with VT_ROCM_PINNED_H2D_MIB=0 wedged in svm_range_set_attr for 135 of 197 wchan samples and was killed at 1200 s (rc job 672093bc). diff --git a/.agents/specs/rocm-chunked-pinned-h2d.md b/.agents/specs/rocm-chunked-pinned-h2d.md index 9036acc5b..c8a0b92c7 100644 --- a/.agents/specs/rocm-chunked-pinned-h2d.md +++ b/.agents/specs/rocm-chunked-pinned-h2d.md @@ -258,15 +258,64 @@ row has already shipped two selectors that matched nothing and reported `Status: SUCCESS!`; a selector that matches nothing is indistinguishable from one that passes. -| binary | selector | cases | assertions | +Measured on `strix:gpu0` (gfx1151, ROCm 7.2.4) under `rc` job +`e8bf3b66-eb80-4659-8e6e-48167eb6bf60`, from a clean clone built in +`/tmp/vllmcpp-chunked-h2d` with +`cmake -G Ninja -DCMAKE_BUILD_TYPE=Release -DVLLM_CPP_HIP=ON` +`-DVLLM_CPP_HIP_ARCHITECTURES=gfx1151 -DROCM_PATH=/opt/rocm` +`-DVLLM_CPP_BUILD_TESTS=ON`, `ninja -j 4`, run with +`LD_LIBRARY_PATH=/opt/rocm-7.2.4/lib` — without which the binary exits 127 +having measured nothing. + +| binary | selector | cases | assertions | result | +|---|---|---|---|---| +| `test_rocm_pinned_h2d` | (none — whole binary) | 9 | 84 | SUCCESS | +| `test_backend_cross_device` | `-tc=*pinned bounce*` | 1 | 8 | SUCCESS | +| `test_backend_cross_device` | (none — whole binary) | 61 | 84841 | SUCCESS | +| `test_backend_cross_device` | `-tc=*DSA*` | 2 | 273 | SUCCESS | + +Baseline at `98e2cd7da`: `test_backend_cross_device` whole binary **60 cases / +84833 assertions**, `-tc=*DSA*` **2 / 273**. The whole binary therefore grew by +exactly one case and eight assertions, which is this spec's case and nothing +else, and `-tc=*DSA*` is unmoved. + +**RED, at the test commit `e1bf7fd1d`, same binary, same board.** +`-tc=*pinned bounce*` reported **1 case, 0 passed, 1 failed / 8 assertions, 3 +passed, 5 failed**, `Status: FAILURE!`, exit 1, and printed +`staged=0 direct=0 chunks=0 max_chunk=0 ring_bytes=0`. At the implementation +commit `81f91b600` the same selector on the same board printed +`staged=1 direct=2 chunks=3 max_chunk=67108864 ring_bytes=268435456` and +`Status: SUCCESS!`. Both runs printed `pinned H2D case ran on a ROCm board: 1`, +so neither is the shape where a selector matched nothing. + +**MUTATIONS, on the board, each with the binary proven changed by md5.** The +clean implementation build is `88efb430a22c65c4` (`test_backend_cross_device`) +and `fc26b58276f3e6f9` (`test_rocm_pinned_h2d`). + +| mutation | binaries | device selector | unit binary | |---|---|---|---| -| `test_rocm_pinned_h2d` | (none — whole binary) | RECORDED AT §7 | RECORDED AT §7 | -| `test_backend_cross_device` | `-tc=*pinned bounce*` | RECORDED AT §7 | RECORDED AT §7 | -| `test_backend_cross_device` | (none — whole binary) | RECORDED AT §7 | RECORDED AT §7 | -| `test_backend_cross_device` | `-tc=*DSA*` | RECORDED AT §7 | RECORDED AT §7 | - -Baseline to hold at `98e2cd7da`: `test_backend_cross_device` whole binary -**60 cases / 84833 assertions**, `-tc=*DSA*` **2 / 273**. +| delete `if (StagedCopy(...)) return;` from `RocmBackend::Copy` | `006c453cce2b25fb` / `8f82b9f94db909aa` | FAILURE, 4 of 8 assertions failed, `staged=0 direct=3 chunks=0 ring_bytes=0` | n/a | +| widen the chunk to the whole buffer | `74949cfa593d0843` / `f3720497fdbe96ca` | FAILURE, SIGSEGV (exit 139) before any assertion | FAILURE, 10 of 19 assertions failed, then SIGSEGV | + +The second mutation convicts by CRASH rather than by assertion, and that is +reported as what it is rather than dressed up: with the chunk widened, both the +fake ring's slot buffers and the real pinned slots are one chunk long and the +copy reads past them, so the process dies. It is still caused by the mutation and +still red, and the arm that matters — the first one — fails cleanly on the +instrument. + +**The tree was restored byte for byte, and the proof is the binary, not +`git status`.** The rebuild after both mutations produced md5 +`88efb430a22c65c4` and `fc26b58276f3e6f9` — identical to the clean +implementation build — and re-ran 9/84 and 1/8 green. `git status` reported one +dirty path after the second mutation; that is the `core` file the SIGSEGV +dumped into the checkout, not a source edit, and the identical binary hashes are +what settle it. + +Eight further mutations of the pure header were run off-board before the branch +was pushed, each convicting and each with a distinct binary md5: the chunk size +widened, the slot wait deleted, each of the four predicate terms deleted, the +cross-call in-flight state reset, and the knob parse made to answer 0 on a typo. ```sh python3 scripts/check-agent-record.py @@ -315,9 +364,97 @@ The artifact's compiled feature set is asserted before it is timed: `ldd` for we want, and the last error is cleared with `hipGetLastError()` immediately so it cannot poison the next `Check`. -## 7. Outcome - -RECORDED ON LANDING. +## 7. Outcome — THE MODEL GATE IS MET, AND THE RING IS WHAT MEETS IT + +Measured 2026-09-13 on `strix:gpu0` (gfx1151, Radeon 8060S, ROCm 7.2.4, 30 GiB +host, 32 CPUs), box exclusively leased, under `rc` job +`672093bc-932b-4e54-b319-e15529f70256`. Built on the worker in +`/tmp/vllmcpp-chunked-h2d` from a clean clone at `81f91b6005c7e8` with +`cmake -G Ninja -DCMAKE_BUILD_TYPE=Release -DVLLM_CPP_HIP=ON` +`-DVLLM_CPP_HIP_ARCHITECTURES=gfx1151 -DROCM_PATH=/opt/rocm -DVLLM_CPP_SERVER=ON`, +`ninja -j 4 vllm-cli vllm-server`, run with +`LD_LIBRARY_PATH=/opt/rocm-7.2.4/lib`. `VT_ROCM_MANAGED_ALLOC` unset. + +**The artifact was asserted before it was timed.** `ldd` on `vllm-cli` resolves +`libamdhip64.so.7`, `libhsa-runtime64.so.1`, `libhipblaslt.so.1`, +`libhipblas.so.3` and `librocblas.so.5`, all from `/opt/rocm-7.2.4/lib`, and +CMake reported `ROCm backend: ENABLED for arch(es) [gfx1151]`. `vllm-cli` md5 +`364f11fddc6392feb79cdc23f8d69cf5`, `vllm-server` md5 +`aab04a1549ad3a5a543b1a6012834054`. + +Workload: `examples/vllm-cli --device auto --max-tokens 32 --temperature 0 +--max-num-seqs 1 --repeat 3 --prompt "Say hello"` over +`/workspace/ckpt/qwen4exp-flash-next-iq1s/Qwen3.8-Flash-Next-UD-IQ1_S-00001-of-00003.gguf` +(shard 1 of 3; 10,946,624 + 49,990,818,368 + 22,544,696,352 bytes). + +### The A/B, ONE BINARY, ONE BOOT, THE KNOB READ AT RUNTIME + +| | ARM A — ring ON | ARM B — `VT_ROCM_PINNED_H2D_MIB=0` | +|---|---|---| +| token | **YES — 3 x 32, `finish_reason=length`** | **NONE**, killed at the 1200 s deadline (exit 137) | +| wall | 771 s, exited 0 | 1218 s, `killed_at_deadline=1` | +| `[vt load] weights` | 61.187 s | 35.909 s (page cache warm from arm A) | +| peak `VmHWM` | 20,714,504 kB (19.75 GiB) | 27,076,580 kB (25.82 GiB) | +| peak `RssFile` | 14,516,792 kB (13.84 GiB) | 21,093,296 kB (20.12 GiB) | +| peak `RssAnon` | 6,190,884 kB (5.90 GiB) | 5,807,020 kB (5.54 GiB) | +| peak device `mem_info_vram_used` | **77,271,658,496 B** | 31,873,912,832 B | +| `wchan` over D-state threads | 115 `folio_wait_bit_common`, 1 `wait_for_response`, **0 `svm_range_set_attr`**, 122 samples | **135 `svm_range_set_attr`**, 39 `folio_wait_bit_common`, 20 `do_mprotect_pkey`, 8 `exit_mm`, 1 `wait_for_response`, 197 samples | + +Arm B is `rocm-host-residency-after-upload.md` §6a reproduced to within noise — +27.08 vs 27.25/27.32 GB `VmHWM`, 21.09 vs 21.21/21.34 GB `RssFile`, +31,873,912,832 vs 31,878,860,800/31,880,183,808 B of device memory, and the same +`svm_range_set_attr` majority. That is what makes arm A attributable: the two +arms are the SAME BINARY minutes apart on the same board, and the only thing that +differs is one environment variable this change reads. The ring is the cause. + +**The device memory is the tell.** Arm B stops at 29.69 GiB and stays there, +exactly as §6a recorded. Arm A reaches 71.96 GiB, which is the whole checkpoint. +So the wedge was never the model failing to fit; it was the upload never +finishing. (The 71.96 GiB figure exceeds the 33.27 GB `hipMemGetInfo` total §6a +quotes; the 96 GiB VRAM carve on this part is the obvious explanation and this +spec does NOT resolve the discrepancy, it reports both raw numbers.) + +### The generations + +| run | tokens | seconds | tok/s | +|---|---|---|---| +| 1 | 32 | 695.871 | 0.046 | +| 2 | 32 | 6.048 | 5.291 | +| 3 | 32 | 6.069 | 5.273 | + +**Run 1 is not a decode number and must not be quoted as one.** Weight staging +is lazy — `dense_attn::ResidentWeight` uploads on first use, behind the `d_dev` +memo — so run 1 carries the one-time 72 GiB host-to-device transfer of the whole +checkpoint. The steady-state pair is 5.291 and 5.273 tok/s, a spread of 0.34% +over two samples. TWO warm samples is not three, and this is recorded as thin +rather than dressed up: `--repeat 3` gives three generations of which exactly one +is cold. gfx1151 fails about two runs in five with an illegal GPU memory access +(`ISSUE-LOCAL-01M2BY2M2ATNVR3XQKV2DB1BJD`); neither arm here hit that signature, +and arm A's three generations all completed with `finish_reason=length`. + +No TTFT is recorded: `vllm-cli` in blocking mode reports whole-generation +seconds, not first-token latency, and no number is invented from it. + +### What this does NOT claim + +- No throughput COMPARISON. There is no llama.cpp or vLLM denominator here, and + 5.28 tok/s is this engine's first number on this part, not a ratio. +- No load-time verdict for the ring. Arm A's 61.187 s and arm B's 35.909 s are + not comparable: arm A read the checkpoint cold off CIFS and arm B read it with + the page cache already warm. A load-time A/B needs interleaved repeats from a + dropped cache and was not taken. +- Nothing about any other family. §3b names them; only Qwen4-Exp was run. + +## 7a. Where the remaining time goes, so nobody reads 0.046 as the answer + +Arm A spends 695.871 s inside its first generation and 6.05 s inside each +subsequent one. That 690 s difference is the lazy device staging of a 72 GiB +checkpoint, which is about 104 MiB/s through the ring — slow, and the next +question for this row rather than this spec's. The `wchan` histogram says where +it is: 115 of 122 samples in `folio_wait_bit_common`, which is page-cache read +wait on the CIFS mount (`//192.168.68.102/Data`), not device work. §6a named that +mount as an unseparated confound; arm C of this job stages the checkpoint to +local disk to separate it. ## 8. Stop conditions diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index 96d2f7231..2f949a298 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -353,6 +353,7 @@ Setting it does nothing, and the row is gone rather than caveated | `VT_GEMMA4_HOST_EXPERT_MB` | `512` | Host-side expert staging budget (MiB) for non-resident paths | | `VT_GEMMA4_LAYER_TRACE` | off | `=1` layer GPU-synced phase timers; `=2` per-layer heartbeats | | `VLLM_CPP_HTTP_FIXED_POOL` | `1` (fixed) | `=0` reverts the HTTP worker pool to the legacy dynamic mode. Production uses the capacity-derived fixed pool; the opt-out exists for same-binary A/B attribution | +| `VT_ROCM_PINNED_H2D_MIB` | `64` | Chunk size, in MiB, of the ROCm bounded pinned bounce ring. A host-to-device copy of at least this many bytes, from host storage the HIP runtime does not know, to device memory, on a stream that is not capturing, is staged through four pinned buffers of this size with an event each instead of one `hipMemcpyAsync` off the pageable source — llama.cpp's shape at `10bf611e5`, `src/llama-model-loader.cpp:1440` and `:1449`. Host residency for the transfer is then `4 x` this value for a model of any size. `=0` disables the ring and restores the single call, which is the same-binary A/B. Anything that does not parse takes the default rather than `0`, because a disabled ring and a working one differ only in host residency. On `strix:gpu0` (gfx1151) the default is what lets the 67.56 GiB Qwen3.8-Flash-Next UD-IQ1_S produce a token at all: with `=0` the same binary wedges in `svm_range_set_attr` and is killed at a 1200 s deadline (`.agents/specs/rocm-chunked-pinned-h2d.md` §7). Every other backend ignores it | | `VT_ROCM_MANAGED_ALLOC` | unset | Which allocator the ROCm backend uses. Unset: `hipMallocManaged` only on an integrated, managed-capable device that ALSO reports `PageableMemoryAccess = 1`, i.e. one that can take a recoverable page fault; everything else gets plain `hipMalloc`. `=0` never takes the managed branch. `=1` takes it wherever the managed attributes allow, ignoring `PageableMemoryAccess` — the pre-[#2511](https://github.com/mudler/vllm.cpp/issues/2511) behaviour, restored whole (allocator and unified claim), for a single-binary A/B. On an XNACK-less APU (gfx1151, gfx1103) the managed branch measured 17 GPU faults in 21 legs against 0 in 21 without it, so `=1` re-arms that; a discrete card ignores all three values. Note `=1` and unset also decide `UnifiedMemory()` there, and with it whether the CPU reference tier can serve an op ROCm has no kernel for | | `VT_ROCM_ATTN_CPU_REF` | unset | `=1` routes ROCm paged attention through the CPU reference kernel instead of the HIP kernel — a correctness A/B for the ROCm attention bring-up | | `VT_ATTN_DECODE_GQA4` | off | `=1` routes f32-query decode through `PagedAttnDecodeGqaF32Q` (QG=4 fused q-heads per KV group, warp-strided walk) for bf16 or fp8-e4m3 KV at d=128/256, hq=16, kv=4; the default falls to the `PagedAttnOnline` reference. The arm's reduction order differs from the reference's, so greedy anchors can move at exact ties — opt-in until near-tie adjudication lands | diff --git a/docs/USAGE.md b/docs/USAGE.md index 512d8ec85..ff24aaa5e 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -892,6 +892,17 @@ skips with that refusal quoted. - On ROCm, decode-shaped GEMMs (batch of 4 or fewer, bf16) run on a split-K skinny-GEMM kernel rather than the tiled BLAS path. Set `VT_ROCM_SKINNY=0` to restore the BLAS path when you want to compare the two. +- On ROCm, a host-to-device copy of 64 MiB or more out of ordinary host storage + is staged through four pinned 64 MiB buffers rather than handed to the driver + as one pageable transfer, so host residency for the copy is 256 MiB whatever + the model's size. Every smaller copy, every readback, every device-to-device + copy, an already-pinned or managed source, and anything inside a graph capture + take the previous single call unchanged. Set `VT_ROCM_PINNED_H2D_MIB=0` to + restore that single call for every copy, or to another number of MiB to change + the chunk. On an integrated part with no pageable-memory access this is what + makes a large GGUF checkpoint load at all: with the ring off, a 67.56 GiB + artifact on gfx1151 stops at 29.69 GiB of device memory and never produces a + token. See [Environment variables](ENVIRONMENT.md). - On ROCm, Gemma-4 FP8 mixture-of-experts decode uses the device-indexed expert gate for batches up to 63 tokens; wider batches use the prefill-batch path. Set `VT_GEMMA4_DECODE_INDEXED_MAX_T=1` to restore the From d4a0388c66ade17c03ae265cbd636e27b477c5e0 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sun, 13 Sep 2026 06:58:05 +0000 Subject: [PATCH 05/10] record(MODEL-MM-QWEN4-EXP): separate the CIFS confound, and give the decode number four samples MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The third arm ran the same binary and the same ring over a local copy of the same three shards, and it answers the question §6a of the residency spec left open. The first generation falls from 695.871 s to 38.438 s and the D-state histogram collapses from 115 page-cache waits to nothing, so 657 of those 696 seconds were reading the file over CIFS and the ring moves the 72 GiB checkpoint onto the board in about 38 s. The confound was never the cause of the wedge -- the off arm wedges on the same mount the on arm succeeds on -- but it was almost all of what was left. It also fixes the thin part of the previous record. The steady-state decode number is now four samples across two independent process launches on two different source filesystems, 5.291 / 5.273 / 5.274 / 5.275 tok/s, a 0.34% spread. Three repetitions inside one handle would have satisfied the row's rule more weakly than this does. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code] --- .agents/specs/rocm-chunked-pinned-h2d.md | 93 +++++++++++++++--------- 1 file changed, 59 insertions(+), 34 deletions(-) diff --git a/.agents/specs/rocm-chunked-pinned-h2d.md b/.agents/specs/rocm-chunked-pinned-h2d.md index c8a0b92c7..ef56b6171 100644 --- a/.agents/specs/rocm-chunked-pinned-h2d.md +++ b/.agents/specs/rocm-chunked-pinned-h2d.md @@ -389,16 +389,22 @@ Workload: `examples/vllm-cli --device auto --max-tokens 32 --temperature 0 ### The A/B, ONE BINARY, ONE BOOT, THE KNOB READ AT RUNTIME -| | ARM A — ring ON | ARM B — `VT_ROCM_PINNED_H2D_MIB=0` | -|---|---|---| -| token | **YES — 3 x 32, `finish_reason=length`** | **NONE**, killed at the 1200 s deadline (exit 137) | -| wall | 771 s, exited 0 | 1218 s, `killed_at_deadline=1` | -| `[vt load] weights` | 61.187 s | 35.909 s (page cache warm from arm A) | -| peak `VmHWM` | 20,714,504 kB (19.75 GiB) | 27,076,580 kB (25.82 GiB) | -| peak `RssFile` | 14,516,792 kB (13.84 GiB) | 21,093,296 kB (20.12 GiB) | -| peak `RssAnon` | 6,190,884 kB (5.90 GiB) | 5,807,020 kB (5.54 GiB) | -| peak device `mem_info_vram_used` | **77,271,658,496 B** | 31,873,912,832 B | -| `wchan` over D-state threads | 115 `folio_wait_bit_common`, 1 `wait_for_response`, **0 `svm_range_set_attr`**, 122 samples | **135 `svm_range_set_attr`**, 39 `folio_wait_bit_common`, 20 `do_mprotect_pkey`, 8 `exit_mm`, 1 `wait_for_response`, 197 samples | +| | ARM A — ring ON, CIFS | ARM B — `VT_ROCM_PINNED_H2D_MIB=0`, CIFS | ARM C — ring ON, LOCAL copy | +|---|---|---|---| +| token | **YES — 3 x 32, `finish_reason=length`** | **NONE**, killed at the 1200 s deadline (exit 137) | **YES — 3 x 32** | +| wall | 771 s, exited 0 | 1218 s, `killed_at_deadline=1` | **62 s**, exited 0 | +| `[vt load] mmap+header` | 0.568 s | 0.201 s | 0.023 s | +| `[vt load] weights` | 61.187 s | 35.909 s (page cache warm from arm A) | **8.770 s** | +| peak `VmHWM` | 20,714,504 kB (19.75 GiB) | 27,076,580 kB (25.82 GiB) | 26,920,604 kB (25.67 GiB) | +| peak `RssFile` | 14,516,792 kB (13.84 GiB) | 21,093,296 kB (20.12 GiB) | 20,728,476 kB (19.77 GiB) | +| peak `RssAnon` | 6,190,884 kB (5.90 GiB) | 5,807,020 kB (5.54 GiB) | 6,190,900 kB (5.90 GiB) | +| peak device `mem_info_vram_used` | **77,271,658,496 B** | 31,873,912,832 B | **77,271,609,344 B** | +| `wchan` over D-state threads | 115 `folio_wait_bit_common`, 1 `wait_for_response`, **0 `svm_range_set_attr`**, 122 samples | **135 `svm_range_set_attr`**, 39 `folio_wait_bit_common`, 20 `do_mprotect_pkey`, 8 `exit_mm`, 1 `wait_for_response`, 197 samples | 1 `wait_for_response`, 1 unresolved, **0 `svm_range_set_attr`**, 11 samples | + +Arm C is arm A with the three shards copied to the worker's local disk first +(`/tmp/ckpt-iq1s`, same bytes, same sizes), so it answers the confound §6a left +open: the checkpoint lives on a CIFS mount and page-cache read wait is not device +work. `VT_ROCM_PINNED_H2D_MIB` is unset in both. Arm B is `rocm-host-residency-after-upload.md` §6a reproduced to within noise — 27.08 vs 27.25/27.32 GB `VmHWM`, 21.09 vs 21.21/21.34 GB `RssFile`, @@ -416,24 +422,34 @@ spec does NOT resolve the discrepancy, it reports both raw numbers.) ### The generations -| run | tokens | seconds | tok/s | -|---|---|---|---| -| 1 | 32 | 695.871 | 0.046 | -| 2 | 32 | 6.048 | 5.291 | -| 3 | 32 | 6.069 | 5.273 | - -**Run 1 is not a decode number and must not be quoted as one.** Weight staging -is lazy — `dense_attn::ResidentWeight` uploads on first use, behind the `d_dev` -memo — so run 1 carries the one-time 72 GiB host-to-device transfer of the whole -checkpoint. The steady-state pair is 5.291 and 5.273 tok/s, a spread of 0.34% -over two samples. TWO warm samples is not three, and this is recorded as thin -rather than dressed up: `--repeat 3` gives three generations of which exactly one -is cold. gfx1151 fails about two runs in five with an illegal GPU memory access -(`ISSUE-LOCAL-01M2BY2M2ATNVR3XQKV2DB1BJD`); neither arm here hit that signature, -and arm A's three generations all completed with `finish_reason=length`. +| arm | run | tokens | seconds | tok/s | +|---|---|---|---|---| +| A (CIFS) | 1 | 32 | 695.871 | 0.046 | +| A | 2 | 32 | 6.048 | 5.291 | +| A | 3 | 32 | 6.069 | 5.273 | +| C (local) | 1 | 32 | 38.438 | 0.833 | +| C | 2 | 32 | 6.067 | 5.274 | +| C | 3 | 32 | 6.067 | 5.275 | + +**Run 1 of each arm is not a decode number and must not be quoted as one.** +Weight staging is lazy — `dense_attn::ResidentWeight` uploads on first use, +behind the `d_dev` memo — so run 1 carries the one-time 72 GiB host-to-device +transfer of the whole checkpoint. + +**The steady-state number is 5.27-5.29 tok/s, and it is FOUR samples across TWO +independent process launches on two different source filesystems:** 5.291, +5.273, 5.274, 5.275. Spread max-to-min 0.34%. That satisfies the row's +three-repetition rule and it does so across process boundaries rather than three +times inside one handle, which is the stronger shape. gfx1151 fails about two +runs in five with an illegal GPU memory access +(`ISSUE-LOCAL-01M2BY2M2ATNVR3XQKV2DB1BJD`); no arm here hit that signature, and +all six generations completed with `finish_reason=length`. No TTFT is recorded: `vllm-cli` in blocking mode reports whole-generation -seconds, not first-token latency, and no number is invented from it. +seconds, not first-token latency, and no number is invented from it. What arm C +does give is a cold-start-to-first-answer figure on local storage: +0.023 s header + 8.770 s weights + 38.438 s first generation, 62 s of wall for +the whole process. ### What this does NOT claim @@ -445,16 +461,25 @@ seconds, not first-token latency, and no number is invented from it. dropped cache and was not taken. - Nothing about any other family. §3b names them; only Qwen4-Exp was run. -## 7a. Where the remaining time goes, so nobody reads 0.046 as the answer +## 7a. Where the first generation's time goes — MEASURED, not hypothesised Arm A spends 695.871 s inside its first generation and 6.05 s inside each -subsequent one. That 690 s difference is the lazy device staging of a 72 GiB -checkpoint, which is about 104 MiB/s through the ring — slow, and the next -question for this row rather than this spec's. The `wchan` histogram says where -it is: 115 of 122 samples in `folio_wait_bit_common`, which is page-cache read -wait on the CIFS mount (`//192.168.68.102/Data`), not device work. §6a named that -mount as an unseparated confound; arm C of this job stages the checkpoint to -local disk to separate it. +subsequent one, and the obvious misreading is that the ring stages at about +104 MiB/s. It does not. Arm A's `wchan` histogram is 115 of 122 samples in +`folio_wait_bit_common`, which is page-cache read wait on the CIFS mount +(`//192.168.68.102/Data`), not device work. + +**Arm C settles it.** The same binary, the same checkpoint, the same ring, with +the three shards copied to local disk first: the first generation takes +**38.438 s instead of 695.871**, an 18.1x reduction, and the D-state histogram +collapses to 11 samples with nothing in it. Weight load falls the same way, +8.770 s against 61.187. So 657 of arm A's 696 s were reading the file over CIFS, +and the ring itself moves the 72 GiB checkpoint onto the board in about 38 s, +which is roughly 1.9 GiB/s. + +**§6a's CIFS confound is therefore separated and closed for this row.** It was +never the cause of the wedge — arm B wedges off the same mount arm A succeeds +on — and it accounts for essentially all of the remaining first-generation cost. ## 8. Stop conditions From d83c4b1c6f6818891c4a415057a08b5e35a33627 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sun, 13 Sep 2026 07:28:07 +0000 Subject: [PATCH 06/10] test(MODEL-MM-QWEN4-EXP): read the pinned ring's own residency number on the managed arm The cross-device pinned-bounce case already printed `ring_bytes` on both arms and asserted it on only one. On the managed arm -- gfx1103, or any board run with `VT_ROCM_MANAGED_ALLOC=1` -- the destination is already device addressable, every copy is direct, and the case was green while the instrument printed `ring_bytes=268435456`. A ring that is allocated and never used changes no byte, so nothing else in this file can see it. It is not free on the board that shows it. `.agents/environment.md:95-100` measures gfx1151's managed ceiling as bounded by HOST RAM rather than by the 96 GiB carve: 27 GiB reached against 29.3 GiB available. Pinned host memory comes out of exactly that bound, so 256 MiB of it is charged to the resource that decides whether the model fits, on the arm that never stages. This case fails until `RocmBackend::StagedCopy` stops allocating the ring before it knows the copy will use it. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code] --- tests/vt/test_backend_cross_device.cpp | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/tests/vt/test_backend_cross_device.cpp b/tests/vt/test_backend_cross_device.cpp index d0d152290..370b85efb 100644 --- a/tests/vt/test_backend_cross_device.cpp +++ b/tests/vt/test_backend_cross_device.cpp @@ -317,6 +317,16 @@ TEST_CASE("a large pageable H2D takes the pinned bounce ring, and a small one do // deliberately not engaged and every copy is direct. CHECK(after_small.staged_copies == before.staged_copies); CHECK(after_small.direct_copies > before.direct_copies); + // AND NOT ONE PINNED BYTE IS ALLOCATED. A ring that is built and then + // never used is invisible to every other assertion here -- the bytes are + // right, the copies are direct, the case is green -- and it costs 256 MiB + // of PINNED HOST memory on exactly the boards that never stage. That is + // not free on this part: .agents/environment.md:95-100 measures gfx1151's + // managed ceiling as bounded by HOST RAM (27 GiB reached against 29.3 GiB + // available), so pinned host is the binding resource. The instrument + // printed ring_bytes=268435456 on this arm for a whole review cycle and + // nothing read it; this is the line that reads it. + CHECK(after_small.ring_bytes == before.ring_bytes); } else { // The staging arm. This is the assertion the whole change is for. CHECK(after_big.staged_copies == before.staged_copies + 1); From a6621c93739745b456656b9c15e0d15e6548d86f Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sun, 13 Sep 2026 07:30:44 +0000 Subject: [PATCH 07/10] fix(MODEL-MM-QWEN4-EXP): allocate the pinned ring only for a copy that will use it `StagedCopy` evaluated `in.ring_available = EnsureRing(chunk)` before it asked `ShouldStageH2D`, so the first copy of 64 MiB or more on a non-capturing stream built the ring whatever the pointers turned out to be. Any ROCm process that performs one such copy permanently allocated 256 MiB of pinned host memory and four events even when the copy then took the direct path -- D2D, an already-pinned source, or a managed destination. Measured rather than reasoned: `VT_ROCM_MANAGED_ALLOC=1` on gfx1151 printed `staged=0 direct=3 chunks=0 max_chunk=0 ring_bytes=268435456` and the case still passed, because the managed arm asserted nothing about `ring_bytes`. The instrument printed the defect and nothing read it. That assertion now exists and is what this commit turns green. It is not a free 256 MiB on the boards that show it. `.agents/environment.md` measures gfx1151's managed ceiling as bounded by HOST RAM, 27 GiB against 29.3 GiB available, so pinned host memory is charged to the resource that decides whether a checkpoint fits, on the arm that never stages. `ShouldStageH2D` stays the single authority on the decision. It is now spelled as `StagingTermsExceptRing(in) && in.ring_available`, production asks the cheap half first and calls `EnsureRing` only when it passes, and a new truth-table case asserts the two agree over 480 inputs so the split cannot drift. The same restructure un-deadens `stream_capturing`. It was hardcoded false while a separate early return did the work, so the term the truth table gates was a value production never supplied. The probe now feeds the field and the predicate is the guard, which is what the spec always said it was. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code] --- include/vt/rocm/rocm_pinned_h2d.h | 33 ++++++++++++++++-- src/vt/rocm/rocm_backend.hip | 41 +++++++++++++++++----- tests/vt/test_rocm_pinned_h2d.cpp | 57 +++++++++++++++++++++++++++++++ 3 files changed, 120 insertions(+), 11 deletions(-) diff --git a/include/vt/rocm/rocm_pinned_h2d.h b/include/vt/rocm/rocm_pinned_h2d.h index e029051d6..8fc4bf4fe 100644 --- a/include/vt/rocm/rocm_pinned_h2d.h +++ b/include/vt/rocm/rocm_pinned_h2d.h @@ -61,6 +61,32 @@ struct StagedH2DInputs { bool ring_available = false; }; +// THE FOUR TERMS THAT COST NOTHING TO ANSWER, split out so the caller can ask +// them BEFORE it allocates anything. +// +// `ring_available` is the one term a caller cannot answer by looking: on the +// first qualifying copy, answering it MEANS allocating 256 MiB of pinned host +// memory and four events. Evaluating it eagerly -- as this file's first +// implementation did -- charges that allocation to every ROCm process that ever +// hands `Copy` one buffer of 64 MiB or more, including the processes that then +// take the direct path because the source is already pinned, the destination is +// managed, or the copy is D2D. Measured: `VT_ROCM_MANAGED_ALLOC=1` on gfx1151 +// printed `staged=0 direct=3 chunks=0 ring_bytes=268435456`, a ring built and +// never touched. `.agents/environment.md:95-100` measures that board's managed +// ceiling as bounded by HOST RAM, so those bytes come out of the binding +// resource on the arm that never stages. +// +// So the production path asks THIS first and calls the allocator only when it +// passes. `ShouldStageH2D` stays the single authority on the decision and the +// thing the truth table gates; this is that expression with `ring_available` +// held true, and a case in tests/vt/test_rocm_pinned_h2d.cpp asserts the two +// agree over the whole table so the split cannot drift. +constexpr bool StagingTermsExceptRing(const StagedH2DInputs& in) { + return in.chunk_bytes != 0 && !in.stream_capturing && + in.dst == PtrKind::kDevice && in.src == PtrKind::kUnregisteredHost && + in.bytes >= in.chunk_bytes; +} + // FIVE terms, all required. Spelled as one expression so the truth table in // tests/vt/test_rocm_pinned_h2d.cpp can flip exactly one at a time. // @@ -81,10 +107,11 @@ struct StagedH2DInputs { // * `bytes >= chunk_bytes` — below one chunk the ring is strictly one extra // copy of the bytes with no overlap and no second slot // ever used. A 4 KiB norm weight must not pay a bounce. +// +// The first four are `StagingTermsExceptRing` above, which is what production +// evaluates before it lets `EnsureRing` allocate anything. constexpr bool ShouldStageH2D(const StagedH2DInputs& in) { - return in.chunk_bytes != 0 && in.ring_available && !in.stream_capturing && - in.dst == PtrKind::kDevice && in.src == PtrKind::kUnregisteredHost && - in.bytes >= in.chunk_bytes; + return StagingTermsExceptRing(in) && in.ring_available; } // VT_ROCM_PINNED_H2D_MIB, in MiB. Absent or empty takes the default; an diff --git a/src/vt/rocm/rocm_backend.hip b/src/vt/rocm/rocm_backend.hip index 8b8704b82..18a742b83 100644 --- a/src/vt/rocm/rocm_backend.hip +++ b/src/vt/rocm/rocm_backend.hip @@ -579,17 +579,29 @@ class RocmBackend final : public Backend { // src/llama-model-loader.cpp:1440 (n_buffers = 4), :1449 (64 MiB), :1591-1642 // (wait on the slot's event, fill it, upload from it, record, advance). // - // Returns true when it handled the copy. The four cheap terms are checked - // before any HIP probe, deliberately: about 1,361 ResidentWeight calls per - // forward step reach Copy on this row's checkpoint, and a 4 KiB norm weight - // must not pay two hipPointerGetAttributes calls to learn that it is small. - // Those same two terms are then part of ShouldStageH2D, which stays the single - // authority on the decision and is the thing the truth table gates. + // Returns true when it handled the copy. THE ORDER OF THE TERMS IS THE + // DESIGN, cheapest first, and each step is paid for only by the copies that + // got past the previous one: + // + // 1. the two size terms, with no HIP call at all -- about 1,361 + // ResidentWeight calls per forward step reach Copy on this row's + // checkpoint, and a 4 KiB norm weight must not pay two + // hipPointerGetAttributes calls to learn that it is small; + // 2. one hipStreamIsCapturing and two hipPointerGetAttributes; + // 3. EnsureRing, which is not a probe but an ALLOCATION. + // + // Those terms are the same terms ShouldStageH2D spells, which stays the + // single authority on the decision and is the thing the truth table gates. bool StagedCopy(Queue& q, void* dst, const void* src, size_t bytes) { const size_t chunk = vt::rocm::PinnedH2DChunkBytes(); if (chunk == 0 || bytes < chunk) return false; hipStream_t stream = AsStream(q); - if (StreamIsCapturing(stream)) return false; + // Probed OUTSIDE the lock, and then handed to the predicate as term 4 + // rather than being a second, separate guard. An early `return false` here + // would leave `in.stream_capturing` hardcoded false, which makes the term + // the truth table gates a value production never supplies -- the guard + // would be real but it would not be the guard the table describes. + const bool capturing = StreamIsCapturing(stream); std::lock_guard lock(h2d_mutex_); vt::rocm::StagedH2DInputs in; @@ -597,7 +609,20 @@ class RocmBackend final : public Backend { in.dst = ClassifyPointer(dst); in.bytes = bytes; in.chunk_bytes = chunk; - in.stream_capturing = false; // probed above, outside the lock + in.stream_capturing = capturing; + // THE RING IS ALLOCATED LAST, AND ONLY WHEN EVERY OTHER TERM ALREADY HOLDS. + // `EnsureRing` is not a question, it is 256 MiB of pinned host memory and + // four events. Asking it first -- as the first implementation did -- built + // the ring in every ROCm process that ever copies 64 MiB, including the ones + // that then take the direct path anyway because the source is already + // pinned, the destination is managed, or the copy is D2D. + // `VT_ROCM_MANAGED_ALLOC=1` on gfx1151 measured exactly that: + // `staged=0 direct=3 chunks=0 ring_bytes=268435456`. On that board the + // managed ceiling is bounded by HOST RAM (.agents/environment.md:95-100), + // so the wasted bytes come out of the resource that decides whether the + // model fits. + in.ring_available = false; + if (!vt::rocm::StagingTermsExceptRing(in)) return false; in.ring_available = EnsureRing(chunk); if (!vt::rocm::ShouldStageH2D(in)) return false; diff --git a/tests/vt/test_rocm_pinned_h2d.cpp b/tests/vt/test_rocm_pinned_h2d.cpp index 14245df86..6739e5b58 100644 --- a/tests/vt/test_rocm_pinned_h2d.cpp +++ b/tests/vt/test_rocm_pinned_h2d.cpp @@ -33,6 +33,7 @@ using vt::rocm::RunStagedH2D; using vt::rocm::ShouldStageH2D; using vt::rocm::StagedH2DInputs; using vt::rocm::StagedH2DRing; +using vt::rocm::StagingTermsExceptRing; // A fake device: N pinned slots, an ordered event log, and a reassembled // destination. Everything the real .hip does, minus HIP. @@ -299,3 +300,59 @@ TEST_CASE("the ring size and chunk size are llama.cpp's, and bound the residency // The number the whole change exists to bound, for a model of ANY size. CHECK(vt::rocm::kPinnedH2DRingBytes == (static_cast(256) << 20)); } + +// --------------------------------------------------------------------------- +// 7. The split between the cheap terms and the ALLOCATING one. +// +// Production cannot evaluate `ring_available` without allocating 256 MiB of +// pinned host memory, so it asks StagingTermsExceptRing first and calls +// EnsureRing only when that passes. That is two expressions where the spec +// describes one decision, and two expressions drift. This case is what stops +// them: over the whole input space this file's truth table walks, +// StagingTermsExceptRing must be exactly ShouldStageH2D with `ring_available` +// held true -- no more, no less. Deleting a term from either one fails here. +// --------------------------------------------------------------------------- +TEST_CASE("the cheap terms are exactly the decision minus the allocating term") { + const PtrKind kinds[] = {PtrKind::kUnregisteredHost, PtrKind::kPinnedHost, + PtrKind::kDevice, PtrKind::kOther}; + const size_t chunks[] = {0, 1024, kPinnedH2DChunkBytesDefault}; + const size_t sizes[] = {0, 1023, 1024, kPinnedH2DChunkBytesDefault, + kPinnedH2DChunkBytesDefault * 3}; + size_t staged = 0; + size_t total = 0; + for (PtrKind s : kinds) { + for (PtrKind d : kinds) { + for (size_t c : chunks) { + for (size_t b : sizes) { + for (bool cap : {false, true}) { + StagedH2DInputs in; + in.src = s; + in.dst = d; + in.chunk_bytes = c; + in.bytes = b; + in.stream_capturing = cap; + + // The ring is the ONLY term the split holds back. + in.ring_available = true; + REQUIRE(StagingTermsExceptRing(in) == ShouldStageH2D(in)); + if (ShouldStageH2D(in)) ++staged; + + // And with no ring, the decision is always no, while the cheap + // terms are unmoved -- which is the whole point: production learns + // the answer is no WITHOUT paying for the ring to find out. + in.ring_available = false; + CHECK_FALSE(ShouldStageH2D(in)); + in.ring_available = true; + CHECK(StagingTermsExceptRing(in) == ShouldStageH2D(in)); + ++total; + } + } + } + } + } + // The table is not degenerate: some rows stage and most do not. A helper that + // returned a constant would satisfy the equality above and fail here. + CHECK(total == 480); + CHECK(staged > 0); + CHECK(staged < total); +} From 933039a2eaa9d5c353d852ea4bd582e630e008b4 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sun, 13 Sep 2026 07:38:27 +0000 Subject: [PATCH 08/10] record(MODEL-MM-QWEN4-EXP): say what the ROCm arm produced, and correct three figures it got wrong `docs/FEATURES.md` still read "NO TOKEN HAS COME OUT OF A GPU FOR THIS MODEL and none is claimed". That is false for ROCm since 2026-09-13: `--device auto` on gfx1151 returns 32 tokens with `finish_reason=length` on three launches out of three, fluent and prompt-dependent. What is claimed is LIVENESS and nothing more. This architecture has no GPU oracle at all -- llama.cpp aborts in `build_delta_net_chunking` before reading a byte and no vLLM revision implements `qwen4_exp` -- so there is no token gate, no denominator, and no admissible speed comparison, and the sentence says so. The CUDA half of the old claim is untouched and still owed by QSADEV and #2423. `docs/USAGE.md`'s checkpoint cell enumerated `--device cpu` and `--device cuda` and now names the ROCm arm on the same terms. Three figures are corrected rather than carried: The decode number. This spec quoted "5.27-5.29 tok/s, spread 0.34%" off four samples across two process launches. A fresh review launched it three more times and read 5.002, 5.183 and 5.097. Nine samples across six launches run 5.002 to 5.291, a 5.8% spread, so the honest quotation is the range 5.0-5.3 tok/s with the launch count beside it. The board's memory total. `rocm-host-residency-after-upload.md` section 6a said 31.88 GB climbed toward "the board's 33.27 GB total" and this spec's section 7 inherited the number and declined to resolve the discrepancy it created. `.agents/environment.md` already resolved it: since the 2026-09-11 firmware change this part reports 96.000 GiB of VRAM, and 33,270,497,280 B is the box's HOST RAM. Both specs now say so, and 71.96 GiB of 96.000 GiB needs no explanation. The two guarantees the repair moved. Section 3b term 3 promised `VT_ROCM_MANAGED_ALLOC=1` was unchanged while the ring was still built on that arm, and section 6 placed the capture guard in a truth-table term production never fed. Both now describe the code, and section 6 states plainly what is still not measured on a board. Section 5a records the repair wave's red, its green, its mutation and the one confirmation run that the model still produces tokens. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code] --- .../ISSUE-LOCAL-01M2CV6HSEM3QD2H37G9DTBJ0S.md | 19 +++ .agents/specs/rocm-chunked-pinned-h2d.md | 124 ++++++++++++++++-- .../specs/rocm-host-residency-after-upload.md | 16 ++- docs/FEATURES.md | 2 +- docs/USAGE.md | 7 +- 5 files changed, 153 insertions(+), 15 deletions(-) create mode 100644 .agents/issues/MODEL-MM-QWEN4-EXP/ISSUE-LOCAL-01M2CV6HSEM3QD2H37G9DTBJ0S.md diff --git a/.agents/issues/MODEL-MM-QWEN4-EXP/ISSUE-LOCAL-01M2CV6HSEM3QD2H37G9DTBJ0S.md b/.agents/issues/MODEL-MM-QWEN4-EXP/ISSUE-LOCAL-01M2CV6HSEM3QD2H37G9DTBJ0S.md new file mode 100644 index 000000000..8092048f3 --- /dev/null +++ b/.agents/issues/MODEL-MM-QWEN4-EXP/ISSUE-LOCAL-01M2CV6HSEM3QD2H37G9DTBJ0S.md @@ -0,0 +1,19 @@ +ID: ISSUE-LOCAL-01M2CV6HSEM3QD2H37G9DTBJ0S +Title: ROCm pinned H2D ring is allocated before the decision that would use it +Row: MODEL-MM-QWEN4-EXP +State: CLOSED +Kind: bug +GitHub: - +Mirror: PENDING +Availability: FULL +Created: 2026-09-13 +Updated: 2026-09-13 +Closed: 2026-09-13 + +## Problem + +RocmBackend::StagedCopy evaluated in.ring_available = EnsureRing(chunk) before ShouldStageH2D, so the first copy of 64 MiB or more on a non-capturing stream allocated 256 MiB of pinned host memory and four events whatever the pointer kinds turned out to be. Measured on gfx1151 with VT_ROCM_MANAGED_ALLOC=1: staged=0 direct=3 chunks=0 ring_bytes=268435456, a ring built and never touched, on the arm whose managed ceiling .agents/environment.md measures as bounded by host RAM. The device case printed the number and asserted nothing about it. A second, narrower defect rides with it: in.stream_capturing was hardcoded false while a separate early return did the work, so the truth table gated a value production never supplied. + +## Resolution + +Fixed 2026-09-13. EnsureRing is now the LAST term StagedCopy evaluates: the header splits the four cheap terms out as StagingTermsExceptRing, production asks those first and calls the allocator only when they pass, and ShouldStageH2D is spelled as StagingTermsExceptRing(in) && in.ring_available so it stays the single authority the truth table gates. A new 480-input case in tests/vt/test_rocm_pinned_h2d.cpp asserts the two agree over the whole table so the split cannot drift. The capture probe now feeds in.stream_capturing instead of a separate early return, so the guard is the predicate term the table describes; StreamIsCapturing's own return value stays ungated on a board and .agents/specs/rocm-chunked-pinned-h2d.md section 6 says why. RED on strix:gpu0 at c794b5dda with VT_ROCM_MANAGED_ALLOC=1: FAILURE, 1 of 4 assertions failed, ring_bytes=268435456. GREEN at 33a1eaaf8: 4/4, ring_bytes=0. Gates: test_rocm_pinned_h2d 10/1527, cross-device 61/84841, -tc=*DSA* 2/273, -tc=*pinned bounce* 1/8, every selector printing the ROCm-board line. Evidence /workspace/vtchunked/repair-20260913-073233/. diff --git a/.agents/specs/rocm-chunked-pinned-h2d.md b/.agents/specs/rocm-chunked-pinned-h2d.md index ef56b6171..76d6396f8 100644 --- a/.agents/specs/rocm-chunked-pinned-h2d.md +++ b/.agents/specs/rocm-chunked-pinned-h2d.md @@ -147,6 +147,18 @@ hold, and anything else takes the existing single call unchanged: (`hipMemoryTypeManaged` / `Unified`) never bounces either, because it is already device-addressable — which is exactly the `VT_ROCM_MANAGED_ALLOC=1` configuration, so that knob's behaviour is unchanged. + + **THIS TERM WAS TRUE OF THE DECISION AND FALSE OF THE PROCESS, AND A REVIEW + CAUGHT IT.** The first implementation evaluated `ring_available = + EnsureRing(chunk)` BEFORE `ShouldStageH2D`, so a managed destination still + built the ring: `VT_ROCM_MANAGED_ALLOC=1` on gfx1151 measured + `staged=0 direct=3 chunks=0 ring_bytes=268435456`, and the device case + printed that line while asserting nothing about it. The copy was direct, as + this term promises, and 256 MiB of pinned host memory was allocated anyway -- + on exactly the boards that never stage, and out of the resource + `.agents/environment.md` measures as gfx1151's real ceiling. `EnsureRing` is + now the LAST term evaluated and the managed arm of the device case asserts + `ring_bytes` is unmoved. 4. The stream is NOT capturing a graph. `hipEventSynchronize` inside a capture region aborts the capture. The capture contract documented at `rocm_backend.hip:334-348` already forbids "host<->device blocking copies" @@ -274,6 +286,71 @@ having measured nothing. | `test_backend_cross_device` | (none — whole binary) | 61 | 84841 | SUCCESS | | `test_backend_cross_device` | `-tc=*DSA*` | 2 | 273 | SUCCESS | +### 5a. THE REPAIR WAVE, AND THE ONE ASSERTION THAT WAS MISSING + +A fresh review returned FAIL on the eager `EnsureRing` (§3b term 3) and on the +dead `stream_capturing` field (§6). Both are repaired above. The gate was re-run +on `strix:gpu0` (gfx1151, ROCm 7.2.4) in the same `/tmp/vllmcpp-chunked-h2d` +clone, `LD_LIBRARY_PATH=/opt/rocm-7.2.4/lib`, evidence archived at +`/workspace/vtchunked/repair-20260913-073233/` (`rc` logs age out within a day, +so the share holds them). + +**The SHAs below are the ones the worker checked out and the evidence names.** +The branch was rebased onto `origin/main` twice afterwards, so `c794b5dda` is +now `d83c4b1c6` and `33a1eaaf8` is now `a6621c937`. The four files the gate +reads -- `include/vt/rocm/rocm_pinned_h2d.h`, `src/vt/rocm/rocm_backend.hip`, +`tests/vt/test_rocm_pinned_h2d.cpp` and +`tests/vt/test_backend_cross_device.cpp` -- are byte-identical across both +rewrites, and neither `43622bc37` nor `ee0644eab` touches any of them, so the +measurement carries. + +**RED FIRST, and the selector that produces it is the one nobody had run.** At +the test commit `c794b5dda` — the new managed-arm assertion, no fix — +`VT_ROCM_MANAGED_ALLOC=1 test_backend_cross_device -tc='*pinned bounce*'` +reported **1 case, 0 passed, 1 failed / 4 assertions, 3 passed, 1 failed**, +`Status: FAILURE!`, exit 1, and printed +`pinned H2D: managed_alloc=1 staged=0 direct=3 chunks=0 max_chunk=0 ring_bytes=268435456` +with `pinned H2D case ran on a ROCm board: 1`. That is the defect: a ring +allocated, never used, on the arm that never stages. Binaries +`fc26b58276f3e6f9a2cda2a550b0596a` (`test_rocm_pinned_h2d`) and +`3cc2b9a5a58c76f02807dd6ad42ac007` (`test_backend_cross_device`). + +**GREEN AFTER, at the fix commit `33a1eaaf8`**, binaries proven changed: +`eba82c4b2d73198179ebf29a0fd37461` and `45bacea1a99edb77d508f590532cb03d`. Every +selector prints its counts and its board line. + +| binary | selector | cases | assertions | result | +|---|---|---|---|---| +| `test_rocm_pinned_h2d` | (none — whole binary) | 10 | 1527 | SUCCESS | +| `test_backend_cross_device` | `-tc=*pinned bounce*` | 1 | 8 | SUCCESS | +| `test_backend_cross_device` | `-tc=*pinned bounce*`, `VT_ROCM_MANAGED_ALLOC=1` | 1 | 4 | SUCCESS, `ring_bytes=0` | +| `test_backend_cross_device` | (none — whole binary) | 61 | 84841 | SUCCESS | +| `test_backend_cross_device` | `-tc=*DSA*` | 2 | 273 | SUCCESS | + +The unit binary grew by one case and 1443 assertions. That one case is +"the cheap terms are exactly the decision minus the allocating term", which +walks 480 inputs and asserts `StagingTermsExceptRing` equals `ShouldStageH2D` +with `ring_available` held true, so the split the repair introduces cannot +drift from the decision the truth table gates. Deleting the `dst == kDevice` +term from the helper fails it (3 assertions, binaries `0f5b26d660f20f7e` clean +vs `7e301ff417b3d850` mutated, restored build back to `0f5b26d660f20f7e`), run +off-board because the header is HIP-free. The cross-device binary is unmoved at +61 / 84841: the repair adds one assertion to an arm that is SKIPPED in the +default configuration, which is why the managed selector had to be run +explicitly and is now a declared gate line. + +**THE MODEL STILL PRODUCES TOKENS AFTER THE REPAIR**, confirmed once rather +than re-measured, because §7's result is accepted and the staging decision is +unchanged for the arm that stages. At `33a1eaaf8`, `/tmp/ckpt-iq1s` (worker-local +shards), `--device auto --max-tokens 32 --temperature 0 --max-num-seqs 1 +--repeat 3`: three runs, each `finish_reason=length completion_tokens=32`, at +42.399 s / 6.086 s / 6.074 s, decode **5.258 and 5.269 tok/s**, exit 0. Evidence +at `/workspace/vtchunked/repair-model-20260913-073358/`. `vllm-cli` reports md5 +`364f11fddc6392feb79cdc23f8d69cf5`, which is byte-identical to §7's — and that +is NOT a stale build: the executable is a 26,720-byte shim and the code links +through `libvllm.so.0.0.3`, relinked at 07:33:59 from a `rocm_backend.hip.o` +rebuilt at 07:32:52 on the fix commit. + Baseline at `98e2cd7da`: `test_backend_cross_device` whole binary **60 cases / 84833 assertions**, `-tc=*DSA*` **2 / 273**. The whole binary therefore grew by exactly one case and eight assertions, which is this spec's case and nothing @@ -358,7 +435,21 @@ The artifact's compiled feature set is asserted before it is timed: `ldd` for oracle's own bound, and it replaces an unbounded pinned range the driver was creating per copy. - **A qualifying copy inside a graph capture.** Guarded by term 4, and the guard - is in the truth table. + is in the truth table. **A review found that this was only half true of the + first implementation.** `in.stream_capturing` was hardcoded `false` and a + separate early `return false` did the work, so the term the truth table gates + was a value production never supplied -- the program was safe and the + guarantee was in the wrong place. The probe now feeds the field and the + predicate IS the guard. + + **WHAT IS STILL NOT MEASURED ON A BOARD, STATED RATHER THAN GLOSSED:** + `StreamIsCapturing`'s own return value. No case in + `tests/vt/test_backend_cross_device.cpp` opens a capture region, and one was + considered and NOT written, because a pageable asynchronous H2D is itself + illegal inside a capture region -- such a case would measure HIP's refusal + rather than this guard, and a green would not distinguish the two. The term + it feeds is gated in the truth table; the probe is not. That is the honest + extent of it. - **`hipPointerGetAttributes` leaves a sticky error for an unregistered host pointer.** It returns `hipErrorInvalidValue` for one, which is the very answer we want, and the last error is cleared with `hipGetLastError()` immediately so @@ -416,9 +507,15 @@ differs is one environment variable this change reads. The ring is the cause. **The device memory is the tell.** Arm B stops at 29.69 GiB and stays there, exactly as §6a recorded. Arm A reaches 71.96 GiB, which is the whole checkpoint. So the wedge was never the model failing to fit; it was the upload never -finishing. (The 71.96 GiB figure exceeds the 33.27 GB `hipMemGetInfo` total §6a -quotes; the 96 GiB VRAM carve on this part is the obvious explanation and this -spec does NOT resolve the discrepancy, it reports both raw numbers.) +finishing. + +**THERE IS NO DISCREPANCY, AND THIS PARAGRAPH USED TO SAY THERE WAS.** It read +that 71.96 GiB "exceeds the 33.27 GB `hipMemGetInfo` total §6a quotes" and +declined to resolve it. `.agents/environment.md` §"strix" already resolved it: +since the 2026-09-11 firmware change this board reports `mem_info_vram_total` +and `hipMemGetInfo` total of **96.000 GiB**, while 33,270,497,280 B is the +box's HOST RAM. §6a read one for the other and this spec propagated it. Both +specs now say so; 71.96 GiB of 96.000 GiB is an unremarkable number. ### The generations @@ -436,11 +533,18 @@ Weight staging is lazy — `dense_attn::ResidentWeight` uploads on first use, behind the `d_dev` memo — so run 1 carries the one-time 72 GiB host-to-device transfer of the whole checkpoint. -**The steady-state number is 5.27-5.29 tok/s, and it is FOUR samples across TWO -independent process launches on two different source filesystems:** 5.291, -5.273, 5.274, 5.275. Spread max-to-min 0.34%. That satisfies the row's -three-repetition rule and it does so across process boundaries rather than three -times inside one handle, which is the stronger shape. gfx1151 fails about two +**The steady-state number is 5.0-5.3 tok/s.** This spec first wrote +"5.27-5.29 tok/s ... spread max-to-min 0.34%" off FOUR samples across TWO +process launches — 5.291, 5.273, 5.274, 5.275 — and **0.34% understates the +spread of this measurement.** A fresh review launched the process three more +times on the reviewed head and read 5.002, 5.183 and 5.097 tok/s, a 3.6% spread +on its own; the repair wave's confirmation run (§5a) read 5.258 and 5.269. Nine +samples across six independent launches run from **5.002 to 5.291 tok/s**, 5.8% +max-to-min. Quote the range and the launch count; 5.29 is not reproducible to +that precision and must not be written as if it were. The row's +three-repetition rule is satisfied several times over, and it is satisfied +across process boundaries rather than three times inside one handle, which is +the stronger shape. gfx1151 fails about two runs in five with an illegal GPU memory access (`ISSUE-LOCAL-01M2BY2M2ATNVR3XQKV2DB1BJD`); no arm here hit that signature, and all six generations completed with `finish_reason=length`. @@ -454,7 +558,7 @@ the whole process. ### What this does NOT claim - No throughput COMPARISON. There is no llama.cpp or vLLM denominator here, and - 5.28 tok/s is this engine's first number on this part, not a ratio. + 5.0-5.3 tok/s is this engine's first number on this part, not a ratio. - No load-time verdict for the ring. Arm A's 61.187 s and arm B's 35.909 s are not comparable: arm A read the checkpoint cold off CIFS and arm B read it with the page cache already warm. A load-time A/B needs interleaved repeats from a diff --git a/.agents/specs/rocm-host-residency-after-upload.md b/.agents/specs/rocm-host-residency-after-upload.md index b62ea2766..6ca946185 100644 --- a/.agents/specs/rocm-host-residency-after-upload.md +++ b/.agents/specs/rocm-host-residency-after-upload.md @@ -412,8 +412,20 @@ result is the same failure in both and the axis is not a number. **Device memory is no longer UNVERIFIED.** `rocm-smi` is not on `PATH` in the leased container, so this is read from `/sys/class/drm/card0/device/mem_info_vram_used`: 154,816,512 B at rest, climbing -to 31.88 GB of the board's 33.27 GB total. The earlier "717 MB of 103 GB" figure -does not describe this board and is withdrawn rather than carried forward. +to 31.88 GB. The earlier "717 MB of 103 GB" figure does not describe this board +and is withdrawn rather than carried forward. + +**CORRECTION, 2026-09-13.** This paragraph said "of the board's 33.27 GB total" +and that was a MISREAD. 33,270,497,280 B is this box's HOST RAM, not its VRAM +carve: `.agents/environment.md` §"strix" records that since the 2026-09-11 +firmware change `mem_info_vram_total` and `hipMemGetInfo` both report +**96.000 GiB**, against host RAM total / available of 33,270,497,280 B / +29,304,037,376 B. So 31.88 GB was never near a device ceiling, and nothing in +this measurement should be read as the model failing to fit. The header line +above quotes the same number correctly, as `30 GiB host`. The device figure +`.agents/specs/rocm-chunked-pinned-h2d.md` §7 reaches with the ring on -- +71.96 GiB -- sits comfortably inside 96.000 GiB and needs no explanation beyond +this one. **Where it blocks.** Sampling `/proc//{stat,wchan}` every 6 s over both runs, the uninterruptible thread is in `svm_range_set_attr` for 153 of 196 diff --git a/docs/FEATURES.md b/docs/FEATURES.md index cbeb6c39b..f3df90376 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -160,7 +160,7 @@ speed-pending, which [BENCHMARKS.md](BENCHMARKS.md) tracks. |---|---|---|---| | `Qwen3_5ForConditionalGeneration` | Qwen3.6-27B NVFP4 (`unsloth` @`890bdef7`, `nvidia` @`0893e160`); Qwen3.5-4B BF16; **Qwen3.8-27B BF16** @`1d4bf0f2` | 27B strict 235/235 text + 32/32 image/video; 4B cached 3/3; Qwen3.8-27B 4/7 strict, 3 exact fp32 ties in band (#915) | `unsloth` 27B at/above vLLM, ModelOpt 0.85x; 4B 1.021x; 3.8-27B c4 **0.963x**, c1/c8 absolutes (#915). Loads BF16/per-tensor FP8/NVFP4 (CT+ModelOpt); `modelopt_mixed` FP8 tower NATIVE (#164), GDN qkvz merged. CUDA/CPU | | `Qwen3_5MoeForConditionalGeneration` | Qwen3.6-35B-A3B (NVFP4 text; published BF16 text + vision tower) | NVFP4 strict 315/315 vs vLLM 0.25.0; published BF16 6/7 prompts strict 16/16 vs the pin, 7th an exact tie (#910). Image/video IMPLEMENTED, NOT GATED (#891): the tower loads and runs, mm gate OWED | gate model: 0.93x to 1.03x grid; NO BF16 or mm speed claim | -| `Qwen4ExpForConditionalGeneration` | GGUF (`qwen4exp`) — **LOADS, DECODES, AND SERVES ONE SEQUENCE AT A TIME ON `--device cpu`** (W5a, W5f, W5g, W5k, W5L, #2031) | **THE GGUF LOAD IS GATED, AND SO IS THE LAYER LOOP; WHAT DOES NOT EXIST IS A SECOND STEP.** A `qwen4exp` file reaches the architecture's own config builder through the GGUF dispatch, the registry resolves the class, and `load_weights` materializes the whole text tower — **on CPU **and, since KGATHER, on CUDA** ([spec](../.agents/specs/cuda-quant-gather.md)). The 51.2 G-parameter n-gram table would expand from 26.822 GiB of IQ4_NL to 95.368 GiB of bf16 on a device that cannot gather blocks, which the on-disk #1123 device-fit guard cannot see, so the load REFUSES BY NAME on such a device ahead of any tensor I/O (#2083). `DeviceQuantGatherSupported` is no longer a device list at all: it asks `OpRegistered(kEmbeddingQuant, dev)`, CPU and CUDA register that op, and METAL, VULKAN, ROCM and TENSTORRENT do not and are still refused ([#2394](https://github.com/mudler/vllm.cpp/issues/2394)). So this row's earlier clause that the gate 'is true for CPU alone' is FALSE and is replaced — every convert-time transform inverted (the `+1` fold on every norm gamma with `ssm_norm` the one exception, `ssm_a` back to `log(-x)`, and the V-head reorder on every Gated DeltaNet tensor), gated in both directions against a committed 1224-tensor manifest of the shipped `unsloth/Qwen3.8-Flash-Next-GGUF UD-IQ1_S` and value-wise against a synthetic file. `Qwen4ExpTextModel::Forward` now exists (W5f, #2336) and `ModelRegistry::Forward` reaches it: the 48-layer loop composes all four block seams and is gated END TO END against the lane-pinned transformers 5.16.0 oracle at a tiny config, max|diff| 0.00982 against a 0.03 bf16-vs-f32 bound, with seven mutations separating by 0.78 to 2.02. **ONE TOKEN NOW COMES OUT OF A PRODUCTION SEAM, and the claim is exactly that and no more** (W5g): on a model loaded by `ModelRegistry::Load` from a synthetic `qwen4exp` GGUF, `ModelRegistry::Forward` returns `[1, vocab]` f32 logits with every element finite and `vt::GreedyArgmax` samples an id from them, on CPU. It is a REACH and a SAMPLE, not a token gate — the fixture's weights are a deterministic ramp, so the id is compared against no reference, and the tower's arithmetic is gated separately by the oracle case above. W5g is what made the prefill complete at all: `Qwen4ExpPleLayout` derived the n-gram head vocabulary from a DEFAULTED `ngram_vocab_size_base` and refused when a file's stated sizes disagreed, which a `qwen4exp` GGUF cannot avoid because llama.cpp #27742's converter writes the resolved arrays and no base — so the check held for exactly one artifact in existence and refused every other file with correctly loaded weights. The stated set is now the authority for the layout, as `NgramTableRows` already treated it, and the cross-check runs only where the source stated the base. **THE ENGINE'S CACHE CHANNEL NOW REACHES THIS MODEL** (W5j, #2031, #2353). This row said the `multi_kv` channel "is refused for every model by `ModelRegistry::Forward`"; that guard is now a MODEL-DECLARED capability, `ModelFactory::consumes_multi_kv`, landed with its first consumer, and this architecture is that consumer. A step carrying all three published groups reaches the hook, which resolves every one of its five caches BY NAME through `MultiKvCacheIndex::Resolve` — including the recurrent members, which `ENG-MULTIKV-BYNAME` made addressable — and reads the QSA indexer side cache out of the engine's own group-2 pages through group 2's own gathered block table. The guard still refuses `DeepseekV4ForCausalLM`, `Glm5NextForConditionalGeneration` and every architecture that declares nothing, and clearing the bit drives that refusal red in the gate. **IT NOW DECODES, AND IT NOW SERVES** (W5k and W5L, #2031). This row said "it still decodes NO token, and the reason is now the MODEL and not the engine", and named a dtype and a residency: the recurrent group publishes the PLE conv ring at the model dtype while `RunQwen4ExpPleBlock` required f32, and publishes the n-gram token history as a device i64 state while the same block read it through a HOST pointer. W5k settled both against the RUNNING lane pin — transformers 5.16.0, `modeling_qwen4_exp.py` sha256 `77fec77d…c459`, confirmed by regenerating this row's committed forward golden byte-identically. Upstream types each cache slot from the tensor that first reaches it (`cache_utils.py:1019-1023`), so the ring carries the MODEL dtype and the history lives on the compute device: the PUBLISHER was right twice and both requirements moved to the block. `ModelRegistry::Forward` then runs a prefill and a `past_len > 0` DECODE over one set of persistent caches. **W5L drives the engine itself.** A real `GPUModelRunner` allocates all three published groups, gathers every group's block table, publishes the five-name by-name index and runs `execute_model` / `sample_tokens` for a prefill and then a decode; `LoadedEngine::FromModelDir` loads a `qwen4exp` GGUF and `generate` returns tokens; and `examples/server` answers `POST /v1/completions` on CPU. The cross-step gate is the PLE n-gram history read out of the RUNNER's own recurrent state at the slot the runner assigned — int64 token ids, which cannot saturate as this fixture's bf16 activations do. **WHAT SERVES IS EXACTLY THIS AND NO MORE: `--device cpu`, ONE SEQUENCE AT A TIME, over a GGUF.** `num_reqs > 1` is refused by name — `RunQwen4ExpQsaBlockPaged` takes a block table of one sequence — and because an EngineCore that meets that refusal dies rather than degrades, this factory sets `ModelFactory::serves_one_sequence_per_step` and `LoadedEngine::ResolveMaxNumSeqs` clamps `--max-num-seqs` to 1 and says so on stderr; concurrent clients are accepted and served in sequence. MEASURED before that clamp existed: three overlapping `/v1/completions` calls at `--max-num-seqs 4` each returned a 500 carrying this hook's own message, and the engine never served again. The quant arms are the loader's — IQ4_NL, Q5_0 and the dequantizing gather (#1989), whose CUDA arm landed with KGATHER. **ALL SIX `qwen4_exp` ops PLUS `vt::RmsNormGroup` NOW HAVE CUDA ARMS** (W6-CUDA and W6-CUDA-B, #2031) and this row's earlier sentence that "no CUDA arm exists for any `qwen4_exp` op" is false and is replaced. **THAT LEAVES THE GATHER, AND KGATHER LANDED IT**, so the sentence this row carried — that the one remaining reason is `EmbeddingKernelCuda` refusing a block-quantized table — is FALSE and is replaced. `vt::Embedding` on a CUDA queue decodes a block row across all 18 encodings `vt::cpu::BlockToFloat` decodes, measured bit-exact against the CPU arm on a GPU. **WITH BOTH LANDED, WHAT BLOCKS A CUDA FORWARD IS NEITHER OP REGISTRATION NOR THE LOADER, AND THIS IS A PREDICTION RATHER THAN A MEASUREMENT:** the expected shape is partial dispatch through the PLE, then a NAMED REFUSAL at the first QSA layer, because `qwen4_exp_qsa_block.cpp` still reads three operands on the HOST — `CheckRopeLayoutsAgree`, `IndexerRows` on the block table, and `Qwen4ExpQsaIndex` on `kv_lens` — which is owned by the QSADEV wave and NOT by this row. A second wall the synthetic fixture never reaches: `IsCudaKeepQuantSupported` still excludes IQ4_NL and Q5_0, which the released UD-IQ1_S uses, owned by [#2423](https://github.com/mudler/vllm.cpp/issues/2423). **NO TOKEN HAS COME OUT OF A GPU FOR THIS MODEL and none is claimed.** **NO TOKEN NUMBER AND NO SPEED NUMBER**: everything above ran on a synthetic fixture whose weights are a deterministic ramp, and the safetensors arm refuses because every published safetensors artifact exceeds every device this project owns. **THE CLAUSE 'no published `qwen4exp` checkpoint has been served' IS NOW HALF FALSE AND IS REPLACED BY A MEASUREMENT.** The released `unsloth/Qwen3.8-Flash-Next-GGUF` UD-IQ1_S (67.564 GiB, 3 shards, 1224 tensors) was driven through `examples/server` on `thor:gpu0` on 2026-08-30 (`rc` job `0f188dd1`, [evidence](bench-evidence/qwen4exp-released-checkpoint-serve-20260830.md)): **it LOADED and the server LISTENED, in 4446 s at 69.206 GiB peak RSS, keeping every one of its nine encodings quantized -- and it produced ZERO TOKENS.** `POST /v1/completions` returned 500 because the forward refused the artifact by name: `vt: qwen4_exp_gated_residual: input_mix_weight_down must be float (f32/bf16 for outputs)`. The file stores all **194** hyper-connection mix weights as Q8_0, the loader correctly keeps them quantized, and `vt::Qwen4ExpGatedResidual` accepted only float; every arm of the synthetic fixture wrote those same tensors as F32, which is why every gate on this row was green and none of them could see it. **W5p REMOVED THAT REFUSAL AT ITS SOURCE** (#2031): `mix_down`, `mix_up` and `block_inject` now accept a block-quantized `[N,K]` weight and route through `vt::MatmulBT`, which dispatches the keep-quant GEMM `kMatmulBTQuant` -- mirroring llama.cpp, which merged this architecture on 2026-08-27 (`6c84c7d5d`, first tag `b10660`), declares all six of these projections `GGML_OP_MUL_MAT` and never dequantizes one. The ELEMENTWISE operands did not move: a block-typed `hc_*_norm` gamma is still refused by name, which is llama.cpp's own split (`GGML_OP_MUL` for the norm, with an explicit f32 cast where a file-typed weight meets an elementwise multiply). The synthetic fixture grew the arm that was missing (`FixtureOpts::hc_mix_q8_0`), and `ModelRegistry::Forward` runs a prefill and a second prompt over a Q8_0-mix file; restoring the old contract reds that case with the verbatim string above, which is what makes the reach measured. **W5q RE-RAN THE RELEASED CHECKPOINT THROUGH THE REPAIRED PATH, AND THE REFUSAL IS GONE WHILE THE OUTPUT IS DEGENERATE** ([evidence](bench-evidence/qwen4exp-released-checkpoint-serve-20260831.md)): on `thor:gpu0` `--device cpu`, staged to worker-local disk, the artifact loads in **61 s** (against 4446 s from the CIFS share) at `VmHWM` 73.935 GiB, a 5-token prefill and eight decode steps run with nothing thrown, and `POST /v1/completions` returns **HTTP 200** with 8 tokens where W5n got a 500. **But every one of those tokens is id 0 — `!` in this file's own vocabulary — and the answer is BYTE-IDENTICAL for two different prompts.** So the forward is degenerate and prompt-independent on the real weights, no usable token has yet come out of a published `qwen4exp` checkpoint, and the DECODES and SERVES claims at the head of this row remain true OF THE FIXTURE. **W5s THEN GOT REAL TOKENS OUT OF IT, AND THE CAUSE WAS THE REPACK MARKER** ([evidence](bench-evidence/qwen4exp-released-checkpoint-tokens-20260831.md)): on `origin/main` `52f7ccbfc`, which carries W5r as well as W5p, the same artifact on the same box answers `" Paris. Given this fact, what is"` and `" 100°C at sea level"` — two different prompts, two different prompt-dependent completions, eight distinct token ids none of them 0. W5q's tree predated W5r, so on `thor` (aarch64 i8mm, where `vt::cpu::QuantRepackActive()` is TRUE) `dense_attn::ResidentWeight` was still dropping the repack marker and `kMatmulBTQuant` read `block_q8_0x4` buffers as flat `q8_0` on every hyper-connection mix weight; a read-only per-stage probe puts a NaN in `stream.after_layer_0` (`nan=51200`) collapsing to an all-zero `LOGITS` row (`zero=248320`), and `argmax` over a row with no maximum returns index 0. Post-W5r that stage is `nan=0` and the logit row is `min -9.89818 max 15.7873` with argmax id 11751 = the `" Paris"` token; `VT_CPU_QUANT_REPACK=0` is byte-identical to the default, which is what a correct performance transform must be. **WHAT RUNS IS EXACTLY THIS: `--device cpu`, ONE SEQUENCE AT A TIME, the UD-IQ1_S GGUF arm, and no more.** **IT IS NOT A TOKEN GATE** — no oracle decoded these prompts, llama.cpp aborts in `build_delta_net_chunking` before loading a byte, the other six published quants are unrun, and there is no speed number. The repaired route is also a per-TOKEN matvec where llama.cpp batches the projection over the whole prefill; batching it is owed and unmeasured. The shipped GGUF is TEXT-ONLY (1224 tensors, no `v.blk.*`), so the multimodal arm has no artifact to load either. **CONFIG LAYER GATED as well.** The config resolves and validates against a RUNNING transformers 5.16.0 oracle (it imports without torch, so `validate_architecture` executes): a 39-case two-direction sweep agrees on 35 and differs on 4, all four being local guards stricter than upstream, never looser. All 15 upstream `validate_architecture` rejections are implemented and tabulated against their upstream line. The forward and the KV-cache spec REFUSE BY NAME, each naming the wave that owes it. vLLM implements `qwen4_exp` at NO revision, so the algorithm oracle is transformers **5.16.0** under an accepted lane exception; `gateable = no` because nothing published fits a fleet device — `Qwen/Qwen3.8-Flash-Next` is ~360 GB bf16, ~180 GB FP8, ~128 GB NVFP4 against ~119.6 GiB usable on GB10 | none, and no speed claim is admissible from this row until a token gate exists | +| `Qwen4ExpForConditionalGeneration` | GGUF (`qwen4exp`) — **LOADS, DECODES, AND SERVES ONE SEQUENCE AT A TIME ON `--device cpu`** (W5a, W5f, W5g, W5k, W5L, #2031) | **THE GGUF LOAD IS GATED, AND SO IS THE LAYER LOOP; WHAT DOES NOT EXIST IS A SECOND STEP.** A `qwen4exp` file reaches the architecture's own config builder through the GGUF dispatch, the registry resolves the class, and `load_weights` materializes the whole text tower — **on CPU **and, since KGATHER, on CUDA** ([spec](../.agents/specs/cuda-quant-gather.md)). The 51.2 G-parameter n-gram table would expand from 26.822 GiB of IQ4_NL to 95.368 GiB of bf16 on a device that cannot gather blocks, which the on-disk #1123 device-fit guard cannot see, so the load REFUSES BY NAME on such a device ahead of any tensor I/O (#2083). `DeviceQuantGatherSupported` is no longer a device list at all: it asks `OpRegistered(kEmbeddingQuant, dev)`, CPU and CUDA register that op, and METAL, VULKAN, ROCM and TENSTORRENT do not and are still refused ([#2394](https://github.com/mudler/vllm.cpp/issues/2394)). So this row's earlier clause that the gate 'is true for CPU alone' is FALSE and is replaced — every convert-time transform inverted (the `+1` fold on every norm gamma with `ssm_norm` the one exception, `ssm_a` back to `log(-x)`, and the V-head reorder on every Gated DeltaNet tensor), gated in both directions against a committed 1224-tensor manifest of the shipped `unsloth/Qwen3.8-Flash-Next-GGUF UD-IQ1_S` and value-wise against a synthetic file. `Qwen4ExpTextModel::Forward` now exists (W5f, #2336) and `ModelRegistry::Forward` reaches it: the 48-layer loop composes all four block seams and is gated END TO END against the lane-pinned transformers 5.16.0 oracle at a tiny config, max|diff| 0.00982 against a 0.03 bf16-vs-f32 bound, with seven mutations separating by 0.78 to 2.02. **ONE TOKEN NOW COMES OUT OF A PRODUCTION SEAM, and the claim is exactly that and no more** (W5g): on a model loaded by `ModelRegistry::Load` from a synthetic `qwen4exp` GGUF, `ModelRegistry::Forward` returns `[1, vocab]` f32 logits with every element finite and `vt::GreedyArgmax` samples an id from them, on CPU. It is a REACH and a SAMPLE, not a token gate — the fixture's weights are a deterministic ramp, so the id is compared against no reference, and the tower's arithmetic is gated separately by the oracle case above. W5g is what made the prefill complete at all: `Qwen4ExpPleLayout` derived the n-gram head vocabulary from a DEFAULTED `ngram_vocab_size_base` and refused when a file's stated sizes disagreed, which a `qwen4exp` GGUF cannot avoid because llama.cpp #27742's converter writes the resolved arrays and no base — so the check held for exactly one artifact in existence and refused every other file with correctly loaded weights. The stated set is now the authority for the layout, as `NgramTableRows` already treated it, and the cross-check runs only where the source stated the base. **THE ENGINE'S CACHE CHANNEL NOW REACHES THIS MODEL** (W5j, #2031, #2353). This row said the `multi_kv` channel "is refused for every model by `ModelRegistry::Forward`"; that guard is now a MODEL-DECLARED capability, `ModelFactory::consumes_multi_kv`, landed with its first consumer, and this architecture is that consumer. A step carrying all three published groups reaches the hook, which resolves every one of its five caches BY NAME through `MultiKvCacheIndex::Resolve` — including the recurrent members, which `ENG-MULTIKV-BYNAME` made addressable — and reads the QSA indexer side cache out of the engine's own group-2 pages through group 2's own gathered block table. The guard still refuses `DeepseekV4ForCausalLM`, `Glm5NextForConditionalGeneration` and every architecture that declares nothing, and clearing the bit drives that refusal red in the gate. **IT NOW DECODES, AND IT NOW SERVES** (W5k and W5L, #2031). This row said "it still decodes NO token, and the reason is now the MODEL and not the engine", and named a dtype and a residency: the recurrent group publishes the PLE conv ring at the model dtype while `RunQwen4ExpPleBlock` required f32, and publishes the n-gram token history as a device i64 state while the same block read it through a HOST pointer. W5k settled both against the RUNNING lane pin — transformers 5.16.0, `modeling_qwen4_exp.py` sha256 `77fec77d…c459`, confirmed by regenerating this row's committed forward golden byte-identically. Upstream types each cache slot from the tensor that first reaches it (`cache_utils.py:1019-1023`), so the ring carries the MODEL dtype and the history lives on the compute device: the PUBLISHER was right twice and both requirements moved to the block. `ModelRegistry::Forward` then runs a prefill and a `past_len > 0` DECODE over one set of persistent caches. **W5L drives the engine itself.** A real `GPUModelRunner` allocates all three published groups, gathers every group's block table, publishes the five-name by-name index and runs `execute_model` / `sample_tokens` for a prefill and then a decode; `LoadedEngine::FromModelDir` loads a `qwen4exp` GGUF and `generate` returns tokens; and `examples/server` answers `POST /v1/completions` on CPU. The cross-step gate is the PLE n-gram history read out of the RUNNER's own recurrent state at the slot the runner assigned — int64 token ids, which cannot saturate as this fixture's bf16 activations do. **WHAT SERVES IS EXACTLY THIS AND NO MORE: `--device cpu`, ONE SEQUENCE AT A TIME, over a GGUF.** `num_reqs > 1` is refused by name — `RunQwen4ExpQsaBlockPaged` takes a block table of one sequence — and because an EngineCore that meets that refusal dies rather than degrades, this factory sets `ModelFactory::serves_one_sequence_per_step` and `LoadedEngine::ResolveMaxNumSeqs` clamps `--max-num-seqs` to 1 and says so on stderr; concurrent clients are accepted and served in sequence. MEASURED before that clamp existed: three overlapping `/v1/completions` calls at `--max-num-seqs 4` each returned a 500 carrying this hook's own message, and the engine never served again. The quant arms are the loader's — IQ4_NL, Q5_0 and the dequantizing gather (#1989), whose CUDA arm landed with KGATHER. **ALL SIX `qwen4_exp` ops PLUS `vt::RmsNormGroup` NOW HAVE CUDA ARMS** (W6-CUDA and W6-CUDA-B, #2031) and this row's earlier sentence that "no CUDA arm exists for any `qwen4_exp` op" is false and is replaced. **THAT LEAVES THE GATHER, AND KGATHER LANDED IT**, so the sentence this row carried — that the one remaining reason is `EmbeddingKernelCuda` refusing a block-quantized table — is FALSE and is replaced. `vt::Embedding` on a CUDA queue decodes a block row across all 18 encodings `vt::cpu::BlockToFloat` decodes, measured bit-exact against the CPU arm on a GPU. **WITH BOTH LANDED, WHAT BLOCKS A CUDA FORWARD IS NEITHER OP REGISTRATION NOR THE LOADER, AND THIS IS A PREDICTION RATHER THAN A MEASUREMENT:** the expected shape is partial dispatch through the PLE, then a NAMED REFUSAL at the first QSA layer, because `qwen4_exp_qsa_block.cpp` still reads three operands on the HOST — `CheckRopeLayoutsAgree`, `IndexerRows` on the block table, and `Qwen4ExpQsaIndex` on `kv_lens` — which is owned by the QSADEV wave and NOT by this row. A second wall the synthetic fixture never reaches: `IsCudaKeepQuantSupported` still excludes IQ4_NL and Q5_0, which the released UD-IQ1_S uses, owned by [#2423](https://github.com/mudler/vllm.cpp/issues/2423). **A GPU HAS NOW PRODUCED TOKENS FOR THIS MODEL, ON ROCm, AND THE CLAIM IS LIVENESS AND NOT CORRECTNESS.** Measured 2026-09-13 on `strix:gpu0` (gfx1151, Radeon 8060S, ROCm 7.2.4) inside an `rc` lease ([spec](../.agents/specs/rocm-chunked-pinned-h2d.md)): the released `unsloth/Qwen3.8-Flash-Next-GGUF` UD-IQ1_S loads through `--device auto` and `examples/vllm-cli` returns 32 tokens with `finish_reason=length` on three of three launches -- fluent, coherent, prompt-dependent output, byte-identical across the three. Steady-state decode is **5.0-5.3 tok/s** -- nine samples across six independent process launches, 5.002 to 5.291 tok/s, 5.8% max-to-min. It is a RANGE: an earlier two-launch reading quoted 5.27-5.29 at a 0.34% spread and that precision does not reproduce. **IT IS NOT A TOKEN GATE AND IT IS NOT A PARITY CLAIM.** This architecture has no GPU oracle at all: llama.cpp aborts in `build_delta_net_chunking` before it reads a byte and no vLLM revision implements `qwen4_exp`, so nothing decoded these prompts beside us, there is no denominator, and no speed comparison is admissible. What reaches a token is the bounded pinned host-to-device ring; the same binary with `VT_ROCM_PINNED_H2D_MIB=0` stops at 29.69 GiB of device memory and produces nothing in 1200 s. **NO TOKEN HAS COME OUT OF A CUDA DEVICE FOR THIS MODEL and none is claimed** -- that half of the previous sentence is unchanged, and it is owed by the QSADEV host-operand wave and [#2423](https://github.com/mudler/vllm.cpp/issues/2423), not by the ROCm row. **NO TOKEN NUMBER AND NO SPEED NUMBER**: everything above ran on a synthetic fixture whose weights are a deterministic ramp, and the safetensors arm refuses because every published safetensors artifact exceeds every device this project owns. **THE CLAUSE 'no published `qwen4exp` checkpoint has been served' IS NOW HALF FALSE AND IS REPLACED BY A MEASUREMENT.** The released `unsloth/Qwen3.8-Flash-Next-GGUF` UD-IQ1_S (67.564 GiB, 3 shards, 1224 tensors) was driven through `examples/server` on `thor:gpu0` on 2026-08-30 (`rc` job `0f188dd1`, [evidence](bench-evidence/qwen4exp-released-checkpoint-serve-20260830.md)): **it LOADED and the server LISTENED, in 4446 s at 69.206 GiB peak RSS, keeping every one of its nine encodings quantized -- and it produced ZERO TOKENS.** `POST /v1/completions` returned 500 because the forward refused the artifact by name: `vt: qwen4_exp_gated_residual: input_mix_weight_down must be float (f32/bf16 for outputs)`. The file stores all **194** hyper-connection mix weights as Q8_0, the loader correctly keeps them quantized, and `vt::Qwen4ExpGatedResidual` accepted only float; every arm of the synthetic fixture wrote those same tensors as F32, which is why every gate on this row was green and none of them could see it. **W5p REMOVED THAT REFUSAL AT ITS SOURCE** (#2031): `mix_down`, `mix_up` and `block_inject` now accept a block-quantized `[N,K]` weight and route through `vt::MatmulBT`, which dispatches the keep-quant GEMM `kMatmulBTQuant` -- mirroring llama.cpp, which merged this architecture on 2026-08-27 (`6c84c7d5d`, first tag `b10660`), declares all six of these projections `GGML_OP_MUL_MAT` and never dequantizes one. The ELEMENTWISE operands did not move: a block-typed `hc_*_norm` gamma is still refused by name, which is llama.cpp's own split (`GGML_OP_MUL` for the norm, with an explicit f32 cast where a file-typed weight meets an elementwise multiply). The synthetic fixture grew the arm that was missing (`FixtureOpts::hc_mix_q8_0`), and `ModelRegistry::Forward` runs a prefill and a second prompt over a Q8_0-mix file; restoring the old contract reds that case with the verbatim string above, which is what makes the reach measured. **W5q RE-RAN THE RELEASED CHECKPOINT THROUGH THE REPAIRED PATH, AND THE REFUSAL IS GONE WHILE THE OUTPUT IS DEGENERATE** ([evidence](bench-evidence/qwen4exp-released-checkpoint-serve-20260831.md)): on `thor:gpu0` `--device cpu`, staged to worker-local disk, the artifact loads in **61 s** (against 4446 s from the CIFS share) at `VmHWM` 73.935 GiB, a 5-token prefill and eight decode steps run with nothing thrown, and `POST /v1/completions` returns **HTTP 200** with 8 tokens where W5n got a 500. **But every one of those tokens is id 0 — `!` in this file's own vocabulary — and the answer is BYTE-IDENTICAL for two different prompts.** So the forward is degenerate and prompt-independent on the real weights, no usable token has yet come out of a published `qwen4exp` checkpoint, and the DECODES and SERVES claims at the head of this row remain true OF THE FIXTURE. **W5s THEN GOT REAL TOKENS OUT OF IT, AND THE CAUSE WAS THE REPACK MARKER** ([evidence](bench-evidence/qwen4exp-released-checkpoint-tokens-20260831.md)): on `origin/main` `52f7ccbfc`, which carries W5r as well as W5p, the same artifact on the same box answers `" Paris. Given this fact, what is"` and `" 100°C at sea level"` — two different prompts, two different prompt-dependent completions, eight distinct token ids none of them 0. W5q's tree predated W5r, so on `thor` (aarch64 i8mm, where `vt::cpu::QuantRepackActive()` is TRUE) `dense_attn::ResidentWeight` was still dropping the repack marker and `kMatmulBTQuant` read `block_q8_0x4` buffers as flat `q8_0` on every hyper-connection mix weight; a read-only per-stage probe puts a NaN in `stream.after_layer_0` (`nan=51200`) collapsing to an all-zero `LOGITS` row (`zero=248320`), and `argmax` over a row with no maximum returns index 0. Post-W5r that stage is `nan=0` and the logit row is `min -9.89818 max 15.7873` with argmax id 11751 = the `" Paris"` token; `VT_CPU_QUANT_REPACK=0` is byte-identical to the default, which is what a correct performance transform must be. **WHAT RUNS IS EXACTLY THIS: `--device cpu`, ONE SEQUENCE AT A TIME, the UD-IQ1_S GGUF arm, and no more.** **IT IS NOT A TOKEN GATE** — no oracle decoded these prompts, llama.cpp aborts in `build_delta_net_chunking` before loading a byte, the other six published quants are unrun, and there is no speed number. The repaired route is also a per-TOKEN matvec where llama.cpp batches the projection over the whole prefill; batching it is owed and unmeasured. The shipped GGUF is TEXT-ONLY (1224 tensors, no `v.blk.*`), so the multimodal arm has no artifact to load either. **CONFIG LAYER GATED as well.** The config resolves and validates against a RUNNING transformers 5.16.0 oracle (it imports without torch, so `validate_architecture` executes): a 39-case two-direction sweep agrees on 35 and differs on 4, all four being local guards stricter than upstream, never looser. All 15 upstream `validate_architecture` rejections are implemented and tabulated against their upstream line. The forward and the KV-cache spec REFUSE BY NAME, each naming the wave that owes it. vLLM implements `qwen4_exp` at NO revision, so the algorithm oracle is transformers **5.16.0** under an accepted lane exception; `gateable = no` because nothing published fits a fleet device — `Qwen/Qwen3.8-Flash-Next` is ~360 GB bf16, ~180 GB FP8, ~128 GB NVFP4 against ~119.6 GiB usable on GB10 | none, and no speed claim is admissible from this row until a token gate exists | | `Qwen3_5ForCausalLM`, `Qwen3_5MoeForCausalLM` | none: no text-only Qwen3.5 checkpoint fits this hardware | **NO RUN GATE, OWED.** Gated on `test_qwen3_8_text_only.cpp`; NO token claim. Loader reads stacked BF16 experts (#740) plus BF16 towers, shared expert and `lm_head` (#864), so both published indices satisfy the load plan | not measured | | `Qwen3ForCausalLM` | Qwen3 dense 0.6B/1.7B/4B/32B, NVFP4A16 | near-tie strict 16/16 vs vLLM 0.25.0 | c1 every-axis parity, c8 decode residual | | `Qwen3MoeForCausalLM` | Qwen3-Coder-30B-A3B | strict 6/6 vs vLLM 0.25.0 | 11/16 grid cells at or above graphed vLLM | diff --git a/docs/USAGE.md b/docs/USAGE.md index ff24aaa5e..15d283c98 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -897,7 +897,10 @@ skips with that refusal quoted. as one pageable transfer, so host residency for the copy is 256 MiB whatever the model's size. Every smaller copy, every readback, every device-to-device copy, an already-pinned or managed source, and anything inside a graph capture - take the previous single call unchanged. Set `VT_ROCM_PINNED_H2D_MIB=0` to + take the previous single call, and a process that only ever makes those + copies allocates no pinned buffer at all -- the four buffers are created on + the first copy that will actually use them, not on the first copy that is + merely large enough. Set `VT_ROCM_PINNED_H2D_MIB=0` to restore that single call for every copy, or to another number of MiB to change the chunk. On an integrated part with no pageable-memory access this is what makes a large GGUF checkpoint load at all: with the ring off, a 67.56 GiB @@ -1039,7 +1042,7 @@ repository in this project's history. | dots3-note vision tower | `model-vision.safetensors` | 13,742,557,056 bytes | `dots-studio/dots3-note-prev` @ `1e1e7b0cd37a3a48a6c8d7fa55d5f9d14377006b` | Owed, as above | none | **W6a ported the DENSE half and W6b ([#2613](https://github.com/mudler/vllm.cpp/issues/2613)) the PYRAMID one, so this checkpoint's tower LOADS.** All 2195 of its `vision_encoder.*` tensors are read: 235 dense (`patch_embed`, blocks 0-24, the `patch_merger` adapter) plus the 1960 the pyramid adds (17 routed blocks x 8, and 608 routed experts x 3). The 17 F32 tensors in this otherwise all-BF16 file are exactly the `mlp.router_bias` buffers, one per routed block, which is upstream's own `dtype=torch.float32` (`vision.py:152-155` @ `9035151d6`) and is asserted in both directions at load. An image request against this tower is SERVED, and since W6c ([#2537](https://github.com/mudler/vllm.cpp/issues/2537)) it no longer has to be a multiple of 28 on each side: `PilResizeBicubicRgb` ports Pillow's own `Image.Resampling.BICUBIC` — support scaled by `max(1, in/out)` on a downscale, weights normalized per output pixel, 22-bit fixed point across two passes over a uint8 intermediate — so the resample upstream always performs now happens here too instead of being refused. What still refuses BY NAME is the blockwise-FP8 sibling below (W9); no bf16 dots3-note `vision_config` published so far selects the `use_bias` arm ([#2616](https://github.com/mudler/vllm.cpp/issues/2616)) or the softmax / top-k-below-2 router arms ([#2615](https://github.com/mudler/vllm.cpp/issues/2615)), which are the only bf16 configurations left owed. **W9d ([#2881](https://github.com/mudler/vllm.cpp/issues/2881)) read `vision_config.enable_fp8_moe` for the first time and found that upstream defaults it TRUE** (`vision.py:69` @ `9035151d6`), so upstream selects `MoESwiGLUFFNFP8` for all 17 pyramid blocks on this very file — and that class RUNS on it. `_per_block_cast_to_fp8_padded` rounds each expert shard up to a multiple of 128 (`vision.py:230-235`) and `per_block_cast_to_fp8` slices only back to the shape it was handed, which is that padded one (`deep_gemm/utils/math.py:61` @ DeepGEMM `e21c821f`, vendored as `vllm.third_party.deep_gemm`). So `moe_intermediate_size` 2112 becomes a 2176-row shard, `w13` is 4352 wide, `activated_size` is 2176, and `per_token_group_quant_fp8`'s `assert x.shape[-1] % group_size == 0` (`fp8_utils.py:563-566`) passes. The pad is also what makes the `w13` concat representable: 17 + 17 scale rows against `cdiv(4352,128)` = 34, where the unpadded reading gives 34 against 33. This file therefore serves an image through `layers::Fp8BlockMlpGateUpMethod` + `layers::Fp8BlockLinearMethod` with the `clamp_min` denominator. An earlier version of this entry said the class cannot run here; that was FALSE and [PR #2947](https://github.com/mudler/vllm.cpp/pull/2947) corrected it. What still takes the bf16 arm is a `vision_config` whose `embed_dim` is not a multiple of 128 — the first quantization is over the activation the tower hands in (`vision_moe.py:77-81`) and no pad reaches it — and such a tower says so on stderr at load rather than pretending. No published checkpoint carries one. Nothing about any of this is compared against vLLM: no oracle for this model runs on any host this project reaches | | dots3-note audio tower | `model-audio.safetensors` | 1,772,399,360 bytes | `dots-studio/dots3-note-prev` @ `1e1e7b0cd37a3a48a6c8d7fa55d5f9d14377006b` | Owed, as above | `added_tokens.json` sha256 `1aa71a4e0dbab80a72fd925389fd6c9cc52d1cb9da5dee8282784c15c6fa789b`; `tokenizer.json` sha256 `7f4e21a1d9fa472439f70201b4849977da5ec11e73df5a36552ab5ee99af554b` | **W7a ([#2703](https://github.com/mudler/vllm.cpp/issues/2703)) puts this tower on a SERVED request.** All 430 `audio_encoder.*` tensors are read and every one is BF16 — not one F32 in the file. An OpenAI `input_audio` part reaches the 32-layer `dots` speech encoder through `ApiServer::handle_chat_completions` on the default configuration. The three audio markers are resolved BY STRING from this checkpoint's own tokenizer, which carries them as special added tokens `<|audio_comp_start|>` 151718, `<|audio_comp_end|>` 151719, `<|audio_comp_pad|>` 151720 — note `pad == start + 2`, not `start + 1`. **W7b ([#2797](https://github.com/mudler/vllm.cpp/issues/2797)) removed the `chunk_seconds` ceiling**: a recording of any length is sliced into `chunk_seconds` segments, each padded to `chunk_samples` and log-melled on its own, run through the tower at ITS OWN valid sample length, and concatenated IN ORDER (`nvidia/audio.py:193-234` @ `9035151d6`). On this checkpoint `chunk_samples` is 960000 = 750 token strides, so the tower's per-segment row sum and the prompt side's single `ceil(total/stride)` agree for every waveform; a checkpoint whose `chunk_samples` is NOT a whole number of `token_stride`s is refused per request past one chunk, at the chat seam. **W7c-1 ([#2813](https://github.com/mudler/vllm.cpp/issues/2813)) lifted the MONO restriction**: a multi-channel PCM16 WAV already at 16 kHz is served, reduced to mono by the per-sample mean over its channels, which is upstream's own reduction (`load_audio(..., mono=True)` -> `np.mean`, `vllm/multimodal/media/audio.py:207-208`, `:220` @ `9035151d6`; and `ChannelReduction.MEAN` with `AudioSpec.target_channels = 1`, which dots3-note selects at `vllm/models/dots3_note/common/processor.py:523-525`). **W7c-2 ([#2828](https://github.com/mudler/vllm.cpp/issues/2828)) lifted the 16 kHz restriction**: a PCM16 WAV at any sampling rate is served, resampled to this checkpoint's `audio_config.sampling_rate` before the front end. This is the row's one RECORDED DIVERGENCE. Upstream's default resampler is `pyav`/libswresample, which is not bit-identical to itself across CPU dispatch on one binary and one input, so a bit-exact gate against it is impossible in principle; what is ported is `resample_audio_scipy` (`vllm/multimodal/audio.py:232-250` @ `9035151d6`), an arm of upstream's own `AudioResampler` switch that vLLM ships in production for phi4mm, and the gate is a consistency gate against `scipy.signal.resample_poly` with committed goldens. **Refused BY NAME**: a non-positive sample rate, and a reduced polyphase ratio whose `max(up, down)` exceeds 100000 — a deliberate divergence, because the rate is named by the request's own WAV header and the anti-alias filter is `20 * max(up, down) + 1` taps, while every ordinary rate reduces far below the bound (44100 -> 441, 48000 -> 3, 22050 -> 441, 8000 -> 2); any container but RIFF/WAVE PCM16 — `mp3`, `flac`, `ogg` need a demuxer this tree does not vendor, and that is owned by the shared codec brick [#2814](https://github.com/mudler/vllm.cpp/issues/2814) rather than by this row; and the unshipped `audio_config` arms (`use_causal`, `use_conv1d_stem`, `use_latent_input`, `merge_factor != 1`, a non-`dots` `encoder_type`), none of which this checkpoint selects. The blockwise-FP8 sibling below is still W9 | | dots3-note blockwise-FP8 sibling | `model-000{01..131}-of-00131.safetensors` plus the two tower files | 298,673,280,504 bytes total (278.16 GiB) across 133 safetensors, read 2026-08-28 | `dots-studio/dots3-note-prev-fp8` @ `7c14222e22423d6df6848eb0d1c5c3a88a00311a` | Owed: only `config.json` and `model.safetensors.index.json` were read | none | **Refused BY NAME at the forward, naming W9.** Its `quantization_config` is `{"quant_method": "fp8", "fmt": "e4m3", "activation_scheme": "dynamic", "weight_block_size": [128, 128]}` and its index (73,029 entries) ships a `weight_scale_inv` beside every projection — at the routed experts' `[1536, 5120]` that scale is `[12, 40]`. This port's bf16 loaders read a per-tensor or per-output-ROW `_scale` and nothing else, so without the named refusal the load would fail with a bare "tensor not found". It does not fit either: 278.16 GiB against the same 122 GiB ceiling | -| Qwen3.8-Flash-Next GGUF | `Qwen3.8-Flash-Next-UD-IQ1_S-0000{1..3}-of-00003.gguf` | 72,546,461,344 bytes total (67.564 GiB) across three shards (10,946,624 + 49,990,818,368 + 22,544,696,352); 1224 tensors | `unsloth/Qwen3.8-Flash-Next-GGUF` @ `8bdc666649440e9bdc97e16f3f75782c98478ff5`, path `UD-IQ1_S` | `88a1420825a9304063e882ada29d438263617f51ac8923d438d927496693bafd` (shard 1); `3a62e35bbf9add4733bd1438ebd3a67649d5edd6cb0e72bb78e33c913992b2b6` (shard 2); `0e25ceaeb89b8a80aa973c6c0c7448943682f7408c2855b2ebd016b7643a861a` (shard 3). Shard 1's digest was recomputed TWICE for this row -- on the development box and again INSIDE the `thor` lease against the bytes the server actually opened. Shards 2 and 3 carry the digests recorded in [the ladder-arm evidence file](bench-evidence/qwen4exp-llamacpp-ladder-arm-20260829.md), which recomputed all three on the staged copy on 29 August 2026; **this wave did not re-derive those two**, because the hash was killed mid-run for reading the same CIFS share as the load being measured | **LOADS on `--device cpu`, and the server LISTENS -- it produces NO TOKEN.** Measured on `thor:gpu0` 2026-08-30 (`rc` job `0f188dd1`, [evidence](bench-evidence/qwen4exp-released-checkpoint-serve-20260830.md)): all three shards load through `LoadedEngine::FromModelDir`, the engine sizes all three published cache groups, the tokenizer and the 9993-character chat template come out of the GGUF's own metadata, and `examples/server` answers on `/health`. **Load wall time 4446 s (74.1 min); peak RSS `VmHWM` 69.206 GiB against a 67.564 GiB artifact.** Residency is keep-quant: anonymous memory moved 4 -> 11 GiB across a load whose n-gram table alone would have added 95.368 GiB there, so all nine encodings in the file (F32, Q8_0, Q4_K, Q5_K, Q6_K, IQ2_XXS, IQ1_S, IQ4_NL, BF16) keep their blocks. `POST /v1/completions` then returns **500** and zero tokens | **THE FORWARD REFUSED THIS ARTIFACT BY NAME ON THAT RUN, AND W5p REMOVED THE REFUSAL**: `vt: qwen4_exp_gated_residual: input_mix_weight_down must be float (f32/bf16 for outputs)`. The file stores all **194** hyper-connection mix weights (`blk.N.hc_{attn,ffn}_{down,up}.weight` and `output_hc_{down,up}.weight`) as **Q8_0**; our loader correctly keeps them quantized (`qwen4_exp_weights.cpp` -> `LoadMatmul`), and `vt::Qwen4ExpGatedResidual` accepted only float, while every arm of the synthetic fixture wrote those same names as ggml type 0 (F32) -- so every prior wave gated the float case only and none could see this. Since W5p the three PROJECTION operands (`mix_down`, `mix_up`, `block_inject`) accept a block-quantized `[N,K]` weight and route through `vt::MatmulBT`/`kMatmulBTQuant`, mirroring llama.cpp, which merged this architecture on 2026-08-27 (`6c84c7d5d`, first tag `b10660`) and declares all six of them `GGML_OP_MUL_MAT`; the ELEMENTWISE `hc_*_norm` gamma is still refused by name, which is llama.cpp's own split. `FixtureOpts::hc_mix_q8_0` is the fixture arm that was missing. **W5q RE-RAN THIS ARTIFACT ON 2026-08-31** ([evidence](bench-evidence/qwen4exp-released-checkpoint-serve-20260831.md)): staged to worker-local disk it loads in **61 s** rather than 4446 s, `VmHWM` 73.935 GiB, the prefill and eight decode steps complete with nothing thrown, and `POST /v1/completions` returns **200** with 8 tokens. **Every token was id 0 (`!`) and two different prompts returned a byte-identical answer.** **W5s RE-RAN IT ON 2026-08-31 ON `origin/main` `52f7ccbfc`, WHICH CARRIES W5r, AND THE TOKENS ARE REAL** ([evidence](bench-evidence/qwen4exp-released-checkpoint-tokens-20260831.md)): `"The capital of France is"` -> `" Paris. Given this fact, what is"` and `"Water boils at"` -> `" 100°C at sea level"`, eight distinct token ids none of them 0, loaded in 60 s from the staged copy at `VmHWM` 73.93 GiB with system `used` flat at 11 GiB. **The cause of W5q's degeneracy was the dropped repack marker W5r fixed**: on this aarch64 i8mm box `kMatmulBTQuant` had been reading `block_q8_0x4` buffers as flat `q8_0`, putting a NaN in layer 0 that collapsed to an all-zero logit row, and `argmax` over a row with no maximum returns index 0. `VT_CPU_QUANT_REPACK=0` now gives byte-identical output to the default. **WHAT RUNS IS EXACTLY THIS AND NO MORE: `--device cpu`, ONE SEQUENCE AT A TIME, the UD-IQ1_S arm.** It is **NOT a token gate** — no oracle decoded these prompts, and there is no speed number. ISSUE OWED (this account is suspended for GitHub **API** writes -- `gh issue create` returns `HTTP 403: Sorry. Your account was suspended`, while `git push` over SSH succeeds, which is how this row reached `main`); scoped under `## Owed` in [the spec](../.agents/specs/qwen4-exp-flash-next.md). **Also refused or absent:** the other six published quants (UD-IQ1_M, UD-Q2_K_XL, UD-IQ3_XXS, UD-Q3_K_XL, UD-IQ4_XS, UD-Q4_K_XL) are staged but **none has been run**; every safetensors artifact (~360 GB bf16, ~180 GB FP8, ~128 GB NVFP4) exceeds the 122.80 GiB of the largest box in this fleet; the n-gram table stays HOST-side on every arm, because `DeviceQuantGatherSupported` is true for `kCPU` alone and moving it would expand it from 26.822 GiB to 95.368 GiB ([#2083](https://github.com/mudler/vllm.cpp/issues/2083)) -- the clause here previously said that "any non-CPU device refuses by name ahead of tensor I/O", which the CUDA run below falsifies; **`--device cuda` NOW SERVES THIS ARTIFACT, FLUENTLY BUT NOT TOKEN-EXACTLY** (the sentence here previously read that `ModelRegistry::Forward` is all-or-nothing and "no `qwen4_exp` step reaches a CUDA queue", which stopped being true once every op on the path had a device arm): on `thor:gpu0` `sm_110` the same binary answers `The capital of France is` with `11751 13 15767 411 1928 11 628 567` against the CPU control's `11751 13 15767 411 2029 11 1092 369` -- five of eight ids, both continuations grammatical English. **BOTH ARMS NOW MIRROR vLLM: since [#2612](https://github.com/mudler/vllm.cpp/issues/2612) the CPU GDN prefill runs vLLM's chunked decomposition too, by default, so the CPU/CUDA distinction this cell used to draw is GONE.** **BOTH SEQUENCES ABOVE WERE RE-MEASURED UNDER THE NEW DEFAULT ON 2026-09-04 AND NEITHER MOVED** ([evidence](bench-evidence/qwen4exp-gdn-chunked-token-ids-20260904.md), [#2858](https://github.com/mudler/vllm.cpp/issues/2858)): same artifact, same prompt, production configuration, `11751 13 15767 411 2029 11 1092 369` on `--device cpu` and `11751 13 15767 411 1928 11 628 567` on `--device cuda`, still five of eight. This cell previously said those CPU ids were "expected to move and have NOT been re-measured"; that expectation is falsified, and `VT_GDN_CHUNKED=0` now emits the same eight ids as the default rather than a different set. **The chunked arm is not inert** -- on the same binary it moves decoder layer 0's Gated DeltaNet block output by `3.702e-04` -- it simply flips none of the eight argmaxes, which is what the row's spec says to expect from an argmax over near-ties. This cell also claimed the port "brings the CPU arm to `1.772e-05` of CUDA where the sequential arm sat `3.525e-04` away, a 19.9x reduction", that the whole-model gap is "now dominated by the undiagnosed MoE residue", and that `out` "improves only 1.23x"; **all three are withdrawn by the annotation below**, which reads the same three measurements under one framing instead of a different framing per row. **ANNOTATION 2026-09-04 ([#2877](https://github.com/mudler/vllm.cpp/issues/2877)): NO `VT_Q4EXP_LAYER_FP` FINGERPRINT HAS YET OBSERVED A STEP WHERE THE TWO ARMS DISAGREE, AND THE METRIC BEHIND EVERY RATIO ABOVE CANNOT RANK THE STEPS IT DID OBSERVE.** **CORRECTION 2026-09-05 ([#2969](https://github.com/mudler/vllm.cpp/issues/2969)): the annotation above said "NO INSTRUMENT", and that is wrong for the other tap.** `VT_MOE_SEL_FP` counts MoE block invocations rather than model forwards, so its window does reach the disagreeing forwards. MOEDIV's reading there is VOID, because that run's CUDA arm answered the degenerate pre-#2550 sequence, and a non-degenerate arm was what was missing rather than a wider window. **READING TAKEN 2026-09-06 ([#2998](https://github.com/mudler/vllm.cpp/issues/2998), [evidence](bench-evidence/qwen4exp-moe-selection-fwd-20260906.md)): forwards 4, 6 and 7 have now been read on a non-degenerate pair, and the expert selection DOES flip -- 296 of 576 slots over 8 forwards, against a negative control that flips none. IT EXPLAINS NOTHING ABOUT THE THREE IDS.** Forward 5 flips 48 of 48 expert slots and its id AGREES; forward 7 flips 48 of 48 and its id DISAGREES; forward 4 disagrees while flipping 33 of 48. A flip count separates the two groups in neither direction, so no attribution of the three disagreeing ids is claimed anywhere in this table. The sentences are kept so the shape of the error stays visible. **First and decisively:** `LayerFp` returns early on `s.step >= s.budget` (`qwen4_exp_forward.cpp:118`), so `VT_Q4EXP_LAYER_FP=3` fingerprints model forwards 0, 1 and 2 -- tokens `11751 13 15767`, **which AGREE on both arms** -- while the three disagreeing ids are emitted at forwards 4, 6 and 7, outside the window; forwards 0-2 are causally upstream of forward 4 through the recurrent state, so those taps are not irrelevant to the disagreement, they simply never observe it. **Second:** `rel(sum|x|)` is a difference of NORMS, not a norm of DIFFERENCES, so its zero means "equal L1 norm" and not "equal tensors". At this tap's `n = 12800`, over 400 seeds of a hermetic control (committed as `MetricSpread` in `tests/scripts/test_q4exp_layerfp_diff.py`), it under-reports a zero-mean perturbation by a MEDIAN 75x-140x with a p05..p95 of 34..1500 -- a distribution, where this cell previously quoted one seed draw as `~122x` -- and at a **fixed** true divergence two readings differ by a median **2.1x** and by **24x** at p95. Every statistical figure in this cell is quoted to the two significant figures 400 draws support, and that control READS THIS CELL: its `test_the_PUBLISHER_reproduces_docs_USAGE_md` case compares each figure printed here to the value it draws, so this cell and the control cannot move apart; the three-figure set this cell carried before [#2879](https://github.com/mudler/vllm.cpp/pull/2879) came from a script that was never committed and did not reproduce from it. **Third:** applied consistently, that spread leaves nothing ranked. The "19.9x reduction" above is CPU-sequential vs CUDA-**chunked** compared against CPU-chunked vs CUDA-chunked -- a change of algorithm on one side -- and the same three measurements read as **algorithm-matched CPU-vs-CUDA pairs** say `L00 blk` moved **16.7x FURTHER** (1.062e-06 seq/seq -> 1.772e-05 chunked/chunked), which is what the chunked decomposition's larger reassociation freedom predicts. Those two ratios sit at 6% and 7% of the metric's own no-change distribution; no change at all produces a ratio at least as large as 1.80x, 2.02x, 2.34x and 3.15x in 59%, 52%, 45% and 33% of draws -- this cell first called that "between its 33rd and 59th percentile", which states the complement and inverts the ranking: 3.15x sits at the **67th** percentile of no change, not the 33rd. So "the residue grew" and "the residue did not grow" are equally unsupported, and so is a 19.9x or a 16.7x at the block. The residue's mechanism was already named by [#2552](https://github.com/mudler/vllm.cpp/issues/2552) -- the keep-quant grouped expert GEMM's scale-sum reassociation plus a bimodal top-k term at a 32.9% exact-bf16-tie rate, both FLOORS that do not scale with input distance, and both faithful mirrors of vLLM and llama.cpp rather than defects. It is **still not closed, for a narrower reason since 2026-09-06**: this sentence said the matched pair's `4.324e-05` lands inside #2552's layer-0 flip bracket and that `VT_MOE_SEL_FP` was never run on it; **it has been run, and layer 0 does NOT flip** ([evidence](bench-evidence/qwen4exp-moe-selection-fwd-20260906.md)), so the no-flip bound lifts from `2.139e-05` to `4.324e-05` and the layer-0 residue at this pair is the keep-quant expert GEMM's reassociation rather than the top-k term. That tie rate figure is prefill-only on a different pair; the same run reads 24.8% (CPU) and 23.1% (CUDA) of 576 boundaries across all eight forwards on the matched pair, which is a different population and not a discrepancy. What stays open is the CAUSE of the three disagreeing ids ([#2999](https://github.com/mudler/vllm.cpp/issues/2999)). The sequential arm is still the more ACCURATE one -- it lands `1.15e-08` from the exact answer where vLLM's own chunked kernel lands `2.29e-04` ([decomposition](bench-evidence/gdn-chunked-decomposition-20260902.md)) -- and it is retained for exactly that, but accuracy and faithfulness are different things and the ids vLLM would emit are the chunked arm's. The two arms are not a defect apart: the first tensor that differs is decoder layer 0's Gated DeltaNet block output, from a bit-identical input, because the CUDA arm ran vLLM's chunked prefill decomposition and the CPU arm an exact sequential recurrence ([evidence](bench-evidence/qwen4exp-cuda-prefill-divergence-20260902.md), [#2547](https://github.com/mudler/vllm.cpp/issues/2547)) -- that divergence SOURCE is closed by [#2612](https://github.com/mudler/vllm.cpp/issues/2612), which put both arms on the chunked decomposition, though the row makes no token-agreement claim and none should be read into it; `num_reqs > 1` is refused by name; MTP is absent (**zero** `nextn`/`mtp` tensors of 1224 against 31 in the safetensors repo, [#1993](https://github.com/mudler/vllm.cpp/issues/1993)); and the file is TEXT-ONLY (no `v.blk.*`), so the multimodal arm has no artifact | +| Qwen3.8-Flash-Next GGUF | `Qwen3.8-Flash-Next-UD-IQ1_S-0000{1..3}-of-00003.gguf` | 72,546,461,344 bytes total (67.564 GiB) across three shards (10,946,624 + 49,990,818,368 + 22,544,696,352); 1224 tensors | `unsloth/Qwen3.8-Flash-Next-GGUF` @ `8bdc666649440e9bdc97e16f3f75782c98478ff5`, path `UD-IQ1_S` | `88a1420825a9304063e882ada29d438263617f51ac8923d438d927496693bafd` (shard 1); `3a62e35bbf9add4733bd1438ebd3a67649d5edd6cb0e72bb78e33c913992b2b6` (shard 2); `0e25ceaeb89b8a80aa973c6c0c7448943682f7408c2855b2ebd016b7643a861a` (shard 3). Shard 1's digest was recomputed TWICE for this row -- on the development box and again INSIDE the `thor` lease against the bytes the server actually opened. Shards 2 and 3 carry the digests recorded in [the ladder-arm evidence file](bench-evidence/qwen4exp-llamacpp-ladder-arm-20260829.md), which recomputed all three on the staged copy on 29 August 2026; **this wave did not re-derive those two**, because the hash was killed mid-run for reading the same CIFS share as the load being measured | **LOADS on `--device cpu`, and the server LISTENS -- it produces NO TOKEN.** Measured on `thor:gpu0` 2026-08-30 (`rc` job `0f188dd1`, [evidence](bench-evidence/qwen4exp-released-checkpoint-serve-20260830.md)): all three shards load through `LoadedEngine::FromModelDir`, the engine sizes all three published cache groups, the tokenizer and the 9993-character chat template come out of the GGUF's own metadata, and `examples/server` answers on `/health`. **Load wall time 4446 s (74.1 min); peak RSS `VmHWM` 69.206 GiB against a 67.564 GiB artifact.** Residency is keep-quant: anonymous memory moved 4 -> 11 GiB across a load whose n-gram table alone would have added 95.368 GiB there, so all nine encodings in the file (F32, Q8_0, Q4_K, Q5_K, Q6_K, IQ2_XXS, IQ1_S, IQ4_NL, BF16) keep their blocks. `POST /v1/completions` then returns **500** and zero tokens | **THE FORWARD REFUSED THIS ARTIFACT BY NAME ON THAT RUN, AND W5p REMOVED THE REFUSAL**: `vt: qwen4_exp_gated_residual: input_mix_weight_down must be float (f32/bf16 for outputs)`. The file stores all **194** hyper-connection mix weights (`blk.N.hc_{attn,ffn}_{down,up}.weight` and `output_hc_{down,up}.weight`) as **Q8_0**; our loader correctly keeps them quantized (`qwen4_exp_weights.cpp` -> `LoadMatmul`), and `vt::Qwen4ExpGatedResidual` accepted only float, while every arm of the synthetic fixture wrote those same names as ggml type 0 (F32) -- so every prior wave gated the float case only and none could see this. Since W5p the three PROJECTION operands (`mix_down`, `mix_up`, `block_inject`) accept a block-quantized `[N,K]` weight and route through `vt::MatmulBT`/`kMatmulBTQuant`, mirroring llama.cpp, which merged this architecture on 2026-08-27 (`6c84c7d5d`, first tag `b10660`) and declares all six of them `GGML_OP_MUL_MAT`; the ELEMENTWISE `hc_*_norm` gamma is still refused by name, which is llama.cpp's own split. `FixtureOpts::hc_mix_q8_0` is the fixture arm that was missing. **W5q RE-RAN THIS ARTIFACT ON 2026-08-31** ([evidence](bench-evidence/qwen4exp-released-checkpoint-serve-20260831.md)): staged to worker-local disk it loads in **61 s** rather than 4446 s, `VmHWM` 73.935 GiB, the prefill and eight decode steps complete with nothing thrown, and `POST /v1/completions` returns **200** with 8 tokens. **Every token was id 0 (`!`) and two different prompts returned a byte-identical answer.** **W5s RE-RAN IT ON 2026-08-31 ON `origin/main` `52f7ccbfc`, WHICH CARRIES W5r, AND THE TOKENS ARE REAL** ([evidence](bench-evidence/qwen4exp-released-checkpoint-tokens-20260831.md)): `"The capital of France is"` -> `" Paris. Given this fact, what is"` and `"Water boils at"` -> `" 100°C at sea level"`, eight distinct token ids none of them 0, loaded in 60 s from the staged copy at `VmHWM` 73.93 GiB with system `used` flat at 11 GiB. **The cause of W5q's degeneracy was the dropped repack marker W5r fixed**: on this aarch64 i8mm box `kMatmulBTQuant` had been reading `block_q8_0x4` buffers as flat `q8_0`, putting a NaN in layer 0 that collapsed to an all-zero logit row, and `argmax` over a row with no maximum returns index 0. `VT_CPU_QUANT_REPACK=0` now gives byte-identical output to the default. **WHAT RUNS IS EXACTLY THIS AND NO MORE: `--device cpu`, ONE SEQUENCE AT A TIME, the UD-IQ1_S arm.** It is **NOT a token gate** — no oracle decoded these prompts, and there is no speed number. ISSUE OWED (this account is suspended for GitHub **API** writes -- `gh issue create` returns `HTTP 403: Sorry. Your account was suspended`, while `git push` over SSH succeeds, which is how this row reached `main`); scoped under `## Owed` in [the spec](../.agents/specs/qwen4-exp-flash-next.md). **Also refused or absent:** the other six published quants (UD-IQ1_M, UD-Q2_K_XL, UD-IQ3_XXS, UD-Q3_K_XL, UD-IQ4_XS, UD-Q4_K_XL) are staged but **none has been run**; every safetensors artifact (~360 GB bf16, ~180 GB FP8, ~128 GB NVFP4) exceeds the 122.80 GiB of the largest box in this fleet; the n-gram table stays HOST-side on every arm, because `DeviceQuantGatherSupported` is true for `kCPU` alone and moving it would expand it from 26.822 GiB to 95.368 GiB ([#2083](https://github.com/mudler/vllm.cpp/issues/2083)) -- the clause here previously said that "any non-CPU device refuses by name ahead of tensor I/O", which the CUDA run below falsifies; **`--device cuda` NOW SERVES THIS ARTIFACT, FLUENTLY BUT NOT TOKEN-EXACTLY** (the sentence here previously read that `ModelRegistry::Forward` is all-or-nothing and "no `qwen4_exp` step reaches a CUDA queue", which stopped being true once every op on the path had a device arm): on `thor:gpu0` `sm_110` the same binary answers `The capital of France is` with `11751 13 15767 411 1928 11 628 567` against the CPU control's `11751 13 15767 411 2029 11 1092 369` -- five of eight ids, both continuations grammatical English. **BOTH ARMS NOW MIRROR vLLM: since [#2612](https://github.com/mudler/vllm.cpp/issues/2612) the CPU GDN prefill runs vLLM's chunked decomposition too, by default, so the CPU/CUDA distinction this cell used to draw is GONE.** **BOTH SEQUENCES ABOVE WERE RE-MEASURED UNDER THE NEW DEFAULT ON 2026-09-04 AND NEITHER MOVED** ([evidence](bench-evidence/qwen4exp-gdn-chunked-token-ids-20260904.md), [#2858](https://github.com/mudler/vllm.cpp/issues/2858)): same artifact, same prompt, production configuration, `11751 13 15767 411 2029 11 1092 369` on `--device cpu` and `11751 13 15767 411 1928 11 628 567` on `--device cuda`, still five of eight. This cell previously said those CPU ids were "expected to move and have NOT been re-measured"; that expectation is falsified, and `VT_GDN_CHUNKED=0` now emits the same eight ids as the default rather than a different set. **The chunked arm is not inert** -- on the same binary it moves decoder layer 0's Gated DeltaNet block output by `3.702e-04` -- it simply flips none of the eight argmaxes, which is what the row's spec says to expect from an argmax over near-ties. This cell also claimed the port "brings the CPU arm to `1.772e-05` of CUDA where the sequential arm sat `3.525e-04` away, a 19.9x reduction", that the whole-model gap is "now dominated by the undiagnosed MoE residue", and that `out` "improves only 1.23x"; **all three are withdrawn by the annotation below**, which reads the same three measurements under one framing instead of a different framing per row. **ANNOTATION 2026-09-04 ([#2877](https://github.com/mudler/vllm.cpp/issues/2877)): NO `VT_Q4EXP_LAYER_FP` FINGERPRINT HAS YET OBSERVED A STEP WHERE THE TWO ARMS DISAGREE, AND THE METRIC BEHIND EVERY RATIO ABOVE CANNOT RANK THE STEPS IT DID OBSERVE.** **CORRECTION 2026-09-05 ([#2969](https://github.com/mudler/vllm.cpp/issues/2969)): the annotation above said "NO INSTRUMENT", and that is wrong for the other tap.** `VT_MOE_SEL_FP` counts MoE block invocations rather than model forwards, so its window does reach the disagreeing forwards. MOEDIV's reading there is VOID, because that run's CUDA arm answered the degenerate pre-#2550 sequence, and a non-degenerate arm was what was missing rather than a wider window. **READING TAKEN 2026-09-06 ([#2998](https://github.com/mudler/vllm.cpp/issues/2998), [evidence](bench-evidence/qwen4exp-moe-selection-fwd-20260906.md)): forwards 4, 6 and 7 have now been read on a non-degenerate pair, and the expert selection DOES flip -- 296 of 576 slots over 8 forwards, against a negative control that flips none. IT EXPLAINS NOTHING ABOUT THE THREE IDS.** Forward 5 flips 48 of 48 expert slots and its id AGREES; forward 7 flips 48 of 48 and its id DISAGREES; forward 4 disagrees while flipping 33 of 48. A flip count separates the two groups in neither direction, so no attribution of the three disagreeing ids is claimed anywhere in this table. The sentences are kept so the shape of the error stays visible. **First and decisively:** `LayerFp` returns early on `s.step >= s.budget` (`qwen4_exp_forward.cpp:118`), so `VT_Q4EXP_LAYER_FP=3` fingerprints model forwards 0, 1 and 2 -- tokens `11751 13 15767`, **which AGREE on both arms** -- while the three disagreeing ids are emitted at forwards 4, 6 and 7, outside the window; forwards 0-2 are causally upstream of forward 4 through the recurrent state, so those taps are not irrelevant to the disagreement, they simply never observe it. **Second:** `rel(sum|x|)` is a difference of NORMS, not a norm of DIFFERENCES, so its zero means "equal L1 norm" and not "equal tensors". At this tap's `n = 12800`, over 400 seeds of a hermetic control (committed as `MetricSpread` in `tests/scripts/test_q4exp_layerfp_diff.py`), it under-reports a zero-mean perturbation by a MEDIAN 75x-140x with a p05..p95 of 34..1500 -- a distribution, where this cell previously quoted one seed draw as `~122x` -- and at a **fixed** true divergence two readings differ by a median **2.1x** and by **24x** at p95. Every statistical figure in this cell is quoted to the two significant figures 400 draws support, and that control READS THIS CELL: its `test_the_PUBLISHER_reproduces_docs_USAGE_md` case compares each figure printed here to the value it draws, so this cell and the control cannot move apart; the three-figure set this cell carried before [#2879](https://github.com/mudler/vllm.cpp/pull/2879) came from a script that was never committed and did not reproduce from it. **Third:** applied consistently, that spread leaves nothing ranked. The "19.9x reduction" above is CPU-sequential vs CUDA-**chunked** compared against CPU-chunked vs CUDA-chunked -- a change of algorithm on one side -- and the same three measurements read as **algorithm-matched CPU-vs-CUDA pairs** say `L00 blk` moved **16.7x FURTHER** (1.062e-06 seq/seq -> 1.772e-05 chunked/chunked), which is what the chunked decomposition's larger reassociation freedom predicts. Those two ratios sit at 6% and 7% of the metric's own no-change distribution; no change at all produces a ratio at least as large as 1.80x, 2.02x, 2.34x and 3.15x in 59%, 52%, 45% and 33% of draws -- this cell first called that "between its 33rd and 59th percentile", which states the complement and inverts the ranking: 3.15x sits at the **67th** percentile of no change, not the 33rd. So "the residue grew" and "the residue did not grow" are equally unsupported, and so is a 19.9x or a 16.7x at the block. The residue's mechanism was already named by [#2552](https://github.com/mudler/vllm.cpp/issues/2552) -- the keep-quant grouped expert GEMM's scale-sum reassociation plus a bimodal top-k term at a 32.9% exact-bf16-tie rate, both FLOORS that do not scale with input distance, and both faithful mirrors of vLLM and llama.cpp rather than defects. It is **still not closed, for a narrower reason since 2026-09-06**: this sentence said the matched pair's `4.324e-05` lands inside #2552's layer-0 flip bracket and that `VT_MOE_SEL_FP` was never run on it; **it has been run, and layer 0 does NOT flip** ([evidence](bench-evidence/qwen4exp-moe-selection-fwd-20260906.md)), so the no-flip bound lifts from `2.139e-05` to `4.324e-05` and the layer-0 residue at this pair is the keep-quant expert GEMM's reassociation rather than the top-k term. That tie rate figure is prefill-only on a different pair; the same run reads 24.8% (CPU) and 23.1% (CUDA) of 576 boundaries across all eight forwards on the matched pair, which is a different population and not a discrepancy. What stays open is the CAUSE of the three disagreeing ids ([#2999](https://github.com/mudler/vllm.cpp/issues/2999)). The sequential arm is still the more ACCURATE one -- it lands `1.15e-08` from the exact answer where vLLM's own chunked kernel lands `2.29e-04` ([decomposition](bench-evidence/gdn-chunked-decomposition-20260902.md)) -- and it is retained for exactly that, but accuracy and faithfulness are different things and the ids vLLM would emit are the chunked arm's. The two arms are not a defect apart: the first tensor that differs is decoder layer 0's Gated DeltaNet block output, from a bit-identical input, because the CUDA arm ran vLLM's chunked prefill decomposition and the CPU arm an exact sequential recurrence ([evidence](bench-evidence/qwen4exp-cuda-prefill-divergence-20260902.md), [#2547](https://github.com/mudler/vllm.cpp/issues/2547)) -- that divergence SOURCE is closed by [#2612](https://github.com/mudler/vllm.cpp/issues/2612), which put both arms on the chunked decomposition, though the row makes no token-agreement claim and none should be read into it; `num_reqs > 1` is refused by name; MTP is absent (**zero** `nextn`/`mtp` tensors of 1224 against 31 in the safetensors repo, [#1993](https://github.com/mudler/vllm.cpp/issues/1993)); and the file is TEXT-ONLY (no `v.blk.*`), so the multimodal arm has no artifact **AND `--device auto` NOW SERVES IT ON ROCm — A LIVENESS RESULT, NOT A CORRECTNESS ONE.** Measured 2026-09-13 on `strix:gpu0` (gfx1151, Radeon 8060S, ROCm 7.2.4, box exclusively leased) with the three shards staged to worker-local disk: `examples/vllm-cli --device auto --max-tokens 32 --temperature 0 --max-num-seqs 1` returns 32 tokens with `finish_reason=length`, three launches out of three, fluent and prompt-dependent, and steady-state decode is **5.0-5.3 tok/s** -- nine samples across six independent process launches, 5.002 to 5.291 tok/s, 5.8% max-to-min, quoted as a range because an earlier two-launch reading's 0.34% spread does not reproduce. **NOTHING HERE IS A TOKEN-EXACTNESS OR PARITY CLAIM**: llama.cpp aborts on this architecture before it reads a byte and no vLLM revision implements `qwen4_exp`, so this arm has NO oracle and no denominator, and the CPU-vs-CUDA id disagreement above is a separate, still-open question this run does not speak to. What made the ROCm arm reach a token at all is the bounded pinned H2D ring ([spec](../.agents/specs/rocm-chunked-pinned-h2d.md)): the same binary with `VT_ROCM_PINNED_H2D_MIB=0` stops at 29.69 GiB of device memory and produces nothing in 1200 s. | ### Convert a GLM-5.3-Flash checkpoint to GGUF From 309f626c4a6a9fd3d1c99b42c6ad26ab0c16f045 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sun, 13 Sep 2026 08:00:20 +0000 Subject: [PATCH 09/10] record(BACKEND-GATE-ROCM-VLLM): file the low-disk failure mode of the Strix oracle suite `scripts/agent-preflight.sh` fails its "tools suites" gate on this dev box with 49 failures and 1 error from `tests/tools/test_strix_vllm_oracle.py`. None of them mention disk. The cause is `tools/bench/strix_vllm_oracle/worker.py:257` raising `ValueError("disk headroom exhausted")`, which the test helper turns into an `assertEqual(returncode, 0)` whose message carries the traceback and whose name does not. It reproduces on a PRISTINE `git archive` of `origin/main` `ee0644eab`, to the same 49 and 1, on a host at 97% full. So it is an environment condition an unrelated row's preflight reads as its own red, which is the shape this repository calls an instrument whose failure looks like a result. A headroom guard that cannot run should SKIP with its reason named, as preflight already does for its five argument-starved gates. This file rides in the ROCm chunked-H2D pull request because that is the flow that found it, and because a filed gap is worth more tracked than left untracked in a worktree. It changes no code and belongs to no other change on this branch; drop the commit if it is preferred as its own pull request. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code] --- .../ISSUE-LOCAL-01M2CWFFYH3N4HSD6RR3EC3M7Z.md | 19 +++++++++++++++++++ 1 file changed, 19 insertions(+) create mode 100644 .agents/issues/BACKEND-GATE-ROCM-VLLM/ISSUE-LOCAL-01M2CWFFYH3N4HSD6RR3EC3M7Z.md diff --git a/.agents/issues/BACKEND-GATE-ROCM-VLLM/ISSUE-LOCAL-01M2CWFFYH3N4HSD6RR3EC3M7Z.md b/.agents/issues/BACKEND-GATE-ROCM-VLLM/ISSUE-LOCAL-01M2CWFFYH3N4HSD6RR3EC3M7Z.md new file mode 100644 index 000000000..435f07fc5 --- /dev/null +++ b/.agents/issues/BACKEND-GATE-ROCM-VLLM/ISSUE-LOCAL-01M2CWFFYH3N4HSD6RR3EC3M7Z.md @@ -0,0 +1,19 @@ +ID: ISSUE-LOCAL-01M2CWFFYH3N4HSD6RR3EC3M7Z +Title: the Strix vLLM oracle suite turns a low-disk box into 49 opaque test failures +Row: BACKEND-GATE-ROCM-VLLM +State: OPEN +Kind: bug +GitHub: - +Mirror: PENDING +Availability: FULL +Created: 2026-09-13 +Updated: 2026-09-13 +Closed: - + +## Problem + +On a host with little free disk, tests/tools/test_strix_vllm_oracle.py reports 49 failures and 1 error, and scripts/agent-preflight.sh fails its 'tools suites' gate with them. The cause is not the tree: tools/bench/strix_vllm_oracle/worker.py:257 raises ValueError('disk headroom exhausted') from session.headroom(), the test's built() helper asserts returncode == 0, and every case that builds a session fails with that traceback buried in an assertEqual message. Reproduced 2026-09-13 on a dev box at 97% full (15 GiB free of 447 GiB) against a PRISTINE git archive of origin/main ee0644eab, 49 failures and 1 error, byte-identical to what the same suite reports on row/MODEL-MM-QWEN4-EXP-ROCM-CHUNKED-H2D, whose diff touches no file under tools/ or tests/tools/. So this is an environment condition that an unrelated row's preflight reads as its own red. A headroom guard that cannot run should SKIP with the named reason, the way preflight already reports its five argument-starved gates, rather than fail 49 cases whose messages do not say 'disk'. Found while repairing MODEL-MM-QWEN4-EXP-ROCM-CHUNKED-H2D; not fixed there, because the fix is this row's tooling and needs its own scope. + +## Resolution + +- From 92d3305cdea3d34e0b64d514372b69a7ba473d56 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sun, 13 Sep 2026 08:14:32 +0000 Subject: [PATCH 10/10] test(MODEL-MM-QWEN4-EXP): make the split-drift case convict a deleted term, and correct what it was said to prove The 480-input case that landed with the eager-`EnsureRing` repair asserted `StagingTermsExceptRing(in) == ShouldStageH2D(in)` with `ring_available` held true. The decision is DEFINED as the helper AND that flag, so that assertion is `X == (X && true)`: true for any definition of the helper, deleted terms included. A fresh review measured it on `strix:gpu0`: deleting `in.dst == PtrKind::kDevice` from the helper, and separately deleting `in.bytes >= in.chunk_bytes`, each left the case 1443/1443 SUCCESS. Both mutants died to the PRE-EXISTING five-term truth table instead, three and two assertions elsewhere in the binary, so a drift guard that cannot detect drift was being cited as protection. The case now computes its expectation in the test file from the four inputs, as a sequence of refusals rather than as a conjunction, and checks BOTH expressions against it. A term deleted from either one now fails this case itself. The spec, the header comment and the pull request body each claimed the old case proved that; all three are corrected, and the spec's attribution of a whole-binary failure count to the case just added is withdrawn by name. `docs/FEATURES.md` contradicted itself inside one cell: it recorded that a GPU has produced tokens for this model on ROCm and five sentences later still said "WHAT RUNS IS EXACTLY THIS: `--device cpu` ... and no more". The sentence now names both arms and keeps the ROCm one as liveness and never parity, because this architecture has no GPU oracle and there is no token gate. The spec's section 6 still read "a 33.27 GB board" after the commit that claimed to correct that misread everywhere. 33,270,497,280 B is this box's HOST RAM; `hipMemGetInfo` reports 96.000 GiB since the 2026-09-11 firmware change (`.agents/environment.md:89-92`). Fixed, and both specs were swept: the number now appears only where it is labelled host RAM or inside a correction paragraph that names the misread. No release, staging or predicate logic changes here. The guarantee itself was already gated, by the five-term truth table this case sits beside. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code] --- .agents/specs/rocm-chunked-pinned-h2d.md | 53 +++++++++++++++++----- docs/FEATURES.md | 2 +- include/vt/rocm/rocm_pinned_h2d.h | 8 +++- tests/vt/test_rocm_pinned_h2d.cpp | 56 ++++++++++++++++++------ 4 files changed, 91 insertions(+), 28 deletions(-) diff --git a/.agents/specs/rocm-chunked-pinned-h2d.md b/.agents/specs/rocm-chunked-pinned-h2d.md index 76d6396f8..84614a5e4 100644 --- a/.agents/specs/rocm-chunked-pinned-h2d.md +++ b/.agents/specs/rocm-chunked-pinned-h2d.md @@ -321,20 +321,49 @@ selector prints its counts and its board line. | binary | selector | cases | assertions | result | |---|---|---|---|---| -| `test_rocm_pinned_h2d` | (none — whole binary) | 10 | 1527 | SUCCESS | +| `test_rocm_pinned_h2d` | (none — whole binary) | 10 | 2007 | SUCCESS | +| `test_rocm_pinned_h2d` | `-tc=*cheap terms*` | 1 | 1923 | SUCCESS | | `test_backend_cross_device` | `-tc=*pinned bounce*` | 1 | 8 | SUCCESS | | `test_backend_cross_device` | `-tc=*pinned bounce*`, `VT_ROCM_MANAGED_ALLOC=1` | 1 | 4 | SUCCESS, `ring_bytes=0` | | `test_backend_cross_device` | (none — whole binary) | 61 | 84841 | SUCCESS | | `test_backend_cross_device` | `-tc=*DSA*` | 2 | 273 | SUCCESS | -The unit binary grew by one case and 1443 assertions. That one case is -"the cheap terms are exactly the decision minus the allocating term", which -walks 480 inputs and asserts `StagingTermsExceptRing` equals `ShouldStageH2D` -with `ring_available` held true, so the split the repair introduces cannot -drift from the decision the truth table gates. Deleting the `dst == kDevice` -term from the helper fails it (3 assertions, binaries `0f5b26d660f20f7e` clean -vs `7e301ff417b3d850` mutated, restored build back to `0f5b26d660f20f7e`), run -off-board because the header is HIP-free. The cross-device binary is unmoved at +The unit binary grew by one case, "the cheap terms are exactly the decision +minus the allocating term", which walks 480 inputs. + +**A review measured that case's FIRST version as a tautology, and it is +replaced.** That version asserted `StagingTermsExceptRing(in) == +ShouldStageH2D(in)` with `ring_available` held true. The decision is DEFINED as +the helper AND that flag, so the assertion is `X == (X && true)` -- true for any +definition of the helper, deleted terms included. Measured on `strix:gpu0`: +deleting `dst == kDevice` from the helper, and separately deleting +`bytes >= chunk_bytes`, each left that case 1443/1443 SUCCESS. Both mutants died +to the PRE-EXISTING five-term truth table instead (3 and 2 failed assertions +respectively, elsewhere in the binary), so this spec's earlier sentence -- that +deleting `dst == kDevice` "fails it (3 assertions)" -- attributed a whole-binary +count to the case just added. That is the `touched the symbol is not caused the +failure` error, and the sentence is WITHDRAWN. + +The case now computes the expected answer in the test file from the four +inputs, as a sequence of refusals rather than as a conjunction, and checks BOTH +expressions against it, so a term deleted from either one fails the case itself. +It is run off-board because the header is HIP-free. + +**RE-MEASURED, and the new case now convicts.** Same box, same clone +(`/tmp/vllmcpp-chunked-h2d`), `rc` job `ecf7b0b6-6e6f-4123-b6d2-50d5178a32dd`, +at `8e1f0ac03`. Clean unit binary `32b4cade5a4bba7d2060f6086a681dde`, 10 cases / +2007 assertions SUCCESS, and the case alone 1 / 1923 SUCCESS. + +| mutation in `StagingTermsExceptRing` | binary md5 | `-tc=*cheap terms*` | whole unit binary | +|---|---|---|---| +| delete `in.dst == PtrKind::kDevice &&` | `c7f00ede469d373502f961347e5b7361` | **FAILURE**, 1 case failed, first failed assertion at `test_rocm_pinned_h2d.cpp:363` | FAILURE, 2 cases / 4 assertions failed | +| delete `in.bytes >= in.chunk_bytes` | `c1e2d3d36df481a2cc49c7f091c1210b` | **FAILURE**, 1 case failed, same line | FAILURE, 2 cases / 3 assertions failed | + +Both restored builds hash back to `32b4cade5a4bba7d2060f6086a681dde` and the +restored tree runs 10 / 2007 SUCCESS, which is what proves the restoration +rather than `git status`. A `REQUIRE` aborts its case, so a mutant run reports +fewer assertions than the clean one; the verdict is the failure, not the count. +The cross-device binary is unmoved at 61 / 84841: the repair adds one assertion to an arm that is SKIPPED in the default configuration, which is why the managed selector had to be run explicitly and is now a declared gate line. @@ -424,8 +453,10 @@ The artifact's compiled feature set is asserted before it is timed: `ldd` for - **The stall survives this too.** Then the trigger is neither the residency of the source nor the shape of the transfer, and the next hypothesis is the - allocation side — 29.69 GiB of `hipMalloc` on a 33.27 GB board behind a 96 GiB - carve, or the CIFS mount, which §6a already named as an unseparated confound. + allocation side — 29.69 GiB of `hipMalloc` against this board's + `hipMemGetInfo` total of **96.000 GiB** (`.agents/environment.md:89-92`; + 33,270,497,280 B is the box's HOST RAM, not its VRAM carve), or the CIFS + mount, which §6a already named as an unseparated confound. A negative result with a `wchan` distribution is the reportable outcome, not a failure of the change. - **An extra `memcpy` per 64 MiB slows the load.** Load seconds are recorded on diff --git a/docs/FEATURES.md b/docs/FEATURES.md index f3df90376..1cb0e2d49 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -160,7 +160,7 @@ speed-pending, which [BENCHMARKS.md](BENCHMARKS.md) tracks. |---|---|---|---| | `Qwen3_5ForConditionalGeneration` | Qwen3.6-27B NVFP4 (`unsloth` @`890bdef7`, `nvidia` @`0893e160`); Qwen3.5-4B BF16; **Qwen3.8-27B BF16** @`1d4bf0f2` | 27B strict 235/235 text + 32/32 image/video; 4B cached 3/3; Qwen3.8-27B 4/7 strict, 3 exact fp32 ties in band (#915) | `unsloth` 27B at/above vLLM, ModelOpt 0.85x; 4B 1.021x; 3.8-27B c4 **0.963x**, c1/c8 absolutes (#915). Loads BF16/per-tensor FP8/NVFP4 (CT+ModelOpt); `modelopt_mixed` FP8 tower NATIVE (#164), GDN qkvz merged. CUDA/CPU | | `Qwen3_5MoeForConditionalGeneration` | Qwen3.6-35B-A3B (NVFP4 text; published BF16 text + vision tower) | NVFP4 strict 315/315 vs vLLM 0.25.0; published BF16 6/7 prompts strict 16/16 vs the pin, 7th an exact tie (#910). Image/video IMPLEMENTED, NOT GATED (#891): the tower loads and runs, mm gate OWED | gate model: 0.93x to 1.03x grid; NO BF16 or mm speed claim | -| `Qwen4ExpForConditionalGeneration` | GGUF (`qwen4exp`) — **LOADS, DECODES, AND SERVES ONE SEQUENCE AT A TIME ON `--device cpu`** (W5a, W5f, W5g, W5k, W5L, #2031) | **THE GGUF LOAD IS GATED, AND SO IS THE LAYER LOOP; WHAT DOES NOT EXIST IS A SECOND STEP.** A `qwen4exp` file reaches the architecture's own config builder through the GGUF dispatch, the registry resolves the class, and `load_weights` materializes the whole text tower — **on CPU **and, since KGATHER, on CUDA** ([spec](../.agents/specs/cuda-quant-gather.md)). The 51.2 G-parameter n-gram table would expand from 26.822 GiB of IQ4_NL to 95.368 GiB of bf16 on a device that cannot gather blocks, which the on-disk #1123 device-fit guard cannot see, so the load REFUSES BY NAME on such a device ahead of any tensor I/O (#2083). `DeviceQuantGatherSupported` is no longer a device list at all: it asks `OpRegistered(kEmbeddingQuant, dev)`, CPU and CUDA register that op, and METAL, VULKAN, ROCM and TENSTORRENT do not and are still refused ([#2394](https://github.com/mudler/vllm.cpp/issues/2394)). So this row's earlier clause that the gate 'is true for CPU alone' is FALSE and is replaced — every convert-time transform inverted (the `+1` fold on every norm gamma with `ssm_norm` the one exception, `ssm_a` back to `log(-x)`, and the V-head reorder on every Gated DeltaNet tensor), gated in both directions against a committed 1224-tensor manifest of the shipped `unsloth/Qwen3.8-Flash-Next-GGUF UD-IQ1_S` and value-wise against a synthetic file. `Qwen4ExpTextModel::Forward` now exists (W5f, #2336) and `ModelRegistry::Forward` reaches it: the 48-layer loop composes all four block seams and is gated END TO END against the lane-pinned transformers 5.16.0 oracle at a tiny config, max|diff| 0.00982 against a 0.03 bf16-vs-f32 bound, with seven mutations separating by 0.78 to 2.02. **ONE TOKEN NOW COMES OUT OF A PRODUCTION SEAM, and the claim is exactly that and no more** (W5g): on a model loaded by `ModelRegistry::Load` from a synthetic `qwen4exp` GGUF, `ModelRegistry::Forward` returns `[1, vocab]` f32 logits with every element finite and `vt::GreedyArgmax` samples an id from them, on CPU. It is a REACH and a SAMPLE, not a token gate — the fixture's weights are a deterministic ramp, so the id is compared against no reference, and the tower's arithmetic is gated separately by the oracle case above. W5g is what made the prefill complete at all: `Qwen4ExpPleLayout` derived the n-gram head vocabulary from a DEFAULTED `ngram_vocab_size_base` and refused when a file's stated sizes disagreed, which a `qwen4exp` GGUF cannot avoid because llama.cpp #27742's converter writes the resolved arrays and no base — so the check held for exactly one artifact in existence and refused every other file with correctly loaded weights. The stated set is now the authority for the layout, as `NgramTableRows` already treated it, and the cross-check runs only where the source stated the base. **THE ENGINE'S CACHE CHANNEL NOW REACHES THIS MODEL** (W5j, #2031, #2353). This row said the `multi_kv` channel "is refused for every model by `ModelRegistry::Forward`"; that guard is now a MODEL-DECLARED capability, `ModelFactory::consumes_multi_kv`, landed with its first consumer, and this architecture is that consumer. A step carrying all three published groups reaches the hook, which resolves every one of its five caches BY NAME through `MultiKvCacheIndex::Resolve` — including the recurrent members, which `ENG-MULTIKV-BYNAME` made addressable — and reads the QSA indexer side cache out of the engine's own group-2 pages through group 2's own gathered block table. The guard still refuses `DeepseekV4ForCausalLM`, `Glm5NextForConditionalGeneration` and every architecture that declares nothing, and clearing the bit drives that refusal red in the gate. **IT NOW DECODES, AND IT NOW SERVES** (W5k and W5L, #2031). This row said "it still decodes NO token, and the reason is now the MODEL and not the engine", and named a dtype and a residency: the recurrent group publishes the PLE conv ring at the model dtype while `RunQwen4ExpPleBlock` required f32, and publishes the n-gram token history as a device i64 state while the same block read it through a HOST pointer. W5k settled both against the RUNNING lane pin — transformers 5.16.0, `modeling_qwen4_exp.py` sha256 `77fec77d…c459`, confirmed by regenerating this row's committed forward golden byte-identically. Upstream types each cache slot from the tensor that first reaches it (`cache_utils.py:1019-1023`), so the ring carries the MODEL dtype and the history lives on the compute device: the PUBLISHER was right twice and both requirements moved to the block. `ModelRegistry::Forward` then runs a prefill and a `past_len > 0` DECODE over one set of persistent caches. **W5L drives the engine itself.** A real `GPUModelRunner` allocates all three published groups, gathers every group's block table, publishes the five-name by-name index and runs `execute_model` / `sample_tokens` for a prefill and then a decode; `LoadedEngine::FromModelDir` loads a `qwen4exp` GGUF and `generate` returns tokens; and `examples/server` answers `POST /v1/completions` on CPU. The cross-step gate is the PLE n-gram history read out of the RUNNER's own recurrent state at the slot the runner assigned — int64 token ids, which cannot saturate as this fixture's bf16 activations do. **WHAT SERVES IS EXACTLY THIS AND NO MORE: `--device cpu`, ONE SEQUENCE AT A TIME, over a GGUF.** `num_reqs > 1` is refused by name — `RunQwen4ExpQsaBlockPaged` takes a block table of one sequence — and because an EngineCore that meets that refusal dies rather than degrades, this factory sets `ModelFactory::serves_one_sequence_per_step` and `LoadedEngine::ResolveMaxNumSeqs` clamps `--max-num-seqs` to 1 and says so on stderr; concurrent clients are accepted and served in sequence. MEASURED before that clamp existed: three overlapping `/v1/completions` calls at `--max-num-seqs 4` each returned a 500 carrying this hook's own message, and the engine never served again. The quant arms are the loader's — IQ4_NL, Q5_0 and the dequantizing gather (#1989), whose CUDA arm landed with KGATHER. **ALL SIX `qwen4_exp` ops PLUS `vt::RmsNormGroup` NOW HAVE CUDA ARMS** (W6-CUDA and W6-CUDA-B, #2031) and this row's earlier sentence that "no CUDA arm exists for any `qwen4_exp` op" is false and is replaced. **THAT LEAVES THE GATHER, AND KGATHER LANDED IT**, so the sentence this row carried — that the one remaining reason is `EmbeddingKernelCuda` refusing a block-quantized table — is FALSE and is replaced. `vt::Embedding` on a CUDA queue decodes a block row across all 18 encodings `vt::cpu::BlockToFloat` decodes, measured bit-exact against the CPU arm on a GPU. **WITH BOTH LANDED, WHAT BLOCKS A CUDA FORWARD IS NEITHER OP REGISTRATION NOR THE LOADER, AND THIS IS A PREDICTION RATHER THAN A MEASUREMENT:** the expected shape is partial dispatch through the PLE, then a NAMED REFUSAL at the first QSA layer, because `qwen4_exp_qsa_block.cpp` still reads three operands on the HOST — `CheckRopeLayoutsAgree`, `IndexerRows` on the block table, and `Qwen4ExpQsaIndex` on `kv_lens` — which is owned by the QSADEV wave and NOT by this row. A second wall the synthetic fixture never reaches: `IsCudaKeepQuantSupported` still excludes IQ4_NL and Q5_0, which the released UD-IQ1_S uses, owned by [#2423](https://github.com/mudler/vllm.cpp/issues/2423). **A GPU HAS NOW PRODUCED TOKENS FOR THIS MODEL, ON ROCm, AND THE CLAIM IS LIVENESS AND NOT CORRECTNESS.** Measured 2026-09-13 on `strix:gpu0` (gfx1151, Radeon 8060S, ROCm 7.2.4) inside an `rc` lease ([spec](../.agents/specs/rocm-chunked-pinned-h2d.md)): the released `unsloth/Qwen3.8-Flash-Next-GGUF` UD-IQ1_S loads through `--device auto` and `examples/vllm-cli` returns 32 tokens with `finish_reason=length` on three of three launches -- fluent, coherent, prompt-dependent output, byte-identical across the three. Steady-state decode is **5.0-5.3 tok/s** -- nine samples across six independent process launches, 5.002 to 5.291 tok/s, 5.8% max-to-min. It is a RANGE: an earlier two-launch reading quoted 5.27-5.29 at a 0.34% spread and that precision does not reproduce. **IT IS NOT A TOKEN GATE AND IT IS NOT A PARITY CLAIM.** This architecture has no GPU oracle at all: llama.cpp aborts in `build_delta_net_chunking` before it reads a byte and no vLLM revision implements `qwen4_exp`, so nothing decoded these prompts beside us, there is no denominator, and no speed comparison is admissible. What reaches a token is the bounded pinned host-to-device ring; the same binary with `VT_ROCM_PINNED_H2D_MIB=0` stops at 29.69 GiB of device memory and produces nothing in 1200 s. **NO TOKEN HAS COME OUT OF A CUDA DEVICE FOR THIS MODEL and none is claimed** -- that half of the previous sentence is unchanged, and it is owed by the QSADEV host-operand wave and [#2423](https://github.com/mudler/vllm.cpp/issues/2423), not by the ROCm row. **NO TOKEN NUMBER AND NO SPEED NUMBER**: everything above ran on a synthetic fixture whose weights are a deterministic ramp, and the safetensors arm refuses because every published safetensors artifact exceeds every device this project owns. **THE CLAUSE 'no published `qwen4exp` checkpoint has been served' IS NOW HALF FALSE AND IS REPLACED BY A MEASUREMENT.** The released `unsloth/Qwen3.8-Flash-Next-GGUF` UD-IQ1_S (67.564 GiB, 3 shards, 1224 tensors) was driven through `examples/server` on `thor:gpu0` on 2026-08-30 (`rc` job `0f188dd1`, [evidence](bench-evidence/qwen4exp-released-checkpoint-serve-20260830.md)): **it LOADED and the server LISTENED, in 4446 s at 69.206 GiB peak RSS, keeping every one of its nine encodings quantized -- and it produced ZERO TOKENS.** `POST /v1/completions` returned 500 because the forward refused the artifact by name: `vt: qwen4_exp_gated_residual: input_mix_weight_down must be float (f32/bf16 for outputs)`. The file stores all **194** hyper-connection mix weights as Q8_0, the loader correctly keeps them quantized, and `vt::Qwen4ExpGatedResidual` accepted only float; every arm of the synthetic fixture wrote those same tensors as F32, which is why every gate on this row was green and none of them could see it. **W5p REMOVED THAT REFUSAL AT ITS SOURCE** (#2031): `mix_down`, `mix_up` and `block_inject` now accept a block-quantized `[N,K]` weight and route through `vt::MatmulBT`, which dispatches the keep-quant GEMM `kMatmulBTQuant` -- mirroring llama.cpp, which merged this architecture on 2026-08-27 (`6c84c7d5d`, first tag `b10660`), declares all six of these projections `GGML_OP_MUL_MAT` and never dequantizes one. The ELEMENTWISE operands did not move: a block-typed `hc_*_norm` gamma is still refused by name, which is llama.cpp's own split (`GGML_OP_MUL` for the norm, with an explicit f32 cast where a file-typed weight meets an elementwise multiply). The synthetic fixture grew the arm that was missing (`FixtureOpts::hc_mix_q8_0`), and `ModelRegistry::Forward` runs a prefill and a second prompt over a Q8_0-mix file; restoring the old contract reds that case with the verbatim string above, which is what makes the reach measured. **W5q RE-RAN THE RELEASED CHECKPOINT THROUGH THE REPAIRED PATH, AND THE REFUSAL IS GONE WHILE THE OUTPUT IS DEGENERATE** ([evidence](bench-evidence/qwen4exp-released-checkpoint-serve-20260831.md)): on `thor:gpu0` `--device cpu`, staged to worker-local disk, the artifact loads in **61 s** (against 4446 s from the CIFS share) at `VmHWM` 73.935 GiB, a 5-token prefill and eight decode steps run with nothing thrown, and `POST /v1/completions` returns **HTTP 200** with 8 tokens where W5n got a 500. **But every one of those tokens is id 0 — `!` in this file's own vocabulary — and the answer is BYTE-IDENTICAL for two different prompts.** So the forward is degenerate and prompt-independent on the real weights, no usable token has yet come out of a published `qwen4exp` checkpoint, and the DECODES and SERVES claims at the head of this row remain true OF THE FIXTURE. **W5s THEN GOT REAL TOKENS OUT OF IT, AND THE CAUSE WAS THE REPACK MARKER** ([evidence](bench-evidence/qwen4exp-released-checkpoint-tokens-20260831.md)): on `origin/main` `52f7ccbfc`, which carries W5r as well as W5p, the same artifact on the same box answers `" Paris. Given this fact, what is"` and `" 100°C at sea level"` — two different prompts, two different prompt-dependent completions, eight distinct token ids none of them 0. W5q's tree predated W5r, so on `thor` (aarch64 i8mm, where `vt::cpu::QuantRepackActive()` is TRUE) `dense_attn::ResidentWeight` was still dropping the repack marker and `kMatmulBTQuant` read `block_q8_0x4` buffers as flat `q8_0` on every hyper-connection mix weight; a read-only per-stage probe puts a NaN in `stream.after_layer_0` (`nan=51200`) collapsing to an all-zero `LOGITS` row (`zero=248320`), and `argmax` over a row with no maximum returns index 0. Post-W5r that stage is `nan=0` and the logit row is `min -9.89818 max 15.7873` with argmax id 11751 = the `" Paris"` token; `VT_CPU_QUANT_REPACK=0` is byte-identical to the default, which is what a correct performance transform must be. **WHAT RUNS IS EXACTLY THIS: `--device cpu`, ONE SEQUENCE AT A TIME, the UD-IQ1_S GGUF arm, and no more.** **IT IS NOT A TOKEN GATE** — no oracle decoded these prompts, llama.cpp aborts in `build_delta_net_chunking` before loading a byte, the other six published quants are unrun, and there is no speed number. The repaired route is also a per-TOKEN matvec where llama.cpp batches the projection over the whole prefill; batching it is owed and unmeasured. The shipped GGUF is TEXT-ONLY (1224 tensors, no `v.blk.*`), so the multimodal arm has no artifact to load either. **CONFIG LAYER GATED as well.** The config resolves and validates against a RUNNING transformers 5.16.0 oracle (it imports without torch, so `validate_architecture` executes): a 39-case two-direction sweep agrees on 35 and differs on 4, all four being local guards stricter than upstream, never looser. All 15 upstream `validate_architecture` rejections are implemented and tabulated against their upstream line. The forward and the KV-cache spec REFUSE BY NAME, each naming the wave that owes it. vLLM implements `qwen4_exp` at NO revision, so the algorithm oracle is transformers **5.16.0** under an accepted lane exception; `gateable = no` because nothing published fits a fleet device — `Qwen/Qwen3.8-Flash-Next` is ~360 GB bf16, ~180 GB FP8, ~128 GB NVFP4 against ~119.6 GiB usable on GB10 | none, and no speed claim is admissible from this row until a token gate exists | +| `Qwen4ExpForConditionalGeneration` | GGUF (`qwen4exp`) — **LOADS, DECODES, AND SERVES ONE SEQUENCE AT A TIME ON `--device cpu`** (W5a, W5f, W5g, W5k, W5L, #2031) | **THE GGUF LOAD IS GATED, AND SO IS THE LAYER LOOP; WHAT DOES NOT EXIST IS A SECOND STEP.** A `qwen4exp` file reaches the architecture's own config builder through the GGUF dispatch, the registry resolves the class, and `load_weights` materializes the whole text tower — **on CPU **and, since KGATHER, on CUDA** ([spec](../.agents/specs/cuda-quant-gather.md)). The 51.2 G-parameter n-gram table would expand from 26.822 GiB of IQ4_NL to 95.368 GiB of bf16 on a device that cannot gather blocks, which the on-disk #1123 device-fit guard cannot see, so the load REFUSES BY NAME on such a device ahead of any tensor I/O (#2083). `DeviceQuantGatherSupported` is no longer a device list at all: it asks `OpRegistered(kEmbeddingQuant, dev)`, CPU and CUDA register that op, and METAL, VULKAN, ROCM and TENSTORRENT do not and are still refused ([#2394](https://github.com/mudler/vllm.cpp/issues/2394)). So this row's earlier clause that the gate 'is true for CPU alone' is FALSE and is replaced — every convert-time transform inverted (the `+1` fold on every norm gamma with `ssm_norm` the one exception, `ssm_a` back to `log(-x)`, and the V-head reorder on every Gated DeltaNet tensor), gated in both directions against a committed 1224-tensor manifest of the shipped `unsloth/Qwen3.8-Flash-Next-GGUF UD-IQ1_S` and value-wise against a synthetic file. `Qwen4ExpTextModel::Forward` now exists (W5f, #2336) and `ModelRegistry::Forward` reaches it: the 48-layer loop composes all four block seams and is gated END TO END against the lane-pinned transformers 5.16.0 oracle at a tiny config, max|diff| 0.00982 against a 0.03 bf16-vs-f32 bound, with seven mutations separating by 0.78 to 2.02. **ONE TOKEN NOW COMES OUT OF A PRODUCTION SEAM, and the claim is exactly that and no more** (W5g): on a model loaded by `ModelRegistry::Load` from a synthetic `qwen4exp` GGUF, `ModelRegistry::Forward` returns `[1, vocab]` f32 logits with every element finite and `vt::GreedyArgmax` samples an id from them, on CPU. It is a REACH and a SAMPLE, not a token gate — the fixture's weights are a deterministic ramp, so the id is compared against no reference, and the tower's arithmetic is gated separately by the oracle case above. W5g is what made the prefill complete at all: `Qwen4ExpPleLayout` derived the n-gram head vocabulary from a DEFAULTED `ngram_vocab_size_base` and refused when a file's stated sizes disagreed, which a `qwen4exp` GGUF cannot avoid because llama.cpp #27742's converter writes the resolved arrays and no base — so the check held for exactly one artifact in existence and refused every other file with correctly loaded weights. The stated set is now the authority for the layout, as `NgramTableRows` already treated it, and the cross-check runs only where the source stated the base. **THE ENGINE'S CACHE CHANNEL NOW REACHES THIS MODEL** (W5j, #2031, #2353). This row said the `multi_kv` channel "is refused for every model by `ModelRegistry::Forward`"; that guard is now a MODEL-DECLARED capability, `ModelFactory::consumes_multi_kv`, landed with its first consumer, and this architecture is that consumer. A step carrying all three published groups reaches the hook, which resolves every one of its five caches BY NAME through `MultiKvCacheIndex::Resolve` — including the recurrent members, which `ENG-MULTIKV-BYNAME` made addressable — and reads the QSA indexer side cache out of the engine's own group-2 pages through group 2's own gathered block table. The guard still refuses `DeepseekV4ForCausalLM`, `Glm5NextForConditionalGeneration` and every architecture that declares nothing, and clearing the bit drives that refusal red in the gate. **IT NOW DECODES, AND IT NOW SERVES** (W5k and W5L, #2031). This row said "it still decodes NO token, and the reason is now the MODEL and not the engine", and named a dtype and a residency: the recurrent group publishes the PLE conv ring at the model dtype while `RunQwen4ExpPleBlock` required f32, and publishes the n-gram token history as a device i64 state while the same block read it through a HOST pointer. W5k settled both against the RUNNING lane pin — transformers 5.16.0, `modeling_qwen4_exp.py` sha256 `77fec77d…c459`, confirmed by regenerating this row's committed forward golden byte-identically. Upstream types each cache slot from the tensor that first reaches it (`cache_utils.py:1019-1023`), so the ring carries the MODEL dtype and the history lives on the compute device: the PUBLISHER was right twice and both requirements moved to the block. `ModelRegistry::Forward` then runs a prefill and a `past_len > 0` DECODE over one set of persistent caches. **W5L drives the engine itself.** A real `GPUModelRunner` allocates all three published groups, gathers every group's block table, publishes the five-name by-name index and runs `execute_model` / `sample_tokens` for a prefill and then a decode; `LoadedEngine::FromModelDir` loads a `qwen4exp` GGUF and `generate` returns tokens; and `examples/server` answers `POST /v1/completions` on CPU. The cross-step gate is the PLE n-gram history read out of the RUNNER's own recurrent state at the slot the runner assigned — int64 token ids, which cannot saturate as this fixture's bf16 activations do. **WHAT SERVES IS EXACTLY THIS AND NO MORE: `--device cpu`, ONE SEQUENCE AT A TIME, over a GGUF.** `num_reqs > 1` is refused by name — `RunQwen4ExpQsaBlockPaged` takes a block table of one sequence — and because an EngineCore that meets that refusal dies rather than degrades, this factory sets `ModelFactory::serves_one_sequence_per_step` and `LoadedEngine::ResolveMaxNumSeqs` clamps `--max-num-seqs` to 1 and says so on stderr; concurrent clients are accepted and served in sequence. MEASURED before that clamp existed: three overlapping `/v1/completions` calls at `--max-num-seqs 4` each returned a 500 carrying this hook's own message, and the engine never served again. The quant arms are the loader's — IQ4_NL, Q5_0 and the dequantizing gather (#1989), whose CUDA arm landed with KGATHER. **ALL SIX `qwen4_exp` ops PLUS `vt::RmsNormGroup` NOW HAVE CUDA ARMS** (W6-CUDA and W6-CUDA-B, #2031) and this row's earlier sentence that "no CUDA arm exists for any `qwen4_exp` op" is false and is replaced. **THAT LEAVES THE GATHER, AND KGATHER LANDED IT**, so the sentence this row carried — that the one remaining reason is `EmbeddingKernelCuda` refusing a block-quantized table — is FALSE and is replaced. `vt::Embedding` on a CUDA queue decodes a block row across all 18 encodings `vt::cpu::BlockToFloat` decodes, measured bit-exact against the CPU arm on a GPU. **WITH BOTH LANDED, WHAT BLOCKS A CUDA FORWARD IS NEITHER OP REGISTRATION NOR THE LOADER, AND THIS IS A PREDICTION RATHER THAN A MEASUREMENT:** the expected shape is partial dispatch through the PLE, then a NAMED REFUSAL at the first QSA layer, because `qwen4_exp_qsa_block.cpp` still reads three operands on the HOST — `CheckRopeLayoutsAgree`, `IndexerRows` on the block table, and `Qwen4ExpQsaIndex` on `kv_lens` — which is owned by the QSADEV wave and NOT by this row. A second wall the synthetic fixture never reaches: `IsCudaKeepQuantSupported` still excludes IQ4_NL and Q5_0, which the released UD-IQ1_S uses, owned by [#2423](https://github.com/mudler/vllm.cpp/issues/2423). **A GPU HAS NOW PRODUCED TOKENS FOR THIS MODEL, ON ROCm, AND THE CLAIM IS LIVENESS AND NOT CORRECTNESS.** Measured 2026-09-13 on `strix:gpu0` (gfx1151, Radeon 8060S, ROCm 7.2.4) inside an `rc` lease ([spec](../.agents/specs/rocm-chunked-pinned-h2d.md)): the released `unsloth/Qwen3.8-Flash-Next-GGUF` UD-IQ1_S loads through `--device auto` and `examples/vllm-cli` returns 32 tokens with `finish_reason=length` on three of three launches -- fluent, coherent, prompt-dependent output, byte-identical across the three. Steady-state decode is **5.0-5.3 tok/s** -- nine samples across six independent process launches, 5.002 to 5.291 tok/s, 5.8% max-to-min. It is a RANGE: an earlier two-launch reading quoted 5.27-5.29 at a 0.34% spread and that precision does not reproduce. **IT IS NOT A TOKEN GATE AND IT IS NOT A PARITY CLAIM.** This architecture has no GPU oracle at all: llama.cpp aborts in `build_delta_net_chunking` before it reads a byte and no vLLM revision implements `qwen4_exp`, so nothing decoded these prompts beside us, there is no denominator, and no speed comparison is admissible. What reaches a token is the bounded pinned host-to-device ring; the same binary with `VT_ROCM_PINNED_H2D_MIB=0` stops at 29.69 GiB of device memory and produces nothing in 1200 s. **NO TOKEN HAS COME OUT OF A CUDA DEVICE FOR THIS MODEL and none is claimed** -- that half of the previous sentence is unchanged, and it is owed by the QSADEV host-operand wave and [#2423](https://github.com/mudler/vllm.cpp/issues/2423), not by the ROCm row. **NO TOKEN NUMBER AND NO SPEED NUMBER**: everything above ran on a synthetic fixture whose weights are a deterministic ramp, and the safetensors arm refuses because every published safetensors artifact exceeds every device this project owns. **THE CLAUSE 'no published `qwen4exp` checkpoint has been served' IS NOW HALF FALSE AND IS REPLACED BY A MEASUREMENT.** The released `unsloth/Qwen3.8-Flash-Next-GGUF` UD-IQ1_S (67.564 GiB, 3 shards, 1224 tensors) was driven through `examples/server` on `thor:gpu0` on 2026-08-30 (`rc` job `0f188dd1`, [evidence](bench-evidence/qwen4exp-released-checkpoint-serve-20260830.md)): **it LOADED and the server LISTENED, in 4446 s at 69.206 GiB peak RSS, keeping every one of its nine encodings quantized -- and it produced ZERO TOKENS.** `POST /v1/completions` returned 500 because the forward refused the artifact by name: `vt: qwen4_exp_gated_residual: input_mix_weight_down must be float (f32/bf16 for outputs)`. The file stores all **194** hyper-connection mix weights as Q8_0, the loader correctly keeps them quantized, and `vt::Qwen4ExpGatedResidual` accepted only float; every arm of the synthetic fixture wrote those same tensors as F32, which is why every gate on this row was green and none of them could see it. **W5p REMOVED THAT REFUSAL AT ITS SOURCE** (#2031): `mix_down`, `mix_up` and `block_inject` now accept a block-quantized `[N,K]` weight and route through `vt::MatmulBT`, which dispatches the keep-quant GEMM `kMatmulBTQuant` -- mirroring llama.cpp, which merged this architecture on 2026-08-27 (`6c84c7d5d`, first tag `b10660`), declares all six of these projections `GGML_OP_MUL_MAT` and never dequantizes one. The ELEMENTWISE operands did not move: a block-typed `hc_*_norm` gamma is still refused by name, which is llama.cpp's own split (`GGML_OP_MUL` for the norm, with an explicit f32 cast where a file-typed weight meets an elementwise multiply). The synthetic fixture grew the arm that was missing (`FixtureOpts::hc_mix_q8_0`), and `ModelRegistry::Forward` runs a prefill and a second prompt over a Q8_0-mix file; restoring the old contract reds that case with the verbatim string above, which is what makes the reach measured. **W5q RE-RAN THE RELEASED CHECKPOINT THROUGH THE REPAIRED PATH, AND THE REFUSAL IS GONE WHILE THE OUTPUT IS DEGENERATE** ([evidence](bench-evidence/qwen4exp-released-checkpoint-serve-20260831.md)): on `thor:gpu0` `--device cpu`, staged to worker-local disk, the artifact loads in **61 s** (against 4446 s from the CIFS share) at `VmHWM` 73.935 GiB, a 5-token prefill and eight decode steps run with nothing thrown, and `POST /v1/completions` returns **HTTP 200** with 8 tokens where W5n got a 500. **But every one of those tokens is id 0 — `!` in this file's own vocabulary — and the answer is BYTE-IDENTICAL for two different prompts.** So the forward is degenerate and prompt-independent on the real weights, no usable token has yet come out of a published `qwen4exp` checkpoint, and the DECODES and SERVES claims at the head of this row remain true OF THE FIXTURE. **W5s THEN GOT REAL TOKENS OUT OF IT, AND THE CAUSE WAS THE REPACK MARKER** ([evidence](bench-evidence/qwen4exp-released-checkpoint-tokens-20260831.md)): on `origin/main` `52f7ccbfc`, which carries W5r as well as W5p, the same artifact on the same box answers `" Paris. Given this fact, what is"` and `" 100°C at sea level"` — two different prompts, two different prompt-dependent completions, eight distinct token ids none of them 0. W5q's tree predated W5r, so on `thor` (aarch64 i8mm, where `vt::cpu::QuantRepackActive()` is TRUE) `dense_attn::ResidentWeight` was still dropping the repack marker and `kMatmulBTQuant` read `block_q8_0x4` buffers as flat `q8_0` on every hyper-connection mix weight; a read-only per-stage probe puts a NaN in `stream.after_layer_0` (`nan=51200`) collapsing to an all-zero `LOGITS` row (`zero=248320`), and `argmax` over a row with no maximum returns index 0. Post-W5r that stage is `nan=0` and the logit row is `min -9.89818 max 15.7873` with argmax id 11751 = the `" Paris"` token; `VT_CPU_QUANT_REPACK=0` is byte-identical to the default, which is what a correct performance transform must be. **WHAT RUNS ON A PUBLISHED CHECKPOINT IS EXACTLY THIS: the UD-IQ1_S GGUF arm, ONE SEQUENCE AT A TIME, on `--device cpu` -- and, for LIVENESS ONLY, on ROCm `--device auto` on one gfx1151 board as the paragraph above records -- and no more.** **IT IS NOT A TOKEN GATE** — no oracle decoded these prompts, llama.cpp aborts in `build_delta_net_chunking` before loading a byte, the other six published quants are unrun, and there is no speed number. The repaired route is also a per-TOKEN matvec where llama.cpp batches the projection over the whole prefill; batching it is owed and unmeasured. The shipped GGUF is TEXT-ONLY (1224 tensors, no `v.blk.*`), so the multimodal arm has no artifact to load either. **CONFIG LAYER GATED as well.** The config resolves and validates against a RUNNING transformers 5.16.0 oracle (it imports without torch, so `validate_architecture` executes): a 39-case two-direction sweep agrees on 35 and differs on 4, all four being local guards stricter than upstream, never looser. All 15 upstream `validate_architecture` rejections are implemented and tabulated against their upstream line. The forward and the KV-cache spec REFUSE BY NAME, each naming the wave that owes it. vLLM implements `qwen4_exp` at NO revision, so the algorithm oracle is transformers **5.16.0** under an accepted lane exception; `gateable = no` because nothing published fits a fleet device — `Qwen/Qwen3.8-Flash-Next` is ~360 GB bf16, ~180 GB FP8, ~128 GB NVFP4 against ~119.6 GiB usable on GB10 | none, and no speed claim is admissible from this row until a token gate exists | | `Qwen3_5ForCausalLM`, `Qwen3_5MoeForCausalLM` | none: no text-only Qwen3.5 checkpoint fits this hardware | **NO RUN GATE, OWED.** Gated on `test_qwen3_8_text_only.cpp`; NO token claim. Loader reads stacked BF16 experts (#740) plus BF16 towers, shared expert and `lm_head` (#864), so both published indices satisfy the load plan | not measured | | `Qwen3ForCausalLM` | Qwen3 dense 0.6B/1.7B/4B/32B, NVFP4A16 | near-tie strict 16/16 vs vLLM 0.25.0 | c1 every-axis parity, c8 decode residual | | `Qwen3MoeForCausalLM` | Qwen3-Coder-30B-A3B | strict 6/6 vs vLLM 0.25.0 | 11/16 grid cells at or above graphed vLLM | diff --git a/include/vt/rocm/rocm_pinned_h2d.h b/include/vt/rocm/rocm_pinned_h2d.h index 8fc4bf4fe..7a6ff1125 100644 --- a/include/vt/rocm/rocm_pinned_h2d.h +++ b/include/vt/rocm/rocm_pinned_h2d.h @@ -79,8 +79,12 @@ struct StagedH2DInputs { // So the production path asks THIS first and calls the allocator only when it // passes. `ShouldStageH2D` stays the single authority on the decision and the // thing the truth table gates; this is that expression with `ring_available` -// held true, and a case in tests/vt/test_rocm_pinned_h2d.cpp asserts the two -// agree over the whole table so the split cannot drift. +// held true. A case in tests/vt/test_rocm_pinned_h2d.cpp walks 480 inputs and +// checks BOTH expressions against an expectation spelled out in that file from +// the four inputs, so a term deleted from either one fails there. Checking the +// two against EACH OTHER would not: that is `X == (X && true)`, true for any +// definition of this helper, and it is what an earlier version of that case +// did. constexpr bool StagingTermsExceptRing(const StagedH2DInputs& in) { return in.chunk_bytes != 0 && !in.stream_capturing && in.dst == PtrKind::kDevice && in.src == PtrKind::kUnregisteredHost && diff --git a/tests/vt/test_rocm_pinned_h2d.cpp b/tests/vt/test_rocm_pinned_h2d.cpp index 6739e5b58..219edbb60 100644 --- a/tests/vt/test_rocm_pinned_h2d.cpp +++ b/tests/vt/test_rocm_pinned_h2d.cpp @@ -307,11 +307,35 @@ TEST_CASE("the ring size and chunk size are llama.cpp's, and bound the residency // Production cannot evaluate `ring_available` without allocating 256 MiB of // pinned host memory, so it asks StagingTermsExceptRing first and calls // EnsureRing only when that passes. That is two expressions where the spec -// describes one decision, and two expressions drift. This case is what stops -// them: over the whole input space this file's truth table walks, -// StagingTermsExceptRing must be exactly ShouldStageH2D with `ring_available` -// held true -- no more, no less. Deleting a term from either one fails here. +// describes one decision, and two expressions drift. +// +// This case walks 480 inputs against an expectation computed HERE, from the +// four inputs, as a sequence of refusals -- never by calling either expression +// under test. An earlier version of this case asserted +// `StagingTermsExceptRing(in) == ShouldStageH2D(in)` with `ring_available` +// held true, which is `X == (X && true)`: true for ANY definition of the +// helper, including a deleted term. It convicted nothing. Measured: deleting +// `dst == kDevice`, and separately `bytes >= chunk_bytes`, from the helper left +// that version 1443/1443 SUCCESS. The independent expectation below is what +// makes a deleted term fail HERE and not only in the truth table above. // --------------------------------------------------------------------------- +namespace { + +// The four cheap terms, written out independently of the header. A refusal +// sequence rather than a conjunction, so this is not a copy of the expression +// it checks. +bool ExpectedCheapTerms(PtrKind src, PtrKind dst, size_t bytes, + size_t chunk_bytes, bool capturing) { + if (chunk_bytes == 0) return false; + if (capturing) return false; + if (dst != PtrKind::kDevice) return false; + if (src != PtrKind::kUnregisteredHost) return false; + if (bytes < chunk_bytes) return false; + return true; +} + +} // namespace + TEST_CASE("the cheap terms are exactly the decision minus the allocating term") { const PtrKind kinds[] = {PtrKind::kUnregisteredHost, PtrKind::kPinnedHost, PtrKind::kDevice, PtrKind::kOther}; @@ -332,26 +356,30 @@ TEST_CASE("the cheap terms are exactly the decision minus the allocating term") in.bytes = b; in.stream_capturing = cap; - // The ring is the ONLY term the split holds back. + const bool want = ExpectedCheapTerms(s, d, b, c, cap); + + // The cheap terms are the four, independently of the ring. in.ring_available = true; - REQUIRE(StagingTermsExceptRing(in) == ShouldStageH2D(in)); - if (ShouldStageH2D(in)) ++staged; + REQUIRE(StagingTermsExceptRing(in) == want); + // And the decision is those four AND the ring, so with a ring it + // is the same answer -- proved against `want`, not against the + // helper, so a term deleted from EITHER expression fails here. + REQUIRE(ShouldStageH2D(in) == want); + if (want) ++staged; - // And with no ring, the decision is always no, while the cheap - // terms are unmoved -- which is the whole point: production learns - // the answer is no WITHOUT paying for the ring to find out. + // With no ring, the decision is always no, while the cheap terms + // are unmoved -- which is the whole point: production learns the + // answer is no WITHOUT paying for the ring to find out. in.ring_available = false; CHECK_FALSE(ShouldStageH2D(in)); - in.ring_available = true; - CHECK(StagingTermsExceptRing(in) == ShouldStageH2D(in)); + CHECK(StagingTermsExceptRing(in) == want); ++total; } } } } } - // The table is not degenerate: some rows stage and most do not. A helper that - // returned a constant would satisfy the equality above and fail here. + // The table is not degenerate: some rows stage and most do not. CHECK(total == 480); CHECK(staged > 0); CHECK(staged < total);