[AMDGPU] generates v_cndmask/lshlrev for uniform select between 0 and a power of 2

[Engininja2]

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]

[llvmbot]

@llvm/issue-subscribers-backend-amdgpu

Author: None (Engininja2)

Here is my test kernel: ```cpp 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.
```asm
	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]

[llvmbot]

Hi!

This issue may be a good introductory issue for people new to working on LLVM. If you would like to work on this issue, your first steps are:

1. Check that no other contributor has already been assigned to this issue. If you believe that no one is actually working on it despite an assignment, ping the person. After one week without a response, the assignee may be changed.
2. In the comments of this issue, request for it to be assigned to you, or just create a pull request after following the steps below. Mention this issue in the description of the pull request.
3. Fix the issue locally.
4. Run the test suite locally. Remember that the subdirectories under test/ create fine-grained testing targets, so you can e.g. use make check-clang-ast to only run Clang's AST tests.
5. Create a Git commit.
6. Run git clang-format HEAD~1 to format your changes.
7. Open a pull request to the upstream repository on GitHub. Detailed instructions can be found in GitHub's documentation. Mention this issue in the description of the pull request.

If you have any further questions about this issue, don't hesitate to ask via a comment in the thread below.

[llvmbot]

@llvm/issue-subscribers-good-first-issue

Author: None (Engininja2)

Here is my test kernel: ```cpp 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.
```asm
	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]

[doraeneko]

I'd like to look into that issue, if I may (would be my first).

[doraeneko]

A first observation (I'm trying to grasp the transformation steps from the AMDGPU target on the way):
the special transformation of select that occurs comparing for 0 and powers of 2 seems not to be there in the "DAG-Combine1" phase, but appears in the "DAG-Combine2" phase (checked with llc ... -view-dag-combine1-dags and llc ... -view-dag-combine2-dags).

[doraeneko]

Variant for 32:

Variant for 33:

[doraeneko]

Ok, in SIFixSGPRCopies there seems to be a code transformation changing the "S_LSHL_B32" for the 32-case into a
V_LSHLREV_B32_e64, and propagates the target to the inline assembler part:

  %17:vgpr_32 = V_CNDMASK_B32_e64 0, 0, 0, 1, killed %16:sreg_64_xexec, implicit $exec
  %18:sreg_32 = S_MOV_B32 5
  %19:sreg_32 = COPY %17:vgpr_32
  %13:sreg_32 = S_LSHL_B32 killed %19:sreg_32, killed %18:sreg_32, implicit-def dead $scc
  INLINEASM &"s_mov_b32 $0, $1" [sideeffect] [isconvergent] [attdialect], $0:[regdef:SReg_32], def %12:sreg_32, $1:[reguse:SReg_32], %13:sreg_32, !8

==>

  %17:vgpr_32 = V_CNDMASK_B32_e64 0, 0, 0, 1, killed %16:sreg_64_xexec, implicit $exec
  %18:sreg_32 = S_MOV_B32 5
  %19:sreg_32 = IMPLICIT_DEF
  %20:vgpr_32 = V_LSHLREV_B32_e64 killed %18:sreg_32, %17:vgpr_32, implicit $exec
  INLINEASM &"s_mov_b32 $0, $1" [sideeffect] [isconvergent] [attdialect], $0:[regdef:SReg_32], def %12:sreg_32, $1:[reguse:SReg_32], %20:vgpr_32, !8

Maybe inline assembler is not correctly considered in there?

[arsenm]

SIFixSGPRCopies is a hack, and it does mishandle inline assembly (and arbitrary calls). This case is quite similar to the one I fixed in c08d7b3. In principle we can just insert readfirstlane calls on known uniform SGPR asm arguments.

However in this case, we do not need a readfirstlane. The issue is we seem to assume i1 -> i32 zero_extend is VALU:

llvm-project/llvm/lib/Target/AMDGPU/SIInstructions.td

Line 2339 in 9e831d5

def : Ext32Pat <zext>;

We need a uniform variant of the pattern, similar to https://github.com/llvm/llvm-project/blob/9e831d50a0d67c6779048742657edb437818b681/llvm/lib/Target/AMDGPU/SIInstructions.td#L15C2-L15C3 for sext_inreg.