[Vop3p]

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.

[Niko1221]

Heads-up: today we cleaned up this repository's commit history, and GitHub closed this PR automatically when main was force-pushed (that is the "base branch was force-pushed" event above). It is not a rejection.

GitHub cannot reopen it, because its branch is built on the old history. If you want to keep it going: start a new branch from the new main, cherry-pick your own commits onto it (git fetch upstream && git switch -c my-fix upstream/main && git cherry-pick <your commits>), and open a new PR. Or simply leave it: when we take your change for a release we port it onto the new main ourselves and keep you as its author.

Sorry for the noise, and thank you for the contribution.

[Niko1221]

Thank you, xjc10. The gfx103x attention kernel with 8 cells per step and DPP lane exchanges is in 0.1.40.2 with your authorship kept, with your bench notes. The HIP build compiles for gfx1030 and gfx1031. The PR was closed by the history rewrite; the commit is the credit.