fix(vulkan): the 27B load held the model TWICE on a unified box (#203)

.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 |

.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.

.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 |

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 |

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 |

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

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 +

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 `<tokenizer.json | model.gguf> <corpus.txt>`.
 
+### 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

include/vllm/model_executor/models/dense_attn_block.h

@@ -193,6 +193,10 @@ inline Tensor ResidentWeight(Dev d, const OwnedTensor& w, std::vector<int64_t> s
     d.b.Copy(d.q, p, w.bytes.data(), nb);
     Backend* bk = &d.b;
     w.d_dev = std::shared_ptr<void>(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);
 }