feat(rocm): unified memory true by construction on integrated APUs (issue #41 F6, approach (b))

.agents/NOW.md

@@ -27,7 +27,7 @@ Working head: `row/backend-rocm-w0` (#41). Prior: benchmark checkpoint
 | 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 |
-| `BACKEND-ROCM` W0 | Skeleton in; **HIP never compiled** (no AMD HW) | #41 contributors build it; a compile error IS the deliverable |
+| `BACKEND-ROCM` | **(b) fix in, unverified; W0 green 4 archs** | compile + 2 new ctests + M2 ([spec](specs/rocm-unified-memory-b.md)) |
 | TP spike #287 (PR #143) | **LANDED** ([spec](specs/tensor-parallelism-spike.md)); DSpark rider grounded | dispatch TP-W1 (CPU-able) |
 | Release | SPIKE; 30/30 | #129 |
 | Surface coverage (`ARCH-ONE-SURFACE`) | **ROW 8 LANDED; #139 repair CPU-GREEN**: ABI v14 stable; registry-resolved named platform; DSR 39→32; execution guard 52/52 | Fresh re-review #139; CUDA A/B residual |

.agents/porting-inventory.md

@@ -977,6 +977,23 @@ Examples: `examples/cli` ✅ (C-API client), `examples/server` ✅ (OpenAI serve
     the sole possible source. **Method rule this earns: a grep against the
     installed package is not evidence about upstream. Record the version you
     measured, and check `main` before writing "no upstream" into the record.**
+14. **ROCm integrated-APU managed allocation (`BACKEND-ROCM` W1, approach (b)
+    from issue #41 F6, maintainer-ratified 2026-08-08).** On a device probing
+    `hipDeviceAttributeIntegrated=1` + `ManagedMemory=1` +
+    `ConcurrentManagedAccess=1`, `RocmBackend::Alloc` uses
+    `hipMallocManaged(hipMemAttachGlobal)` and `UnifiedMemory()` returns true
+    exactly then, so the CPU reference tier's host-dereference contract is
+    API-guaranteed on XNACK-less RDNA3 APUs (gfx1151/gfx1103 measure
+    `PageableMemoryAccess=0`, vetoing the CUDA-shaped W0 probe even though the
+    aliasing holds). **No upstream analog exists to mirror:** allocation is
+    torch's job in vLLM, `vllm/platforms/rocm.py` knows the APUs only as
+    device-name map entries (`rocm.py:75-77`) plus `is_navi`
+    (`rocm.py:909-910`), and `csrc/` has no `hipMallocManaged` call at the
+    pin — so this is ADDITIVE, grounded in the issue-41 measurements
+    (community F6 report), not in an upstream file. Discrete devices are
+    byte-identical to W0 (`Integrated=0` kills the branch). Spec:
+    `specs/rocm-unified-memory-b.md`; blind-written, community compile+ctest
+    evidence PENDING.
 
 ## 10. E2E test suites (T0 deliverable)
 

.agents/specs/rocm-unified-memory-b.md

@@ -0,0 +1,120 @@
+# ROCm unified memory by construction — approach (b) (BACKEND-ROCM W1)
+
+Row: `BACKEND-ROCM` (backend-matrix). Task #286, resolving issue #41's F6 fork.
+Implemented blind (no AMD hardware, no hipcc here) under this lane's standing
+policy: board owners provide compile evidence.
+
+## The decision, verbatim
+
+Maintainer call on #41 (2026-08-08), quoted in full because it is the binding
+statement this change implements:
+
+> Maintainer call on the F6 fork (@jimmykarily, @arch-btw): **approach (b)** —
+> make unified memory true by construction. On a device reporting
+> hipDeviceAttributeIntegrated=1 (and ManagedMemory=1 +
+> ConcurrentManagedAccess=1), Backend::Alloc in the ROCm backend uses
+> hipMallocManaged instead of hipMalloc, and UnifiedMemory() returns true
+> exactly then — host access becomes API-guaranteed rather than architecturally
+> incidental, which is the standard this gate exists to hold. The fix lives
+> entirely in src/vt/rocm/ (the shared gate and the CUDA/GB10 path stay
+> byte-untouched), costs one allocation-path branch, and (a) remains available
+> as a fallback note if managed allocations measure slower on gfx1151 —
+> measure, don't assume. A PR from a board owner with compile+ctest evidence
+> (the M0/M1 tables you have both already posted are exactly the right shape)
+> lands it; the reference-tier e2e (M2) should unblock immediately after. On
+> the discrete lanes this decision changes nothing: UnifiedMemory()=false stays
+> honest there and native kernels remain the path (#140 is doing exactly that
+> for gfx1201).
+
+## The measurements it rests on (jimmykarily, gfx1151; arch-btw, gfx1103)
+
+F6, measured on Strix Halo and confirmed byte-for-byte in direction on the
+780M:
+
+```
+hipDeviceAttributeIntegrated                             : 1
+hipDeviceAttributePageableMemoryAccess                   : 0   <-- the W0 veto
+hipDeviceAttributePageableMemoryAccessUsesHostPageTables : 0
+hipDeviceAttributeManagedMemory                          : 1
+hipDeviceAttributeConcurrentManagedAccess                : 1
+hipDeviceAttributeUnifiedAddressing                      : 1
+```
+
+PageableMemoryAccess needs XNACK; RDNA3 has none (`XNACK enabled: NO`, and
+`HSA_XNACK=1` changes nothing), so the W0 conjunction (integrated AND pageable,
+carried over from CudaBackend) reads false on every RDNA3 APU — while the
+aliasing the reference tier needs was separately measured to HOLD there (kernel
+writes to a hipMalloc pointer, host reads it back with no memcpy). The
+attribute answers the opposite direction (device reading pageable host memory)
+from the one the tier requires (host reading device allocations); the two
+coincide on NVIDIA integrated parts (GB10/Jetson report both 1) and come apart
+on RDNA. hipPointerGetAttributes advertises no host alias for that working
+pointer (type=device, hostPointer=nil), which is why "it worked in my test" was
+not accepted as the probe and (b) was ratified over (a).
+
+## What was built
+
+- `src/vt/rocm/rocm_backend.hip`: `ProbeDevice` additionally probes
+  `hipDeviceAttributeManagedMemory` + `hipDeviceAttributeConcurrentManagedAccess`
+  (failed probes default to 0 — the safe direction, never unregistering the
+  device). `UseManagedAlloc(caps) = integrated && managed &&
+  concurrent_managed` selects the branch; `Backend::Alloc` then uses
+  `hipMallocManaged(hipMemAttachGlobal)` (bytes==0 stays on `hipMalloc` to
+  keep the zero-size contract byte-identical); `UnifiedMemory() =
+  managed_branch || (integrated && pageable)` — the W0 conjunction kept intact
+  as ground 1. `Free` stays `hipFree` for both branches: the HIP runtime API
+  documents it as the release call for `hipMalloc` and `hipMallocManaged`
+  alike (mirroring `cudaFree`; there is no `hipFreeManaged`). Certainty HIGH,
+  API-documented, stated in the code as an assumption a board can falsify.
+- Allocation-site audit (every site in the backend): `Alloc` branches;
+  `Free` single-path by API contract; inherited `Backend::AllocPinned/
+  FreePinned` delegate to `Alloc`/`Free` (`src/vt/backend.cpp:19-20`) and so
+  ride the same branch — coherent with the pinned contract
+  (`include/vt/backend.h:76-78`, "ordinary host memory via Alloc, correct on
+  unified memory"); `rocm_rmsnorm.hip` has zero allocation sites; the skeleton
+  has no pool paths. Discrete devices take `Integrated=0`, so the managed
+  branch is provably dead and the discrete path is byte-identical to W0.
+- Introspection seam (HIP-free): `vt::rocm::ManagedAllocActive(index)` /
+  `IntegratedDevice(index)` in `include/vt/rocm/rocm_runtime.h`, so the tests
+  and board reports can name the active path without a device header.
+- Tests (`tests/vt/test_rocm_backend.cpp`, compile everywhere via the
+  syntax-check object target, run only under `VLLM_CPP_HIP` with a device):
+  the alloc-path/UnifiedMemory coupling case (discrete: managed branch dead
+  AND unified false; integrated: managed active AND unified true, loud
+  failure with the probe triple if a board class outside the fix appears),
+  and F6's decisive experiment as a standing gate (host-writes inputs with no
+  Copy, native RmsNorm kernel reads them, host-reads the kernel-written
+  output with no Copy, against the golden row).
+- CMake F1/F3 absorption (`CMakeLists.txt`, HIP branch, before
+  `check_language(HIP)`): when `ROCM_PATH` exists on disk, derive
+  `CMAKE_HIP_COMPILER_ROCM_ROOT`, seed `--rocm-path=${ROCM_PATH}` into
+  `CMAKE_HIP_FLAGS`, and export `ROCM_PATH` into the environment — each ONLY
+  when unset, so explicit flags/env win and non-ROCm machines are
+  byte-identical. F2 is downstream of the unidentified compiler and needs no
+  separate change.
+
+## Upstream grounding
+
+Upstream vLLM has NO analog: allocation is torch's job there, and
+`vllm/platforms/rocm.py` knows the APUs only as device-name map entries
+(`rocm.py:75-77`, Strix Point/Halo) plus `is_navi` (`rocm.py:909-910`);
+`grep -rn hipMallocManaged csrc/` over the pin comes back empty. This is
+therefore a recorded ADDITIVE deviation (porting-inventory §9), grounded in the
+issue-41 measurements above rather than in an upstream file. What IS mirrored:
+the `ROCM_PATH` build handling (`vllm/CMakeLists.txt:54-60` @ pin `555967922`)
+and the W0 probe's CUDA lineage (`src/vt/cuda/cuda_backend.cu:295-303`).
+
+## Verification the community owes (all PENDING, no AMD hardware here)
+
+| Gate | Board | What to post on #41 |
+|---|---|---|
+| Configure with NO manual flags (F1/F3 absorbed) | Arch/TheRock: gfx1151, gfx1103 | configure log; any flag still needed is a miss |
+| `.hip` compile of the (b) delta | any | first error text, or "clean" |
+| `ctest -R 'rocm\|cross_device'` incl. the two new cases | all four boards | pass/fail + the printed `integrated/managed-alloc/UnifiedMemory` triple |
+| Discrete branch provably dead | gfx1100, gfx1201 | the same triple: `0/0/false` |
+| M2: small dense model, greedy parity vs `--device cpu` | gfx1151, gfx1103 | tokens + `VT_OP_PROVIDER_STATS=1` fallback list |
+| (a)-fallback check: managed-alloc speed | gfx1151 | only if M5-era numbers regress vs hipMalloc — measure, don't assume |
+
+STATUS/BENCHMARKS: docs/STATUS.md and docs/BENCHMARKS.md carry the
+IMPLEMENTED-UNVERIFIED / PENDING-community rows for this change; docs/ROCM.md
+§3.1 and §5.2 are the contributor-facing statement.

.agents/state.md

@@ -43085,3 +43085,35 @@ F79-4 remain open on the landed tree; the review's merge-and-fix map is the
 binding description. Pi concurrency, BF16 GEMM/speed closure (W6) stay open
 as the lane's own next steps.
 
+
+## 2026-08-08 — ROCm approach-(b): unified memory true by construction on integrated APUs
+<!-- state: 2026-08-08T21:30 -->
+
+Implements the maintainer-ratified approach (b) from issue #41 F6 (decision
+comment 2026-08-08, quoted verbatim in `specs/rocm-unified-memory-b.md`): on a
+ROCm device probing `Integrated=1 && ManagedMemory=1 &&
+ConcurrentManagedAccess=1`, `RocmBackend::Alloc` allocates through
+`hipMallocManaged(hipMemAttachGlobal)` and `UnifiedMemory()` returns true
+exactly then — the host-dereference contract the CPU reference tier needs
+becomes API-guaranteed on XNACK-less RDNA3 APUs (gfx1151/gfx1103), where the
+W0 CUDA-shaped probe read `PageableMemoryAccess=0` and blocked M2. The W0
+conjunction stays as ground 1; `Free` stays `hipFree` (API-documented for both
+alloc paths); inherited `AllocPinned` rides the same branch via its
+delegate-to-`Alloc` base; discrete devices are byte-identical (branch provably
+dead at `Integrated=0`, gated by a new runtime test through the new HIP-free
+probes `ManagedAllocActive`/`IntegratedDevice`). F6's decisive experiment
+(kernel write → host read, no copy) is now a standing test. Also absorbed:
+issue #41 F1/F3 — when `ROCM_PATH` exists, the configure derives
+`CMAKE_HIP_COMPILER_ROCM_ROOT`, seeds `--rocm-path` into `CMAKE_HIP_FLAGS`,
+and exports `ROCM_PATH`, each only when unset, making Arch/TheRock layouts
+configure flag-free (F2 was downstream of the unidentified compiler). NO
+upstream analog to mirror (vLLM allocates via torch; `rocm.py:75-77,909-910`
+only names the APUs): recorded as ADDITIVE deviation, porting-inventory §9.14.
+
+Blind-written under this lane's policy — no AMD hardware or hipcc here; the
+CPU gates (preflight, `vllm_rocm_platform_syntax_check` -Werror
+object-compile of the platform + test TUs) are the only ones run. Community
+owes: the (b) delta's first compile, the two new ctest cases, a flag-free
+Arch/TheRock configure log, and the unblocked M2 on gfx1151/gfx1103 —
+itemized in the spec's verification table and on the #41 handoff comment
+(PR #144, `row/ROCM-UNIFIED-MEMORY-B`).

CMakeLists.txt

@@ -251,6 +251,52 @@ endif()
 if(VLLM_CPP_HIP STREQUAL "AUTO")
   set(VLLM_CPP_HIP OFF)
 elseif(VLLM_CPP_HIP)
+  # F1/F3 ABSORPTION (issue #41, jimmykarily's gfx1151 findings; layout also
+  # confirmed by arch-btw on gfx1103/TheRock). Arch and TheRock dist layouts
+  # install the compiler at ${ROCM_PATH}/lib/llvm/bin instead of
+  # ${ROCM_PATH}/llvm/bin, so clang's walk-up-from-InstalledDir autodetection
+  # derives the ROCm root as ${ROCM_PATH}/lib and finds neither the HIP headers
+  # (F1) nor the hip-lang CMake package (F3: CMakeDetermineHIPCompiler
+  # hard-fails). The unidentified-compiler fallout also broke the LINKER:
+  # prefix expansion on the --whole-archive link options further down (F2) —
+  # one root cause, three unrelated-looking failures. The report's finding:
+  # ROCM_PATH was advertised by our own FATAL_ERROR below but only ever used
+  # for find_library, so on those layouts the documented flag did not do what
+  # its error message promised.
+  #
+  # Fix: derive the three hints the report set by hand from ROCM_PATH — and
+  # ONLY where unset, so an explicit -DCMAKE_HIP_FLAGS /
+  # -DCMAKE_HIP_COMPILER_ROCM_ROOT / exported HIPFLAGS/ROCM_PATH environment
+  # always wins, and a machine whose ROCM_PATH does not exist on disk (hipcc
+  # found elsewhere on PATH) configures byte-identically to before. On the
+  # standard /opt/rocm layout the values filled are exactly what clang/CMake
+  # autodetect anyway. Upstream vLLM honours ROCM_PATH the same way
+  # (vllm/CMakeLists.txt:54-60 @ pin 555967922). All three must be set BEFORE
+  # check_language/enable_language: the report measured that the ROOT hint
+  # alone regresses the compiler identification, because the identification
+  # try-compile then has no --rocm-path.
+  if(EXISTS "${ROCM_PATH}")
+    # CMake's own ROCm-root probe (F3), consulted by CMakeDetermineHIPCompiler.
+    if(NOT DEFINED CMAKE_HIP_COMPILER_ROCM_ROOT)
+      set(CMAKE_HIP_COMPILER_ROCM_ROOT "${ROCM_PATH}")
+    endif()
+    # clang's own ROCm-root flag (F1). Rides CMAKE_HIP_FLAGS so the
+    # compiler-identification try-compile sees it too; pre-seeded as the cache
+    # entry enable_language(HIP) would otherwise create empty. Skipped when the
+    # user passed -DCMAKE_HIP_FLAGS or exported HIPFLAGS, which must keep
+    # winning.
+    if(NOT DEFINED CACHE{CMAKE_HIP_FLAGS} AND NOT DEFINED ENV{HIPFLAGS})
+      set(CMAKE_HIP_FLAGS "--rocm-path=${ROCM_PATH}" CACHE STRING
+        "Flags used by the HIP compiler during all build types.")
+    endif()
+    # hipcc/hipconfig read ROCM_PATH from the environment (the mechanism the
+    # report validated: `ROCM_PATH=/opt/rocm hipcc` compiles where bare hipcc
+    # cannot), and the environment also reaches check_language's sub-configure
+    # below, which forwards no cache variables.
+    if(NOT DEFINED ENV{ROCM_PATH})
+      set(ENV{ROCM_PATH} "${ROCM_PATH}")
+    endif()
+  endif()
   # CMake's own HIP language support (>= 3.21; we require 3.24) understands .hip
   # sources, which is why upstream calls enable_language(HIP) explicitly rather
   # than leaning on torch's setup (vllm/CMakeLists.txt:196-202).

docs/BENCHMARKS.md

@@ -346,7 +346,7 @@ built on it rather than keeping the flattering one.
 | 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 |
 | Vulkan vs llama.cpp Vulkan (`BENCH-VK-LLAMA`) | **NOT APPLICABLE: nothing measured, claimed or owed.** 22 NATIVE kernels (+6 GDN glue, CPU-oracle gated, no speed); 65 host-tier. opt-125m e2e token-exact on llvmpipe. [Detail](../.agents/specs/vulkan-full-support.md) | `VK-C` coopmat A/B on Thor (`VT_VULKAN_COOPMAT=0` A/Bs it): **11.1x-32.9x** vs our UNTILED scalar kernel, not vs a competent GEMM. `VK-E`: llama.cpp `-DGGML_VULKAN=ON` at `237ad9b96` on dgx, same GGUF, three columns |
-| ROCm (`BACKEND-GATE-ROCM-VLLM` / `-SGLANG`) | **NOT APPLICABLE: no number measured, claimed or owed.** The W0 skeleton registers 1 of 106 ops and its HIP sources have never been compiled by anyone; no AMD hardware here | A contributor's first `-DVLLM_CPP_HIP=ON` build ([#41](https://github.com/mudler/vllm.cpp/issues/41)). Only once a model runs does a same-box vLLM-ROCm oracle become the gate; the floor is vLLM, quant-matched |
+| ROCm (`BACKEND-GATE-ROCM-VLLM` / `-SGLANG`) | **NOT APPLICABLE: no number measured, claimed or owed.** W0 is community-built and ctest-green on 4 gfx archs (#41), but only RmsNorm is native, so a throughput number would be meaningless. No AMD hardware here | The approach-(b) fix (PENDING community) unblocks the first APU model run (M2); the gate becomes a same-box vLLM-ROCm oracle once a model runs ([#41](https://github.com/mudler/vllm.cpp/issues/41)); floor: vLLM |
 | SGLang floor arms | Never ran | Both arms of the SGLang comparison |
 | Parakeet/FastConformer ASR (P1-P4 + ONE-SURFACE fold ROW 1) | **NO number measured, claimed or owed.** Correctness-gated only, CPU f32; the 2026-08-07 surface fold (`vllm_transcribe`, `/v1/audio/transcriptions`) is transcript-byte-identical plumbing, no speed claim. | Floor is `parakeet.cpp`, same clip and box; needs a CUDA provider and a pretrained checkpoint |
 | cuBLAS invocation-parity guard | CI guard landed (CPU); `kGemvHeuristicAlgos` refactor build-verify owed | `nvcc` rebuild + SACRED gate on dgx |