AMDGPU/SI: Implement readlane/readfirstlane intrinsics to expose the instructions.

[cfang]

Summary

This patch defines a new register class (VM0_32 = VReg_32 + M0) and implements readlane and readfirstlane intrinsics.

[arsenm]

include/llvm/IR/IntrinsicsAMDGPU.td

395 ↗

(On Diff #64421)

Comment typo and not needed

400 ↗

(On Diff #64421)

Ditto

lib/Target/AMDGPU/SIInstructions.td

2364–2365 ↗

(On Diff #64421)

This is a very simple pattern so should go with the instruction definition's pattern

2371–2373 ↗

(On Diff #64421)

Ditto

test/CodeGen/AMDGPU/llvm.amdgcn.readfirstlane.ll

7–11 ↗

(On Diff #64421)

Should also include a test which has an immediate source to make sure that it is moved into a register. Another that uses inline asm to put a value in m0 would also be useful (same for the other intrinsic too)

test/CodeGen/AMDGPU/llvm.amdgcn.readlane.ll

8 ↗

(On Diff #64421)

Attributes not needed on call site

15 ↗

(On Diff #64421)

Use attribute group for the nounwind also

[cfang]

update based on Matt's comments:

1. Remove unnecessary comments in intrinsic definitions;
2. move the pattern defs into instructions;
3. add imm src test cases;
4. remove unnecessary attributes at call sites;
5. put "nounwind" into a new attribute group.

NOTE: M0 test case is not added in this update.

[arsenm]

I think you should still add the m0 testcase even though it isn't folded

[cfang]

Add "m0" testcase based on Matt's request.

TODO: shoudl fold m0 into the src operand of readlane/readfirstlane. This is at a lower priority and can be done in a separate patch.

[• tstellarAMD]

lib/Target/AMDGPU/SIInstructions.td

1225 ↗

(On Diff #64711)

You should use the type in the pattern rather than the register class.

1598 ↗

(On Diff #64711)

Same here.

2361 ↗

(On Diff #64711)

Extra whitespace.

[nhaehnle]

Do you plan to add code to DivergenceAnalysis to recognize these intrinsics? We might eventually want to use a readfirstlane intrinsic from Mesa to hint when certain indices can be assumed to be uniform.

(Right now, LLVM ends up deciding where to insert the readfirstlane in those cases, since they are situations where buffer resource descriptors are loaded which must end up in SGPRs anyway. By giving a readfirstlane hint, we could perhaps generate more efficient code.)

[cfang]

Update based on Tom's Comments:

1. Use the type in the pattern instead of register class.
2. remove extra white space.

[cfang]

In D22489#497604, @nhaehnle wrote:

Do you plan to add code to DivergenceAnalysis to recognize these intrinsics? We might eventually want to use a readfirstlane intrinsic from Mesa to hint when certain indices can be assumed to be uniform.

(Right now, LLVM ends up deciding where to insert the readfirstlane in those cases, since they are situations where buffer resource descriptors are loaded which must end up in SGPRs anyway. By giving a readfirstlane hint, we could perhaps generate more efficient code.)

I do not have this plan right now. But we will definitely consider your suggestion.

[cfang]

Hi, Matt and Tom:

Is there any further suggestions on this patch? Thanks.

[cfang]

In D22489#502282, @cfang wrote:

Hi, Matt and Tom:

Is there any further suggestions on this patch? Thanks.

Need further suggestion what to do on this. I think readlane and readfirstlane are very important.

[• tstellarAMD]

llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td

 def int_amdgcn_icmp :
   Intrinsic<[llvm_i64_ty], [llvm_anyint_ty, LLVMMatchType<0>, llvm_i32_ty],
             [IntrNoMem, IntrConvergent]>;
 
 def int_amdgcn_fcmp :
   Intrinsic<[llvm_i64_ty], [llvm_anyfloat_ty, LLVMMatchType<0>, llvm_i32_ty],
             [IntrNoMem, IntrConvergent]>;
 
+def int_amdgcn_readfirstlane :
+  GCCBuiltin<"__builtin_amdgcn_readfirstlane">,
+  Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
+
+def int_amdgcn_readlane :
+  GCCBuiltin<"__builtin_amdgcn_readlane">,
+  Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
+
 //===----------------------------------------------------------------------===//
 // CI+ Intrinsics
 //===----------------------------------------------------------------------===//
 
 def int_amdgcn_s_dcache_inv_vol :
   GCCBuiltin<"__builtin_amdgcn_s_dcache_inv_vol">,
   Intrinsic<[], [], []>;
 

llvm/trunk/lib/Target/AMDGPU/SIInstructions.td

 
 let Uses = [EXEC] in {
 
 // FIXME: Specify SchedRW for READFIRSTLANE_B32
 
 def V_READFIRSTLANE_B32 : VOP1 <
   0x00000002,
   (outs SReg_32:$vdst),
-  (ins VS_32:$src0),
+  (ins VGPR_32:$src0),
   "v_readfirstlane_b32 $vdst, $src0",
-  []
+  [(set i32:$vdst, (int_amdgcn_readfirstlane i32:$src0))]
 > {
   let isConvergent = 1;
 }
 
 }
 
 let SchedRW = [WriteQuarterRate32] in {
 

 
 // These are special and do not read the exec mask.
 let isConvergent = 1, Uses = []<Register> in {
 
 defm V_READLANE_B32 : VOP2SI_3VI_m <
   vop3 <0x001, 0x289>,
   "v_readlane_b32",
   (outs SReg_32:$vdst),
-  (ins VS_32:$src0, SCSrc_32:$src1),
-  "v_readlane_b32 $vdst, $src0, $src1"
+  (ins VGPR_32:$src0, SCSrc_32:$src1),
+  "v_readlane_b32 $vdst, $src0, $src1",
+  [(set i32:$vdst, (int_amdgcn_readlane i32:$src0, i32:$src1))]
 >;
 
 defm V_WRITELANE_B32 : VOP2SI_3VI_m <
   vop3 <0x002, 0x28a>,
   "v_writelane_b32",
   (outs VGPR_32:$vdst),
   (ins SReg_32:$src0, SCSrc_32:$src1),
   "v_writelane_b32 $vdst, $src0, $src1"

llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.readfirstlane.ll

+; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=fiji -verify-machineinstrs < %s | FileCheck %s
+
+declare i32 @llvm.amdgcn.readfirstlane(i32) #0
+
+; CHECK-LABEL: {{^}}test_readfirstlane:
+; CHECK: v_readfirstlane_b32 s{{[0-9]+}}, v{{[0-9]+}}
+define void @test_readfirstlane(i32 addrspace(1)* %out, i32 %src) #1 {
+  %readfirstlane = call i32 @llvm.amdgcn.readfirstlane(i32 %src)
+  store i32 %readfirstlane, i32 addrspace(1)* %out, align 4
+  ret void
+}
+
+; CHECK-LABEL: {{^}}test_readfirstlane_imm:
+; CHECK: v_mov_b32_e32 [[VVAL:v[0-9]]], 32
+; CHECK: v_readfirstlane_b32 s{{[0-9]+}}, [[VVAL]]
+define void @test_readfirstlane_imm(i32 addrspace(1)* %out) #1 {
+  %readfirstlane = call i32 @llvm.amdgcn.readfirstlane(i32 32)
+  store i32 %readfirstlane, i32 addrspace(1)* %out, align 4
+  ret void
+}
+
+; TODO: m0 should be folded.
+; CHECK-LABEL: {{^}}test_readfirstlane_m0:
+; CHECK: s_mov_b32 m0, -1
+; CHECK: v_mov_b32_e32 [[VVAL:v[0-9]]], m0
+; CHECK: v_readfirstlane_b32 s{{[0-9]+}}, [[VVAL]]
+define void @test_readfirstlane_m0(i32 addrspace(1)* %out) #1 {
+  %m0 = call i32 asm "s_mov_b32 m0, -1", "={M0}"()
+  %readfirstlane = call i32 @llvm.amdgcn.readfirstlane(i32 %m0)
+  store i32 %readfirstlane, i32 addrspace(1)* %out, align 4
+  ret void
+}
+
+attributes #0 = { nounwind readnone convergent }
+attributes #1 = { nounwind }

llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.readlane.ll

+; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=fiji -verify-machineinstrs < %s | FileCheck %s
+
+declare i32 @llvm.amdgcn.readlane(i32, i32) #0
+
+; CHECK-LABEL: {{^}}test_readlane_sreg:
+; CHECK: v_readlane_b32 s{{[0-9]+}}, v{{[0-9]+}}, s{{[0-9]+}}
+define void @test_readlane_sreg(i32 addrspace(1)* %out, i32 %src0, i32 %src1) #1 {
+  %readlane = call i32 @llvm.amdgcn.readlane(i32 %src0, i32 %src1)
+  store i32 %readlane, i32 addrspace(1)* %out, align 4
+  ret void
+}
+
+; CHECK-LABEL: {{^}}test_readlane_imm_sreg:
+; CHECK: v_mov_b32_e32 [[VVAL:v[0-9]]], 32
+; CHECK: v_readlane_b32 s{{[0-9]+}}, [[VVAL]], s{{[0-9]+}}
+define void @test_readlane_imm_sreg(i32 addrspace(1)* %out, i32 %src1) #1 {
+  %readlane = call i32 @llvm.amdgcn.readlane(i32 32, i32 %src1)
+  store i32 %readlane, i32 addrspace(1)* %out, align 4
+  ret void
+}
+
+; TODO: m0 should be folded.
+; CHECK-LABEL: {{^}}test_readlane_m0_sreg:
+; CHECK: s_mov_b32 m0, -1
+; CHECK: v_mov_b32_e32 [[VVAL:v[0-9]]], m0
+; CHECK: v_readlane_b32 s{{[0-9]+}}, [[VVAL]], s{{[0-9]+}}
+define void @test_readlane_m0_sreg(i32 addrspace(1)* %out, i32 %src1) #1 {
+  %m0 = call i32 asm "s_mov_b32 m0, -1", "={M0}"()
+  %readlane = call i32 @llvm.amdgcn.readlane(i32 %m0, i32 %src1)
+  store i32 %readlane, i32 addrspace(1)* %out, align 4
+  ret void
+}
+
+; CHECK-LABEL: {{^}}test_readlane_imm:
+; CHECK: v_readlane_b32 s{{[0-9]+}}, v{{[0-9]+}}, 32
+define void @test_readlane_imm(i32 addrspace(1)* %out, i32 %src0) #1 {
+  %readlane = call i32 @llvm.amdgcn.readlane(i32 %src0, i32 32) #0
+  store i32 %readlane, i32 addrspace(1)* %out, align 4
+  ret void
+}
+
+attributes #0 = { nounwind readnone convergent }
+attributes #1 = { nounwind }