build(nix): add rocwmma to the ROCm dev shell so gfx12 builds again (#444) - #638
Merged
Conversation
…udler#444) Since 9302732, rocm_paged_attn.hip includes <rocwmma/rocwmma.hpp> 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 0f2b12e 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 mudler#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: mudler#444 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code]
5 tasks
Contributor
|
@mudler @richiejp This focused Nix fix looks good to merge. It adds the header-only rocWMMA package to both the overlay inputs and shell packages, matching the gfx12 compile path that includes |
Collaborator
|
LGTM |
…s PR was split out to close (mudler#444) Review repair on top of joral's commit; theirs is untouched. This PR exists because mudler#473's review found mudler#444 cited in a commit but absent from the roadmap table and from every file under `.agents/`. mudler#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 (mudler#201, mudler#132) in the same sentence, and it already documents `nix develop .#rocm-shell`. mudler#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, mudler#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]
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
4 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Split out of #473 per
localai-org-maint-bot's review there: PR #473 (thevt::BackendhipGraph capture seam, W1 of #332) is meant to stay scoped to the graph-capture seam, and this dev-shell fix is unrelated scope that belongs with #444 instead.What changed
Since
9302732f,rocm_paged_attn.hipincludes<rocwmma/rocwmma.hpp>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 asROCM_PATH— does not ship it, sonix develop .#rocm-shellcould not buildmainat all on a gfx12 board:Reproduced on pristine
upstream/mainat0f2b12edin 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:
rocmOverlayInputssymlinks each input'sinclude/into$ROCM_OVERLAY/include, which is already onCPATH. Adding it to the shell's packages list too keeps the two lists reading the same. Note the overlay is cached behind a.completesentinel, so an existing shell needs$ROCM_OVERLAYremoved 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 missing 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), butVT_ROCWMMA_OKalso selects the kernel BODY —PagedAttnPrefillWmmaWaveatrocm_paged_attn.hip:900casts 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.Evidence
With this change
nix develop .#rocm-shellbuilds the HIP target on gfx1200 with 0 warnings, andctest -R 'rocm|cross_device'is 4/4.Issue: #444
🤖 Generated with Claude Code