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

Comment typo and not needed

400

Ditto

lib/Target/AMDGPU/SIInstructions.td

2364–2365

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

2371–2373

Ditto

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

7–11

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

Attributes not needed on call site

15

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

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

1598

Same here.

2361

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]

include/llvm/IR/IntrinsicsAMDGPU.td

 // llvm.amdgcn.ds.swizzle src offset
 def int_amdgcn_ds_swizzle :
   GCCBuiltin<"__builtin_amdgcn_ds_swizzle">,
   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
 
 // llvm.amdgcn.lerp
 def int_amdgcn_lerp :
   GCCBuiltin<"__builtin_amdgcn_lerp">,
   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
 
+// llvm.amdgcn.readfirlane src
+def int_amdgcn_readfirstlane :
+  GCCBuiltin<"__builtin_amdgcn_readfirstlane">,
+  Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
+
+// llvm.amdgcn.readlane
+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<[], [], []>;
 
 def int_amdgcn_buffer_wbinvl1_vol :
   GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_vol">,

lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.h

 
     MCOperand errOperand(unsigned V, const llvm::Twine& ErrMsg) const;
 
     DecodeStatus tryDecodeInst(const uint8_t* Table,
                                MCInst &MI,
                                uint64_t Inst,
                                uint64_t Address) const;
 
     MCOperand decodeOperand_VGPR_32(unsigned Val) const;
     MCOperand decodeOperand_VS_32(unsigned Val) const;
+    MCOperand decodeOperand_VM0_32(unsigned Val) const;
     MCOperand decodeOperand_VS_64(unsigned Val) const;
 
     MCOperand decodeOperand_VReg_64(unsigned Val) const;
     MCOperand decodeOperand_VReg_96(unsigned Val) const;
     MCOperand decodeOperand_VReg_128(unsigned Val) const;
 
     MCOperand decodeOperand_SReg_32(unsigned Val) const;
     MCOperand decodeOperand_SReg_32_XM0(unsigned Val) const;
     MCOperand decodeOperand_SReg_64(unsigned Val) const;
     MCOperand decodeOperand_SReg_128(unsigned Val) const;

lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp

                                                     uint64_t /*Addr*/, \
                                                     const void *Decoder) { \
   auto DAsm = static_cast<const AMDGPUDisassembler*>(Decoder); \
   return addOperand(Inst, DAsm->decodeOperand_##DecName(Imm)); \
 }
 
 #define DECODE_OPERAND(RegClass) DECODE_OPERAND2(RegClass, RegClass)
 
 DECODE_OPERAND(VGPR_32)
 DECODE_OPERAND(VS_32)
+DECODE_OPERAND(VM0_32)
 DECODE_OPERAND(VS_64)
 
 DECODE_OPERAND(VReg_64)
 DECODE_OPERAND(VReg_96)
 DECODE_OPERAND(VReg_128)
 
 DECODE_OPERAND(SReg_32)
 DECODE_OPERAND(SReg_32_XM0)
 DECODE_OPERAND(SReg_64)
 DECODE_OPERAND(SReg_128)

   if (Val % (1 << shift))
     *CommentStream << "Warning: " << getRegClassName(SRegClassID)
                    << ": scalar reg isn't aligned " << Val;
   return createRegOperand(SRegClassID, Val >> shift);
 }
 
 MCOperand AMDGPUDisassembler::decodeOperand_VS_32(unsigned Val) const {
   return decodeSrcOp(OPW32, Val);
 }
 
+MCOperand AMDGPUDisassembler::decodeOperand_VM0_32(unsigned Val) const {
+  return decodeSrcOp(OPW32, Val);
+}
+
 MCOperand AMDGPUDisassembler::decodeOperand_VS_64(unsigned Val) const {
   return decodeSrcOp(OPW64, Val);
 }
 
 MCOperand AMDGPUDisassembler::decodeOperand_VGPR_32(unsigned Val) const {
   return createRegOperand(AMDGPU::VGPR_32RegClassID, Val);
 }
 
 MCOperand AMDGPUDisassembler::decodeOperand_VReg_64(unsigned Val) const {
   return createRegOperand(AMDGPU::VReg_64RegClassID, Val);

lib/Target/AMDGPU/SIInstructions.td

 defm V_MOV_B32 : VOP1Inst <vop1<0x1>, "v_mov_b32", VOP_I32_I32>;
 } // End isMoveImm = 1
 
 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 VM0_32:$src0),
   "v_readfirstlane_b32 $vdst, $src0",
   []
 > {
   let isConvergent = 1;
 }
 
 }
 
 let SchedRW = [WriteQuarterRate32] in {
 

 
 } // End isCommutable = 1
 
 // 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),
+  (ins VM0_32:$src0, SCSrc_32:$src1),
   "v_readlane_b32 $vdst, $src0, $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"
 >;

 
 //===----------------------------------------------------------------------===//
 // DS_SWIZZLE Intrinsic Pattern.
 //===----------------------------------------------------------------------===//
 def : Pat <
   (int_amdgcn_ds_swizzle i32:$src, imm:$offset16),
   (DS_SWIZZLE_B32 $src, (as_i16imm $offset16), (i1 0))
 >;
 
 //===----------------------------------------------------------------------===//
+// READFIRSTLANE Intrinsic Pattern.
+//===----------------------------------------------------------------------===//
+def : Pat <
+  (int_amdgcn_readfirstlane i32:$src),
+  (V_READFIRSTLANE_B32 $src)
+>;
+
+//===----------------------------------------------------------------------===//
+// READLANE Intrinsic Pattern.
+//===----------------------------------------------------------------------===//
+def : Pat <
+  (int_amdgcn_readlane i32:$src0, i32:$src1),
+  (V_READLANE_B32 $src0, $src1)
+>;
+
+//===----------------------------------------------------------------------===//
 // SMRD Patterns
 //===----------------------------------------------------------------------===//
 
 multiclass SMRD_Pattern <string Instr, ValueType vt> {
 
   // 1. IMM offset
   def : Pat <
     (smrd_load (SMRDImm i64:$sbase, i32:$offset)),
     (vt (!cast<SMRD>(Instr#"_IMM") $sbase, $offset))
   >;

lib/Target/AMDGPU/SIRegisterInfo.td

 def SCSrc_32 : RegInlineOperand<SReg_32> {
   let ParserMatchClass = RegImmMatcher<"SCSrc32">;
 }
 
 //===----------------------------------------------------------------------===//
 //  VSrc_* Operands with an SGPR, VGPR or a 32-bit immediate
 //===----------------------------------------------------------------------===//
 
 def VS_32 : RegisterClass<"AMDGPU", [i32, f32], 32, (add VGPR_32, SReg_32)>;
 
+def VM0_32 : RegisterClass<"AMDGPU", [i32, f32], 32, (add VGPR_32, M0)> {
+  let isAllocatable = 0;
+}
+
 def VS_64 : RegisterClass<"AMDGPU", [i64, f64], 32, (add VReg_64, SReg_64)> {
   let CopyCost = 2;
 }
 
 def VSrc_32 : RegisterOperand<VS_32> {
   let OperandNamespace = "AMDGPU";
   let OperandType = "OPERAND_REG_IMM32";
   let ParserMatchClass = RegImmMatcher<"VSrc32">;
 }
 

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) nounwind {
+  %readfirstlane = call i32 @llvm.amdgcn.readfirstlane(i32 %src) #0
+  store i32 %readfirstlane, i32 addrspace(1)* %out, align 4
+  ret void
+}
+
+attributes #0 = { nounwind readnone convergent }

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: {{^}}readlane_sreg:
+; CHECK: v_readlane_b32 s{{[0-9]+}}, v{{[0-9]+}}, s{{[0-9]+}}
+define void @readlane_sreg(i32 addrspace(1)* %out, i32 %src0, i32 %src1) nounwind {
+  %readlane = call i32 @llvm.amdgcn.readlane(i32 %src0, i32 %src1) #0
+  store i32 %readlane, i32 addrspace(1)* %out, align 4
+  ret void
+}
+
+; CHECK-LABEL: {{^}}readlane_imm:
+; CHECK: v_readlane_b32 s{{[0-9]+}}, v{{[0-9]+}}, 32
+define void @readlane_imm(i32 addrspace(1)* %out, i32 %src0) nounwind {
+  %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 }