llvm/llvm-project
Voir sur GitHub[AMDGPU] generates v_cndmask/lshlrev for uniform select between 0 and a power of 2
Open
#87 938 ouverte le 7 avr. 2024
backend:AMDGPUgood first issue
Métriques du dépôt
- Stars
- (26 378 stars)
- Métriques de merge PR
- (Merge moyen 1j 2h) (1 000 PRs mergées en 30 j)
Description
Here is my test kernel:
static __global__ void test(const int32_t input, int32_t * result) {
const int data = input ? 32 : 0;
int output;
asm volatile("s_mov_b32 %0, %1" : "=s"(output) : "s"(data));
result[threadIdx.x] = output;
}
and here's the asm it compiles to for gfx900 on rocm 6.0.0, where it fails to assemble because of an invalid operand.
s_load_dword s2, s[4:5], 0x0
s_load_dwordx2 s[0:1], s[4:5], 0x8
v_lshlrev_b32_e32 v0, 2, v0
s_waitcnt lgkmcnt(0)
s_cmp_lg_u32 s2, 0
s_cselect_b64 s[2:3], -1, 0
v_cndmask_b32_e64 v1, 0, 1, s[2:3]
v_lshlrev_b32_e32 v1, 5, v1
;;#ASMSTART
s_mov_b32 s2, v1
;;#ASMEND
v_mov_b32_e32 v1, s2
global_store_dword v0, v1, s[0:1]
If either of the values being selected is changed so that either one isn't 0, or the other isn't a power of 2, instead it compiles to this, which works.
s_load_dword s2, s[4:5], 0x0
s_load_dwordx2 s[0:1], s[4:5], 0x8
v_lshlrev_b32_e32 v0, 2, v0
s_waitcnt lgkmcnt(0)
s_cmp_eq_u32 s2, 0
s_cselect_b32 s2, 0, 33
;;#ASMSTART
s_mov_b32 s2, s2
;;#ASMEND
v_mov_b32_e32 v1, s2
global_store_dword v0, v1, s[0:1]