[morousg]

Phase 3 of the RDNA4 port: the scalar attention DPPs and the FP8 KV cache run on HIP (RX 9060, gfx1200).

Stacked on #394 and #395: this branch contains their commits. Review those first; once they merge, this diff shrinks to the last commit (7c78320).

Changes

• Backends: softmax.h, flash_attention.h and flash_decode.h now build their DPPs under __NVCC__ || __HIPCC__. HIP launches use getHIPStream()/hipGetLastError(), the CUDA launch lines are unchanged in #else, and HIP kernels call attnRequireWave32(). flash_attention_mma.h stays CUDA-only and is still harmless to include on HIP.
• Wave size: the 2 shuffle sites in these files now use attnShflXor, and the 32s that mean "lanes per wave" use ATTN_WAVE_SIZE. Every remaining 32 is a tile size (BLOCK_N) and is untouched. The 5 shuffle sites in the MMA file are unchanged.
• FP8 (OCP e4m3 on both backends): HIP gets its own branch first (__HIPCC__ && __has_include(<hip/hip_fp8.h>) → __hip_fp8_e4m3, not the FNUZ __hip_fp8_e4m3_fnuz of gfx942); the CUDA condition is still exactly __has_include(<cuda_fp8.h>). The type is fk::attn_fp8_e4m3, decoded by a single attnFp8ToF32() used by Fp8TokenDequantRead, flash decode and the tests. attnFp8IsOcpE4M3() checks known bytes (0x7E = 448, 0x38 = 1, 0x80 = -0). On this ROCm (10.1), device and host decoding and encoding give OCP values, and encoding saturates at 448 like CUDA's satfinite.
• bf16 and IOp traits: Bf16AttentionRead, its makeAttentionRead overload and the isRawBf16Read/isFp8KVRead/isInt8KVRead traits moved from flash_attention_mma.h to flash_attention.h, so decode doesn't need the CUDA-only header. They have the same names and code, with bf16 spelled attn_bf16/attn_bf162 (__nv_* under nvcc). flash_decode.h still includes the MMA header.
• Tests: utest_flash_attention, utest_softmax, utest_flash_decode and utest_attention_wave now use Ptr instead of raw cudaMalloc/cudaMemcpy[2D], and are marked with the new __ONLY_GPU__.
  • cmake/tests/discover_tests.cmake: an ONLY_GPU substring suppresses only the _cpp target. No existing file contains that substring.
  • AGENTS.md documents the marker and the HIP status of attention/.
  • utest_flash_attention adds an fp8 format check (testFp8Format).
  • Test cases and tolerances are unchanged; HIP uses the same tolerances as CUDA.

CUDA-visible changes (same types and behavior; all compiled with nvcc 13.4)

They don't change behavior, but they aren't token-identical:

1. The entry points' stream parameter is spelled Stream_<defaultParArch>&, which is Stream_<ParArch::GPU_NVIDIA>& under nvcc.
2. low_precision::attnToF32/attnFromF32 are now FK_HOST_DEVICE_CNST instead of FK_DEVICE_CNST. The host packing helpers (quantizeKVCacheHost, quantizeKVCacheFp8Host) already called them from host code. nvcc accepts that for a constexpr __device__ function, but HIP's clang doesn't.
3. Fp8TokenDequantRead::exec decodes through attnFp8ToF32, which sets __x and converts. The old device branch reinterpreted the byte as __nv_fp8_e4m3; the old host branch already used the new code.
4. The wave-size spellings (% ATTN_WAVE_SIZE instead of & 31, and so on) and attnShflXor, which is __shfl_xor_sync(0xffffffffu, …) on CUDA.
5. The moved bf16 read and traits (see above).

Verified on this machine (RX 9060, gfx1200, ROCm 10.1)

• Full HIP suite with default arches (gfx1200;gfx1201): 51/51, the 47 on main plus utest_attention_wave_hip, utest_softmax_hip, utest_flash_attention_hip and utest_flash_decode_hip. Same unique warnings as main.

• Measured max errors on HIP:

  Test	Max error	Tolerance
  Scalar FA (dense, int8 KV, fp8 KV, prologue/epilogue)	6e-8 to 7e-7	5e-6 / 1e-6
  Softmax	≤ 8e-8	1e-6
  Decode, every case	≤ 5e-8	1e-5

  Decode covers fp32, bf16 (including d32/d256), int8 and fp8, with empty splits and repeated calls.

• Mutation checks: making the shuffle reduction stop one step early fails every dense case (max error 0.05 to 0.28). Decoding fp8 at twice the value fails the format check. The KV cases can't catch that one, because the oracle decodes with the same function, which is why the format check exists.

• nvcc 13.4 + g++-13, sm_86 (PyPI wheels, no GPU): full build with 0 failures and 0 warnings; the 49 _cpp tests pass, and the _cpp list is the same as RDNA4 phase 2: CUDA tests for FP8 KV, flash decode and the mma kernel (CUDA-side, unverified locally) #395's.

Needs CUDA CI

Running utest_flash_attention_cu, utest_softmax_cu, utest_flash_decode_cu, utest_flash_attention_mma_cu and utest_attention_wave_cu on the A2000, after the changes above.

🤖 Generated with Claude Code

https://claude.ai/code/session_014d2rht2WMSopuodhtXYZzx

[Copilot]

🟡 Changes recommended

The MMA test can break older CUDA architecture builds, and numerical and workspace-reuse checks can falsely pass.

3 open issues

• Guard unsupported MMA kernels from compilation · New
• Reject NaN differences in error accumulation · New
• Reset output before verifying counter reset · New

What changed in this PR

Ports scalar attention kernels to HIP/RDNA4 while preserving CUDA support, including the stacked architecture and test changes from #394 and #395.

Changes:

• Adds HIP attention launches, wave32 helpers, and shared low-precision types.
• Expands attention tests and replaces CUDA-specific buffer handling.
• Updates HIP architecture selection, GPU-only test discovery, and documentation.

File	Description
utests/​algorithm/​attention/​utest_softmax.h	Uses portable buffers and enables HIP testing.
utests/​algorithm/​attention/​utest_flash_decode.h	Adds decode format and workspace-reuse tests.
utests/​algorithm/​attention/​utest_flash_attention.h	Adds FP8 coverage and portable buffers.
utests/​algorithm/​attention/​utest_flash_attention_mma.h	Adds CUDA bf16 MMA tests.
utests/​algorithm/​attention/​utest_attention_wave.h	Tests wave reductions and exchanges.
utests/​algorithm/​attention/​attention_test_common.h	Shares reference calculations and buffer helpers.
include/​fused_kernel/​algorithms/​attention/​softmax.h	Adds HIP launches and host-callable conversions.
include/​fused_kernel/​algorithms/​attention/​flash_decode.h	Enables HIP and fixes unsupported packed-load widths.
include/​fused_kernel/​algorithms/​attention/​flash_attention.h	Adds HIP support and shared FP8/bf16 types.
include/​fused_kernel/​algorithms/​attention/​flash_attention_mma.h	Moves shared reads and traits out.
include/​fused_kernel/​algorithms/​attention/​attention_wave.h	Adds wave32 helpers and enforcement.
cmake/​tests/​discover_tests.cmake	Supports GPU-only test markers.
cmake/​libs/​hip/​hip.cmake	Aligns HIP package architecture targets.
cmake/​libs/​hip/​hip_init.cmake	Seeds architectures before HIP initialization.
cmake/​libs/​hip/​archs.cmake	Applies the selected HIP architectures.
AGENTS.md	Documents HIP attention and test configuration.

_{Review effort: Balanced}