[AMDGPU] gfx1010 wave32 icmp/fcmp intrinsic changes for wave32

[rampitec]

Summary

Another independent part of wave32 support to reduce D63204

[arsenm]

lib/Target/AMDGPU/AMDGPUAtomicOptimizer.cpp

252

This one should query the wavesize?

lib/Target/AMDGPU/SIInstructions.td

606–614

I don't think these should ever get here? These should have been turned into AMDGPUSsetcc

[rampitec]

lib/Target/AMDGPU/AMDGPUAtomicOptimizer.cpp

252

They actually don't, they are anyint and 64 fits it.

lib/Target/AMDGPU/SIInstructions.td

606–614

That's an inplace replacement. We can explore if it still used or not, but keeping just wave64 version is clearly wrong.

[arsenm]

LGTM

[hakzsam]

HI,

This change introduces a bunch of failures with CTS on GFX8/GFX9. Looks like the IR validation fail now, see below:

$ RADV_DEBUG=checkir,nocache ./deqp-vk --deqp-case=dEQP-VK.glsl.derivate.dfdxfine.dynamic_loop.vec2_highp
Writing test log into TestResults.qpa
dEQP Core git-c2a38f956feeb2e141c76829c6ecb41c9f68d253 (0xc2a38f95) starting..

target implementation = 'Default'

Test case 'dEQP-VK.glsl.derivate.dfdxfine.dynamic_loop.vec2_highp'..
Intrinsic name not mangled correctly for type arguments! Should be: llvm.amdgcn.icmp.i64.i32
i64 (i32, i32, i32)* @llvm.amdgcn.icmp.i32
in function main
LLVM ERROR: Broken function found, compilation aborted!

Can you investigate please?
Thanks!

[rampitec]

In D63301#1543215, @hakzsam wrote:

HI,

This change introduces a bunch of failures with CTS on GFX8/GFX9. Looks like the IR validation fail now, see below:

$ RADV_DEBUG=checkir,nocache ./deqp-vk --deqp-case=dEQP-VK.glsl.derivate.dfdxfine.dynamic_loop.vec2_highp
Writing test log into TestResults.qpa
dEQP Core git-c2a38f956feeb2e141c76829c6ecb41c9f68d253 (0xc2a38f95) starting..

target implementation = 'Default'

Test case 'dEQP-VK.glsl.derivate.dfdxfine.dynamic_loop.vec2_highp'..
Intrinsic name not mangled correctly for type arguments! Should be: llvm.amdgcn.icmp.i64.i32
i64 (i32, i32, i32)* @llvm.amdgcn.icmp.i32
in function main
LLVM ERROR: Broken function found, compilation aborted!

Can you investigate please?
Thanks!

Intrinsic mangling has changed. Normally llvm can handle it gracefully. How do you load a module? Can you attach a testcase please? I assume autoupgrade code may be needed, but we need to understand the workflow of RADV.

[hakzsam]

Do we need to update the intrinsic name from mesa?
Btw, it's not related to RADV, RadeonSI is probably affected too because the LLVM backend is common.

[hakzsam]

Sorry for the noise.
Should be fixed with https://patchwork.freedesktop.org/patch/310419/?series=62097&rev=1

[rampitec]

In D63301#1543248, @hakzsam wrote:

Sorry for the noise.
Should be fixed with https://patchwork.freedesktop.org/patch/310419/?series=62097&rev=1

Yes, thanks! If module with old name is created on the fly it has no chance to upgrade. Thus clang was updated and Mesa too.

include/llvm/IR/IntrinsicsAMDGPU.td

 // Set EXEC according to a thread count packed in an SGPR input:
 //    thread_count = (input >> bitoffset) & 0x7f;
 // This is always moved to the beginning of the basic block.
 def int_amdgcn_init_exec_from_input : Intrinsic<[],
   [llvm_i32_ty,       // 32-bit SGPR input
    llvm_i32_ty],      // bit offset of the thread count
   [IntrConvergent]>;
 
+def int_amdgcn_wavefrontsize :
+  GCCBuiltin<"__builtin_amdgcn_wavefrontsize">,
+  Intrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable]>;
+
 
 //===----------------------------------------------------------------------===//
 // Instruction Intrinsics
 //===----------------------------------------------------------------------===//
 
 // The first parameter is s_sendmsg immediate (i16),
 // the second one is copied to m0
 def int_amdgcn_s_sendmsg : GCCBuiltin<"__builtin_amdgcn_s_sendmsg">,

 
 def int_amdgcn_cvt_pk_u8_f32 :
   GCCBuiltin<"__builtin_amdgcn_cvt_pk_u8_f32">,
   Intrinsic<[llvm_i32_ty], [llvm_float_ty, llvm_i32_ty, llvm_i32_ty],
   [IntrNoMem, IntrSpeculatable]
 >;
 
 def int_amdgcn_icmp :
-  Intrinsic<[llvm_i64_ty], [llvm_anyint_ty, LLVMMatchType<0>, llvm_i32_ty],
+  Intrinsic<[llvm_anyint_ty], [llvm_anyint_ty, LLVMMatchType<1>, llvm_i32_ty],
             [IntrNoMem, IntrConvergent, ImmArg<2>]>;
 
 def int_amdgcn_fcmp :
-  Intrinsic<[llvm_i64_ty], [llvm_anyfloat_ty, LLVMMatchType<0>, llvm_i32_ty],
+  Intrinsic<[llvm_anyint_ty], [llvm_anyfloat_ty, LLVMMatchType<1>, llvm_i32_ty],
             [IntrNoMem, IntrConvergent, ImmArg<2>]>;
 
 def int_amdgcn_readfirstlane :
   GCCBuiltin<"__builtin_amdgcn_readfirstlane">,
   Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
 
 // The lane argument must be uniform across the currently active threads of the
 // current wave. Otherwise, the result is undefined.
Intrinsic<
     ],
     [IntrNoMem, IntrSpeculatable, ImmArg<3>]
   >;
 
 //===----------------------------------------------------------------------===//
 // Special Intrinsics for backend internal use only. No frontend
 // should emit calls to these.
 // ===----------------------------------------------------------------------===//
-def int_amdgcn_if : Intrinsic<[llvm_i1_ty, llvm_i64_ty],
+def int_amdgcn_if : Intrinsic<[llvm_i1_ty, llvm_anyint_ty],
   [llvm_i1_ty], [IntrConvergent]
 >;
 
-def int_amdgcn_else : Intrinsic<[llvm_i1_ty, llvm_i64_ty],
-  [llvm_i64_ty], [IntrConvergent]
+def int_amdgcn_else : Intrinsic<[llvm_i1_ty, llvm_anyint_ty],
+  [llvm_anyint_ty], [IntrConvergent]
 >;
 
-def int_amdgcn_if_break : Intrinsic<[llvm_i64_ty],
-  [llvm_i1_ty, llvm_i64_ty], [IntrNoMem, IntrConvergent]
+def int_amdgcn_if_break : Intrinsic<[llvm_anyint_ty],
+  [llvm_i1_ty, llvm_anyint_ty], [IntrNoMem, IntrConvergent]
 >;
 
 def int_amdgcn_loop : Intrinsic<[llvm_i1_ty],
-  [llvm_i64_ty], [IntrConvergent]
+  [llvm_anyint_ty], [IntrConvergent]
 >;
 
-def int_amdgcn_end_cf : Intrinsic<[], [llvm_i64_ty], [IntrConvergent]>;
+def int_amdgcn_end_cf : Intrinsic<[], [llvm_anyint_ty], [IntrConvergent]>;
 
 // Represent unreachable in a divergent region.
 def int_amdgcn_unreachable : Intrinsic<[], [], [IntrConvergent]>;
 
 // Emit 2.5 ulp, no denormal division. Should only be inserted by
 // pass based on !fpmath metadata.
 def int_amdgcn_fdiv_fast : Intrinsic<
   [llvm_float_ty], [llvm_float_ty, llvm_float_ty],
   [IntrNoMem, IntrSpeculatable]
 >;
 }

lib/Target/AMDGPU/AMDGPUAtomicOptimizer.cpp

void AMDGPUAtomicOptimizer::optimizeAtomic(Instruction &I,
 
   // This is the value in the atomic operation we need to combine in order to
   // reduce the number of atomic operations.
   Value *const V = I.getOperand(ValIdx);
 
   // We need to know how many lanes are active within the wavefront, and we do
   // this by doing a ballot of active lanes.
   CallInst *const Ballot =
-      B.CreateIntrinsic(Intrinsic::amdgcn_icmp, {B.getInt32Ty()},
+      B.CreateIntrinsic(Intrinsic::amdgcn_icmp,
+                        {B.getInt64Ty(), B.getInt32Ty()},
                         {B.getInt32(1), B.getInt32(0), B.getInt32(33)});
 
   // We need to know how many lanes are active within the wavefront that are
   // below us. If we counted each lane linearly starting from 0, a lane is
   // below us only if its associated index was less than ours. We do this by
   // using the mbcnt intrinsic.
   Value *const BitCast = B.CreateBitCast(Ballot, VecTy);
   Value *const ExtractLo = B.CreateExtractElement(BitCast, B.getInt32(0));

lib/Target/AMDGPU/AMDGPUSubtarget.h

public:
   /// subtarget's specifications, or does not meet number of waves per execution
   /// unit requirement.
   unsigned getMaxNumVGPRs(const MachineFunction &MF) const;
 
   void getPostRAMutations(
       std::vector<std::unique_ptr<ScheduleDAGMutation>> &Mutations)
       const override;
 
+  bool isWave32() const {
+    return WavefrontSize == 32;
+  }
+
   /// \returns Maximum number of work groups per compute unit supported by the
   /// subtarget and limited by given \p FlatWorkGroupSize.
   unsigned getMaxWorkGroupsPerCU(unsigned FlatWorkGroupSize) const override {
     return AMDGPU::IsaInfo::getMaxWorkGroupsPerCU(this, FlatWorkGroupSize);
   }
 
   /// \returns Minimum flat work group size supported by the subtarget.
   unsigned getMinFlatWorkGroupSize() const override {

lib/Target/AMDGPU/SIAnnotateControlFlow.cpp

 //===- SIAnnotateControlFlow.cpp ------------------------------------------===//
 //
 // 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
 /// Annotates the control flow with hardware specific intrinsics.
 //
 //===----------------------------------------------------------------------===//
 
 #include "AMDGPU.h"
+#include "AMDGPUSubtarget.h"
 #include "llvm/ADT/DepthFirstIterator.h"
 #include "llvm/ADT/STLExtras.h"
 #include "llvm/ADT/SmallVector.h"
 #include "llvm/Analysis/LegacyDivergenceAnalysis.h"
 #include "llvm/Analysis/LoopInfo.h"
+#include "llvm/CodeGen/TargetPassConfig.h"
 #include "llvm/IR/BasicBlock.h"
 #include "llvm/IR/CFG.h"
 #include "llvm/IR/Constant.h"
 #include "llvm/IR/Constants.h"
 #include "llvm/IR/DerivedTypes.h"
 #include "llvm/IR/Dominators.h"
 #include "llvm/IR/Function.h"
 #include "llvm/IR/Instruction.h"

 using StackEntry = std::pair<BasicBlock *, Value *>;
 using StackVector = SmallVector<StackEntry, 16>;
 
 class SIAnnotateControlFlow : public FunctionPass {
   LegacyDivergenceAnalysis *DA;
 
   Type *Boolean;
   Type *Void;
-  Type *Int64;
+  Type *IntMask;
   Type *ReturnStruct;
 
   ConstantInt *BoolTrue;
   ConstantInt *BoolFalse;
   UndefValue *BoolUndef;
-  Constant *Int64Zero;
+  Constant *IntMaskZero;
 
   Function *If;
   Function *Else;
   Function *IfBreak;
   Function *Loop;
   Function *EndCf;
 
   DominatorTree *DT;
   StackVector Stack;
 
   LoopInfo *LI;
 
+  void initialize(Module &M, const GCNSubtarget &ST);
+
   bool isUniform(BranchInst *T);
 
   bool isTopOfStack(BasicBlock *BB);
 
   Value *popSaved();
 
   void push(BasicBlock *BB, Value *Saved);
 
class SIAnnotateControlFlow : public FunctionPass {
 
   void closeControlFlow(BasicBlock *BB);
 
 public:
   static char ID;
 
   SIAnnotateControlFlow() : FunctionPass(ID) {}
 
-  bool doInitialization(Module &M) override;
-
   bool runOnFunction(Function &F) override;
 
   StringRef getPassName() const override { return "SI annotate control flow"; }
 
   void getAnalysisUsage(AnalysisUsage &AU) const override {
     AU.addRequired<LoopInfoWrapperPass>();
     AU.addRequired<DominatorTreeWrapperPass>();
     AU.addRequired<LegacyDivergenceAnalysis>();
     AU.addPreserved<DominatorTreeWrapperPass>();
+    AU.addRequired<TargetPassConfig>();
     FunctionPass::getAnalysisUsage(AU);
   }
 };
 
 } // end anonymous namespace
 
 INITIALIZE_PASS_BEGIN(SIAnnotateControlFlow, DEBUG_TYPE,
                       "Annotate SI Control Flow", false, false)
 INITIALIZE_PASS_DEPENDENCY(DominatorTreeWrapperPass)
 INITIALIZE_PASS_DEPENDENCY(LegacyDivergenceAnalysis)
+INITIALIZE_PASS_DEPENDENCY(TargetPassConfig)
 INITIALIZE_PASS_END(SIAnnotateControlFlow, DEBUG_TYPE,
                     "Annotate SI Control Flow", false, false)
 
 char SIAnnotateControlFlow::ID = 0;
 
 /// Initialize all the types and constants used in the pass
-bool SIAnnotateControlFlow::doInitialization(Module &M) {
+void SIAnnotateControlFlow::initialize(Module &M, const GCNSubtarget &ST) {
   LLVMContext &Context = M.getContext();
 
   Void = Type::getVoidTy(Context);
   Boolean = Type::getInt1Ty(Context);
-  Int64 = Type::getInt64Ty(Context);
-  ReturnStruct = StructType::get(Boolean, Int64);
+  IntMask = ST.isWave32() ? Type::getInt32Ty(Context)
+                           : Type::getInt64Ty(Context);
+  ReturnStruct = StructType::get(Boolean, IntMask);
 
   BoolTrue = ConstantInt::getTrue(Context);
   BoolFalse = ConstantInt::getFalse(Context);
   BoolUndef = UndefValue::get(Boolean);
-  Int64Zero = ConstantInt::get(Int64, 0);
+  IntMaskZero = ConstantInt::get(IntMask, 0);
 
-  If = Intrinsic::getDeclaration(&M, Intrinsic::amdgcn_if);
-  Else = Intrinsic::getDeclaration(&M, Intrinsic::amdgcn_else);
-  IfBreak = Intrinsic::getDeclaration(&M, Intrinsic::amdgcn_if_break);
-  Loop = Intrinsic::getDeclaration(&M, Intrinsic::amdgcn_loop);
-  EndCf = Intrinsic::getDeclaration(&M, Intrinsic::amdgcn_end_cf);
-  return false;
+  If = Intrinsic::getDeclaration(&M, Intrinsic::amdgcn_if, { IntMask });
+  Else = Intrinsic::getDeclaration(&M, Intrinsic::amdgcn_else,
+                                   { IntMask, IntMask });
+  IfBreak = Intrinsic::getDeclaration(&M, Intrinsic::amdgcn_if_break,
+                                      { IntMask, IntMask });
+  Loop = Intrinsic::getDeclaration(&M, Intrinsic::amdgcn_loop, { IntMask });
+  EndCf = Intrinsic::getDeclaration(&M, Intrinsic::amdgcn_end_cf, { IntMask });
 }
 
 /// Is the branch condition uniform or did the StructurizeCFG pass
 /// consider it as such?
 bool SIAnnotateControlFlow::isUniform(BranchInst *T) {
   return DA->isUniform(T) ||
          T->getMetadata("structurizecfg.uniform") != nullptr;
 }
if (isUniform(Term))
     return;
 
   BasicBlock *BB = Term->getParent();
   llvm::Loop *L = LI->getLoopFor(BB);
   if (!L)
     return;
 
   BasicBlock *Target = Term->getSuccessor(1);
-  PHINode *Broken = PHINode::Create(Int64, 0, "phi.broken", &Target->front());
+  PHINode *Broken = PHINode::Create(IntMask, 0, "phi.broken", &Target->front());
 
   Value *Cond = Term->getCondition();
   Term->setCondition(BoolTrue);
   Value *Arg = handleLoopCondition(Cond, Broken, L, Term);
 
   for (BasicBlock *Pred : predecessors(Target)) {
-    Value *PHIValue = Int64Zero;
+    Value *PHIValue = IntMaskZero;
     if (Pred == BB) // Remember the value of the previous iteration.
       PHIValue = Arg;
     // If the backedge from Pred to Target could be executed before the exit
     // of the loop at BB, it should not reset or change "Broken", which keeps
     // track of the number of threads exited the loop at BB.
     else if (L->contains(Pred) && DT->dominates(Pred, BB))
       PHIValue = Broken;
     Broken->addIncoming(PHIValue, Pred);

 }
 
 /// Annotate the control flow with intrinsics so the backend can
 /// recognize if/then/else and loops.
 bool SIAnnotateControlFlow::runOnFunction(Function &F) {
   DT = &getAnalysis<DominatorTreeWrapperPass>().getDomTree();
   LI = &getAnalysis<LoopInfoWrapperPass>().getLoopInfo();
   DA = &getAnalysis<LegacyDivergenceAnalysis>();
+  TargetPassConfig &TPC = getAnalysis<TargetPassConfig>();
+  const TargetMachine &TM = TPC.getTM<TargetMachine>();
+
+  initialize(*F.getParent(), TM.getSubtarget<GCNSubtarget>(F));
 
   for (df_iterator<BasicBlock *> I = df_begin(&F.getEntryBlock()),
        E = df_end(&F.getEntryBlock()); I != E; ++I) {
     BasicBlock *BB = *I;
     BranchInst *Term = dyn_cast<BranchInst>(BB->getTerminator());
 
     if (!Term || Term->isUnconditional()) {
       if (isTopOfStack(BB))

lib/Target/AMDGPU/SIISelLowering.cpp

static SDValue lowerICMPIntrinsic(const SITargetLowering &TLI,
   const auto *CD = cast<ConstantSDNode>(N->getOperand(3));
   int CondCode = CD->getSExtValue();
   if (CondCode < ICmpInst::Predicate::FIRST_ICMP_PREDICATE ||
       CondCode > ICmpInst::Predicate::LAST_ICMP_PREDICATE)
     return DAG.getUNDEF(VT);
 
   ICmpInst::Predicate IcInput = static_cast<ICmpInst::Predicate>(CondCode);
 
-
   SDValue LHS = N->getOperand(1);
   SDValue RHS = N->getOperand(2);
 
   SDLoc DL(N);
 
   EVT CmpVT = LHS.getValueType();
   if (CmpVT == MVT::i16 && !TLI.isTypeLegal(MVT::i16)) {
     unsigned PromoteOp = ICmpInst::isSigned(IcInput) ?
       ISD::SIGN_EXTEND : ISD::ZERO_EXTEND;
     LHS = DAG.getNode(PromoteOp, DL, MVT::i32, LHS);
     RHS = DAG.getNode(PromoteOp, DL, MVT::i32, RHS);
   }
 
   ISD::CondCode CCOpcode = getICmpCondCode(IcInput);
 
-  return DAG.getNode(AMDGPUISD::SETCC, DL, VT, LHS, RHS,
+  unsigned WavefrontSize = TLI.getSubtarget()->getWavefrontSize();
+  EVT CCVT = EVT::getIntegerVT(*DAG.getContext(), WavefrontSize);
+
+  SDValue SetCC = DAG.getNode(AMDGPUISD::SETCC, DL, CCVT, LHS, RHS,
                               DAG.getCondCode(CCOpcode));
+  if (VT.bitsEq(CCVT))
+    return SetCC;
+  return DAG.getZExtOrTrunc(SetCC, DL, VT);
 }
 
 static SDValue lowerFCMPIntrinsic(const SITargetLowering &TLI,
                                   SDNode *N, SelectionDAG &DAG) {
   EVT VT = N->getValueType(0);
   const auto *CD = cast<ConstantSDNode>(N->getOperand(3));
 
   int CondCode = CD->getSExtValue();
static SDValue lowerFCMPIntrinsic(const SITargetLowering &TLI,
 
   if (CmpVT == MVT::f16 && !TLI.isTypeLegal(CmpVT)) {
     Src0 = DAG.getNode(ISD::FP_EXTEND, SL, MVT::f32, Src0);
     Src1 = DAG.getNode(ISD::FP_EXTEND, SL, MVT::f32, Src1);
   }
 
   FCmpInst::Predicate IcInput = static_cast<FCmpInst::Predicate>(CondCode);
   ISD::CondCode CCOpcode = getFCmpCondCode(IcInput);
-  return DAG.getNode(AMDGPUISD::SETCC, SL, VT, Src0,
+  unsigned WavefrontSize = TLI.getSubtarget()->getWavefrontSize();
+  EVT CCVT = EVT::getIntegerVT(*DAG.getContext(), WavefrontSize);
+  SDValue SetCC = DAG.getNode(AMDGPUISD::SETCC, SL, CCVT, Src0,
                               Src1, DAG.getCondCode(CCOpcode));
+  if (VT.bitsEq(CCVT))
+    return SetCC;
+  return DAG.getZExtOrTrunc(SetCC, SL, VT);
 }
 
 void SITargetLowering::ReplaceNodeResults(SDNode *N,
                                           SmallVectorImpl<SDValue> &Results,
                                           SelectionDAG &DAG) const {
   switch (N->getOpcode()) {
   case ISD::INSERT_VECTOR_ELT: {
     if (SDValue Res = lowerINSERT_VECTOR_ELT(SDValue(N, 0), DAG))
case Intrinsic::r600_read_tidig_y:
     return loadInputValue(DAG, &AMDGPU::VGPR_32RegClass, MVT::i32,
                           SDLoc(DAG.getEntryNode()),
                           MFI->getArgInfo().WorkItemIDY);
   case Intrinsic::amdgcn_workitem_id_z:
   case Intrinsic::r600_read_tidig_z:
     return loadInputValue(DAG, &AMDGPU::VGPR_32RegClass, MVT::i32,
                           SDLoc(DAG.getEntryNode()),
                           MFI->getArgInfo().WorkItemIDZ);
+  case Intrinsic::amdgcn_wavefrontsize:
+    return DAG.getConstant(MF.getSubtarget<GCNSubtarget>().getWavefrontSize(),
+                           SDLoc(Op), MVT::i32);
   case Intrinsic::amdgcn_s_buffer_load: {
     unsigned Cache = cast<ConstantSDNode>(Op.getOperand(3))->getZExtValue();
     return lowerSBuffer(VT, DL, Op.getOperand(1), Op.getOperand(2),
                         DAG.getTargetConstant(Cache & 1, DL, MVT::i1), DAG);
   }
   case Intrinsic::amdgcn_fdiv_fast:
     return lowerFDIV_FAST(Op, DAG);
   case Intrinsic::amdgcn_interp_mov: {
SDValue SITargetLowering::LowerINTRINSIC_WO_CHAIN(SDValue Op,
   case Intrinsic::amdgcn_wwm: {
     SDValue Src = Op.getOperand(1);
     return SDValue(DAG.getMachineNode(AMDGPU::WWM, DL, Src.getValueType(), Src),
                    0);
   }
   case Intrinsic::amdgcn_fmad_ftz:
     return DAG.getNode(AMDGPUISD::FMAD_FTZ, DL, VT, Op.getOperand(1),
                        Op.getOperand(2), Op.getOperand(3));
+
+  case Intrinsic::amdgcn_if_break:
+    return SDValue(DAG.getMachineNode(AMDGPU::SI_IF_BREAK, DL, VT,
+                                      Op->getOperand(1), Op->getOperand(2)), 0);
+
   default:
     if (const AMDGPU::ImageDimIntrinsicInfo *ImageDimIntr =
             AMDGPU::getImageDimIntrinsicInfo(IntrinsicID))
       return lowerImage(Op, ImageDimIntr, DAG);
 
     return Op;
   }
 }
case Intrinsic::amdgcn_struct_buffer_store_format: {
     EVT VDataType = VData.getValueType().getScalarType();
     if (VDataType == MVT::i8 || VDataType == MVT::i16)
       return handleByteShortBufferStores(DAG, VDataType, DL, Ops, M);
 
     return DAG.getMemIntrinsicNode(Opc, DL, Op->getVTList(), Ops,
                                    M->getMemoryVT(), M->getMemOperand());
   }
 
+  case Intrinsic::amdgcn_end_cf:
+    return SDValue(DAG.getMachineNode(AMDGPU::SI_END_CF, DL, MVT::Other,
+                                      Op->getOperand(2), Chain), 0);
+
   default: {
     if (const AMDGPU::ImageDimIntrinsicInfo *ImageDimIntr =
             AMDGPU::getImageDimIntrinsicInfo(IntrinsicID))
       return lowerImage(Op, ImageDimIntr, DAG);
 
     return Op;
   }
   }

lib/Target/AMDGPU/SIInstructions.td

 
 def : Pat <
   (int_amdgcn_kill (i1 (setcc f32:$src, InlineFPImm<f32>:$imm, cond:$cond))),
   (SI_KILL_F32_COND_IMM_PSEUDO $src, (bitcast_fpimm_to_i32 $imm), (cond_as_i32imm $cond))
 >;
 
   // TODO: we could add more variants for other types of conditionals
 
 def : Pat <
-  (int_amdgcn_icmp i1:$src, (i1 0), (i32 33)),
+  (i64 (int_amdgcn_icmp i1:$src, (i1 0), (i32 33))),
+  (COPY $src) // Return the SGPRs representing i1 src
+>;
+
+def : Pat <
+  (i32 (int_amdgcn_icmp i1:$src, (i1 0), (i32 33))),
   (COPY $src) // Return the SGPRs representing i1 src
 >;
 
 //===----------------------------------------------------------------------===//
 // VOP1 Patterns
 //===----------------------------------------------------------------------===//
 
 let OtherPredicates = [UnsafeFPMath] in {
 
 //def : RcpPat<V_RCP_F64_e32, f64>;

lib/Transforms/InstCombine/InstCombineCalls.cpp

if (match(Src1, m_Zero()) &&
             SrcLHS = Builder.CreateZExt(SrcLHS, CmpTy);
             SrcRHS = Builder.CreateZExt(SrcRHS, CmpTy);
           }
         }
       } else if (!Ty->isFloatTy() && !Ty->isDoubleTy() && !Ty->isHalfTy())
         break;
 
       Function *NewF =
-          Intrinsic::getDeclaration(II->getModule(), NewIID, SrcLHS->getType());
+          Intrinsic::getDeclaration(II->getModule(), NewIID,
+                                    { II->getType(),
+                                      SrcLHS->getType() });
       Value *Args[] = { SrcLHS, SrcRHS,
                         ConstantInt::get(CC->getType(), SrcPred) };
       CallInst *NewCall = Builder.CreateCall(NewF, Args);
       NewCall->takeName(II);
       return replaceInstUsesWith(*II, NewCall);
     }
 
     break;

test/CodeGen/AMDGPU/diverge-switch-default.ll

 ; is sensitive to optimizations; so we just ensure that the relevant
 ; operations in the block body are indeed in the same block.
 
 ; CHECK: [[PHI:%[a-zA-Z0-9._]+]]  = phi i64
 ; CHECK-NOT: {{ br }}
 ; CHECK: load i8
 ; CHECK-NOT: {{ br }}
 ; CHECK: [[ICMP:%[a-zA-Z0-9._]+]] = icmp eq
-; CHECK: [[IF:%[a-zA-Z0-9._]+]]   = call i64 @llvm.amdgcn.if.break(i1 [[ICMP]], i64 [[PHI]])
-; CHECK: [[LOOP:%[a-zA-Z0-9._]+]] = call i1 @llvm.amdgcn.loop(i64 [[IF]])
+; CHECK: [[IF:%[a-zA-Z0-9._]+]]   = call i64 @llvm.amdgcn.if.break.i64.i64(i1 [[ICMP]], i64 [[PHI]])
+; CHECK: [[LOOP:%[a-zA-Z0-9._]+]] = call i1 @llvm.amdgcn.loop.i64(i64 [[IF]])
 ; CHECK: br i1 [[LOOP]]
 
 sw.while:
   %p = phi i8 addrspace(1)* [ %gep_in, %sw.epilog ], [ %incdec.ptr, %sw.while ]
   %count = phi i32 [ 0, %sw.epilog ], [ %count.inc, %sw.while ]
   %char = load i8, i8 addrspace(1)* %p, align 1
   %tobool = icmp eq i8 %char, 0
   %incdec.ptr = getelementptr inbounds i8, i8 addrspace(1)* %p, i64 1

test/CodeGen/AMDGPU/loop_break.ll

 
 ; OPT: bb4:
 ; OPT: load volatile
 ; OPT: icmp slt i32
 ; OPT: xor i1 %cmp1
 ; OPT: br label %Flow
 
 ; OPT: Flow:
-; OPT: call i64 @llvm.amdgcn.if.break(
-; OPT: call i1 @llvm.amdgcn.loop(i64
+; OPT: call i64 @llvm.amdgcn.if.break.i64.i64(
+; OPT: call i1 @llvm.amdgcn.loop.i64(i64
 ; OPT: br i1 %{{[0-9]+}}, label %bb9, label %bb1
 
 ; OPT: bb9:
-; OPT: call void @llvm.amdgcn.end.cf(i64
+; OPT: call void @llvm.amdgcn.end.cf.i64(i64
 
 ; GCN-LABEL: {{^}}break_loop:
 ; GCN:      s_mov_b64         [[OUTER_MASK:s\[[0-9]+:[0-9]+\]]], 0{{$}}
 
 ; GCN: [[LOOP_ENTRY:BB[0-9]+_[0-9]+]]: ; %bb1
 ; GCN:      v_cmp_lt_i32_e32  vcc, -1
 ; GCN:      s_and_b64         vcc, exec, vcc
 ; GCN:      s_or_b64          [[INNER_MASK:s\[[0-9]+:[0-9]+\]]], [[INNER_MASK]], exec

 ; OPT: bb4:
 ; OPT-NEXT: %load = load volatile i32, i32 addrspace(1)* undef, align 4
 ; OPT-NEXT: %cmp1 = icmp sge i32 %tmp, %load
 ; OPT-NEXT: br label %Flow
 
 ; OPT: Flow:
 ; OPT-NEXT: %tmp2 = phi i32 [ %lsr.iv.next, %bb4 ], [ undef, %bb1 ]
 ; OPT-NEXT: %tmp3 = phi i1 [ %cmp1, %bb4 ], [ undef, %bb1 ]
-; OPT-NEXT: %0 = call i64 @llvm.amdgcn.if.break(i1 %tmp3, i64 %phi.broken)
-; OPT-NEXT: %1 = call i1 @llvm.amdgcn.loop(i64 %0)
+; OPT-NEXT: %0 = call i64 @llvm.amdgcn.if.break.i64.i64(i1 %tmp3, i64 %phi.broken)
+; OPT-NEXT: %1 = call i1 @llvm.amdgcn.loop.i64(i64 %0)
 ; OPT-NEXT: br i1 %1, label %bb9, label %bb1
 
 ; OPT: bb9:                                              ; preds = %Flow
-; OPT-NEXT: call void @llvm.amdgcn.end.cf(i64 %0)
+; OPT-NEXT: call void @llvm.amdgcn.end.cf.i64(i64 %0)
 ; OPT-NEXT: store volatile i32 7
 ; OPT-NEXT: ret void
 define amdgpu_kernel void @undef_phi_cond_break_loop(i32 %arg) #0 {
 bb:
   %id = call i32 @llvm.amdgcn.workitem.id.x()
   %tmp = sub i32 %id, %arg
   br label %bb1
 

 ; OPT: bb4:
 ; OPT-NEXT: %load = load volatile i32, i32 addrspace(1)* undef, align 4
 ; OPT-NEXT: %cmp1 = icmp sge i32 %tmp, %load
 ; OPT-NEXT: br label %Flow
 
 ; OPT: Flow:
 ; OPT-NEXT: %tmp2 = phi i32 [ %lsr.iv.next, %bb4 ], [ undef, %bb1 ]
 ; OPT-NEXT: %tmp3 = phi i1 [ %cmp1, %bb4 ], [ icmp ne (i32 addrspace(3)* inttoptr (i32 4 to i32 addrspace(3)*), i32 addrspace(3)* @lds), %bb1 ]
-; OPT-NEXT: %0 = call i64 @llvm.amdgcn.if.break(i1 %tmp3, i64 %phi.broken)
-; OPT-NEXT: %1 = call i1 @llvm.amdgcn.loop(i64 %0)
+; OPT-NEXT: %0 = call i64 @llvm.amdgcn.if.break.i64.i64(i1 %tmp3, i64 %phi.broken)
+; OPT-NEXT: %1 = call i1 @llvm.amdgcn.loop.i64(i64 %0)
 ; OPT-NEXT: br i1 %1, label %bb9, label %bb1
 
 ; OPT: bb9:                                              ; preds = %Flow
-; OPT-NEXT: call void @llvm.amdgcn.end.cf(i64 %0)
+; OPT-NEXT: call void @llvm.amdgcn.end.cf.i64(i64 %0)
 ; OPT-NEXT: store volatile i32 7
 ; OPT-NEXT: ret void
 define amdgpu_kernel void @constexpr_phi_cond_break_loop(i32 %arg) #0 {
 bb:
   %id = call i32 @llvm.amdgcn.workitem.id.x()
   %tmp = sub i32 %id, %arg
   br label %bb1
 

 ; OPT: bb4:
 ; OPT-NEXT: %load = load volatile i32, i32 addrspace(1)* undef, align 4
 ; OPT-NEXT: %cmp1 = icmp sge i32 %tmp, %load
 ; OPT-NEXT: br label %Flow
 
 ; OPT: Flow:
 ; OPT-NEXT: %tmp2 = phi i32 [ %lsr.iv.next, %bb4 ], [ undef, %bb1 ]
 ; OPT-NEXT: %tmp3 = phi i1 [ %cmp1, %bb4 ], [ true, %bb1 ]
-; OPT-NEXT: %0 = call i64 @llvm.amdgcn.if.break(i1 %tmp3, i64 %phi.broken)
-; OPT-NEXT: %1 = call i1 @llvm.amdgcn.loop(i64 %0)
+; OPT-NEXT: %0 = call i64 @llvm.amdgcn.if.break.i64.i64(i1 %tmp3, i64 %phi.broken)
+; OPT-NEXT: %1 = call i1 @llvm.amdgcn.loop.i64(i64 %0)
 ; OPT-NEXT: br i1 %1, label %bb9, label %bb1
 
 ; OPT: bb9:                                              ; preds = %Flow
-; OPT-NEXT: call void @llvm.amdgcn.end.cf(i64 %0)
+; OPT-NEXT: call void @llvm.amdgcn.end.cf.i64(i64 %0)
 ; OPT-NEXT: store volatile i32 7
 ; OPT-NEXT: ret void
 define amdgpu_kernel void @true_phi_cond_break_loop(i32 %arg) #0 {
 bb:
   %id = call i32 @llvm.amdgcn.workitem.id.x()
   %tmp = sub i32 %id, %arg
   br label %bb1
 

 ; OPT: bb4:
 ; OPT-NEXT: %load = load volatile i32, i32 addrspace(1)* undef, align 4
 ; OPT-NEXT: %cmp1 = icmp sge i32 %tmp, %load
 ; OPT-NEXT: br label %Flow
 
 ; OPT: Flow:
 ; OPT-NEXT: %tmp2 = phi i32 [ %lsr.iv.next, %bb4 ], [ undef, %bb1 ]
 ; OPT-NEXT: %tmp3 = phi i1 [ %cmp1, %bb4 ], [ false, %bb1 ]
-; OPT-NEXT: %0 = call i64 @llvm.amdgcn.if.break(i1 %tmp3, i64 %phi.broken)
-; OPT-NEXT: %1 = call i1 @llvm.amdgcn.loop(i64 %0)
+; OPT-NEXT: %0 = call i64 @llvm.amdgcn.if.break.i64.i64(i1 %tmp3, i64 %phi.broken)
+; OPT-NEXT: %1 = call i1 @llvm.amdgcn.loop.i64(i64 %0)
 ; OPT-NEXT: br i1 %1, label %bb9, label %bb1
 
 ; OPT: bb9:                                              ; preds = %Flow
-; OPT-NEXT: call void @llvm.amdgcn.end.cf(i64 %0)
+; OPT-NEXT: call void @llvm.amdgcn.end.cf.i64(i64 %0)
 ; OPT-NEXT: store volatile i32 7
 ; OPT-NEXT: ret void
 define amdgpu_kernel void @false_phi_cond_break_loop(i32 %arg) #0 {
 bb:
   %id = call i32 @llvm.amdgcn.workitem.id.x()
   %tmp = sub i32 %id, %arg
   br label %bb1
 

 ; OPT-NEXT: %load = load volatile i32, i32 addrspace(1)* undef, align 4
 ; OPT-NEXT: %cmp1 = icmp sge i32 %tmp, %load
 ; OPT-NEXT: br label %Flow
 
 ; OPT: Flow:
 ; OPT-NEXT: %tmp2 = phi i32 [ %lsr.iv.next, %bb4 ], [ undef, %bb1 ]
 ; OPT-NEXT: %tmp3 = phi i1 [ %cmp1, %bb4 ], [ true, %bb1 ]
 ; OPT-NEXT: %0 = xor i1 %tmp3, true
-; OPT-NEXT: %1 = call i64 @llvm.amdgcn.if.break(i1 %0, i64 %phi.broken)
-; OPT-NEXT: %2 = call i1 @llvm.amdgcn.loop(i64 %1)
+; OPT-NEXT: %1 = call i64 @llvm.amdgcn.if.break.i64.i64(i1 %0, i64 %phi.broken)
+; OPT-NEXT: %2 = call i1 @llvm.amdgcn.loop.i64(i64 %1)
 ; OPT-NEXT: br i1 %2, label %bb9, label %bb1
 
 ; OPT: bb9:
-; OPT-NEXT: call void @llvm.amdgcn.end.cf(i64 %1)
+; OPT-NEXT: call void @llvm.amdgcn.end.cf.i64(i64 %1)
 ; OPT-NEXT: store volatile i32 7, i32 addrspace(3)* undef
 ; OPT-NEXT: ret void
 define amdgpu_kernel void @invert_true_phi_cond_break_loop(i32 %arg) #0 {
 bb:
   %id = call i32 @llvm.amdgcn.workitem.id.x()
   %tmp = sub i32 %id, %arg
   br label %bb1
 

test/CodeGen/AMDGPU/multi-divergent-exit-region.ll

 ; RUN: opt -mtriple=amdgcn-- -S -amdgpu-unify-divergent-exit-nodes -verify -structurizecfg -verify -si-annotate-control-flow %s | FileCheck -check-prefix=IR %s
 ; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
 
 ; Add an extra verifier runs. There were some cases where invalid IR
 ; was produced but happened to be fixed by the later passes.
 
 ; Make sure divergent control flow with multiple exits from a region
 ; is properly handled. UnifyFunctionExitNodes should be run before
 ; StructurizeCFG.
 
 ; IR-LABEL: @multi_divergent_region_exit_ret_ret(
-; IR: %1 = call { i1, i64 } @llvm.amdgcn.if(i1 %0)
+; IR: %1 = call { i1, i64 } @llvm.amdgcn.if.i64(i1 %0)
 ; IR: %2 = extractvalue { i1, i64 } %1, 0
 ; IR: %3 = extractvalue { i1, i64 } %1, 1
 ; IR: br i1 %2, label %LeafBlock1, label %Flow
 
 ; IR: Flow:
 ; IR: %4 = phi i1 [ true, %LeafBlock1 ], [ false, %entry ]
 ; IR: %5 = phi i1 [ %10, %LeafBlock1 ], [ false, %entry ]
-; IR: %6 = call { i1, i64 } @llvm.amdgcn.else(i64 %3)
+; IR: %6 = call { i1, i64 } @llvm.amdgcn.else.i64.i64(i64 %3)
 ; IR: %7 = extractvalue { i1, i64 } %6, 0
 ; IR: %8 = extractvalue { i1, i64 } %6, 1
 ; IR: br i1 %7, label %LeafBlock, label %Flow1
 
 ; IR: LeafBlock:
 ; IR: br label %Flow1
 
 ; IR: LeafBlock1:
 ; IR: br label %Flow{{$}}
 
 ; IR:  Flow2:
 ; IR: %11 = phi i1 [ false, %exit1 ], [ %15, %Flow1 ]
-; IR: call void @llvm.amdgcn.end.cf(i64 %19)
-; IR: %12 = call { i1, i64 } @llvm.amdgcn.if(i1 %11)
+; IR: call void @llvm.amdgcn.end.cf.i64(i64 %19)
+; IR: %12 = call { i1, i64 } @llvm.amdgcn.if.i64(i1 %11)
 ; IR: %13 = extractvalue { i1, i64 } %12, 0
 ; IR: %14 = extractvalue { i1, i64 } %12, 1
 ; IR: br i1 %13, label %exit0, label %UnifiedReturnBlock
 
 ; IR: exit0:
 ; IR: store volatile i32 9, i32 addrspace(1)* undef
 ; IR: br label %UnifiedReturnBlock
 
 ; IR: Flow1:
 ; IR: %15 = phi i1 [ %SwitchLeaf, %LeafBlock ], [ %4, %Flow ]
 ; IR: %16 = phi i1 [ %9, %LeafBlock ], [ %5, %Flow ]
-; IR: call void @llvm.amdgcn.end.cf(i64 %8)
-; IR: %17 = call { i1, i64 } @llvm.amdgcn.if(i1 %16)
+; IR: call void @llvm.amdgcn.end.cf.i64(i64 %8)
+; IR: %17 = call { i1, i64 } @llvm.amdgcn.if.i64(i1 %16)
 ; IR: %18 = extractvalue { i1, i64 } %17, 0
 ; IR: %19 = extractvalue { i1, i64 } %17, 1
 ; IR: br i1 %18, label %exit1, label %Flow2
 
 ; IR: exit1:
 ; IR: store volatile i32 17, i32 addrspace(3)* undef
 ; IR:  br label %Flow2
 
 ; IR: UnifiedReturnBlock:
-; IR: call void @llvm.amdgcn.end.cf(i64 %14)
+; IR: call void @llvm.amdgcn.end.cf.i64(i64 %14)
 ; IR: ret void
 
 
 ; GCN-LABEL: {{^}}multi_divergent_region_exit_ret_ret:
 
 ; GCN-DAG:  s_mov_b64           [[EXIT1:s\[[0-9]+:[0-9]+\]]], 0
 ; GCN-DAG:  v_cmp_lt_i32_e32    vcc, 1,
 ; GCN-DAG:  s_mov_b64           [[EXIT0:s\[[0-9]+:[0-9]+\]]], 0
exit0:                                     ; preds = %LeafBlock, %LeafBlock1
   ret void
 
 exit1:                                     ; preds = %LeafBlock, %LeafBlock1
   store volatile i32 17, i32 addrspace(3)* undef
   ret void
 }
 
 ; IR-LABEL: @multi_divergent_region_exit_unreachable_unreachable(
-; IR: %1 = call { i1, i64 } @llvm.amdgcn.if(i1 %0)
+; IR: %1 = call { i1, i64 } @llvm.amdgcn.if.i64(i1 %0)
 
-; IR: %6 = call { i1, i64 } @llvm.amdgcn.else(i64 %3)
+; IR: %6 = call { i1, i64 } @llvm.amdgcn.else.i64.i64(i64 %3)
 
 ; IR: %11 = phi i1 [ false, %exit1 ], [ %15, %Flow1 ]
-; IR: call void @llvm.amdgcn.end.cf(i64 %19)
-; IR: %12 = call { i1, i64 } @llvm.amdgcn.if(i1 %11)
+; IR: call void @llvm.amdgcn.end.cf.i64(i64 %19)
+; IR: %12 = call { i1, i64 } @llvm.amdgcn.if.i64(i1 %11)
 ; IR: br i1 %13, label %exit0, label %UnifiedUnreachableBlock
 
 
 ; IR: UnifiedUnreachableBlock:
 ; IR-NEXT: unreachable
 
 
 ; FIXME: Probably should insert an s_endpgm anyway.

 ; IR-LABEL: @multi_exit_region_divergent_ret_uniform_ret(
 ; IR: %divergent.cond0 = icmp slt i32 %tmp16, 2
 ; IR: llvm.amdgcn.if
 ; IR: br i1
 
 ; IR: {{^}}Flow:
 ; IR: %4 = phi i1 [ true, %LeafBlock1 ], [ false, %entry ]
 ; IR: %5 = phi i1 [ %10, %LeafBlock1 ], [ false, %entry ]
-; IR: %6 = call { i1, i64 } @llvm.amdgcn.else(i64 %3)
+; IR: %6 = call { i1, i64 } @llvm.amdgcn.else.i64.i64(i64 %3)
 ; IR: br i1 %7, label %LeafBlock, label %Flow1
 
 ; IR: {{^}}LeafBlock:
 ; IR: %divergent.cond1 = icmp eq i32 %tmp16, 1
 ; IR: %9 = xor i1 %divergent.cond1, true
 ; IR: br label %Flow1
 
 ; IR: LeafBlock1:
 ; IR: %uniform.cond0 = icmp eq i32 %arg3, 2
 ; IR: %10 = xor i1 %uniform.cond0, true
 ; IR: br label %Flow
 
 ; IR: Flow2:
 ; IR: %11 = phi i1 [ false, %exit1 ], [ %15, %Flow1 ]
-; IR: call void @llvm.amdgcn.end.cf(i64 %19)
-; IR: %12 = call { i1, i64 } @llvm.amdgcn.if(i1 %11)
+; IR: call void @llvm.amdgcn.end.cf.i64(i64 %19)
+; IR: %12 = call { i1, i64 } @llvm.amdgcn.if.i64(i1 %11)
 ; IR: br i1 %13, label %exit0, label %UnifiedReturnBlock
 
 ; IR: exit0:
 ; IR: store volatile i32 9, i32 addrspace(1)* undef
 ; IR: br label %UnifiedReturnBlock
 
 ; IR: {{^}}Flow1:
 ; IR: %15 = phi i1 [ %divergent.cond1, %LeafBlock ], [ %4, %Flow ]
 ; IR: %16 = phi i1 [ %9, %LeafBlock ], [ %5, %Flow ]
-; IR: call void @llvm.amdgcn.end.cf(i64 %8)
-; IR: %17 = call { i1, i64 } @llvm.amdgcn.if(i1 %16)
+; IR: call void @llvm.amdgcn.end.cf.i64(i64 %8)
+; IR: %17 = call { i1, i64 } @llvm.amdgcn.if.i64(i1 %16)
 ; IR: %18 = extractvalue { i1, i64 } %17, 0
 ; IR: %19 = extractvalue { i1, i64 } %17, 1
 ; IR: br i1 %18, label %exit1, label %Flow2
 
 ; IR: exit1:
 ; IR: store volatile i32 17, i32 addrspace(3)* undef
 ; IR: br label %Flow2
 
 ; IR: UnifiedReturnBlock:
-; IR: call void @llvm.amdgcn.end.cf(i64 %14)
+; IR: call void @llvm.amdgcn.end.cf.i64(i64 %14)
 ; IR: ret void
 define amdgpu_kernel void @multi_exit_region_divergent_ret_uniform_ret(i32 addrspace(1)* nocapture %arg0, i32 addrspace(1)* nocapture %arg1, i32 addrspace(1)* nocapture %arg2, i32 %arg3) #0 {
 entry:
   %tmp = tail call i32 @llvm.amdgcn.workitem.id.x() #1
   %tmp1 = add i32 0, %tmp
   %tmp2 = zext i32 %tmp1 to i64
   %tmp3 = add i64 0, %tmp2
   %tmp4 = shl i64 %tmp3, 32
exit0:                                     ; preds = %LeafBlock, %LeafBlock1
   ret void
 
 exit1:                                     ; preds = %LeafBlock, %LeafBlock1
   store volatile i32 17, i32 addrspace(3)* undef
   ret void
 }
 
 ; IR-LABEL: @multi_exit_region_uniform_ret_divergent_ret(
-; IR: %1 = call { i1, i64 } @llvm.amdgcn.if(i1 %0)
+; IR: %1 = call { i1, i64 } @llvm.amdgcn.if.i64(i1 %0)
 ; IR: br i1 %2, label %LeafBlock1, label %Flow
 
 ; IR: Flow:
 ; IR: %4 = phi i1 [ true, %LeafBlock1 ], [ false, %entry ]
 ; IR: %5 = phi i1 [ %10, %LeafBlock1 ], [ false, %entry ]
-; IR: %6 = call { i1, i64 } @llvm.amdgcn.else(i64 %3)
+; IR: %6 = call { i1, i64 } @llvm.amdgcn.else.i64.i64(i64 %3)
 
 ; IR: %11 = phi i1 [ false, %exit1 ], [ %15, %Flow1 ]
-; IR: call void @llvm.amdgcn.end.cf(i64 %19)
-; IR: %12 = call { i1, i64 } @llvm.amdgcn.if(i1 %11)
+; IR: call void @llvm.amdgcn.end.cf.i64(i64 %19)
+; IR: %12 = call { i1, i64 } @llvm.amdgcn.if.i64(i1 %11)
 
 define amdgpu_kernel void @multi_exit_region_uniform_ret_divergent_ret(i32 addrspace(1)* nocapture %arg0, i32 addrspace(1)* nocapture %arg1, i32 addrspace(1)* nocapture %arg2, i32 %arg3) #0 {
 entry:
   %tmp = tail call i32 @llvm.amdgcn.workitem.id.x() #1
   %tmp1 = add i32 0, %tmp
   %tmp2 = zext i32 %tmp1 to i64
   %tmp3 = add i64 0, %tmp2
   %tmp4 = shl i64 %tmp3, 32

 exit1:                                     ; preds = %LeafBlock, %LeafBlock1
   store volatile i32 17, i32 addrspace(3)* undef
   ret void
 }
 
 ; IR-LABEL: @multi_divergent_region_exit_ret_ret_return_value(
 ; IR: Flow2:
 ; IR: %11 = phi i1 [ false, %exit1 ], [ %15, %Flow1 ]
-; IR: call void @llvm.amdgcn.end.cf(i64 %19)
+; IR: call void @llvm.amdgcn.end.cf.i64(i64 %19)
 
 ; IR: UnifiedReturnBlock:
 ; IR: %UnifiedRetVal = phi float [ 2.000000e+00, %Flow2 ], [ 1.000000e+00, %exit0 ]
-; IR: call void @llvm.amdgcn.end.cf(i64 %14)
+; IR: call void @llvm.amdgcn.end.cf.i64(i64 %14)
 ; IR: ret float %UnifiedRetVal
 define amdgpu_ps float @multi_divergent_region_exit_ret_ret_return_value(i32 %vgpr) #0 {
 entry:
   %Pivot = icmp slt i32 %vgpr, 2
   br i1 %Pivot, label %LeafBlock, label %LeafBlock1
 
 LeafBlock:                                        ; preds = %entry
   %SwitchLeaf = icmp eq i32 %vgpr, 1
exit0:                                     ; preds = %LeafBlock, %LeafBlock1
   ret float 1.0
 
 exit1:                                     ; preds = %LeafBlock, %LeafBlock1
   store i32 17, i32 addrspace(3)* undef
   ret float 2.0
 }
 
 ; IR-LABEL: @multi_divergent_region_exit_ret_unreachable(
-; IR: %1 = call { i1, i64 } @llvm.amdgcn.if(i1 %0)
+; IR: %1 = call { i1, i64 } @llvm.amdgcn.if.i64(i1 %0)
 
 ; IR: Flow:
 ; IR: %4 = phi i1 [ true, %LeafBlock1 ], [ false, %entry ]
 ; IR: %5 = phi i1 [ %10, %LeafBlock1 ], [ false, %entry ]
-; IR: %6 = call { i1, i64 } @llvm.amdgcn.else(i64 %3)
+; IR: %6 = call { i1, i64 } @llvm.amdgcn.else.i64.i64(i64 %3)
 
 ; IR: Flow2:
 ; IR: %11 = phi i1 [ false, %exit1 ], [ %15, %Flow1 ]
-; IR: call void @llvm.amdgcn.end.cf(i64 %19)
-; IR: %12 = call { i1, i64 } @llvm.amdgcn.if(i1 %11)
+; IR: call void @llvm.amdgcn.end.cf.i64(i64 %19)
+; IR: %12 = call { i1, i64 } @llvm.amdgcn.if.i64(i1 %11)
 ; IR: br i1 %13, label %exit0, label %UnifiedReturnBlock
 
 ; IR: exit0:
 ; IR-NEXT: store volatile i32 17, i32 addrspace(3)* undef
 ; IR-NEXT: br label %UnifiedReturnBlock
 
 ; IR: Flow1:
 ; IR: %15 = phi i1 [ %SwitchLeaf, %LeafBlock ], [ %4, %Flow ]
 ; IR: %16 = phi i1 [ %9, %LeafBlock ], [ %5, %Flow ]
-; IR: call void @llvm.amdgcn.end.cf(i64 %8)
-; IR: %17 = call { i1, i64 } @llvm.amdgcn.if(i1 %16)
+; IR: call void @llvm.amdgcn.end.cf.i64(i64 %8)
+; IR: %17 = call { i1, i64 } @llvm.amdgcn.if.i64(i1 %16)
 ; IR: %18 = extractvalue { i1, i64 } %17, 0
 ; IR: %19 = extractvalue { i1, i64 } %17, 1
 ; IR: br i1 %18, label %exit1, label %Flow2
 
 ; IR: exit1:
 ; IR-NEXT: store volatile i32 9, i32 addrspace(1)* undef
 ; IR-NEXT: call void @llvm.amdgcn.unreachable()
 ; IR-NEXT: br label %Flow2
 
 ; IR: UnifiedReturnBlock:
-; IR-NEXT: call void @llvm.amdgcn.end.cf(i64 %14)
+; IR-NEXT: call void @llvm.amdgcn.end.cf.i64(i64 %14)
 ; IR-NEXT: ret void
 define amdgpu_kernel void @multi_divergent_region_exit_ret_unreachable(i32 addrspace(1)* nocapture %arg0, i32 addrspace(1)* nocapture %arg1, i32 addrspace(1)* nocapture %arg2) #0 {
 entry:
   %tmp = tail call i32 @llvm.amdgcn.workitem.id.x() #1
   %tmp1 = add i32 0, %tmp
   %tmp2 = zext i32 %tmp1 to i64
   %tmp3 = add i64 0, %tmp2
   %tmp4 = shl i64 %tmp3, 32

 ; IR: indirect.exit1:
 ; IR: %load = load volatile i32, i32 addrspace(1)* undef
 ; IR: store volatile i32 %load, i32 addrspace(1)* undef
 ; IR: store volatile i32 9, i32 addrspace(1)* undef
 ; IR: call void @llvm.amdgcn.unreachable()
 ; IR-NEXT: br label %Flow2
 
 ; IR: UnifiedReturnBlock:                               ; preds = %exit0, %Flow2
-; IR-NEXT: call void @llvm.amdgcn.end.cf(i64 %14)
+; IR-NEXT: call void @llvm.amdgcn.end.cf.i64(i64 %14)
 ; IR-NEXT: ret void
 define amdgpu_kernel void @indirect_multi_divergent_region_exit_ret_unreachable(i32 addrspace(1)* nocapture %arg0, i32 addrspace(1)* nocapture %arg1, i32 addrspace(1)* nocapture %arg2) #0 {
 entry:
   %tmp = tail call i32 @llvm.amdgcn.workitem.id.x() #1
   %tmp1 = add i32 0, %tmp
   %tmp2 = zext i32 %tmp1 to i64
   %tmp3 = add i64 0, %tmp2
   %tmp4 = shl i64 %tmp3, 32

 ; IR: %8 = phi i1 [ false, %uniform.ret1 ], [ true, %uniform.multi.exit.region ]
 ; IR: br i1 %8, label %uniform.if, label %Flow2
 
 ; IR: Flow:                                             ; preds = %uniform.then, %uniform.if
 ; IR: %11 = phi i1 [ %10, %uniform.then ], [ %9, %uniform.if ]
 ; IR: br i1 %11, label %uniform.endif, label %uniform.ret0
 
 ; IR: UnifiedReturnBlock:                               ; preds = %Flow3, %Flow2
-; IR-NEXT: call void @llvm.amdgcn.end.cf(i64 %6)
+; IR-NEXT: call void @llvm.amdgcn.end.cf.i64(i64 %6)
 ; IR-NEXT: ret void
 define amdgpu_kernel void @uniform_complex_multi_ret_nest_in_divergent_triangle(i32 %arg0) #0 {
 entry:
   %id.x = tail call i32 @llvm.amdgcn.workitem.id.x()
   %divergent.cond0 = icmp eq i32 %id.x, 0
   br i1 %divergent.cond0, label %uniform.multi.exit.region, label %divergent.ret
 
 uniform.multi.exit.region:

 }
 
 ; IR-LABEL: @multi_divergent_unreachable_exit(
 ; IR: UnifiedUnreachableBlock:
 ; IR-NEXT: call void @llvm.amdgcn.unreachable()
 ; IR-NEXT: br label %UnifiedReturnBlock
 
 ; IR: UnifiedReturnBlock:
-; IR-NEXT: call void @llvm.amdgcn.end.cf(i64
+; IR-NEXT: call void @llvm.amdgcn.end.cf.i64(i64
 ; IR-NEXT: ret void
 define amdgpu_kernel void @multi_divergent_unreachable_exit() #0 {
 bb:
   %tmp = tail call i32 @llvm.amdgcn.workitem.id.x()
   switch i32 %tmp, label %bb3 [
     i32 2, label %bb1
     i32 0, label %bb2
   ]

test/CodeGen/AMDGPU/multilevel-break.ll

 ; RUN: opt -S -mtriple=amdgcn-- -structurizecfg -si-annotate-control-flow < %s | FileCheck -check-prefix=OPT %s
 ; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
 
 ; OPT-LABEL: {{^}}define amdgpu_vs void @multi_else_break(
 ; OPT: main_body:
 ; OPT: LOOP.outer:
 ; OPT: LOOP:
-; OPT:     [[if:%[0-9]+]] = call { i1, i64 } @llvm.amdgcn.if(
+; OPT:     [[if:%[0-9]+]] = call { i1, i64 } @llvm.amdgcn.if.i64(
 ; OPT:     [[if_exec:%[0-9]+]] = extractvalue { i1, i64 } [[if]], 1
 ;
 ; OPT: Flow:
 ;
 ; Ensure two if.break calls, for both the inner and outer loops
 
 ; OPT:        call void @llvm.amdgcn.end.cf
-; OPT-NEXT:   call i64 @llvm.amdgcn.if.break(i1
-; OPT-NEXT:   call i1 @llvm.amdgcn.loop(i64
-; OPT-NEXT:   call i64 @llvm.amdgcn.if.break(i1
+; OPT-NEXT:   call i64 @llvm.amdgcn.if.break.i64.i64(i1
+; OPT-NEXT:   call i1 @llvm.amdgcn.loop.i64(i64
+; OPT-NEXT:   call i64 @llvm.amdgcn.if.break.i64.i64(i1
 ;
 ; OPT: Flow1:
 
 ; GCN-LABEL: {{^}}multi_else_break:
 
 ; GCN: ; %main_body
 ; GCN:      s_mov_b64           [[LEFT_OUTER:s\[[0-9]+:[0-9]+\]]], 0{{$}}
 

test/CodeGen/AMDGPU/nested-loop-conditions.ll

 ; RUN: opt -mtriple=amdgcn-- -S -structurizecfg -si-annotate-control-flow %s | FileCheck -check-prefix=IR %s
 ; RUN: llc -march=amdgcn -mcpu=hawaii -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
 
 ; After structurizing, there are 3 levels of loops. The i1 phi
 ; conditions mutually depend on each other, so it isn't safe to delete
 ; the condition that appears to have no uses until the loop is
 ; completely processed.
 
 
 ; IR-LABEL: @reduced_nested_loop_conditions(
 
 ; IR: bb5:
 ; IR-NEXT: %phi.broken = phi i64 [ %3, %bb10 ], [ 0, %bb ]
 ; IR-NEXT: %tmp6 = phi i32 [ 0, %bb ], [ %tmp11, %bb10 ]
 ; IR-NEXT: %tmp7 = icmp eq i32 %tmp6, 1
-; IR-NEXT: %0 = call { i1, i64 } @llvm.amdgcn.if(i1 %tmp7)
+; IR-NEXT: %0 = call { i1, i64 } @llvm.amdgcn.if.i64(i1 %tmp7)
 ; IR-NEXT: %1 = extractvalue { i1, i64 } %0, 0
 ; IR-NEXT: %2 = extractvalue { i1, i64 } %0, 1
 ; IR-NEXT: br i1 %1, label %bb8, label %Flow
 
 ; IR: bb8:
 ; IR-NEXT: br label %bb13
 
 ; IR: bb10:
 ; IR-NEXT: %tmp11 = phi i32 [ %6, %Flow ]
 ; IR-NEXT: %tmp12 = phi i1 [ %5, %Flow ]
-; IR-NEXT: %3 = call i64 @llvm.amdgcn.if.break(i1 %tmp12, i64 %phi.broken)
-; IR-NEXT: %4 = call i1 @llvm.amdgcn.loop(i64 %3)
+; IR-NEXT: %3 = call i64 @llvm.amdgcn.if.break.i64.i64(i1 %tmp12, i64 %phi.broken)
+; IR-NEXT: %4 = call i1 @llvm.amdgcn.loop.i64(i64 %3)
 ; IR-NEXT: br i1 %4, label %bb23, label %bb5
 
 ; IR: Flow:
 ; IR-NEXT: %5 = phi i1 [ %tmp22, %bb4 ], [ true, %bb5 ]
 ; IR-NEXT: %6 = phi i32 [ %tmp21, %bb4 ], [ undef, %bb5 ]
-; IR-NEXT: call void @llvm.amdgcn.end.cf(i64 %2)
+; IR-NEXT: call void @llvm.amdgcn.end.cf.i64(i64 %2)
 ; IR-NEXT: br label %bb10
 
 ; IR: bb13:
 ; IR-NEXT: %tmp14 = phi i1 [ %tmp22, %bb3 ], [ true, %bb8 ]
 ; IR-NEXT: %tmp15 = bitcast i64 %tmp2 to <2 x i32>
 ; IR-NEXT: br i1 %tmp14, label %bb16, label %bb20
 
 ; IR: bb16:
 ; IR-NEXT: %tmp17 = extractelement <2 x i32> %tmp15, i64 1
 ; IR-NEXT: %tmp18 = getelementptr inbounds i32, i32 addrspace(3)* undef, i32 %tmp17
 ; IR-NEXT: %tmp19 = load volatile i32, i32 addrspace(3)* %tmp18
 ; IR-NEXT: br label %bb20
 
 ; IR: bb20:
 ; IR-NEXT: %tmp21 = phi i32 [ %tmp19, %bb16 ], [ 0, %bb13 ]
 ; IR-NEXT: %tmp22 = phi i1 [ false, %bb16 ], [ %tmp14, %bb13 ]
 ; IR-NEXT: br label %bb9
 
 ; IR: bb23:
-; IR-NEXT: call void @llvm.amdgcn.end.cf(i64 %3)
+; IR-NEXT: call void @llvm.amdgcn.end.cf.i64(i64 %3)
 ; IR-NEXT: ret void
 
 ; GCN-LABEL: {{^}}reduced_nested_loop_conditions:
 
 ; GCN: s_cmp_lg_u32 s{{[0-9]+}}, 1
 ; GCN-NEXT: s_cbranch_scc0
 
 ; FIXME: Should fold to unconditional branch?

 bb23:                                             ; preds = %bb10
   ret void
 }
 
 ; Earlier version of above, before a run of the structurizer.
 ; IR-LABEL: @nested_loop_conditions(
 
 ; IR: Flow3:
-; IR-NEXT: call void @llvm.amdgcn.end.cf(i64 %21)
-; IR-NEXT: %0 = call { i1, i64 } @llvm.amdgcn.if(i1 %14)
+; IR-NEXT: call void @llvm.amdgcn.end.cf.i64(i64 %21)
+; IR-NEXT: %0 = call { i1, i64 } @llvm.amdgcn.if.i64(i1 %14)
 ; IR-NEXT: %1 = extractvalue { i1, i64 } %0, 0
 ; IR-NEXT: %2 = extractvalue { i1, i64 } %0, 1
 ; IR-NEXT: br i1 %1, label %bb4.bb13_crit_edge, label %Flow4
 
 ; IR: Flow4:
 ; IR-NEXT: %3 = phi i1 [ true, %bb4.bb13_crit_edge ], [ false, %Flow3 ]
-; IR-NEXT: call void @llvm.amdgcn.end.cf(i64 %2)
+; IR-NEXT: call void @llvm.amdgcn.end.cf.i64(i64 %2)
 ; IR-NEXT: br label %Flow
 
 ; IR: Flow:
 ; IR-NEXT:  %4 = phi i1 [ %3, %Flow4 ], [ true, %bb ]
-; IR-NEXT:  %5 = call { i1, i64 } @llvm.amdgcn.if(i1 %4)
+; IR-NEXT:  %5 = call { i1, i64 } @llvm.amdgcn.if.i64(i1 %4)
 ; IR-NEXT:  %6 = extractvalue { i1, i64 } %5, 0
 ; IR-NEXT:  %7 = extractvalue { i1, i64 } %5, 1
 ; IR-NEXT:  br i1 %6, label %bb13, label %bb31
 
 ; IR: bb14:
 ; IR: %tmp15 = icmp eq i32 %tmp1037, 1
-; IR-NEXT: %8 = call { i1, i64 } @llvm.amdgcn.if(i1 %tmp15)
+; IR-NEXT: %8 = call { i1, i64 } @llvm.amdgcn.if.i64(i1 %tmp15)
 
 ; IR: Flow1:
 ; IR-NEXT: %11 = phi <4 x i32> [ %tmp9, %bb21 ], [ undef, %bb14 ]
 ; IR-NEXT: %12 = phi i32 [ %tmp10, %bb21 ], [ undef, %bb14 ]
 ; IR-NEXT: %13 = phi i1 [ %18, %bb21 ], [ true, %bb14 ]
 ; IR-NEXT: %14 = phi i1 [ %18, %bb21 ], [ false, %bb14 ]
 ; IR-NEXT: %15 = phi i1 [ false, %bb21 ], [ true, %bb14 ]
-; IR-NEXT: call void @llvm.amdgcn.end.cf(i64 %10)
-; IR-NEXT: %16 = call i64 @llvm.amdgcn.if.break(i1 %13, i64 %phi.broken)
-; IR-NEXT: %17 = call i1 @llvm.amdgcn.loop(i64 %16)
+; IR-NEXT: call void @llvm.amdgcn.end.cf.i64(i64 %10)
+; IR-NEXT: %16 = call i64 @llvm.amdgcn.if.break.i64.i64(i1 %13, i64 %phi.broken)
+; IR-NEXT: %17 = call i1 @llvm.amdgcn.loop.i64(i64 %16)
 ; IR-NEXT: br i1 %17, label %Flow2, label %bb14
 
 ; IR: bb21:
 ; IR: %tmp12 = icmp slt i32 %tmp11, 9
 ; IR-NEXT: %18 = xor i1 %tmp12, true
 ; IR-NEXT: br label %Flow1
 
 ; IR: Flow2:
-; IR-NEXT: call void @llvm.amdgcn.end.cf(i64 %16)
-; IR-NEXT: %19 = call { i1, i64 } @llvm.amdgcn.if(i1 %15)
+; IR-NEXT: call void @llvm.amdgcn.end.cf.i64(i64 %16)
+; IR-NEXT: %19 = call { i1, i64 } @llvm.amdgcn.if.i64(i1 %15)
 ; IR-NEXT: %20 = extractvalue { i1, i64 } %19, 0
 ; IR-NEXT: %21 = extractvalue { i1, i64 } %19, 1
 ; IR-NEXT: br i1 %20, label %bb31.loopexit, label %Flow3
 
 ; IR: bb31:
-; IR-NEXT: call void @llvm.amdgcn.end.cf(i64 %7)
+; IR-NEXT: call void @llvm.amdgcn.end.cf.i64(i64 %7)
 ; IR-NEXT: store volatile i32 0, i32 addrspace(1)* undef
 ; IR-NEXT: ret void
 
 
 ; GCN-LABEL: {{^}}nested_loop_conditions:
 
 ; GCN: v_cmp_lt_i32_e32 vcc, 8, v
 ; GCN: s_and_b64 vcc, exec, vcc

test/CodeGen/AMDGPU/si-annotate-cf-unreachable.ll

 ; RUN: opt -mtriple=amdgcn-- -S -structurizecfg -si-annotate-control-flow %s | FileCheck -check-prefix=OPT %s
 ; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
 
 
 ; OPT-LABEL: @annotate_unreachable(
-; OPT: call { i1, i64 } @llvm.amdgcn.if(
-; OPT-NOT: call void @llvm.amdgcn.end.cf(
+; OPT: call { i1, i64 } @llvm.amdgcn.if.i64(
+; OPT-NOT: call void @llvm.amdgcn.end.cf
 
 
 ; GCN-LABEL: {{^}}annotate_unreachable:
 ; GCN: s_and_saveexec_b64
 ; GCN-NOT: s_endpgm
 ; GCN: .Lfunc_end0
 define amdgpu_kernel void @annotate_unreachable(<4 x float> addrspace(1)* noalias nocapture readonly %arg) #0 {
 bb:

test/CodeGen/AMDGPU/si-annotatecfg-multiple-backedges.ll

 ; OPT-NEXT:    [[TMP2:%.*]] = shl nsw i32 [[ARG:%.*]], 1
 ; OPT-NEXT:    br label [[LOOP:%.*]]
 ; OPT:       loop:
 ; OPT-NEXT:    [[PHI_BROKEN1:%.*]] = phi i64 [ [[TMP7:%.*]], [[LOOP_END:%.*]] ], [ [[PHI_BROKEN1]], [[LOOP]] ], [ 0, [[ENTRY:%.*]] ]
 ; OPT-NEXT:    [[PHI_BROKEN:%.*]] = phi i64 [ 0, [[LOOP_END]] ], [ [[TMP0:%.*]], [[LOOP]] ], [ 0, [[ENTRY]] ]
 ; OPT-NEXT:    [[TMP4:%.*]] = phi i32 [ 0, [[ENTRY]] ], [ [[TMP5:%.*]], [[LOOP]] ], [ 0, [[LOOP_END]] ]
 ; OPT-NEXT:    [[TMP5]] = add nsw i32 [[TMP4]], [[TMP]]
 ; OPT-NEXT:    [[TMP6:%.*]] = icmp slt i32 [[ARG]], [[TMP5]]
-; OPT-NEXT:    [[TMP0]] = call i64 @llvm.amdgcn.if.break(i1 [[TMP6]], i64 [[PHI_BROKEN]])
-; OPT-NEXT:    [[TMP1:%.*]] = call i1 @llvm.amdgcn.loop(i64 [[TMP0]])
+; OPT-NEXT:    [[TMP0]] = call i64 @llvm.amdgcn.if.break.i64.i64(i1 [[TMP6]], i64 [[PHI_BROKEN]])
+; OPT-NEXT:    [[TMP1:%.*]] = call i1 @llvm.amdgcn.loop.i64(i64 [[TMP0]])
 ; OPT-NEXT:    br i1 [[TMP1]], label [[LOOP_END]], label [[LOOP]]
 ; OPT:       loop_end:
-; OPT-NEXT:    call void @llvm.amdgcn.end.cf(i64 [[TMP0]])
+; OPT-NEXT:    call void @llvm.amdgcn.end.cf.i64(i64 [[TMP0]])
 ; OPT-NEXT:    [[EXIT:%.*]] = icmp sgt i32 [[TMP5]], [[TMP2]]
-; OPT-NEXT:    [[TMP7]] = call i64 @llvm.amdgcn.if.break(i1 [[EXIT]], i64 [[PHI_BROKEN1]])
-; OPT-NEXT:    [[TMP3:%.*]] = call i1 @llvm.amdgcn.loop(i64 [[TMP7]])
+; OPT-NEXT:    [[TMP7]] = call i64 @llvm.amdgcn.if.break.i64.i64(i1 [[EXIT]], i64 [[PHI_BROKEN1]])
+; OPT-NEXT:    [[TMP3:%.*]] = call i1 @llvm.amdgcn.loop.i64(i64 [[TMP7]])
 ; OPT-NEXT:    br i1 [[TMP3]], label [[LOOP_EXIT:%.*]], label [[LOOP]]
 ; OPT:       loop_exit:
-; OPT-NEXT:    call void @llvm.amdgcn.end.cf(i64 [[TMP7]])
+; OPT-NEXT:    call void @llvm.amdgcn.end.cf.i64(i64 [[TMP7]])
 ; OPT-NEXT:    [[TMP12:%.*]] = zext i32 [[TMP]] to i64
 ; OPT-NEXT:    [[TMP13:%.*]] = getelementptr inbounds i32, i32* [[ARG1:%.*]], i64 [[TMP12]]
 ; OPT-NEXT:    [[TMP14:%.*]] = addrspacecast i32* [[TMP13]] to i32 addrspace(1)*
 ; OPT-NEXT:    store i32 [[TMP5]], i32 addrspace(1)* [[TMP14]], align 4
 ; OPT-NEXT:    ret void
 ;
 entry:
   %tmp = tail call i32 @llvm.amdgcn.workitem.id.x()

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

;
   %med = call float @llvm.amdgcn.fmed3.f32(float 0.0, float 1.0, float undef)
   ret float %med
 }
 
 ; --------------------------------------------------------------------
 ; llvm.amdgcn.icmp
 ; --------------------------------------------------------------------
 
-declare i64 @llvm.amdgcn.icmp.i32(i32, i32, i32 immarg) nounwind readnone convergent
-declare i64 @llvm.amdgcn.icmp.i64(i64, i64, i32 immarg) nounwind readnone convergent
-declare i64 @llvm.amdgcn.icmp.i1(i1, i1, i32 immarg) nounwind readnone convergent
+declare i64 @llvm.amdgcn.icmp.i64.i32(i32, i32, i32 immarg) nounwind readnone convergent
+declare i64 @llvm.amdgcn.icmp.i64.i64(i64, i64, i32 immarg) nounwind readnone convergent
+declare i64 @llvm.amdgcn.icmp.i64.i1(i1, i1, i32 immarg) nounwind readnone convergent
 
 define i64 @invalid_icmp_code(i32 %a, i32 %b) {
 ; CHECK-LABEL: @invalid_icmp_code(
-; CHECK-NEXT:    [[UNDER:%.*]] = call i64 @llvm.amdgcn.icmp.i32(i32 [[A:%.*]], i32 [[B:%.*]], i32 31)
-; CHECK-NEXT:    [[OVER:%.*]] = call i64 @llvm.amdgcn.icmp.i32(i32 [[A]], i32 [[B]], i32 42)
+; CHECK-NEXT:    [[UNDER:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i32(i32 [[A:%.*]], i32 [[B:%.*]], i32 31)
+; CHECK-NEXT:    [[OVER:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i32(i32 [[A]], i32 [[B]], i32 42)
 ; CHECK-NEXT:    [[OR:%.*]] = or i64 [[UNDER]], [[OVER]]
 ; CHECK-NEXT:    ret i64 [[OR]]
 ;
-  %under = call i64 @llvm.amdgcn.icmp.i32(i32 %a, i32 %b, i32 31)
-  %over = call i64 @llvm.amdgcn.icmp.i32(i32 %a, i32 %b, i32 42)
+  %under = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %a, i32 %b, i32 31)
+  %over = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %a, i32 %b, i32 42)
   %or = or i64 %under, %over
   ret i64 %or
 }
 
 define i64 @icmp_constant_inputs_false() {
 ; CHECK-LABEL: @icmp_constant_inputs_false(
 ; CHECK-NEXT:    ret i64 0
 ;
-  %result = call i64 @llvm.amdgcn.icmp.i32(i32 9, i32 8, i32 32)
+  %result = call i64 @llvm.amdgcn.icmp.i64.i32(i32 9, i32 8, i32 32)
   ret i64 %result
 }
 
 define i64 @icmp_constant_inputs_true() {
 ; CHECK-LABEL: @icmp_constant_inputs_true(
 ; CHECK-NEXT:    [[RESULT:%.*]] = call i64 @llvm.read_register.i64(metadata !0) #5
 ; CHECK-NEXT:    ret i64 [[RESULT]]
 ;
-  %result = call i64 @llvm.amdgcn.icmp.i32(i32 9, i32 8, i32 34)
+  %result = call i64 @llvm.amdgcn.icmp.i64.i32(i32 9, i32 8, i32 34)
   ret i64 %result
 }
 
 define i64 @icmp_constant_to_rhs_slt(i32 %x) {
 ; CHECK-LABEL: @icmp_constant_to_rhs_slt(
-; CHECK-NEXT:    [[RESULT:%.*]] = call i64 @llvm.amdgcn.icmp.i32(i32 [[X:%.*]], i32 9, i32 38)
+; CHECK-NEXT:    [[RESULT:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i32(i32 [[X:%.*]], i32 9, i32 38)
 ; CHECK-NEXT:    ret i64 [[RESULT]]
 ;
-  %result = call i64 @llvm.amdgcn.icmp.i32(i32 9, i32 %x, i32 40)
+  %result = call i64 @llvm.amdgcn.icmp.i64.i32(i32 9, i32 %x, i32 40)
   ret i64 %result
 }
 
 define i64 @fold_icmp_ne_0_zext_icmp_eq_i32(i32 %a, i32 %b) {
 ; CHECK-LABEL: @fold_icmp_ne_0_zext_icmp_eq_i32(
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i32(i32 [[A:%.*]], i32 [[B:%.*]], i32 32)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i32(i32 [[A:%.*]], i32 [[B:%.*]], i32 32)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp eq i32 %a, %b
   %zext.cmp = zext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %zext.cmp, i32 0, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %zext.cmp, i32 0, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_ne_0_zext_icmp_ne_i32(i32 %a, i32 %b) {
 ; CHECK-LABEL: @fold_icmp_ne_0_zext_icmp_ne_i32(
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i32(i32 [[A:%.*]], i32 [[B:%.*]], i32 33)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i32(i32 [[A:%.*]], i32 [[B:%.*]], i32 33)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp ne i32 %a, %b
   %zext.cmp = zext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %zext.cmp, i32 0, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %zext.cmp, i32 0, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_ne_0_zext_icmp_sle_i32(i32 %a, i32 %b) {
 ; CHECK-LABEL: @fold_icmp_ne_0_zext_icmp_sle_i32(
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i32(i32 [[A:%.*]], i32 [[B:%.*]], i32 41)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i32(i32 [[A:%.*]], i32 [[B:%.*]], i32 41)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp sle i32 %a, %b
   %zext.cmp = zext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %zext.cmp, i32 0, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %zext.cmp, i32 0, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_ne_0_zext_icmp_ugt_i64(i64 %a, i64 %b) {
 ; CHECK-LABEL: @fold_icmp_ne_0_zext_icmp_ugt_i64(
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64(i64 [[A:%.*]], i64 [[B:%.*]], i32 34)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i64(i64 [[A:%.*]], i64 [[B:%.*]], i32 34)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp ugt i64 %a, %b
   %zext.cmp = zext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %zext.cmp, i32 0, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %zext.cmp, i32 0, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_ne_0_zext_icmp_ult_swap_i64(i64 %a, i64 %b) {
 ; CHECK-LABEL: @fold_icmp_ne_0_zext_icmp_ult_swap_i64(
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64(i64 [[A:%.*]], i64 [[B:%.*]], i32 34)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i64(i64 [[A:%.*]], i64 [[B:%.*]], i32 34)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp ugt i64 %a, %b
   %zext.cmp = zext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 0, i32 %zext.cmp, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 0, i32 %zext.cmp, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_ne_0_zext_fcmp_oeq_f32(float %a, float %b) {
 ; CHECK-LABEL: @fold_icmp_ne_0_zext_fcmp_oeq_f32(
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.fcmp.f32(float [[A:%.*]], float [[B:%.*]], i32 1)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.fcmp.i64.f32(float [[A:%.*]], float [[B:%.*]], i32 1)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = fcmp oeq float %a, %b
   %zext.cmp = zext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %zext.cmp, i32 0, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %zext.cmp, i32 0, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_ne_0_zext_fcmp_une_f32(float %a, float %b) {
 ; CHECK-LABEL: @fold_icmp_ne_0_zext_fcmp_une_f32(
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.fcmp.f32(float [[A:%.*]], float [[B:%.*]], i32 14)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.fcmp.i64.f32(float [[A:%.*]], float [[B:%.*]], i32 14)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = fcmp une float %a, %b
   %zext.cmp = zext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %zext.cmp, i32 0, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %zext.cmp, i32 0, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_ne_0_zext_fcmp_olt_f64(double %a, double %b) {
 ; CHECK-LABEL: @fold_icmp_ne_0_zext_fcmp_olt_f64(
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.fcmp.f64(double [[A:%.*]], double [[B:%.*]], i32 4)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.fcmp.i64.f64(double [[A:%.*]], double [[B:%.*]], i32 4)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = fcmp olt double %a, %b
   %zext.cmp = zext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %zext.cmp, i32 0, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %zext.cmp, i32 0, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_sext_icmp_ne_0_i32(i32 %a, i32 %b) {
 ; CHECK-LABEL: @fold_icmp_sext_icmp_ne_0_i32(
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i32(i32 [[A:%.*]], i32 [[B:%.*]], i32 32)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i32(i32 [[A:%.*]], i32 [[B:%.*]], i32 32)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp eq i32 %a, %b
   %sext.cmp = sext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %sext.cmp, i32 0, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %sext.cmp, i32 0, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_eq_0_zext_icmp_eq_i32(i32 %a, i32 %b) {
 ; CHECK-LABEL: @fold_icmp_eq_0_zext_icmp_eq_i32(
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i32(i32 [[A:%.*]], i32 [[B:%.*]], i32 33)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i32(i32 [[A:%.*]], i32 [[B:%.*]], i32 33)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp eq i32 %a, %b
   %zext.cmp = zext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %zext.cmp, i32 0, i32 32)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %zext.cmp, i32 0, i32 32)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_eq_0_zext_icmp_slt_i32(i32 %a, i32 %b) {
 ; CHECK-LABEL: @fold_icmp_eq_0_zext_icmp_slt_i32(
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i32(i32 [[A:%.*]], i32 [[B:%.*]], i32 39)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i32(i32 [[A:%.*]], i32 [[B:%.*]], i32 39)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp slt i32 %a, %b
   %zext.cmp = zext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %zext.cmp, i32 0, i32 32)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %zext.cmp, i32 0, i32 32)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_eq_0_zext_fcmp_oeq_f32(float %a, float %b) {
 ; CHECK-LABEL: @fold_icmp_eq_0_zext_fcmp_oeq_f32(
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.fcmp.f32(float [[A:%.*]], float [[B:%.*]], i32 14)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.fcmp.i64.f32(float [[A:%.*]], float [[B:%.*]], i32 14)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = fcmp oeq float %a, %b
   %zext.cmp = zext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %zext.cmp, i32 0, i32 32)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %zext.cmp, i32 0, i32 32)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_eq_0_zext_fcmp_ule_f32(float %a, float %b) {
 ; CHECK-LABEL: @fold_icmp_eq_0_zext_fcmp_ule_f32(
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.fcmp.f32(float [[A:%.*]], float [[B:%.*]], i32 2)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.fcmp.i64.f32(float [[A:%.*]], float [[B:%.*]], i32 2)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = fcmp ule float %a, %b
   %zext.cmp = zext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %zext.cmp, i32 0, i32 32)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %zext.cmp, i32 0, i32 32)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_eq_0_zext_fcmp_ogt_f32(float %a, float %b) {
 ; CHECK-LABEL: @fold_icmp_eq_0_zext_fcmp_ogt_f32(
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.fcmp.f32(float [[A:%.*]], float [[B:%.*]], i32 13)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.fcmp.i64.f32(float [[A:%.*]], float [[B:%.*]], i32 13)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = fcmp ogt float %a, %b
   %zext.cmp = zext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %zext.cmp, i32 0, i32 32)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %zext.cmp, i32 0, i32 32)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_zext_icmp_eq_1_i32(i32 %a, i32 %b) {
 ; CHECK-LABEL: @fold_icmp_zext_icmp_eq_1_i32(
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i32(i32 [[A:%.*]], i32 [[B:%.*]], i32 32)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i32(i32 [[A:%.*]], i32 [[B:%.*]], i32 32)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp eq i32 %a, %b
   %zext.cmp = zext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %zext.cmp, i32 1, i32 32)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %zext.cmp, i32 1, i32 32)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_zext_argi1_eq_1_i32(i1 %cond) {
 ; CHECK-LABEL: @fold_icmp_zext_argi1_eq_1_i32(
 ; CHECK-NEXT:    [[ZEXT_COND:%.*]] = zext i1 [[COND:%.*]] to i32
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i32(i32 [[ZEXT_COND]], i32 0, i32 33)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i32(i32 [[ZEXT_COND]], i32 0, i32 33)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %zext.cond = zext i1 %cond to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %zext.cond, i32 1, i32 32)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %zext.cond, i32 1, i32 32)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_zext_argi1_eq_neg1_i32(i1 %cond) {
 ; CHECK-LABEL: @fold_icmp_zext_argi1_eq_neg1_i32(
 ; CHECK-NEXT:    [[ZEXT_COND:%.*]] = zext i1 [[COND:%.*]] to i32
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i32(i32 [[ZEXT_COND]], i32 -1, i32 32)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i32(i32 [[ZEXT_COND]], i32 -1, i32 32)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %zext.cond = zext i1 %cond to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %zext.cond, i32 -1, i32 32)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %zext.cond, i32 -1, i32 32)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_sext_argi1_eq_1_i32(i1 %cond) {
 ; CHECK-LABEL: @fold_icmp_sext_argi1_eq_1_i32(
 ; CHECK-NEXT:    [[SEXT_COND:%.*]] = sext i1 [[COND:%.*]] to i32
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i32(i32 [[SEXT_COND]], i32 1, i32 32)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i32(i32 [[SEXT_COND]], i32 1, i32 32)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %sext.cond = sext i1 %cond to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %sext.cond, i32 1, i32 32)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %sext.cond, i32 1, i32 32)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_sext_argi1_eq_neg1_i32(i1 %cond) {
 ; CHECK-LABEL: @fold_icmp_sext_argi1_eq_neg1_i32(
 ; CHECK-NEXT:    [[SEXT_COND:%.*]] = sext i1 [[COND:%.*]] to i32
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i32(i32 [[SEXT_COND]], i32 0, i32 33)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i32(i32 [[SEXT_COND]], i32 0, i32 33)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %sext.cond = sext i1 %cond to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %sext.cond, i32 -1, i32 32)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %sext.cond, i32 -1, i32 32)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_sext_argi1_eq_neg1_i64(i1 %cond) {
 ; CHECK-LABEL: @fold_icmp_sext_argi1_eq_neg1_i64(
 ; CHECK-NEXT:    [[SEXT_COND:%.*]] = sext i1 [[COND:%.*]] to i64
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64(i64 [[SEXT_COND]], i64 0, i32 33)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i64(i64 [[SEXT_COND]], i64 0, i32 33)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %sext.cond = sext i1 %cond to i64
-  %mask = call i64 @llvm.amdgcn.icmp.i64(i64 %sext.cond, i64 -1, i32 32)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i64(i64 %sext.cond, i64 -1, i32 32)
   ret i64 %mask
 }
 
 ; TODO: Should be able to fold to false
 define i64 @fold_icmp_sext_icmp_eq_1_i32(i32 %a, i32 %b) {
 ; CHECK-LABEL: @fold_icmp_sext_icmp_eq_1_i32(
 ; CHECK-NEXT:    [[CMP:%.*]] = icmp eq i32 [[A:%.*]], [[B:%.*]]
 ; CHECK-NEXT:    [[SEXT_CMP:%.*]] = sext i1 [[CMP]] to i32
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i32(i32 [[SEXT_CMP]], i32 1, i32 32)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i32(i32 [[SEXT_CMP]], i32 1, i32 32)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp eq i32 %a, %b
   %sext.cmp = sext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %sext.cmp, i32 1, i32 32)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %sext.cmp, i32 1, i32 32)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_sext_icmp_eq_neg1_i32(i32 %a, i32 %b) {
 ; CHECK-LABEL: @fold_icmp_sext_icmp_eq_neg1_i32(
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i32(i32 [[A:%.*]], i32 [[B:%.*]], i32 32)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i32(i32 [[A:%.*]], i32 [[B:%.*]], i32 32)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp eq i32 %a, %b
   %sext.cmp = sext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %sext.cmp, i32 -1, i32 32)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %sext.cmp, i32 -1, i32 32)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_sext_icmp_sge_neg1_i32(i32 %a, i32 %b) {
 ; CHECK-LABEL: @fold_icmp_sext_icmp_sge_neg1_i32(
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i32(i32 [[A:%.*]], i32 [[B:%.*]], i32 39)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i32(i32 [[A:%.*]], i32 [[B:%.*]], i32 39)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp sge i32 %a, %b
   %sext.cmp = sext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %sext.cmp, i32 -1, i32 32)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %sext.cmp, i32 -1, i32 32)
   ret i64 %mask
 }
 
 define i64 @fold_not_icmp_ne_0_zext_icmp_sle_i32(i32 %a, i32 %b) {
 ; CHECK-LABEL: @fold_not_icmp_ne_0_zext_icmp_sle_i32(
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i32(i32 [[A:%.*]], i32 [[B:%.*]], i32 38)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i32(i32 [[A:%.*]], i32 [[B:%.*]], i32 38)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp sle i32 %a, %b
   %not = xor i1 %cmp, true
   %zext.cmp = zext i1 %not to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %zext.cmp, i32 0, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %zext.cmp, i32 0, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_ne_0_zext_icmp_eq_i4(i4 %a, i4 %b) {
 ; CHECK-LABEL: @fold_icmp_ne_0_zext_icmp_eq_i4(
 ; CHECK-NEXT:    [[TMP1:%.*]] = zext i4 [[A:%.*]] to i16
 ; CHECK-NEXT:    [[TMP2:%.*]] = zext i4 [[B:%.*]] to i16
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i16(i16 [[TMP1]], i16 [[TMP2]], i32 32)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i16(i16 [[TMP1]], i16 [[TMP2]], i32 32)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp eq i4 %a, %b
   %zext.cmp = zext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %zext.cmp, i32 0, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %zext.cmp, i32 0, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_ne_0_zext_icmp_eq_i8(i8 %a, i8 %b) {
 ; CHECK-LABEL: @fold_icmp_ne_0_zext_icmp_eq_i8(
 ; CHECK-NEXT:    [[TMP1:%.*]] = zext i8 [[A:%.*]] to i16
 ; CHECK-NEXT:    [[TMP2:%.*]] = zext i8 [[B:%.*]] to i16
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i16(i16 [[TMP1]], i16 [[TMP2]], i32 32)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i16(i16 [[TMP1]], i16 [[TMP2]], i32 32)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp eq i8 %a, %b
   %zext.cmp = zext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %zext.cmp, i32 0, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %zext.cmp, i32 0, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_ne_0_zext_icmp_eq_i16(i16 %a, i16 %b) {
 ; CHECK-LABEL: @fold_icmp_ne_0_zext_icmp_eq_i16(
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i16(i16 [[A:%.*]], i16 [[B:%.*]], i32 32)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i16(i16 [[A:%.*]], i16 [[B:%.*]], i32 32)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp eq i16 %a, %b
   %zext.cmp = zext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %zext.cmp, i32 0, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %zext.cmp, i32 0, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_ne_0_zext_icmp_eq_i36(i36 %a, i36 %b) {
 ; CHECK-LABEL: @fold_icmp_ne_0_zext_icmp_eq_i36(
 ; CHECK-NEXT:    [[TMP1:%.*]] = zext i36 [[A:%.*]] to i64
 ; CHECK-NEXT:    [[TMP2:%.*]] = zext i36 [[B:%.*]] to i64
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64(i64 [[TMP1]], i64 [[TMP2]], i32 32)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i64(i64 [[TMP1]], i64 [[TMP2]], i32 32)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp eq i36 %a, %b
   %zext.cmp = zext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %zext.cmp, i32 0, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %zext.cmp, i32 0, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_ne_0_zext_icmp_eq_i128(i128 %a, i128 %b) {
 ; CHECK-LABEL: @fold_icmp_ne_0_zext_icmp_eq_i128(
 ; CHECK-NEXT:    [[CMP:%.*]] = icmp eq i128 [[A:%.*]], [[B:%.*]]
 ; CHECK-NEXT:    [[ZEXT_CMP:%.*]] = zext i1 [[CMP]] to i32
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i32(i32 [[ZEXT_CMP]], i32 0, i32 33)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i32(i32 [[ZEXT_CMP]], i32 0, i32 33)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp eq i128 %a, %b
   %zext.cmp = zext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %zext.cmp, i32 0, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %zext.cmp, i32 0, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_ne_0_zext_fcmp_oeq_f16(half %a, half %b) {
 ; CHECK-LABEL: @fold_icmp_ne_0_zext_fcmp_oeq_f16(
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.fcmp.f16(half [[A:%.*]], half [[B:%.*]], i32 1)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.fcmp.i64.f16(half [[A:%.*]], half [[B:%.*]], i32 1)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = fcmp oeq half %a, %b
   %zext.cmp = zext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %zext.cmp, i32 0, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %zext.cmp, i32 0, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_ne_0_zext_fcmp_oeq_f128(fp128 %a, fp128 %b) {
-;
 ; CHECK-LABEL: @fold_icmp_ne_0_zext_fcmp_oeq_f128(
 ; CHECK-NEXT:    [[CMP:%.*]] = fcmp oeq fp128 [[A:%.*]], [[B:%.*]]
 ; CHECK-NEXT:    [[ZEXT_CMP:%.*]] = zext i1 [[CMP]] to i32
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i32(i32 [[ZEXT_CMP]], i32 0, i32 33)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i32(i32 [[ZEXT_CMP]], i32 0, i32 33)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = fcmp oeq fp128 %a, %b
   %zext.cmp = zext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %zext.cmp, i32 0, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %zext.cmp, i32 0, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_ne_0_zext_icmp_slt_i4(i4 %a, i4 %b) {
 ; CHECK-LABEL: @fold_icmp_ne_0_zext_icmp_slt_i4(
 ; CHECK-NEXT:    [[TMP1:%.*]] = sext i4 [[A:%.*]] to i16
 ; CHECK-NEXT:    [[TMP2:%.*]] = sext i4 [[B:%.*]] to i16
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i16(i16 [[TMP1]], i16 [[TMP2]], i32 40)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i16(i16 [[TMP1]], i16 [[TMP2]], i32 40)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp slt i4 %a, %b
   %zext.cmp = zext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %zext.cmp, i32 0, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %zext.cmp, i32 0, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_ne_0_zext_icmp_slt_i8(i8 %a, i8 %b) {
 ; CHECK-LABEL: @fold_icmp_ne_0_zext_icmp_slt_i8(
 ; CHECK-NEXT:    [[TMP1:%.*]] = sext i8 [[A:%.*]] to i16
 ; CHECK-NEXT:    [[TMP2:%.*]] = sext i8 [[B:%.*]] to i16
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i16(i16 [[TMP1]], i16 [[TMP2]], i32 40)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i16(i16 [[TMP1]], i16 [[TMP2]], i32 40)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp slt i8 %a, %b
   %zext.cmp = zext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %zext.cmp, i32 0, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %zext.cmp, i32 0, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_ne_0_zext_icmp_slt_i16(i16 %a, i16 %b) {
 ; CHECK-LABEL: @fold_icmp_ne_0_zext_icmp_slt_i16(
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i16(i16 [[A:%.*]], i16 [[B:%.*]], i32 40)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i16(i16 [[A:%.*]], i16 [[B:%.*]], i32 40)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp slt i16 %a, %b
   %zext.cmp = zext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %zext.cmp, i32 0, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %zext.cmp, i32 0, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_ne_0_zext_icmp_ult_i4(i4 %a, i4 %b) {
 ; CHECK-LABEL: @fold_icmp_ne_0_zext_icmp_ult_i4(
 ; CHECK-NEXT:    [[TMP1:%.*]] = zext i4 [[A:%.*]] to i16
 ; CHECK-NEXT:    [[TMP2:%.*]] = zext i4 [[B:%.*]] to i16
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i16(i16 [[TMP1]], i16 [[TMP2]], i32 36)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i16(i16 [[TMP1]], i16 [[TMP2]], i32 36)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp ult i4 %a, %b
   %zext.cmp = zext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %zext.cmp, i32 0, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %zext.cmp, i32 0, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_ne_0_zext_icmp_ult_i8(i8 %a, i8 %b) {
 ; CHECK-LABEL: @fold_icmp_ne_0_zext_icmp_ult_i8(
 ; CHECK-NEXT:    [[TMP1:%.*]] = zext i8 [[A:%.*]] to i16
 ; CHECK-NEXT:    [[TMP2:%.*]] = zext i8 [[B:%.*]] to i16
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i16(i16 [[TMP1]], i16 [[TMP2]], i32 36)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i16(i16 [[TMP1]], i16 [[TMP2]], i32 36)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp ult i8 %a, %b
   %zext.cmp = zext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %zext.cmp, i32 0, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %zext.cmp, i32 0, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_ne_0_zext_icmp_ult_i16(i16 %a, i16 %b) {
 ; CHECK-LABEL: @fold_icmp_ne_0_zext_icmp_ult_i16(
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i16(i16 [[A:%.*]], i16 [[B:%.*]], i32 36)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i16(i16 [[A:%.*]], i16 [[B:%.*]], i32 36)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp ult i16 %a, %b
   %zext.cmp = zext i1 %cmp to i32
-  %mask = call i64 @llvm.amdgcn.icmp.i32(i32 %zext.cmp, i32 0, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %zext.cmp, i32 0, i32 33)
   ret i64 %mask
 }
 
 ; 1-bit NE comparisons
 
 define i64 @fold_icmp_i1_ne_0_icmp_eq_i1(i32 %a, i32 %b) {
 ; CHECK-LABEL: @fold_icmp_i1_ne_0_icmp_eq_i1(
 ; CHECK-NEXT:    [[CMP:%.*]] = icmp eq i32 [[A:%.*]], [[B:%.*]]
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i1(i1 [[CMP]], i1 false, i32 33)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i1(i1 [[CMP]], i1 false, i32 33)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp eq i32 %a, %b
-  %mask = call i64 @llvm.amdgcn.icmp.i1(i1 %cmp, i1 false, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i1(i1 %cmp, i1 false, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_i1_ne_0_icmp_ne_i1(i32 %a, i32 %b) {
 ; CHECK-LABEL: @fold_icmp_i1_ne_0_icmp_ne_i1(
 ; CHECK-NEXT:    [[CMP:%.*]] = icmp ne i32 [[A:%.*]], [[B:%.*]]
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i1(i1 [[CMP]], i1 false, i32 33)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i1(i1 [[CMP]], i1 false, i32 33)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp ne i32 %a, %b
-  %mask = call i64 @llvm.amdgcn.icmp.i1(i1 %cmp, i1 false, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i1(i1 %cmp, i1 false, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_i1_ne_0_icmp_sle_i1(i32 %a, i32 %b) {
 ; CHECK-LABEL: @fold_icmp_i1_ne_0_icmp_sle_i1(
 ; CHECK-NEXT:    [[CMP:%.*]] = icmp sle i32 [[A:%.*]], [[B:%.*]]
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i1(i1 [[CMP]], i1 false, i32 33)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i1(i1 [[CMP]], i1 false, i32 33)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp sle i32 %a, %b
-  %mask = call i64 @llvm.amdgcn.icmp.i1(i1 %cmp, i1 false, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i1(i1 %cmp, i1 false, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_i1_ne_0_icmp_ugt_i64(i64 %a, i64 %b) {
 ; CHECK-LABEL: @fold_icmp_i1_ne_0_icmp_ugt_i64(
 ; CHECK-NEXT:    [[CMP:%.*]] = icmp ugt i64 [[A:%.*]], [[B:%.*]]
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i1(i1 [[CMP]], i1 false, i32 33)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i1(i1 [[CMP]], i1 false, i32 33)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp ugt i64 %a, %b
-  %mask = call i64 @llvm.amdgcn.icmp.i1(i1 %cmp, i1 false, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i1(i1 %cmp, i1 false, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_i1_ne_0_icmp_ult_swap_i64(i64 %a, i64 %b) {
 ; CHECK-LABEL: @fold_icmp_i1_ne_0_icmp_ult_swap_i64(
 ; CHECK-NEXT:    [[CMP:%.*]] = icmp ugt i64 [[A:%.*]], [[B:%.*]]
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i1(i1 [[CMP]], i1 false, i32 33)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i1(i1 [[CMP]], i1 false, i32 33)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp ugt i64 %a, %b
-  %mask = call i64 @llvm.amdgcn.icmp.i1(i1 false, i1 %cmp, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i1(i1 false, i1 %cmp, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_i1_ne_0_fcmp_oeq_f32(float %a, float %b) {
 ; CHECK-LABEL: @fold_icmp_i1_ne_0_fcmp_oeq_f32(
 ; CHECK-NEXT:    [[CMP:%.*]] = fcmp oeq float [[A:%.*]], [[B:%.*]]
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i1(i1 [[CMP]], i1 false, i32 33)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i1(i1 [[CMP]], i1 false, i32 33)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = fcmp oeq float %a, %b
-  %mask = call i64 @llvm.amdgcn.icmp.i1(i1 %cmp, i1 false, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i1(i1 %cmp, i1 false, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_i1_ne_0_fcmp_une_f32(float %a, float %b) {
 ; CHECK-LABEL: @fold_icmp_i1_ne_0_fcmp_une_f32(
 ; CHECK-NEXT:    [[CMP:%.*]] = fcmp une float [[A:%.*]], [[B:%.*]]
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i1(i1 [[CMP]], i1 false, i32 33)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i1(i1 [[CMP]], i1 false, i32 33)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = fcmp une float %a, %b
-  %mask = call i64 @llvm.amdgcn.icmp.i1(i1 %cmp, i1 false, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i1(i1 %cmp, i1 false, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_i1_ne_0_fcmp_olt_f64(double %a, double %b) {
 ; CHECK-LABEL: @fold_icmp_i1_ne_0_fcmp_olt_f64(
 ; CHECK-NEXT:    [[CMP:%.*]] = fcmp olt double [[A:%.*]], [[B:%.*]]
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i1(i1 [[CMP]], i1 false, i32 33)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i1(i1 [[CMP]], i1 false, i32 33)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = fcmp olt double %a, %b
-  %mask = call i64 @llvm.amdgcn.icmp.i1(i1 %cmp, i1 false, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i1(i1 %cmp, i1 false, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_i1_ne_0_icmp_eq_i4(i4 %a, i4 %b) {
 ; CHECK-LABEL: @fold_icmp_i1_ne_0_icmp_eq_i4(
 ; CHECK-NEXT:    [[CMP:%.*]] = icmp eq i4 [[A:%.*]], [[B:%.*]]
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i1(i1 [[CMP]], i1 false, i32 33)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i1(i1 [[CMP]], i1 false, i32 33)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp eq i4 %a, %b
-  %mask = call i64 @llvm.amdgcn.icmp.i1(i1 %cmp, i1 false, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i1(i1 %cmp, i1 false, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_i1_ne_0_icmp_eq_i8(i8 %a, i8 %b) {
 ; CHECK-LABEL: @fold_icmp_i1_ne_0_icmp_eq_i8(
 ; CHECK-NEXT:    [[CMP:%.*]] = icmp eq i8 [[A:%.*]], [[B:%.*]]
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i1(i1 [[CMP]], i1 false, i32 33)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i1(i1 [[CMP]], i1 false, i32 33)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp eq i8 %a, %b
-  %mask = call i64 @llvm.amdgcn.icmp.i1(i1 %cmp, i1 false, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i1(i1 %cmp, i1 false, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_i1_ne_0_icmp_eq_i16(i16 %a, i16 %b) {
 ; CHECK-LABEL: @fold_icmp_i1_ne_0_icmp_eq_i16(
 ; CHECK-NEXT:    [[CMP:%.*]] = icmp eq i16 [[A:%.*]], [[B:%.*]]
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i1(i1 [[CMP]], i1 false, i32 33)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i1(i1 [[CMP]], i1 false, i32 33)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp eq i16 %a, %b
-  %mask = call i64 @llvm.amdgcn.icmp.i1(i1 %cmp, i1 false, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i1(i1 %cmp, i1 false, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_i1_ne_0_icmp_eq_i36(i36 %a, i36 %b) {
 ; CHECK-LABEL: @fold_icmp_i1_ne_0_icmp_eq_i36(
 ; CHECK-NEXT:    [[CMP:%.*]] = icmp eq i36 [[A:%.*]], [[B:%.*]]
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i1(i1 [[CMP]], i1 false, i32 33)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i1(i1 [[CMP]], i1 false, i32 33)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp eq i36 %a, %b
-  %mask = call i64 @llvm.amdgcn.icmp.i1(i1 %cmp, i1 false, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i1(i1 %cmp, i1 false, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_i1_ne_0_icmp_eq_i128(i128 %a, i128 %b) {
 ; CHECK-LABEL: @fold_icmp_i1_ne_0_icmp_eq_i128(
 ; CHECK-NEXT:    [[CMP:%.*]] = icmp eq i128 [[A:%.*]], [[B:%.*]]
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i1(i1 [[CMP]], i1 false, i32 33)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i1(i1 [[CMP]], i1 false, i32 33)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp eq i128 %a, %b
-  %mask = call i64 @llvm.amdgcn.icmp.i1(i1 %cmp, i1 false, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i1(i1 %cmp, i1 false, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_i1_ne_0_fcmp_oeq_f16(half %a, half %b) {
 ; CHECK-LABEL: @fold_icmp_i1_ne_0_fcmp_oeq_f16(
 ; CHECK-NEXT:    [[CMP:%.*]] = fcmp oeq half [[A:%.*]], [[B:%.*]]
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i1(i1 [[CMP]], i1 false, i32 33)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i1(i1 [[CMP]], i1 false, i32 33)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = fcmp oeq half %a, %b
-  %mask = call i64 @llvm.amdgcn.icmp.i1(i1 %cmp, i1 false, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i1(i1 %cmp, i1 false, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_i1_ne_0_fcmp_oeq_f128(fp128 %a, fp128 %b) {
-;
 ; CHECK-LABEL: @fold_icmp_i1_ne_0_fcmp_oeq_f128(
 ; CHECK-NEXT:    [[CMP:%.*]] = fcmp oeq fp128 [[A:%.*]], [[B:%.*]]
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i1(i1 [[CMP]], i1 false, i32 33)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i1(i1 [[CMP]], i1 false, i32 33)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = fcmp oeq fp128 %a, %b
-  %mask = call i64 @llvm.amdgcn.icmp.i1(i1 %cmp, i1 false, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i1(i1 %cmp, i1 false, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_i1_ne_0_icmp_slt_i4(i4 %a, i4 %b) {
 ; CHECK-LABEL: @fold_icmp_i1_ne_0_icmp_slt_i4(
 ; CHECK-NEXT:    [[CMP:%.*]] = icmp slt i4 [[A:%.*]], [[B:%.*]]
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i1(i1 [[CMP]], i1 false, i32 33)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i1(i1 [[CMP]], i1 false, i32 33)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp slt i4 %a, %b
-  %mask = call i64 @llvm.amdgcn.icmp.i1(i1 %cmp, i1 false, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i1(i1 %cmp, i1 false, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_i1_ne_0_icmp_slt_i8(i8 %a, i8 %b) {
 ; CHECK-LABEL: @fold_icmp_i1_ne_0_icmp_slt_i8(
 ; CHECK-NEXT:    [[CMP:%.*]] = icmp slt i8 [[A:%.*]], [[B:%.*]]
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i1(i1 [[CMP]], i1 false, i32 33)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i1(i1 [[CMP]], i1 false, i32 33)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp slt i8 %a, %b
-  %mask = call i64 @llvm.amdgcn.icmp.i1(i1 %cmp, i1 false, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i1(i1 %cmp, i1 false, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_i1_ne_0_icmp_slt_i16(i16 %a, i16 %b) {
 ; CHECK-LABEL: @fold_icmp_i1_ne_0_icmp_slt_i16(
 ; CHECK-NEXT:    [[CMP:%.*]] = icmp slt i16 [[A:%.*]], [[B:%.*]]
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i1(i1 [[CMP]], i1 false, i32 33)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i1(i1 [[CMP]], i1 false, i32 33)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp slt i16 %a, %b
-  %mask = call i64 @llvm.amdgcn.icmp.i1(i1 %cmp, i1 false, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i1(i1 %cmp, i1 false, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_i1_ne_0_icmp_ult_i4(i4 %a, i4 %b) {
 ; CHECK-LABEL: @fold_icmp_i1_ne_0_icmp_ult_i4(
 ; CHECK-NEXT:    [[CMP:%.*]] = icmp ult i4 [[A:%.*]], [[B:%.*]]
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i1(i1 [[CMP]], i1 false, i32 33)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i1(i1 [[CMP]], i1 false, i32 33)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp ult i4 %a, %b
-  %mask = call i64 @llvm.amdgcn.icmp.i1(i1 %cmp, i1 false, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i1(i1 %cmp, i1 false, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_i1_ne_0_icmp_ult_i8(i8 %a, i8 %b) {
 ; CHECK-LABEL: @fold_icmp_i1_ne_0_icmp_ult_i8(
 ; CHECK-NEXT:    [[CMP:%.*]] = icmp ult i8 [[A:%.*]], [[B:%.*]]
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i1(i1 [[CMP]], i1 false, i32 33)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i1(i1 [[CMP]], i1 false, i32 33)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp ult i8 %a, %b
-  %mask = call i64 @llvm.amdgcn.icmp.i1(i1 %cmp, i1 false, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i1(i1 %cmp, i1 false, i32 33)
   ret i64 %mask
 }
 
 define i64 @fold_icmp_i1_ne_0_icmp_ult_i16(i16 %a, i16 %b) {
 ; CHECK-LABEL: @fold_icmp_i1_ne_0_icmp_ult_i16(
 ; CHECK-NEXT:    [[CMP:%.*]] = icmp ult i16 [[A:%.*]], [[B:%.*]]
-; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i1(i1 [[CMP]], i1 false, i32 33)
+; CHECK-NEXT:    [[MASK:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i1(i1 [[CMP]], i1 false, i32 33)
 ; CHECK-NEXT:    ret i64 [[MASK]]
 ;
   %cmp = icmp ult i16 %a, %b
-  %mask = call i64 @llvm.amdgcn.icmp.i1(i1 %cmp, i1 false, i32 33)
+  %mask = call i64 @llvm.amdgcn.icmp.i64.i1(i1 %cmp, i1 false, i32 33)
   ret i64 %mask
 }
 
 ; --------------------------------------------------------------------
 ; llvm.amdgcn.fcmp
 ; --------------------------------------------------------------------
 
-declare i64 @llvm.amdgcn.fcmp.f32(float, float, i32 immarg) nounwind readnone convergent
+declare i64 @llvm.amdgcn.fcmp.i64.f32(float, float, i32 immarg) nounwind readnone convergent
 
 define i64 @invalid_fcmp_code(float %a, float %b) {
 ; CHECK-LABEL: @invalid_fcmp_code(
-; CHECK-NEXT:    [[UNDER:%.*]] = call i64 @llvm.amdgcn.fcmp.f32(float [[A:%.*]], float [[B:%.*]], i32 -1)
-; CHECK-NEXT:    [[OVER:%.*]] = call i64 @llvm.amdgcn.fcmp.f32(float [[A]], float [[B]], i32 16)
+; CHECK-NEXT:    [[UNDER:%.*]] = call i64 @llvm.amdgcn.fcmp.i64.f32(float [[A:%.*]], float [[B:%.*]], i32 -1)
+; CHECK-NEXT:    [[OVER:%.*]] = call i64 @llvm.amdgcn.fcmp.i64.f32(float [[A]], float [[B]], i32 16)
 ; CHECK-NEXT:    [[OR:%.*]] = or i64 [[UNDER]], [[OVER]]
 ; CHECK-NEXT:    ret i64 [[OR]]
 ;
-  %under = call i64 @llvm.amdgcn.fcmp.f32(float %a, float %b, i32 -1)
-  %over = call i64 @llvm.amdgcn.fcmp.f32(float %a, float %b, i32 16)
+  %under = call i64 @llvm.amdgcn.fcmp.i64.f32(float %a, float %b, i32 -1)
+  %over = call i64 @llvm.amdgcn.fcmp.i64.f32(float %a, float %b, i32 16)
   %or = or i64 %under, %over
   ret i64 %or
 }
 
 define i64 @fcmp_constant_inputs_false() {
 ; CHECK-LABEL: @fcmp_constant_inputs_false(
 ; CHECK-NEXT:    ret i64 0
 ;
-  %result = call i64 @llvm.amdgcn.fcmp.f32(float 2.0, float 4.0, i32 1)
+  %result = call i64 @llvm.amdgcn.fcmp.i64.f32(float 2.0, float 4.0, i32 1)
   ret i64 %result
 }
 
 define i64 @fcmp_constant_inputs_true() {
 ; CHECK-LABEL: @fcmp_constant_inputs_true(
 ; CHECK-NEXT:    [[RESULT:%.*]] = call i64 @llvm.read_register.i64(metadata !0) #5
 ; CHECK-NEXT:    ret i64 [[RESULT]]
 ;
-  %result = call i64 @llvm.amdgcn.fcmp.f32(float 2.0, float 4.0, i32 4)
+  %result = call i64 @llvm.amdgcn.fcmp.i64.f32(float 2.0, float 4.0, i32 4)
   ret i64 %result
 }
 
 define i64 @fcmp_constant_to_rhs_olt(float %x) {
 ; CHECK-LABEL: @fcmp_constant_to_rhs_olt(
-; CHECK-NEXT:    [[RESULT:%.*]] = call i64 @llvm.amdgcn.fcmp.f32(float [[X:%.*]], float 4.000000e+00, i32 2)
+; CHECK-NEXT:    [[RESULT:%.*]] = call i64 @llvm.amdgcn.fcmp.i64.f32(float [[X:%.*]], float 4.000000e+00, i32 2)
 ; CHECK-NEXT:    ret i64 [[RESULT]]
 ;
-  %result = call i64 @llvm.amdgcn.fcmp.f32(float 4.0, float %x, i32 4)
+  %result = call i64 @llvm.amdgcn.fcmp.i64.f32(float 4.0, float %x, i32 4)
   ret i64 %result
 }
 
 ; --------------------------------------------------------------------
 ; llvm.amdgcn.wqm.vote
 ; --------------------------------------------------------------------
 
 declare i1 @llvm.amdgcn.wqm.vote(i1)

test/Verifier/AMDGPU/intrinsic-immarg.ll

define void @exp_compr_invalid_inputs(i32 %tgt, i32 %en, i1 %bool) {
 
   ; CHECK: immarg operand has non-immediate parameter
   ; CHECK-NEXT: i1 %bool
   ; CHECK-NEXT: call void @llvm.amdgcn.exp.compr.v2f16(i32 0, i32 5, <2 x half> <half 0xH3C00, half 0xH4000>, <2 x half> <half 0xH3800, half 0xH4400>, i1 false, i1 %bool)
   call void @llvm.amdgcn.exp.compr.v2f16(i32 0, i32 5, <2 x half> <half 1.0, half 2.0>, <2 x half> <half 0.5, half 4.0>, i1 false, i1 %bool)
   ret void
 }
 
-declare i64 @llvm.amdgcn.icmp.i32(i32, i32, i32)
+declare i64 @llvm.amdgcn.icmp.i64.i32(i32, i32, i32)
 
 define i64 @invalid_nonconstant_icmp_code(i32 %a, i32 %b, i32 %c) {
   ; CHECK: immarg operand has non-immediate parameter
   ; CHECK-NEXT: i32 %c
-  ; CHECK-NEXT: %result = call i64 @llvm.amdgcn.icmp.i32(i32 %a, i32 %b, i32 %c)
-  %result = call i64 @llvm.amdgcn.icmp.i32(i32 %a, i32 %b, i32 %c)
+  ; CHECK-NEXT: %result = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %a, i32 %b, i32 %c)
+  %result = call i64 @llvm.amdgcn.icmp.i64.i32(i32 %a, i32 %b, i32 %c)
   ret i64 %result
 }
 
-declare i64 @llvm.amdgcn.fcmp.f32(float, float, i32)
+declare i64 @llvm.amdgcn.fcmp.i64.f32(float, float, i32)
 define i64 @invalid_nonconstant_fcmp_code(float %a, float %b, i32 %c) {
   ; CHECK: immarg operand has non-immediate parameter
   ; CHECK-NEXT: i32 %c
-  ; CHECK-NEXT: %result = call i64 @llvm.amdgcn.fcmp.f32(float %a, float %b, i32 %c)
-  %result = call i64 @llvm.amdgcn.fcmp.f32(float %a, float %b, i32 %c)
+  ; CHECK-NEXT: %result = call i64 @llvm.amdgcn.fcmp.i64.f32(float %a, float %b, i32 %c)
+  %result = call i64 @llvm.amdgcn.fcmp.i64.f32(float %a, float %b, i32 %c)
   ret i64 %result
 }
 
 declare i32 @llvm.amdgcn.atomic.inc.i32.p3i32(i32 addrspace(3)* nocapture, i32, i32, i32, i1)
 define amdgpu_kernel void @invalid_atomic_inc(i32 addrspace(1)* %out, i32 addrspace(3)* %ptr, i32 %var, i1 %bool) {
   ; CHECK: immarg operand has non-immediate parameter
   ; CHECK-NEXT: i32 %var
   ; CHECK-NEXT: %result0 = call i32 @llvm.amdgcn.atomic.inc.i32.p3i32(i32 addrspace(3)* %ptr, i32 42, i32 %var, i32 0, i1 false)