llvm/llvm-project
在 GitHub 查看[AMDGPU] generates v_cndmask/lshlrev for uniform select between 0 and a power of 2
Open
#87,938 建立於 2024年4月7日
backend:AMDGPUgood first issue
倉庫指標
- Star
- (26,378 star)
- PR 合併指標
- (平均合併 1天 2小時) (30 天內合併 1,000 個 PR)
描述
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]