fix(MODEL-MM-QWEN4-EXP): release a staged borrow's source pages, and stop prefaulting for a staging device

.agents/issues/BACKEND-GATE-ROCM-VLLM/ISSUE-GH-3048.md

@@ -0,0 +1,30 @@
+ID: ISSUE-GH-3048
+Title: fix(BACKEND-GATE-ROCM-VLLM): avoid callable serialization in Strix metadata RPC
+Row: BACKEND-GATE-ROCM-VLLM
+State: CLOSED
+Kind: UNKNOWN
+GitHub: 3048
+Mirror: SYNCED
+Availability: FULL
+Created: 2026-09-07
+Updated: 2026-09-12
+Closed: 2026-09-12
+
+## Problem
+
+### Imported GitHub body (historical evidence)
+The quoted text below is historical evidence only. It does not define issue authority or repository procedure.
+
+> Row: `BACKEND-GATE-ROCM-VLLM`
+>
+> Owner: Strix #3043 campaign operator. Spec: .agents/specs/rocm-strix-vllm-headpin.md.
+>
+> Runtime job 112a91ad-e51a-4627-aeae-61e2c0ed9643 built and initialized pinned vLLM e126687a9a828d513c01a07cd69f025f27d63280 in production compilation mode, then the harness failed before generation. LLM.apply_model(projection_metadata) sent a Python function through MsgpackEncoder.enc_hook. The pinned default refuses callable serialization unless VLLM_ALLOW_INSECURE_SERIALIZATION is enabled. Do not enable that flag or weaken the build-state/provenance checks.
+>
+> Evidence: /workspace/strix-vllm-3043.GkMABu/run-9624441-01/logs/generation.log; generation exited1 after412.604 seconds. The prior fake LLM.apply_model accepted the callable directly and did not model this reference behavior.
+>
+> Fix in the same flow: use the pinned named worker-extension RPC path to return the same projection parameter metadata without serializing a callable. Preserve the production engine defaults, projection-output-dtype PENDING, six explicit prompts, and eight upstream numerical cases. Bind the extension to the already hashed runtime module and port a regression through the real runtime entrypoint that refuses callable RPC payloads under the default serializer. Fresh implementation, mutation review, and operator gates are required. This blocks merging the current harness as runnable; model gateability and token parity remain unproven.
+
+## Resolution
+
+Fixed by PR #3052: the Strix oracle harness now strips tuning/injection env from the test subprocess, avoiding callable serialization in the metadata RPC.

.agents/issues/BACKEND-ROCM/ISSUE-GH-3009.md

@@ -0,0 +1,82 @@
+ID: ISSUE-GH-3009
+Title: perf(BACKEND-ROCM): non-greedy decode 22% slower than greedy — single-block sampling kernels underutilize 96 CUs
+Row: BACKEND-ROCM
+State: CLOSED
+Kind: UNKNOWN
+GitHub: 3009
+Mirror: SYNCED
+Availability: FULL
+Created: 2026-09-06
+Updated: 2026-09-12
+Closed: 2026-09-12
+
+## Problem
+
+### Imported GitHub body (historical evidence)
+The quoted text below is historical evidence only. It does not define issue authority or repository procedure.
+
+> Row: `BACKEND-ROCM`
+>
+> ## Problem
+>
+> Non-greedy decode (temperature + top-p + softmax + Gumbel-max sample) on ROCm runs **22% slower** than greedy decode at single-request batch size. On a Qwen3.5-4B Q4_K_M model (vocab ~152K) with fp8 KV cache on an RX 7900 XTX (gfx1100, 96 CUs):
+>
+> | Path | tok/s | ms/tok |
+> |---|---|---|
+> | Greedy (temp=0) | 97.8 | 10.2 |
+> | Non-greedy (temp=0.7, top_p=0.9) | 75.6 | 13.2 |
+> | **Sampling overhead** | — | **3.0 ms/tok** |
+>
+> ## Root cause
+>
+> Three per-row sampling kernels in `src/vt/rocm/rocm_sample.hip` launch as `<<<1, 256>>>` — one block of 256 threads (4 wavefronts) on **1 CU out of 96**, scanning the full ~152K vocab:
+>
+> | Kernel | Avg µs/call | Bottleneck |
+> |---|---|---|
+> | `ApplyTopKTopPRowK` | 1270 | Ternary search, up to 64 iterations over full vocab |
+> | `RandomSampleK` | 1665 | `GumbelScore` → `ExpNoise` computes `log(u)` in double precision per element — compute-bound, 148 serial evaluations per thread |
+> | `SoftmaxK` | 406 | 3-pass reduce (max, sum, normalize) |
+> | **Total** | **3341** | 3.3 ms/tok on 1 CU |
+>
+> The greedy `ArgmaxK` has the same single-block shape but is cheaper per element (one compare vs one `log`), so the gap is specific to the non-greedy path.
+>
+> ## Evidence
+>
+> rocprofv3 kernel trace, 32 decode tokens, Qwen3.5-4B Q4_K_M, fp8 KV, gfx1100:
+>
+> ```
+> ApplyTemperatureK   32 calls  avg=     3.8 us
+> ApplyTopKTopPRowK   32 calls  avg=  1270.0 us
+> SoftmaxK            32 calls  avg=   406.0 us
+> RandomSampleK       32 calls  avg=  1665.0 us
+> ```
+>
+> The CUDA backend (`src/vt/cuda/cuda_sample.cu`) already has a multi-block split-phase greedy argmax (`ArgBlocksPerRow`, `ArgmaxPartialKernel` + `ArgmaxFinalKernel`), but the ROCm random sample has no equivalent.
+>
+> ## Proposed fix
+>
+> Two changes, both in `src/vt/rocm/rocm_sample.hip`:
+>
+> 1. **Widen per-row sampling block from 256 to 1024 threads** (`kVocabBlock = 1024`). 16 wavefronts vs 4 gives 4× more latency hiding on the same CU. The `BlockRed*` helpers use `blockDim.x` instead of the compile-time constant. Top-p drops 5.5×, softmax 3.8×.
+>
+> 2. **Split-phase random sample** (`VT_SAMPLE_SPLIT=1`, default ON). Spread the Gumbel-score argmax across `kSampleSplitBlocks = 128` blocks (all 96 CUs), each thread handling 1–2 elements instead of 148. Phase A writes per-block `(score, index)` partials to scratch; Phase B reduces them. `ArgReduce` is associative and order-independent (same property the greedy split relies on), so the result is **bit-identical** to the single-block kernel.
+>
+> ## Measured result
+>
+> | Path | Before | After | Change |
+> |---|---|---|---|
+> | Greedy | 97.2 tok/s | 97.8 tok/s | unchanged |
+> | Non-greedy | 75.6 tok/s | 96.5 tok/s | **+27.6%** |
+> | Gap | 22% | 3.4% | |
+>
+> ## Correctness
+>
+> - `test_ops_sample`: 29/29 pass (2 CUDA-only skipped), 278K assertions
+> - `test_sampler`: 21/21 pass
+> - A/B test: 20 seed/temperature combinations (seeds 1–12345, temps 0.3–1.5) produce **bit-identical** output with `VT_SAMPLE_SPLIT=0` vs `VT_SAMPLE_SPLIT=1`
+> - Determinism: 3 consecutive runs with same seed produce identical output
+> - Full test suite: 225/225 pass, 3 consecutive runs, zero failures
+
+## Resolution
+
+Fixed by PR #3010: sampling block widen and split-phase implementation.

.agents/issues/BACKEND-ROCM/ISSUE-GH-3062.md

@@ -0,0 +1,29 @@
+ID: ISSUE-GH-3062
+Title: fix(BACKEND-ROCM): isolate split-sampling scratch by queue and device
+Row: BACKEND-ROCM
+State: CLOSED
+Kind: UNKNOWN
+GitHub: 3062
+Mirror: SYNCED
+Availability: FULL
+Created: 2026-09-08
+Updated: 2026-09-12
+Closed: 2026-09-12
+
+## Problem
+
+### Imported GitHub body (historical evidence)
+The quoted text below is historical evidence only. It does not define issue authority or repository procedure.
+
+> Row: `BACKEND-ROCM`
+>
+> PR #3010 stores split-sampling partial buffers and capacity in process-wide statics. Concurrent queues can overwrite another queue's partials between the two kernels. Calls on another device can reuse allocations from the first device. Host pointer and capacity updates also race.
+>
+> The repair belongs in #3010. Use the existing device-and-stream scratch ownership mechanism and preserve captured graph lifetimes. Add regression coverage for overlapping queues and allocation ownership. Retaining allocations for captured graphs must follow the existing backend lifetime contract.
+>
+> Acceptance requires focused tests, an independent mutation review, and the operator's ROCm gate. CPU-only skips do not satisfy the device gate.
+>
+
+## Resolution
+
+Fixed by PR #3010: sampling block widen and split-phase implementation.

.agents/issues/KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT/ISSUE-GH-3067.md

@@ -0,0 +1,29 @@
+ID: ISSUE-GH-3067
+Title: test(KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT): seal every device table entry
+Row: KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT
+State: CLOSED
+Kind: UNKNOWN
+GitHub: 3067
+Mirror: SYNCED
+Availability: FULL
+Created: 2026-09-08
+Updated: 2026-09-12
+Closed: 2026-09-12
+
+## Problem
+
+### Imported GitHub body (historical evidence)
+The quoted text below is historical evidence only. It does not define issue authority or repository procedure.
+
+> Row: `KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT`
+>
+> PR #3029 copies I-quant lookup tables into ROCm constant memory. Its owning spec, `.agents/specs/kernel-quant-ciq-gemm-rocm-iquant.md`, requires a complete table seal. The new header also explicitly requires a device snapshot and byte comparison. The changed tests do not implement that gate.
+>
+> Fresh review of integration head `b2ee9d8389caf974f3178613a6313e788dd93c4b` found that all four host-parsed arrays currently match their CPU references. This does not verify the executing device bytes or pin every entry in a committed test.
+>
+> Add the ROCm equivalent of the existing CUDA device-table snapshot and compare every entry against the pinned oracle reference. Preserve the current table values, kernel arithmetic, and dispatch. Mutating one entry in each table must make the focused device seal fail. Run independent review and the operator's leased ROCm gate before merging #3029. Keep the other formats owed by #1940 open.
+>
+
+## Resolution
+
+Fixed by PR #3029: IQ4_XS and IQ3_XXS now have ROCm DotIQ4XS/DotIQ3XXS kernels in rocm_grouped_gemm.hip, admitted in DeviceKeepQuantSupported.

.agents/issues/MODEL-MM-QWEN4-EXP/ISSUE-LOCAL-01M2BZ5DZ2710201WMH9TNSVH3.md

@@ -0,0 +1,19 @@
+ID: ISSUE-LOCAL-01M2BZ5DZ2710201WMH9TNSVH3
+Title: gfx1151: the whole GGUF is prefaulted and never released, so a 67.56 GiB load wedges in svm_range_set_attr
+Row: MODEL-MM-QWEN4-EXP
+State: OPEN
+Kind: bug
+GitHub: -
+Mirror: PENDING
+Availability: FULL
+Created: 2026-09-12
+Updated: 2026-09-12
+Closed: -
+
+## Problem
+
+On gfx1151 the 67.56 GiB Qwen3.8-Flash-Next UD-IQ1_S loads with zero op refusals and then wedges forever in svm_range_set_attr (uninterruptible, gpu_busy 0) on a ~31 GiB host. Two host-residency defects compound. (1) Every borrowed GGUF span is synchronously prefaulted -- madvise(WILLNEED) plus a one-byte-per-page touch loop over the whole span (PrefaultBorrowedSpan, qwen3_5_gguf_weights.cpp) -- faulting in 65.488 GiB that is read exactly once. (2) Nothing releases those pages after the device upload: ResidentWeight's staging arm copies the borrow to the device and leaves every source page mapped for the process lifetime, and OwnedTensor::ReleaseHost() refuses to madvise a BORROWED buffer on the argument that clean file-backed pages need no help. That argument holds against the page reclaimer and fails against the KFD's resident-system-memory accounting. llama.cpp (pin 10bf611e5) does the opposite on both counts: MAP_POPULATE plus an advisory posix_madvise with no synchronous touch loop, and unmap_fragment(0, mmap_used.first) after load, which for a fully offloaded model munmaps the entire mapping.
+
+## Resolution
+
+-

.agents/issues/MODEL-MM-QWEN4-EXP/ISSUE-LOCAL-01M2BZ5QK4XRETK48CXKSHKRDW.md

@@ -0,0 +1,19 @@
+ID: ISSUE-LOCAL-01M2BZ5QK4XRETK48CXKSHKRDW
+Title: OWED: chunked H2D through a pinned bounce buffer, llama.cpp's 4 x 64 MiB shape
+Row: MODEL-MM-QWEN4-EXP
+State: OPEN
+Kind: task
+GitHub: -
+Mirror: PENDING
+Availability: FULL
+Created: 2026-09-12
+Updated: 2026-09-12
+Closed: -
+
+## Problem
+
+ResidentWeight stages every weight with one hipMemcpyAsync straight out of a PAGEABLE, file-backed mmap. llama.cpp does not: llama-model-loader.cpp uses a set of pinned host buffers (4 x 64 MiB) and copies the weight through them in chunks, so the driver never has to pin or stage an arbitrarily large pageable range itself. If the gfx1151 svm_range_set_attr stall survives the host-residency fixes in .agents/specs/rocm-host-residency-after-upload.md, the pageable file-backed source is itself the trigger and this is the next change. It is deliberately NOT built there: it touches every staged weight on every backend and needs its own measurement.
+
+## Resolution
+
+-

.agents/issues/MODEL-MM-QWEN4-EXP/ISSUE-LOCAL-01M2CCNA0S74WT5WBV50B3VD0W.md

@@ -0,0 +1,19 @@
+ID: ISSUE-LOCAL-01M2CCNA0S74WT5WBV50B3VD0W
+Title: OWED: the DeviceMemoryIsHostAddressable and Synchronize terms of MaybeReleaseStagedBorrowSource are UNGATED
+Row: MODEL-MM-QWEN4-EXP
+State: OPEN
+Kind: task
+GitHub: -
+Mirror: PENDING
+Availability: FULL
+Created: 2026-09-13
+Updated: 2026-09-13
+Closed: -
+
+## Problem
+
+Deleting backend.DeviceMemoryIsHostAddressable() from MaybeReleaseStagedBorrowSource (qwen3_5_weights.cpp) leaves the focused suite 25/25 green with the binary proven changed, and so does deleting backend.Synchronize(queue). Both survive because of the harness, not because the code is wrong. The fake backend in tests/vllm/model_executor/test_resident_weight_host_addressable.cpp answers DeviceMemoryIsHostAddressable() false unconditionally, so only the platform half of the host-addressability pair is ever exercised; and HostBackend::Copy is a synchronous memcpy that does not override Synchronize, so no case in this tree can express a DMA still reading the source pages when they are dropped. Convicting the first needs a backend answering true while the platform answers false, a combination no fleet device presents and that the file's process-global registrar cannot hold beside the existing one. Convicting the second needs a backend whose Copy defers. Recorded as owed rather than repaired by weakening or contorting a passing test. See .agents/specs/rocm-host-residency-after-upload.md, its Ungated-guarantees section.
+
+## Resolution
+
+-

.agents/issues/MODEL-MM-QWEN4-EXP/ISSUE-LOCAL-01M2CKN5516AKE7W2JVDV86Z8X.md

@@ -0,0 +1,19 @@
+ID: ISSUE-LOCAL-01M2CKN5516AKE7W2JVDV86Z8X
+Title: OWED: dense_attn::ResidentWeight's staged-borrow release is UNCOVERED on the four other GGUF model families that execute it
+Row: MODEL-MM-QWEN4-EXP
+State: OPEN
+Kind: task
+GitHub: -
+Mirror: PENDING
+Availability: FULL
+Created: 2026-09-13
+Updated: 2026-09-13
+Closed: -
+
+## Problem
+
+Fix 1 of .agents/specs/rocm-host-residency-after-upload.md put MaybeReleaseStagedBorrowSource inside dense_attn::ResidentWeight (include/vllm/model_executor/models/dense_attn_block.h:246-274), which is a SHARED seam. Four production model families besides qwen4_exp reach it with weights that satisfy both of the release's predicates (a BORROWED span with mmap_fd >= 0, which only the GGUF keep-quant borrow producers set, and a platform plus backend that cannot dereference host memory): GLM-MoE-DSA (glm_moe_dsa_forward.cpp:162-186, :269, :334), GLM5-Next (glm5_next_moe.cpp:243-245), Muse-Glimmer (muse_glimmer.cpp:156, 251, 259, 273, 295, 309, 315, 326, 395, 421, 434, 435 and muse_glimmer_mm.cpp:264) and the Qwen3.5 DFlash draft head sharing a GGUF target's kept-F16 embedding table (qwen3_dflash.cpp:597, 906, 1851, 1956). The release therefore fires on dgx, thor, orin and strix for all five families, and NO test covers any family but qwen4_exp. The focused harness cannot reach the others today: its fake backend registers memory operations only, so every entry point above the seam refuses on a missing op before residency is asked about, and section 4a of the spec records this rather than contorting a case. What is owed is a harness that can drive a second family's production entry point on a non-host-addressable fake platform, and one case per family that convicts the release at that family's own call site. See .agents/specs/rocm-host-residency-after-upload.md section 4a for the verified enumeration and for the two families checked and excluded (DeepSeek-V4, Laguna).
+
+## Resolution
+
+-

.agents/issues/MODEL-MM-QWEN4-EXP/ISSUE-LOCAL-01M2CKNHVKJXT29XN8WM11KF6W.md

@@ -0,0 +1,19 @@
+ID: ISSUE-LOCAL-01M2CKNHVKJXT29XN8WM11KF6W
+Title: OWED: three latent second-reader hazards around the staged-borrow release, including the header ResidentWeight's missing expert_streamed refusal
+Row: MODEL-MM-QWEN4-EXP
+State: OPEN
+Kind: task
+GitHub: -
+Mirror: PENDING
+Availability: FULL
+Created: 2026-09-13
+Updated: 2026-09-13
+Closed: -
+
+## Problem
+
+A fresh review of the staged-borrow release found three hazards that are latent today and that this row deliberately did NOT re-engineer, because each repair is a change to a shared seam with a blast radius nothing measures (see ISSUE-LOCAL-01M2CKN5516AKE7W2JVDV86Z8X). ONE. dense_attn::ResidentWeightF32 (dense_attn_block.h:389) memoizes on d_dev_f32 INDEPENDENTLY of d_dev and reads w.bytes to upcast, so a weight that ResidentWeight has staged and released and that is then upcast would read dropped pages. It is unreachable only because every ResidentWeightF32 target today is a LoadNormBf16/ExpandBf16 owned tensor whose mmap_fd is -1, which is an accident of the current loaders and not an invariant either function states. TWO. The header ResidentWeight carries no VT_CHECK(!w.expert_streamed); its translation-unit-local twin at qwen3_5.cpp:1181 does, so nothing stops a streamed tower being staged and released on the header path. This was NOT closed in the same flow on purpose: adding the refusal would make the header seam REFUSE a weight it accepts today, and GLM5-Next stages three expert banks through it at glm5_next_moe.cpp:243-245 with no test on any staging device, so a one-line VT_CHECK could remove a working path rather than close a hazard. It needs the cross-family harness first. THREE. glm5_next_moe.cpp:292 and :300 are a host-fallback arm that reads the same OwnedTensor bytes, reachable in a partial-stage window because :243-245 are three sequential, un-transactional uploads. It re-faults, so it is correct, but the cost is a full re-read of an expert bank off the file. The window is narrower than it first reads: that arm VT_CHECKs its queue is CPU (glm5_next_moe.cpp:271) while the release only fires on a non-CPU device, so reaching it needs a CPU queue paired with a separate device pointer whose banks have already staged. Recorded with that qualification rather than dropped, because the re-read is real whenever the pairing occurs.
+
+## Resolution
+
+-