[llvm] [AMDGPU] Add an SCC def to the wave.reduce pseudos (PR #228723)

Mian Miftah via llvm-commits llvm-commits at lists.llvm.org
Sat Oct 3 08:37:05 PDT 2026


mmiftahx wrote:

hi @arsenm @easyonaadit could you please review ?

Reproducer (OpenCL, gfx1100, -O2):

```c
kernel void k(global int *out, int a, int b, int u) {
  bool c = a == b;
  int r = __builtin_amdgcn_wave_reduce_add_u32(u, 0);
  *out = c ? r : 1;
}
```

On main this stores `r` even when `a != b`, because the `s_cselect` reads SCC from `s_bcnt1` instead of from a compare.

The first commit adds the test, so the second one shows the change.

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


More information about the llvm-commits mailing list