[llvm] [AMDGPU] Fix a wave.reduce crash on a move of a non-immediate (PR #228819)

[Mian Miftah via llvm-commits]

mmiftahx wrote:

@arsenm @easyonaadit could you please review ?

Reproducer (HIP, gfx1100, -O2), which aborts on main:

```cpp
__global__ void k(unsigned *p, unsigned x) {
  __hip_atomic_fetch_add(p, __builtin_amdgcn_wave_reduce_min_u32(x, 0),
                         __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
```

The atomic optimizer reduces the umin result with an add reduction, and the uniform umin is lowered to a move of `x`.

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