AMDGPU/SI: Add llvm.amdgcn.mbcnt.* intrinsics

[• tstellarAMD]

Summary

These are meant to be used instead of the llvm.SI.tid intrinsic which will
be deprecated at some point.

[arsenm]

include/llvm/IR/IntrinsicsAMDGPU.td

150–151

I don't think the type should be included in the name. For instructions that have multiple types, it is redundant with the llvm name mangling. For ones that don't work on multiple types, it isn't necessary information. I would have this be int_amdgcn_mbcnt_lo/int_amdgcn_mbcnt_hi

test/CodeGen/AMDGPU/llvm.SI.tid.ll

16–33

The test should be in its own mbcnt file

[• tstellarAMD]

Keep old intrinsics.

[arsenm]

LGTM

include/llvm/IR/IntrinsicsAMDGPU.td

 
 // __builtin_amdgcn_v_interp_p2 <p1>, <j>, <attr_chan>, <attr>, <m0>
 def int_amdgcn_v_interp_p2 :
   GCCBuiltin<"__builtin_amdgcn_v_interp_p2">,
   Intrinsic<[llvm_float_ty],
             [llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
             [IntrNoMem]>;  // See int_amdgcn_v_interp_p1 for why this is
                            // IntrNoMem.
+
+def int_amdgcn_v_mbcnt_lo_u32_b32 :
+  GCCBuiltin<"__builtin_amdgcn_v_mbcnt_lo_u32_b32">,
+  Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
+
+def int_amdgcn_v_mbcnt_hi_u32_b32 :
+  GCCBuiltin<"__builtin_amdgcn_v_mbcnt_hi_u32_b32">,
+  Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
 }

lib/Target/AMDGPU/AMDGPUIntrinsicExpander.cpp

 namespace {
 
 class AMDGPUIntrinsicExpander : public FunctionPass,
                                 public InstVisitor<AMDGPUIntrinsicExpander> {
 private:
   Module *Mod;
 
   void expandSIFsInterp(IntrinsicInst &I);
+  void expandSITid(IntrinsicInst &I);
 
 public:
   static char ID;
 
   AMDGPUIntrinsicExpander() : FunctionPass(ID) { }
 
   bool doInitialization(Module &M) override;
   bool runOnFunction(Function &F) override;
void AMDGPUIntrinsicExpander::expandSIFsInterp(IntrinsicInst &I) {
 
   Value *InterpP2Args[] = {P1, JCoord, AttrChan, Attr, M0};
   Value *P2 = Builder.CreateCall(InterpP2, InterpP2Args);
 
   I.replaceAllUsesWith(P2);
   I.eraseFromParent();
 }
 
+void AMDGPUIntrinsicExpander::expandSITid(IntrinsicInst &I) {
+  IRBuilder<> Builder(&I);
+
+  Function *MbcntLo =
+    Intrinsic::getDeclaration(Mod, Intrinsic::amdgcn_v_mbcnt_lo_u32_b32);
+  Function *MbcntHi =
+    Intrinsic::getDeclaration(Mod, Intrinsic::amdgcn_v_mbcnt_hi_u32_b32);
+
+  IntegerType *I32 = IntegerType::get(I.getContext(), 32);
+  Constant *Mask = ConstantInt::get(I32, APInt::getAllOnesValue(32));
+  Constant *Zero = ConstantInt::get(I32, 0);
+
+  Value *LoArgs[] = {Mask, Zero};
+  Value *Lo = Builder.CreateCall(MbcntLo, LoArgs);
+
+  Value *HiArgs[] = {Mask, Lo};
+  Value *Hi = Builder.CreateCall(MbcntHi, HiArgs);
+
+  I.replaceAllUsesWith(Hi);
+  I.eraseFromParent();
+}
+
 void AMDGPUIntrinsicExpander::visitCallInst(CallInst &I) {
   IntrinsicInst *Intr = dyn_cast<IntrinsicInst>(&I);
   if (!Intr)
     return;
 
   if (Intr->getCalledFunction()->getName() == "llvm.SI.fs.interp")
     expandSIFsInterp(*Intr);
+  else if (Intr->getCalledFunction()->getName() == "llvm.SI.tid")
+    expandSITid(*Intr);
+
 }
 
 bool AMDGPUIntrinsicExpander::runOnFunction(Function &F) {
   bool Changed = false;
   visit(F);
 
   return Changed;
 }
 
 FunctionPass *llvm::createAMDGPUIntrinsicExpanderPass() {
   return new AMDGPUIntrinsicExpander();
 }

lib/Target/AMDGPU/SIInstructions.td

 
 defm V_BFM_B32 : VOP2_VI3_Inst <vop23<0x1e, 0x293>, "v_bfm_b32",
   VOP_I32_I32_I32
 >;
 defm V_BCNT_U32_B32 : VOP2_VI3_Inst <vop23<0x22, 0x28b>, "v_bcnt_u32_b32",
   VOP_I32_I32_I32
 >;
 defm V_MBCNT_LO_U32_B32 : VOP2_VI3_Inst <vop23<0x23, 0x28c>, "v_mbcnt_lo_u32_b32",
-  VOP_I32_I32_I32
+  VOP_I32_I32_I32, int_amdgcn_v_mbcnt_lo_u32_b32
 >;
 defm V_MBCNT_HI_U32_B32 : VOP2_VI3_Inst <vop23<0x24, 0x28d>, "v_mbcnt_hi_u32_b32",
-  VOP_I32_I32_I32
+  VOP_I32_I32_I32, int_amdgcn_v_mbcnt_hi_u32_b32
 >;
 defm V_LDEXP_F32 : VOP2_VI3_Inst <vop23<0x2b, 0x288>, "v_ldexp_f32",
   VOP_F32_F32_I32, AMDGPUldexp
 >;
 
 defm V_CVT_PKACCUM_U8_F32 : VOP2_VI3_Inst <vop23<0x2c, 0x1f0>, "v_cvt_pkaccum_u8_f32",
   VOP_I32_F32_I32>; // TODO: set "Uses = dst"
 

 // The multiplication scales from [0,1] to the unsigned integer range
 def : Pat <
   (AMDGPUurecip i32:$src0),
   (V_CVT_U32_F32_e32
     (V_MUL_F32_e32 CONST.FP_UINT_MAX_PLUS_1,
                    (V_RCP_IFLAG_F32_e32 (V_CVT_F32_U32_e32 $src0))))
 >;
 
-def : Pat <
-  (int_SI_tid),
-  (V_MBCNT_HI_U32_B32_e64 0xffffffff,
-                          (V_MBCNT_LO_U32_B32_e64 0xffffffff, 0))
->;
-
 //===----------------------------------------------------------------------===//
 // VOP3 Patterns
 //===----------------------------------------------------------------------===//
 
 def : IMad24Pat<V_MAD_I32_I24>;
 def : UMad24Pat<V_MAD_U32_U24>;
 
 def : Pat <

lib/Target/AMDGPU/SIIntrinsics.td

 //===-- SIIntrinsics.td - SI Intrinsic defs ----------------*- tablegen -*-===//
 //
 //                     The LLVM Compiler Infrastructure
 //
 // This file is distributed under the University of Illinois Open Source
 // License. See LICENSE.TXT for details.
 //
 //===----------------------------------------------------------------------===//
 //
 // SI Intrinsic Definitions
 //
 //===----------------------------------------------------------------------===//
 
 
 let TargetPrefix = "SI", isTarget = 1 in {
 
-  def int_SI_tid : Intrinsic <[llvm_i32_ty], [], [IntrNoMem]>;
   def int_SI_packf16 : Intrinsic <[llvm_i32_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem]>;
   def int_SI_export : Intrinsic <[], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], []>;
   def int_SI_load_const : Intrinsic <[llvm_float_ty], [llvm_anyint_ty, llvm_i32_ty], [IntrNoMem]>;
   def int_SI_vs_load_input : Intrinsic <[llvm_v4f32_ty], [llvm_anyint_ty, llvm_i16_ty, llvm_i32_ty], [IntrNoMem]> ;
 
   // Fully-flexible TBUFFER_STORE_FORMAT_* except for the ADDR64 bit, which is not exposed
   def int_SI_tbuffer_store : Intrinsic <
     [],

test/CodeGen/AMDGPU/llvm.SI.tid.ll

 ;RUN: llc < %s -march=amdgcn -mcpu=verde -verify-machineinstrs | FileCheck --check-prefix=SI --check-prefix=GCN %s
 ;RUN: llc < %s -march=amdgcn -mcpu=tonga -verify-machineinstrs | FileCheck --check-prefix=VI --check-prefix=GCN %s
 
 ;GCN: v_mbcnt_lo_u32_b32_e64
 ;SI: v_mbcnt_hi_u32_b32_e32
 ;VI: v_mbcnt_hi_u32_b32_e64
 
 define void @main(<16 x i8> addrspace(2)* inreg, <16 x i8> addrspace(2)* inreg, <32 x i8> addrspace(2)* inreg, i32 inreg) "ShaderType"="0" {
 main_body:
   %4 = call i32 @llvm.SI.tid()
   %5 = bitcast i32 %4 to float
   call void @llvm.SI.export(i32 15, i32 1, i32 1, i32 0, i32 1, float %5, float %5, float %5, float %5)
   ret void
 }
 
+;GCN-LABEL: {{^}}mbcnt_intrinsics:
+;GCN: v_mbcnt_lo_u32_b32_e64 [[LO:v[0-9]+]], -1, 0
+;SI: v_mbcnt_hi_u32_b32_e32 {{v[0-9]+}}, -1, [[LO]]
+;VI: v_mbcnt_hi_u32_b32_e64 {{v[0-9]+}}, -1, [[LO]]
+
+define void @mbcnt_intrinsics(<16 x i8> addrspace(2)* inreg, <16 x i8> addrspace(2)* inreg, <32 x i8> addrspace(2)* inreg, i32 inreg) "ShaderType"="0" {
+main_body:
+  %lo = call i32 @llvm.amdgcn.v.mbcnt.lo.u32.b32(i32 -1, i32 0)
+  %hi = call i32 @llvm.amdgcn.v.mbcnt.hi.u32.b32(i32 -1, i32 %lo)
+  %4 = bitcast i32 %hi to float
+  call void @llvm.SI.export(i32 15, i32 1, i32 1, i32 0, i32 1, float %4, float %4, float %4, float %4)
+  ret void
+}
+
+declare i32 @llvm.amdgcn.v.mbcnt.lo.u32.b32(i32, i32) readnone
+
+declare i32 @llvm.amdgcn.v.mbcnt.hi.u32.b32(i32, i32) readnone
+
 declare i32 @llvm.SI.tid() readnone
 
 declare void @llvm.SI.export(i32, i32, i32, i32, i32, float, float, float, float)