From 14af81bd36a89fa831076b6cadfc15e0f7bcf75c Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sat, 8 Aug 2026 07:36:40 +0000 Subject: [PATCH 1/2] =?UTF-8?q?docs(rocm):=20record=20the=20ratified=20F6?= =?UTF-8?q?=20decision=20=E2=80=94=20approach=20(b),=20unified=20memory=20?= =?UTF-8?q?by=20construction?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Issue #41's F6 fork is decided: on integrated managed-capable devices the backend will allocate through hipMallocManaged so host access is API-guaranteed, and UnifiedMemory() returns true exactly then (maintainer comment, 2026-08-08, quoted in docs/ROCM.md §3.1). This commit is the contributor-facing statement: the decision and its behavior table (integrated vs discrete), the community-verified W0 state (M0/M1 MET on gfx1151/gfx1103/gfx1100/gfx1201), the two teardown known-issues, and the exact post-fix sequence for board owners (§5.2). The implementation follows in the next commit. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-fable-5 [ClaudeCode] --- docs/BUILD.md | 29 +++++--- docs/ROCM.md | 189 ++++++++++++++++++++++++++++++++++++++++---------- 2 files changed, 170 insertions(+), 48 deletions(-) diff --git a/docs/BUILD.md b/docs/BUILD.md index 15f5d3d27..19f6e342b 100644 --- a/docs/BUILD.md +++ b/docs/BUILD.md @@ -109,17 +109,18 @@ cmake -S . -B build-vulkan -DVLLM_CPP_VULKAN=ON cmake --build build-vulkan -j ``` -## ROCm build (AMD GPUs) — never yet compiled +## ROCm build (AMD GPUs) — community-verified W0, blind F6 fix -> **Read this before you file a bug.** The HIP sources in this tree have **never -> been compiled by anyone.** There is no AMD GPU and no ROCm toolchain on any -> machine the maintainers use, so unlike every other backend here this one has no -> build report at all — not even "it compiles". If it fails for you, that is the -> expected first outcome and the most useful thing you can report. Please do, -> on [issue #41](https://github.com/mudler/vllm.cpp/issues/41). +> The W0 HIP sources compiled clean and passed `ctest -R 'rocm|cross_device'` +> on four community boards — gfx1151, gfx1103, gfx1100, gfx1201 +> ([issue #41](https://github.com/mudler/vllm.cpp/issues/41)). The +> unified-memory fix on top of them (approach (b), +> [docs/ROCM.md §3.1](ROCM.md)) was again written with **no AMD GPU or ROCm +> toolchain on any maintainer machine**, so a compile error in it is expected, +> useful, and belongs on #41. ```sh -cmake -S . -B build-hip -DVLLM_CPP_HIP=ON -DVLLM_CPP_HIP_ARCHITECTURES=gfx1100 +cmake -S . -B build-hip -DVLLM_CPP_HIP=ON -DCMAKE_BUILD_TYPE=Release cmake --build build-hip -j ctest --test-dir build-hip -R 'rocm|cross_device' ``` @@ -128,14 +129,22 @@ ctest --test-dir build-hip -R 'rocm|cross_device' installed GPU, which is what you want when building on the machine you will run on. The validated names are upstream vLLM's `HIP_SUPPORTED_ARCHS`; anything else configures with a warning and is passed to hipcc anyway. If ROCm lives outside -`/opt/rocm`, point at it with `-DROCM_PATH=`. +`/opt/rocm`, point at it with `-DROCM_PATH=`. When `ROCM_PATH` names a +real install, the configure now derives the compiler hints from it +(`CMAKE_HIP_COMPILER_ROCM_ROOT`, `--rocm-path` in `CMAKE_HIP_FLAGS`, the +`ROCM_PATH` environment variable — each only if you have not set it), which is +what makes Arch and TheRock dist-tarball layouts configure without the manual +flags issue #41's gfx1151 report needed. An explicit `-DCMAKE_BUILD_TYPE` (any +optimizing one) matters on ROCm: a `-O0` device build trips a CLR teardown race +([#132](https://github.com/mudler/vllm.cpp/issues/132)). `-DVLLM_CPP_HIP=ON` **fails the configure** when no HIP compiler is found rather than quietly producing a CPU-only build, for the same reason the CUTLASS note above exists: a silent downgrade is indistinguishable from success. What exists today is the W0 skeleton — the `vt::Backend`, the `Platform`, one -registered kernel (RmsNorm), and the tests that gate them. What that does and +registered kernel (RmsNorm), and the tests that gate them — plus the approach-(b) +unified-memory branch for integrated APUs. What that does and does not get you, and where to start on your specific board, is [docs/ROCM.md](ROCM.md). diff --git a/docs/ROCM.md b/docs/ROCM.md index e81e14aa5..29cf22d36 100644 --- a/docs/ROCM.md +++ b/docs/ROCM.md @@ -1,12 +1,19 @@ # ROCm (AMD GPU) backend — contributor guide -**State today: the W0 skeleton is committed, and no HIP source in it has ever -been compiled.** `kROCM` exists, `-DVLLM_CPP_HIP=ON` exists, and there is a -`vt::Backend`, a `Platform`, and exactly one registered kernel (RmsNorm). The -plain-C++ parts compile and are tested here; the three `.hip` files have not been -built by anyone, because no maintainer machine has an AMD GPU. **Your first HIP -compile is genuinely the first**, and a failure is the expected outcome rather -than a sign you did something wrong. +**State today: the W0 skeleton is community-verified on four architectures, and +the F6 unified-memory fix (approach (b)) is committed but unverified.** +[Issue #41](https://github.com/mudler/vllm.cpp/issues/41) board owners compiled +the HIP sources clean and ran the ctest gates on gfx1151 (Strix Halo), gfx1103 +(Radeon 780M), gfx1100 (4x 7900 XTX) and gfx1201 (2x R9700) — M0 and M1 MET on +all four, with two runtime-teardown caveats recorded in §7. Their headline +finding, F6: `UnifiedMemory()` probed **false** on the RDNA3 APUs, because +XNACK-less RDNA reports `PageableMemoryAccess=0`, so the zero-kernel reference +tier the whole unified-memory plan rests on did not install and M2 was blocked. +The ratified fix (§3.1) allocates through `hipMallocManaged` on integrated +managed-capable devices so host access is API-guaranteed — written blind like +the skeleton before it, so **the (b) branch owes the same community compile+run +evidence W0 already earned**; a compile error in it is useful data, not a +mistake on your side. This page exists because several people offered hardware in [issue #41](https://github.com/mudler/vllm.cpp/issues/41), and it answers the @@ -14,9 +21,11 @@ three questions that decide whether that goes anywhere: what a backend actually *is* in this codebase, what to write first on the hardware you own, and what "done" means. The design record behind the skeleton, including what was deliberately left out, is -[.agents/specs/rocm-backend-w0.md](../.agents/specs/rocm-backend-w0.md). +[.agents/specs/rocm-backend-w0.md](../.agents/specs/rocm-backend-w0.md); the +unified-memory decision record is +[.agents/specs/rocm-unified-memory-b.md](../.agents/specs/rocm-unified-memory-b.md). -Everything here is checked against the tree on 2026-08-06. Where a number is +Everything here is checked against the tree on 2026-08-08. Where a number is counted, the command that counts it is given, because these numbers drift. ## 1. Why ROCm is the cheapest backend to add @@ -56,13 +65,13 @@ on a real machine, and what has not. |---|---|---| | Device enum | [`include/vt/device.h`](../include/vt/device.h) | ✅ compiled; the enum forced exactly one switch site tree-wide | | — | [`include/vt/rocm/rocm_arch.h`](../include/vt/rocm/rocm_arch.h) — gfx name → `(major, minor)`, ported 1:1 from `rocm.py:223` | ✅ **unit-tested**, 40 assertions, no GPU needed | -| Runtime backend | [`src/vt/rocm/rocm_backend.hip`](../src/vt/rocm/rocm_backend.hip) — the 6 `vt::Backend` virtuals | ❌ **never compiled** | -| Op table | [`src/vt/rocm/rocm_ops.hip`](../src/vt/rocm/rocm_ops.hip) — one `RegisterOp` line | ❌ never compiled | -| Kernel | [`src/vt/rocm/rocm_rmsnorm.hip`](../src/vt/rocm/rocm_rmsnorm.hip) | ❌ never compiled | -| Platform | [`src/vllm/platforms/rocm.cpp`](../src/vllm/platforms/rocm.cpp) — mirrors `vllm/platforms/rocm.py` | ✅ compiles `-Werror` (plain C++, object-compiled in every build as a bit-rot guard); never *run* | +| Runtime backend | [`src/vt/rocm/rocm_backend.hip`](../src/vt/rocm/rocm_backend.hip) — the 6 `vt::Backend` virtuals | ✅ W0 compiled + ctest-run on gfx1151/1103/1100/1201 (#41) — ❌ the approach-(b) delta is **unbuilt** | +| Op table | [`src/vt/rocm/rocm_ops.hip`](../src/vt/rocm/rocm_ops.hip) — one `RegisterOp` line | ✅ compiled + run on the same four boards | +| Kernel | [`src/vt/rocm/rocm_rmsnorm.hip`](../src/vt/rocm/rocm_rmsnorm.hip) | ✅ NMSE ≤ 5e-4 vs the CPU oracle on all four boards (F5: RDNA is wave32, the wave64 hazard moves to a future gfx9 board) | +| Platform | [`src/vllm/platforms/rocm.cpp`](../src/vllm/platforms/rocm.cpp) — mirrors `vllm/platforms/rocm.py` | ✅ compiles `-Werror` everywhere; ✅ run on the four boards | | Attention | *(none yet — `get_attn_backend_priority()` returns empty)* | — | -| Build | `VLLM_CPP_HIP` in [`CMakeLists.txt`](../CMakeLists.txt) | ✅ the OFF path and the fail-without-hipcc path | -| Test | [`tests/vt/test_rocm_backend.cpp`](../tests/vt/test_rocm_backend.cpp) | ✅ compiles (same guard) — ❌ never run | +| Build | `VLLM_CPP_HIP` in [`CMakeLists.txt`](../CMakeLists.txt) | ✅ ON-path configure+build on the four boards (Arch/TheRock now auto-hinted, §5 M0); ✅ the OFF path and the fail-without-hipcc path | +| Test | [`tests/vt/test_rocm_backend.cpp`](../tests/vt/test_rocm_backend.cpp) | ✅ W0 cases run green (1044 assertions in the #41 tables) — ❌ the two approach-(b) cases never run | So the shape is decided and the parts that hold a *decision* are tested; what you are validating is the API glue. Adding your own op is one line in @@ -113,13 +122,60 @@ Two rules that keep this honest: `VT_OP_PROVIDER_STATS=1` prints the first time each `(op, device)` falls back, and `GetReferenceTierHits()` **must be 0 in any performance measurement**. A non-zero value means you benchmarked the CPU. +### 3.1 The F6 fix: unified memory true by construction (approach (b)) + +Issue #41's headline finding (F6, measured on gfx1151, confirmed on gfx1103): +XNACK-less RDNA3 APUs report `hipDeviceAttributeIntegrated=1` but +`hipDeviceAttributePageableMemoryAccess=0`, so the W0 probe (CUDA's own +conjunction, integrated AND pageable) answered `UnifiedMemory() == false` on the +very boards the zero-kernel plan was written for — even though host dereference +of `hipMalloc` memory demonstrably worked there. The attribute answers the +opposite question (device reading pageable host memory) from the one the tier +needs (host reading device allocations); the two coincide on NVIDIA integrated +parts and come apart on RDNA. + +The maintainer decision (#41, 2026-08-08) is approach **(b)** — make the claim +true by construction rather than gate on an architectural accident: + +> 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. + +What each device class gets: + +| Device class | `Backend::Alloc` | `UnifiedMemory()` | Reference tier | M2 | +|---|---|---|---|---| +| Integrated + managed-capable (gfx1151, gfx1103: all three attributes probed 1) | `hipMallocManaged(hipMemAttachGlobal)` | **true**, by construction | installs — a model runs with one native kernel | unblocked | +| Integrated, NOT managed-capable (no known board; the probes default to 0 on error) | `hipMalloc` | false unless the W0 conjunction holds | does not install | blocked — post the probe triple on #41 | +| Discrete (gfx1100, gfx1201, MI50...) | `hipMalloc` — the managed branch is provably dead (`Integrated=0`) | false | never installs (memory-safety gate) | native kernels required, unchanged | + +The free path is `hipFree` for both branches: the HIP runtime API documents it +as the release call for `hipMalloc` and `hipMallocManaged` allocations alike, +mirroring `cudaFree`. `Backend::AllocPinned` inherits the base delegation to +`Alloc` (`src/vt/backend.cpp:19`), so pinned blocks ride the same branch and +stay host-accessible — coherent with its contract (`include/vt/backend.h:76-78`). +Introspection for tests and bug reports: +`vt::rocm::ManagedAllocActive(index)` / `IntegratedDevice(index)` in +[`include/vt/rocm/rocm_runtime.h`](../include/vt/rocm/rocm_runtime.h) report +which path the silicon took, and `tests/vt/test_rocm_backend.cpp` gates that the +alloc path and the `UnifiedMemory()` claim move together — including F6's +decisive experiment (kernel writes, host reads back, **no copy**) as a standing +test. + +Approach (a) — gate the tier on `Integrated` alone — remains the recorded +fallback **if managed allocations measure slower on gfx1151: measure, don't +assume** (the maintainer decision, verbatim). Decision record: +[.agents/specs/rocm-unified-memory-b.md](../.agents/specs/rocm-unified-memory-b.md). + ## 4. Pick your first task from your hardware | Hardware | Arch | Memory | Start here | |---|---|---|---| -| Strix Halo / GTR9 Pro 128GB | gfx1151 | unified | **Build M0/M1, then M2.** Once it compiles, the reference tier means a model runs with no further kernel written. Closest analogue to GB10, so the residency-policy question in §6 is yours | -| Radeon 780M iGPU | gfx1103 | shared | **Build M0/M1**, same path, smaller models. Best position to find every place a "CUDA" assumption is really an "NVIDIA" assumption. A vLLM-ROCm oracle is unlikely on this board, so M4 stays PENDING there — fine, and to be said rather than papered over | -| 4x 7900 XTX | gfx1100 | discrete | **Build M0/M1, then the kernel path**, since the reference tier cannot install on a dGPU and a model needs real kernels. The only board that can host a vLLM-ROCm oracle for M4 and, later, multi-GPU TP — the backend already registers all four at `Device{kROCM, i}` | +| Strix Halo / GTR9 Pro 128GB | gfx1151 | unified | M0/M1 **MET** (#41). Now: **verify the §3.1 fix, then M2** (§5.2) — the reference tier means a model runs with no further kernel written. Closest analogue to GB10, so the residency-policy question in §6 is yours | +| Radeon 780M iGPU | gfx1103 | shared | M0/M1 **MET** (#41). Same §5.2 path, smaller models. Best position to find every place a "CUDA" assumption is really an "NVIDIA" assumption. A vLLM-ROCm oracle is unlikely on this board, so M4 stays PENDING there — fine, and to be said rather than papered over | +| 4x 7900 XTX | gfx1100 | discrete | M0/M1 **MET** (#41, with the #132 caveat). Now **the kernel path**, since the reference tier cannot install on a dGPU and a model needs real kernels. The only board class that can host a vLLM-ROCm oracle for M4 and, later, multi-GPU TP — the backend already registers all four at `Device{kROCM, i}`. gfx1201 (2x R9700) is on the same discrete lane via PR #140 | These do not collide. Two people can be on M0/M1/M2 on unified parts while a third does the hipify pass, and the discrete board is what turns the result into @@ -127,26 +183,83 @@ a gated backend. ## 5. Milestones as concrete PRs -**M0 — build. WRITTEN, unverified.** Tri-state `VLLM_CPP_HIP`, hipcc detection -that fails loudly, `VLLM_CPP_HIP_ARCHITECTURES`, `ROCM_PATH`. What remains is -for someone to run it. Acceptance: `cmake -DVLLM_CPP_HIP=ON` configures and -`cmake --build` produces a binary. **This is the open task.** - -**M1 — platform + backend. WRITTEN, unverified.** `rocm_backend.hip`, -`platforms/rocm.cpp`, the `kROCM` enum, the capability parse (tested), and one -registered op. Acceptance: `ctest -R 'rocm|cross_device'` green on the device — -which also means the RmsNorm kernel matched the CPU oracle at NMSE ≤ 5e-4, so -seam 3 is proven end to end. - -Expect M0/M1 to need fixes. A compile error in `rocm_backend.hip` is the single -most valuable thing anyone can report right now, and it belongs in this repo -rather than in a fork. - -**M2 — first model end to end.** On a unified part this is mostly free: assert -`ReferenceTierEligible(kROCM)` and run a small dense model. Acceptance: greedy -token parity against the **CPU backend** on the same build, plus the -`VT_OP_PROVIDER_STATS=1` output showing which ops fell back, which is your -kernel to-do list, sorted by real usage rather than by guesswork. +**M0 — build. MET** on gfx1151, gfx1103, gfx1100 and gfx1201 (#41 tables). +Tri-state `VLLM_CPP_HIP`, hipcc detection that fails loudly, +`VLLM_CPP_HIP_ARCHITECTURES`, `ROCM_PATH`. The Arch/TheRock findings F1/F3 are +now absorbed into the configure: when `ROCM_PATH` points at a real install +(default `/opt/rocm`), CMake derives `CMAKE_HIP_COMPILER_ROCM_ROOT`, seeds +`--rocm-path` into `CMAKE_HIP_FLAGS`, and exports `ROCM_PATH` into the +environment — each only when you have not set it yourself. The manual +three-flag workaround from the gfx1151 report +(`-DCMAKE_HIP_COMPILER_ROCM_ROOT=... -DCMAKE_HIP_FLAGS=--rocm-path=...` on top +of `-DROCM_PATH`) is therefore **legacy**: still honoured if passed, no longer +required. F2 (raw `--whole-archive` reaching the clang driver) was downstream +of the unidentified compiler and disappears with F1/F3 — if you still see it, +your compiler identification failed and that configure log is the thing to +post. *The absorption itself is untested on a real Arch/TheRock layout — a +configure log from one, with no manual flags, is wanted evidence on #41.* + +**M1 — platform + backend. MET** on the same four boards: `ctest -R +'rocm|cross_device'` green, RmsNorm within NMSE ≤ 5e-4 of the CPU oracle on +real silicon — with two runtime caveats, both teardown-related, in §5.1 below. + +**M2 — first model end to end. UNBLOCKED-UNVERIFIED on unified parts** by the +§3.1 fix: assert `ReferenceTierEligible(kROCM)` and run a small dense model. +Acceptance: greedy token parity against the **CPU backend** on the same build, +plus the `VT_OP_PROVIDER_STATS=1` output showing which ops fell back, which is +your kernel to-do list, sorted by real usage rather than by guesswork. + +### 5.1 Known runtime issues on the #41 boards + +- **TheRock nightly teardown hang (gfx1103).** All three test binaries print + `Status: SUCCESS!` and then never exit, so `ctest` times out waiting. + arch-btw's GDB backtrace pins it inside `libamdhip64.so.7` during + `__cxa_finalize`/`_dl_fini`, waiting on HSA `AsyncEventsLoop` threads stuck + in an `ioctl` wait — an upstream runtime teardown deadlock in the TheRock + nightly (`10.1.0a20260731`), not a vllm.cpp bug. Treat "SUCCESS printed, then + hang" as a PASS of the test body plus this known issue; report the ROCm + build you saw it on. +- **`-O0` hostcall teardown race ([#132](https://github.com/mudler/vllm.cpp/issues/132), + gfx1100).** A no-build-type compile leaves the RmsNorm kernels with + hidden-hostcall metadata at `-O0`; ROCm CLR's listener handshake can then + deadlock the process finalizer under CPU saturation (intermittent, 14/20 at + 48 threads). Validated avoidance: build with an optimization level + (`-DCMAKE_BUILD_TYPE=Release`), which removes the hostcall path entirely. + +### 5.2 The sequence for board owners, post-F6-fix + +Everything below assumes the tree at or after the approach-(b) change. On +Arch/TheRock, no compiler flags beyond `ROCM_PATH` should now be needed — if +that is false, the configure log is finding number one. + +```sh +cmake -S . -B build-hip -DVLLM_CPP_HIP=ON -DCMAKE_BUILD_TYPE=Release \ + -DROCM_PATH=/opt/rocm # or your TheRock dist prefix +cmake --build build-hip -j +ctest --test-dir build-hip -R 'rocm|cross_device' --output-on-failure +``` + +What to report on #41, in the M0/M1 table shape already in use there: + +1. The configure/compile/link results, and whether any manual flag was still + required (that would mean §5's F1/F3 absorption missed your layout). +2. The new probe triple printed by `test_rocm_backend` — `integrated`, + `managed-alloc`, `UnifiedMemory()` — plus the pass/fail of the two + approach-(b) cases ("alloc path and UnifiedMemory() move together", + "kernel-written value is host-readable with no copy"). +3. On an APU, the M2 attempt: a small dense model through the CLI. `--device` + has no `rocm` literal yet; `auto` (the default) selects ROCm on an AMD box + with no CUDA via the platform priority walk. E.g.: + + ```sh + VT_OP_PROVIDER_STATS=1 ./build-hip/examples/vllm-cli \ + --model \ + --prompt 'The capital of France is' --max-tokens 8 --temperature 0 + ``` + + Post the generated tokens, whether they match the same command with + `--device cpu`, and the `VT_OP_PROVIDER_STATS` fallback list — that list is + the prioritized M3 kernel to-do for your board. **M3 — kernels + attention.** Hipify `src/vt/cuda/` family by family, starting with what M2's fallback log actually hit: layernorm, rope, activations, glue, From 5c8969f23b2ee67e622d0e2b48e8ea3507a9e2ca Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sat, 8 Aug 2026 07:43:57 +0000 Subject: [PATCH 2/2] =?UTF-8?q?feat(rocm):=20approach=20(b)=20=E2=80=94=20?= =?UTF-8?q?unified=20memory=20true=20by=20construction=20on=20integrated?= =?UTF-8?q?=20APUs?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Implements the maintainer-ratified approach (b) from issue #41 F6. On a device probing hipDeviceAttributeIntegrated=1 + ManagedMemory=1 + ConcurrentManagedAccess=1, RocmBackend::Alloc allocates through hipMallocManaged(hipMemAttachGlobal) and UnifiedMemory() returns true exactly then: host access to every Backend::Alloc block becomes API-guaranteed rather than architecturally incidental, which is what the CPU reference tier's host-dereference contract needs on XNACK-less RDNA3 APUs (gfx1151/gfx1103 measure PageableMemoryAccess=0, vetoing the W0 CUDA-shaped probe even though the aliasing demonstrably holds there). - The W0 conjunction (integrated AND pageable) stays as ground 1; Free stays hipFree, API-documented for both alloc paths; the inherited AllocPinned delegates to Alloc (backend.cpp:19) and rides the branch, coherent with its unified-memory contract (backend.h:76-78). - Discrete devices byte-identical to W0: the branch is provably dead at Integrated=0, asserted by a new runtime test through the new HIP-free probes ManagedAllocActive/IntegratedDevice (rocm_runtime.h). - F6's decisive experiment is now a standing test: host-write inputs with no Copy, native RmsNorm kernel reads them, host-read the kernel-written output with no Copy. - CMake F1/F3 absorption: when ROCM_PATH exists, derive CMAKE_HIP_COMPILER_ROCM_ROOT, seed --rocm-path into CMAKE_HIP_FLAGS and export ROCM_PATH, each only when unset, before check_language(HIP) — Arch/TheRock layouts configure flag-free (F2 was downstream of the unidentified compiler; no separate change). - No upstream analog to mirror (vLLM allocates via torch; rocm.py:75-77, 909-910 only name the APUs): recorded as ADDITIVE deviation, porting-inventory §9.14. Spec: .agents/specs/rocm-unified-memory-b.md (decision quoted verbatim + the community verification table). Lands via row/ROCM-UNIFIED-MEMORY-B (PR #144). BLIND-WRITTEN: no AMD GPU or hipcc on the authoring machine; the .hip delta has never been compiled. CPU gates green (preflight EXIT=0; -Werror object-compile of the platform + test TUs; compile-mutation of the new probe API red then green on restore). Community evidence owed per the spec table; PR #144 carries the ask. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-fable-5 [ClaudeCode] --- .agents/NOW.md | 2 +- .agents/backend-matrix.md | 2 +- .agents/porting-inventory.md | 17 +++ .agents/specs/rocm-unified-memory-b.md | 120 +++++++++++++++++++++ .agents/state.md | 32 ++++++ CMakeLists.txt | 46 ++++++++ docs/BENCHMARKS.md | 2 +- docs/FEATURES.md | 4 +- docs/STATUS.md | 6 +- include/vt/rocm/rocm_runtime.h | 31 ++++-- scripts/check-public-doc-tables.py | 9 +- src/vt/rocm/rocm_backend.hip | 144 ++++++++++++++++++++----- tests/vt/test_rocm_backend.cpp | 104 ++++++++++++++++-- 13 files changed, 472 insertions(+), 47 deletions(-) create mode 100644 .agents/specs/rocm-unified-memory-b.md diff --git a/.agents/NOW.md b/.agents/NOW.md index bca56be40..e8f2be4cf 100644 --- a/.agents/NOW.md +++ 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 | diff --git a/.agents/backend-matrix.md b/.agents/backend-matrix.md index 8d16fff8e..ab80c0420 100644 --- a/.agents/backend-matrix.md +++ b/.agents/backend-matrix.md @@ -225,7 +225,7 @@ on every listed target. | `BACKEND-PLATFORM` | Platform capability/memory-model seam (`is_cuda`/`is_cpu`, `is_unified_memory`, `has_device_capability`, `supported_dtypes`, `residency_policy`, `supports_graph_capture`; `get_attn_backend_priority` REALIZED by `BACKEND-ATTN-REGISTRY` item 4; **+ the S3 capability predicates `supports_fp8`/`cutlass_fp4_supported`/`opaque_attention_op`/`is_integrated_gpu`/`support_static_graph_mode`/`is_device_capability_family`, mirroring `interface.py:914,933,977,1058,441-476` + `nvfp4_utils.py:56`, added by `CLAIM-BACKEND-SEAM-S3-1` and consumed by the fp4/fp8/graph gates in `qwen3_5.cpp`**); self-registered per `DeviceType` via `CurrentPlatform()` | `vllm/platforms/interface.py:134-229` (`class Platform`), `:409-439` (get/has_device_capability), `:181-187` (supported_dtypes) @ pin `e24d1b24` | interface [interface.h:56](../include/vllm/platforms/interface.h#L56), registry/CurrentPlatform [platform.cpp:57](../src/vllm/platforms/platform.cpp#L57), CPU [cpu.cpp:11](../src/vllm/platforms/cpu.cpp#L11), CUDA [cuda.cpp:14](../src/vllm/platforms/cuda.cpp#L14); composes `vt::Backend` [backend.h:22](../include/vt/backend.h#L22) | [platform tests](../tests/vllm/platforms/test_platform.cpp#L28) (registration + CPU capabilities + has_device_capability), clean CPU `-Werror` + full CPU CTest; 7 memory-model/residency sites migrated (`runner.cpp`, `model_registry.cpp`, `qwen3_5.cpp`) — **per-tensor sites key on `GetPlatform(.device.type).is_cuda()` (the OBJECT's device), NOT the process-global `CurrentPlatform()`, which is accelerator-first and would mis-route a CPU queue/tensor on a GPU box (regression fixed 2026-07-18; red DGX CPU tests `test_platform`/`test_qwen27_dense_forward`); `CurrentPlatform()` is reserved for process-level "which accelerator" questions only**. DGX behavior-preserving model gates 27B 235/235 + 35B 315/315 CONFIRMED @ (CUDA -Werror-clean after a GCC13 dangling-pointer FP fix). **Item-2 residency_policy() CONSUMPTION LANDED (2026-07-19, `CLAIM-BACKEND-PLATFORM-2`):** the host-free / load-stream / DevicePool-cap decisions in `qwen3_5.cpp` now READ `GetPlatform(.device.type).residency_policy()` (per-device) instead of an inline `device.type`/env gate — host-free `qwen3_5.cpp:BuildMoeMarlinResident` via `ShouldReleaseHostWeights`, load-stream `Qwen3_5Model::PrepareMarlinResident` via `ShouldInterleaveLoadStream` (both `interface.h`), DevicePool soft cap via `residency_policy().device_pool_cap_bytes` (`DBuf`/`DevicePool`). `CudaPlatform::residency_policy().release_host_weights_after_upload` flipped **false→true** (now CONSUMED ⇒ reproduces today's GB10 host-free-after-Marlin-build EXACTLY); `MarlinMoeEnabled()` stays the orthogonal KERNEL-PATH gate; `VT_MOE_HOST_FREE`/`VT_MOE_LOADSTREAM` env stay as overrides. Consumption unit-tested `test_platform.cpp` (7 cases/43 assertions). A new (discrete) GPU sets `residency_policy()` values ⇒ ZERO model edit. **DGX gate PASSED @ `62fc0e0` (production flags CUTLASS sm120a + FA2 sm_121a + Triton AOT, one flock): clean CUDA `-Werror` 0 warn, 27B 235/235 + 35B 315/315 token-exact, 35B VmHWM ≈ 4.0 GiB (load-stream win preserved), memcheck 0 errors.** | [Platform seam plan](specs/extensibility-platform-seam-2026-07-18.md) | `ANCHOR-BACKFILL` — item-1 extraction DONE + DGX-CONFIRMED; item-2 residency consumption LANDED + DGX-PROVEN (`CLAIM-BACKEND-PLATFORM-2`); kernel-shape dispatch branches deliberately left for items 4/5 | CLAIM-BACKEND-PLATFORM-1 / CLAIM-BACKEND-PLATFORM-2 | | `BACKEND-CPU` | CPU correctness and production path | `platforms/cpu.py:42-125`, CPU ops rooted at `csrc/cpu/torch_bindings.cpp:123-139`; llama.cpp `ggml-cpu.c:471-610,3024-3390` owns the added native pool reference | [cpu_backend.cpp:11](../src/vt/cpu/cpu_backend.cpp#L11), [threadpool.cpp:78](../src/vt/cpu/cpu_threadpool.cpp#L78), [chunked ops](../src/vt/cpu/cpu_ops.cpp#L88) | [backend tests](../tests/vt/test_backend.cpp#L10), [op parity](../tests/parity/test_op_parity.cpp#L34), [threadpool/upstream-test port](../tests/vt/test_cpu_threadpool.cpp#L63); 1/3/20 full CPU suites + TSAN pass. The binding 20-core Arm/i8mm Qwen3.5-2B Q8 single-stream gate is at or ahead of same-file llama.cpp on prefill/decode/RSS; server concurrency remains open. **Raspberry Pi 5 / Cortex-A76 lane GATING 2026-08-06:** R0-R3 pin the exact Q8_K_XL vehicle, local buildx/QEMU build and physical-Pi correctness/PMU baseline. **R4-R5 GREEN:** [KERNEL-CPU-A76-Q8-DOT](kernel-matrix.md) adds exact-order compiler SDOT and a scheduled AAPCS64 leaf; Pi operation and 64-token model outputs are exact. Assembly beats compiler SDOT 3.66-5.08% on M1/T1 and M128/T1/T4, lowers Qwen TTFT 1.55% and E2E 0.13%, so A76+DotProd auto-selects it. M1/T4 regresses 2.43%. **The separate Pi llama.cpp floor is now MEASURED/NOT MET on speed:** vllm.cpp is 0.461x prefill / 0.653x decode+E2E, while using 24.2% less RSS; exact-prompt 64-token output matches. No 20-core/i8mm result transfers. Thread partition, BF16 GEMM/speed closure and concurrency stay open. **macOS UNBLOCKED 2026-07-22 (`CLAIM-BACKEND-FANOUT-1`, W0):** the CPU registrar force-link is fixed; `test_backend` is **7/7 (18/18)** and the whole tree builds `-Werror`-clean on M4. Two unrelated macOS test gaps remain recorded in [environment.md](environment.md) | [threadpool leaf](specs/gguf-cpu-threadpool.md); [RPi5 Cortex-A76 campaign](specs/rpi5-cortex-a76-cpu-optimization.md); [Pi competitor evidence](../docs/bench-evidence/rpi5-a76-llamacpp-20260806.md) | `PARTIAL` | PR #65 | | `BACKEND-CPU-ZEN` | AMD Zen CPU with ZenDNN/zentorch dispatch and weight prepack | `platforms/zen_cpu.py:12-32`; detection `platforms/__init__.py:153-192`; `tests/test_zen_cpu_platform_detection.py:8-37` | generic CPU backend only; no Zen-specific dispatch | - | [CUDA inventory](specs/cuda-architecture-inventory.md); leaf spike required | `INVENTORIED` | - | -| `BACKEND-ROCM` | AMD ROCm/HIP | `platforms/__init__.py:110-128`, `platforms/rocm.py:43-125` (incl. `_capability_from_gcn_arch:223-291`, `_get_backend_priorities:407`), `CMakeLists.txt:20-59,196-211` (HIP build), ROCm ops rooted at `csrc/rocm/moe_q_gemm_rdna3.cu:1` @ pin `555967922` | Lands through the realized seams: a `platforms/rocm.cpp` Platform + a `kROCM` attention TU self-registering via the [attn-registry seam](../include/vllm/v1/attention/registry.h#L44) + a `rocm.cpp` `get_attn_backend_priority()` slot (mirror `rocm.py`'s ROCM_ATTN/AITER ordering); ZERO selector/model/runner edit **AS-BUILT:** **W0 skeleton landed 2026-08-06; the HIP TUs are NEVER COMPILED (no AMD hardware here — weaker than `build-supported`, which means it compiles). The two plain-C++ legs ARE compile-gated (`vllm_rocm_platform_syntax_check`, object-only, never linked).** [device.h](../include/vt/device.h) `kROCM`; [rocm_arch.h](../include/vt/rocm/rocm_arch.h) (capability parse, HIP-free, 1:1 from `rocm.py:223-291`); [rocm_runtime.h](../include/vt/rocm/rocm_runtime.h); [rocm_backend.hip](../src/vt/rocm/rocm_backend.hip) (6 virtuals + multi-device registrar); [rocm_rmsnorm.hip](../src/vt/rocm/rocm_rmsnorm.hip) (port of `cuda_ops.cu:96-126`); [rocm_ops.hip](../src/vt/rocm/rocm_ops.hip) (1 of 106 OpIds); [platforms/rocm.cpp](../src/vllm/platforms/rocm.cpp) (empty attn priority — no kernel, so no claim); `VLLM_CPP_HIP` + `VLLM_CPP_HIP_ARCHITECTURES` + `ROCM_PATH` in [CMakeLists.txt](../CMakeLists.txt); and the ONE non-additive site this uncovered — [platform.cpp](../src/vllm/platforms/platform.cpp) `kCurrentPriority` (a platform absent from that walk registers fine and is NEVER selected, with no compiler diagnostic), now exposed via `CurrentPlatformPriority()` and gated | **VERIFIED on the CPU tier (no GPU):** [test_rocm_arch](../tests/vt/test_rocm_arch.cpp) 7 cases / 40 assertions incl. upstream's own worked examples + the gfx9 feature-suffix strip + `static_assert`s; [test_platform](../tests/vllm/platforms/test_platform.cpp) new case gating that EVERY `DeviceType` is in the selection walk and CPU is last; clean `-Werror` CPU build (the enum forced exactly ONE switch site tree-wide, `test_backend_cross_device.cpp:59`); full `ctest` green; `check-device-leakage` unchanged at 32. **PENDING, no hardware:** any HIP compile, [test_rocm_backend](../tests/vt/test_rocm_backend.cpp), and the cross-device RmsNorm-vs-CPU-oracle comparison (`test_backend_cross_device` picks kROCM up automatically at NMSE <= 5e-4) | [ROCm W0 spec](specs/rocm-backend-w0.md), [contributor guide](../docs/ROCM.md) | `ACTIVE` | `CLAIM-ROCM-W0-1` | +| `BACKEND-ROCM` | AMD ROCm/HIP | `platforms/__init__.py:110-128`, `platforms/rocm.py:43-125` (incl. `_capability_from_gcn_arch:223-291`, `_get_backend_priorities:407`), `CMakeLists.txt:20-59,196-211` (HIP build), ROCm ops rooted at `csrc/rocm/moe_q_gemm_rdna3.cu:1` @ pin `555967922` | Lands through the realized seams: a `platforms/rocm.cpp` Platform + a `kROCM` attention TU self-registering via the [attn-registry seam](../include/vllm/v1/attention/registry.h#L44) + a `rocm.cpp` `get_attn_backend_priority()` slot (mirror `rocm.py`'s ROCM_ATTN/AITER ordering); ZERO selector/model/runner edit **AS-BUILT:** **W0 skeleton landed 2026-08-06 and was COMMUNITY-VERIFIED (issue #41): compiled clean + `ctest -R 'rocm\|cross_device'` green on gfx1151, gfx1103, gfx1100 and gfx1201 — M0/M1 MET on all four (teardown caveats: TheRock-nightly exit hang on gfx1103; #132 `-O0` CLR race on gfx1100). W1 approach-(b) landed 2026-08-08, IMPLEMENTED-UNVERIFIED (blind, no AMD hardware here): managed-alloc branch on integrated managed-capable devices (`hipMallocManaged` + `UnifiedMemory()=true` by construction, resolving F6 where XNACK-less RDNA3 probes `PageableMemoryAccess=0`), introspection probes `ManagedAllocActive`/`IntegratedDevice`, two runtime-gated tests, and the F1/F3 `ROCM_PATH`→HIP-compiler-hints CMake absorption for Arch/TheRock layouts. Discrete path byte-identical (branch provably dead at `Integrated=0`). The two plain-C++ legs stay compile-gated everywhere (`vllm_rocm_platform_syntax_check`, object-only, never linked).** [device.h](../include/vt/device.h) `kROCM`; [rocm_arch.h](../include/vt/rocm/rocm_arch.h) (capability parse, HIP-free, 1:1 from `rocm.py:223-291`); [rocm_runtime.h](../include/vt/rocm/rocm_runtime.h); [rocm_backend.hip](../src/vt/rocm/rocm_backend.hip) (6 virtuals + multi-device registrar); [rocm_rmsnorm.hip](../src/vt/rocm/rocm_rmsnorm.hip) (port of `cuda_ops.cu:96-126`); [rocm_ops.hip](../src/vt/rocm/rocm_ops.hip) (1 of 106 OpIds); [platforms/rocm.cpp](../src/vllm/platforms/rocm.cpp) (empty attn priority — no kernel, so no claim); `VLLM_CPP_HIP` + `VLLM_CPP_HIP_ARCHITECTURES` + `ROCM_PATH` in [CMakeLists.txt](../CMakeLists.txt); and the ONE non-additive site this uncovered — [platform.cpp](../src/vllm/platforms/platform.cpp) `kCurrentPriority` (a platform absent from that walk registers fine and is NEVER selected, with no compiler diagnostic), now exposed via `CurrentPlatformPriority()` and gated | **VERIFIED on the CPU tier (no GPU):** [test_rocm_arch](../tests/vt/test_rocm_arch.cpp) 7 cases / 40 assertions incl. upstream's own worked examples + the gfx9 feature-suffix strip + `static_assert`s; [test_platform](../tests/vllm/platforms/test_platform.cpp) new case gating that EVERY `DeviceType` is in the selection walk and CPU is last; clean `-Werror` CPU build (the enum forced exactly ONE switch site tree-wide, `test_backend_cross_device.cpp:59`); full `ctest` green; `check-device-leakage` unchanged at 32. **COMMUNITY-VERIFIED (W0, #41 tables):** [test_rocm_backend](../tests/vt/test_rocm_backend.cpp) 5 cases/1044 assertions and the cross-device RmsNorm-vs-CPU-oracle comparison (NMSE <= 5e-4) green on all four boards. **PENDING-community (W1 (b), no hardware here):** the (b) delta's first HIP compile, the two new test cases (alloc-path/UnifiedMemory coupling; kernel-write→host-read no-copy), the flag-free Arch/TheRock configure, and the unblocked M2 reference-tier e2e on gfx1151/gfx1103 | [ROCm W0 spec](specs/rocm-backend-w0.md), [unified-memory (b) decision record](specs/rocm-unified-memory-b.md), [contributor guide](../docs/ROCM.md) | `ACTIVE` | `CLAIM-ROCM-W0-1` | | `BACKEND-XPU` | Intel XPU loyal port; gating exploration E4 ([backends.md](backends.md) kernel sourcing: SYCL-vs-Level-Zero call + whether upstream XPU attention contracts translate 1:1; explore when T2 scheduling begins) | `platforms/__init__.py:131-150`, `platforms/xpu.py:103-125` | enum slot only [device.h:11](../include/vt/device.h#L11) | [unavailable-backend test](../tests/vt/test_backend.cpp#L37); **SPIKE-ONLY verdict, MEASURED 2026-07-22: HW-BLOCKED — no Intel GPU on any box we own, and no acquisition path recorded.** Compounding finding: the loyalty target is INCOMPLETE upstream — `vllm/platforms/xpu.py:11-13` imports the EXTERNAL `vllm_xpu_kernels._C`/`._moe_C`/`._xpu_C`, and a search of the pinned tree for `*.sycl`/`*.dp.cpp` returns NOTHING, so there is no in-tree SYCL kernel source to mirror 1:1. What stays doable with zero Intel hardware: the attention-selector/priority POLICY port (pure data, unit-testable), oneAPI DPC++ compile coverage, and SYCL kernels executed on the oneAPI OpenCL CPU device for unit-level numerics. **No e2e or performance gate is proposed, because none can be run.** No XPU backend code exists | [backend fan-out](specs/backend-fanout-metal-vulkan-xpu.md); [CUDA inventory](specs/cuda-architecture-inventory.md) | `SPIKE` | `CLAIM-BACKEND-FANOUT-1` | | `BACKEND-TPU` | vLLM TPU parity surface | `platforms/__init__.py:35-56,202-208`, `platforms/tpu.py:9-20` | - | - | [CUDA inventory](specs/cuda-architecture-inventory.md) | `INVENTORIED` | - | | `BACKEND-ACCEL-PROVIDER` | **The acceleration-PROVIDER seam** — two or more implementations of ONE `vt::` op on ONE `DeviceType` coexisting, selected DETERMINISTICALLY and observably. Answers the user's standing requirement "build it so we can extend acceleration easily to other platforms", which is a question about the SEAM, not about any one backend | no upstream mirror (vllm.cpp original). Mirrors the SHAPE of the runtime tactic/heuristic dispatch every provider in vLLM's chain uses instead of compile-time pinning: flashinfer's per-arch tactic registry (`flashinfer/gemm/fp4_gemm_cutlass_template_sm120.h:187-220`), cuBLASLt/CUTLASS per-call heuristics | `vt::OpProvider` + device-neutral `vt::ProviderCaps` [op_provider.h](../include/vt/op_provider.h); registry, deterministic selection, decline-and-fall-back and stats [op_provider.cpp](../src/vt/op_provider.cpp). `RegisterOp`/`GetOp`/`OpRegistered` MOVED OUT of [ops.cpp](../src/vt/ops.cpp) with **identical signatures and semantics** — all ~70 op wrappers in that file are byte-unchanged, which is what "zero call-site edits" means. GENERALIZED FROM [cuda_arch_tactics.h](../src/vt/cuda/cuda_arch_tactics.h) (capacity-bounded static storage, capability predicate, decline-by-return, selection stats), lifted out of `vt::cuda` and keyed on (OpId, DeviceType). First consumer: the MLX GEMM provider on Metal [metal_mlx_provider.mm](../src/vt/metal/metal_mlx_provider.mm) | **THE DEFECT FIXED, STATED PRECISELY:** the old table held ONE `void*` per (OpId, DeviceType) and `RegisterOp` overwrote it with no check and no warning, so two providers of one op resolved by STATIC-INIT ORDER ACROSS TUs — unspecified by the standard, i.e. a nondeterministic BUILD. Selection is now `(priority DESC, name ASC by strcmp)`, both compile-time constants of the registering TU, hence a pure function of WHICH providers are linked. **PROVEN, not asserted:** [test_op_provider.cpp](../tests/vt/test_op_provider.cpp) registers the SAME three providers in OPPOSITE orders on two slots and requires the same winner AND the same full order (11 cases / 47 assertions), plus equal-priority name tie-break, duplicate-name rejection, capability-predicate skip, caps re-publication re-resolution, decline-and-fall-back down a 3-deep stack, the `declines` counter, per-call `selections` stats, and the `VT_OP_PROVIDER_DISABLE` same-binary A/B lever. **END-TO-END on a real accelerator (M4):** MLX and the native MSL GEMM coexist on `kMatmul`/`kMatmulBT`, MLX wins by priority, and an interior-pointer activation makes MLX DECLINE exactly once and fall through to ours with the right answer ([test_metal_backend.cpp](../tests/vt/test_metal_backend.cpp), 9 cases / 108 assertions with MLX ON). clean `-Werror` 0 warnings on all three toolchains (AppleClang 21 CLT-only macOS 26.5.2 Metal ON and Metal+MLX ON; GCC Linux CPU; nvcc 13.0 sm_121a on dgx with `VLLM_CPP_TRITON=ON`). **REGRESSION-SAFE on the hottest shared file:** `GetOp` steady state is one relaxed atomic load of a resolved-selection cache (was one array load); negative resolution is memoized so `OpRegistered`, which the fused-recipe ladder calls per step for ops a backend lacks, stays O(1); the provider-disable lookup short-circuits lock-free when nothing is disabled. dgx regression set ALL UNCHANGED, each STANDALONE (see the state log entry) — anchor `tests/vt/test_op_provider.cpp:64` | [Metal/MLX reuse study §6](specs/metal-mlx-reuse-study.md) (which specced it, work row `W0b-2`); reconciled with — not rivalling — [drop-in kernel ABI](specs/dropin-kernel-abi.md), which is the ARGUMENT half for raw-C launchers while this is the SELECTION half | `ACTIVE` — the mechanism is landed and gated with two real providers on one op; it is deliberately not closed, because the CUDA (cuBLASLt/CUTLASS/flashinfer), CPU (llama.cpp `vec_dot`) and Vulkan (coopmat) rows of the §6.1 table are DESIGNED FOR but not yet POPULATED, and the `QuantTypeTraits` split (study §3.4, work row `W0b-3`) that keys on the same predicate is not landed | `CLAIM-BACKEND-ACCEL-PROVIDER-1` | diff --git a/.agents/porting-inventory.md b/.agents/porting-inventory.md index c9e08e75b..f1f362015 100644 --- a/.agents/porting-inventory.md +++ b/.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) diff --git a/.agents/specs/rocm-unified-memory-b.md b/.agents/specs/rocm-unified-memory-b.md new file mode 100644 index 000000000..230c7d196 --- /dev/null +++ b/.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. diff --git a/.agents/state.md b/.agents/state.md index a890b51e1..c6efa4205 100644 --- a/.agents/state.md +++ b/.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 + + +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`). diff --git a/CMakeLists.txt b/CMakeLists.txt index 473e07085..63a5c1af2 100644 --- a/CMakeLists.txt +++ 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). diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index bdbbc4ff2..7a089d2a9 100644 --- a/docs/BENCHMARKS.md +++ b/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 | diff --git a/docs/FEATURES.md b/docs/FEATURES.md index 2c3810c54..ed3e55eef 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -202,7 +202,7 @@ the registered engine forward. | CPU (x86, Arm i8mm; A76 assembly correct/default, llama speed gate open) | ✅ | ◐ | ☐ | ✅ | | Metal (Apple Silicon) | ✅ | ☐ | ☐ | ✅ | | Vulkan | ◐ | ☐ | ☐ | ✅ | -| ROCm | ☐ (W0 skeleton, HIP never compiled) | ✅ | ✅ | ✅ | +| ROCm | ◐ (W0 community-verified on 4 gfx archs, #41; APU unified-memory fix landed, unverified) | ✅ | ✅ | ✅ | | XPU / TPU | ☐ | ✅ | ◐ | ☐ | CUDA runtime-verified on GB10 (sm_121a), Jetson Thor (sm_110) and Jetson AGX @@ -279,7 +279,7 @@ CPU elementwise GEMM (f32/f16/bf16) runs AVX2 and AVX-512 tiers on x86 where the | LoRA end to end | CPU brick landed | Unwired standalone; not usable through the server | | Multimodal over HTTP | Architecturally blocked | Vision tower lives outside the registered engine forward | | Embedding / reranking models | Engine side only | Pooler and runner path landed, no model architecture registered | -| ROCm | W0 skeleton, unbuilt | Backend + platform + 1 op (RmsNorm); the HIP sources have never been compiled by anyone (no AMD board here). Open: [ROCM.md](ROCM.md), [#41](https://github.com/mudler/vllm.cpp/issues/41) | +| ROCm | W0 verified by community, model e2e pending | Backend + platform + 1 op, ctest-green on gfx1151/1103/1100/1201 ([#41](https://github.com/mudler/vllm.cpp/issues/41)). APU UnifiedMemory fix in (managed allocs, unverified); M2 unblocks with it. [ROCM.md](ROCM.md) | | XPU, TPU | Not started | CUDA, CPU, Metal and Vulkan are the built backends | | Custom logits processors on CUDA | Open, not root-caused | Segfaults in a CUDA build, 232/232 green on CPU | | Memory budgeting (`ROAD-V1-MEM`, #83) | Scoped, spike owed | No profiling; KV pool is a hand-typed `--num-blocks`. Target: auto-size to the declared workload, optional total-footprint cap, refuse before allocating | diff --git a/docs/STATUS.md b/docs/STATUS.md index 57b59f9a2..2b3ff18f9 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -412,9 +412,9 @@ 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, GEMV 1.8x; 22 native, +6 GDN glue gated, recurrences host-tier; qwen3_5 #125 fixed, VERIFIED on 27B Vulkan -[campaign](../.agents/specs/vulkan-full-support.md)), ROCm (W0 skeleton: -backend, platform, 1 of 106 ops; **its HIP sources have never been compiled by -anyone** — weaker than "build-supported", since no AMD board is here; +[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 unblocks on verification; [guide](ROCM.md)), and the full tool-calling template surface. **Scale-out / distributed execution is scoped, with two legs landed CPU-gated** (2026-07-28): one `vt::` collective / process-group abstraction with backend transports (NCCL / RDMA / MLX-ring) mirrors vLLM's diff --git a/include/vt/rocm/rocm_runtime.h b/include/vt/rocm/rocm_runtime.h index 7709035e2..ebc3b3cf3 100644 --- a/include/vt/rocm/rocm_runtime.h +++ b/include/vt/rocm/rocm_runtime.h @@ -1,10 +1,11 @@ -// HIP-free declarations of the two probes the ROCm backend exposes upward -// (BACKEND-ROCM, W0). Mirrors the role of src/vt/vulkan/vulkan_context.h in the -// Vulkan skeleton: the engine-side platform TU asks "is there a device?" without -// ever including , which is what keeps src/vllm/ free of -// vendor headers and lets the platform leg be read as plain C++. +// HIP-free declarations of the probes the ROCm backend exposes upward +// (BACKEND-ROCM, W0 + the W1 approach-(b) introspection pair). Mirrors the role +// of src/vt/vulkan/vulkan_context.h in the Vulkan skeleton: the engine-side +// platform TU asks "is there a device?" without ever including +// , which is what keeps src/vllm/ free of vendor headers and +// lets the platform leg be read as plain C++. // -// Both are defined in src/vt/rocm/rocm_backend.hip and both are noexcept: they +// All are defined in src/vt/rocm/rocm_backend.hip and all are noexcept: they // are called from static-init registrars, where throwing would abort the process // at load time on a machine that merely happens to have HIP installed. #pragma once @@ -25,4 +26,22 @@ bool DeviceAvailable() noexcept; // actionable, one that says "an AMD GPU" is not. std::string DeviceArchName(int index) noexcept; +// The raw hipDeviceAttributeIntegrated probe for `index`; false when the device +// is absent or the probe fails. HIP-free so tests can branch integrated vs +// discrete without a device header. +bool IntegratedDevice(int index) noexcept; + +// Which allocation path Backend::Alloc takes for device `index` — the +// approach-(b) introspection seam (issue #41 F6, maintainer-ratified +// 2026-08-08). True: every Backend::Alloc block is +// hipMallocManaged(hipMemAttachGlobal), which is what makes UnifiedMemory() +// true BY CONSTRUCTION on integrated, managed-capable devices (Strix Halo +// gfx1151, Radeon 780M gfx1103). False: plain hipMalloc — every discrete card, +// byte-identical to the W0 behavior — or no device at `index`. +// +// Exists so (a) tests/vt/test_rocm_backend.cpp can assert the discrete branch +// is provably dead without a HIP header, and (b) a board owner can report which +// path their silicon took without reading driver internals. +bool ManagedAllocActive(int index) noexcept; + } // namespace vt::rocm diff --git a/scripts/check-public-doc-tables.py b/scripts/check-public-doc-tables.py index cf23d414b..80f38fdb7 100755 --- a/scripts/check-public-doc-tables.py +++ b/scripts/check-public-doc-tables.py @@ -382,7 +382,14 @@ def features_errors(text: str) -> list[str]: # mapping ... behavioural CPU gate ... remaining work" narrative appeared # twice back to back; the longer, newer version stays). Nothing was lost, # only de-duplicated. Net -92. - "chars": 279150, + # + # 279130 since 2026-08-08 (measured 279120 after the #145/#151 rebases): the ROCm + # entry in "Not supported yet" swapped "HIP never compiled" (falsified by + # the four #41 community build reports) for the current binding state — W0 + # community-green, approach-(b) F6 fix in blind-unverified, M2 on + # verification — paid for inside the same parenthetical; the board/arch + # detail lives in docs/ROCM.md and .agents/backend-matrix.md. Net -12. + "chars": 279130, "h2_sections": 11, "long_paragraphs": 89, "oversized_cells": 47, diff --git a/src/vt/rocm/rocm_backend.hip b/src/vt/rocm/rocm_backend.hip index ab0099b21..11fd3db9b 100644 --- a/src/vt/rocm/rocm_backend.hip +++ b/src/vt/rocm/rocm_backend.hip @@ -1,13 +1,15 @@ -// ROCm leg of the vt::Backend seam (BACKEND-ROCM, W0 skeleton). +// ROCm leg of the vt::Backend seam (BACKEND-ROCM, W0 skeleton + W1 approach-(b) +// unified-memory fix). // -// **UNBUILT.** No AMD GPU exists on the authoring machine, so this translation -// unit has never been compiled by anyone. That is weaker than the project's -// "build-supported" label (which means it compiles and emits real machine code) -// and it must not be quoted as such until a contributor reports a build. It is -// written to be the smallest possible thing a first compile can be wrong about: -// the one piece with a DECISION in it (the gfx capability parse) was lifted out -// into include/vt/rocm/rocm_arch.h, which IS compiled and unit-tested on a -// CPU-only box (tests/vt/test_rocm_arch.cpp, 40 assertions). +// BUILD STATE, kept honest per issue #41: the W0 version of this TU was +// compiled clean AND ran its ctest gates on community boards — gfx1151 (Strix +// Halo, jimmykarily), gfx1103 (Radeon 780M, arch-btw), gfx1100 (4x 7900 XTX, +// VikashLoomba) and gfx1201 (2x R9700, bakon11). The approach-(b) DELTA in this +// revision is again written blind (no AMD GPU on the authoring machine) and +// owes the same community compile evidence before it may be called anything +// stronger. The one piece with a DECISION in it (the gfx capability parse) +// remains lifted out into include/vt/rocm/rocm_arch.h, which IS compiled and +// unit-tested on a CPU-only box (tests/vt/test_rocm_arch.cpp, 40 assertions). // // SHAPE: src/vt/cuda/cuda_backend.cu, virtual for virtual, including its // unified-memory reasoning (see UnifiedMemory below) and its multi-device @@ -50,13 +52,16 @@ void Check(hipError_t err, const char* what) { hipStream_t AsStream(const Queue& q) { return static_cast(q.handle); } -// One probe per device, cached: the registrar needs the residency attributes and -// the capability together, and hipGetDeviceProperties is not free. Mirrors -// vt/cuda/cuda_device_caps.h, kept local because there is no second consumer yet. +// One probe gathers the residency attributes and the capability together. +// Mirrors vt/cuda/cuda_device_caps.h in shape but is NOT cached (every consumer +// is init-time or test-time, never per-step), and stays file-local: the public +// probes below re-run it rather than exporting the struct. struct DeviceCaps { bool valid = false; bool integrated = false; bool pageable_memory_access = false; + bool managed_memory = false; + bool concurrent_managed_access = false; int major = 0; int minor = 0; std::string gcn_arch; @@ -89,27 +94,94 @@ DeviceCaps ProbeDevice(int index) noexcept { } caps.integrated = integrated != 0; caps.pageable_memory_access = pageable != 0; + // The MANAGED-memory attributes decide the approach-(b) allocation branch + // below. A failed probe of either must not unregister the device the way the + // two probes above do: both default to 0, which disables the managed branch + // and the unified claim — the safe direction (an APU then behaves like W0: + // registered, discrete-shaped, no CPU fallback tier). + int managed = 0; + int concurrent = 0; + if (hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory, index) != hipSuccess) { + managed = 0; + } + if (hipDeviceGetAttribute(&concurrent, hipDeviceAttributeConcurrentManagedAccess, index) != + hipSuccess) { + concurrent = 0; + } + caps.managed_memory = managed != 0; + caps.concurrent_managed_access = concurrent != 0; caps.valid = true; return caps; } +// APPROACH (b) FROM ISSUE #41 F6, maintainer-ratified 2026-08-08: "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." +// +// Why the W0 probe alone was not enough (measured on gfx1151, confirmed on +// gfx1103): XNACK-less RDNA3 APUs report Integrated=1 but +// PageableMemoryAccess=0, so the CUDA-shaped conjunction below read false even +// though host dereference of hipMalloc memory demonstrably worked — an +// architectural accident this project refuses to gate memory safety on. The +// managed route makes the aliasing a documented API property instead. All three +// attributes measured 1 on gfx1151 (Strix Halo) and gfx1103 (Radeon 780M). +// +// A discrete card reports Integrated=0, so this returns false — the branch is +// provably dead — on every dGPU, and the discrete path stays byte-identical to +// W0. +bool UseManagedAlloc(const DeviceCaps& caps) noexcept { + return caps.integrated && caps.managed_memory && caps.concurrent_managed_access; +} + class RocmBackend final : public Backend { public: // Non-throwing by design: constructed during static init by the registrar, // which probes the device beforehand. Same contract as CudaBackend. - RocmBackend(int device, bool unified_memory, int major, int minor) noexcept - : device_(device), unified_memory_(unified_memory), major_(major), minor_(minor) {} + RocmBackend(int device, bool unified_memory, bool managed_alloc, int major, int minor) noexcept + : device_(device), + unified_memory_(unified_memory), + managed_alloc_(managed_alloc), + major_(major), + minor_(minor) {} int DeviceCapabilityMajor() const override { return major_; } int DeviceCapabilityMinor() const override { return minor_; } - // hipMalloc's alignment is at least 256 bytes, as cudaMalloc's is, which - // satisfies the >= 64B contract Backend::Alloc owes StepArena. + // Both branches satisfy the >= 64B contract Backend::Alloc owes StepArena: + // hipMalloc aligns to at least 256 bytes as cudaMalloc does, and managed + // allocations are at least page-granular. void* Alloc(size_t bytes) override { void* p = nullptr; - Check(hipMalloc(&p, bytes), "hipMalloc"); + if (managed_alloc_ && bytes > 0) { + // Approach (b), issue #41 F6: on an integrated managed-capable device, + // every Backend::Alloc block is a MANAGED allocation, so host access to + // it is guaranteed by the API (ManagedMemory=1 + + // ConcurrentManagedAccess=1 are preconditions of this branch), which is + // exactly the property UnifiedMemory()==true advertises and the CPU + // reference tier depends on. hipMemAttachGlobal (the API default) keeps + // the block accessible from every stream and every device. + // + // The bytes==0 case stays on hipMalloc so the zero-size contract + // (success, nullptr out) is byte-identical to W0 on every device rather + // than resting on hipMallocManaged's less-documented zero-size behavior. + Check(hipMallocManaged(&p, bytes, hipMemAttachGlobal), "hipMallocManaged"); + } else { + Check(hipMalloc(&p, bytes), "hipMalloc"); + } return p; } + // ONE free path for both Alloc branches, deliberately: the HIP runtime API + // documents hipFree as the release call for hipMalloc AND hipMallocManaged + // allocations (there is no hipFreeManaged; hipHostFree is for hipHostMalloc + // only), mirroring cudaFree, whose documentation likewise covers + // cudaMallocManaged memory. Certainty: HIGH — API-documented on both sides — + // but stated rather than assumed because this file is written blind: if a + // board ever reports hipErrorInvalidValue from this call on the managed + // path, this assumption is the first suspect, and the fix would be tracking + // the branch per-pointer, not a global switch. void Free(void* p) override { Check(hipFree(p), "hipFree"); } void Memset(Queue& q, void* p, int value, size_t bytes) override { Check(hipMemsetAsync(p, value, bytes, AsStream(q)), "hipMemsetAsync"); @@ -141,16 +213,24 @@ class RocmBackend final : public Backend { // would be memory corruption, so the tier must stay off and unimplemented ops // must throw instead. // - // The conjunction is deliberate and is CudaBackend's own (cuda_backend.cu:295- - // 303): pageable-memory-access ALONE is not enough, because a discrete part can - // report it while the driver services host pointers through HMM page migration. - // That is not the zero-copy aliasing this contract needs, so an INTEGRATED GPU - // is additionally required. Probed, never inferred from the gfx name. + // TRUE on exactly two grounds, both API-anchored, computed by the registrar: + // 1. The W0 conjunction, CudaBackend's own (cuda_backend.cu:295-303): + // Integrated AND PageableMemoryAccess. Pageable-memory-access ALONE is + // not enough, because a discrete part can report it while the driver + // services host pointers through HMM page migration — not the zero-copy + // aliasing this contract needs. Kept byte-for-byte for any part that + // reports both (NVIDIA-style integrated semantics). + // 2. The approach-(b) managed branch (UseManagedAlloc above): every + // Backend::Alloc block on this device is hipMallocManaged memory, so + // host access is guaranteed by construction. This is the route + // XNACK-less RDNA3 APUs take, where ground 1 reads false (issue #41 F6). + // Probed, never inferred from the gfx name. bool UnifiedMemory() const override { return unified_memory_; } private: int device_ = 0; bool unified_memory_ = false; + bool managed_alloc_ = false; int major_ = 0; int minor_ = 0; }; @@ -176,8 +256,11 @@ struct Registrar { for (int i = 0; i < count; ++i) { const DeviceCaps caps = ProbeDevice(i); if (!caps.valid) continue; - backends.push_back(std::make_unique( - i, caps.pageable_memory_access && caps.integrated, caps.major, caps.minor)); + const bool managed_alloc = UseManagedAlloc(caps); + const bool unified = + managed_alloc || (caps.pageable_memory_access && caps.integrated); + backends.push_back( + std::make_unique(i, unified, managed_alloc, caps.major, caps.minor)); RegisterBackend(Device{DeviceType::kROCM, i}, backends.back().get()); } } @@ -197,4 +280,17 @@ std::string DeviceArchName(int index) noexcept { return caps.gcn_arch; } +bool IntegratedDevice(int index) noexcept { + const DeviceCaps caps = ProbeDevice(index); + return caps.valid && caps.integrated; +} + +bool ManagedAllocActive(int index) noexcept { + // Answers from the SAME pure predicate over the SAME probe the registrar + // used, so this cannot drift from what the registered backend decided at + // static init: the attributes are immutable hardware properties. + const DeviceCaps caps = ProbeDevice(index); + return caps.valid && UseManagedAlloc(caps); +} + } // namespace vt::rocm diff --git a/tests/vt/test_rocm_backend.cpp b/tests/vt/test_rocm_backend.cpp index 4bb1a2e63..a59ff0a52 100644 --- a/tests/vt/test_rocm_backend.cpp +++ b/tests/vt/test_rocm_backend.cpp @@ -2,14 +2,16 @@ // C++ ROCm backend tests to port. Mirrors tests/vt/test_metal_backend.cpp, which // mirrors tests/vt/test_backend.cpp, so all three read side by side. // -// **THIS FILE HAS NEVER RUN.** It is LINKED into a test binary only in a HIP -// build (tests/CMakeLists.txt gates it on VLLM_CPP_HIP) and no AMD GPU exists on -// the authoring machine. It is COMPILED everywhere, though: a non-HIP build -// object-compiles it as a bit-rot guard (see the CMake block next to the ROCm -// sources), so its types are checked on CI even with no ROCm installed. Compiled -// is not run. If you are the first person to run it: a failure here is -// far more likely to be a bug in the skeleton than a bug in your setup, and the -// most useful thing you can do is paste the output into +// RUN STATE, per issue #41: the W0 cases in this file ran green on community +// boards — gfx1151, gfx1103, gfx1100 and gfx1201 (5 cases, 1044 assertions in +// the posted tables). The two approach-(b) cases (alloc path / host-readable) +// have NEVER RUN: no AMD GPU exists on the authoring machine. The file is +// LINKED into a test binary only in a HIP build (tests/CMakeLists.txt gates it +// on VLLM_CPP_HIP) but COMPILED everywhere: a non-HIP build object-compiles it +// as a bit-rot guard (see the CMake block next to the ROCm sources), so its +// types are checked on CI even with no ROCm installed. Compiled is not run. If +// a new case fails on your board, that is far more likely a bug in the blind +// change than in your setup — paste the output into // https://github.com/mudler/vllm.cpp/issues/41 with the arch it printed. // // Deliberately plain C++ with no HIP header: every assertion goes through the @@ -126,6 +128,92 @@ TEST_CASE("the reference tier follows UnifiedMemory, which is the memory-safety } } +TEST_CASE("approach (b): the alloc path and UnifiedMemory() move together") { + if (NoDevice()) return; + Backend& rocm = vt::GetBackend(DeviceType::kROCM); + const bool integrated = vt::rocm::IntegratedDevice(0); + const bool managed = vt::rocm::ManagedAllocActive(0); + const bool unified = rocm.UnifiedMemory(); + // Printed unconditionally: this triple is the first thing a bring-up report + // on issue #41 should carry. + MESSAGE("ROCm device 0 integrated: ", integrated, " managed-alloc: ", managed, + " UnifiedMemory(): ", unified); + + if (!integrated) { + // DISCRETE (7900 XTX, R9700): the managed branch must be provably dead and + // the unified claim false — the byte-identical-to-W0 half of the (b) + // decision. A CPU fallback here would be memory corruption, not a slow + // path, so these two CHECKs are the memory-safety gate itself. + CHECK_FALSE(managed); + CHECK_FALSE(unified); + return; + } + // INTEGRATED. Every board measured on issue #41 (gfx1151 F6 attribute table, + // gfx1103 confirmation) reports ManagedMemory=1 + ConcurrentManagedAccess=1, + // so the managed branch is active and UnifiedMemory() is true by + // construction. An integrated device that probes NOT managed-capable would + // fail here: that is a hardware class the (b) fix does not cover, and a loud + // failure carrying the triple above is more useful than a silent skip — + // please post it on https://github.com/mudler/vllm.cpp/issues/41. + CHECK_MESSAGE(managed, + "integrated device without the managed-alloc branch: " + "ManagedMemory or ConcurrentManagedAccess probed 0 — post the " + "triple above on issue #41"); + CHECK_MESSAGE(unified == managed, + "UnifiedMemory() must be true EXACTLY when the managed branch " + "is active on an XNACK-less integrated part"); +} + +TEST_CASE("unified path: a kernel-written value is host-readable with no copy") { + if (NoDevice()) return; + Backend& rocm = vt::GetBackend(DeviceType::kROCM); + // On a discrete card a host dereference of Backend::Alloc memory is + // undefined behavior, so this gate only exists where UnifiedMemory() claims + // it is safe — which is exactly the claim under test. + if (!rocm.UnifiedMemory()) return; + + // Issue #41 F6's decisive experiment ("a kernel writes ...; the host reads + // ... back directly, no hipMemcpy"), turned into the standing gate. The + // kernel is the one op this backend registers (RmsNorm), so the file stays + // free of HIP: host WRITES the inputs directly (what a reference-tier CPU + // kernel does), the native device kernel reads them, and the host READS the + // device-written output directly. Same golden row as the native-RmsNorm case + // below, so a numeric mismatch here isolates COHERENCE, not arithmetic. + Queue q = rocm.CreateQueue(); + float* dx = static_cast(rocm.Alloc(2 * sizeof(float))); + float* dw = static_cast(rocm.Alloc(2 * sizeof(float))); + float* dout = static_cast(rocm.Alloc(2 * sizeof(float))); + REQUIRE(dx != nullptr); + REQUIRE(dw != nullptr); + REQUIRE(dout != nullptr); + + // Host writes, no Copy staging. + dx[0] = 3.0f; + dx[1] = 4.0f; + dw[0] = 2.0f; + dw[1] = 0.5f; + dout[0] = -1.0f; + dout[1] = -1.0f; + + const Device dev{DeviceType::kROCM, 0}; + Tensor tx = Tensor::Contiguous(dx, DType::kF32, dev, {1, 2}); + Tensor tw = Tensor::Contiguous(dw, DType::kF32, dev, {2}); + Tensor to = Tensor::Contiguous(dout, DType::kF32, dev, {1, 2}); + vt::RmsNorm(q, to, tx, tw, vt::RmsNormArgs{0.0f, false}); + rocm.Synchronize(q); + + // Host reads the device-written output directly, no Copy back. If this + // faults or reads the -1.0f sentinels, UnifiedMemory() lied — the exact + // failure mode approach (b) exists to make impossible. + CHECK(dout[0] == doctest::Approx(1.697056f)); + CHECK(dout[1] == doctest::Approx(0.565685f)); + + rocm.Free(dx); + rocm.Free(dw); + rocm.Free(dout); + rocm.DestroyQueue(q); +} + TEST_CASE("RmsNorm is registered natively, and the tier does not displace it") { if (NoDevice()) return; // Seam 3: the op table. One op today (src/vt/rocm/rocm_ops.hip).