[llvm] [AMDGPU] Fix wave.reduce.fsub with the DPP strategy (PR #228817)

[Mian Miftah via llvm-commits]

mmiftahx wrote:

hi @arsenm @easyonaadit could you please review this ?

On gfx1201 in wave64, `__builtin_amdgcn_wave_reduce_fsub_f32(v, 2)` with lane values 1 to 64 returns 0 instead of -2080; with the patch it returns -2080.

I also ran all 26 `__builtin_amdgcn_wave_reduce_*` builtins on gfx1201 against a host reference, with the Iterative and DPP strategies, in wave32 and wave64, with full and partial masks. Before the patch only float fsub with DPP was wrong; with it, all 416 cases match.

https://github.com/llvm/llvm-project/pull/228817