AMDGPU: Add type mangling to llvm.amdgcn.readfirstlane

[arsenm]

Summary

I've thought this should at minimum work for all the legal 32-bit
types for a long time. Arguably we should also legalize other types,
or at least multiples of 32.

This fixes selection for non-s32 types when readfirstlane is inserted
for SGPR return values.

This will require API users to add the type mangling when getting the
declaration. I think the way mesa constructs intrinsics will happen to
get away without needing to change anything.

[arsenm]

Also not sure if I should change readlane/writelane at the same time

[nhaehnle]

I think this makes sense, but it is a flag-day change for both Mesa and LLPC. Intrinsics created with incorrect mangling don't fall over completely (as evidenced by the fact that your change *doesn't* change the test files), but certain things like attributes tend to not work quite correctly anymore.

So I think we can still do this change, but please do all of readlane, readfirstlane, and writelane in a single change so that the pain is minimized.

You do have have to fix all test cases to use the correct mangling.

[nhaehnle]

On second thought, can we make this not be a flag-day change?

Also, are you currently working on this? If not, I may be interested in picking it up.

[arsenm]

In D84639#2218645, @nhaehnle wrote:

On second thought, can we make this not be a flag-day change?

Also, are you currently working on this? If not, I may be interested in picking it up.

I'm not actively working on this. This only fixes 1 lit test which I switch the default to GlobalISel, so it's not particularly high on my priority list

[foad]

llvm/lib/Target/AMDGPU/AMDGPUAtomicOptimizer.cpp

609

Hoist this above the "if" on line 605?

[nhaehnle]

D86154 is an attempt to do essentially this, on all lane intrinsics simultaneously, giving them a new name to avoid a flag day change.

[arsenm]

Obsoleted by D147732

clang/lib/CodeGen/CGBuiltin.cpp

if (ProcessOrderScopeAMDGCN(EmitScalarExpr(E->getArg(2)),
       bool Volatile =
           PtrTy->castAs<PointerType>()->getPointeeType().isVolatileQualified();
       Value *IsVolatile = Builder.getInt1(static_cast<bool>(Volatile));
 
       return Builder.CreateCall(F, {Ptr, Val, MemOrder, MemScope, IsVolatile});
     }
     LLVM_FALLTHROUGH;
   }
+  case AMDGPU::BI__builtin_amdgcn_readfirstlane:
+    return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_readfirstlane);
   default:
     return nullptr;
   }
 }
 
 /// Handle a SystemZ function in which the final argument is a pointer
 /// to an int that receives the post-instruction CC value.  At the LLVM level
 /// this is represented as a function that returns a {result, cc} pair.

clang/test/CodeGenOpenCL/builtins-amdgcn.cl

 
 // CHECK-LABEL: @test_ds_bpermute
 // CHECK: call i32 @llvm.amdgcn.ds.bpermute(i32 %a, i32 %b)
 void test_ds_bpermute(global int* out, int a, int b)
 {
   *out = __builtin_amdgcn_ds_bpermute(a, b);
 }
 
-// CHECK-LABEL: @test_readfirstlane
-// CHECK: call i32 @llvm.amdgcn.readfirstlane(i32 %a)
+// CHECK-LABEL: @test_readfirstlane(
+// CHECK: call i32 @llvm.amdgcn.readfirstlane.i32(i32 %a)
 void test_readfirstlane(global int* out, int a)
 {
   *out = __builtin_amdgcn_readfirstlane(a);
 }
 
 // CHECK-LABEL: @test_readlane
 // CHECK: call i32 @llvm.amdgcn.readlane(i32 %a, i32 %b)
 void test_readlane(global int* out, int a, int b)

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

Intrinsic<[llvm_anyint_ty], [llvm_anyfloat_ty, LLVMMatchType<1>, llvm_i32_ty],
             [IntrNoMem, IntrConvergent, IntrWillReturn,
              ImmArg<ArgIndex<2>>]>;
 
 def int_amdgcn_ballot :
   Intrinsic<[llvm_anyint_ty], [llvm_i1_ty],
             [IntrNoMem, IntrConvergent, IntrWillReturn]>;
 
 def int_amdgcn_readfirstlane :
-  GCCBuiltin<"__builtin_amdgcn_readfirstlane">,
-  Intrinsic<[llvm_i32_ty], [llvm_i32_ty],
+  Intrinsic<[llvm_any_ty], [LLVMMatchType<0>],
             [IntrNoMem, IntrConvergent, IntrWillReturn]>;
 
 // The lane argument must be uniform across the currently active threads of the
 // current wave. Otherwise, the result is undefined.
 def int_amdgcn_readlane :
   GCCBuiltin<"__builtin_amdgcn_readlane">,
   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
             [IntrNoMem, IntrConvergent, IntrWillReturn]>;

llvm/lib/Target/AMDGPU/AMDGPUAtomicOptimizer.cpp

 private:
   SmallVector<ReplacementInfo, 8> ToReplace;
   const LegacyDivergenceAnalysis *DA;
   const DataLayout *DL;
   DominatorTree *DT;
   const GCNSubtarget *ST;
   bool IsPixelShader;
 
+  Function *Readfirstlane = nullptr;
+
   Value *buildScan(IRBuilder<> &B, AtomicRMWInst::BinOp Op, Value *V,
                    Value *const Identity) const;
   Value *buildShiftRight(IRBuilder<> &B, Value *V, Value *const Identity) const;
   void optimizeAtomic(Instruction &I, AtomicRMWInst::BinOp Op, unsigned ValIdx,
-                      bool ValDivergent) const;
+                      bool ValDivergent);
 
 public:
   static char ID;
 
   AMDGPUAtomicOptimizer() : FunctionPass(ID) {}
 
   bool runOnFunction(Function &F) override;
 
   void getAnalysisUsage(AnalysisUsage &AU) const override {
     AU.addPreserved<DominatorTreeWrapperPass>();
     AU.addRequired<LegacyDivergenceAnalysis>();
     AU.addRequired<TargetPassConfig>();
   }
 
+  Function *getReadfirstlane(Module *M) {
+    if (!Readfirstlane) {
+      Readfirstlane =
+          Intrinsic::getDeclaration(M, Intrinsic::amdgcn_readfirstlane,
+                                    {Type::getInt32Ty(M->getContext())});
+    }
+
+    return Readfirstlane;
+  }
+
   void visitAtomicRMWInst(AtomicRMWInst &I);
   void visitIntrinsicInst(IntrinsicInst &I);
 };
 
 } // namespace
 
 char AMDGPUAtomicOptimizer::ID = 0;
 
bool AMDGPUAtomicOptimizer::runOnFunction(Function &F) {
   visit(F);
 
   const bool Changed = !ToReplace.empty();
 
   for (ReplacementInfo &Info : ToReplace) {
     optimizeAtomic(*Info.I, Info.Op, Info.ValIdx, Info.ValDivergent);
   }
 
+  Readfirstlane = nullptr;
   ToReplace.clear();
 
   return Changed;
 }
 
 void AMDGPUAtomicOptimizer::visitAtomicRMWInst(AtomicRMWInst &I) {
   // Early exit for unhandled address space atomic instructions.
   switch (I.getPointerAddressSpace()) {
case AtomicRMWInst::Max:
     return APInt::getSignedMinValue(BitWidth);
   case AtomicRMWInst::Min:
     return APInt::getSignedMaxValue(BitWidth);
   }
 }
 
 void AMDGPUAtomicOptimizer::optimizeAtomic(Instruction &I,
                                            AtomicRMWInst::BinOp Op,
-                                           unsigned ValIdx,
-                                           bool ValDivergent) const {
+                                           unsigned ValIdx, bool ValDivergent) {
   // Start building just before the instruction.
   IRBuilder<> B(&I);
 
   // If we are in a pixel shader, because of how we have to mask out helper
   // lane invocations, we need to record the entry and exit BB's.
   BasicBlock *PixelEntryBB = nullptr;
   BasicBlock *PixelExitBB = nullptr;
 
if (NeedResult) {
     PHINode *const PHI = B.CreatePHI(Ty, 2);
     PHI->addIncoming(UndefValue::get(Ty), EntryBB);
     PHI->addIncoming(NewI, SingleLaneTerminator->getParent());
 
     // We need to broadcast the value who was the lowest active lane (the first
     // lane) to all other lanes in the wavefront. We use an intrinsic for this,
     // but have to handle 64-bit broadcasts with two calls to this intrinsic.
     Value *BroadcastI = nullptr;
+    Module *M = B.GetInsertBlock()->getModule();
 
     if (TyBitWidth == 64) {
       Value *const ExtractLo = B.CreateTrunc(PHI, B.getInt32Ty());
       Value *const ExtractHi =
           B.CreateTrunc(B.CreateLShr(PHI, 32), B.getInt32Ty());
-      CallInst *const ReadFirstLaneLo =
-          B.CreateIntrinsic(Intrinsic::amdgcn_readfirstlane, {}, ExtractLo);
-      CallInst *const ReadFirstLaneHi =
-          B.CreateIntrinsic(Intrinsic::amdgcn_readfirstlane, {}, ExtractHi);
+      Function *Readfirstlane = getReadfirstlane(M);
+
+      CallInst *const ReadFirstLaneLo = B.CreateCall(Readfirstlane, ExtractLo);
+      CallInst *const ReadFirstLaneHi = B.CreateCall(Readfirstlane, ExtractHi);
       Value *const PartialInsert = B.CreateInsertElement(
           UndefValue::get(VecTy), ReadFirstLaneLo, B.getInt32(0));
       Value *const Insert =
           B.CreateInsertElement(PartialInsert, ReadFirstLaneHi, B.getInt32(1));
       BroadcastI = B.CreateBitCast(Insert, Ty);
     } else if (TyBitWidth == 32) {
-
-      BroadcastI = B.CreateIntrinsic(Intrinsic::amdgcn_readfirstlane, {}, PHI);
+      BroadcastI = B.CreateCall(getReadfirstlane(M), PHI);
     } else {
       llvm_unreachable("Unhandled atomic bit width");
     }
 
     // Now that we have the result of our single atomic operation, we need to
     // get our individual lane's slice into the result. We use the lane offset
     // we previously calculated combined with the atomic result value we got
     // from the first lane, to get our lane's index into the atomic result.

llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp

if (match(Src,
       return IC.replaceInstUsesWith(II, Src);
     }
 
     if (IID == Intrinsic::amdgcn_readfirstlane) {
       // readfirstlane (readlane x, y) -> readlane x, y
       if (match(Src, PatternMatch::m_Intrinsic<Intrinsic::amdgcn_readlane>())) {
         return IC.replaceInstUsesWith(II, Src);
       }
+
+      // readfirstlane (bitcast x) -> bitcast (readfirstlane x)
+      Value *BitcastInput = nullptr;
+      if (match(Src,
+                PatternMatch::m_BitCast(PatternMatch::m_Value(BitcastInput)))) {
+        CallInst *NewCall =
+            IC.Builder.CreateIntrinsic(Intrinsic::amdgcn_readfirstlane,
+                                       {BitcastInput->getType()}, BitcastInput);
+        Value *NewCast = IC.Builder.CreateBitCast(NewCall, II.getType());
+        return IC.replaceInstUsesWith(II, NewCast);
+      }
     } else {
       // readlane (readlane x, y), y -> readlane x, y
       if (match(Src, PatternMatch::m_Intrinsic<Intrinsic::amdgcn_readlane>(
                          PatternMatch::m_Value(),
                          PatternMatch::m_Specific(II.getArgOperand(1))))) {
         return IC.replaceInstUsesWith(II, Src);
       }
     }

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

case Intrinsic::amdgcn_cvt_pk_u16: {
       if (isTypeLegal(VT))
         Results.push_back(DAG.getNode(Opcode, SL, VT, Src0, Src1));
       else {
         SDValue Cvt = DAG.getNode(Opcode, SL, MVT::i32, Src0, Src1);
         Results.push_back(DAG.getNode(ISD::BITCAST, SL, MVT::v2i16, Cvt));
       }
       return;
     }
+    case Intrinsic::amdgcn_readfirstlane: {
+      // Hack out illegal types.
+      EVT VT = N->getValueType(0);
+      if (VT.getSizeInBits() != 32)
+        return;
+
+      SDLoc DL(N);
+      SDValue Cast = DAG.getNode(ISD::BITCAST, DL, MVT::i32, N->getOperand(1));
+      SDValue Readlane = DAG.getNode(ISD::INTRINSIC_WO_CHAIN, DL, MVT::i32,
+                                     N->getOperand(0), Cast);
+      Results.push_back(DAG.getNode(ISD::BITCAST, DL, VT, Readlane));
+      return;
+    }
     }
     break;
   }
   case ISD::INTRINSIC_W_CHAIN: {
     if (SDValue Res = LowerINTRINSIC_W_CHAIN(SDValue(N, 0), DAG)) {
       if (Res.getOpcode() == ISD::MERGE_VALUES) {
         // FIXME: Hacky
         Results.push_back(Res.getOperand(0));

llvm/lib/Target/AMDGPU/SIInstructions.td

def : GCNPat<
   let SubtargetPredicate = NotHasAddNoCarryInsts;
 }
 
 
 // Avoid pointlessly materializing a constant in VGPR.
 // FIXME: Should also do this for readlane, but tablegen crashes on
 // the ignored src1.
 def : GCNPat<
-  (int_amdgcn_readfirstlane (i32 imm:$src)),
+  (i32 (int_amdgcn_readfirstlane (i32 imm:$src))),
   (S_MOV_B32 SReg_32:$src)
 >;
 
+// FIXME: Should have matcher for any size
+foreach vt = Reg32Types.types in {
+  def : GCNPat<
+    (vt (int_amdgcn_readfirstlane (vt VRegOrLds_32:$src))),
+    (V_READFIRSTLANE_B32 VRegOrLds_32:$src)
+  >;
+}
+
 multiclass BFMPatterns <ValueType vt, InstSI BFM, InstSI MOV> {
   def : GCNPat <
     (vt (shl (vt (add (vt (shl 1, vt:$a)), -1)), vt:$b)),
     (BFM $a, $b)
   >;
 
   def : GCNPat <
     (vt (add (vt (shl 1, vt:$a)), -1)),

llvm/test/CodeGen/AMDGPU/GlobalISel/inst-select-amdgcn.readfirstlane.mir

 # NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py
-# RUN: llc -march=amdgcn -mcpu=tahiti -run-pass=instruction-select -verify-machineinstrs -global-isel-abort=2 -pass-remarks-missed='gisel*' %s -o - 2> %t | FileCheck -check-prefix=GCN %s
+# RUN: llc -march=amdgcn -mcpu=tahiti -run-pass=instruction-select -verify-machineinstrs -global-isel-abort=2 -pass-remarks-missed='gisel.*' %s -o - 2> %t | FileCheck -check-prefix=GCN %s
 # RUN: FileCheck -check-prefix=ERR %s < %t
 
-# ERR: remark: <unknown>:0:0: cannot select: %1:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), %0:sgpr(s32) (in function: readfirstlane_s)
+# ERR: remark: <unknown>:0:0: cannot select: %1:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), %0:sgpr(s32) (in function: readfirstlane_s32_s)
 
 ---
-name: readfirstlane_v
+name: readfirstlane_s32_v
 legalized: true
 regBankSelected: true
 tracksRegLiveness: true
 
 body: |
   bb.0:
     liveins: $vgpr0
-    ; GCN-LABEL: name: readfirstlane_v
+    ; GCN-LABEL: name: readfirstlane_s32_v
     ; GCN: liveins: $vgpr0
     ; GCN: [[COPY:%[0-9]+]]:vgpr_32 = COPY $vgpr0
     ; GCN: [[V_READFIRSTLANE_B32_:%[0-9]+]]:sreg_32 = V_READFIRSTLANE_B32 [[COPY]], implicit $exec
     ; GCN: S_ENDPGM 0, implicit [[V_READFIRSTLANE_B32_]]
     %0:vgpr(s32) = COPY $vgpr0
     %1:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), %0
     S_ENDPGM 0, implicit %1
 ...
 
 ---
-name: readfirstlane_v_imm
+name: readfirstlane_v_s32_imm
 legalized: true
 regBankSelected: true
 tracksRegLiveness: true
 
 body: |
   bb.0:
 
-    ; GCN-LABEL: name: readfirstlane_v_imm
+    ; GCN-LABEL: name: readfirstlane_v_s32_imm
     ; GCN: [[V_MOV_B32_e32_:%[0-9]+]]:vgpr_32 = V_MOV_B32_e32 123, implicit $exec
     ; GCN: [[COPY:%[0-9]+]]:sreg_32 = COPY [[V_MOV_B32_e32_]]
     ; GCN: [[S_MOV_B32_:%[0-9]+]]:sreg_32 = S_MOV_B32 [[COPY]]
     ; GCN: S_ENDPGM 0, implicit [[S_MOV_B32_]]
     %0:vgpr(s32) = G_CONSTANT i32 123
     %1:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), %0
     S_ENDPGM 0, implicit %1
 ...
 
+---
+name: readfirstlane_v2s16_v
+legalized: true
+regBankSelected: true
+tracksRegLiveness: true
+
+body: |
+  bb.0:
+    liveins: $vgpr0
+    ; GCN-LABEL: name: readfirstlane_v2s16_v
+    ; GCN: liveins: $vgpr0
+    ; GCN: [[COPY:%[0-9]+]]:vgpr_32 = COPY $vgpr0
+    ; GCN: [[V_READFIRSTLANE_B32_:%[0-9]+]]:sreg_32 = V_READFIRSTLANE_B32 [[COPY]], implicit $exec
+    ; GCN: S_ENDPGM 0, implicit [[V_READFIRSTLANE_B32_]]
+    %0:vgpr(<2 x s16>) = COPY $vgpr0
+    %1:sgpr(<2 x s16>) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), %0
+    S_ENDPGM 0, implicit %1
+...
+
+---
+name: readfirstlane_p3_v
+legalized: true
+regBankSelected: true
+tracksRegLiveness: true
+
+body: |
+  bb.0:
+    liveins: $vgpr0
+    ; GCN-LABEL: name: readfirstlane_p3_v
+    ; GCN: liveins: $vgpr0
+    ; GCN: [[COPY:%[0-9]+]]:vgpr_32 = COPY $vgpr0
+    ; GCN: [[V_READFIRSTLANE_B32_:%[0-9]+]]:sreg_32 = V_READFIRSTLANE_B32 [[COPY]], implicit $exec
+    ; GCN: S_ENDPGM 0, implicit [[V_READFIRSTLANE_B32_]]
+    %0:vgpr(p3) = COPY $vgpr0
+    %1:sgpr(p3) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), %0
+    S_ENDPGM 0, implicit %1
+...
+
+---
+name: readfirstlane_p5_v
+legalized: true
+regBankSelected: true
+tracksRegLiveness: true
+
+body: |
+  bb.0:
+    liveins: $vgpr0
+    ; GCN-LABEL: name: readfirstlane_p5_v
+    ; GCN: liveins: $vgpr0
+    ; GCN: [[COPY:%[0-9]+]]:vgpr_32 = COPY $vgpr0
+    ; GCN: [[V_READFIRSTLANE_B32_:%[0-9]+]]:sreg_32 = V_READFIRSTLANE_B32 [[COPY]], implicit $exec
+    ; GCN: S_ENDPGM 0, implicit [[V_READFIRSTLANE_B32_]]
+    %0:vgpr(p5) = COPY $vgpr0
+    %1:sgpr(p5) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), %0
+    S_ENDPGM 0, implicit %1
+...
+
+---
+name: readfirstlane_p2_v
+legalized: true
+regBankSelected: true
+tracksRegLiveness: true
+
+body: |
+  bb.0:
+    liveins: $vgpr0
+    ; GCN-LABEL: name: readfirstlane_p2_v
+    ; GCN: liveins: $vgpr0
+    ; GCN: [[COPY:%[0-9]+]]:vgpr_32 = COPY $vgpr0
+    ; GCN: [[V_READFIRSTLANE_B32_:%[0-9]+]]:sreg_32 = V_READFIRSTLANE_B32 [[COPY]], implicit $exec
+    ; GCN: S_ENDPGM 0, implicit [[V_READFIRSTLANE_B32_]]
+    %0:vgpr(p2) = COPY $vgpr0
+    %1:sgpr(p2) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), %0
+    S_ENDPGM 0, implicit %1
+...
+
 # Make sure this fails to select
 ---
-name: readfirstlane_s
+name: readfirstlane_s32_s
 legalized: true
 regBankSelected: true
 tracksRegLiveness: true
 
 body: |
   bb.0:
     liveins: $sgpr0
-    ; GCN-LABEL: name: readfirstlane_s
+    ; GCN-LABEL: name: readfirstlane_s32_s
     ; GCN: liveins: $sgpr0
     ; GCN: [[COPY:%[0-9]+]]:sgpr(s32) = COPY $sgpr0
     ; GCN: [[INT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY]](s32)
     ; GCN: S_ENDPGM 0, implicit [[INT]](s32)
     %0:sgpr(s32) = COPY $sgpr0
     %1:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), %0
     S_ENDPGM 0, implicit %1
 ...

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

+; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=hawaii -verify-machineinstrs < %s | FileCheck -enable-var-scope %s
 ; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=fiji -verify-machineinstrs < %s | FileCheck -enable-var-scope %s
 
-declare i32 @llvm.amdgcn.readfirstlane(i32) #0
+declare i32 @llvm.amdgcn.readfirstlane.i32(i32) #0
+declare float @llvm.amdgcn.readfirstlane.f32(float) #0
+declare <2 x half> @llvm.amdgcn.readfirstlane.v2f16(<2 x half>) #0
+declare <2 x i16> @llvm.amdgcn.readfirstlane.v2i16(<2 x i16>) #0
+declare i8 addrspace(3)* @llvm.amdgcn.readfirstlane.p3i8(i8 addrspace(3)*) #0
 
-; CHECK-LABEL: {{^}}test_readfirstlane:
+; CHECK-LABEL: {{^}}test_readfirstlane_i32:
 ; CHECK: v_readfirstlane_b32 s{{[0-9]+}}, v2
-define void @test_readfirstlane(i32 addrspace(1)* %out, i32 %src) #1 {
-  %readfirstlane = call i32 @llvm.amdgcn.readfirstlane(i32 %src)
+define void @test_readfirstlane_i32(i32 addrspace(1)* %out, i32 %src) #1 {
+  %readfirstlane = call i32 @llvm.amdgcn.readfirstlane.i32(i32 %src)
   store i32 %readfirstlane, i32 addrspace(1)* %out, align 4
   ret void
 }
 
 ; CHECK-LABEL: {{^}}test_readfirstlane_imm:
 ; CHECK: s_mov_b32 [[SGPR_VAL:s[0-9]]], 32
 ; CHECK-NOT: [[SGPR_VAL]]
 ; CHECK: ; use [[SGPR_VAL]]
 define amdgpu_kernel void @test_readfirstlane_imm(i32 addrspace(1)* %out) #1 {
-  %readfirstlane = call i32 @llvm.amdgcn.readfirstlane(i32 32)
+  %readfirstlane = call i32 @llvm.amdgcn.readfirstlane.i32(i32 32)
   call void asm sideeffect "; use $0", "s"(i32 %readfirstlane)
   ret void
 }
 
 ; CHECK-LABEL: {{^}}test_readfirstlane_imm_fold:
 ; CHECK: v_mov_b32_e32 [[VVAL:v[0-9]]], 32
 ; CHECK-NOT: [[VVAL]]
 ; CHECK: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[VVAL]]
 define amdgpu_kernel void @test_readfirstlane_imm_fold(i32 addrspace(1)* %out) #1 {
-  %readfirstlane = call i32 @llvm.amdgcn.readfirstlane(i32 32)
+  %readfirstlane = call i32 @llvm.amdgcn.readfirstlane.i32(i32 32)
   store i32 %readfirstlane, i32 addrspace(1)* %out, align 4
   ret void
 }
 
 ; CHECK-LABEL: {{^}}test_readfirstlane_m0:
 ; CHECK: s_mov_b32 m0, -1
 ; CHECK: v_mov_b32_e32 [[VVAL:v[0-9]]], m0
 ; CHECK: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[VVAL]]
 define amdgpu_kernel 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)
+  %readfirstlane = call i32 @llvm.amdgcn.readfirstlane.i32(i32 %m0)
   store i32 %readfirstlane, i32 addrspace(1)* %out, align 4
   ret void
 }
 
 ; CHECK-LABEL: {{^}}test_readfirstlane_copy_from_sgpr:
 ; CHECK: ;;#ASMSTART
 ; CHECK-NEXT: s_mov_b32 [[SGPR:s[0-9]+]]
 ; CHECK: ;;#ASMEND
 ; CHECK-NOT: [[SGPR]]
 ; CHECK-NOT: readfirstlane
 ; CHECK: v_mov_b32_e32 [[VCOPY:v[0-9]+]], [[SGPR]]
 ; CHECK: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[VCOPY]]
 define amdgpu_kernel void @test_readfirstlane_copy_from_sgpr(i32 addrspace(1)* %out) #1 {
   %sgpr = call i32 asm "s_mov_b32 $0, 0", "=s"()
-  %readfirstlane = call i32 @llvm.amdgcn.readfirstlane(i32 %sgpr)
+  %readfirstlane = call i32 @llvm.amdgcn.readfirstlane.i32(i32 %sgpr)
   store i32 %readfirstlane, i32 addrspace(1)* %out, align 4
   ret void
 }
 
 ; Make sure this doesn't crash.
 ; CHECK-LABEL: {{^}}test_readfirstlane_fi:
 ; CHECK: s_mov_b32 [[FIVAL:s[0-9]]], 4
 define amdgpu_kernel void @test_readfirstlane_fi(i32 addrspace(1)* %out) #1 {
   %alloca = alloca i32, addrspace(5)
   %int = ptrtoint i32 addrspace(5)* %alloca to i32
-  %readfirstlane = call i32 @llvm.amdgcn.readfirstlane(i32 %int)
+  %readfirstlane = call i32 @llvm.amdgcn.readfirstlane.i32(i32 %int)
   call void asm sideeffect "; use $0", "s"(i32 %readfirstlane)
   ret void
 }
 
+; CHECK-LABEL: {{^}}test_readfirstlane_f32:
+; CHECK: v_readfirstlane_b32 s{{[0-9]+}}, v2
+define void @test_readfirstlane_f32(float addrspace(1)* %out, float %src) #1 {
+  %readfirstlane = call float @llvm.amdgcn.readfirstlane.f32(float %src)
+  store float %readfirstlane, float addrspace(1)* %out, align 4
+  ret void
+}
+
+; CHECK-LABEL: {{^}}test_readfirstlane_v2f16:
+; CHECK: v_readfirstlane_b32 s{{[0-9]+}}, v2
+define void @test_readfirstlane_v2f16(<2 x half> addrspace(1)* %out, <2 x half> %src) #1 {
+  %readfirstlane = call <2 x half> @llvm.amdgcn.readfirstlane.v2f16(<2 x half> %src)
+  store <2 x half> %readfirstlane, <2 x half> addrspace(1)* %out, align 4
+  ret void
+}
+
+; CHECK-LABEL: {{^}}test_readfirstlane_v2i16:
+; CHECK: v_readfirstlane_b32 s{{[0-9]+}}, v2
+define void @test_readfirstlane_v2i16(<2 x i16> addrspace(1)* %out, <2 x i16> %src) #1 {
+  %readfirstlane = call <2 x i16> @llvm.amdgcn.readfirstlane.v2i16(<2 x i16> %src)
+  store <2 x i16> %readfirstlane, <2 x i16> addrspace(1)* %out, align 4
+  ret void
+}
+
+; CHECK-LABEL: {{^}}test_readfirstlane_p3:
+; CHECK: v_readfirstlane_b32 s{{[0-9]+}}, v2
+define void @test_readfirstlane_p3(i8 addrspace(3)* addrspace(1)* %out, i8 addrspace(3)* %src) #1 {
+  %readfirstlane = call i8 addrspace(3)* @llvm.amdgcn.readfirstlane.p3i8(i8 addrspace(3)* %src)
+  store i8 addrspace(3)* %readfirstlane, i8 addrspace(3)* addrspace(1)* %out, align 4
+  ret void
+}
+
 attributes #0 = { nounwind readnone convergent }
 attributes #1 = { nounwind }

llvm/test/Transforms/InstCombine/AMDGPU/amdgcn-intrinsics.ll

 ; --------------------------------------------------------------------
 
 declare i32 @llvm.amdgcn.readfirstlane(i32)
 
 @gv = constant i32 0
 
 define amdgpu_kernel void @readfirstlane_constant(i32 %arg) {
 ; CHECK-LABEL: @readfirstlane_constant(
-; CHECK-NEXT:    [[VAR:%.*]] = call i32 @llvm.amdgcn.readfirstlane(i32 [[ARG:%.*]])
+; CHECK-NEXT:    [[VAR:%.*]] = call i32 @llvm.amdgcn.readfirstlane.i32(i32 [[ARG:%.*]])
 ; CHECK-NEXT:    store volatile i32 [[VAR]], i32* undef, align 4
 ; CHECK-NEXT:    store volatile i32 0, i32* undef, align 4
 ; CHECK-NEXT:    store volatile i32 123, i32* undef, align 4
 ; CHECK-NEXT:    store volatile i32 ptrtoint (i32* @gv to i32), i32* undef, align 4
 ; CHECK-NEXT:    store volatile i32 undef, i32* undef, align 4
 ; CHECK-NEXT:    ret void
 ;
   %var = call i32 @llvm.amdgcn.readfirstlane(i32 %arg)
   %zero = call i32 @llvm.amdgcn.readfirstlane(i32 0)
   %imm = call i32 @llvm.amdgcn.readfirstlane(i32 123)
   %constexpr = call i32 @llvm.amdgcn.readfirstlane(i32 ptrtoint (i32* @gv to i32))
   %undef = call i32 @llvm.amdgcn.readfirstlane(i32 undef)
   store volatile i32 %var, i32* undef
   store volatile i32 %zero, i32* undef
   store volatile i32 %imm, i32* undef
   store volatile i32 %constexpr, i32* undef
   store volatile i32 %undef, i32* undef
   ret void
 }
 
 define i32 @readfirstlane_idempotent(i32 %arg) {
 ; CHECK-LABEL: @readfirstlane_idempotent(
-; CHECK-NEXT:    [[READ0:%.*]] = call i32 @llvm.amdgcn.readfirstlane(i32 [[ARG:%.*]])
+; CHECK-NEXT:    [[READ0:%.*]] = call i32 @llvm.amdgcn.readfirstlane.i32(i32 [[ARG:%.*]])
 ; CHECK-NEXT:    ret i32 [[READ0]]
 ;
   %read0 = call i32 @llvm.amdgcn.readfirstlane(i32 %arg)
   %read1 = call i32 @llvm.amdgcn.readfirstlane(i32 %read0)
   %read2 = call i32 @llvm.amdgcn.readfirstlane(i32 %read1)
   ret i32 %read2
 }
 
 define i32 @readfirstlane_readlane(i32 %arg) {
 ; CHECK-LABEL: @readfirstlane_readlane(
-; CHECK-NEXT:    [[READ0:%.*]] = call i32 @llvm.amdgcn.readfirstlane(i32 [[ARG:%.*]])
+; CHECK-NEXT:    [[READ0:%.*]] = call i32 @llvm.amdgcn.readfirstlane.i32(i32 [[ARG:%.*]])
 ; CHECK-NEXT:    ret i32 [[READ0]]
 ;
   %read0 = call i32 @llvm.amdgcn.readfirstlane(i32 %arg)
   %read1 = call i32 @llvm.amdgcn.readlane(i32 %read0, i32 0)
   ret i32 %read1
 }
 
 define i32 @readfirstlane_readfirstlane_different_block(i32 %arg) {
 ; CHECK-LABEL: @readfirstlane_readfirstlane_different_block(
 ; CHECK-NEXT:  bb0:
-; CHECK-NEXT:    [[READ0:%.*]] = call i32 @llvm.amdgcn.readfirstlane(i32 [[ARG:%.*]])
+; CHECK-NEXT:    [[READ0:%.*]] = call i32 @llvm.amdgcn.readfirstlane.i32(i32 [[ARG:%.*]])
 ; CHECK-NEXT:    br label [[BB1:%.*]]
 ; CHECK:       bb1:
-; CHECK-NEXT:    [[READ1:%.*]] = call i32 @llvm.amdgcn.readfirstlane(i32 [[READ0]])
+; CHECK-NEXT:    [[READ1:%.*]] = call i32 @llvm.amdgcn.readfirstlane.i32(i32 [[READ0]])
 ; CHECK-NEXT:    ret i32 [[READ1]]
 ;
 bb0:
   %read0 = call i32 @llvm.amdgcn.readfirstlane(i32 %arg)
   br label %bb1
 
 bb1:
   %read1 = call i32 @llvm.amdgcn.readfirstlane(i32 %read0)
   ret i32 %read1
 }
 
 define i32 @readfirstlane_readlane_different_block(i32 %arg) {
 ; CHECK-LABEL: @readfirstlane_readlane_different_block(
 ; CHECK-NEXT:  bb0:
 ; CHECK-NEXT:    [[READ0:%.*]] = call i32 @llvm.amdgcn.readlane(i32 [[ARG:%.*]], i32 0)
 ; CHECK-NEXT:    br label [[BB1:%.*]]
 ; CHECK:       bb1:
-; CHECK-NEXT:    [[READ1:%.*]] = call i32 @llvm.amdgcn.readfirstlane(i32 [[READ0]])
+; CHECK-NEXT:    [[READ1:%.*]] = call i32 @llvm.amdgcn.readfirstlane.i32(i32 [[READ0]])
 ; CHECK-NEXT:    ret i32 [[READ1]]
 ;
 bb0:
   %read0 = call i32 @llvm.amdgcn.readlane(i32 %arg, i32 0)
   br label %bb1
 
 bb1:
   %read1 = call i32 @llvm.amdgcn.readfirstlane(i32 %read0)
   ret i32 %read1
 }
 
+define i32 @readfirstlane_bitcast(float %arg) {
+; CHECK-LABEL: @readfirstlane_bitcast(
+; CHECK-NEXT:    [[TMP1:%.*]] = call float @llvm.amdgcn.readfirstlane.f32(float [[ARG:%.*]])
+; CHECK-NEXT:    [[TMP2:%.*]] = bitcast float [[TMP1]] to i32
+; CHECK-NEXT:    ret i32 [[TMP2]]
+;
+  %bitcast.arg = bitcast float %arg to i32
+  %read = call i32 @llvm.amdgcn.readfirstlane(i32 %bitcast.arg)
+  ret i32 %read
+}
+
+define float @bitcast_readfirstlane_bitcast(float %arg) {
+; CHECK-LABEL: @bitcast_readfirstlane_bitcast(
+; CHECK-NEXT:    [[TMP1:%.*]] = call float @llvm.amdgcn.readfirstlane.f32(float [[ARG:%.*]])
+; CHECK-NEXT:    ret float [[TMP1]]
+;
+  %bitcast.arg = bitcast float %arg to i32
+  %read = call i32 @llvm.amdgcn.readfirstlane(i32 %bitcast.arg)
+  %cast.read = bitcast i32 %read to float
+  ret float %cast.read
+}
+
+define i32 @readfirstlane_bitcast_multi_use(float %arg) {
+; CHECK-LABEL: @readfirstlane_bitcast_multi_use(
+; CHECK-NEXT:    store float [[ARG:%.*]], float* undef, align 4
+; CHECK-NEXT:    [[TMP1:%.*]] = call float @llvm.amdgcn.readfirstlane.f32(float [[ARG]])
+; CHECK-NEXT:    [[TMP2:%.*]] = bitcast float [[TMP1]] to i32
+; CHECK-NEXT:    ret i32 [[TMP2]]
+;
+  %bitcast.arg = bitcast float %arg to i32
+  store i32 %bitcast.arg, i32* undef
+  %read = call i32 @llvm.amdgcn.readfirstlane(i32 %bitcast.arg)
+  ret i32 %read
+}
+
 ; --------------------------------------------------------------------
 ; llvm.amdgcn.readlane
 ; --------------------------------------------------------------------
 
 declare i32 @llvm.amdgcn.readlane(i32, i32)
 
 define amdgpu_kernel void @readlane_constant(i32 %arg, i32 %lane) {
 ; CHECK-LABEL: @readlane_constant(

 ;
   %read0 = call i32 @llvm.amdgcn.readlane(i32 %arg, i32 %lane0)
   %read1 = call i32 @llvm.amdgcn.readlane(i32 %read0, i32 %lane1)
   ret i32 %read1
 }
 
 define i32 @readlane_readfirstlane(i32 %arg) {
 ; CHECK-LABEL: @readlane_readfirstlane(
-; CHECK-NEXT:    [[READ0:%.*]] = call i32 @llvm.amdgcn.readfirstlane(i32 [[ARG:%.*]])
+; CHECK-NEXT:    [[READ0:%.*]] = call i32 @llvm.amdgcn.readfirstlane.i32(i32 [[ARG:%.*]])
 ; CHECK-NEXT:    ret i32 [[READ0]]
 ;
   %read0 = call i32 @llvm.amdgcn.readfirstlane(i32 %arg)
   %read1 = call i32 @llvm.amdgcn.readlane(i32 %read0, i32 0)
   ret i32 %read1
 }
 
 define i32 @readlane_idempotent_different_block(i32 %arg, i32 %lane) {
bb1:
   %read1 = call i32 @llvm.amdgcn.readlane(i32 %read0, i32 %lane)
   ret i32 %read1
 }
 
 
 define i32 @readlane_readfirstlane_different_block(i32 %arg) {
 ; CHECK-LABEL: @readlane_readfirstlane_different_block(
 ; CHECK-NEXT:  bb0:
-; CHECK-NEXT:    [[READ0:%.*]] = call i32 @llvm.amdgcn.readfirstlane(i32 [[ARG:%.*]])
+; CHECK-NEXT:    [[READ0:%.*]] = call i32 @llvm.amdgcn.readfirstlane.i32(i32 [[ARG:%.*]])
 ; CHECK-NEXT:    br label [[BB1:%.*]]
 ; CHECK:       bb1:
 ; CHECK-NEXT:    [[READ1:%.*]] = call i32 @llvm.amdgcn.readlane(i32 [[READ0]], i32 0)
 ; CHECK-NEXT:    ret i32 [[READ1]]
 ;
 bb0:
   %read0 = call i32 @llvm.amdgcn.readfirstlane(i32 %arg)
   br label %bb1