ROCm: main does not build for gfx1200/gfx1201 — rocm_paged_attn.hip includes rocwmma unconditionally on arch, not availability

[joral]

Summary

main does not build for ROCm in this repo's own nix develop .#rocm-shell
when the target arch is gfx1200/gfx1201. The build stops at:

src/vt/rocm/rocm_paged_attn.hip:8:10: fatal error: 'rocwmma/rocwmma.hpp' file not found
    8 | #include <rocwmma/rocwmma.hpp>

Introduced by 9302732f ("feat(gemma4/rocm): FP8 resident MoE + SharedK-WMMA",
2026-08-11). Reproduced on pristine upstream/main at 0f2b12ed in a clean
worktree with no other commits applied, so this is not a local-branch artifact.

Reproduce

git worktree add /tmp/mainhip upstream/main
cd /tmp/mainhip
nix develop .#rocm-shell -c cmake -S . -B build-hip -DVLLM_CPP_HIP=ON \
  -DVLLM_CPP_HIP_ARCHITECTURES=gfx1200 -DROCM_PATH=$ROCM_PATH -DCMAKE_BUILD_TYPE=Release
nix develop .#rocm-shell -c cmake --build build-hip --target vllm -j8

Board: AMD Radeon RX 9060 XT, gfx1200, Navi 44, RDNA4, discrete. ROCm 7.2.3,
hipClang/Clang 22.0.0, NixOS.

Why it happens

The include is gated on architecture, not on availability
(src/vt/rocm/rocm_paged_attn.hip:6-10):

// rocWMMA only on gfx12 (R9700). Fat builds also emit gfx1036 — must not include there.
#if defined(__gfx1200__) || defined(__gfx1201__)
#include <rocwmma/rocwmma.hpp>
#define VT_ROCWMMA_OK 1
#endif

Targeting gfx1200 satisfies the guard, so the include fires whether or not
rocWMMA exists. On a stock ROCm install it ships alongside HIP and this is
invisible. In this repo's dev shell it is absent — the clr derivation the shell
provides carries no rocwmma/, nothing in /nix/store matches *rocwmma*, and
flake.nix does not mention it. Nothing in CMakeLists.txt or cmake/ detects
or requires it either, so configure succeeds and the failure lands mid-build.

docs/ROCM.md, flake.nix and .agents/environment.md do not list rocWMMA as a
prerequisite, so a contributor following the documented ROCm path has no warning.

Please do NOT fix this with __has_include alone

__has_include(<rocwmma/rocwmma.hpp>) does work here — verified under hipcc,
it correctly reports absent — but swapping the arch guard for it would turn a
loud build failure into a silent correctness bug.

VT_ROCWMMA_OK also selects the kernel body. At
rocm_paged_attn.hip:900, PagedAttnPrefillWmmaWave compiles to a no-op when
the macro is undefined — every parameter cast to (void) and an immediate
return; (that path exists for fat builds that also emit gfx1036, where the
kernel is never dispatched). If availability rather than arch decided the macro,
a gfx1200 build without rocWMMA would compile cleanly and produce a paged
attention kernel that writes nothing.

grep -rn VT_ROCWMMA_OK src/ shows the macro is confined to
rocm_paged_attn.hip, so no host-side dispatch code can currently gate on it
either.

Safer shapes, in rough order of preference:

1. Detect in CMake and fail configure with a message naming the package,
   when a gfx12 target is requested and rocWMMA is missing. Loud, early, and
   keeps the arch guard's meaning intact.
2. Add rocWMMA to flake.nix's rocm-shell so the documented path works out
   of the box. Complements (1) rather than replacing it — it fixes this shell,
   not other environments.
3. Detect availability and route dispatch away from the WMMA kernel, so the
   no-op body is unreachable rather than merely unused. More work, and only worth
   it if a rocWMMA-free gfx12 build is meant to be supported at all.

At minimum, documenting rocWMMA as a gfx12 build prerequisite in docs/ROCM.md
would stop the next person losing the same hour.

Scope

• Only src/vt/rocm/rocm_paged_attn.hip includes rocWMMA (grep -rln rocwmma src/).
• Blocks any ROCm build targeting gfx1200/gfx1201 in an environment without
  rocWMMA, including this repo's own dev shell.
• Other gfx targets are unaffected: the guard is false, so the include never
  fires.
• No open issue mentions rocWMMA (searched rocwmma in:title,body).

[VikashLoomba]

Data point isolating this to RDNA4: current main (7020de93, which includes the rocwmma-carrying paged-attention code) builds clean end-to-end on gfx1100 (4× RX 7900 XTX, ROCm 7.14, hipcc auto-targets gfx1100) — full tree including rocm_paged_attn.hip, 0 errors, vllm-cli links. So the unconditional rocwmma include is tolerated on gfx1100's ROCm 7.14 toolchain headers but not on your gfx1200/gfx1201 setup — consistent with the availability-not-arch guard being the right fix (rocwmma headers/guard macros differ across the RDNA3/RDNA4 + ROCm-version matrix). Happy to run a candidate branch on gfx1100 as a regression check once there's a PR.

[VikashLoomba]

Correction to my SHA label: the clean gfx1100 build was on current main e6f73b31 (not 7020de93 as I wrote — my checkout had fast-forwarded and I cited a stale SHA). Substance unchanged: the tree containing the rocwmma-carrying rocm_paged_attn.hip builds clean end-to-end on gfx1100 / ROCm 7.14.

[joral]

Partial fix in #473 — it adds rocwmma to flake.nix's rocm-shell, which is
option 2 above. That PR is about the hipGraph capture seam (W1 of #332); the
flake change rides along because it is what made a gfx1200 build possible at all
to run W1's gates. Happy to split it into its own PR if that reads cleaner.

This does not close the issue. It fixes one dev shell. Any other
rocWMMA-free environment targeting gfx1200/gfx1201 still fails at the same
include, because the guard is still on architecture rather than availability.
Option 1 — CMake detection that fails configure with a message naming the
package — is still the fix worth having, and the caution above stands: gating
the include on __has_include alone would compile the no-op
PagedAttnPrefillWmmaWave body and produce attention that silently writes
nothing.

Verified after the change: nix develop .#rocm-shell builds the HIP target on
gfx1200 with 0 warnings, and ctest -R 'rocm|cross_device' is 4/4. Note the
overlay caches behind a .complete sentinel, so an existing shell needs
$ROCM_OVERLAY removed once to pick it up.