#849 HIP gfx103x: PR #540's attention kernel with 8 cells per step and DPP lane exchanges (bit-exact, prompt +6-12% on RX 6900 XT)

本文

## Summary

`attn_chunk_kernel_pre75` (PR #540, bit for bit the default attention kernel) is built only into the experimental CUDA
build and run below sm_75. On HIP the default kernel runs, whose score loop is bound by the LDS pipe on gfx1030: its
12 heads × 5 `__shfl_xor_sync` per cell all compile to `ds_bpermute_b32` (ROCm 10.0's hipcc emits no DPP for them), sharing the LDS
pipe with the shared-memory q reads.

This builds the pre75 kernel for HIP wave32 (not the gfx906 wave64 backend) and runs it on gfx103x by default, with two
template parameters whose defaults keep CUDA exactly as it is:
- `NC_` (cells per step): CUDA 2 as before; HIP 8. On gfx1030 every step up in NC was faster although occupancy went
  down from NC 1 (NC = 1 / 2 / 4 / 8: 101 / 142 / 149 / 144 VGPRs; 9 / 7 / 6 / 7 waves/SIMD), so the q reads, not
  occupancy, bound it there - the opposite of the V100, where 4 was slower than 2.
- `DPP_`: the butterfly's lane exchanges (in `reduce12` and the softmax's max / sum) in registers - DPP `row_xmask`
  within a 16-lane row, `v_permlanex16` across the two rows - instead of `ds_bpermute`. The same values move in the same
  pairing order, so every sum is unchanged bit for bit. CUDA keeps `__shfl_xor_sync`.

`pre75_attn()` on HIP: on for gfx103x (per device), `STRATA_ATTN_PRE75=0` off (as before), `=1` also on another wave32
card. The CUDA branch of `pre75_attn()` and the non-experimental CUDA build are unchanged.

## Measurements (one RX 6900 XT, ROCm 10.0, Qwen3.8-Flash-Next GSQ-RCO IQ3_S, KV int8)

Prompt tok/s, cold server per run, same binary with `STRATA_ATTN_PRE75=0` vs default:

| | 9.4K | 34.7K | 105.8K |
| --- | ---: | ---: | ---: |
| default kernel | 441 | 467 | 460 |
| this change | 459 (+4.2%) | 494 (+5.7%) | 488 (+6.1%) |

With #835 (FP16 GEMMs) also applied, attention is a larger share of the prompt and the gain grows:
746 / 917 / 928 -> 801 / 1,022 / 1,041 (+7.4 / +11.4 / +12.2%). On that base, step by step (same three prompts):

| pre75 variant | 9.4K | 34.7K | 105.8K |
| --- | ---: | ---: | ---: |
| NC 1 | 763 | 950 | 960 |
| NC 2 | 779 | 982 | 997 |
| NC 4 | 787 | 997 | 1,008 |
| NC 8 | 792 | 1,006 | 1,025 |
| NC 8 + DPP | 801 | 1,022 | 1,041 |

Decode: measured unchanged within 1.3%, slightly faster (the A/B in "Validation and docs" below).

## Related: #835

#835 runs the prompt path's 16-bit GEMMs FP16 in and out on gfx103x. The two changes are independent (different files,
either applies alone) and compound: with both, a prompt reads at 1.82x / 2.19x / 2.26x today's main on this card
(439 / 466 / 461 -> 801 / 1,022 / 1,041 tok/s at 9.4K / 34.7K / 105.8K). This one is bit-exact; #835 is not (its
distribution check is in its README).

## Exactness

Teacher-forced log-probabilities (`STRATA_LOGPOS`, top 256, 3 chats of 9.4K / 34.7K / 105.8K tokens read by the batched
prompt path, 1,578 scored positions): this change against the default kernel - KL 0, argmax and top-10 agreement 100%,
at every NC above and with DPP. ISA check (`-Rpass-analysis=kernel-resource-usage`, `-S`): the DPP variant has 0
`ds_bpermute` (154 without) and 103 LDS instructions (257).

## Not tested

CUDA builds (no CUDA on this machine; the CUDA path is unchanged by construction: template defaults equal the old
constants and `shfl_x<false>` is `__shfl_xor_sync`). Other wave32 AMD architectures (opt-in with `STRATA_ATTN_PRE75=1`).

## Validation and docs

- **ctest** (`-DSTRATA_BUILD_TESTS=ON`, the same options as the engine build, gfx1030, ROCm 10.0): 60 of 65 pass, the same 60 as `6f32ec0` without this change. Skipped: `hip_prompt_attn_wmma` (no matrix cores on gfx1030), `hip_prefill_hipblaslt_gemm` (no hipBLASLt table). Failed for reasons outside the engine, identically with and without the change: `ple_parity` (a Q2_0 PLE file that is not on this machine), `expert_multi_test` (this Zen 3 CPU has no AVX-512), `platform_memory_test` (`mlock` at the shell's default `ulimit -l`).
- **Decode A/B** (one card, greedy, 256-token answers to three ~4K prompts, `STRATA_DECODE_TIMING=1`, two rounds alternating): the window is 0.3-0.6 ms shorter with this change (56.0 -> 55.7, 42.2 -> 41.9, 42.5 -> 42.0, 44.2 -> 43.6 ms per window; the two rounds of one arm differ by at most 0.2 ms), decode 43-53 -> 43-56 tok/s. Details in the README below.
- **Docs:** `docs/OLDER_GPUS.md`'s AMD table row for RX 6800 / 6900 (the kernel is the default on gfx1030, the switch, the gain) and its A/B-switches line; the measurements, compiler output and exactness check are in `bench/results/2026-10-04-rdna2-pre75-attention/README.md` with the configuration used.

Developed with an AI coding assistant; every number above was measured on one RX 6900 XT (gfx1030) / Ryzen 5 5600X, ROCm 10.0.