mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-07-11 17:51:40 +00:00
a49a273a7f
The packed score-shift (v_pk_fma_f32) + packed alu1 o_acc rescale (v_pk_mul_f32) and the fmha_alu0 rowmax/shift split, all added while bottleneck-hunting after the 1882 milestone, are a net REGRESSION on the canonical fp8 prefill shape, not a win. Same-session, same-harness, GPU2, 3-run median (b1 sq=sk=75600 hq=hk=5 d128 fp8 non-causal): * scalar softmax (this commit): ~1877 TF/s * packed shift + packed alu1: ~1825 TF/s (-3%) The softmax score-shift/rescale are hidden under the 8-wave ping-pong overlap, so collapsing their FMAs/muls does not shorten the critical path -- it only perturbs the scheduler and loses. The earlier "+8%" that motivated the packs was a confounded GPU0 measurement (GPU0 sustains a lower boost clock than GPU2 on this shared box). Reverts the fmha_alu0 rowmax/shift split entirely and flips UA_FA4_PACKED_SHIFT / UA_FA4_PACKED_ALU1_RESCALE to default OFF (code kept behind the macros, now documented as measured regressions, so they are not re-attempted). The production config is unchanged: kv128 + cooperative K/V load + single-sp + wide-MMA + conditional-rescale, scalar softmax -- i.e. the ~1882 baseline. Co-authored-by: Cursor <cursoragent@cursor.com>