RDNA4 phase 1: make HIP_ARCH work, attention wave-size helpers

[morousg]

Phase 1 of the RDNA4 (gfx12) port of algorithms/attention/.

Changes

HIP architectures (cmake/libs/hip/)

• HIP_ARCH didn't do anything on main. enable_language(HIP) runs before archs.cmake and always sets CMAKE_HIP_ARCHITECTURES to the build machine's GPU, so set_target_hip_arch_flags always took that branch. On top of that, hip::device compiled each target's C++ sources (tests/main.cpp) for GPU_TARGETS, which hip-config also detects. Checked on main: -DHIP_ARCH=gfx1201 still produced --offload-arch=gfx1200 on all 165 compile lines.
• Now HIP_ARCH seeds CMAKE_HIP_ARCHITECTURES before enable_language(HIP), and GPU_TARGETS follows CMAKE_HIP_ARCHITECTURES, so main.cpp gets the same arches as the HIP sources.
• Default (after 1cf881b): the GPU of the build machine, as CMake detects it. -DHIP_ARCH="gfx1200;gfx1201" builds for Navi 44 (RX 9060) and Navi 48 (RX 9070/XT), which can't load each other's code objects. HIP_ARCH applies on the first configure of a build directory, and an explicit -DCMAKE_HIP_ARCHITECTURES wins over it. gfx12-generic also works through CMAKE_HIP_ARCHITECTURES.
• Checked after 1cf881b + 7ce748d (docs): default → gfx1200 on all 168 compile lines, HIP_ARCH=gfx1200;gfx1201 → both, a reconfigure with another HIP_ARCH keeps the first value, explicit CMAKE_HIP_ARCHITECTURES=gfx12-generic wins.

Wave size (attention/attention_wave.h, new; nothing includes it yet)

• ATTN_WAVE_SIZE = 32, attnShflXor(value, laneMask) (__shfl_xor_sync(0xffffffffu, …) on CUDA, __shfl_xor(…, 32) on HIP), and attnRequireWave32().
• ROCm 10.1's clang (LLVM 24) no longer defines __AMDGCN_WAVEFRONT_SIZE__, and neither warpSize nor __builtin_amdgcn_wavefrontsize() is a constant expression, so a wave64 compile can't be rejected at build time. The header has a #error for compilers that still define the macro. Otherwise kernels call attnRequireWave32(), which folds to nothing at -O3 in wave32 and becomes one s_trap in wave64. HIP's __shfl_xor_sync would need a 64-bit mask, so the HIP path uses __shfl_xor.
• utests/algorithm/attention/utest_attention_wave.h: 2 waves per block, butterfly sum plus a single xor exchange (lanes must not cross waves). It gets _hip, _cu and _cpp targets; the _cpp target only prints that there's nothing to test on the CPU.

AGENTS.md: the HIP_ARCH description is updated.

Verified on the RX 9060 (gfx1200, ROCm 10.1)

• Full ctest suite: 47/47 on clean main → 48/48 on this branch (the same 47 plus utest_attention_wave_hip), built for gfx1200;gfx1201. Same unique warnings as main.
• -DHIP_ARCH=gfx1201: full build succeeds (it can't run here).
• The test binary bundles both gfx1200 and gfx1201 code objects.
• Wave64 (-DCMAKE_HIP_FLAGS=-mwavefrontsize64): the test traps on the GPU (hardware exception, exit 134) instead of computing wrong results.
• Build time for the full suite with 2 arches: 105 s on this machine (16 threads).

Only CUDA CI can verify

• utest_attention_wave_cu (nvcc build and run on the A2000) and utest_attention_wave_cpp (g++/clang/MSVC). No existing CUDA-visible code changed: the only CUDA-visible additions are the new unused header and this test.

clang-format isn't installed on this machine, so the new files are formatted by hand to .clang-format (4 spaces, at most 120 columns).

🤖 Generated with Claude Code

https://claude.ai/code/session_014d2rht2WMSopuodhtXYZzx

[Copilot]

🟡 Changes recommended

Reconfiguration can overwrite an explicit architecture selection and build binaries for the wrong GPU.

1 open issue

• Reconfiguration overwrites explicit CMAKE_HIP_ARCHITECTURES value · New

What changed in this PR

Prepares FKL’s attention code for RDNA4 by updating HIP target selection and adding standalone wave32 helpers.

Changes:

• Defaults HIP builds to gfx1200 and gfx1201 and aligns C++ and HIP targets.
• Adds wave32 validation and shuffle helpers with GPU tests.
• Documents architecture selection.

File	Description
utests/​algorithm/​attention/​utest_attention_wave.h	Tests reductions and exchanges within two waves.
include/​fused_kernel/​algorithms/​attention/​attention_wave.h	Adds wave32 guards and shuffle helpers.
cmake/​libs/​hip/​hip.cmake	Aligns HIP package GPU targets.
cmake/​libs/​hip/​hip_init.cmake	Seeds architectures before enabling HIP.
cmake/​libs/​hip/​archs.cmake	Simplifies target architecture assignment.
AGENTS.md	Documents HIP architecture options.

_{Review effort: Balanced}