diff --git a/.agents/NOW.md b/.agents/NOW.md index f678416e6..db3c38fdc 100644 --- a/.agents/NOW.md +++ b/.agents/NOW.md @@ -26,7 +26,7 @@ Work: exact-chunks on main `1ce0d662b`; sm_120 measured at `3d2581551`. | CPU levers (`QUANT-GGUF-CIQ-GEMM`) | Profile DONE: decode **47% threadpool sync**, prefill **~39% paged attn**. **G5 not next** | Parakeet encoder; attn dtype hoist | | Supported-models list | **LANDED**: FEATURES arch table CI-bound (33 archs) | — | | `/v1/videos` OpenAI shape | **MERGED** (#71): Sora `model`/`size`/`seconds` + `GET /{id}/content` | `row/SERVE-VIDEOS-REFS` PR open: reference conditioning | -| Vulkan 27B decode | **MET: 4.36 vs llama.cpp 4.35** (7 legs, main). Barriers -19.8%/tok, GPU -1.09 ms, e2e 8/12, OFF | Re-measure before flipping `VT_VULKAN_SMART_BARRIERS` | +| Vulkan 27B | decode **MET 4.36 vs 4.35** (barriers OFF). **LOADMEM: load held the model TWICE, VmRSS 100.759 -> 53.413 GiB** | Load-phase host build is the new peak | | `BACKEND-ROCM` | **(b) fix in; #140 gfx1201 hipBLAS + Gemma-4 MoE landed (contributor, authorship-preserved); W0 green 4 archs** | compile + M2 ([spec](specs/rocm-unified-memory-b.md)) | | TP spike #287 (PR #143) | **TP-W1 LANDED**: rank-group table + TP handle (6/6); DSR leak FIXED (unblocks #127/#154/#155) | TP-W2 (linears + loader) | | Release | **ACTIVE; W5 19/19+10/10; contract 30/30** | #141; artifacts pending | diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md index 161a86f04..045823061 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -16957,3 +16957,111 @@ draining. The hazard sets hold raw handle VALUES and are never dereferenced, so destroyed handle in them is harmless (handle reuse can only manufacture an extra barrier). (3) `docs/ENVIRONMENT.md` documents `VT_VULKAN_RING` as defaulting to 128; `kDescriptorRing` is 256. + +## BACKEND-VULKAN-LOADMEM — the 27B Vulkan load held the model TWICE; 100.759 -> 53.413 GiB VmRSS, device bytes unchanged (2026-08-09, GB10, `row/BACKEND-VULKAN-LOADMEM`) + +**Base:** `2b08dd24`. Build `-DVLLM_CPP_VULKAN=ON`, `CMAKE_CUDA_COMPILER:NOTFOUND`, +Release, at `~/vkloadmem/build-vk` on dgx.casa; source md5-verified against the +worktree after `git archive`. All GPU work under `flock $HOME/gpu.lock`, every +model load behind a MemAvailable guard plus a kill-the-process watchdog. + +### 0. `VLLM_CPP_DEVICE` IS NOT READ ANYWHERE, and the Vulkan gate has been selected by accident + +`grep -rn VLLM_CPP_DEVICE src/ include/ tests/ docs/` returns NOTHING. The engine +picks its device in `src/vllm/entrypoints/model_loader.cpp:81` via +`CurrentPlatform()`, which walks `{kCUDA, kXPU, kVULKAN, kMETAL, kCPU}` and takes +the first backend that probed a device. So `VLLM_CPP_DEVICE=vulkan +test_opt_paged_engine` — the command recorded for this gate throughout the Vulkan +campaign — selects Vulkan only because those builds had no CUDA compiler. +MEASURED both ways on the same source: with `/usr/local/cuda/bin` on `PATH` at +configure time the identical command reports `the engine selected device type 1` +(kCUDA) and passes 6/6; without it, `device type 3` (kVULKAN), 6/6, 0 declines. +`~/vkbar/build-vk/CMakeCache.txt` carries `CMAKE_CUDA_COMPILER:FILEPATH=NOTFOUND`, +which is why the earlier rows were genuinely on Vulkan. The env var is a placebo +and the gate needs a real selector. + +### 1. The attribution: a flat second copy of the model, not allocator excess + +`VT_VULKAN_ALLOC_STATS=1` (added by this row) prints, on every 1 GiB high-water +crossing, the caller-REQUESTED bytes, the driver-COMMITTED bytes +(`VkMemoryRequirements::size`), live buffer count, and the `/proc` context. + +Qwen3.6-27B bf16 (50.89 GiB on disk), `VT_ADOPT_DEVICE_BYTES` A/B on ONE binary: + +| arm | Vulkan live | buffers | VmRSS / VmHWM | MemAvailable floor | outcome | +|---|---|---|---|---|---| +| OFF (old behaviour) | 50.755 GiB | 863 of 894 | **100.759 GiB** | 13.85 GiB (MemFree 1.13) | watchdog KILLED it, still allocating | +| ON (default) | 50.756 GiB | 894 | **53.413 GiB** | 47.33 GiB | completed, TTFT 18.44 s | + +Qwen3-4B bf16 (7.6 GiB): 8.622 GiB Vulkan / 375 buffers in BOTH arms; VmHWM +16.392 -> 9.607 GiB; machine-wide cost 17.1 -> 9.45 GiB. + +Three things this rules out, with numbers rather than reasoning: + +* **No allocator excess.** `requested == committed` EXACTLY in every run + (50.756 GiB over 894 buffers at 27B, 8.622 over 375 at 4B). GB10's driver adds + no per-allocation rounding at these sizes, so there is nothing to win by + suballocating, and `maxMemoryAllocationCount` is nowhere near 894. +* **No transient held too long.** The excess is a FLAT offset, not a spike: on the + 27B OFF arm the RSS-minus-Vulkan gap reads 50.003 / 50.003 / 50.003 / 50.004 GiB + across the last four high-water lines. A staging buffer, a dequant scratch or a + duplicated upload would show as a bump, not a constant equal to the model. +* **The windowed source-page release DOES fire.** During load, RSS holds ONE copy + of the copied bytes, not the mmap as well. + +The offset is the host `OwnedTensor.bytes` mirror. `ResidentWeight` +(`dense_attn_block.h`, and its twin in `qwen3_5.cpp`) uploaded each weight and +kept the host buffer. On a discrete GPU that mirror is free; the Vulkan backend +allocates every buffer `HOST_VISIBLE|HOST_COHERENT` and persistently mapped, so +on GB10 both copies are the same 119.6 GiB of system RAM. +`src/vllm/platforms/vulkan.cpp` had reasoned that "there is exactly one copy of +the bytes"; `dense_attn_block.h` made a second one. + +### 2. Page cache is NOT the OOM cause, and MemFree is the wrong instrument + +Sampled at 2 Hz through the 27B load, `Cached` tracks the copied bytes 1:1 and is +never dropped: `MADV_DONTNEED` on a private file mapping drops the mapping's pages +but leaves the page-cache pages, which needs `POSIX_FADV_DONTNEED` on the fd. It +is reclaimable, so it does not consume MemAvailable — a first watchdog keyed on +MemFree killed a perfectly healthy fixed-arm load at MemFree 11.46 GiB while +MemAvailable was still 60.65 GiB. Recorded as a NEGATIVE for anyone tempted to +guard on MemFree: on a 51 GiB checkpoint, MemFree measures the read, not the risk. + +### 3. Why it took the machine down rather than failing + +At the kill point the OFF arm had 1.13 GiB of MemFree and 13.85 GiB of +MemAvailable with 31 buffers still to allocate, on a box whose Vulkan heap +(89.72 GiB) was never the binding constraint — the MACHINE was. The recorded +reboots ran repeated cold 27B loads with `drop_caches`, and the box also normally +runs `local-ai-worker`, whose vLLM reserves HOST RAM on this unified machine. Two +copies of a 50.89 GiB model plus any other tenant does not fit, and the NVIDIA +driver reports that as `NV_ERR_NO_MEMORY` from `_memdescAllocInternal`. + +### 4. The fix, and what it is not + +`AdoptDeviceBytesAsHost` re-points `bytes` AT the device allocation through the +existing `OwnedBytes` borrow, keyed alive by `d_dev`'s own control block. It is an +ADOPTION, not a release: the bytes survive at the device address, so every +`.bytes` reader (the f32 upcast, the portable CPU reference tier, `View`/`Numel`) +reads the same bytes from the surviving copy, and unlike `ReleaseHost` it needs no +"is the device path committed" proof. Gated on the new +`Backend::DeviceMemoryIsHostAddressable()`, default FALSE, which is deliberately +narrower than `UnifiedMemory()`: CUDA on GB10 is unified yet a `cudaMalloc` +pointer is not host-dereferenceable. + +### 5. Gates, and the mutations that prove the test + +GB10, Vulkan build: `test_opt_paged_engine` **6/6 prompts token-exact (96/96 +tokens), 0 declines, device type 3**; `test_backend_cross_device` **11/11 (132)**; +`test_vulkan_backend` **35/35 (2650/2650)**; `gen-vulkan-spirv.py --check` clean. +The three mechanism cases in `test_qwen36_weights` are 3/3 (20) and go RED under +each mutation: a no-op body (3 assertions), a dropped host-addressable guard +(3), and a dropped keep-alive (2, including a genuine use-after-free that read +`128` where `167` was written). + +### 6. Left open + +Peak is now the LOAD phase — the host `OwnedTensor` build reaches ~51 GiB before +the first upload, and the adoption only acts afterwards. Copying from the mmap +straight into the device buffer at load would cut that too. The unreleased page +cache is a second, independent lever. diff --git a/.agents/roadmap_v1.md b/.agents/roadmap_v1.md index 11e261266..5a9e19a3f 100644 --- a/.agents/roadmap_v1.md +++ b/.agents/roadmap_v1.md @@ -36,6 +36,7 @@ issue is not yet placed. Keyed record: update in place, never append. | Issue | Row | Title | Kind | |---:|---|---|---| +| [#203](https://github.com/mudler/vllm.cpp/issues/203) | `BACKEND-VULKAN` | Vulkan on unified memory holds TWO copies of the weights: 27B peaks at 100.8 GiB RSS and OOM-reboots a Spark | bug | | [#201](https://github.com/mudler/vllm.cpp/issues/201) | `BACKEND-ROCM` | `hipblasGemmEx` overload mismatch in `rocm_matmul_hipblaslt.hip` | bug | | [#199](https://github.com/mudler/vllm.cpp/issues/199) | `BACKEND-METAL-MLX` | macOS MLX build fails on `-Werror` in MLX headers | bug | | [#193](https://github.com/mudler/vllm.cpp/issues/193) | — | A100 (sm_80): crashes and wrong GDN output in fast paths | bug | diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index fd6150144..44fb97f78 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -362,6 +362,7 @@ built on it rather than keeping the flattering one. | MiniMax-H3 encoder quantization (`H3-ENC-BF16-COND-DIFF`) | **Measured Thor (`d1085374`).** Q4_K_M vs bf16 encoder, same 233-token prompt, same forward: rel RMS **0.0340** (0.0685 excl. sink), per-token cosine mean **0.99745** / min 0.909, median rotation **3.5°** | NOT a scale change (best rescale 0.0340->0.0328). Same energy as a ONE-WORD prompt edit but DIFFUSE: 232/233 tokens rotate vs 172/233 untouched. Render A/B owed. Detail: benchmark-record | | MXFP4 Qwen3-8B (W4A16 Marlin) | **`KERNEL-MARLIN-DENSE-EXEC` x3 (dense-ON default): c1 1.020, c2/c4/c8 0.962/0.966/0.969, GPU mem 2.63x less** (beats #51 1.005/0.925/0.939/0.953 EVERY axis); #44 3/3, 32B-NVFP4A16 6/6; -Werror test-guard fixes x2 | **VT_MARLIN_DENSE default-ON** (+951us). `FLASH-PTXAS` #82: cuModule A/B ties our+vLLM PTX across ptxas 13.0/13.2/driver-JIT (~144us); +10us is engine CONTEXT not codegen, no ptxas lever/flip (retires #75) | | Vulkan vs llama.cpp Vulkan (`BENCH-VK-LLAMA`) | **Both arms measured, same weights.** 0.6B @128-in/32-out: llama.cpp Vulkan **11,956** pp / **174.8** tg; ours **575** pp / **66.6** tg | Decode **8.59 -> 91.7 t/s** (**10.7x**), 6/6 exact; **2.62x** off llama.cpp at matched shape. CUDA arm unblocked. 27B: fallbacks **11->5**. paged_attn batching REFUTED ([plan](../.agents/specs/bench-27b-five-way.md)) | +| Vulkan load memory (`BACKEND-VULKAN-LOADMEM`) | **The load held the model TWICE.** 27B bf16, GB10, `VT_ADOPT_DEVICE_BYTES` A/B: **VmRSS 100.759 -> 53.413 GiB**, MemAvailable floor 13.85 -> 47.3 of 119.6. Device bytes identical. [Detail](../.agents/benchmark-record.md) | Load-phase peak (host build), and the page cache that tracks copied bytes 1:1 | | Memory footprint vs declared workload (`ROAD-V1-MEM`, #83) | **Never measured, and not measurable today**: there is no auto-sizing to compare against, because the KV pool is a hand-typed `--num-blocks`, so "what the run actually needed" has no number | Once M1's `MemoryBudget` lands: predicted-vs-actual bytes per allocation class, then peak footprint ours-auto vs vLLM at its 0.9 default on the same model and config | | Startup latency (cold to first `/health`) | **36.51 s vs vLLM 0.25.0's 221.51 s = 6.07x** (medians of 3, 27B-NVFP4, GB10). PROVISIONAL: 3 of 6 legs contended, repeat killed by a host reboot. [Detail](../.agents/benchmark-record.md) | Uncontended 3-rep re-run on a quiet box | | Speculation depth (`ROAD-V1-D3-SPEC-K`, #81) | **Never measured, MTP is k=1** (our port covers vLLM's k=1 branch only), so no acceptance-vs-depth curve exists | k=2..4 three-way greedy gate, then the c1/c>1 A/B + the per-workload (prose vs code) acceptance-vs-depth curve any dynamic or adaptive depth policy needs | diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index 632056aba..8b804a56d 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -105,6 +105,8 @@ portable/reference path. In normal operation leave them unset. | `VT_DFLASH_ATTN_WARP` | off (CUDA) | `=1` falls back to the older per-key warp-reduction block-attention kernel instead of the default chunked reduce-scatter form. Kept for the same-binary A/B that recorded the verdict | | `VT_DFLASH_ATTN_KEYLANE` | off (CUDA) | `=1` selects the one-key-per-lane block-attention form. **MEASURED NEGATIVE and not a tuning knob:** 28.90 s/step against the per-key warp kernel's 18.73 on the same binary (sm_110, MiniMax-H3 512x512/33f, seq 3224), 54% slower, because giving each lane a whole K row makes every K load 32-way scattered. Kept only because it is the experiment that located the real constraint | +| `VT_ADOPT_DEVICE_BYTES` | on (only acts where the backend advertises host-addressable device memory — Vulkan today) | After a weight is uploaded, re-point its host buffer AT the device allocation instead of keeping a second copy. On a unified box the two copies come out of the same RAM. MEASURED on GB10, Qwen3.6-27B bf16 (50.89 GiB): with the mirror the process reaches **VmRSS 100.759 GiB** and drives the machine to MemAvailable 13.85 / MemFree 1.13 GiB of 119.6 GiB before it has even finished allocating; without it the same load completes at **VmHWM 53.413 GiB**. Qwen3-4B: **16.392 -> 9.607 GiB**. Vulkan allocation is byte-identical either way. `0` is the same-binary A/B back to the two-copy behaviour. It is an adoption, not a release — the bytes survive at the device address and every reader sees them — so tokens are unchanged either way (`test_opt_paged_engine` on Vulkan is 6/6 token-exact, 96/96, both arms). No effect on CUDA/CPU/Metal, whose backends do not advertise the property | +| `VT_VULKAN_ALLOC_STATS` | off | `=1` prints a device-memory line on every 1 GiB high-water crossing and a summary at exit: live buffer count, bytes REQUESTED by the caller, bytes COMMITTED by the driver (`VkMemoryRequirements::size`), peak live bytes, and the process/system context (`VmRSS`, `VmHWM`, `MemAvailable`, `Cached`) read from `/proc`. On a unified-memory device the Vulkan heap IS system RAM, so separating "the backend allocated it", "the process allocated it some other way" and "it is page cache" is the whole of a memory attribution. Diagnostic only; it changes no numerics. The counters themselves are always maintained (one relaxed atomic per allocation) and are readable from a test through `vt::vulkan::DeviceAllocStatsSnapshot()`. Vulkan-only | | `VT_VULKAN_DISPATCH_STATS` | off | `=1` traces every Vulkan compute submit to stderr (index, shader, workgroup count) BEFORE its fence wait, prints any wait over 200 ms, reports a running dispatch rate every 100 submits, and dumps a per-shader histogram at exit. Printing before the wait is what makes a HANG visible: a post-wait print never runs if the fence never signals, so the last line names the dispatch that hung. This is how the coopmat out-of-bounds load was found. Diagnostic only; it changes no numerics. Vulkan-only | | `VT_VULKAN_GEMV_UNROLL` | 4 | `=1` forces the un-unrolled decode GEMV body. Four independent accumulators keep four reads per lane in flight instead of one -- memory-level parallelism, not instruction count. It rides a specialization constant, so both arms are the same committed module and A/B in one binary. MEASURED **1.055x, 7 of 8 interleaved pairs**. Worth noting it measured 5/8 and was REVERTED earlier the same day: that test ran while the GPU was only 26% busy, where a 10% GEMV win moves e2e by 1.4% and is unresolvable against this box's noise. After the ring fix made the run GPU-bound the same code reads 7/8. A negative result is regime-dependent. Vulkan-only | | `VT_VULKAN_GEMV_PACK` | 2 | Load width for a 16-bit decode GEMV operand: `0` one element per load (2 B), `1` two through the buffer's 32-bit view (4 B), `2` four through a 64-bit view (8 B). Same bytes, same coalescing (32 lanes still cover 128 or 256 CONSECUTIVE bytes), so this cannot reduce DRAM traffic; it reduces LOAD INSTRUCTIONS. **MEASURED, and the two regimes disagree by 7x, which is the point.** In `benchmarks/vulkan_gemv_ab.cpp` over the 27B's own decode shapes (9 arms x 4 rotated passes, each arm paired against the width-0 baseline measured IN THE SAME PASS, because the box drifted 15.5% peak-to-peak between passes) width 2 reads **1.086x**. In REAL 27B decode it reads **1.012x** and **1.020x**, two GPU-timestamp two-length diffs (output-len 36 minus 4 over 32 tokens), moving `vt_matmul_vec` from **90.0% to 91.1%** of the 273 GB/s roof. The sweep re-reads one 356 MB buffer 320 times, so its DRAM rows and TLB stay hot and instruction issue is visible; decode streams 50 GB once per token, where DRAM is the whole story. e2e: **4.126 -> 4.157 tok/s**, 5 of 6 clean pairs, against a 0.88% clean-leg noise floor. Degrades a width at a time when K or an operand byte offset is not aligned to it, and declines entirely for an f32 operand, which is already one element per 32-bit word. Unlike the row count this axis DOES change the answer's low bits, because it repartitions K across lanes; the opt-125m STRICT gate (6/6 token-exact) is what clears it. Vulkan-only | diff --git a/docs/FEATURES.md b/docs/FEATURES.md index 41c484acb..201b368f8 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -217,11 +217,10 @@ Vulkan **runs a model end to end**: `opt-125m` greedy is STRICT token-exact, 6/6 prompts / 96/96 tokens vs the vLLM 0.25.0 oracle, all nine of that model's ops dispatched natively with **zero provider declines**. Qwen3.6-27B runs too, both GDN recurrences and the fused attention preamble native, its GDN state cache in -place, and its RMSNorm 1024-wide (a batch-1 row is ONE workgroup, so it was -occupancy-bound): **decode 4.24 tok/s vs llama.cpp's 4.35, prefill -21.5x** (GB10). Still partial at 25 native kernels plus 8 GDN, the rest on the -portable CPU tier (`kRopeCosSinCache` stays host-side, mirroring vLLM); -quant/MoE/MLA have none at all. +place, and its RMSNorm 1024-wide: **decode 4.24 tok/s vs llama.cpp's 4.35, +prefill 21.5x** (GB10). A load keeps **one** copy of the weights, not two: 27B +peak RSS 100.8 GiB before, **53.4 GiB** now. Still partial at 25 native kernels +plus 8 GDN, the rest on the portable CPU tier; quant/MoE/MLA have none at all. Build with `-DVLLM_CPP_VULKAN=ON`; off by default. ## Serving, API and operations diff --git a/docs/STATUS.md b/docs/STATUS.md index 467019725..5be0009f1 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -423,7 +423,7 @@ Parakeet ASR (2026-08-07): *CPU-correct, ON THE ONE SURFACE (ROW 1)*. Ids exact LoRA (W1 CPU runtime brick landed; not yet usable end-to-end), multi-GPU, Vulkan (opt-125m exact; 25 native +8 GDN, both recurrences + fused attn preamble; 27B prefill 21.5x, decode -4.36/4.35 MET; barriers -19.8%, GPU -1.09 ms; #125 +4.36/4.35 MET; 27B load 100.8 -> 53.4 GiB; #125 [campaign](../.agents/specs/vulkan-full-support.md)), ROCm (W0 community-green on 4 gfx archs (#41); the ratified (b) APU unified-memory fix is in (**blind-written, unverified**); M2 needs verification; gfx1201 hipBLAS + diff --git a/docs/USAGE.md b/docs/USAGE.md index 6701b5434..fbeb3e3b0 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -51,6 +51,17 @@ Two more example binaries ship alongside it: - `tokenize` ([`examples/tokenize/main.cpp`](../examples/tokenize/main.cpp)), a tokenizer smoke tool taking ` `. +### How much memory a Vulkan load needs + +On a unified-memory device (a DGX Spark) the Vulkan heap and system RAM are the +same bytes, so budget roughly **the checkpoint size plus about 5%**, plus your KV +pool. Measured on GB10: Qwen3.6-27B bf16 (50.89 GiB on disk) peaks at 53.4 GiB of +process RSS. Reading the checkpoint also fills the page cache with about the file +size; that is reclaimable and does not need to be budgeted, but it does make +`MemFree` look alarming during a load. Use `MemAvailable`, not `MemFree`, to +decide whether a model fits. `VT_VULKAN_ALLOC_STATS=1` prints the running device +total and the `/proc` context if you need to see where it goes. + A Vulkan build (`-DVLLM_CPP_VULKAN=ON`) adds three kernel-measurement binaries. They exist so a Vulkan tuning knob can be A/B'd in ONE binary, which is this project's benchmark protocol, and each one prints WHICH kernel variant it ran so diff --git a/include/vllm/model_executor/models/dense_attn_block.h b/include/vllm/model_executor/models/dense_attn_block.h index 6da5ecd6a..d9328a31a 100644 --- a/include/vllm/model_executor/models/dense_attn_block.h +++ b/include/vllm/model_executor/models/dense_attn_block.h @@ -193,6 +193,10 @@ inline Tensor ResidentWeight(Dev d, const OwnedTensor& w, std::vector s d.b.Copy(d.q, p, w.bytes.data(), nb); Backend* bk = &d.b; w.d_dev = std::shared_ptr(p, [bk](void* q) { bk->Free(q); }); + // The host mirror is now redundant wherever device memory is host- + // addressable (Vulkan). See AdoptDeviceBytesAsHost — this is what keeps a + // unified-memory box from holding the whole model twice. + AdoptDeviceBytesAsHost(d.b, w); } return MakeTensor(w.d_dev.get(), w.dtype, d.q.device, shape); } diff --git a/include/vllm/model_executor/models/qwen3_5_weights.h b/include/vllm/model_executor/models/qwen3_5_weights.h index c623dfff5..75a1771b9 100644 --- a/include/vllm/model_executor/models/qwen3_5_weights.h +++ b/include/vllm/model_executor/models/qwen3_5_weights.h @@ -32,6 +32,10 @@ #include "vt/dtype.h" #include "vt/tensor.h" +namespace vt { +class Backend; +} // namespace vt + namespace vllm { // One owned, contiguous, host tensor: heap bytes + shape/dtype. View() builds a @@ -105,6 +109,35 @@ struct OwnedTensor { mutable std::shared_ptr d_dev_f32; }; +// ADOPT the device-resident copy AS the host buffer, where the backend says its +// allocations are host-addressable (vt::Backend::DeviceMemoryIsHostAddressable). +// +// THE DEFECT THIS CLOSES (BACKEND-VULKAN-LOADMEM). `ResidentWeight` uploads a +// weight and keeps `bytes` as well, so on Vulkan the model became resident +// TWICE. That is invisible on a discrete GPU -- the two copies are in different +// memories -- but GB10 is unified, so both come out of the same 119 GiB of +// system RAM. MEASURED on Qwen3-4B (7.6 GiB on disk): 8.622 GiB of Vulkan +// allocation and 16.392 GiB of process VmHWM, i.e. a whole second copy, and +// 17.1 GiB off MemAvailable. Extrapolated to the 27B (50.89 GiB) that is over +// 100 GiB, which is why loading it could take the machine down rather than fail. +// `src/vllm/platforms/vulkan.cpp` even REASONED that there is "exactly one copy +// of the bytes"; this is what makes that true. +// +// It is an ADOPTION, not a free: `bytes` is re-pointed at the device allocation +// (persistently mapped, host-coherent) and keeps it alive through `d_dev`, so +// every existing `.bytes` reader -- `ResidentWeightF32`'s upcast, the portable +// CPU reference tier, `Numel`/`View` -- reads the SAME bytes it read before, +// from the surviving copy. Nothing is dropped that anyone could still want, so +// unlike `ReleaseHost()` this needs no "is the device path committed" proof. +// +// A no-op unless the backend opts in, and a no-op on an already-BORROWED buffer +// (a GGUF mmap or a tied-weight expansion): those own no anonymous pages, and a +// tied pair must keep sharing one keep-alive. +// +// `VT_ADOPT_DEVICE_BYTES=0` is the same-binary A/B back to the two-copy +// behavior (house convention for a default-on residency change). +void AdoptDeviceBytesAsHost(vt::Backend& backend, const OwnedTensor& w); + // Device-resident NVFP4 W4A16 weight (M2.2b). The modelopt packed fp4 codes + // fp8-e4m3 group scales + per-tensor scale, kept RAW in the ORIGINAL torch // [N=out_features, K=in_features] orientation vt::MatmulNvfp4 expects (NOT diff --git a/include/vt/backend.h b/include/vt/backend.h index 36e975324..b90c0626b 100644 --- a/include/vt/backend.h +++ b/include/vt/backend.h @@ -51,6 +51,30 @@ class Backend { // True when host and device share one memory space (CPU, GB10, Apple). virtual bool UnifiedMemory() const = 0; + // True when a pointer returned by Alloc() may be DEREFERENCED BY THE HOST + // directly -- loaded, stored, memcpy'd -- with no map/unmap call and no + // staging bounce. + // + // This is STRICTLY NARROWER than UnifiedMemory(), and the difference is the + // whole reason it exists. CUDA on GB10 reports unified memory because host and + // device address the same physical RAM, yet a plain `cudaMalloc` pointer is + // still not host-dereferenceable. Vulkan here allocates every buffer + // HOST_VISIBLE|HOST_COHERENT and keeps it persistently mapped, so its pointers + // are ordinary host memory that the GPU also reads -- which is already what + // this backend's Copy/Memset (plain memcpy/memset) and the portable CPU + // reference tier depend on. + // + // MEASURED consequence (BACKEND-VULKAN-LOADMEM): where this is true, a weight + // that has been uploaded needs NO host mirror, because the device allocation + // IS a host buffer. Keeping one costs a second full copy of the model -- + // 16.392 GiB of process RSS for a 7.6 GiB Qwen3-4B, against 8.622 GiB of + // Vulkan allocation -- and on a unified box that second copy comes out of the + // same RAM the first one does. + // + // Default false: a backend must OPT IN, because being wrong here hands a + // device pointer to a host memcpy and segfaults. + virtual bool DeviceMemoryIsHostAddressable() const { return false; } + // --- Device compute capability (BACKEND-CUDA-ARCH-ADDITIVITY seam-gap #4) --- // The architecture the backend is actually running on, as the familiar // `(major, minor)` pair (GB10/sm_121 -> {12, 1}). Before this, the capability diff --git a/src/vllm/model_executor/models/qwen3_5.cpp b/src/vllm/model_executor/models/qwen3_5.cpp index eb5f43173..31d5790da 100644 --- a/src/vllm/model_executor/models/qwen3_5.cpp +++ b/src/vllm/model_executor/models/qwen3_5.cpp @@ -851,6 +851,10 @@ Tensor ResidentWeight(Dev d, const OwnedTensor& w, std::vector shape = d.b.Copy(d.q, p, w.bytes.data(), nb); Backend* bk = &d.b; w.d_dev = std::shared_ptr(p, [bk](void* q) { bk->Free(q); }); + // Same adoption as the dense block's ResidentWeight: on a host-addressable + // device the uploaded buffer IS the host buffer, so keeping the mirror + // costs a second full copy of the model out of the same unified RAM. + AdoptDeviceBytesAsHost(d.b, w); } return MakeTensor(w.d_dev.get(), w.dtype, d.q.device, shape); } diff --git a/src/vllm/model_executor/models/qwen3_5_weights.cpp b/src/vllm/model_executor/models/qwen3_5_weights.cpp index d9b94e93a..4276867e2 100644 --- a/src/vllm/model_executor/models/qwen3_5_weights.cpp +++ b/src/vllm/model_executor/models/qwen3_5_weights.cpp @@ -17,6 +17,7 @@ #endif #include "vllm/model_executor/model_loader/nvfp4_dequant.h" +#include "vt/backend.h" #include "vt/dtype.h" namespace vllm { @@ -73,6 +74,47 @@ void OwnedTensor::ReleaseHost() const { self.host_released = true; } +void AdoptDeviceBytesAsHost(vt::Backend& backend, const OwnedTensor& w) { + if (!backend.DeviceMemoryIsHostAddressable()) return; + // A BORROWED buffer owns no anonymous pages (a GGUF mmap is clean and + // file-backed; a shared expansion is a tied pair's single copy), so adopting + // would reclaim nothing and would break the tie. Same reasoning as + // ReleaseHost's borrowed branch. + if (w.d_dev == nullptr || w.bytes.empty() || w.bytes.borrowed()) return; + if (const char* v = std::getenv("VT_ADOPT_DEVICE_BYTES"); v != nullptr && v[0] == '0') { + return; + } + auto& self = *const_cast(&w); + const size_t nb = self.bytes.size(); +#if defined(__unix__) || defined(__APPLE__) + // Drop the resident anonymous pages BEFORE the vector is destroyed, for + // exactly the reason ReleaseHost above does it: glibc raises its dynamic mmap + // threshold as large blocks are freed, so free() alone leaves many weight + // buffers on the sbrk arena free-list with their pages still resident, and + // the whole point here is the RSS. Interior whole pages only, so free()'s + // boundary metadata is untouched. + { + const long ps_l = ::sysconf(_SC_PAGESIZE); + const auto ps = static_cast(ps_l > 0 ? ps_l : 4096); + const auto begin = reinterpret_cast(self.bytes.data()); + const uintptr_t end = begin + nb; + const uintptr_t page_begin = (begin + ps - 1) & ~(ps - 1); + const uintptr_t page_end = end & ~(ps - 1); + if (page_end > page_begin) { + ::madvise(reinterpret_cast(page_begin), + static_cast(page_end - page_begin), MADV_DONTNEED); + } + } +#endif + // The keep-alive is the device allocation's own shared_ptr, so the borrowed + // view cannot outlive the bytes it points at: the aliasing constructor shares + // d_dev's control block, and the buffer is freed through the vt Backend only + // when BOTH the weight's d_dev and this view are gone. + std::shared_ptr keep(w.d_dev, static_cast(w.d_dev.get())); + self.bytes = OwnedBytes::Borrow(static_cast(w.d_dev.get()), nb, + std::move(keep)); +} + vt::Tensor OwnedTensor::View() const { VT_CHECK(!host_released, "OwnedTensor::View: host bytes were released after device upload"); diff --git a/src/vt/vulkan/vulkan_backend.cpp b/src/vt/vulkan/vulkan_backend.cpp index 24a57d9ec..8b3b65aa4 100644 --- a/src/vt/vulkan/vulkan_backend.cpp +++ b/src/vt/vulkan/vulkan_backend.cpp @@ -127,6 +127,13 @@ class VulkanBackend final : public Backend { bool UnifiedMemory() const override { return VulkanContext::Get().unified_memory(); } + // Every allocation is HOST_VISIBLE|HOST_COHERENT and persistently mapped by + // AllocBuffer, and Copy/Memset above are already a plain host memcpy/memset + // over exactly that pointer. So this is not a new claim -- it NAMES the + // property this backend has always relied on, so the weight loader can rely + // on it too instead of keeping a redundant host mirror. + bool DeviceMemoryIsHostAddressable() const override { return true; } + // vt_causal_conv1d_update binds the state through the dtype-erased 32/16-bit // view pair and rounds once on store, so a bf16 conv_state is read and written // in place. Without this the caller must gather the cache into an f32 working diff --git a/src/vt/vulkan/vulkan_context.cpp b/src/vt/vulkan/vulkan_context.cpp index 5821ca974..dc5d5b7c3 100644 --- a/src/vt/vulkan/vulkan_context.cpp +++ b/src/vt/vulkan/vulkan_context.cpp @@ -29,6 +29,7 @@ #include "vulkan_context.h" #include +#include #include #include #include @@ -46,6 +47,98 @@ namespace vt::vulkan { namespace { +// DEVICE-MEMORY ACCOUNTING, enabled by VT_VULKAN_ALLOC_STATS (BACKEND-VULKAN- +// LOADMEM). +// +// GB10 is a UNIFIED-memory box: the one Vulkan heap and the machine's RAM are +// the SAME 119 GiB, so a Vulkan allocation and a host allocation compete +// directly and an over-allocating load takes the whole machine down rather than +// failing cleanly (`NV_ERR_NO_MEMORY` out of `_memdescAllocInternal`, twice in +// one day). Attributing that needs three numbers that only this layer can +// supply: how many bytes the CALLER asked for, how many the DRIVER actually +// committed (`VkMemoryRequirements::size`, which is rounded up), and how many +// are LIVE at the moment the process peaks. Everything else -- RSS, page cache, +// MemAvailable -- is observable from outside with /proc. +// +// Cost when off: one relaxed atomic add per allocation, which is noise against +// a vkAllocateMemory. Cost when on: additionally a /proc read, but ONLY on a +// new high-water mark, which is O(heap/step) times over a whole run. +const bool kAllocStats = [] { + const char* v = std::getenv("VT_VULKAN_ALLOC_STATS"); + return v != nullptr && std::strcmp(v, "0") != 0; +}(); + +struct AllocAccounting { + std::atomic live_count{0}; + std::atomic total_count{0}; + std::atomic live_requested{0}; // caller bytes, before rounding + std::atomic live_allocated{0}; // VkMemoryRequirements::size + std::atomic total_requested{0}; + std::atomic total_allocated{0}; + std::atomic peak_allocated{0}; + std::atomic peak_count{0}; + std::atomic next_report{0}; // next high-water print threshold +}; + +AllocAccounting& Accounting() { + static AllocAccounting a; + return a; +} + +// Per-allocation sizes, so FreeBuffer can subtract exactly what AllocBuffer +// added. Keyed by the packed VkDeviceMemory, which is unique while it is live. +// A map plus a mutex is free at this frequency: every entry costs one +// vkAllocateMemory or vkFreeMemory, which is orders of magnitude dearer. +std::mutex& AllocSizeMutex() { + static std::mutex m; + return m; +} +std::map>& AllocSizes() { // {requested, allocated} + static std::map> m; + return m; +} + +// One /proc key, in KiB as the kernel reports it, or 0 when absent. Read on the +// slow path only. +uint64_t ProcKiB(const char* path, const char* key) { + std::FILE* f = std::fopen(path, "r"); + if (f == nullptr) return 0; + char line[256]; + const size_t klen = std::strlen(key); + uint64_t out = 0; + while (std::fgets(line, sizeof(line), f) != nullptr) { + if (std::strncmp(line, key, klen) == 0 && line[klen] == ':') { + out = std::strtoull(line + klen + 1, nullptr, 10); + break; + } + } + std::fclose(f); + return out; +} + +constexpr double kToGiB = 1.0 / (1024.0 * 1024.0 * 1024.0); + +// Prints the full picture at one instant: what Vulkan holds, what the process +// holds, and what the machine has left. The three together are the attribution +// -- a Vulkan-only number cannot tell a driver over-allocation apart from a host +// mirror of the same weights, and that distinction is the whole question. +void ReportAllocState(const char* why) { + const AllocAccounting& a = Accounting(); + std::fprintf( + stderr, + "[vt vulkan] alloc %-9s live=%llu bufs req=%.3f GiB alloc=%.3f GiB " + "peak=%.3f GiB | VmRSS=%.3f GiB VmHWM=%.3f GiB | MemAvail=%.3f GiB " + "Cached=%.3f GiB\n", + why, static_cast(a.live_count.load(std::memory_order_relaxed)), + static_cast(a.live_requested.load(std::memory_order_relaxed)) * kToGiB, + static_cast(a.live_allocated.load(std::memory_order_relaxed)) * kToGiB, + static_cast(a.peak_allocated.load(std::memory_order_relaxed)) * kToGiB, + static_cast(ProcKiB("/proc/self/status", "VmRSS")) * 1024.0 * kToGiB, + static_cast(ProcKiB("/proc/self/status", "VmHWM")) * 1024.0 * kToGiB, + static_cast(ProcKiB("/proc/meminfo", "MemAvailable")) * 1024.0 * kToGiB, + static_cast(ProcKiB("/proc/meminfo", "Cached")) * 1024.0 * kToGiB); +} + // Dispatch accounting, enabled by VT_VULKAN_DISPATCH_STATS (VK-E deep dive). const bool kDispatchStats = [] { const char* v = std::getenv("VT_VULKAN_DISPATCH_STATS"); @@ -864,6 +957,23 @@ VulkanContext::VulkanContext() { } } + // VT_VULKAN_ALLOC_STATS=1 dumps the device-memory summary at exit, for the + // same reason the dispatch histogram does: this context is never destroyed. + if (kAllocStats) { + std::atexit([] { + const AllocAccounting& a = Accounting(); + ReportAllocState("exit"); + std::fprintf(stderr, + "[vt vulkan] alloc lifetime allocations=%llu requested=%.3f GiB " + "committed=%.3f GiB peak_live=%.3f GiB peak_bufs=%llu\n", + static_cast(a.total_count.load()), + static_cast(a.total_requested.load()) * kToGiB, + static_cast(a.total_allocated.load()) * kToGiB, + static_cast(a.peak_allocated.load()) * kToGiB, + static_cast(a.peak_count.load())); + }); + } + // VT_VULKAN_DISPATCH_STATS=1 dumps the per-shader histogram at exit. Registered // with atexit rather than printed from a destructor because this context is a // never-destroyed process singleton, and because a run that is KILLED by a @@ -1008,6 +1118,48 @@ void* VulkanContext::AllocBuffer(size_t bytes, void** out_buffer, void** out_mem Check(vk.vkAllocateMemory(device, &mai, nullptr, &memory), "vkAllocateMemory"); Check(vk.vkBindBufferMemory(device, buffer, memory, 0), "vkBindBufferMemory"); + // Account AFTER the allocation succeeded, so a failed allocation never + // inflates the live total. `req.size` is what the driver committed; `len` is + // what we asked for. On GB10 they are equal for every size the model loader + // uses, and the gap -- if a driver ever introduces one -- is exactly the + // "allocated is bigger than the tensor bytes" term this accounting exists to + // separate from a host mirror. + { + { + std::lock_guard g(AllocSizeMutex()); + AllocSizes()[Pack(memory)] = {static_cast(len), static_cast(req.size)}; + } + AllocAccounting& a = Accounting(); + a.live_count.fetch_add(1, std::memory_order_relaxed); + a.total_count.fetch_add(1, std::memory_order_relaxed); + a.live_requested.fetch_add(len, std::memory_order_relaxed); + a.total_requested.fetch_add(len, std::memory_order_relaxed); + a.total_allocated.fetch_add(req.size, std::memory_order_relaxed); + const uint64_t live = + a.live_allocated.fetch_add(req.size, std::memory_order_relaxed) + req.size; + uint64_t peak = a.peak_allocated.load(std::memory_order_relaxed); + while (live > peak && + !a.peak_allocated.compare_exchange_weak(peak, live, std::memory_order_relaxed)) { + } + uint64_t pc = a.peak_count.load(std::memory_order_relaxed); + const uint64_t lc = a.live_count.load(std::memory_order_relaxed); + while (lc > pc && + !a.peak_count.compare_exchange_weak(pc, lc, std::memory_order_relaxed)) { + } + if (kAllocStats) { + // Report on each new 1 GiB high-water mark. A per-allocation line would be + // hundreds of thousands of lines on a 27B and would itself perturb the + // load; the high-water crossings are what a memory attribution needs. + constexpr uint64_t kStep = uint64_t{1} << 30; + uint64_t mark = a.next_report.load(std::memory_order_relaxed); + if (live >= mark && + a.next_report.compare_exchange_strong(mark, ((live / kStep) + 1) * kStep, + std::memory_order_relaxed)) { + ReportAllocState("high-water"); + } + } + } + void* mapped = nullptr; Check(vk.vkMapMemory(device, memory, 0, VK_WHOLE_SIZE, 0, &mapped), "vkMapMemory"); // vt::StepArena depends on >= 64-byte alignment (include/vt/backend.h:26). @@ -1025,11 +1177,41 @@ void* VulkanContext::AllocBuffer(size_t bytes, void** out_buffer, void** out_mem void VulkanContext::FreeBuffer(void* buffer, void* memory) { const VulkanApi& vk = Api(); auto device = Unpack(device_); + { + std::pair sizes{0, 0}; + { + std::lock_guard g(AllocSizeMutex()); + auto& m = AllocSizes(); + auto it = m.find(memory); + if (it != m.end()) { + sizes = it->second; + m.erase(it); + } + } + AllocAccounting& a = Accounting(); + a.live_count.fetch_sub(1, std::memory_order_relaxed); + a.live_requested.fetch_sub(sizes.first, std::memory_order_relaxed); + a.live_allocated.fetch_sub(sizes.second, std::memory_order_relaxed); + } vk.vkUnmapMemory(device, Unpack(memory)); vk.vkDestroyBuffer(device, Unpack(buffer), nullptr); vk.vkFreeMemory(device, Unpack(memory), nullptr); } +DeviceAllocStats DeviceAllocStatsSnapshot() { + const AllocAccounting& a = Accounting(); + DeviceAllocStats s; + s.live_count = a.live_count.load(std::memory_order_relaxed); + s.total_count = a.total_count.load(std::memory_order_relaxed); + s.live_requested = a.live_requested.load(std::memory_order_relaxed); + s.live_allocated = a.live_allocated.load(std::memory_order_relaxed); + s.total_requested = a.total_requested.load(std::memory_order_relaxed); + s.total_allocated = a.total_allocated.load(std::memory_order_relaxed); + s.peak_allocated = a.peak_allocated.load(std::memory_order_relaxed); + s.peak_count = a.peak_count.load(std::memory_order_relaxed); + return s; +} + namespace { // The pipeline cache key: the module name plus its specialization values, which diff --git a/src/vt/vulkan/vulkan_context.h b/src/vt/vulkan/vulkan_context.h index 8c80decf1..5af8c490d 100644 --- a/src/vt/vulkan/vulkan_context.h +++ b/src/vt/vulkan/vulkan_context.h @@ -35,6 +35,31 @@ namespace vt::vulkan { +// DEVICE-MEMORY ACCOUNTING (BACKEND-VULKAN-LOADMEM). Maintained unconditionally +// -- one relaxed atomic add per vkAllocateMemory -- because on a unified-memory +// device the Vulkan heap IS system RAM, so "how many bytes does this backend +// hold" is a question a test and a diagnostic both need to be able to ask at any +// instant, not only under an env flag. `VT_VULKAN_ALLOC_STATS=1` additionally +// prints a line on every 1 GiB high-water crossing and a summary at exit. +// +// `requested` is what the caller asked for (after the 4-byte rounding +// AllocBuffer applies for the 32-bit storage view); `allocated` is what the +// driver committed, `VkMemoryRequirements::size`. They are reported separately +// so a driver-side over-allocation is distinguishable from a caller that simply +// allocates too much -- the two have completely different fixes. +struct DeviceAllocStats { + uint64_t live_count = 0; + uint64_t total_count = 0; + uint64_t live_requested = 0; + uint64_t live_allocated = 0; + uint64_t total_requested = 0; + uint64_t total_allocated = 0; + uint64_t peak_allocated = 0; + uint64_t peak_count = 0; +}; + +DeviceAllocStats DeviceAllocStatsSnapshot(); + // Process-wide Vulkan context. Created on first use, never destroyed (the // process outlives it; matching llama.cpp's `vk_instance` singleton lifetime and // the Metal skeleton's MetalContext). diff --git a/tests/vllm/test_qwen36_weights.cpp b/tests/vllm/test_qwen36_weights.cpp index 4dd2439b7..b013c4f6d 100644 --- a/tests/vllm/test_qwen36_weights.cpp +++ b/tests/vllm/test_qwen36_weights.cpp @@ -406,6 +406,164 @@ TEST_CASE("OwnedTensor::ReleaseHost frees host bytes but preserves logical prese std::runtime_error); } +// --- Host-addressable device adoption: ONE copy of the bytes, not two -------- +// BACKEND-VULKAN-LOADMEM. On a backend whose allocations the host can +// dereference (Vulkan: every buffer HOST_VISIBLE|HOST_COHERENT and persistently +// mapped), a weight that has been uploaded must not ALSO keep a host mirror -- +// on GB10 both copies come out of the same unified RAM, so the model became +// resident twice (MEASURED: 16.392 GiB VmHWM against 8.622 GiB of Vulkan +// allocation for a 7.6 GiB Qwen3-4B). +// +// This pins the MECHANISM on the CPU tier, with no GPU and no checkpoint: what +// the buffer POINTS AT afterwards, that the bytes survive, that the device +// allocation is kept alive by the view, and -- the guard that keeps CUDA +// byte-identical -- that a backend which does NOT advertise host-addressable +// memory is left completely alone. +namespace { + +// Minimal Backend whose "device" memory is plain host memory, so adoption is +// observable without a real accelerator. `host_addressable` is the one axis +// under test; `frees` counts Free() so the keep-alive can be proven. +class FakeHostAddressableBackend final : public vt::Backend { + public: + explicit FakeHostAddressableBackend(bool host_addressable) + : host_addressable_(host_addressable) {} + + void* Alloc(size_t bytes) override { + ++allocs; + return std::malloc(bytes == 0 ? 1 : bytes); + } + void Free(void* p) override { + ++frees; + std::free(p); + } + void Memset(vt::Queue&, void* p, int value, size_t bytes) override { + std::memset(p, value, bytes); + } + void Copy(vt::Queue&, void* dst, const void* src, size_t bytes) override { + std::memcpy(dst, src, bytes); + } + vt::Queue CreateQueue() override { + return vt::Queue{vt::Device{vt::DeviceType::kCPU, 0}, nullptr}; + } + bool UnifiedMemory() const override { return true; } + bool DeviceMemoryIsHostAddressable() const override { return host_addressable_; } + + int allocs = 0; + int frees = 0; + + private: + bool host_addressable_; +}; + +// Build a weight with owned host bytes and upload it exactly the way +// ResidentWeight does, so the test exercises the real post-upload state. +vllm::OwnedTensor UploadedWeight(vt::Backend& b, vt::Queue& q, uint8_t pattern, + size_t nbytes) { + vllm::OwnedTensor w; + w.dtype = vt::DType::kI8; + w.rank = 1; + w.shape[0] = static_cast(nbytes); + w.bytes.resize(nbytes, pattern); + void* p = b.Alloc(nbytes); + b.Copy(q, p, w.bytes.data(), nbytes); + vt::Backend* bk = &b; + w.d_dev = std::shared_ptr(p, [bk](void* x) { bk->Free(x); }); + return w; +} + +} // namespace + +TEST_CASE("AdoptDeviceBytesAsHost leaves ONE copy on a host-addressable device") { + ScopedEnv adopt_default("VT_ADOPT_DEVICE_BYTES", "1"); + FakeHostAddressableBackend b(/*host_addressable=*/true); + vt::Queue q = b.CreateQueue(); + constexpr size_t kBytes = 4096; + + vllm::OwnedTensor w = UploadedWeight(b, q, 0xA7, kBytes); + // Precondition: two distinct copies exist, which is the defect. + REQUIRE(w.d_dev != nullptr); + REQUIRE_FALSE(w.bytes.borrowed()); + REQUIRE(static_cast(w.bytes.data()) != w.d_dev.get()); + + vllm::AdoptDeviceBytesAsHost(b, w); + + // THE MECHANISM: the host buffer now IS the device allocation. Not "freed", + // not "smaller" -- the same address, which is the only assertion that + // distinguishes one copy from two. + CHECK(w.bytes.borrowed()); + CHECK(static_cast(w.bytes.data()) == w.d_dev.get()); + CHECK(w.bytes.size() == kBytes); + // The surviving copy holds the right bytes: every reader that used to read + // the mirror reads these instead. + bool all_match = true; + for (size_t i = 0; i < kBytes; ++i) { + if (w.bytes.data()[i] != 0xA7) { all_match = false; break; } + } + CHECK(all_match); + // Logical presence is untouched -- nothing was released, so nothing may look + // absent to dispatch. + CHECK_FALSE(w.Empty()); + CHECK(w.HasHostBytes()); + CHECK_FALSE(w.host_released); + + // THE LIFETIME: the view keeps the allocation alive. Dropping d_dev alone + // must NOT free the buffer the bytes still point at. + CHECK(b.frees == 0); + w.d_dev.reset(); + CHECK(b.frees == 0); + CHECK(w.bytes.data()[0] == 0xA7); + // Only when the view goes too. + w.bytes.Reset(); + CHECK(b.frees == 1); +} + +TEST_CASE("AdoptDeviceBytesAsHost is a NO-OP where device memory is not host-addressable") { + ScopedEnv adopt_default("VT_ADOPT_DEVICE_BYTES", "1"); + // The guard that keeps every discrete-GPU path byte-identical: a CUDA + // pointer is not host-dereferenceable, so adopting it would hand a device + // address to a host memcpy. + FakeHostAddressableBackend b(/*host_addressable=*/false); + vt::Queue q = b.CreateQueue(); + constexpr size_t kBytes = 2048; + + vllm::OwnedTensor w = UploadedWeight(b, q, 0x31, kBytes); + const void* host_before = w.bytes.data(); + + vllm::AdoptDeviceBytesAsHost(b, w); + + CHECK_FALSE(w.bytes.borrowed()); + CHECK(static_cast(w.bytes.data()) == host_before); + CHECK(static_cast(w.bytes.data()) != w.d_dev.get()); + CHECK(w.bytes.size() == kBytes); +} + +TEST_CASE("AdoptDeviceBytesAsHost leaves a BORROWED buffer alone") { + ScopedEnv adopt_default("VT_ADOPT_DEVICE_BYTES", "1"); + // A borrowed buffer owns no anonymous pages (a GGUF mmap, or the single + // expansion a tied token_embd/lm_head pair shares). Adopting would reclaim + // nothing and would break the tie, so it must be skipped. + FakeHostAddressableBackend b(/*host_addressable=*/true); + vt::Queue q = b.CreateQueue(); + constexpr size_t kBytes = 1024; + + auto shared = std::make_shared>(kBytes, 0x5C); + vllm::OwnedTensor w; + w.dtype = vt::DType::kI8; + w.rank = 1; + w.shape[0] = static_cast(kBytes); + w.bytes = vllm::OwnedBytes::Borrow(shared->data(), kBytes, shared); + void* p = b.Alloc(kBytes); + b.Copy(q, p, w.bytes.data(), kBytes); + vt::Backend* bk = &b; + w.d_dev = std::shared_ptr(p, [bk](void* x) { bk->Free(x); }); + + vllm::AdoptDeviceBytesAsHost(b, w); + + CHECK(static_cast(w.bytes.data()) == shared->data()); + CHECK(static_cast(w.bytes.data()) != w.d_dev.get()); +} + namespace { // The Marlin runtime gate (MarlinMoeEnabled()) caches VT_NVFP4_MARLIN on first // use, so it cannot be flipped mid-process; the effective gate equals the LAUNCH