AMDGPU: Add llvm.amdgcn.{read,readfirst,write}lane2 intrinsics with type overloads

[nhaehnle]

Summary

These intrinsics should work at least with standard integer and floating
point sizes, pointers, and vectors of those.

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

Moving the atomic optimizer pass in the pass pipeline so that it can be
simplified and rely on the more general support of lane intrinsics.

API users should move to these new intrinsics so that we can remove the
old versions.

Change-Id: I1c5e7e7858890e1c30d3b46c8551e74ab7027552
Based-on: https://reviews.llvm.org/D84639

[arsenm]

Do we really have to use worse names here? Keeping the name works even if suboptimal for the attributes

[nhaehnle]

Note that part of my motivation here over D84639 is to support more general types on the lane intrinsics, since they also express some semantic content which would be interesting to be able to express e.g. on descriptors. I wasn't able to bend the SelectionDAG type legalization to my will, so that's why I instead "legalize" the intrinsics in the AMDGPUCodeGenPrepare pass.

[arsenm]

In D86154#2224270, @nhaehnle wrote:

Note that part of my motivation here over D84639 is to support more general types on the lane intrinsics, since they also express some semantic content which would be interesting to be able to express e.g. on descriptors. I wasn't able to bend the SelectionDAG type legalization to my will, so that's why I instead "legalize" the intrinsics in the AMDGPUCodeGenPrepare pass.

Don't you just need to handle this in ReplaceNodeResults the same way?

[nhaehnle]

Don't duplicate the intrinsics. Rely on D86317 to reduce the pain of this
change caused to downstream users.

[nhaehnle]

In D86154#2224272, @arsenm wrote:

In D86154#2224270, @nhaehnle wrote:

Note that part of my motivation here over D84639 is to support more general types on the lane intrinsics, since they also express some semantic content which would be interesting to be able to express e.g. on descriptors. I wasn't able to bend the SelectionDAG type legalization to my will, so that's why I instead "legalize" the intrinsics in the AMDGPUCodeGenPrepare pass.

Don't you just need to handle this in ReplaceNodeResults the same way?

ReplaceNodeResults expects the result type to be changed in semi-magical ways during vector type legalization, which is non-obvious since the method can be called from different places. I think it *could* be made to work with a lot of patience, but it's really a bad interface -- and besides, by doing it in IR we reduce code duplication between SelectionDAG and GlobalISel, which is an added benefit IMO.

[arsenm]

In D86154#2229292, @nhaehnle wrote:

In D86154#2224272, @arsenm wrote:

In D86154#2224270, @nhaehnle wrote:

Note that part of my motivation here over D84639 is to support more general types on the lane intrinsics, since they also express some semantic content which would be interesting to be able to express e.g. on descriptors. I wasn't able to bend the SelectionDAG type legalization to my will, so that's why I instead "legalize" the intrinsics in the AMDGPUCodeGenPrepare pass.

Don't you just need to handle this in ReplaceNodeResults the same way?

ReplaceNodeResults expects the result type to be changed in semi-magical ways during vector type legalization, which is non-obvious since the method can be called from different places. I think it *could* be made to work with a lot of patience, but it's really a bad interface -- and besides, by doing it in IR we reduce code duplication between SelectionDAG and GlobalISel, which is an added benefit IMO.

Well the globalisel handling should be much simpler. We have a lot of stuff that's randomly handled in the IR to work around the DAG which long term should be moved where it belongs in codegen

[nhaehnle]

ReplaceNodeResults expects the result type to be changed in semi-magical ways during vector type legalization, which is non-obvious since the method can be called from different places. I think it *could* be made to work with a lot of patience, but it's really a bad interface -- and besides, by doing it in IR we reduce code duplication between SelectionDAG and GlobalISel, which is an added benefit IMO.

Well the globalisel handling should be much simpler. We have a lot of stuff that's randomly handled in the IR to work around the DAG which long term should be moved where it belongs in codegen

Is the GlobalISel handling simpler than the handling in IR? What would happen if we wanted to extend the handling to struct types?

[arsenm]

Should be obsoleted by D147732

[nhaehnle]

Indeed.

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_readfirstlane2);
+  case AMDGPU::BI__builtin_amdgcn_readlane:
+    return emitBinaryBuiltin(*this, E, Intrinsic::amdgcn_readlane2);
   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.readfirstlane2.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)
+// CHECK: call i32 @llvm.amdgcn.readlane2.i32(i32 %a, i32 %b)
 void test_readlane(global int* out, int a, int b)
 {
   *out = __builtin_amdgcn_readlane(a, b);
 }
 
 // CHECK-LABEL: @test_fcmp_f32
 // CHECK: call i64 @llvm.amdgcn.fcmp.i64.f32(float %a, float %b, i32 5)
 void test_fcmp_f32(global ulong* out, float a, float b)

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

def int_amdgcn_fcmp :
   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],
+def int_amdgcn_readfirstlane2 :
+  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],
+def int_amdgcn_readlane2 :
+  Intrinsic<[llvm_any_ty],
+            [LLVMMatchType<0>,  // data input
+             llvm_i32_ty],      // uniform lane select
             [IntrNoMem, IntrConvergent, IntrWillReturn]>;
 
 // The value to write and lane select arguments must be uniform across the
 // currently active threads of the current wave. Otherwise, the result is
 // undefined.
+def int_amdgcn_writelane2 :
+  Intrinsic<[llvm_any_ty], [
+    LLVMMatchType<0>,  // uniform value to write: returned by the selected lane
+    llvm_i32_ty,       // uniform lane select
+    LLVMMatchType<0>   // returned by all lanes other than the selected one
+  ],
+  [IntrNoMem, IntrConvergent, IntrWillReturn]
+>;
+
+// Non-overloaded versions of readfirstlane2 / readlane2 / writelane2.
+def int_amdgcn_readfirstlane :
+  Intrinsic<[llvm_i32_ty], [llvm_i32_ty],
+            [IntrNoMem, IntrConvergent, IntrWillReturn]>;
+def int_amdgcn_readlane :
+  Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
+            [IntrNoMem, IntrConvergent, IntrWillReturn]>;
 def int_amdgcn_writelane :
-  GCCBuiltin<"__builtin_amdgcn_writelane">,
   Intrinsic<[llvm_i32_ty], [
     llvm_i32_ty,    // uniform value to write: returned by the selected lane
     llvm_i32_ty,    // uniform lane select
     llvm_i32_ty     // returned by all lanes other than the selected one
   ],
   [IntrNoMem, IntrConvergent, IntrWillReturn]
 >;
 

llvm/lib/Target/AMDGPU/AMDGPUAtomicOptimizer.cpp

 
 // Use the builder to create an inclusive scan of V across the wavefront, with
 // all lanes active.
 Value *AMDGPUAtomicOptimizer::buildScan(IRBuilder<> &B, AtomicRMWInst::BinOp Op,
                                         Value *V, Value *const Identity) const {
   Type *const Ty = V->getType();
   Module *M = B.GetInsertBlock()->getModule();
   Function *UpdateDPP =
-      Intrinsic::getDeclaration(M, Intrinsic::amdgcn_update_dpp, Ty);
+      Intrinsic::getDeclaration(M, Intrinsic::amdgcn_update_dpp, {Ty});
   Function *PermLaneX16 =
       Intrinsic::getDeclaration(M, Intrinsic::amdgcn_permlanex16, {});
   Function *ReadLane =
-      Intrinsic::getDeclaration(M, Intrinsic::amdgcn_readlane, {});
+      Intrinsic::getDeclaration(M, Intrinsic::amdgcn_readlane2, {Ty});
 
   for (unsigned Idx = 0; Idx < 4; Idx++) {
     V = buildNonAtomicBinOp(
         B, Op, V,
         B.CreateCall(UpdateDPP,
                      {Identity, V, B.getInt32(DPP::ROW_SHR0 | 1 << Idx),
                       B.getInt32(0xf), B.getInt32(0xf), B.getFalse()}));
   }

 
 // Use the builder to create a shift right of V across the wavefront, with all
 // lanes active, to turn an inclusive scan into an exclusive scan.
 Value *AMDGPUAtomicOptimizer::buildShiftRight(IRBuilder<> &B, Value *V,
                                               Value *const Identity) const {
   Type *const Ty = V->getType();
   Module *M = B.GetInsertBlock()->getModule();
   Function *UpdateDPP =
-      Intrinsic::getDeclaration(M, Intrinsic::amdgcn_update_dpp, Ty);
+      Intrinsic::getDeclaration(M, Intrinsic::amdgcn_update_dpp, {Ty});
   Function *ReadLane =
-      Intrinsic::getDeclaration(M, Intrinsic::amdgcn_readlane, {});
+      Intrinsic::getDeclaration(M, Intrinsic::amdgcn_readlane2, {Ty});
   Function *WriteLane =
-      Intrinsic::getDeclaration(M, Intrinsic::amdgcn_writelane, {});
+      Intrinsic::getDeclaration(M, Intrinsic::amdgcn_writelane2, {Ty});
 
   if (ST->hasDPPWavefrontShifts()) {
     // GFX9 has DPP wavefront shift operations.
     V = B.CreateCall(UpdateDPP,
                      {Identity, V, B.getInt32(DPP::WAVE_SHR1), B.getInt32(0xf),
                       B.getInt32(0xf), B.getFalse()});
   } else {
     // On GFX10 all DPP operations are confined to a single row. To get cross-
const AtomicRMWInst::BinOp ScanOp =
         Op == AtomicRMWInst::Sub ? AtomicRMWInst::Add : Op;
     NewV = buildScan(B, ScanOp, NewV, Identity);
     ExclScan = buildShiftRight(B, NewV, Identity);
 
     // Read the value from the last lane, which has accumlated the values of
     // each active lane in the wavefront. This will be our new value which we
     // will provide to the atomic operation.
     Value *const LastLaneIdx = B.getInt32(ST->getWavefrontSize() - 1);
-    if (TyBitWidth == 64) {
-      Value *const ExtractLo = B.CreateTrunc(NewV, B.getInt32Ty());
-      Value *const ExtractHi =
-          B.CreateTrunc(B.CreateLShr(NewV, 32), B.getInt32Ty());
-      CallInst *const ReadLaneLo = B.CreateIntrinsic(
-          Intrinsic::amdgcn_readlane, {}, {ExtractLo, LastLaneIdx});
-      CallInst *const ReadLaneHi = B.CreateIntrinsic(
-          Intrinsic::amdgcn_readlane, {}, {ExtractHi, LastLaneIdx});
-      Value *const PartialInsert = B.CreateInsertElement(
-          UndefValue::get(VecTy), ReadLaneLo, B.getInt32(0));
-      Value *const Insert =
-          B.CreateInsertElement(PartialInsert, ReadLaneHi, B.getInt32(1));
-      NewV = B.CreateBitCast(Insert, Ty);
-    } else if (TyBitWidth == 32) {
-      NewV = B.CreateIntrinsic(Intrinsic::amdgcn_readlane, {},
-                               {NewV, LastLaneIdx});
-    } else {
-      llvm_unreachable("Unhandled atomic bit width");
-    }
+    NewV = B.CreateIntrinsic(Intrinsic::amdgcn_readlane2, {Ty},
 
     // Finally mark the readlanes in the WWM section.
     NewV = B.CreateIntrinsic(Intrinsic::amdgcn_wwm, Ty, NewV);
   } else {
     switch (Op) {
     default:
       llvm_unreachable("Unhandled atomic op");
 
if (NeedResult) {
     // Create a PHI node to get our new atomic result into the exit block.
     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;
-
-    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);
-      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);
-    } else {
-      llvm_unreachable("Unhandled atomic bit width");
-    }
+    Value *BroadcastI =
+        B.CreateIntrinsic(Intrinsic::amdgcn_readfirstlane2, {Ty}, {PHI});
 
     // 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.
     Value *LaneOffset = nullptr;
     if (ValDivergent) {
       LaneOffset = B.CreateIntrinsic(Intrinsic::amdgcn_wwm, Ty, ExclScan);

llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp

 //===-- llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp - Call lowering -----===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 ///
 /// \file
void assignValueToReg(Register ValVReg, Register PhysReg,
     Register ExtReg = extendRegisterMin32(ValVReg, VA);
 
     // If this is a scalar return, insert a readfirstlane just in case the value
     // ends up in a VGPR.
     // FIXME: Assert this is a shader return.
     const SIRegisterInfo *TRI
       = static_cast<const SIRegisterInfo *>(MRI.getTargetRegisterInfo());
     if (TRI->isSGPRReg(MRI, PhysReg)) {
-      auto ToSGPR = MIRBuilder.buildIntrinsic(Intrinsic::amdgcn_readfirstlane,
+      auto ToSGPR = MIRBuilder.buildIntrinsic(Intrinsic::amdgcn_readfirstlane2,
                                               {MRI.getType(ExtReg)}, false)
         .addReg(ExtReg);
       ExtReg = ToSGPR.getReg(0);
     }
 
     MIRBuilder.buildCopy(PhysReg, ExtReg);
     MIB.addUse(PhysReg, RegState::Implicit);
   }

llvm/lib/Target/AMDGPU/AMDGPUCodeGenPrepare.cpp

public:
   bool visitInstruction(Instruction &I) { return false; }
   bool visitBinaryOperator(BinaryOperator &I);
   bool visitLoadInst(LoadInst &I);
   bool visitICmpInst(ICmpInst &I);
   bool visitSelectInst(SelectInst &I);
 
   bool visitIntrinsicInst(IntrinsicInst &I);
   bool visitBitreverseIntrinsicInst(IntrinsicInst &I);
+  bool visitLaneIntrinsicInst(IntrinsicInst &I, Intrinsic::ID CanonicalIID);
+  Value *buildLegalLaneIntrinsic(IRBuilder<> &B, Intrinsic::ID IID,
+                                 Value *Data0, Value *Lane = nullptr,
+                                 Value *Data1 = nullptr);
 
   bool doInitialization(Module &M) override;
   bool runOnFunction(Function &F) override;
 
   StringRef getPassName() const override { return "AMDGPU IR optimizations"; }
 
   void getAnalysisUsage(AnalysisUsage &AU) const override {
     AU.addRequired<AssumptionCacheTracker>();
bool AMDGPUCodeGenPrepare::visitSelectInst(SelectInst &I) {
 
   return Changed;
 }
 
 bool AMDGPUCodeGenPrepare::visitIntrinsicInst(IntrinsicInst &I) {
   switch (I.getIntrinsicID()) {
   case Intrinsic::bitreverse:
     return visitBitreverseIntrinsicInst(I);
+  case Intrinsic::amdgcn_readfirstlane:
+    return visitLaneIntrinsicInst(I, Intrinsic::amdgcn_readfirstlane2);
+  case Intrinsic::amdgcn_readlane:
+    return visitLaneIntrinsicInst(I, Intrinsic::amdgcn_readlane2);
+  case Intrinsic::amdgcn_writelane:
+    return visitLaneIntrinsicInst(I, Intrinsic::amdgcn_writelane2);
+  case Intrinsic::amdgcn_readfirstlane2:
+  case Intrinsic::amdgcn_readlane2:
+  case Intrinsic::amdgcn_writelane2:
+    return visitLaneIntrinsicInst(I, I.getIntrinsicID());
   default:
     return false;
   }
 }
 
 bool AMDGPUCodeGenPrepare::visitBitreverseIntrinsicInst(IntrinsicInst &I) {
   bool Changed = false;
 
   if (ST->has16BitInsts() && needsPromotionToI32(I.getType()) &&
       DA->isUniform(&I))
     Changed |= promoteUniformBitreverseToI32(I);
 
   return Changed;
 }
 
+Value *AMDGPUCodeGenPrepare::buildLegalLaneIntrinsic(IRBuilder<> &B,
+                                                     Intrinsic::ID IID,
+                                                     Value *Data0, Value *Lane,
+                                                     Value *Data1) {
+  Type *Ty = Data0->getType();
+
+  if (Ty == B.getInt32Ty()) {
+    Value *Args[3] = {Data0, Lane, Data1};
+    unsigned NumArgs = Data1 != nullptr ? 3 : Lane != nullptr ? 2 : 1;
+    return B.CreateIntrinsic(IID, {B.getInt32Ty()}, {Args, NumArgs});
+  }
+
+  if (auto *VecTy = dyn_cast<FixedVectorType>(Ty)) {
+    Type *EltType = VecTy->getElementType();
+    bool is16Bit =
+        (EltType->isIntegerTy() && EltType->getIntegerBitWidth() == 16) ||
+        (EltType->isHalfTy());
+    int EC = VecTy->getElementCount().Min;
+
+    Value *Result = UndefValue::get(Ty);
+    for (int i = 0; i < EC; i += 1 + is16Bit) {
+      Value *EltData0;
+      Value *EltData1 = nullptr;
+
+      if (is16Bit) {
+        int Idxs[2] = {i, i + 1};
+        EltData0 = B.CreateShuffleVector(Data0, UndefValue::get(Ty), Idxs);
+        EltData0 = B.CreateBitCast(EltData0, B.getInt32Ty());
+      } else {
+        EltData0 = B.CreateExtractElement(Data0, i);
+      }
+
+      if (Data1) {
+        if (is16Bit) {
+          int Idxs[2] = {i, i + 1};
+          EltData1 = B.CreateShuffleVector(Data1, UndefValue::get(Ty), Idxs);
+          EltData1 = B.CreateBitCast(EltData1, B.getInt32Ty());
+        } else {
+          EltData1 = B.CreateExtractElement(Data1, i);
+        }
+      }
+
+      Value *EltResult =
+          buildLegalLaneIntrinsic(B, IID, EltData0, Lane, EltData1);
+
+      if (is16Bit) {
+        EltResult =
+            B.CreateBitCast(EltResult, FixedVectorType::get(EltType, 2));
+        for (int j = 0; j < 2; ++j) {
+          if (i + j >= EC)
+            break;
+          Result = B.CreateInsertElement(
+              Result, B.CreateExtractElement(EltResult, j), i + j);
+        }
+      } else {
+        Result = B.CreateInsertElement(Result, EltResult, i);
+      }
+    }
+
+    return Result;
+  }
+
+  unsigned BitWidth = DL->getTypeSizeInBits(Ty);
+  Type *IntTy = Ty;
+
+  if (!Ty->isIntegerTy()) {
+    IntTy = IntegerType::get(Mod->getContext(), BitWidth);
+    Data0 = B.CreateBitOrPointerCast(Data0, IntTy);
+    if (Data1)
+      Data1 = B.CreateBitOrPointerCast(Data1, IntTy);
+  }
+
+  if ((BitWidth % 32) != 0) {
+    Type *ExtendedTy =
+        IntegerType::get(Mod->getContext(), (BitWidth + 31) & ~31);
+    Data0 = B.CreateZExt(Data0, ExtendedTy);
+    if (Data1)
+      Data1 = B.CreateZExt(Data1, ExtendedTy);
+  }
+
+  if (BitWidth > 32) {
+    Type *VecTy = FixedVectorType::get(B.getInt32Ty(), (BitWidth + 31) / 32);
+    Data0 = B.CreateBitCast(Data0, VecTy);
+    if (Data1)
+      Data1 = B.CreateBitCast(Data1, VecTy);
+  }
+
+  Value *Result = buildLegalLaneIntrinsic(B, IID, Data0, Lane, Data1);
+
+  if ((BitWidth % 32) != 0) {
+    if (BitWidth > 32) {
+      Result = B.CreateBitCast(
+          Result, IntegerType::get(Mod->getContext(), (BitWidth + 31) / 32));
+    }
+
+    Result =
+        B.CreateTrunc(Result, IntegerType::get(Mod->getContext(), BitWidth));
+  }
+
+  return B.CreateBitOrPointerCast(Result, Ty);
+}
+
+/// "Legalize" readfirstlane/readlane/writelane to single-dword intrinsics
+/// on i32.
+///
+/// Done during codegen prepare purely because this turned out to be simpler
+/// than doing it in this generality in SelectionDAG.
+bool AMDGPUCodeGenPrepare::visitLaneIntrinsicInst(IntrinsicInst &I,
+                                                  Intrinsic::ID CanonicalIID) {
+  Type *Ty = I.getType();
+  if (I.getIntrinsicID() == CanonicalIID && Ty->isIntegerTy(32) &&
+      Ty->getIntegerBitWidth() == 32)
+    return false; // already legal
+
+  Value *Data0 = I.getArgOperand(0);
+  Value *Lane = nullptr;
+  Value *Data1 = nullptr;
+
+  if (CanonicalIID == Intrinsic::amdgcn_readlane2) {
+    Lane = I.getArgOperand(1);
+  } else if (CanonicalIID == Intrinsic::amdgcn_writelane2) {
+    Lane = I.getArgOperand(1);
+    Data1 = I.getArgOperand(2);
+  }
+
+  IRBuilder<> Builder(&I);
+  Value *Legalized =
+      buildLegalLaneIntrinsic(Builder, CanonicalIID, Data0, Lane, Data1);
+
+  I.replaceAllUsesWith(Legalized);
+  I.eraseFromParent();
+  return true;
+}
+
 bool AMDGPUCodeGenPrepare::doInitialization(Module &M) {
   Mod = &M;
   DL = &Mod->getDataLayout();
   return false;
 }
 
 bool AMDGPUCodeGenPrepare::runOnFunction(Function &F) {
   if (skipFunction(F))

llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp

switch (N->getOpcode()) {
     case ISD::TokenFactor:
       return true;
     case ISD::INTRINSIC_WO_CHAIN:
     {
       unsigned IntrID = cast<ConstantSDNode>(N->getOperand(0))->getZExtValue();
       switch (IntrID) {
         default:
         return false;
+        case Intrinsic::amdgcn_readfirstlane2:
+        case Intrinsic::amdgcn_readlane2:
+          return true;
         case Intrinsic::amdgcn_readfirstlane:
         case Intrinsic::amdgcn_readlane:
-          return true;
+          llvm_unreachable("should have been updated during CodeGenPrepare");
       }
     }
     break;
     case ISD::LOAD:
     {
       if (cast<LoadSDNode>(N)->getMemOperand()->getAddrSpace() ==
           AMDGPUAS::CONSTANT_ADDRESS_32BIT)
         return true;

llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp

case Intrinsic::amdgcn_permlanex16: {
     ConstantInt *FetchInvalid = cast<ConstantInt>(II.getArgOperand(4));
     ConstantInt *BoundCtrl = cast<ConstantInt>(II.getArgOperand(5));
     if (!FetchInvalid->getZExtValue() && !BoundCtrl->getZExtValue())
       break;
 
     return IC.replaceOperand(II, 0, UndefValue::get(VDstIn->getType()));
   }
   case Intrinsic::amdgcn_readfirstlane:
-  case Intrinsic::amdgcn_readlane: {
+  case Intrinsic::amdgcn_readlane:
+  case Intrinsic::amdgcn_writelane: {
+    // Canonicalize to overloaded intrinsics.
+    Intrinsic::ID NewIID;
+    if (IID == Intrinsic::amdgcn_readfirstlane)
+      NewIID = Intrinsic::amdgcn_readfirstlane2;
+    else if (IID == Intrinsic::amdgcn_readlane)
+      NewIID = Intrinsic::amdgcn_readlane2;
+    else
+      NewIID = Intrinsic::amdgcn_writelane2;
+
+    SmallVector<Value *, 3> Args;
+    for (Value *Arg : II.arg_operands())
+      Args.push_back(Arg);
+    CallInst *UpgradedCall = IC.Builder.CreateIntrinsic(
+        NewIID, {II.getArgOperand(0)->getType()}, Args);
+    return IC.replaceInstUsesWith(II, UpgradedCall);
+  }
+  case Intrinsic::amdgcn_readfirstlane2:
+  case Intrinsic::amdgcn_readlane2: {
     // A constant value is trivially uniform.
     if (Constant *C = dyn_cast<Constant>(II.getArgOperand(0))) {
       return IC.replaceInstUsesWith(II, C);
     }
 
     // The rest of these may not be safe if the exec may not be the same between
     // the def and use.
     Value *Src = II.getArgOperand(0);
     Instruction *SrcInst = dyn_cast<Instruction>(Src);
     if (SrcInst && SrcInst->getParent() != II.getParent())
       break;
 
     // readfirstlane (readfirstlane x) -> readfirstlane x
     // readlane (readfirstlane x), y -> readfirstlane x
     if (match(Src,
-              PatternMatch::m_Intrinsic<Intrinsic::amdgcn_readfirstlane>())) {
+              PatternMatch::m_Intrinsic<Intrinsic::amdgcn_readfirstlane2>())) {
       return IC.replaceInstUsesWith(II, Src);
     }
 
-    if (IID == Intrinsic::amdgcn_readfirstlane) {
+    if (IID == Intrinsic::amdgcn_readfirstlane2) {
       // readfirstlane (readlane x, y) -> readlane x, y
-      if (match(Src, PatternMatch::m_Intrinsic<Intrinsic::amdgcn_readlane>())) {
+      if (match(Src,
+                PatternMatch::m_Intrinsic<Intrinsic::amdgcn_readlane2>())) {
         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_readfirstlane2,
+                                       {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>(
+      if (match(Src, PatternMatch::m_Intrinsic<Intrinsic::amdgcn_readlane2>(
                          PatternMatch::m_Value(),
                          PatternMatch::m_Specific(II.getArgOperand(1))))) {
         return IC.replaceInstUsesWith(II, Src);
       }
     }
 
     break;
   }

llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp

bool AMDGPUInstructionSelector::selectG_INTRINSIC(MachineInstr &I) const {
   case Intrinsic::amdgcn_interp_p1_f16:
     return selectInterpP1F16(I);
   case Intrinsic::amdgcn_wqm:
     return constrainCopyLikeIntrin(I, AMDGPU::WQM);
   case Intrinsic::amdgcn_softwqm:
     return constrainCopyLikeIntrin(I, AMDGPU::SOFT_WQM);
   case Intrinsic::amdgcn_wwm:
     return constrainCopyLikeIntrin(I, AMDGPU::WWM);
-  case Intrinsic::amdgcn_writelane:
+  case Intrinsic::amdgcn_writelane2:
     return selectWritelane(I);
   case Intrinsic::amdgcn_div_scale:
     return selectDivScale(I);
   case Intrinsic::amdgcn_icmp:
     return selectIntrinsicIcmp(I);
   case Intrinsic::amdgcn_ballot:
     return selectBallot(I);
   case Intrinsic::amdgcn_reloc_constant:

llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp

AMDGPURegisterBankInfo::addMappingFromTable(
 
   return AltMappings;
 }
 
 RegisterBankInfo::InstructionMappings
 AMDGPURegisterBankInfo::getInstrAlternativeMappingsIntrinsic(
     const MachineInstr &MI, const MachineRegisterInfo &MRI) const {
   switch (MI.getIntrinsicID()) {
-  case Intrinsic::amdgcn_readlane: {
+  case Intrinsic::amdgcn_readlane2: {
     static const OpRegBankEntry<3> Table[2] = {
       // Perfectly legal.
       { { AMDGPU::SGPRRegBankID, AMDGPU::VGPRRegBankID, AMDGPU::SGPRRegBankID }, 1 },
 
       // Need a readfirstlane for the index.
       { { AMDGPU::SGPRRegBankID, AMDGPU::VGPRRegBankID, AMDGPU::VGPRRegBankID }, 2 }
     };
 
     const std::array<unsigned, 3> RegSrcOpIdx = { { 0, 2, 3 } };
     return addMappingFromTable<3>(MI, MRI, RegSrcOpIdx, makeArrayRef(Table));
   }
-  case Intrinsic::amdgcn_writelane: {
+  case Intrinsic::amdgcn_writelane2: {
     static const OpRegBankEntry<4> Table[4] = {
       // Perfectly legal.
       { { AMDGPU::VGPRRegBankID, AMDGPU::SGPRRegBankID, AMDGPU::SGPRRegBankID, AMDGPU::VGPRRegBankID }, 1 },
 
       // Need readfirstlane of first op
       { { AMDGPU::VGPRRegBankID, AMDGPU::VGPRRegBankID, AMDGPU::SGPRRegBankID, AMDGPU::VGPRRegBankID }, 2 },
 
       // Need readfirstlane of second op
case AMDGPU::G_AMDGPU_BUFFER_ATOMIC_CMPSWAP: {
     return;
   }
   case AMDGPU::G_AMDGPU_S_BUFFER_LOAD: {
     applyMappingSBufferLoad(OpdMapper);
     return;
   }
   case AMDGPU::G_INTRINSIC: {
     switch (MI.getIntrinsicID()) {
-    case Intrinsic::amdgcn_readlane: {
+    case Intrinsic::amdgcn_readlane2: {
       substituteSimpleCopyRegs(OpdMapper, 2);
 
       assert(OpdMapper.getVRegs(0).empty());
       assert(OpdMapper.getVRegs(3).empty());
 
       // Make sure the index is an SGPR. It doesn't make sense to run this in a
       // waterfall loop, so assume it's a uniform value.
       constrainOpWithReadfirstlane(MI, MRI, 3); // Index
       return;
     }
-    case Intrinsic::amdgcn_writelane: {
+    case Intrinsic::amdgcn_writelane2: {
       assert(OpdMapper.getVRegs(0).empty());
       assert(OpdMapper.getVRegs(2).empty());
       assert(OpdMapper.getVRegs(3).empty());
 
       substituteSimpleCopyRegs(OpdMapper, 4); // VGPR input val
       constrainOpWithReadfirstlane(MI, MRI, 2); // Source value
       constrainOpWithReadfirstlane(MI, MRI, 3); // Index
       return;
case Intrinsic::amdgcn_fcmp: {
       unsigned DstSize = MRI.getType(MI.getOperand(0).getReg()).getSizeInBits();
       // This is not VCCRegBank because this is not used in boolean contexts.
       OpdsMapping[0] = AMDGPU::getValueMapping(AMDGPU::SGPRRegBankID, DstSize);
       unsigned OpSize = MRI.getType(MI.getOperand(2).getReg()).getSizeInBits();
       OpdsMapping[2] = AMDGPU::getValueMapping(AMDGPU::VGPRRegBankID, OpSize);
       OpdsMapping[3] = AMDGPU::getValueMapping(AMDGPU::VGPRRegBankID, OpSize);
       break;
     }
-    case Intrinsic::amdgcn_readlane: {
+    case Intrinsic::amdgcn_readlane2: {
       // This must be an SGPR, but accept a VGPR.
       Register IdxReg = MI.getOperand(3).getReg();
       unsigned IdxSize = MRI.getType(IdxReg).getSizeInBits();
       unsigned IdxBank = getRegBankID(IdxReg, MRI, AMDGPU::SGPRRegBankID);
       OpdsMapping[3] = AMDGPU::getValueMapping(IdxBank, IdxSize);
       LLVM_FALLTHROUGH;
     }
-    case Intrinsic::amdgcn_readfirstlane: {
+    case Intrinsic::amdgcn_readfirstlane2: {
       unsigned DstSize = MRI.getType(MI.getOperand(0).getReg()).getSizeInBits();
       unsigned SrcSize = MRI.getType(MI.getOperand(2).getReg()).getSizeInBits();
       OpdsMapping[0] = AMDGPU::getValueMapping(AMDGPU::SGPRRegBankID, DstSize);
       OpdsMapping[2] = AMDGPU::getValueMapping(AMDGPU::VGPRRegBankID, SrcSize);
       break;
     }
-    case Intrinsic::amdgcn_writelane: {
+    case Intrinsic::amdgcn_writelane2: {
       unsigned DstSize = MRI.getType(MI.getOperand(0).getReg()).getSizeInBits();
       Register SrcReg = MI.getOperand(2).getReg();
       unsigned SrcSize = MRI.getType(SrcReg).getSizeInBits();
       unsigned SrcBank = getRegBankID(SrcReg, MRI, AMDGPU::SGPRRegBankID);
       Register IdxReg = MI.getOperand(3).getReg();
       unsigned IdxSize = MRI.getType(IdxReg).getSizeInBits();
       unsigned IdxBank = getRegBankID(IdxReg, MRI, AMDGPU::SGPRRegBankID);
       OpdsMapping[0] = AMDGPU::getValueMapping(AMDGPU::VGPRRegBankID, DstSize);

llvm/lib/Target/AMDGPU/AMDGPUSearchableTables.td

 def : SourceOfDivergence<int_amdgcn_ds_swizzle>;
 def : SourceOfDivergence<int_amdgcn_ds_ordered_add>;
 def : SourceOfDivergence<int_amdgcn_ds_ordered_swap>;
 def : SourceOfDivergence<int_amdgcn_permlane16>;
 def : SourceOfDivergence<int_amdgcn_permlanex16>;
 def : SourceOfDivergence<int_amdgcn_mov_dpp>;
 def : SourceOfDivergence<int_amdgcn_mov_dpp8>;
 def : SourceOfDivergence<int_amdgcn_update_dpp>;
+def : SourceOfDivergence<int_amdgcn_writelane2>;
 def : SourceOfDivergence<int_amdgcn_writelane>;
 
 def : SourceOfDivergence<int_amdgcn_mfma_f32_4x4x1f32>;
 def : SourceOfDivergence<int_amdgcn_mfma_f32_4x4x1f32>;
 def : SourceOfDivergence<int_amdgcn_mfma_f32_4x4x4f16>;
 def : SourceOfDivergence<int_amdgcn_mfma_i32_4x4x4i8>;
 def : SourceOfDivergence<int_amdgcn_mfma_f32_4x4x2bf16>;
 def : SourceOfDivergence<int_amdgcn_mfma_f32_16x16x1f32>;

llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp

void AMDGPUPassConfig::addIRPasses() {
   // bitcast calls.
   addPass(createAMDGPUFixFunctionBitcastsPass());
 
   // A call to propagate attributes pass in the backend in case opt was not run.
   addPass(createAMDGPUPropagateAttributesEarlyPass(&TM));
 
   addPass(createAtomicExpandPass());
 
+  if (EnableAtomicOptimizations)
+    addPass(createAMDGPUAtomicOptimizerPass());
 
   addPass(createAMDGPULowerIntrinsicsPass());
 
   // Function calls are not supported, so make sure we inline everything.
   addPass(createAMDGPUAlwaysInlinePass());
   addPass(createAlwaysInlinerLegacyPass());
   // We need to add the barrier noop pass, otherwise adding the function
   // inlining pass will cause all of the PassConfigs passes to be run
ScheduleDAGInstrs *GCNPassConfig::createMachineScheduler(
   if (ST.enableSIScheduler())
     return createSIMachineScheduler(C);
   return createGCNMaxOccupancyMachineScheduler(C);
 }
 
 bool GCNPassConfig::addPreISel() {
   AMDGPUPassConfig::addPreISel();
 
-  if (EnableAtomicOptimizations) {
-    addPass(createAMDGPUAtomicOptimizerPass());
-  }
-
   // FIXME: We need to run a pass to propagate the attributes when calls are
   // supported.
 
   // Merge divergent exit nodes. StructurizeCFG won't recognize the multi-exit
   // regions formed by them.
   addPass(&AMDGPUUnifyDivergentExitNodesID);
   if (!LateCFGStructurize) {
     if (EnableStructurizerWorkarounds) {

llvm/lib/Target/AMDGPU/AMDGPUTargetTransformInfo.cpp

bool GCNTTIImpl::isSourceOfDivergence(const Value *V) const {
   return false;
 }
 
 bool GCNTTIImpl::isAlwaysUniform(const Value *V) const {
   if (const IntrinsicInst *Intrinsic = dyn_cast<IntrinsicInst>(V)) {
     switch (Intrinsic->getIntrinsicID()) {
     default:
       return false;
+    case Intrinsic::amdgcn_readfirstlane2:
+    case Intrinsic::amdgcn_readlane2:
     case Intrinsic::amdgcn_readfirstlane:
     case Intrinsic::amdgcn_readlane:
     case Intrinsic::amdgcn_icmp:
     case Intrinsic::amdgcn_fcmp:
     case Intrinsic::amdgcn_ballot:
     case Intrinsic::amdgcn_if_break:
       return true;
     }

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

#endif
 
   setOperationAction(ISD::INTRINSIC_WO_CHAIN, MVT::Other, Custom);
   setOperationAction(ISD::INTRINSIC_WO_CHAIN, MVT::f32, Custom);
   setOperationAction(ISD::INTRINSIC_WO_CHAIN, MVT::v4f32, Custom);
   setOperationAction(ISD::INTRINSIC_WO_CHAIN, MVT::i16, Custom);
   setOperationAction(ISD::INTRINSIC_WO_CHAIN, MVT::f16, Custom);
   setOperationAction(ISD::INTRINSIC_WO_CHAIN, MVT::v2i16, Custom);
   setOperationAction(ISD::INTRINSIC_WO_CHAIN, MVT::v2f16, Custom);
+  setOperationAction(ISD::INTRINSIC_WO_CHAIN, MVT::v3i16, Custom);
+  setOperationAction(ISD::INTRINSIC_WO_CHAIN, MVT::v3f16, Custom);
 
   setOperationAction(ISD::INTRINSIC_W_CHAIN, MVT::v2f16, Custom);
   setOperationAction(ISD::INTRINSIC_W_CHAIN, MVT::v2i16, Custom);
   setOperationAction(ISD::INTRINSIC_W_CHAIN, MVT::v4f16, Custom);
   setOperationAction(ISD::INTRINSIC_W_CHAIN, MVT::v4i16, Custom);
   setOperationAction(ISD::INTRINSIC_W_CHAIN, MVT::v8f16, Custom);
   setOperationAction(ISD::INTRINSIC_W_CHAIN, MVT::Other, Custom);
   setOperationAction(ISD::INTRINSIC_W_CHAIN, MVT::f16, Custom);

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_readfirstlane2 (i32 imm:$src))),
   (S_MOV_B32 SReg_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)
   >;

llvm/lib/Target/AMDGPU/VOP1Instructions.td

 } // End isMoveImm = 1
 
 // FIXME: Specify SchedRW for READFIRSTLANE_B32
 // TODO: Make profile for this, there is VOP3 encoding also
 def V_READFIRSTLANE_B32 :
   InstSI <(outs SReg_32:$vdst),
     (ins VRegOrLds_32:$src0),
     "v_readfirstlane_b32 $vdst, $src0",
-    [(set i32:$vdst, (int_amdgcn_readfirstlane (i32 VRegOrLds_32:$src0)))]>,
+    [(set i32:$vdst, (int_amdgcn_readfirstlane2 (i32 VRegOrLds_32:$src0)))]>,
   Enc32 {
 
   let isCodeGenOnly = 0;
   let UseNamedOperandTable = 1;
 
   let Size = 4;
   let mayLoad = 0;
   let mayStore = 0;

llvm/lib/Target/AMDGPU/VOP2Instructions.td

 defm V_SUBREV_U32 : VOP2Inst <"v_subrev_u32", VOP_I32_I32_I32_ARITH, null_frag, "v_sub_u32", 1>;
 }
 
 } // End isCommutable = 1
 
 // These are special and do not read the exec mask.
 let isConvergent = 1, Uses = []<Register> in {
 def V_READLANE_B32 : VOP2_Pseudo<"v_readlane_b32", VOP_READLANE,
-  [(set i32:$vdst, (int_amdgcn_readlane i32:$src0, i32:$src1))]>;
+  [(set i32:$vdst, (int_amdgcn_readlane2 i32:$src0, i32:$src1))]>;
 
 let Constraints = "$vdst = $vdst_in", DisableEncoding="$vdst_in" in {
 def V_WRITELANE_B32 : VOP2_Pseudo<"v_writelane_b32", VOP_WRITELANE,
-  [(set i32:$vdst, (int_amdgcn_writelane i32:$src0, i32:$src1, i32:$vdst_in))]>;
+  [(set i32:$vdst, (int_amdgcn_writelane2 i32:$src0, i32:$src1, i32:$vdst_in))]>;
 } // End $vdst = $vdst_in, DisableEncoding $vdst_in
 } // End isConvergent = 1
 
 defm V_BFM_B32 : VOP2Inst <"v_bfm_b32", VOP_NO_EXT<VOP_I32_I32_I32>>;
 defm V_BCNT_U32_B32 : VOP2Inst <"v_bcnt_u32_b32", VOP_NO_EXT<VOP_I32_I32_I32>, add_ctpop>;
 defm V_MBCNT_LO_U32_B32 : VOP2Inst <"v_mbcnt_lo_u32_b32", VOP_NO_EXT<VOP_I32_I32_I32>, int_amdgcn_mbcnt_lo>;
 defm V_MBCNT_HI_U32_B32 : VOP2Inst <"v_mbcnt_hi_u32_b32", VOP_NO_EXT<VOP_I32_I32_I32>, int_amdgcn_mbcnt_hi>;
 defm V_LDEXP_F32 : VOP2Inst <"v_ldexp_f32", VOP_NO_EXT<VOP_F32_F32_I32>, AMDGPUldexp>;

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.readfirstlane2), %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
+    %1:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), %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
+    %1:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), %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: [[INT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY]](s32)
     ; GCN: S_ENDPGM 0, implicit [[INT]](s32)
     %0:sgpr(s32) = COPY $sgpr0
-    %1:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), %0
+    %1:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), %0
     S_ENDPGM 0, implicit %1
 ...

llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-amdgpu_ps.ll

define amdgpu_ps float @vgpr_return(i32 %vgpr) {
   ret float %cast
 }
 
 define amdgpu_ps i32 @sgpr_return_i32(i32 %vgpr) {
   ; CHECK-LABEL: name: sgpr_return_i32
   ; CHECK: bb.1 (%ir-block.0):
   ; CHECK:   liveins: $vgpr0
   ; CHECK:   [[COPY:%[0-9]+]]:_(s32) = COPY $vgpr0
-  ; CHECK:   [[INT:%[0-9]+]]:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY]](s32)
+  ; CHECK:   [[INT:%[0-9]+]]:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY]](s32)
   ; CHECK:   $sgpr0 = COPY [[INT]](s32)
   ; CHECK:   SI_RETURN_TO_EPILOG implicit $sgpr0
   ret i32 %vgpr
 }
 
 define amdgpu_ps i64 @sgpr_return_i64(i64 %vgpr) {
   ; CHECK-LABEL: name: sgpr_return_i64
   ; CHECK: bb.1 (%ir-block.0):
   ; CHECK:   liveins: $vgpr0, $vgpr1
   ; CHECK:   [[COPY:%[0-9]+]]:_(s32) = COPY $vgpr0
   ; CHECK:   [[COPY1:%[0-9]+]]:_(s32) = COPY $vgpr1
   ; CHECK:   [[MV:%[0-9]+]]:_(s64) = G_MERGE_VALUES [[COPY]](s32), [[COPY1]](s32)
   ; CHECK:   [[UV:%[0-9]+]]:_(s32), [[UV1:%[0-9]+]]:_(s32) = G_UNMERGE_VALUES [[MV]](s64)
-  ; CHECK:   [[INT:%[0-9]+]]:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[UV]](s32)
+  ; CHECK:   [[INT:%[0-9]+]]:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[UV]](s32)
   ; CHECK:   $sgpr0 = COPY [[INT]](s32)
-  ; CHECK:   [[INT1:%[0-9]+]]:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[UV1]](s32)
+  ; CHECK:   [[INT1:%[0-9]+]]:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[UV1]](s32)
   ; CHECK:   $sgpr1 = COPY [[INT1]](s32)
   ; CHECK:   SI_RETURN_TO_EPILOG implicit $sgpr0, implicit $sgpr1
   ret i64 %vgpr
 }
 
 define amdgpu_ps <2 x i32> @sgpr_return_v2i32(<2 x i32> %vgpr) {
   ; CHECK-LABEL: name: sgpr_return_v2i32
   ; CHECK: bb.1 (%ir-block.0):
   ; CHECK:   liveins: $vgpr0, $vgpr1
   ; CHECK:   [[COPY:%[0-9]+]]:_(s32) = COPY $vgpr0
   ; CHECK:   [[COPY1:%[0-9]+]]:_(s32) = COPY $vgpr1
   ; CHECK:   [[BUILD_VECTOR:%[0-9]+]]:_(<2 x s32>) = G_BUILD_VECTOR [[COPY]](s32), [[COPY1]](s32)
   ; CHECK:   [[UV:%[0-9]+]]:_(s32), [[UV1:%[0-9]+]]:_(s32) = G_UNMERGE_VALUES [[BUILD_VECTOR]](<2 x s32>)
-  ; CHECK:   [[INT:%[0-9]+]]:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[UV]](s32)
+  ; CHECK:   [[INT:%[0-9]+]]:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[UV]](s32)
   ; CHECK:   $sgpr0 = COPY [[INT]](s32)
-  ; CHECK:   [[INT1:%[0-9]+]]:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[UV1]](s32)
+  ; CHECK:   [[INT1:%[0-9]+]]:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[UV1]](s32)
   ; CHECK:   $sgpr1 = COPY [[INT1]](s32)
   ; CHECK:   SI_RETURN_TO_EPILOG implicit $sgpr0, implicit $sgpr1
   ret <2 x i32> %vgpr
 }
 
 define amdgpu_ps { i32, i32 } @sgpr_struct_return_i32_i32(i32 %vgpr0, i32 %vgpr1) {
   ; CHECK-LABEL: name: sgpr_struct_return_i32_i32
   ; CHECK: bb.1 (%ir-block.0):
   ; CHECK:   liveins: $vgpr0, $vgpr1
   ; CHECK:   [[COPY:%[0-9]+]]:_(s32) = COPY $vgpr0
   ; CHECK:   [[COPY1:%[0-9]+]]:_(s32) = COPY $vgpr1
   ; CHECK:   [[DEF:%[0-9]+]]:_(s32) = G_IMPLICIT_DEF
-  ; CHECK:   [[INT:%[0-9]+]]:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY]](s32)
+  ; CHECK:   [[INT:%[0-9]+]]:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY]](s32)
   ; CHECK:   $sgpr0 = COPY [[INT]](s32)
-  ; CHECK:   [[INT1:%[0-9]+]]:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY1]](s32)
+  ; CHECK:   [[INT1:%[0-9]+]]:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY1]](s32)
   ; CHECK:   $sgpr1 = COPY [[INT1]](s32)
   ; CHECK:   SI_RETURN_TO_EPILOG implicit $sgpr0, implicit $sgpr1
   %insertvalue0 = insertvalue { i32, i32 } undef, i32 %vgpr0, 0
   %value = insertvalue { i32, i32 } %insertvalue0, i32 %vgpr1, 1
   ret { i32, i32 } %value
 }
 
 declare void @llvm.amdgcn.exp.f32(i32 immarg, i32 immarg, float, float, float, float, i1 immarg, i1 immarg)  #0
 
 attributes #0 = { nounwind }
 attributes #1 = { "InitialPSInputAddr"="0x00002" }

llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-amdgpu_vs.ll

 
 define amdgpu_vs <{ i32, i32 }> @ret_struct(i32 inreg %arg0, i32 inreg %arg1) {
   ; CHECK-LABEL: name: ret_struct
   ; CHECK: bb.1.main_body:
   ; CHECK:   liveins: $sgpr2, $sgpr3
   ; CHECK:   [[COPY:%[0-9]+]]:_(s32) = COPY $sgpr2
   ; CHECK:   [[COPY1:%[0-9]+]]:_(s32) = COPY $sgpr3
   ; CHECK:   [[DEF:%[0-9]+]]:_(s32) = G_IMPLICIT_DEF
-  ; CHECK:   [[INT:%[0-9]+]]:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY]](s32)
+  ; CHECK:   [[INT:%[0-9]+]]:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY]](s32)
   ; CHECK:   $sgpr0 = COPY [[INT]](s32)
-  ; CHECK:   [[INT1:%[0-9]+]]:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY1]](s32)
+  ; CHECK:   [[INT1:%[0-9]+]]:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY1]](s32)
   ; CHECK:   $sgpr1 = COPY [[INT1]](s32)
   ; CHECK:   SI_RETURN_TO_EPILOG implicit $sgpr0, implicit $sgpr1
 main_body:
   %tmp0 = insertvalue <{ i32, i32 }> undef, i32 %arg0, 0
   %tmp1 = insertvalue <{ i32, i32 }> %tmp0, i32 %arg1, 1
   ret <{ i32, i32 }> %tmp1
 }
 
 define amdgpu_vs i32 @non_void_ret() {
   ; CHECK-LABEL: name: non_void_ret
   ; CHECK: bb.1 (%ir-block.0):
   ; CHECK:   [[C:%[0-9]+]]:_(s32) = G_CONSTANT i32 0
-  ; CHECK:   [[INT:%[0-9]+]]:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[C]](s32)
+  ; CHECK:   [[INT:%[0-9]+]]:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[C]](s32)
   ; CHECK:   $sgpr0 = COPY [[INT]](s32)
   ; CHECK:   SI_RETURN_TO_EPILOG implicit $sgpr0
   ret i32 0
 }
 
 declare void @llvm.amdgcn.exp.f32(i32, i32, float, float, float, float, i1, i1)  #0
 
 attributes #0 = { nounwind }

llvm/test/CodeGen/AMDGPU/GlobalISel/regbankselect-amdgcn.readfirstlane.mir

 # NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py
 # RUN: llc -march=amdgcn -mcpu=fiji -run-pass=regbankselect -regbankselect-fast -verify-machineinstrs -o - %s | FileCheck %s
 # RUN: llc -march=amdgcn -mcpu=fiji -run-pass=regbankselect -regbankselect-greedy -verify-machineinstrs -o - %s | FileCheck %s
 
 ---
 name: readfirstlane_s
 legalized: true
 
 body: |
   bb.0:
     liveins: $sgpr0
     ; CHECK-LABEL: name: readfirstlane_s
     ; CHECK: [[COPY:%[0-9]+]]:sgpr(s32) = COPY $sgpr0
     ; CHECK: [[COPY1:%[0-9]+]]:vgpr(s32) = COPY [[COPY]](s32)
-    ; CHECK: [[INT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY1]](s32)
+    ; CHECK: [[INT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY1]](s32)
     %0:_(s32) = COPY $sgpr0
-    %1:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), %0
+    %1:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), %0
 ...
 
 ---
 name: readfirstlane_v
 legalized: true
 
 body: |
   bb.0:
     liveins: $vgpr0
     ; CHECK-LABEL: name: readfirstlane_v
     ; CHECK: [[COPY:%[0-9]+]]:vgpr(s32) = COPY $vgpr0
-    ; CHECK: [[INT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY]](s32)
+    ; CHECK: [[INT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY]](s32)
     %0:_(s32) = COPY $vgpr0
-    %1:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), %0
+    %1:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), %0
 ...

llvm/test/CodeGen/AMDGPU/GlobalISel/regbankselect-amdgcn.readlane.mir

 # NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py
 # RUN: llc -march=amdgcn -mcpu=fiji -run-pass=regbankselect %s -verify-machineinstrs -o - -regbankselect-fast | FileCheck %s
 # RUN: llc -march=amdgcn -mcpu=fiji -run-pass=regbankselect %s -verify-machineinstrs -o - -regbankselect-greedy | FileCheck %s
 
 ---
 name: readlane_ss
 legalized: true
 
 body: |
   bb.0:
     liveins: $sgpr0, $sgpr1
     ; CHECK-LABEL: name: readlane_ss
     ; CHECK: [[COPY:%[0-9]+]]:sgpr(s32) = COPY $sgpr0
     ; CHECK: [[COPY1:%[0-9]+]]:sgpr(s32) = COPY $sgpr1
     ; CHECK: [[COPY2:%[0-9]+]]:vgpr(s32) = COPY [[COPY]](s32)
-    ; CHECK: [[INT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readlane), [[COPY2]](s32), [[COPY1]](s32)
+    ; CHECK: [[INT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readlane2), [[COPY2]](s32), [[COPY1]](s32)
     %0:_(s32) = COPY $sgpr0
     %1:_(s32) = COPY $sgpr1
-    %2:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readlane), %0, %1
+    %2:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readlane2), %0, %1
 ...
 
 ---
 name: readlane_vs
 legalized: true
 
 body: |
   bb.0:
     liveins: $vgpr0, $sgpr0
     ; CHECK-LABEL: name: readlane_vs
     ; CHECK: [[COPY:%[0-9]+]]:vgpr(s32) = COPY $vgpr0
     ; CHECK: [[COPY1:%[0-9]+]]:sgpr(s32) = COPY $sgpr0
-    ; CHECK: [[INT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readlane), [[COPY]](s32), [[COPY1]](s32)
+    ; CHECK: [[INT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readlane2), [[COPY]](s32), [[COPY1]](s32)
     %0:_(s32) = COPY $vgpr0
     %1:_(s32) = COPY $sgpr0
-    %2:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readlane), %0, %1
+    %2:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readlane2), %0, %1
 ...
 
 ---
 name: readlane_vv
 legalized: true
 
 body: |
   bb.0:
     liveins: $vgpr0, $vgpr1
     ; CHECK-LABEL: name: readlane_vv
     ; CHECK: [[COPY:%[0-9]+]]:vgpr(s32) = COPY $vgpr0
     ; CHECK: [[COPY1:%[0-9]+]]:vgpr_32(s32) = COPY $vgpr1
     ; CHECK: [[V_READFIRSTLANE_B32_:%[0-9]+]]:sreg_32(s32) = V_READFIRSTLANE_B32 [[COPY1]](s32), implicit $exec
-    ; CHECK: [[INT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readlane), [[COPY]](s32), [[V_READFIRSTLANE_B32_]](s32)
+    ; CHECK: [[INT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readlane2), [[COPY]](s32), [[V_READFIRSTLANE_B32_]](s32)
     %0:_(s32) = COPY $vgpr0
     %1:_(s32) = COPY $vgpr1
-    %2:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readlane), %0, %1
+    %2:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readlane2), %0, %1
 ...
 
 ---
 name: readlane_sv
 legalized: true
 
 body: |
   bb.0:
     liveins: $vgpr0, $sgpr0
     ; CHECK-LABEL: name: readlane_sv
     ; CHECK: [[COPY:%[0-9]+]]:sgpr(s32) = COPY $sgpr0
     ; CHECK: [[COPY1:%[0-9]+]]:vgpr_32(s32) = COPY $vgpr0
     ; CHECK: [[COPY2:%[0-9]+]]:vgpr(s32) = COPY [[COPY]](s32)
     ; CHECK: [[V_READFIRSTLANE_B32_:%[0-9]+]]:sreg_32(s32) = V_READFIRSTLANE_B32 [[COPY1]](s32), implicit $exec
-    ; CHECK: [[INT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readlane), [[COPY2]](s32), [[V_READFIRSTLANE_B32_]](s32)
+    ; CHECK: [[INT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readlane2), [[COPY2]](s32), [[V_READFIRSTLANE_B32_]](s32)
     %0:_(s32) = COPY $sgpr0
     %1:_(s32) = COPY $vgpr0
-    %2:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readlane), %0, %1
+    %2:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readlane2), %0, %1
 ...

llvm/test/CodeGen/AMDGPU/GlobalISel/regbankselect-amdgcn.s.buffer.load.ll

define amdgpu_ps i32 @s_buffer_load_i32(<4 x i32> inreg %rsrc, i32 inreg %soffset) {
   ; CHECK:   [[COPY:%[0-9]+]]:sgpr(s32) = COPY $sgpr2
   ; CHECK:   [[COPY1:%[0-9]+]]:sgpr(s32) = COPY $sgpr3
   ; CHECK:   [[COPY2:%[0-9]+]]:sgpr(s32) = COPY $sgpr4
   ; CHECK:   [[COPY3:%[0-9]+]]:sgpr(s32) = COPY $sgpr5
   ; CHECK:   [[COPY4:%[0-9]+]]:sgpr(s32) = COPY $sgpr6
   ; CHECK:   [[BUILD_VECTOR:%[0-9]+]]:sgpr(<4 x s32>) = G_BUILD_VECTOR [[COPY]](s32), [[COPY1]](s32), [[COPY2]](s32), [[COPY3]](s32)
   ; CHECK:   [[AMDGPU_S_BUFFER_LOAD:%[0-9]+]]:sgpr(s32) = G_AMDGPU_S_BUFFER_LOAD [[BUILD_VECTOR]](<4 x s32>), [[COPY4]](s32), 0 :: (dereferenceable invariant load 4)
   ; CHECK:   [[COPY5:%[0-9]+]]:vgpr(s32) = COPY [[AMDGPU_S_BUFFER_LOAD]](s32)
-  ; CHECK:   [[INT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY5]](s32)
+  ; CHECK:   [[INT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY5]](s32)
   ; CHECK:   $sgpr0 = COPY [[INT]](s32)
   ; CHECK:   SI_RETURN_TO_EPILOG implicit $sgpr0
   %val = call i32 @llvm.amdgcn.s.buffer.load.i32(<4 x i32> %rsrc, i32 %soffset, i32 0)
   ret i32 %val
 }
 
 define amdgpu_ps <2 x i32> @s_buffer_load_v2i32(<4 x i32> inreg %rsrc, i32 inreg %soffset) {
   ; CHECK-LABEL: name: s_buffer_load_v2i32
   ; CHECK: bb.1 (%ir-block.0):
   ; CHECK:   liveins: $sgpr2, $sgpr3, $sgpr4, $sgpr5, $sgpr6
   ; CHECK:   [[COPY:%[0-9]+]]:sgpr(s32) = COPY $sgpr2
   ; CHECK:   [[COPY1:%[0-9]+]]:sgpr(s32) = COPY $sgpr3
   ; CHECK:   [[COPY2:%[0-9]+]]:sgpr(s32) = COPY $sgpr4
   ; CHECK:   [[COPY3:%[0-9]+]]:sgpr(s32) = COPY $sgpr5
   ; CHECK:   [[COPY4:%[0-9]+]]:sgpr(s32) = COPY $sgpr6
   ; CHECK:   [[BUILD_VECTOR:%[0-9]+]]:sgpr(<4 x s32>) = G_BUILD_VECTOR [[COPY]](s32), [[COPY1]](s32), [[COPY2]](s32), [[COPY3]](s32)
   ; CHECK:   [[AMDGPU_S_BUFFER_LOAD:%[0-9]+]]:sgpr(<2 x s32>) = G_AMDGPU_S_BUFFER_LOAD [[BUILD_VECTOR]](<4 x s32>), [[COPY4]](s32), 0 :: (dereferenceable invariant load 8, align 4)
   ; CHECK:   [[UV:%[0-9]+]]:sgpr(s32), [[UV1:%[0-9]+]]:sgpr(s32) = G_UNMERGE_VALUES [[AMDGPU_S_BUFFER_LOAD]](<2 x s32>)
   ; CHECK:   [[COPY5:%[0-9]+]]:vgpr(s32) = COPY [[UV]](s32)
-  ; CHECK:   [[INT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY5]](s32)
+  ; CHECK:   [[INT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY5]](s32)
   ; CHECK:   $sgpr0 = COPY [[INT]](s32)
   ; CHECK:   [[COPY6:%[0-9]+]]:vgpr(s32) = COPY [[UV1]](s32)
-  ; CHECK:   [[INT1:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY6]](s32)
+  ; CHECK:   [[INT1:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY6]](s32)
   ; CHECK:   $sgpr1 = COPY [[INT1]](s32)
   ; CHECK:   SI_RETURN_TO_EPILOG implicit $sgpr0, implicit $sgpr1
   %val = call <2 x i32> @llvm.amdgcn.s.buffer.load.v2i32(<4 x i32> %rsrc, i32 %soffset, i32 0)
   ret <2 x i32> %val
 }
 
 define amdgpu_ps <3 x i32> @s_buffer_load_v3i32(<4 x i32> inreg %rsrc, i32 inreg %soffset) {
   ; CHECK-LABEL: name: s_buffer_load_v3i32
   ; CHECK: bb.1 (%ir-block.0):
   ; CHECK:   liveins: $sgpr2, $sgpr3, $sgpr4, $sgpr5, $sgpr6
   ; CHECK:   [[COPY:%[0-9]+]]:sgpr(s32) = COPY $sgpr2
   ; CHECK:   [[COPY1:%[0-9]+]]:sgpr(s32) = COPY $sgpr3
   ; CHECK:   [[COPY2:%[0-9]+]]:sgpr(s32) = COPY $sgpr4
   ; CHECK:   [[COPY3:%[0-9]+]]:sgpr(s32) = COPY $sgpr5
   ; CHECK:   [[COPY4:%[0-9]+]]:sgpr(s32) = COPY $sgpr6
   ; CHECK:   [[BUILD_VECTOR:%[0-9]+]]:sgpr(<4 x s32>) = G_BUILD_VECTOR [[COPY]](s32), [[COPY1]](s32), [[COPY2]](s32), [[COPY3]](s32)
   ; CHECK:   [[AMDGPU_S_BUFFER_LOAD:%[0-9]+]]:sgpr(<4 x s32>) = G_AMDGPU_S_BUFFER_LOAD [[BUILD_VECTOR]](<4 x s32>), [[COPY4]](s32), 0 :: (dereferenceable invariant load 12, align 4)
   ; CHECK:   [[DEF:%[0-9]+]]:sgpr(<4 x s32>) = G_IMPLICIT_DEF
   ; CHECK:   [[CONCAT_VECTORS:%[0-9]+]]:sgpr(<12 x s32>) = G_CONCAT_VECTORS [[AMDGPU_S_BUFFER_LOAD]](<4 x s32>), [[DEF]](<4 x s32>), [[DEF]](<4 x s32>)
   ; CHECK:   [[UV:%[0-9]+]]:sgpr(<3 x s32>), [[UV1:%[0-9]+]]:sgpr(<3 x s32>), [[UV2:%[0-9]+]]:sgpr(<3 x s32>), [[UV3:%[0-9]+]]:sgpr(<3 x s32>) = G_UNMERGE_VALUES [[CONCAT_VECTORS]](<12 x s32>)
   ; CHECK:   [[UV4:%[0-9]+]]:sgpr(s32), [[UV5:%[0-9]+]]:sgpr(s32), [[UV6:%[0-9]+]]:sgpr(s32) = G_UNMERGE_VALUES [[UV]](<3 x s32>)
   ; CHECK:   [[COPY5:%[0-9]+]]:vgpr(s32) = COPY [[UV4]](s32)
-  ; CHECK:   [[INT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY5]](s32)
+  ; CHECK:   [[INT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY5]](s32)
   ; CHECK:   $sgpr0 = COPY [[INT]](s32)
   ; CHECK:   [[COPY6:%[0-9]+]]:vgpr(s32) = COPY [[UV5]](s32)
-  ; CHECK:   [[INT1:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY6]](s32)
+  ; CHECK:   [[INT1:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY6]](s32)
   ; CHECK:   $sgpr1 = COPY [[INT1]](s32)
   ; CHECK:   [[COPY7:%[0-9]+]]:vgpr(s32) = COPY [[UV6]](s32)
-  ; CHECK:   [[INT2:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY7]](s32)
+  ; CHECK:   [[INT2:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY7]](s32)
   ; CHECK:   $sgpr2 = COPY [[INT2]](s32)
   ; CHECK:   SI_RETURN_TO_EPILOG implicit $sgpr0, implicit $sgpr1, implicit $sgpr2
   %val = call <3 x i32> @llvm.amdgcn.s.buffer.load.v3i32(<4 x i32> %rsrc, i32 %soffset, i32 0)
   ret <3 x i32> %val
 }
 
 define amdgpu_ps <8 x i32> @s_buffer_load_v8i32(<4 x i32> inreg %rsrc, i32 inreg %soffset) {
   ; CHECK-LABEL: name: s_buffer_load_v8i32
   ; CHECK: bb.1 (%ir-block.0):
   ; CHECK:   liveins: $sgpr2, $sgpr3, $sgpr4, $sgpr5, $sgpr6
   ; CHECK:   [[COPY:%[0-9]+]]:sgpr(s32) = COPY $sgpr2
   ; CHECK:   [[COPY1:%[0-9]+]]:sgpr(s32) = COPY $sgpr3
   ; CHECK:   [[COPY2:%[0-9]+]]:sgpr(s32) = COPY $sgpr4
   ; CHECK:   [[COPY3:%[0-9]+]]:sgpr(s32) = COPY $sgpr5
   ; CHECK:   [[COPY4:%[0-9]+]]:sgpr(s32) = COPY $sgpr6
   ; CHECK:   [[BUILD_VECTOR:%[0-9]+]]:sgpr(<4 x s32>) = G_BUILD_VECTOR [[COPY]](s32), [[COPY1]](s32), [[COPY2]](s32), [[COPY3]](s32)
   ; CHECK:   [[AMDGPU_S_BUFFER_LOAD:%[0-9]+]]:sgpr(<8 x s32>) = G_AMDGPU_S_BUFFER_LOAD [[BUILD_VECTOR]](<4 x s32>), [[COPY4]](s32), 0 :: (dereferenceable invariant load 32, align 4)
   ; CHECK:   [[UV:%[0-9]+]]:sgpr(s32), [[UV1:%[0-9]+]]:sgpr(s32), [[UV2:%[0-9]+]]:sgpr(s32), [[UV3:%[0-9]+]]:sgpr(s32), [[UV4:%[0-9]+]]:sgpr(s32), [[UV5:%[0-9]+]]:sgpr(s32), [[UV6:%[0-9]+]]:sgpr(s32), [[UV7:%[0-9]+]]:sgpr(s32) = G_UNMERGE_VALUES [[AMDGPU_S_BUFFER_LOAD]](<8 x s32>)
   ; CHECK:   [[COPY5:%[0-9]+]]:vgpr(s32) = COPY [[UV]](s32)
-  ; CHECK:   [[INT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY5]](s32)
+  ; CHECK:   [[INT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY5]](s32)
   ; CHECK:   $sgpr0 = COPY [[INT]](s32)
   ; CHECK:   [[COPY6:%[0-9]+]]:vgpr(s32) = COPY [[UV1]](s32)
-  ; CHECK:   [[INT1:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY6]](s32)
+  ; CHECK:   [[INT1:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY6]](s32)
   ; CHECK:   $sgpr1 = COPY [[INT1]](s32)
   ; CHECK:   [[COPY7:%[0-9]+]]:vgpr(s32) = COPY [[UV2]](s32)
-  ; CHECK:   [[INT2:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY7]](s32)
+  ; CHECK:   [[INT2:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY7]](s32)
   ; CHECK:   $sgpr2 = COPY [[INT2]](s32)
   ; CHECK:   [[COPY8:%[0-9]+]]:vgpr(s32) = COPY [[UV3]](s32)
-  ; CHECK:   [[INT3:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY8]](s32)
+  ; CHECK:   [[INT3:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY8]](s32)
   ; CHECK:   $sgpr3 = COPY [[INT3]](s32)
   ; CHECK:   [[COPY9:%[0-9]+]]:vgpr(s32) = COPY [[UV4]](s32)
-  ; CHECK:   [[INT4:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY9]](s32)
+  ; CHECK:   [[INT4:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY9]](s32)
   ; CHECK:   $sgpr4 = COPY [[INT4]](s32)
   ; CHECK:   [[COPY10:%[0-9]+]]:vgpr(s32) = COPY [[UV5]](s32)
-  ; CHECK:   [[INT5:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY10]](s32)
+  ; CHECK:   [[INT5:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY10]](s32)
   ; CHECK:   $sgpr5 = COPY [[INT5]](s32)
   ; CHECK:   [[COPY11:%[0-9]+]]:vgpr(s32) = COPY [[UV6]](s32)
-  ; CHECK:   [[INT6:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY11]](s32)
+  ; CHECK:   [[INT6:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY11]](s32)
   ; CHECK:   $sgpr6 = COPY [[INT6]](s32)
   ; CHECK:   [[COPY12:%[0-9]+]]:vgpr(s32) = COPY [[UV7]](s32)
-  ; CHECK:   [[INT7:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY12]](s32)
+  ; CHECK:   [[INT7:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY12]](s32)
   ; CHECK:   $sgpr7 = COPY [[INT7]](s32)
   ; CHECK:   SI_RETURN_TO_EPILOG implicit $sgpr0, implicit $sgpr1, implicit $sgpr2, implicit $sgpr3, implicit $sgpr4, implicit $sgpr5, implicit $sgpr6, implicit $sgpr7
   %val = call <8 x i32> @llvm.amdgcn.s.buffer.load.v8i32(<4 x i32> %rsrc, i32 %soffset, i32 0)
   ret <8 x i32> %val
 }
 
 define amdgpu_ps <16 x i32> @s_buffer_load_v16i32(<4 x i32> inreg %rsrc, i32 inreg %soffset) {
   ; CHECK-LABEL: name: s_buffer_load_v16i32
   ; CHECK: bb.1 (%ir-block.0):
   ; CHECK:   liveins: $sgpr2, $sgpr3, $sgpr4, $sgpr5, $sgpr6
   ; CHECK:   [[COPY:%[0-9]+]]:sgpr(s32) = COPY $sgpr2
   ; CHECK:   [[COPY1:%[0-9]+]]:sgpr(s32) = COPY $sgpr3
   ; CHECK:   [[COPY2:%[0-9]+]]:sgpr(s32) = COPY $sgpr4
   ; CHECK:   [[COPY3:%[0-9]+]]:sgpr(s32) = COPY $sgpr5
   ; CHECK:   [[COPY4:%[0-9]+]]:sgpr(s32) = COPY $sgpr6
   ; CHECK:   [[BUILD_VECTOR:%[0-9]+]]:sgpr(<4 x s32>) = G_BUILD_VECTOR [[COPY]](s32), [[COPY1]](s32), [[COPY2]](s32), [[COPY3]](s32)
   ; CHECK:   [[AMDGPU_S_BUFFER_LOAD:%[0-9]+]]:sgpr(<16 x s32>) = G_AMDGPU_S_BUFFER_LOAD [[BUILD_VECTOR]](<4 x s32>), [[COPY4]](s32), 0 :: (dereferenceable invariant load 64, align 4)
   ; CHECK:   [[UV:%[0-9]+]]:sgpr(s32), [[UV1:%[0-9]+]]:sgpr(s32), [[UV2:%[0-9]+]]:sgpr(s32), [[UV3:%[0-9]+]]:sgpr(s32), [[UV4:%[0-9]+]]:sgpr(s32), [[UV5:%[0-9]+]]:sgpr(s32), [[UV6:%[0-9]+]]:sgpr(s32), [[UV7:%[0-9]+]]:sgpr(s32), [[UV8:%[0-9]+]]:sgpr(s32), [[UV9:%[0-9]+]]:sgpr(s32), [[UV10:%[0-9]+]]:sgpr(s32), [[UV11:%[0-9]+]]:sgpr(s32), [[UV12:%[0-9]+]]:sgpr(s32), [[UV13:%[0-9]+]]:sgpr(s32), [[UV14:%[0-9]+]]:sgpr(s32), [[UV15:%[0-9]+]]:sgpr(s32) = G_UNMERGE_VALUES [[AMDGPU_S_BUFFER_LOAD]](<16 x s32>)
   ; CHECK:   [[COPY5:%[0-9]+]]:vgpr(s32) = COPY [[UV]](s32)
-  ; CHECK:   [[INT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY5]](s32)
+  ; CHECK:   [[INT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY5]](s32)
   ; CHECK:   $sgpr0 = COPY [[INT]](s32)
   ; CHECK:   [[COPY6:%[0-9]+]]:vgpr(s32) = COPY [[UV1]](s32)
-  ; CHECK:   [[INT1:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY6]](s32)
+  ; CHECK:   [[INT1:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY6]](s32)
   ; CHECK:   $sgpr1 = COPY [[INT1]](s32)
   ; CHECK:   [[COPY7:%[0-9]+]]:vgpr(s32) = COPY [[UV2]](s32)
-  ; CHECK:   [[INT2:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY7]](s32)
+  ; CHECK:   [[INT2:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY7]](s32)
   ; CHECK:   $sgpr2 = COPY [[INT2]](s32)
   ; CHECK:   [[COPY8:%[0-9]+]]:vgpr(s32) = COPY [[UV3]](s32)
-  ; CHECK:   [[INT3:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY8]](s32)
+  ; CHECK:   [[INT3:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY8]](s32)
   ; CHECK:   $sgpr3 = COPY [[INT3]](s32)
   ; CHECK:   [[COPY9:%[0-9]+]]:vgpr(s32) = COPY [[UV4]](s32)
-  ; CHECK:   [[INT4:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY9]](s32)
+  ; CHECK:   [[INT4:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY9]](s32)
   ; CHECK:   $sgpr4 = COPY [[INT4]](s32)
   ; CHECK:   [[COPY10:%[0-9]+]]:vgpr(s32) = COPY [[UV5]](s32)
-  ; CHECK:   [[INT5:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY10]](s32)
+  ; CHECK:   [[INT5:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY10]](s32)
   ; CHECK:   $sgpr5 = COPY [[INT5]](s32)
   ; CHECK:   [[COPY11:%[0-9]+]]:vgpr(s32) = COPY [[UV6]](s32)
-  ; CHECK:   [[INT6:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY11]](s32)
+  ; CHECK:   [[INT6:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY11]](s32)
   ; CHECK:   $sgpr6 = COPY [[INT6]](s32)
   ; CHECK:   [[COPY12:%[0-9]+]]:vgpr(s32) = COPY [[UV7]](s32)
-  ; CHECK:   [[INT7:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY12]](s32)
+  ; CHECK:   [[INT7:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY12]](s32)
   ; CHECK:   $sgpr7 = COPY [[INT7]](s32)
   ; CHECK:   [[COPY13:%[0-9]+]]:vgpr(s32) = COPY [[UV8]](s32)
-  ; CHECK:   [[INT8:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY13]](s32)
+  ; CHECK:   [[INT8:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY13]](s32)
   ; CHECK:   $sgpr8 = COPY [[INT8]](s32)
   ; CHECK:   [[COPY14:%[0-9]+]]:vgpr(s32) = COPY [[UV9]](s32)
-  ; CHECK:   [[INT9:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY14]](s32)
+  ; CHECK:   [[INT9:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY14]](s32)
   ; CHECK:   $sgpr9 = COPY [[INT9]](s32)
   ; CHECK:   [[COPY15:%[0-9]+]]:vgpr(s32) = COPY [[UV10]](s32)
-  ; CHECK:   [[INT10:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY15]](s32)
+  ; CHECK:   [[INT10:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY15]](s32)
   ; CHECK:   $sgpr10 = COPY [[INT10]](s32)
   ; CHECK:   [[COPY16:%[0-9]+]]:vgpr(s32) = COPY [[UV11]](s32)
-  ; CHECK:   [[INT11:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY16]](s32)
+  ; CHECK:   [[INT11:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY16]](s32)
   ; CHECK:   $sgpr11 = COPY [[INT11]](s32)
   ; CHECK:   [[COPY17:%[0-9]+]]:vgpr(s32) = COPY [[UV12]](s32)
-  ; CHECK:   [[INT12:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY17]](s32)
+  ; CHECK:   [[INT12:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY17]](s32)
   ; CHECK:   $sgpr12 = COPY [[INT12]](s32)
   ; CHECK:   [[COPY18:%[0-9]+]]:vgpr(s32) = COPY [[UV13]](s32)
-  ; CHECK:   [[INT13:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY18]](s32)
+  ; CHECK:   [[INT13:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY18]](s32)
   ; CHECK:   $sgpr13 = COPY [[INT13]](s32)
   ; CHECK:   [[COPY19:%[0-9]+]]:vgpr(s32) = COPY [[UV14]](s32)
-  ; CHECK:   [[INT14:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY19]](s32)
+  ; CHECK:   [[INT14:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY19]](s32)
   ; CHECK:   $sgpr14 = COPY [[INT14]](s32)
   ; CHECK:   [[COPY20:%[0-9]+]]:vgpr(s32) = COPY [[UV15]](s32)
-  ; CHECK:   [[INT15:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane), [[COPY20]](s32)
+  ; CHECK:   [[INT15:%[0-9]+]]:sgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.readfirstlane2), [[COPY20]](s32)
   ; CHECK:   $sgpr15 = COPY [[INT15]](s32)
   ; CHECK:   SI_RETURN_TO_EPILOG implicit $sgpr0, implicit $sgpr1, implicit $sgpr2, implicit $sgpr3, implicit $sgpr4, implicit $sgpr5, implicit $sgpr6, implicit $sgpr7, implicit $sgpr8, implicit $sgpr9, implicit $sgpr10, implicit $sgpr11, implicit $sgpr12, implicit $sgpr13, implicit $sgpr14, implicit $sgpr15
   %val = call <16 x i32> @llvm.amdgcn.s.buffer.load.v16i32(<4 x i32> %rsrc, i32 %soffset, i32 0)
   ret <16 x i32> %val
 }
 
 ; Check cases that need to be converted to MUBUF due to the offset being a VGPR.
 define amdgpu_ps float @s_buffer_load_f32_vgpr_offset(<4 x i32> inreg %rsrc, i32 %soffset) {

llvm/test/CodeGen/AMDGPU/GlobalISel/regbankselect-amdgcn.writelane.mir

 # NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py
 # RUN: llc -march=amdgcn -mcpu=fiji -run-pass=regbankselect %s -verify-machineinstrs -o - -regbankselect-fast | FileCheck %s
 # RUN: llc -march=amdgcn -mcpu=fiji -run-pass=regbankselect %s -verify-machineinstrs -o - -regbankselect-greedy | FileCheck %s
 
 ---
 name: writelane_sss
 legalized: true
 
 body: |
   bb.0:
     liveins: $sgpr0, $sgpr1, $sgpr2
     ; CHECK-LABEL: name: writelane_sss
     ; CHECK: [[COPY:%[0-9]+]]:sgpr(s32) = COPY $sgpr0
     ; CHECK: [[COPY1:%[0-9]+]]:sgpr(s32) = COPY $sgpr1
     ; CHECK: [[COPY2:%[0-9]+]]:sgpr(s32) = COPY $sgpr2
     ; CHECK: [[COPY3:%[0-9]+]]:vgpr(s32) = COPY [[COPY2]](s32)
-    ; CHECK: [[INT:%[0-9]+]]:vgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.writelane), [[COPY]](s32), [[COPY1]](s32), [[COPY3]](s32)
+    ; CHECK: [[INT:%[0-9]+]]:vgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.writelane2), [[COPY]](s32), [[COPY1]](s32), [[COPY3]](s32)
     %0:_(s32) = COPY $sgpr0
     %1:_(s32) = COPY $sgpr1
     %2:_(s32) = COPY $sgpr2
-    %3:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.writelane), %0, %1, %2
+    %3:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.writelane2), %0, %1, %2
 ...
 
 ---
 name: writelane_ssv
 legalized: true
 
 body: |
   bb.0:
     liveins: $sgpr0, $sgpr1, $vgpr0
     ; CHECK-LABEL: name: writelane_ssv
     ; CHECK: [[COPY:%[0-9]+]]:sgpr(s32) = COPY $sgpr0
     ; CHECK: [[COPY1:%[0-9]+]]:sgpr(s32) = COPY $sgpr1
     ; CHECK: [[COPY2:%[0-9]+]]:vgpr(s32) = COPY $vgpr0
-    ; CHECK: [[INT:%[0-9]+]]:vgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.writelane), [[COPY]](s32), [[COPY1]](s32), [[COPY2]](s32)
+    ; CHECK: [[INT:%[0-9]+]]:vgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.writelane2), [[COPY]](s32), [[COPY1]](s32), [[COPY2]](s32)
     %0:_(s32) = COPY $sgpr0
     %1:_(s32) = COPY $sgpr1
     %2:_(s32) = COPY $vgpr0
-    %3:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.writelane), %0, %1, %2
+    %3:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.writelane2), %0, %1, %2
 ...
 
 ---
 name: writelane_vsv
 legalized: true
 
 body: |
   bb.0:
     liveins: $sgpr0, $vgpr0, $vgpr1
     ; CHECK-LABEL: name: writelane_vsv
     ; CHECK: [[COPY:%[0-9]+]]:vgpr_32(s32) = COPY $vgpr0
     ; CHECK: [[COPY1:%[0-9]+]]:sgpr(s32) = COPY $sgpr0
     ; CHECK: [[COPY2:%[0-9]+]]:vgpr(s32) = COPY $vgpr1
     ; CHECK: [[V_READFIRSTLANE_B32_:%[0-9]+]]:sreg_32(s32) = V_READFIRSTLANE_B32 [[COPY]](s32), implicit $exec
-    ; CHECK: [[INT:%[0-9]+]]:vgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.writelane), [[V_READFIRSTLANE_B32_]](s32), [[COPY1]](s32), [[COPY2]](s32)
+    ; CHECK: [[INT:%[0-9]+]]:vgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.writelane2), [[V_READFIRSTLANE_B32_]](s32), [[COPY1]](s32), [[COPY2]](s32)
     %0:_(s32) = COPY $vgpr0
     %1:_(s32) = COPY $sgpr0
     %2:_(s32) = COPY $vgpr1
-    %3:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.writelane), %0, %1, %2
+    %3:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.writelane2), %0, %1, %2
 ...
 
 ---
 name: writelane_vvv
 legalized: true
 
 body: |
   bb.0:
     liveins: $vgpr0, $vgpr1
     ; CHECK-LABEL: name: writelane_vvv
     ; CHECK: [[COPY:%[0-9]+]]:vgpr_32(s32) = COPY $vgpr0
     ; CHECK: [[COPY1:%[0-9]+]]:vgpr_32(s32) = COPY $vgpr1
     ; CHECK: [[COPY2:%[0-9]+]]:vgpr(s32) = COPY $vgpr2
     ; CHECK: [[V_READFIRSTLANE_B32_:%[0-9]+]]:sreg_32(s32) = V_READFIRSTLANE_B32 [[COPY]](s32), implicit $exec
     ; CHECK: [[V_READFIRSTLANE_B32_1:%[0-9]+]]:sreg_32(s32) = V_READFIRSTLANE_B32 [[COPY1]](s32), implicit $exec
-    ; CHECK: [[INT:%[0-9]+]]:vgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.writelane), [[V_READFIRSTLANE_B32_]](s32), [[V_READFIRSTLANE_B32_1]](s32), [[COPY2]](s32)
+    ; CHECK: [[INT:%[0-9]+]]:vgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.writelane2), [[V_READFIRSTLANE_B32_]](s32), [[V_READFIRSTLANE_B32_1]](s32), [[COPY2]](s32)
     %0:_(s32) = COPY $vgpr0
     %1:_(s32) = COPY $vgpr1
     %2:_(s32) = COPY $vgpr2
-    %3:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.writelane), %0, %1, %2
+    %3:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.writelane2), %0, %1, %2
 ...
 
 ---
 name: writelane_svv
 legalized: true
 
 body: |
   bb.0:
     liveins: $sgpr0, $vgpr0, $vgpr1
     ; CHECK-LABEL: name: writelane_svv
     ; CHECK: [[COPY:%[0-9]+]]:sgpr(s32) = COPY $sgpr0
     ; CHECK: [[COPY1:%[0-9]+]]:vgpr_32(s32) = COPY $vgpr0
     ; CHECK: [[COPY2:%[0-9]+]]:vgpr(s32) = COPY $vgpr1
     ; CHECK: [[V_READFIRSTLANE_B32_:%[0-9]+]]:sreg_32(s32) = V_READFIRSTLANE_B32 [[COPY1]](s32), implicit $exec
-    ; CHECK: [[INT:%[0-9]+]]:vgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.writelane), [[COPY]](s32), [[V_READFIRSTLANE_B32_]](s32), [[COPY2]](s32)
+    ; CHECK: [[INT:%[0-9]+]]:vgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.writelane2), [[COPY]](s32), [[V_READFIRSTLANE_B32_]](s32), [[COPY2]](s32)
     %0:_(s32) = COPY $sgpr0
     %1:_(s32) = COPY $vgpr0
     %2:_(s32) = COPY $vgpr1
-    %3:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.writelane), %0, %1, %2
+    %3:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.writelane2), %0, %1, %2
 ...

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.readfirstlane2.i32(i32) #0
+declare float @llvm.amdgcn.readfirstlane2.f32(float) #0
+declare <2 x half> @llvm.amdgcn.readfirstlane2.v2f16(<2 x half>) #0
+declare <2 x i16> @llvm.amdgcn.readfirstlane2.v2i16(<2 x i16>) #0
+declare i8 addrspace(3)* @llvm.amdgcn.readfirstlane2.p3i8(i8 addrspace(3)*) #0
+declare i16 @llvm.amdgcn.readfirstlane2.i16(i16) #0
+declare half @llvm.amdgcn.readfirstlane2.f16(half) #0
+declare <3 x i16> @llvm.amdgcn.readfirstlane2.v3i16(<3 x i16>) #0
+declare <9 x float> @llvm.amdgcn.readfirstlane2.v9f32(<9 x float>) #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.readfirstlane2.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.readfirstlane2.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.readfirstlane2.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.readfirstlane2.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.readfirstlane2.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.readfirstlane2.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
+; CHECK-NOT: v_readfirstlane_b32
+define void @test_readfirstlane_f32(float addrspace(1)* %out, float %src) #1 {
+  %readfirstlane = call float @llvm.amdgcn.readfirstlane2.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
+; CHECK-NOT: v_readfirstlane_b32
+define void @test_readfirstlane_v2f16(<2 x half> addrspace(1)* %out, <2 x half> %src) #1 {
+  %readfirstlane = call <2 x half> @llvm.amdgcn.readfirstlane2.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
+; CHECK-NOT: v_readfirstlane_b32
+define void @test_readfirstlane_v2i16(<2 x i16> addrspace(1)* %out, <2 x i16> %src) #1 {
+  %readfirstlane = call <2 x i16> @llvm.amdgcn.readfirstlane2.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
+; CHECK-NOT: v_readfirstlane_b32
+define void @test_readfirstlane_p3(i8 addrspace(3)* addrspace(1)* %out, i8 addrspace(3)* %src) #1 {
+  %readfirstlane = call i8 addrspace(3)* @llvm.amdgcn.readfirstlane2.p3i8(i8 addrspace(3)* %src)
+  store i8 addrspace(3)* %readfirstlane, i8 addrspace(3)* addrspace(1)* %out, align 4
+  ret void
+}
+
+; CHECK-LABEL: {{^}}test_readfirstlane_i16:
+; CHECK: v_readfirstlane_b32 s{{[0-9]+}}, v2
+; CHECK-NOT: v_readfirstlane_b32
+define void @test_readfirstlane_i16(i16 addrspace(1)* %out, i16 %src) {
+  %readfirstlane = call i16 @llvm.amdgcn.readfirstlane2.i16(i16 %src)
+  store i16 %readfirstlane, i16 addrspace(1)* %out, align 2
+  ret void
+}
+
+; CHECK-LABEL: {{^}}test_readfirstlane_f16:
+; CHECK: v_readfirstlane_b32 s{{[0-9]+}}, v2
+; CHECK-NOT: v_readfirstlane_b32
+define void @test_readfirstlane_f16(half addrspace(1)* %out, half %src) {
+  %readfirstlane = call half @llvm.amdgcn.readfirstlane2.f16(half %src)
+  store half %readfirstlane, half addrspace(1)* %out, align 2
+  ret void
+}
+
+; CHECK-LABEL: {{^}}test_readfirstlane_v3i16:
+; CHECK-DAG: v_readfirstlane_b32 s{{[0-9]+}},
+; CHECK-DAG: v_readfirstlane_b32 s{{[0-9]+}},
+; CHECK-NOT: v_readfirstlane_b32
+define void @test_readfirstlane_v3i16(<3 x i16> addrspace(1)* %out, <3 x i16> %src) {
+  %readfirstlane = call <3 x i16> @llvm.amdgcn.readfirstlane2.v3i16(<3 x i16> %src)
+  store <3 x i16> %readfirstlane, <3 x i16> addrspace(1)* %out, align 2
+  ret void
+}
+
+; CHECK-LABEL: {{^}}test_readfirstlane_v9f32:
+; CHECK-DAG: v_readfirstlane_b32 s{{[0-9]+}}, v2
+; CHECK-DAG: v_readfirstlane_b32 s{{[0-9]+}}, v3
+; CHECK-DAG: v_readfirstlane_b32 s{{[0-9]+}}, v4
+; CHECK-DAG: v_readfirstlane_b32 s{{[0-9]+}}, v5
+; CHECK-DAG: v_readfirstlane_b32 s{{[0-9]+}}, v6
+; CHECK-DAG: v_readfirstlane_b32 s{{[0-9]+}}, v7
+; CHECK-DAG: v_readfirstlane_b32 s{{[0-9]+}}, v8
+; CHECK-DAG: v_readfirstlane_b32 s{{[0-9]+}}, v9
+; CHECK-DAG: v_readfirstlane_b32 s{{[0-9]+}}, v10
+; CHECK-NOT: v_readfirstlane_b32
+define void @test_readfirstlane_v9f32(<9 x float> addrspace(1)* %out, <9 x float> %src) {
+  %readfirstlane = call <9 x float> @llvm.amdgcn.readfirstlane2.v9f32(<9 x float> %src)
+  store <9 x float> %readfirstlane, <9 x float> addrspace(1)* %out, align 2
+  ret void
+}
+
 attributes #0 = { nounwind readnone convergent }
 attributes #1 = { nounwind }

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

 ; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=fiji -verify-machineinstrs < %s | FileCheck -enable-var-scope %s
 
 declare i32 @llvm.amdgcn.readlane(i32, i32) #0
+declare i8 addrspace(3)* @llvm.amdgcn.readlane2.p3i8(i8 addrspace(3)*, i32) #0
+declare i16 @llvm.amdgcn.readlane2.i16(i16, i32) #0
+declare half @llvm.amdgcn.readlane2.f16(half, i32) #0
+declare <3 x i16> @llvm.amdgcn.readlane2.v3i16(<3 x i16>, i32) #0
+declare <9 x float> @llvm.amdgcn.readlane2.v9f32(<9 x float>, i32) #0
 
 ; CHECK-LABEL: {{^}}test_readlane_sreg_sreg:
 ; CHECK-NOT: v_readlane_b32
 define amdgpu_kernel void @test_readlane_sreg_sreg(i32 %src0, i32 %src1) #1 {
   %readlane = call i32 @llvm.amdgcn.readlane(i32 %src0, i32 %src1)
   call void asm sideeffect "; use $0", "s"(i32 %readlane)
   ret void
 }

 ; CHECK: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[VCOPY]]
 define amdgpu_kernel void @test_readlane_copy_from_sgpr(i32 addrspace(1)* %out) #1 {
   %sgpr = call i32 asm "s_mov_b32 $0, 0", "=s"()
   %readfirstlane = call i32 @llvm.amdgcn.readlane(i32 %sgpr, i32 7)
   store i32 %readfirstlane, i32 addrspace(1)* %out, align 4
   ret void
 }
 
+; CHECK-LABEL: {{^}}test_readlane_p3:
+; CHECK: v_readlane_b32 s{{[0-9]+}}, v2, 15
+; CHECK-NOT: v_readlane_b32
+define void @test_readlane_p3(i8 addrspace(3)* addrspace(1)* %out, i8 addrspace(3)* %src) #1 {
+  %readlane = call i8 addrspace(3)* @llvm.amdgcn.readlane2.p3i8(i8 addrspace(3)* %src, i32 15)
+  store i8 addrspace(3)* %readlane, i8 addrspace(3)* addrspace(1)* %out, align 4
+  ret void
+}
+
+; CHECK-LABEL: {{^}}test_readlane_i16:
+; CHECK: v_readlane_b32 s{{[0-9]+}}, v2, 15
+; CHECK-NOT: v_readlane_b32
+define void @test_readlane_i16(i16 addrspace(1)* %out, i16 %src) {
+  %readlane = call i16 @llvm.amdgcn.readlane2.i16(i16 %src, i32 15)
+  store i16 %readlane, i16 addrspace(1)* %out, align 2
+  ret void
+}
+
+; CHECK-LABEL: {{^}}test_readlane_f16:
+; CHECK: v_readlane_b32 s{{[0-9]+}}, v2, 15
+; CHECK-NOT: v_readlane_b32
+define void @test_readlane_f16(half addrspace(1)* %out, half %src) {
+  %readlane = call half @llvm.amdgcn.readlane2.f16(half %src, i32 15)
+  store half %readlane, half addrspace(1)* %out, align 2
+  ret void
+}
+
+; CHECK-LABEL: {{^}}test_readlane_v3i16:
+; CHECK-DAG: v_readlane_b32 s{{[0-9]+}},
+; CHECK-DAG: v_readlane_b32 s{{[0-9]+}},
+; CHECK-NOT: v_readlane_b32
+define void @test_readlane_v3i16(<3 x i16> addrspace(1)* %out, <3 x i16> %src) {
+  %readlane = call <3 x i16> @llvm.amdgcn.readlane2.v3i16(<3 x i16> %src, i32 15)
+  store <3 x i16> %readlane, <3 x i16> addrspace(1)* %out, align 2
+  ret void
+}
+
+; CHECK-LABEL: {{^}}test_readlane_v9f32:
+; CHECK-DAG: v_readlane_b32 s{{[0-9]+}}, v2, 15
+; CHECK-DAG: v_readlane_b32 s{{[0-9]+}}, v3, 15
+; CHECK-DAG: v_readlane_b32 s{{[0-9]+}}, v4, 15
+; CHECK-DAG: v_readlane_b32 s{{[0-9]+}}, v5, 15
+; CHECK-DAG: v_readlane_b32 s{{[0-9]+}}, v6, 15
+; CHECK-DAG: v_readlane_b32 s{{[0-9]+}}, v7, 15
+; CHECK-DAG: v_readlane_b32 s{{[0-9]+}}, v8, 15
+; CHECK-DAG: v_readlane_b32 s{{[0-9]+}}, v9, 15
+; CHECK-DAG: v_readlane_b32 s{{[0-9]+}}, v10, 15
+; CHECK-NOT: v_readlane_b32
+define void @test_readlane_v9f32(<9 x float> addrspace(1)* %out, <9 x float> %src) {
+  %readlane = call <9 x float> @llvm.amdgcn.readlane2.v9f32(<9 x float> %src, i32 15)
+  store <9 x float> %readlane, <9 x float> addrspace(1)* %out, align 2
+  ret void
+}
+
 declare i32 @llvm.amdgcn.workitem.id.x() #2
 
 attributes #0 = { nounwind readnone convergent }
 attributes #1 = { nounwind }
 attributes #2 = { nounwind readnone }

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

 ; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx700 -verify-machineinstrs < %s | FileCheck -check-prefixes=CHECK,CI,CIGFX9 %s
 ; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx802 -verify-machineinstrs < %s | FileCheck -check-prefixes=CHECK,GFX9,CIGFX9 %s
 ; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx1010 -verify-machineinstrs < %s | FileCheck -check-prefixes=CHECK,GFX10 %s
 
 declare i32 @llvm.amdgcn.writelane(i32, i32, i32) #0
+declare i8 addrspace(3)* @llvm.amdgcn.writelane2.p3i8(i8 addrspace(3)*, i32, i8 addrspace(3)*) #0
+declare i16 @llvm.amdgcn.writelane2.i16(i16, i32, i16) #0
+declare half @llvm.amdgcn.writelane2.f16(half, i32, half) #0
+declare <3 x i16> @llvm.amdgcn.writelane2.v3i16(<3 x i16>, i32, <3 x i16>) #0
+declare <9 x float> @llvm.amdgcn.writelane2.v9f32(<9 x float>, i32, <9 x float>) #0
 
 ; CHECK-LABEL: {{^}}test_writelane_sreg:
 ; CIGFX9: v_writelane_b32 v{{[0-9]+}}, s{{[0-9]+}}, m0
 ; GFX10: v_writelane_b32 v{{[0-9]+}}, s{{[0-9]+}}, s{{[0-9]+}}
 define amdgpu_kernel void @test_writelane_sreg(i32 addrspace(1)* %out, i32 %src0, i32 %src1) #1 {
   %oldval = load i32, i32 addrspace(1)* %out
   %writelane = call i32 @llvm.amdgcn.writelane(i32 %src0, i32 %src1, i32 %oldval)
   store i32 %writelane, i32 addrspace(1)* %out, align 4

 ; CIGFX9: v_writelane_b32 [[OLDVAL]], s{{[0-9]+}}, m0
 ; GFX10: v_writelane_b32 [[OLDVAL]], s{{[0-9]+}}, s{{[0-9]+}}
 define amdgpu_kernel void @test_writelane_imm_oldval(i32 addrspace(1)* %out, i32 %src0, i32 %src1) #1 {
   %writelane = call i32 @llvm.amdgcn.writelane(i32 %src0, i32 %src1, i32 42)
   store i32 %writelane, i32 addrspace(1)* %out, align 4
   ret void
 }
 
+; CHECK-LABEL: {{^}}test_writelane_p3:
+; CHECK: v_writelane_b32 v{{[0-9]+}}, 0, 15
+; CHECK-NOT: v_writelane_b32
+define void @test_writelane_p3(i8 addrspace(3)* addrspace(1)* %out, i8 addrspace(3)* %src) #1 {
+  %writelane = call i8 addrspace(3)* @llvm.amdgcn.writelane2.p3i8(i8 addrspace(3)* null, i32 15, i8 addrspace(3)* %src)
+  store i8 addrspace(3)* %writelane, i8 addrspace(3)* addrspace(1)* %out, align 4
+  ret void
+}
+
+; CHECK-LABEL: {{^}}test_writelane_i16:
+; CHECK: v_writelane_b32 v{{[0-9]+}},
+; CHECK-NOT: v_writelane_b32
+define void @test_writelane_i16(i16 addrspace(1)* %out, i16 %src) {
+  %writelane = call i16 @llvm.amdgcn.writelane2.i16(i16 1234, i32 15, i16 %src)
+  store i16 %writelane, i16 addrspace(1)* %out, align 2
+  ret void
+}
+
+; CHECK-LABEL: {{^}}test_writelane_f16:
+; CHECK: v_writelane_b32 v{{[0-9]+}},
+; CHECK-NOT: v_writelane_b32
+define void @test_writelane_f16(half addrspace(1)* %out, half %src) {
+  %writelane = call half @llvm.amdgcn.writelane2.f16(half 1.0, i32 15, half %src)
+  store half %writelane, half addrspace(1)* %out, align 2
+  ret void
+}
+
+; CHECK-LABEL: {{^}}test_writelane_v3i16:
+; CHECK-DAG: v_writelane_b32 v{{[0-9]+}},
+; CHECK-DAG: v_writelane_b32 v{{[0-9]+}},
+; CHECK-NOT: v_writelane_b32
+define void @test_writelane_v3i16(<3 x i16> addrspace(1)* %out, <3 x i16> %src) {
+  %writelane = call <3 x i16> @llvm.amdgcn.writelane2.v3i16(<3 x i16> zeroinitializer, i32 15, <3 x i16> %src)
+  store <3 x i16> %writelane, <3 x i16> addrspace(1)* %out, align 2
+  ret void
+}
+
+; CHECK-LABEL: {{^}}test_writelane_v9f32:
+; CHECK-DAG: v_writelane_b32 v{{[0-9]+}}, 0, 15
+; CHECK-DAG: v_writelane_b32 v{{[0-9]+}}, 0, 15
+; CHECK-DAG: v_writelane_b32 v{{[0-9]+}}, 0, 15
+; CHECK-DAG: v_writelane_b32 v{{[0-9]+}}, 0, 15
+; CHECK-DAG: v_writelane_b32 v{{[0-9]+}}, 0, 15
+; CHECK-DAG: v_writelane_b32 v{{[0-9]+}}, 0, 15
+; CHECK-DAG: v_writelane_b32 v{{[0-9]+}}, 0, 15
+; CHECK-DAG: v_writelane_b32 v{{[0-9]+}}, 0, 15
+; CHECK-DAG: v_writelane_b32 v{{[0-9]+}}, 0, 15
+; CHECK-NOT: v_writelane_b32
+define void @test_writelane_v9f32(<9 x float> addrspace(1)* %out, <9 x float> %src) {
+  %writelane = call <9 x float> @llvm.amdgcn.writelane2.v9f32(<9 x float> zeroinitializer, i32 15, <9 x float> %src)
+  store <9 x float> %writelane, <9 x float> addrspace(1)* %out, align 2
+  ret void
+}
+
 declare i32 @llvm.amdgcn.workitem.id.x() #2
 
 attributes #0 = { nounwind readnone convergent }
 attributes #1 = { nounwind }
 attributes #2 = { nounwind readnone }

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:    store volatile i32 [[VAR]], i32* undef, align 4
+; CHECK-NEXT:    [[TMP1:%.*]] = call i32 @llvm.amdgcn.readfirstlane2.i32(i32 [[ARG:%.*]])
+; CHECK-NEXT:    store volatile i32 [[TMP1]], 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:    ret i32 [[READ0]]
+; CHECK-NEXT:    [[TMP1:%.*]] = call i32 @llvm.amdgcn.readfirstlane2.i32(i32 [[ARG:%.*]])
+; CHECK-NEXT:    ret i32 [[TMP1]]
 ;
   %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:    ret i32 [[READ0]]
+; CHECK-NEXT:    [[TMP1:%.*]] = call i32 @llvm.amdgcn.readfirstlane2.i32(i32 [[ARG:%.*]])
+; CHECK-NEXT:    ret i32 [[TMP1]]
 ;
   %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:    [[TMP0:%.*]] = call i32 @llvm.amdgcn.readfirstlane2.i32(i32 [[ARG:%.*]])
 ; CHECK-NEXT:    br label [[BB1:%.*]]
 ; CHECK:       bb1:
-; CHECK-NEXT:    [[READ1:%.*]] = call i32 @llvm.amdgcn.readfirstlane(i32 [[READ0]])
-; CHECK-NEXT:    ret i32 [[READ1]]
+; CHECK-NEXT:    [[TMP1:%.*]] = call i32 @llvm.amdgcn.readfirstlane2.i32(i32 [[TMP0]])
+; CHECK-NEXT:    ret i32 [[TMP1]]
 ;
 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:    [[TMP0:%.*]] = call i32 @llvm.amdgcn.readlane2.i32(i32 [[ARG:%.*]], i32 0)
 ; CHECK-NEXT:    br label [[BB1:%.*]]
 ; CHECK:       bb1:
-; CHECK-NEXT:    [[READ1:%.*]] = call i32 @llvm.amdgcn.readfirstlane(i32 [[READ0]])
-; CHECK-NEXT:    ret i32 [[READ1]]
+; CHECK-NEXT:    [[TMP1:%.*]] = call i32 @llvm.amdgcn.readfirstlane2.i32(i32 [[TMP0]])
+; CHECK-NEXT:    ret i32 [[TMP1]]
 ;
 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.readfirstlane2.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.readfirstlane2.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.readfirstlane2.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(
-; CHECK-NEXT:    [[VAR:%.*]] = call i32 @llvm.amdgcn.readlane(i32 [[ARG:%.*]], i32 7)
-; CHECK-NEXT:    store volatile i32 [[VAR]], i32* undef, align 4
+; CHECK-NEXT:    [[TMP1:%.*]] = call i32 @llvm.amdgcn.readlane2.i32(i32 [[ARG:%.*]], i32 7)
+; CHECK-NEXT:    store volatile i32 [[TMP1]], 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.readlane(i32 %arg, i32 7)
   %zero = call i32 @llvm.amdgcn.readlane(i32 0, i32 %lane)
   %imm = call i32 @llvm.amdgcn.readlane(i32 123, i32 %lane)
   %constexpr = call i32 @llvm.amdgcn.readlane(i32 ptrtoint (i32* @gv to i32), i32 %lane)
   %undef = call i32 @llvm.amdgcn.readlane(i32 undef, i32 %lane)
   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 @readlane_idempotent(i32 %arg, i32 %lane) {
 ; CHECK-LABEL: @readlane_idempotent(
-; CHECK-NEXT:    [[READ0:%.*]] = call i32 @llvm.amdgcn.readlane(i32 [[ARG:%.*]], i32 [[LANE:%.*]])
-; CHECK-NEXT:    ret i32 [[READ0]]
+; CHECK-NEXT:    [[TMP1:%.*]] = call i32 @llvm.amdgcn.readlane2.i32(i32 [[ARG:%.*]], i32 [[LANE:%.*]])
+; CHECK-NEXT:    ret i32 [[TMP1]]
 ;
   %read0 = call i32 @llvm.amdgcn.readlane(i32 %arg, i32 %lane)
   %read1 = call i32 @llvm.amdgcn.readlane(i32 %read0, i32 %lane)
   ret i32 %read1
 }
 
 define i32 @readlane_idempotent_different_lanes(i32 %arg, i32 %lane0, i32 %lane1) {
 ; CHECK-LABEL: @readlane_idempotent_different_lanes(
-; CHECK-NEXT:    [[READ0:%.*]] = call i32 @llvm.amdgcn.readlane(i32 [[ARG:%.*]], i32 [[LANE0:%.*]])
-; CHECK-NEXT:    [[READ1:%.*]] = call i32 @llvm.amdgcn.readlane(i32 [[READ0]], i32 [[LANE1:%.*]])
-; CHECK-NEXT:    ret i32 [[READ1]]
+; CHECK-NEXT:    [[TMP1:%.*]] = call i32 @llvm.amdgcn.readlane2.i32(i32 [[ARG:%.*]], i32 [[LANE0:%.*]])
+; CHECK-NEXT:    [[TMP2:%.*]] = call i32 @llvm.amdgcn.readlane2.i32(i32 [[TMP1]], i32 [[LANE1:%.*]])
+; CHECK-NEXT:    ret i32 [[TMP2]]
 ;
   %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:    ret i32 [[READ0]]
+; CHECK-NEXT:    [[TMP1:%.*]] = call i32 @llvm.amdgcn.readfirstlane2.i32(i32 [[ARG:%.*]])
+; CHECK-NEXT:    ret i32 [[TMP1]]
 ;
   %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) {
 ; CHECK-LABEL: @readlane_idempotent_different_block(
 ; CHECK-NEXT:  bb0:
-; CHECK-NEXT:    [[READ0:%.*]] = call i32 @llvm.amdgcn.readlane(i32 [[ARG:%.*]], i32 [[LANE:%.*]])
+; CHECK-NEXT:    [[TMP0:%.*]] = call i32 @llvm.amdgcn.readlane2.i32(i32 [[ARG:%.*]], i32 [[LANE:%.*]])
 ; CHECK-NEXT:    br label [[BB1:%.*]]
 ; CHECK:       bb1:
-; CHECK-NEXT:    [[READ1:%.*]] = call i32 @llvm.amdgcn.readlane(i32 [[READ0]], i32 [[LANE]])
-; CHECK-NEXT:    ret i32 [[READ1]]
+; CHECK-NEXT:    [[TMP1:%.*]] = call i32 @llvm.amdgcn.readlane2.i32(i32 [[TMP0]], i32 [[LANE]])
+; CHECK-NEXT:    ret i32 [[TMP1]]
 ;
 bb0:
   %read0 = call i32 @llvm.amdgcn.readlane(i32 %arg, i32 %lane)
   br label %bb1
 
 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:    [[TMP0:%.*]] = call i32 @llvm.amdgcn.readfirstlane2.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]]
+; CHECK-NEXT:    [[TMP1:%.*]] = call i32 @llvm.amdgcn.readlane2.i32(i32 [[TMP0]], i32 0)
+; CHECK-NEXT:    ret i32 [[TMP1]]
 ;
 bb0:
   %read0 = call i32 @llvm.amdgcn.readfirstlane(i32 %arg)
   br label %bb1
 
 bb1:
   %read1 = call i32 @llvm.amdgcn.readlane(i32 %read0, i32 0)
   ret i32 %read1