From e73cbbaee8626a0c88a758c9fa5ed3bb9a94843d Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sun, 9 Aug 2026 14:59:49 +0000 Subject: [PATCH 1/2] fix(vulkan): the 27B load held the model TWICE on a unified box FOLLOWING_AGENTS_PROTOCOL THE BUG. Loading Qwen3.6-27B bf16 (50.89 GiB) on a Vulkan GB10 could take the machine down rather than fail -- NVRM NV_ERR_NO_MEMORY out of _memdescAllocInternal, twice in one day. MEASURED cause, same binary, VT_ADOPT_DEVICE_BYTES A/B on GB10: OFF VmRSS 100.759 GiB at 50.755 GiB of Vulkan allocation, MemAvailable 13.85 / MemFree 1.13 of 119.6 GiB, and still allocating (863 of 894 buffers) when the harness killed it. ON completes at VmHWM 53.413 GiB, MemAvailable never below 47.3 GiB. The gap is a FLAT ~50.0 GiB across every high-water line -- one whole extra copy of the model. ResidentWeight uploaded each weight and kept `bytes` too. Vulkan allocates every buffer HOST_VISIBLE|HOST_COHERENT and persistently mapped (its Copy/Memset are already a plain host memcpy/memset over that pointer), so on unified memory both copies come out of the same RAM. platforms/vulkan.cpp had REASONED that "there is exactly one copy of the bytes"; this makes that true. WHAT WAS RULED OUT, with numbers. Requested bytes == driver-committed bytes EXACTLY in every run (50.756 GiB over 894 buffers at 27B), so there is no per-allocation rounding, no allocator excess, and no staging or dequant scratch to recover. The windowed source-page release DOES fire: process RSS during load holds one copy, not the mmap as well. Page cache tracks the copied bytes 1:1 and is not dropped, but it is reclaimable -- MemAvailable stayed healthy -- so it is not the OOM cause. THE FIX is an ADOPTION, not a free: `bytes` is re-pointed AT the device allocation through the existing OwnedBytes borrow, keyed alive by d_dev's own control block, so every `.bytes` reader -- the f32 upcast, the portable CPU reference tier, View/Numel -- reads the SAME bytes from the surviving copy. Nothing is dropped, so unlike ReleaseHost this needs no "is the device path committed" proof. Backend::DeviceMemoryIsHostAddressable() is the gate and defaults false, so every discrete-GPU path is byte-identical. It is deliberately narrower than UnifiedMemory(): CUDA on GB10 is unified yet a cudaMalloc pointer is still not host-dereferenceable. Also lands the accounting that made the attribution possible (VT_VULKAN_ALLOC_STATS, counters always maintained so a test can assert on them) and three mutation-checked mechanism tests. GATES on GB10, Vulkan build (VLLM_CPP_VULKAN=ON, CUDA absent so CurrentPlatform resolves kVULKAN -- the engine selected device type 3): test_opt_paged_engine 6/6 prompts token-exact (96/96 tokens), 0 declines; test_backend_cross_device 11/11 (132); test_vulkan_backend 35/35 (2650/2650); test_qwen36_weights adoption cases 3/3 (20), red under all three mutations (no-op body, dropped host-addressable guard, dropped keep-alive -- the last a real use-after-free). Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude-Code:claude-opus-5 [Claude Code] --- .agents/NOW.md | 2 +- .agents/benchmark-record.md | 108 +++++++++++ docs/BENCHMARKS.md | 1 + docs/ENVIRONMENT.md | 2 + docs/FEATURES.md | 9 +- docs/STATUS.md | 2 +- docs/USAGE.md | 11 ++ .../model_executor/models/dense_attn_block.h | 4 + .../model_executor/models/qwen3_5_weights.h | 33 ++++ include/vt/backend.h | 24 +++ src/vllm/model_executor/models/qwen3_5.cpp | 4 + .../model_executor/models/qwen3_5_weights.cpp | 42 ++++ src/vt/vulkan/vulkan_backend.cpp | 7 + src/vt/vulkan/vulkan_context.cpp | 182 ++++++++++++++++++ src/vt/vulkan/vulkan_context.h | 25 +++ tests/vllm/test_qwen36_weights.cpp | 158 +++++++++++++++ 16 files changed, 607 insertions(+), 7 deletions(-) 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/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 From 4947c0a856964ecb269835be1896ceed81085b1f Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sun, 9 Aug 2026 15:25:08 +0000 Subject: [PATCH 2/2] record: link #203 to BACKEND-VULKAN in the issue intake table The new intake rule requires an open issue before a row is claimed, linked from the roadmap, the row's spec and the PR. No open issue covered the Vulkan double-copy at load -- #83 is memory BUDGETING, a different thing -- so #203 was opened for it and is linked here. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude-Code:claude-opus-5 [Claude Code] --- .agents/roadmap_v1.md | 1 + 1 file changed, 1 insertion(+) 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 |