From ce06d47430930b17c5598073006bd7717f52a9d4 Mon Sep 17 00:00:00 2001 From: Justin Card Date: Wed, 12 Aug 2026 08:44:48 -0400 Subject: [PATCH 1/2] build(nix): add rocwmma to the ROCm dev shell so gfx12 builds again (#444) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Since 9302732f, rocm_paged_attn.hip includes whenever the target is gfx1200/gfx1201. The guard is on ARCH, not on availability, so targeting a gfx12 board fires the include whether or not rocWMMA exists. clr — the store path this shell uses as ROCM_PATH — does not ship it, so `nix develop .#rocm-shell` could not build main at all on a gfx12 board: src/vt/rocm/rocm_paged_attn.hip:8:10: fatal error: 'rocwmma/rocwmma.hpp' file not found Reproduced on pristine upstream/main at 0f2b12ed in a clean worktree with no other commits applied, so it is not a local-branch artifact. rocwmma is header-only, so it rides the existing overlay: rocmOverlayInputs symlinks each input's include/ into $ROCM_OVERLAY/include, which is already on CPATH. Adding it to the shell's packages list too keeps the two lists reading the same. Note the overlay is cached behind a .complete sentinel, so an existing shell needs $ROCM_OVERLAY removed once to pick this up. This is the NARROW half of #444 and does not close it. It fixes this shell only; every other rocWMMA-free environment targeting gfx12 still fails the same way. The issue asks for CMake-level detection that fails configure with a message naming the package, which is the real fix and stays open. Deliberately NOT done here: gating the include on __has_include. It works (verified under hipcc, which correctly reports the header absent), but VT_ROCWMMA_OK also selects the kernel BODY — PagedAttnPrefillWmmaWave at rocm_paged_attn.hip:900 casts every parameter to (void) and returns when the macro is undefined. Deciding that macro by availability rather than arch would compile a paged-attention kernel that silently writes nothing on a gfx12 board with no rocWMMA. A build failure is the correct outcome there; this commit supplies the missing dependency instead of hiding the symptom. Verified: with this change `nix develop .#rocm-shell` builds the HIP target on gfx1200 with 0 warnings, and `ctest -R 'rocm|cross_device'` is 4/4. Issue: https://github.com/mudler/vllm.cpp/issues/444 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --- flake.nix | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/flake.nix b/flake.nix index c023a18e6..5d1967ac4 100644 --- a/flake.nix +++ b/flake.nix @@ -67,8 +67,12 @@ # entry (idempotent — skipped if already populated). This is the # one Nix-specific step; a standard /opt/rocm install needs none of # it. + # rocwmma is header-only and rides the same overlay: rocm_paged_attn.hip + # includes whenever the target is gfx1200/gfx1201 + # (arch-gated, not availability-gated), and clr alone does not ship it, + # so a gfx12 build fails at that include without this. See issue #444. rocmOverlayInputs = - [ rocm.hipblas rocm.hipblaslt rocm.hipblas-common ]; + [ rocm.hipblas rocm.hipblaslt rocm.hipblas-common rocm.rocwmma ]; in { default = pkgs.mkShell { packages = commonPackages ++ [ pkgs.gcc ]; @@ -122,6 +126,7 @@ rocm.hipblas rocm.hipblaslt rocm.hipblas-common + rocm.rocwmma rocm.rocminfo ]; From f04a761b9833a8240bc852cf53dcacfc6151d663 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Fri, 14 Aug 2026 08:05:33 +0000 Subject: [PATCH 2/2] docs(BACKEND-ROCM): link #444 from the row's spec, the half this PR was split out to close (#444) Review repair on top of joral's commit; theirs is untouched. This PR exists because #473's review found #444 cited in a commit but absent from the roadmap table and from every file under `.agents/`. #473 landed the roadmap row. Nothing landed the spec half, so AGENTS.md's "three places that must agree" was satisfied in two -- by the very PR split out to fix it. `.agents/specs/rocm-gfx1200-m2-correctness.md` is the right home: it is this board's record, it already names the two sibling ROCm bug issues (#201, #132) in the same sentence, and it already documents `nix develop .#rocm-shell`. #444 was the one ROCm bug issue with no spec presence. Added alongside its siblings, described as what it is -- a build-environment defect on the same board, not a correctness one -- and scoped: this PR is the narrow half, #444 stays open. Two corrections to the same file while in it. The environment note claimed the shell was a "local, uncommitted `flake.nix` addition". Stale since `f93a1290a` (2026-08-10) committed it. A reader was being told the documented ROCm dev path did not exist in the tree. And the caveat that matters operationally: the overlay is cached behind a `$ROCM_OVERLAY/.complete` sentinel, so anyone who has already entered `rocm-shell` keeps a rocWMMA-free overlay after this merges and hits the IDENTICAL `'rocwmma/rocwmma.hpp' file not found` the PR fixes. joral discloses this in the PR body; the tree did not, and a PR description is not a record. The failure mode is indistinguishable from the bug, so it is written where someone hitting the error will look. Not fixed here: the sentinel itself. Keying it on the input list would self-heal, but that changes shell behaviour for every user and belongs in its own change rather than riding a docs repair. Verified: check-agent-record OK, check-public-doc-tables OK. No public-document trigger -- no lifecycle change, and `docs/ROCM.md` documents no nix path. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-opus-5 [Claude Code] --- .agents/specs/rocm-gfx1200-m2-correctness.md | 18 ++++++++++++++---- 1 file changed, 14 insertions(+), 4 deletions(-) diff --git a/.agents/specs/rocm-gfx1200-m2-correctness.md b/.agents/specs/rocm-gfx1200-m2-correctness.md index f38af0d06..f71cc22ec 100644 --- a/.agents/specs/rocm-gfx1200-m2-correctness.md +++ b/.agents/specs/rocm-gfx1200-m2-correctness.md @@ -9,7 +9,12 @@ near-tie prompt is not a defect on either of our backends — see Outcome. gfx1200 M0/M1 report, first M2 attempt on this board. Distinct defect from #201 (`hipblasGemmEx` overload) and #132 (`-O0` teardown deadlock), neither of which reproduced here (and both since fixed on main, `ce134e1d`, picked up -when this investigation rebased onto `origin/main` 2026-08-10). +when this investigation rebased onto `origin/main` 2026-08-10). Also distinct +from [#444](https://github.com/mudler/vllm.cpp/issues/444) (rocWMMA headers +absent from the dev shell, so a gfx12 build configures and then fails midway) — +a build-environment defect on this same board, not a correctness one. Its narrow +half, adding `rocwmma` to `rocm-shell`, lands in PR #638; #444 stays open for +the rest. **Board:** AMD Radeon RX 9060 XT (`gfx1200`, Navi 44, RDNA4, discrete — no reference tier). ROCm 7.2.3 (nixpkgs `rocmPackages.*`), hipClang/Clang 22.0.0. M0/M1 independently MET on this board (separate report to #41). @@ -216,9 +221,14 @@ doesn't hold still on it. ## Environment note (reproducibility) -This board is accessed through `nix develop .#rocm-shell` (local, uncommitted -`flake.nix` addition — see the M0/M1 report) for the ROCm/HIP `build-hip` -build itself; the oracle containers are unrelated to that and need no ROCm +This board is accessed through `nix develop .#rocm-shell` for the ROCm/HIP +`build-hip` build itself. That shell has been committed since `f93a1290a` +(2026-08-10); the "local, uncommitted `flake.nix` addition" this paragraph used +to describe is stale. One caveat if you already have the shell: the overlay is +cached behind a `$ROCM_OVERLAY/.complete` sentinel, so an existing checkout +keeps its rocWMMA-free overlay after #638 merges and fails with the identical +`'rocwmma/rocwmma.hpp' file not found`. Remove `$ROCM_OVERLAY` to +re-materialise it. The oracle containers are unrelated to that and need no ROCm toolchain on the host beyond the kernel driver + `/dev/kfd`/`/dev/dri`. `ROCM_PATH` for `build-hip` is nixpkgs' `clr` output (has `lib/cmake/hip-lang`); hipBLAS/hipBLASLt/hipblas-common are merged into a