[AMDGPU] Stop using make_pair and make_tuple. NFC.

[foad]

Summary

C++17 allows us to call constructors pair and tuple instead of helper
functions make_pair and make_tuple.

[kazu]

Thanks!

llvm/lib/Target/AMDGPU/AMDGPUArgumentUsageInfo.cpp

void AMDGPUArgumentUsageInfo::print(raw_ostream &OS, const Module *M) const {
   }
 }
 
 std::tuple<const ArgDescriptor *, const TargetRegisterClass *, LLT>
 AMDGPUFunctionArgInfo::getPreloadedValue(
     AMDGPUFunctionArgInfo::PreloadedValue Value) const {
   switch (Value) {
   case AMDGPUFunctionArgInfo::PRIVATE_SEGMENT_BUFFER: {
-    return std::make_tuple(PrivateSegmentBuffer ? &PrivateSegmentBuffer
-                                                : nullptr,
-                           &AMDGPU::SGPR_128RegClass, LLT::fixed_vector(4, 32));
+    return std::tuple(PrivateSegmentBuffer ? &PrivateSegmentBuffer : nullptr,
   }
   case AMDGPUFunctionArgInfo::IMPLICIT_BUFFER_PTR:
-    return std::make_tuple(ImplicitBufferPtr ? &ImplicitBufferPtr : nullptr,
-                           &AMDGPU::SGPR_64RegClass,
-                           LLT::pointer(AMDGPUAS::CONSTANT_ADDRESS, 64));
+    return std::tuple(ImplicitBufferPtr ? &ImplicitBufferPtr : nullptr,
   case AMDGPUFunctionArgInfo::WORKGROUP_ID_X:
-    return std::make_tuple(WorkGroupIDX ? &WorkGroupIDX : nullptr,
-                           &AMDGPU::SGPR_32RegClass, LLT::scalar(32));
+    return std::tuple(WorkGroupIDX ? &WorkGroupIDX : nullptr,
   case AMDGPUFunctionArgInfo::WORKGROUP_ID_Y:
-    return std::make_tuple(WorkGroupIDY ? &WorkGroupIDY : nullptr,
-                           &AMDGPU::SGPR_32RegClass, LLT::scalar(32));
+    return std::tuple(WorkGroupIDY ? &WorkGroupIDY : nullptr,
   case AMDGPUFunctionArgInfo::WORKGROUP_ID_Z:
-    return std::make_tuple(WorkGroupIDZ ? &WorkGroupIDZ : nullptr,
-                           &AMDGPU::SGPR_32RegClass, LLT::scalar(32));
+    return std::tuple(WorkGroupIDZ ? &WorkGroupIDZ : nullptr,
   case AMDGPUFunctionArgInfo::LDS_KERNEL_ID:
-    return std::make_tuple(LDSKernelId ? &LDSKernelId : nullptr,
-                           &AMDGPU::SGPR_32RegClass, LLT::scalar(32));
+    return std::tuple(LDSKernelId ? &LDSKernelId : nullptr,
   case AMDGPUFunctionArgInfo::PRIVATE_SEGMENT_WAVE_BYTE_OFFSET:
-    return std::make_tuple(
+    return std::tuple(
         PrivateSegmentWaveByteOffset ? &PrivateSegmentWaveByteOffset : nullptr,
         &AMDGPU::SGPR_32RegClass, LLT::scalar(32));
   case AMDGPUFunctionArgInfo::KERNARG_SEGMENT_PTR:
-    return std::make_tuple(KernargSegmentPtr ? &KernargSegmentPtr : nullptr,
-                           &AMDGPU::SGPR_64RegClass,
-                           LLT::pointer(AMDGPUAS::CONSTANT_ADDRESS, 64));
+    return std::tuple(KernargSegmentPtr ? &KernargSegmentPtr : nullptr,
   case AMDGPUFunctionArgInfo::IMPLICIT_ARG_PTR:
-    return std::make_tuple(ImplicitArgPtr ? &ImplicitArgPtr : nullptr,
-                           &AMDGPU::SGPR_64RegClass,
-                           LLT::pointer(AMDGPUAS::CONSTANT_ADDRESS, 64));
+    return std::tuple(ImplicitArgPtr ? &ImplicitArgPtr : nullptr,
   case AMDGPUFunctionArgInfo::DISPATCH_ID:
-    return std::make_tuple(DispatchID ? &DispatchID : nullptr,
-                           &AMDGPU::SGPR_64RegClass, LLT::scalar(64));
+    return std::tuple(DispatchID ? &DispatchID : nullptr,
   case AMDGPUFunctionArgInfo::FLAT_SCRATCH_INIT:
-    return std::make_tuple(FlatScratchInit ? &FlatScratchInit : nullptr,
-                           &AMDGPU::SGPR_64RegClass, LLT::scalar(64));
+    return std::tuple(FlatScratchInit ? &FlatScratchInit : nullptr,
   case AMDGPUFunctionArgInfo::DISPATCH_PTR:
-    return std::make_tuple(DispatchPtr ? &DispatchPtr : nullptr,
-                           &AMDGPU::SGPR_64RegClass,
-                           LLT::pointer(AMDGPUAS::CONSTANT_ADDRESS, 64));
+    return std::tuple(DispatchPtr ? &DispatchPtr : nullptr,
   case AMDGPUFunctionArgInfo::QUEUE_PTR:
-    return std::make_tuple(QueuePtr ? &QueuePtr : nullptr,
-                           &AMDGPU::SGPR_64RegClass,
-                           LLT::pointer(AMDGPUAS::CONSTANT_ADDRESS, 64));
+    return std::tuple(QueuePtr ? &QueuePtr : nullptr, &AMDGPU::SGPR_64RegClass,
   case AMDGPUFunctionArgInfo::WORKITEM_ID_X:
-    return std::make_tuple(WorkItemIDX ? &WorkItemIDX : nullptr,
-                           &AMDGPU::VGPR_32RegClass, LLT::scalar(32));
+    return std::tuple(WorkItemIDX ? &WorkItemIDX : nullptr,
   case AMDGPUFunctionArgInfo::WORKITEM_ID_Y:
-    return std::make_tuple(WorkItemIDY ? &WorkItemIDY : nullptr,
-                           &AMDGPU::VGPR_32RegClass, LLT::scalar(32));
+    return std::tuple(WorkItemIDY ? &WorkItemIDY : nullptr,
   case AMDGPUFunctionArgInfo::WORKITEM_ID_Z:
-    return std::make_tuple(WorkItemIDZ ? &WorkItemIDZ : nullptr,
-                           &AMDGPU::VGPR_32RegClass, LLT::scalar(32));
+    return std::tuple(WorkItemIDZ ? &WorkItemIDZ : nullptr,
   }
   llvm_unreachable("unexpected preloaded value type");
 }
 
 constexpr AMDGPUFunctionArgInfo AMDGPUFunctionArgInfo::fixedABILayout() {
   AMDGPUFunctionArgInfo AI;
   AI.PrivateSegmentBuffer
     = ArgDescriptor::createRegister(AMDGPU::SGPR0_SGPR1_SGPR2_SGPR3);

llvm/lib/Target/AMDGPU/AMDGPUCodeGenPrepare.cpp

static std::pair<Value*, Value*> getMul64(IRBuilder<> &Builder,
   Type *I64Ty = Builder.getInt64Ty();
 
   Value *LHS_EXT64 = Builder.CreateZExt(LHS, I64Ty);
   Value *RHS_EXT64 = Builder.CreateZExt(RHS, I64Ty);
   Value *MUL64 = Builder.CreateMul(LHS_EXT64, RHS_EXT64);
   Value *Lo = Builder.CreateTrunc(MUL64, I32Ty);
   Value *Hi = Builder.CreateLShr(MUL64, Builder.getInt64(32));
   Hi = Builder.CreateTrunc(Hi, I32Ty);
-  return std::make_pair(Lo, Hi);
+  return std::pair(Lo, Hi);
 }
 
 static Value* getMulHu(IRBuilder<> &Builder, Value *LHS, Value *RHS) {
   return getMul64(Builder, LHS, RHS).second;
 }
 
 /// Figure out how many bits are really needed for this division. \p AtLeast is
 /// an optimization hint to bypass the second ComputeNumSignBits call if we the

llvm/lib/Target/AMDGPU/AMDGPUGlobalISelUtils.cpp

AMDGPU::getBaseWithConstantOffset(MachineRegisterInfo &MRI, Register Reg,
   if (Def->getOpcode() == TargetOpcode::G_CONSTANT) {
     unsigned Offset;
     const MachineOperand &Op = Def->getOperand(1);
     if (Op.isImm())
       Offset = Op.getImm();
     else
       Offset = Op.getCImm()->getZExtValue();
 
-    return std::make_pair(Register(), Offset);
+    return std::pair(Register(), Offset);
   }
 
   int64_t Offset;
   if (Def->getOpcode() == TargetOpcode::G_ADD) {
     // TODO: Handle G_OR used for add case
     if (mi_match(Def->getOperand(2).getReg(), MRI, m_ICst(Offset)))
-      return std::make_pair(Def->getOperand(1).getReg(), Offset);
+      return std::pair(Def->getOperand(1).getReg(), Offset);
 
     // FIXME: matcher should ignore copies
     if (mi_match(Def->getOperand(2).getReg(), MRI, m_Copy(m_ICst(Offset))))
-      return std::make_pair(Def->getOperand(1).getReg(), Offset);
+      return std::pair(Def->getOperand(1).getReg(), Offset);
   }
 
   Register Base;
   if (KnownBits && mi_match(Reg, MRI, m_GOr(m_Reg(Base), m_ICst(Offset))) &&
       KnownBits->maskedValueIsZero(Base, APInt(32, Offset)))
-    return std::make_pair(Base, Offset);
+    return std::pair(Base, Offset);
 
   // Handle G_PTRTOINT (G_PTR_ADD base, const) case
   if (Def->getOpcode() == TargetOpcode::G_PTRTOINT) {
     MachineInstr *Base;
     if (mi_match(Def->getOperand(1).getReg(), MRI,
                  m_GPtrAdd(m_MInstr(Base), m_ICst(Offset)))) {
       // If Base was int converted to pointer, simply return int and offset.
       if (Base->getOpcode() == TargetOpcode::G_INTTOPTR)
-        return std::make_pair(Base->getOperand(1).getReg(), Offset);
+        return std::pair(Base->getOperand(1).getReg(), Offset);
 
       // Register returned here will be of pointer type.
-      return std::make_pair(Base->getOperand(0).getReg(), Offset);
+      return std::pair(Base->getOperand(0).getReg(), Offset);
     }
   }
 
-  return std::make_pair(Reg, 0);
+  return std::pair(Reg, 0);
 }
 
 bool AMDGPU::hasAtomicFaddRtnForTy(const GCNSubtarget &Subtarget,
                                    const LLT &Ty) {
   if (Ty == LLT::scalar(32))
     return Subtarget.hasAtomicFaddRtnInsts();
   if (Ty == LLT::fixed_vector(2, 16) || Ty == LLT::scalar(64))
     return Subtarget.hasGFX90AInsts();
   return false;
 }

llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp

static std::pair<Type *, Align> getArgumentTypeAlign(const Argument &Arg,
   if (Arg.hasByRefAttr()) {
     Ty = Arg.getParamByRefType();
     ArgAlign = Arg.getParamAlign();
   }
 
   if (!ArgAlign)
     ArgAlign = DL.getABITypeAlign(Ty);
 
-  return std::make_pair(Ty, *ArgAlign);
+  return std::pair(Ty, *ArgAlign);
 }
 
 namespace llvm {
 
 static cl::opt<bool> DumpHSAMetadata(
     "amdgpu-dump-hsa-metadata",
     cl::desc("Dump AMDGPU HSA Metadata"));
 static cl::opt<bool> VerifyHSAMetadata(

llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp

void PipelineSolver::convertSyncMapsToArrays() {
   }
 
   int PipelineIDx = SyncedInstrs.size() - 1;
   PipelineInstrs.resize(SyncedInstrs.size());
   for (auto &SyncInstrMap : SyncedInstrs) {
     for (auto &SUsToCandSGs : SyncInstrMap.second) {
       if (PipelineInstrs[PipelineIDx].size() == 0) {
         PipelineInstrs[PipelineIDx].push_back(
-            std::make_pair(SUsToCandSGs.first, SUsToCandSGs.second));
+            std::pair(SUsToCandSGs.first, SUsToCandSGs.second));
         continue;
       }
       auto SortPosition = PipelineInstrs[PipelineIDx].begin();
       // Insert them in sorted order -- this allows for good parsing order in
       // the greedy algorithm
       while (SortPosition != PipelineInstrs[PipelineIDx].end() &&
              SUsToCandSGs.first->NodeNum > SortPosition->first->NodeNum)
         ++SortPosition;
       PipelineInstrs[PipelineIDx].insert(
-          SortPosition,
-          std::make_pair(SUsToCandSGs.first, SUsToCandSGs.second));
+          SortPosition, std::pair(SUsToCandSGs.first, SUsToCandSGs.second));
     }
     --PipelineIDx;
   }
 }
 
 void PipelineSolver::makePipeline() {
   // Preserve the order of barrier for subsequent SchedGroupBarrier mutations
   for (auto &SyncPipeline : BestPipeline) {
for (; I != E; ++I) {
     SchedGroup *Match;
     for (auto &SG : SyncPipeline) {
       if (SG.getSGID() == CandSGID)
         Match = &SG;
     }
 
     if (UseCostHeur) {
       if (Match->isFull()) {
-        ReadyList.push_back(std::make_pair(*I, MissPenalty));
+        ReadyList.push_back(std::pair(*I, MissPenalty));
         continue;
       }
 
       int TempCost = addEdges(SyncPipeline, CurrSU.first, CandSGID, AddedEdges);
-      ReadyList.push_back(std::make_pair(*I, TempCost));
+      ReadyList.push_back(std::pair(*I, TempCost));
       removeEdges(AddedEdges);
     } else
-      ReadyList.push_back(std::make_pair(*I, -1));
+      ReadyList.push_back(std::pair(*I, -1));
   }
 
   if (UseCostHeur) {
     std::sort(ReadyList.begin(), ReadyList.end(),
               [](std::pair<int, int> A, std::pair<int, int> B) {
                 return A.second < B.second;
               });
   }
if (MakePred)
       std::swap(A, B);
 
     if (DAG->IsReachable(B, A))
       continue;
     // tryAddEdge returns false if there is a dependency that makes adding
     // the A->B edge impossible, otherwise it returns true;
     bool Added = tryAddEdge(A, B);
     if (Added)
-      AddedEdges.push_back(std::make_pair(A, B));
+      AddedEdges.push_back(std::pair(A, B));
     else
       ++MissedEdges;
   }
 
   return MissedEdges;
 }
 
 void SchedGroup::link(SUnit &SU, bool MakePred) {

llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp

std::pair<SDValue, SDValue> AMDGPUDAGToDAGISel::foldFrameIndex(SDValue N) const {
   auto *FI = dyn_cast<FrameIndexSDNode>(N);
   SDValue TFI =
       FI ? CurDAG->getTargetFrameIndex(FI->getIndex(), FI->getValueType(0)) : N;
 
   // We rebase the base address into an absolute stack address and hence
   // use constant 0 for soffset. This value must be retained until
   // frame elimination and eliminateFrameIndex will choose the appropriate
   // frame register if need be.
-  return std::make_pair(TFI, CurDAG->getTargetConstant(0, DL, MVT::i32));
+  return std::pair(TFI, CurDAG->getTargetConstant(0, DL, MVT::i32));
 }
 
 bool AMDGPUDAGToDAGISel::SelectMUBUFScratchOffen(SDNode *Parent,
                                                  SDValue Addr, SDValue &Rsrc,
                                                  SDValue &VAddr, SDValue &SOffset,
                                                  SDValue &ImmOffset) const {
 
   SDLoc DL(Addr);

llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp

AMDGPUTargetLowering::split64BitValue(SDValue Op, SelectionDAG &DAG) const {
   SDValue Vec = DAG.getNode(ISD::BITCAST, SL, MVT::v2i32, Op);
 
   const SDValue Zero = DAG.getConstant(0, SL, MVT::i32);
   const SDValue One = DAG.getConstant(1, SL, MVT::i32);
 
   SDValue Lo = DAG.getNode(ISD::EXTRACT_VECTOR_ELT, SL, MVT::i32, Vec, Zero);
   SDValue Hi = DAG.getNode(ISD::EXTRACT_VECTOR_ELT, SL, MVT::i32, Vec, One);
 
-  return std::make_pair(Lo, Hi);
+  return std::pair(Lo, Hi);
 }
 
 SDValue AMDGPUTargetLowering::getLoHalf64(SDValue Op, SelectionDAG &DAG) const {
   SDLoc SL(Op);
 
   SDValue Vec = DAG.getNode(ISD::BITCAST, SL, MVT::v2i32, Op);
   const SDValue Zero = DAG.getConstant(0, SL, MVT::i32);
   return DAG.getNode(ISD::EXTRACT_VECTOR_ELT, SL, MVT::i32, Vec, Zero);
AMDGPUTargetLowering::getSplitDestVTs(const EVT &VT, SelectionDAG &DAG) const {
   EVT LoVT, HiVT;
   EVT EltVT = VT.getVectorElementType();
   unsigned NumElts = VT.getVectorNumElements();
   unsigned LoNumElts = PowerOf2Ceil((NumElts + 1) / 2);
   LoVT = EVT::getVectorVT(*DAG.getContext(), EltVT, LoNumElts);
   HiVT = NumElts - LoNumElts == 1
              ? EltVT
              : EVT::getVectorVT(*DAG.getContext(), EltVT, NumElts - LoNumElts);
-  return std::make_pair(LoVT, HiVT);
+  return std::pair(LoVT, HiVT);
 }
 
 // Split a vector value into two parts of types LoVT and HiVT. HiVT could be
 // scalar.
 std::pair<SDValue, SDValue>
 AMDGPUTargetLowering::splitVector(const SDValue &N, const SDLoc &DL,
                                   const EVT &LoVT, const EVT &HiVT,
                                   SelectionDAG &DAG) const {
   assert(LoVT.getVectorNumElements() +
                  (HiVT.isVector() ? HiVT.getVectorNumElements() : 1) <=
              N.getValueType().getVectorNumElements() &&
          "More vector elements requested than available!");
   SDValue Lo = DAG.getNode(ISD::EXTRACT_SUBVECTOR, DL, LoVT, N,
                            DAG.getVectorIdxConstant(0, DL));
   SDValue Hi = DAG.getNode(
       HiVT.isVector() ? ISD::EXTRACT_SUBVECTOR : ISD::EXTRACT_VECTOR_ELT, DL,
       HiVT, N, DAG.getVectorIdxConstant(LoVT.getVectorNumElements(), DL));
-  return std::make_pair(Lo, Hi);
+  return std::pair(Lo, Hi);
 }
 
 SDValue AMDGPUTargetLowering::SplitVectorLoad(const SDValue Op,
                                               SelectionDAG &DAG) const {
   LoadSDNode *Load = cast<LoadSDNode>(Op);
   EVT VT = Op.getValueType();
   SDLoc SL(Op);
 

llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp

if (IdxBaseReg == AMDGPU::NoRegister) {
     IdxBaseReg = IdxReg;
   }
 
   ArrayRef<int16_t> SubRegs = TRI.getRegSplitParts(SuperRC, EltSize);
 
   // Skip out of bounds offsets, or else we would end up using an undefined
   // register.
   if (static_cast<unsigned>(Offset) >= SubRegs.size())
-    return std::make_pair(IdxReg, SubRegs[0]);
-  return std::make_pair(IdxBaseReg, SubRegs[Offset]);
+    return std::pair(IdxReg, SubRegs[0]);
+  return std::pair(IdxBaseReg, SubRegs[Offset]);
 }
 
 bool AMDGPUInstructionSelector::selectG_EXTRACT_VECTOR_ELT(
   MachineInstr &MI) const {
   Register DstReg = MI.getOperand(0).getReg();
   Register SrcReg = MI.getOperand(1).getReg();
   Register IdxReg = MI.getOperand(2).getReg();
 
std::pair<Register, unsigned> AMDGPUInstructionSelector::selectVOP3ModsImpl(
   if (AllowAbs && MI->getOpcode() == AMDGPU::G_FABS) {
     Src = MI->getOperand(1).getReg();
     Mods |= SISrcMods::ABS;
   }
 
   if (OpSel)
     Mods |= SISrcMods::OP_SEL_0;
 
-  return std::make_pair(Src, Mods);
+  return std::pair(Src, Mods);
 }
 
 Register AMDGPUInstructionSelector::copyToVGPRIfSrcFolded(
     Register Src, unsigned Mods, MachineOperand Root, MachineInstr *InsertPt,
     bool ForceVGPR) const {
   if ((Mods != 0 || ForceVGPR) &&
       RBI.getRegBank(Src, *MRI, TRI)->getID() != AMDGPU::VGPRRegBankID) {
 
AMDGPUInstructionSelector::selectVOP3PModsImpl(
   }
 
   // TODO: Match op_sel through g_build_vector_trunc and g_shuffle_vector.
   (void)IsDOT; // DOTs do not use OPSEL on gfx940+, check ST.hasDOTOpSelHazard()
 
   // Packed instructions do not have abs modifiers.
   Mods |= SISrcMods::OP_SEL_1;
 
-  return std::make_pair(Src, Mods);
+  return std::pair(Src, Mods);
 }
 
 InstructionSelector::ComplexRendererFns
 AMDGPUInstructionSelector::selectVOP3PMods(MachineOperand &Root) const {
   MachineRegisterInfo &MRI
     = Root.getParent()->getParent()->getParent()->getRegInfo();
 
   Register Src;
return {{[=](MachineInstrBuilder &MIB) { MIB.addReg(Base); },
            [=](MachineInstrBuilder &MIB) { MIB.addImm(Offset); }}};
 }
 
 std::pair<Register, int>
 AMDGPUInstructionSelector::selectFlatOffsetImpl(MachineOperand &Root,
                                                 uint64_t FlatVariant) const {
   MachineInstr *MI = Root.getParent();
 
-  auto Default = std::make_pair(Root.getReg(), 0);
+  auto Default = std::pair(Root.getReg(), 0);
 
   if (!STI.hasFlatInstOffsets())
     return Default;
 
   Register PtrBase;
   int64_t ConstOffset;
   std::tie(PtrBase, ConstOffset) =
       getPtrBaseWithConstantOffset(Root.getReg(), *MRI);
   if (ConstOffset == 0)
     return Default;
 
   unsigned AddrSpace = (*MI->memoperands_begin())->getAddrSpace();
   if (!TII.isLegalFLATOffset(ConstOffset, AddrSpace, FlatVariant))
     return Default;
 
-  return std::make_pair(PtrBase, ConstOffset);
+  return std::pair(PtrBase, ConstOffset);
 }
 
 InstructionSelector::ComplexRendererFns
 AMDGPUInstructionSelector::selectFlatOffset(MachineOperand &Root) const {
   auto PtrWithOffset = selectFlatOffsetImpl(Root, SIInstrFlags::FLAT);
 
   return {{
       [=](MachineInstrBuilder &MIB) { MIB.addReg(PtrWithOffset.first); },
return {{
       [=](MachineInstrBuilder &MIB) { MIB.addImm(Offset); } // offset
   }};
 }
 
 std::pair<Register, unsigned>
 AMDGPUInstructionSelector::selectDS1Addr1OffsetImpl(MachineOperand &Root) const {
   const MachineInstr *RootDef = MRI->getVRegDef(Root.getReg());
   if (!RootDef)
-    return std::make_pair(Root.getReg(), 0);
+    return std::pair(Root.getReg(), 0);
 
   int64_t ConstAddr = 0;
 
   Register PtrBase;
   int64_t Offset;
   std::tie(PtrBase, Offset) =
     getPtrBaseWithConstantOffset(Root.getReg(), *MRI);
 
   if (Offset) {
     if (isDSOffsetLegal(PtrBase, Offset)) {
       // (add n0, c0)
-      return std::make_pair(PtrBase, Offset);
+      return std::pair(PtrBase, Offset);
     }
   } else if (RootDef->getOpcode() == AMDGPU::G_SUB) {
     // TODO
 
 
   } else if (mi_match(Root.getReg(), *MRI, m_ICst(ConstAddr))) {
     // TODO
 
   }
 
-  return std::make_pair(Root.getReg(), 0);
+  return std::pair(Root.getReg(), 0);
 }
 
 InstructionSelector::ComplexRendererFns
 AMDGPUInstructionSelector::selectDS1Addr1Offset(MachineOperand &Root) const {
   Register Reg;
   unsigned Offset;
   std::tie(Reg, Offset) = selectDS1Addr1OffsetImpl(Root);
   return {{
return {{
     }};
 }
 
 std::pair<Register, unsigned>
 AMDGPUInstructionSelector::selectDSReadWrite2Impl(MachineOperand &Root,
                                                   unsigned Size) const {
   const MachineInstr *RootDef = MRI->getVRegDef(Root.getReg());
   if (!RootDef)
-    return std::make_pair(Root.getReg(), 0);
+    return std::pair(Root.getReg(), 0);
 
   int64_t ConstAddr = 0;
 
   Register PtrBase;
   int64_t Offset;
   std::tie(PtrBase, Offset) =
     getPtrBaseWithConstantOffset(Root.getReg(), *MRI);
 
   if (Offset) {
     int64_t OffsetValue0 = Offset;
     int64_t OffsetValue1 = Offset + Size;
     if (isDSOffset2Legal(PtrBase, OffsetValue0, OffsetValue1, Size)) {
       // (add n0, c0)
-      return std::make_pair(PtrBase, OffsetValue0 / Size);
+      return std::pair(PtrBase, OffsetValue0 / Size);
     }
   } else if (RootDef->getOpcode() == AMDGPU::G_SUB) {
     // TODO
 
   } else if (mi_match(Root.getReg(), *MRI, m_ICst(ConstAddr))) {
     // TODO
 
   }
 
-  return std::make_pair(Root.getReg(), 0);
+  return std::pair(Root.getReg(), 0);
 }
 
 /// If \p Root is a G_PTR_ADD with a G_CONSTANT on the right hand side, return
 /// the base value with the constant offset. There may be intervening copies
 /// between \p Root and the identified constant. Returns \p Root, 0 if this does
 /// not match the pattern.
 std::pair<Register, int64_t>
 AMDGPUInstructionSelector::getPtrBaseWithConstantOffset(

llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp

return [=](const LegalityQuery &Query) {
     return EltTy.getSizeInBits() == 16 && Ty.getNumElements() > 2;
   };
 }
 
 static LegalizeMutation oneMoreElement(unsigned TypeIdx) {
   return [=](const LegalityQuery &Query) {
     const LLT Ty = Query.Types[TypeIdx];
     const LLT EltTy = Ty.getElementType();
-    return std::make_pair(TypeIdx,
-                          LLT::fixed_vector(Ty.getNumElements() + 1, EltTy));
+    return std::pair(TypeIdx,
   };
 }
 
 static LegalizeMutation fewerEltsToSize64Vector(unsigned TypeIdx) {
   return [=](const LegalityQuery &Query) {
     const LLT Ty = Query.Types[TypeIdx];
     const LLT EltTy = Ty.getElementType();
     unsigned Size = Ty.getSizeInBits();
     unsigned Pieces = (Size + 63) / 64;
     unsigned NewNumElts = (Ty.getNumElements() + 1) / Pieces;
-    return std::make_pair(
-        TypeIdx,
-        LLT::scalarOrVector(ElementCount::getFixed(NewNumElts), EltTy));
+    return std::pair(TypeIdx, LLT::scalarOrVector(
+                                  ElementCount::getFixed(NewNumElts), EltTy));
   };
 }
 
 // Increase the number of vector elements to reach the next multiple of 32-bit
 // type.
 static LegalizeMutation moreEltsToNext32Bit(unsigned TypeIdx) {
   return [=](const LegalityQuery &Query) {
     const LLT Ty = Query.Types[TypeIdx];
 
     const LLT EltTy = Ty.getElementType();
     const int Size = Ty.getSizeInBits();
     const int EltSize = EltTy.getSizeInBits();
     const int NextMul32 = (Size + 31) / 32;
 
     assert(EltSize < 32);
 
     const int NewNumElts = (32 * NextMul32 + EltSize - 1) / EltSize;
-    return std::make_pair(TypeIdx, LLT::fixed_vector(NewNumElts, EltTy));
+    return std::pair(TypeIdx, LLT::fixed_vector(NewNumElts, EltTy));
   };
 }
 
 static LLT getBitcastRegisterType(const LLT Ty) {
   const unsigned Size = Ty.getSizeInBits();
 
   if (Size <= 32) {
     // <2 x s8> -> s16
     // <4 x s8> -> s32
     return LLT::scalar(Size);
   }
 
   return LLT::scalarOrVector(ElementCount::getFixed(Size / 32), 32);
 }
 
 static LegalizeMutation bitcastToRegisterType(unsigned TypeIdx) {
   return [=](const LegalityQuery &Query) {
     const LLT Ty = Query.Types[TypeIdx];
-    return std::make_pair(TypeIdx, getBitcastRegisterType(Ty));
+    return std::pair(TypeIdx, getBitcastRegisterType(Ty));
   };
 }
 
 static LegalizeMutation bitcastToVectorElement32(unsigned TypeIdx) {
   return [=](const LegalityQuery &Query) {
     const LLT Ty = Query.Types[TypeIdx];
     unsigned Size = Ty.getSizeInBits();
     assert(Size % 32 == 0);
-    return std::make_pair(
+    return std::pair(
         TypeIdx, LLT::scalarOrVector(ElementCount::getFixed(Size / 32), 32));
   };
 }
 
 static LegalityPredicate vectorSmallerThan(unsigned TypeIdx, unsigned Size) {
   return [=](const LegalityQuery &Query) {
     const LLT QueryTy = Query.Types[TypeIdx];
     return QueryTy.isVector() && QueryTy.getSizeInBits() < Size;
getActionDefinitionsBuilder({G_SMIN, G_SMAX, G_UMIN, G_UMAX, G_ABS})
       .legalFor({S32})
       .minScalar(0, S32)
       .widenScalarToNextPow2(0)
       .scalarize(0)
       .lower();
   }
 
   getActionDefinitionsBuilder(G_INTTOPTR)
       // List the common cases
       .legalForCartesianProduct(AddrSpaces64, {S64})
       .legalForCartesianProduct(AddrSpaces32, {S32})
       .scalarize(0)
       // Accept any address space as long as the size matches
       .legalIf(sameSize(0, 1))
       .widenScalarIf(smallerThan(1, 0),
                      [](const LegalityQuery &Query) {
-        return std::make_pair(1, LLT::scalar(Query.Types[0].getSizeInBits()));
+                       return std::pair(
+                           1, LLT::scalar(Query.Types[0].getSizeInBits()));
                      })
-    .narrowScalarIf(largerThan(1, 0),
-      [](const LegalityQuery &Query) {
-        return std::make_pair(1, LLT::scalar(Query.Types[0].getSizeInBits()));
+      .narrowScalarIf(largerThan(1, 0), [](const LegalityQuery &Query) {
+        return std::pair(1, LLT::scalar(Query.Types[0].getSizeInBits()));
       });
 
   getActionDefinitionsBuilder(G_PTRTOINT)
       // List the common cases
       .legalForCartesianProduct(AddrSpaces64, {S64})
       .legalForCartesianProduct(AddrSpaces32, {S32})
       .scalarize(0)
       // Accept any address space as long as the size matches
       .legalIf(sameSize(0, 1))
       .widenScalarIf(smallerThan(0, 1),
                      [](const LegalityQuery &Query) {
-        return std::make_pair(0, LLT::scalar(Query.Types[1].getSizeInBits()));
+                       return std::pair(
+                           0, LLT::scalar(Query.Types[1].getSizeInBits()));
                      })
-    .narrowScalarIf(
-      largerThan(0, 1),
-      [](const LegalityQuery &Query) {
-        return std::make_pair(0, LLT::scalar(Query.Types[1].getSizeInBits()));
+      .narrowScalarIf(largerThan(0, 1), [](const LegalityQuery &Query) {
+        return std::pair(0, LLT::scalar(Query.Types[1].getSizeInBits()));
       });
 
   getActionDefinitionsBuilder(G_ADDRSPACE_CAST)
     .scalarize(0)
     .custom();
 
   const auto needToSplitMemOp = [=](const LegalityQuery &Query,
                                     bool IsLoad) -> bool {
Actions
               const LLT DstTy = Query.Types[0];
               const LLT PtrTy = Query.Types[1];
 
               const unsigned DstSize = DstTy.getSizeInBits();
               unsigned MemSize = Query.MMODescrs[0].MemoryTy.getSizeInBits();
 
               // Split extloads.
               if (DstSize > MemSize)
-                return std::make_pair(0, LLT::scalar(MemSize));
+                return std::pair(0, LLT::scalar(MemSize));
 
               unsigned MaxSize = maxSizeForAddrSpace(ST,
                                                      PtrTy.getAddressSpace(),
                                                      Op == G_LOAD);
               if (MemSize > MaxSize)
-                return std::make_pair(0, LLT::scalar(MaxSize));
+                return std::pair(0, LLT::scalar(MaxSize));
 
               uint64_t Align = Query.MMODescrs[0].AlignInBits;
-              return std::make_pair(0, LLT::scalar(Align));
+              return std::pair(0, LLT::scalar(Align));
             })
         .fewerElementsIf(
             [=](const LegalityQuery &Query) -> bool {
               return Query.Types[0].isVector() &&
                      needToSplitMemOp(Query, Op == G_LOAD);
             },
             [=](const LegalityQuery &Query) -> std::pair<unsigned, LLT> {
               const LLT DstTy = Query.Types[0];
Actions
 
               // Split if it's too large for the address space.
               unsigned MemSize = Query.MMODescrs[0].MemoryTy.getSizeInBits();
               if (MemSize > MaxSize) {
                 unsigned NumElts = DstTy.getNumElements();
                 unsigned EltSize = EltTy.getSizeInBits();
 
                 if (MaxSize % EltSize == 0) {
-                  return std::make_pair(
+                  return std::pair(
                       0, LLT::scalarOrVector(
                              ElementCount::getFixed(MaxSize / EltSize), EltTy));
                 }
 
                 unsigned NumPieces = MemSize / MaxSize;
 
                 // FIXME: Refine when odd breakdowns handled
                 // The scalars will need to be re-legalized.
                 if (NumPieces == 1 || NumPieces >= NumElts ||
                     NumElts % NumPieces != 0)
-                  return std::make_pair(0, EltTy);
+                  return std::pair(0, EltTy);
 
-                return std::make_pair(
-                    0, LLT::fixed_vector(NumElts / NumPieces, EltTy));
+                return std::pair(0,
+                                 LLT::fixed_vector(NumElts / NumPieces, EltTy));
               }
 
               // FIXME: We could probably handle weird extending loads better.
               if (DstTy.getSizeInBits() > MemSize)
-                return std::make_pair(0, EltTy);
+                return std::pair(0, EltTy);
 
               unsigned EltSize = EltTy.getSizeInBits();
               unsigned DstSize = DstTy.getSizeInBits();
               if (!isPowerOf2_32(DstSize)) {
                 // We're probably decomposing an odd sized store. Try to split
                 // to the widest type. TODO: Account for alignment. As-is it
                 // should be OK, since the new parts will be further legalized.
                 unsigned FloorSize = PowerOf2Floor(DstSize);
-                return std::make_pair(
+                return std::pair(
                     0, LLT::scalarOrVector(
                            ElementCount::getFixed(FloorSize / EltSize), EltTy));
               }
 
               // May need relegalization for the scalars.
-              return std::make_pair(0, EltTy);
+              return std::pair(0, EltTy);
             })
     .minScalar(0, S32)
     .narrowScalarIf(isWideScalarExtLoadTruncStore(0), changeTo(0, S32))
     .widenScalarToNextPow2(0)
     .moreElementsIf(vectorSmallerThan(0, 32), moreEltsToNext32Bit(0))
     .lower();
   }
 
getActionDefinitionsBuilder(Op)
           // element vector since we may be able to do better indexing
           // if this is scalar. If not, fall back to 32.
           const LLT EltTy = Query.Types[EltTypeIdx];
           const LLT VecTy = Query.Types[VecTypeIdx];
           const unsigned DstEltSize = EltTy.getSizeInBits();
           const unsigned VecSize = VecTy.getSizeInBits();
 
           const unsigned TargetEltSize = DstEltSize % 64 == 0 ? 64 : 32;
-          return std::make_pair(
+          return std::pair(
               VecTypeIdx,
               LLT::fixed_vector(VecSize / TargetEltSize, TargetEltSize));
         })
       .clampScalar(EltTypeIdx, S32, S64)
       .clampScalar(VecTypeIdx, S32, S64)
       .clampScalar(IdxTypeIdx, S32, S32)
       .clampMaxNumElements(VecTypeIdx, S32, 32)
       // TODO: Clamp elements for 64-bit vectors?
Builder.widenScalarIf(
         // Whichever is smaller.
         const LLT &Ty = Query.Types[BigTyIdx];
         unsigned NewSizeInBits = 1 << Log2_32_Ceil(Ty.getSizeInBits() + 1);
         if (NewSizeInBits >= 256) {
           unsigned RoundedTo = alignTo<64>(Ty.getSizeInBits() + 1);
           if (RoundedTo < NewSizeInBits)
             NewSizeInBits = RoundedTo;
         }
-        return std::make_pair(BigTyIdx, LLT::scalar(NewSizeInBits));
+        return std::pair(BigTyIdx, LLT::scalar(NewSizeInBits));
       })
       // Any vectors left are the wrong size. Scalarize them.
       .scalarize(0)
       .scalarize(1);
   }
 
   // S64 is only legal on SALU, and needs to be broken into 32-bit elements in
   // RegBankSelect.
if (!BaseReg) {
       auto OverflowVal = B.buildConstant(S32, Overflow);
       BaseReg = B.buildAdd(S32, BaseReg, OverflowVal).getReg(0);
     }
   }
 
   if (!BaseReg)
     BaseReg = B.buildConstant(S32, 0).getReg(0);
 
-  return std::make_pair(BaseReg, ImmOffset);
+  return std::pair(BaseReg, ImmOffset);
 }
 
 /// Update \p MMO based on the offset inputs to a raw/struct buffer intrinsic.
 void AMDGPULegalizerInfo::updateBufferMMO(MachineMemOperand *MMO,
                                           Register VOffset, Register SOffset,
                                           unsigned ImmOffset, Register VIndex,
                                           MachineRegisterInfo &MRI) const {
   std::optional<ValueAndVReg> MaybeVOffsetVal =

llvm/lib/Target/AMDGPU/AMDGPUMachineCFGStructurizer.cpp

for (MachineBasicBlock *Succ : EndMBB->successors()) {
         Succs.insert(
             std::pair<MachineBasicBlock *, MachineBasicBlock *>(EndMBB, Succ));
       }
     }
   }
 
   for (MachineBasicBlock *Pred : StartMBB->predecessors())
     if (Pred != EndMBB)
-      Succs.insert(std::make_pair(Pred, StartMBB));
+      Succs.insert(std::pair(Pred, StartMBB));
 
   for (auto SI : Succs) {
     std::pair<MachineBasicBlock *, MachineBasicBlock *> Edge = SI;
     LLVM_DEBUG(dbgs() << "Removing edge: " << printMBBReference(*Edge.first)
                       << " -> " << printMBBReference(*Edge.second) << "\n");
     Edge.first->removeSuccessor(Edge.second);
   }
 }
for (auto DRI = PHIInfo.dests_begin(), DE = PHIInfo.dests_end(); DRI != DE;
     // is always executed when the region is executed.
     if (MBBContainsPHISource) {
       for (auto SRI = PHIInfo.sources_begin(DestReg); SRI != SE; ++SRI) {
         PHILinearize::PHISourceT Source = *SRI;
         unsigned SourceReg = Source.first;
         MachineBasicBlock *SourceMBB = Source.second;
         MachineOperand *Def = &(*(MRI->def_begin(SourceReg)));
         if (Def->getParent()->getParent() != MBB) {
-          ElimiatedSources.push_back(
-              std::make_tuple(DestReg, SourceReg, SourceMBB));
+          ElimiatedSources.push_back(std::tuple(DestReg, SourceReg, SourceMBB));
         }
       }
     }
   }
 
   // Remove the PHI sources that are in the given MBB
   for (auto &SourceInfo : ElimiatedSources) {
     PHIInfo.removeSource(std::get<0>(SourceInfo), std::get<1>(SourceInfo),

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp

AMDGPUMachineFunction::AMDGPUMachineFunction(const MachineFunction &MF)
   CallingConv::ID CC = F.getCallingConv();
   if (CC == CallingConv::AMDGPU_KERNEL || CC == CallingConv::SPIR_KERNEL)
     ExplicitKernArgSize = ST.getExplicitKernArgSize(F, MaxKernArgAlign);
 }
 
 unsigned AMDGPUMachineFunction::allocateLDSGlobal(const DataLayout &DL,
                                                   const GlobalVariable &GV,
                                                   Align Trailing) {
-  auto Entry = LocalMemoryObjects.insert(std::make_pair(&GV, 0));
+  auto Entry = LocalMemoryObjects.insert(std::pair(&GV, 0));
   if (!Entry.second)
     return Entry.first->second;
 
   Align Alignment =
       DL.getValueOrABITypeAlignment(GV.getAlign(), GV.getValueType());
 
   unsigned Offset;
   if (GV.getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS) {

llvm/lib/Target/AMDGPU/AMDGPUPromoteAlloca.cpp

Function *LocalSizeZFn
       = Intrinsic::getDeclaration(Mod, Intrinsic::r600_read_local_size_z);
 
     CallInst *LocalSizeY = Builder.CreateCall(LocalSizeYFn, {});
     CallInst *LocalSizeZ = Builder.CreateCall(LocalSizeZFn, {});
 
     ST.makeLIDRangeMetadata(LocalSizeY);
     ST.makeLIDRangeMetadata(LocalSizeZ);
 
-    return std::make_pair(LocalSizeY, LocalSizeZ);
+    return std::pair(LocalSizeY, LocalSizeZ);
   }
 
   // We must read the size out of the dispatch pointer.
   assert(IsAMDGCN);
 
   // We are indexing into this struct, and want to extract the workgroup_size_*
   // fields.
   //
AMDGPUPromoteAllocaImpl::getLocalSizeYZ(IRBuilder<> &Builder) {
   MDNode *MD = MDNode::get(Mod->getContext(), std::nullopt);
   LoadXY->setMetadata(LLVMContext::MD_invariant_load, MD);
   LoadZU->setMetadata(LLVMContext::MD_invariant_load, MD);
   ST.makeLIDRangeMetadata(LoadZU);
 
   // Extract y component. Upper half of LoadZU should be zero already.
   Value *Y = Builder.CreateLShr(LoadXY, 16);
 
-  return std::make_pair(Y, LoadZU);
+  return std::pair(Y, LoadZU);
 }
 
 Value *AMDGPUPromoteAllocaImpl::getWorkitemID(IRBuilder<> &Builder,
                                               unsigned N) {
   Function *F = Builder.GetInsertBlock()->getParent();
   const AMDGPUSubtarget &ST = AMDGPUSubtarget::get(TM, *F);
   Intrinsic::ID IntrID = Intrinsic::not_intrinsic;
   StringRef AttrName;

llvm/lib/Target/AMDGPU/AMDGPUPropagateAttributes.cpp

for (auto &F : M.functions()) {
             break;
           }
 
           NewF = cloneWithProperties(F, NewProps);
           Clones.push_back(Clone(CallerProps, &F, NewF));
           NewRoots.insert(NewF);
         }
 
-        ToReplace.push_back(std::make_pair(CI, NewF));
+        ToReplace.push_back(std::pair(CI, NewF));
         Replaced.insert(&F);
 
         Changed = true;
       }
 
       while (!ToReplace.empty()) {
         auto R = ToReplace.pop_back_val();
         R.first->setCalledFunction(R.second);

llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp

for (MachineOperand &Op : MI.uses()) {
           CondReg = B.buildAnd(S1, CondReg, CmpReg).getReg(0);
           MRI.setRegBank(CondReg, AMDGPU::VCCRegBank);
         }
       }
 
       Op.setReg(CurrentLaneReg);
 
       // Make sure we don't re-process this register again.
-      WaterfalledRegMap.insert(std::make_pair(OldReg, Op.getReg()));
+      WaterfalledRegMap.insert(std::pair(OldReg, Op.getReg()));
     }
   }
 
   // The ballot becomes a no-op during instruction selection.
   CondReg = B.buildIntrinsic(Intrinsic::amdgcn_ballot,
                              {LLT::scalar(Subtarget.isWave32() ? 32 : 64)},
                              false)
                 .addReg(CondReg)

 static std::pair<Register, Register>
 unpackV2S16ToS32(MachineIRBuilder &B, Register Src, unsigned ExtOpcode) {
   const LLT S32 = LLT::scalar(32);
   auto Bitcast = B.buildBitcast(S32, Src);
 
   if (ExtOpcode == TargetOpcode::G_SEXT) {
     auto ExtLo = B.buildSExtInReg(S32, Bitcast, 16);
     auto ShiftHi = B.buildAShr(S32, Bitcast, B.buildConstant(S32, 16));
-    return std::make_pair(ExtLo.getReg(0), ShiftHi.getReg(0));
+    return std::pair(ExtLo.getReg(0), ShiftHi.getReg(0));
   }
 
   auto ShiftHi = B.buildLShr(S32, Bitcast, B.buildConstant(S32, 16));
   if (ExtOpcode == TargetOpcode::G_ZEXT) {
     auto ExtLo = B.buildAnd(S32, Bitcast, B.buildConstant(S32, 0xffff));
-    return std::make_pair(ExtLo.getReg(0), ShiftHi.getReg(0));
+    return std::pair(ExtLo.getReg(0), ShiftHi.getReg(0));
   }
 
   assert(ExtOpcode == TargetOpcode::G_ANYEXT);
-  return std::make_pair(Bitcast.getReg(0), ShiftHi.getReg(0));
+  return std::pair(Bitcast.getReg(0), ShiftHi.getReg(0));
 }
 
 // For cases where only a single copy is inserted for matching register banks.
 // Replace the register in the instruction operand
 static bool substituteSimpleCopyRegs(
   const AMDGPURegisterBankInfo::OperandsMapper &OpdMapper, unsigned OpIdx) {
   SmallVector<unsigned, 1> SrcReg(OpdMapper.getVRegs(OpIdx));
   if (!SrcReg.empty()) {
Register AMDGPURegisterBankInfo::handleD16VData(MachineIRBuilder &B,
 
   return B.buildMerge(LLT::fixed_vector(NumElts, S32), WideRegs).getReg(0);
 }
 
 static std::pair<Register, unsigned>
 getBaseWithConstantOffset(MachineRegisterInfo &MRI, Register Reg) {
   int64_t Const;
   if (mi_match(Reg, MRI, m_ICst(Const)))
-    return std::make_pair(Register(), Const);
+    return std::pair(Register(), Const);
 
   Register Base;
   if (mi_match(Reg, MRI, m_GAdd(m_Reg(Base), m_ICst(Const))))
-    return std::make_pair(Base, Const);
+    return std::pair(Base, Const);
 
   // TODO: Handle G_OR used for add case
-  return std::make_pair(Reg, 0);
+  return std::pair(Reg, 0);
 }
 
 std::pair<Register, unsigned>
 AMDGPURegisterBankInfo::splitBufferOffsets(MachineIRBuilder &B,
                                            Register OrigOffset) const {
   const unsigned MaxImm = 4095;
   Register BaseReg;
   unsigned ImmOffset;

llvm/lib/Target/AMDGPU/AMDGPUReplaceLDSUseWithPointer.cpp

bool shouldIgnorePointerReplacement(GlobalVariable *GV) {
 
     // FIXME: When GV is used within all (or within most of the kernels), then
     // it does not make sense to create a pointer for it.
   }
 
   // Insert new global LDS pointer which points to LDS.
   GlobalVariable *createLDSPointer(GlobalVariable *GV) {
     // LDS pointer which points to LDS is already created? Return it.
-    auto PointerEntry = LDSToPointer.insert(std::make_pair(GV, nullptr));
+    auto PointerEntry = LDSToPointer.insert(std::pair(GV, nullptr));
     if (!PointerEntry.second)
       return PointerEntry.first->second;
 
     // We need to create new LDS pointer which points to LDS.
     //
     // Each CU owns at max 64K of LDS memory, so LDS address ranges from 0 to
     // 2^16 - 1. Hence 16 bit pointer is enough to hold the LDS address.
     auto *I16Ty = Type::getInt16Ty(Ctx);
GlobalVariable *createLDSPointer(GlobalVariable *GV) {
     return LDSPointer;
   }
 
   // Split entry basic block in such a way that only lane 0 of each wave does
   // the LDS pointer initialization, and return newly created basic block.
   BasicBlock *activateLaneZero(Function *K) {
     // If the entry basic block of kernel K is already split, then return
     // newly created basic block.
-    auto BasicBlockEntry = KernelToInitBB.insert(std::make_pair(K, nullptr));
+    auto BasicBlockEntry = KernelToInitBB.insert(std::pair(K, nullptr));
     if (!BasicBlockEntry.second)
       return BasicBlockEntry.first->second;
 
     // Split entry basic block of kernel K.
     auto *EI = &(*(K->getEntryBlock().getFirstInsertionPt()));
     IRBuilder<> Builder(EI);
 
     Value *Mbcnt =
BasicBlock *activateLaneZero(Function *K) {
     return NBB;
   }
 
   // Within given kernel, initialize given LDS pointer to point to given LDS.
   void initializeLDSPointer(Function *K, GlobalVariable *GV,
                             GlobalVariable *LDSPointer) {
     // If LDS pointer is already initialized within K, then nothing to do.
     auto PointerEntry = KernelToLDSPointers.insert(
-        std::make_pair(K, SmallPtrSet<GlobalVariable *, 8>()));
+        std::pair(K, SmallPtrSet<GlobalVariable *, 8>()));
     if (!PointerEntry.second)
       if (PointerEntry.first->second.contains(LDSPointer))
         return;
 
     // Insert instructions at EI which initialize LDS pointer to point-to LDS
     // within kernel K.
     //
     // That is, convert pointer type of GV to i16, and then store this converted
class ReplaceLDSUseImpl {
 
   // Create a set of replacement instructions which together replace LDS within
   // non-kernel function F by accessing LDS indirectly using LDS pointer.
   Value *getReplacementInst(Function *F, GlobalVariable *GV,
                             GlobalVariable *LDSPointer) {
     // If the instruction which replaces LDS within F is already created, then
     // return it.
     auto LDSEntry = FunctionToLDSToReplaceInst.insert(
-        std::make_pair(F, DenseMap<GlobalVariable *, Value *>()));
+        std::pair(F, DenseMap<GlobalVariable *, Value *>()));
     if (!LDSEntry.second) {
       auto ReplaceInstEntry =
-          LDSEntry.first->second.insert(std::make_pair(GV, nullptr));
+          LDSEntry.first->second.insert(std::pair(GV, nullptr));
       if (!ReplaceInstEntry.second)
         return ReplaceInstEntry.first->second;
     }
 
     // Get the instruction insertion point within the beginning of the entry
     // block of current non-kernel function.
     auto *EI = &(*(F->getEntryBlock().getFirstInsertionPt()));
     IRBuilder<> Builder(EI);
if (CollectKernelInsts) {
         continue;
       }
     } else {
       if (llvm::AMDGPU::isKernelCC(F)) {
         continue;
       }
     }
 
-    FunctionToInsts.insert(std::make_pair(F, SmallPtrSet<Instruction *, 8>()));
+    FunctionToInsts.insert(std::pair(F, SmallPtrSet<Instruction *, 8>()));
     FunctionToInsts[F].insert(I);
   }
 
   return FunctionToInsts;
 }
 
 } // namespace AMDGPU
 

llvm/lib/Target/AMDGPU/AMDGPUResourceUsageAnalysis.cpp

bool AMDGPUResourceUsageAnalysis::runOnModule(Module &M) {
   for (auto IT = po_begin(&CG); IT != End; ++IT) {
     Function *F = IT->getFunction();
     if (!F || F->isDeclaration())
       continue;
 
     MachineFunction *MF = MMI.getMachineFunction(*F);
     assert(MF && "function must have been generated already");
 
-    auto CI = CallGraphResourceInfo.insert(
-        std::make_pair(F, SIFunctionResourceInfo()));
+    auto CI =
+        CallGraphResourceInfo.insert(std::pair(F, SIFunctionResourceInfo()));
     SIFunctionResourceInfo &Info = CI.first->second;
     assert(CI.second && "should only be called once per function");
     Info = analyzeResourceUsage(*MF, TM);
     HasIndirectCall |= Info.HasIndirectCall;
   }
 
   // It's possible we have unreachable functions in the module which weren't
   // visited by the PO traversal. Make sure we have some resource counts to
   // report.
   for (const auto &IT : CG) {
     const Function *F = IT.first;
     if (!F || F->isDeclaration())
       continue;
 
-    auto CI = CallGraphResourceInfo.insert(
-      std::make_pair(F, SIFunctionResourceInfo()));
+    auto CI =
+        CallGraphResourceInfo.insert(std::pair(F, SIFunctionResourceInfo()));
     if (!CI.second) // Skip already visited functions
       continue;
 
     SIFunctionResourceInfo &Info = CI.first->second;
     MachineFunction *MF = MMI.getMachineFunction(*F);
     assert(MF && "function must have been generated already");
     Info = analyzeResourceUsage(*MF, TM);
     HasIndirectCall |= Info.HasIndirectCall;

llvm/lib/Target/AMDGPU/AMDGPURewriteOutArguments.cpp

 ///  int foo(int a, int b, int* out) {
 ///     *out = bar();
 ///     return a + b;
 /// }
 ///
 /// into something like this:
 ///
 ///  std::pair<int, int> foo(int a, int b) {
-///     return std::make_pair(a + b, bar());
+///     return std::pair(a + b, bar());
 /// }
 ///
 /// Typically the incoming pointer is a simple alloca for a temporary variable
 /// to use the API, which if replaced with a struct return will be easily SROA'd
 /// out when the stub function we create is inlined
 ///
 /// This pass introduces the struct return, but leaves the unused pointer
 /// arguments and introduces a new stub function calling the struct returning

llvm/lib/Target/AMDGPU/AMDGPUSubtarget.cpp

 AMDGPUSubtarget::getDefaultFlatWorkGroupSize(CallingConv::ID CC) const {
   switch (CC) {
   case CallingConv::AMDGPU_VS:
   case CallingConv::AMDGPU_LS:
   case CallingConv::AMDGPU_HS:
   case CallingConv::AMDGPU_ES:
   case CallingConv::AMDGPU_GS:
   case CallingConv::AMDGPU_PS:
-    return std::make_pair(1, getWavefrontSize());
+    return std::pair(1, getWavefrontSize());
   default:
-    return std::make_pair(1u, getMaxFlatWorkGroupSize());
+    return std::pair(1u, getMaxFlatWorkGroupSize());
   }
 }
 
 std::pair<unsigned, unsigned> AMDGPUSubtarget::getFlatWorkGroupSizes(
   const Function &F) const {
   // Default minimum/maximum flat work group sizes.
   std::pair<unsigned, unsigned> Default =
     getDefaultFlatWorkGroupSize(F.getCallingConv());

llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp

unsigned AMDGPUTargetMachine::getAssumedAddrSpace(const Value *V) const {
   return AMDGPUAS::GLOBAL_ADDRESS;
 }
 
 std::pair<const Value *, unsigned>
 AMDGPUTargetMachine::getPredicatedAddrSpace(const Value *V) const {
   if (auto *II = dyn_cast<IntrinsicInst>(V)) {
     switch (II->getIntrinsicID()) {
     case Intrinsic::amdgcn_is_shared:
-      return std::make_pair(II->getArgOperand(0), AMDGPUAS::LOCAL_ADDRESS);
+      return std::pair(II->getArgOperand(0), AMDGPUAS::LOCAL_ADDRESS);
     case Intrinsic::amdgcn_is_private:
-      return std::make_pair(II->getArgOperand(0), AMDGPUAS::PRIVATE_ADDRESS);
+      return std::pair(II->getArgOperand(0), AMDGPUAS::PRIVATE_ADDRESS);
     default:
       break;
     }
-    return std::make_pair(nullptr, -1);
+    return std::pair(nullptr, -1);
   }
   // Check the global pointer predication based on
   // (!is_share(p) && !is_private(p)). Note that logic 'and' is commutative and
   // the order of 'is_shared' and 'is_private' is not significant.
   Value *Ptr;
   if (match(
           const_cast<Value *>(V),
           m_c_And(m_Not(m_Intrinsic<Intrinsic::amdgcn_is_shared>(m_Value(Ptr))),
                   m_Not(m_Intrinsic<Intrinsic::amdgcn_is_private>(
                       m_Deferred(Ptr))))))
-    return std::make_pair(Ptr, AMDGPUAS::GLOBAL_ADDRESS);
+    return std::pair(Ptr, AMDGPUAS::GLOBAL_ADDRESS);
 
-  return std::make_pair(nullptr, -1);
+  return std::pair(nullptr, -1);
 }
 
 unsigned
 AMDGPUTargetMachine::getAddressSpaceForPseudoSourceKind(unsigned Kind) const {
   switch (Kind) {
   case PseudoSourceValue::Stack:
   case PseudoSourceValue::FixedStack:
     return AMDGPUAS::PRIVATE_ADDRESS;

llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp

bool AMDGPUAsmParser::ParseDirectiveISAVersion() {
   return false;
 }
 
 bool AMDGPUAsmParser::ParseDirectiveHSAMetadata() {
   const char *AssemblerDirectiveBegin;
   const char *AssemblerDirectiveEnd;
   std::tie(AssemblerDirectiveBegin, AssemblerDirectiveEnd) =
       isHsaAbiVersion3AndAbove(&getSTI())
-          ? std::make_tuple(HSAMD::V3::AssemblerDirectiveBegin,
-                            HSAMD::V3::AssemblerDirectiveEnd)
-          : std::make_tuple(HSAMD::AssemblerDirectiveBegin,
-                            HSAMD::AssemblerDirectiveEnd);
+          ? std::tuple(HSAMD::V3::AssemblerDirectiveBegin,
+          : std::tuple(HSAMD::AssemblerDirectiveBegin,
 
   if (getSTI().getTargetTriple().getOS() != Triple::AMDHSA) {
     return Error(getLoc(),
                  (Twine(AssemblerDirectiveBegin) + Twine(" directive is "
                  "not available on non-amdhsa OSes")).str());
   }
 
   std::string HSAMetadataString;

llvm/lib/Target/AMDGPU/GCNNSAReassign.cpp

bool GCNNSAReassign::runOnMachineFunction(MachineFunction &MF) {
   using Candidate = std::pair<const MachineInstr*, bool>;
   SmallVector<Candidate, 32> Candidates;
   for (const MachineBasicBlock &MBB : MF) {
     for (const MachineInstr &MI : MBB) {
       switch (CheckNSA(MI)) {
       default:
         continue;
       case NSA_Status::CONTIGUOUS:
-        Candidates.push_back(std::make_pair(&MI, true));
+        Candidates.push_back(std::pair(&MI, true));
         break;
       case NSA_Status::NON_CONTIGUOUS:
-        Candidates.push_back(std::make_pair(&MI, false));
+        Candidates.push_back(std::pair(&MI, false));
         ++NumNSAInstructions;
         break;
       }
     }
   }
 
   bool Changed = false;
   for (auto &C : Candidates) {

llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp

GCNScheduleDAGMILive::createSchedStage(GCNSchedStageID SchedStageID) {
   }
 
   llvm_unreachable("Unknown SchedStageID.");
 }
 
 void GCNScheduleDAGMILive::schedule() {
   // Collect all scheduling regions. The actual scheduling is performed in
   // GCNScheduleDAGMILive::finalizeSchedule.
-  Regions.push_back(std::make_pair(RegionBegin, RegionEnd));
+  Regions.push_back(std::pair(RegionBegin, RegionEnd));
 }
 
 GCNRegPressure
 GCNScheduleDAGMILive::getRealRegPressure(unsigned RegionIdx) const {
   GCNDownwardRPTracker RPTracker(*LIS);
   RPTracker.advance(begin(), end(), &LiveIns[RegionIdx]);
   return RPTracker.moveMaxPressure();
 }
void GCNSchedStage::setupNewBlock() {
   DAG.startBlock(CurrentMBB);
   // Get real RP for the region if it hasn't be calculated before. After the
   // initial schedule stage real RP will be collected after scheduling.
   if (StageID == GCNSchedStageID::OccInitialSchedule)
     DAG.computeBlockPressure(RegionIdx, CurrentMBB);
 }
 
 void GCNSchedStage::finalizeGCNRegion() {
-  DAG.Regions[RegionIdx] = std::make_pair(DAG.RegionBegin, DAG.RegionEnd);
+  DAG.Regions[RegionIdx] = std::pair(DAG.RegionBegin, DAG.RegionEnd);
   DAG.RescheduleRegions[RegionIdx] = false;
   if (S.HasHighPressure)
     DAG.RegionsWithHighRP[RegionIdx] = true;
 
   // Revert scheduling if we have dropped occupancy or there is some other
   // reason that the original schedule is better.
   checkScheduling();
 
for (MachineInstr *MI : Unsched) {
       break;
     }
   }
 
   // Then move the debug instructions back into their correct place and set
   // RegionBegin and RegionEnd if needed.
   DAG.placeDebugValues();
 
-  DAG.Regions[RegionIdx] = std::make_pair(DAG.RegionBegin, DAG.RegionEnd);
+  DAG.Regions[RegionIdx] = std::pair(DAG.RegionBegin, DAG.RegionEnd);
 }
 
 void PreRARematStage::collectRematerializableInstructions() {
   const SIRegisterInfo *SRI = static_cast<const SIRegisterInfo *>(DAG.TRI);
   for (unsigned I = 0, E = DAG.MRI.getNumVirtRegs(); I != E; ++I) {
     Register Reg = Register::index2VirtReg(I);
     if (!DAG.LIS->hasInterval(Reg))
       continue;
for (; I != E; ++I) {
     if (MI->getParent() != RegionBoundaries[I].first->getParent())
       return;
 
     if (Removing && MI == RegionBoundaries[I].first &&
         MI == RegionBoundaries[I].second) {
       // MI is in a region with size 1, after removing, the region will be
       // size 0, set RegionBegin and RegionEnd to pass end of block iterator.
       RegionBoundaries[I] =
-          std::make_pair(MI->getParent()->end(), MI->getParent()->end());
+          std::pair(MI->getParent()->end(), MI->getParent()->end());
       return;
     }
     if (MI == RegionBoundaries[I].first) {
       if (Removing)
         RegionBoundaries[I] =
-            std::make_pair(std::next(MI), RegionBoundaries[I].second);
+            std::pair(std::next(MI), RegionBoundaries[I].second);
       else
         // Inserted NewMI in front of region, set new RegionBegin to NewMI
-        RegionBoundaries[I] = std::make_pair(MachineBasicBlock::iterator(NewMI),
-                                             RegionBoundaries[I].second);
+        RegionBoundaries[I] = std::pair(MachineBasicBlock::iterator(NewMI),
       return;
     }
     if (Removing && MI == RegionBoundaries[I].second) {
-      RegionBoundaries[I] =
-          std::make_pair(RegionBoundaries[I].first, std::prev(MI));
+      RegionBoundaries[I] = std::pair(RegionBoundaries[I].first, std::prev(MI));
       return;
     }
   }
 }
 
 static bool hasIGLPInstrs(ScheduleDAGInstrs *DAG) {
   return std::any_of(
       DAG->begin(), DAG->end(), [](MachineBasicBlock::iterator MI) {

llvm/lib/Target/AMDGPU/R600InstrInfo.cpp

if (MI.getOpcode() == R600::DOT_4) {
     };
 
     for (const auto &Op : OpTable) {
       MachineOperand &MO = MI.getOperand(getOperandIdx(MI.getOpcode(), Op[0]));
       Register Reg = MO.getReg();
       if (Reg == R600::ALU_CONST) {
         MachineOperand &Sel =
             MI.getOperand(getOperandIdx(MI.getOpcode(), Op[1]));
-        Result.push_back(std::make_pair(&MO, Sel.getImm()));
+        Result.push_back(std::pair(&MO, Sel.getImm()));
         continue;
       }
     }
     return Result;
   }
 
   static const unsigned OpTable[3][2] = {
     {R600::OpName::src0, R600::OpName::src0_sel},
     {R600::OpName::src1, R600::OpName::src1_sel},
     {R600::OpName::src2, R600::OpName::src2_sel},
   };
 
   for (const auto &Op : OpTable) {
     int SrcIdx = getOperandIdx(MI.getOpcode(), Op[0]);
     if (SrcIdx < 0)
       break;
     MachineOperand &MO = MI.getOperand(SrcIdx);
     Register Reg = MO.getReg();
     if (Reg == R600::ALU_CONST) {
       MachineOperand &Sel = MI.getOperand(getOperandIdx(MI.getOpcode(), Op[1]));
-      Result.push_back(std::make_pair(&MO, Sel.getImm()));
+      Result.push_back(std::pair(&MO, Sel.getImm()));
       continue;
     }
     if (Reg == R600::ALU_LITERAL_X) {
       MachineOperand &Operand =
           MI.getOperand(getOperandIdx(MI.getOpcode(), R600::OpName::literal));
       if (Operand.isImm()) {
-        Result.push_back(std::make_pair(&MO, Operand.getImm()));
+        Result.push_back(std::pair(&MO, Operand.getImm()));
         continue;
       }
       assert(Operand.isGlobal());
     }
-    Result.push_back(std::make_pair(&MO, 0));
+    Result.push_back(std::pair(&MO, 0));
   }
   return Result;
 }
 
 std::vector<std::pair<int, unsigned>>
 R600InstrInfo::ExtractSrcs(MachineInstr &MI,
                            const DenseMap<unsigned, unsigned> &PV,
                            unsigned &ConstCount) const {
   ConstCount = 0;
   const std::pair<int, unsigned> DummyPair(-1, 0);
   std::vector<std::pair<int, unsigned>> Result;
   unsigned i = 0;
   for (const auto &Src : getSrcs(MI)) {
     ++i;
     Register Reg = Src.first->getReg();
     int Index = RI.getEncodingValue(Reg) & 0xff;
     if (Reg == R600::OQAP) {
-      Result.push_back(std::make_pair(Index, 0U));
+      Result.push_back(std::pair(Index, 0U));
     }
     if (PV.find(Reg) != PV.end()) {
       // 255 is used to tells its a PS/PV reg
-      Result.push_back(std::make_pair(255, 0U));
+      Result.push_back(std::pair(255, 0U));
       continue;
     }
     if (Index > 127) {
       ConstCount++;
       Result.push_back(DummyPair);
       continue;
     }
     unsigned Chan = RI.getHWRegChan(Reg);
-    Result.push_back(std::make_pair(Index, Chan));
+    Result.push_back(std::pair(Index, Chan));
   }
   for (; i < 3; ++i)
     Result.push_back(DummyPair);
   return Result;
 }
 
 static std::vector<std::pair<int, unsigned>>
 Swizzle(std::vector<std::pair<int, unsigned>> Src,

llvm/lib/Target/AMDGPU/R600OpenCLImageTypeLoweringPass.cpp

for (unsigned i = 0; i < FT->getNumParams(); ++i) {
       // Add format implicit argument.
       ArgTypes.push_back(ImageFormatType);
       ArgMD[2] = ArgMD[3] = MDString::get(*Context, ImageFormatArgMDType);
       PushArgMD(NewArgMDs, ArgMD);
 
       Modified = true;
     }
     if (!Modified) {
-      return std::make_tuple(nullptr, nullptr);
+      return std::tuple(nullptr, nullptr);
     }
 
     // Create function with new signature and clone the old body into it.
     auto NewFT = FunctionType::get(FT->getReturnType(), ArgTypes, false);
     auto NewF = Function::Create(NewFT, F->getLinkage(), F->getName());
     ValueToValueMapTy VMap;
     auto NewFArgIt = NewF->arg_begin();
     for (auto &Arg: F->args()) {
addImplicitArgs(Function *F, MDNode *KernelMDNode) {
 
     // Build new MDNode.
     SmallVector<Metadata *, 6> KernelMDArgs;
     KernelMDArgs.push_back(ConstantAsMetadata::get(NewF));
     for (const MDVector &MDV : NewArgMDs.ArgVector)
       KernelMDArgs.push_back(MDNode::get(*Context, MDV));
     MDNode *NewMDNode = MDNode::get(*Context, KernelMDArgs);
 
-    return std::make_tuple(NewF, NewMDNode);
+    return std::tuple(NewF, NewMDNode);
   }
 
   bool transformKernels(Module &M) {
     NamedMDNode *KernelsMDNode = M.getNamedMetadata(KernelsMDNodeName);
     if (!KernelsMDNode)
       return false;
 
     bool Modified = false;

llvm/lib/Target/AMDGPU/SIAnnotateControlFlow.cpp

 
 /// Pop the last saved value from the control flow stack
 Value *SIAnnotateControlFlow::popSaved() {
   return Stack.pop_back_val().second;
 }
 
 /// Push a BB and saved value to the control flow stack
 void SIAnnotateControlFlow::push(BasicBlock *BB, Value *Saved) {
-  Stack.push_back(std::make_pair(BB, Saved));
+  Stack.push_back(std::pair(BB, Saved));
 }
 
 /// Can the condition represented by this PHI node treated like
 /// an "Else" block?
 bool SIAnnotateControlFlow::isElse(PHINode *Phi) {
   BasicBlock *IDom = DT->getNode(Phi->getParent())->getIDom()->getBlock();
   for (unsigned i = 0, e = Phi->getNumIncomingValues(); i != e; ++i) {
     if (Phi->getIncomingBlock(i) == IDom) {

llvm/lib/Target/AMDGPU/SIFixSGPRCopies.cpp

getCopyRegClasses(const MachineInstr &Copy,
 
   // We don't really care about the subregister here.
   // SrcRC = TRI.getSubRegClass(SrcRC, Copy.getOperand(1).getSubReg());
 
   const TargetRegisterClass *DstRC = DstReg.isVirtual()
                                          ? MRI.getRegClass(DstReg)
                                          : TRI.getPhysRegClass(DstReg);
 
-  return std::make_pair(SrcRC, DstRC);
+  return std::pair(SrcRC, DstRC);
 }
 
 static bool isVGPRToSGPRCopy(const TargetRegisterClass *SrcRC,
                              const TargetRegisterClass *DstRC,
                              const SIRegisterInfo &TRI) {
   return SrcRC != &AMDGPU::VReg_1RegClass && TRI.isSGPRClass(DstRC) &&
          TRI.hasVectorRegisters(SrcRC);
 }
bool SIFixSGPRCopies::needToBeConvertedToVALU(V2SCopyInfo *Info) {
   for (auto J : Info->Siblings) {
     auto InfoIt = V2SCopies.find(J);
     if (InfoIt != V2SCopies.end()) {
       MachineInstr *SiblingCopy = InfoIt->getSecond().Copy;
       if (SiblingCopy->isImplicitDef())
         // the COPY has already been MoveToVALUed
         continue;
 
-      SrcRegs.insert(std::make_pair(SiblingCopy->getOperand(1).getReg(),
-                                    SiblingCopy->getOperand(1).getSubReg()));
+      SrcRegs.insert(std::pair(SiblingCopy->getOperand(1).getReg(),
     }
   }
   Info->SiblingPenalty = SrcRegs.size();
 
   unsigned Penalty =
       Info->NumSVCopies + Info->SiblingPenalty + Info->NumReadfirstlanes;
   unsigned Profit = Info->SChain.size();
   Info->Score = Penalty > Profit ? 0 : Profit - Penalty;

llvm/lib/Target/AMDGPU/SIFoldOperands.cpp

SIFoldOperands::isOMod(const MachineInstr &MI) const {
   case AMDGPU::V_MUL_F32_e64:
   case AMDGPU::V_MUL_F16_t16_e64:
   case AMDGPU::V_MUL_F16_e64: {
     // If output denormals are enabled, omod is ignored.
     if ((Op == AMDGPU::V_MUL_F32_e64 && MFI->getMode().FP32OutputDenormals) ||
         ((Op == AMDGPU::V_MUL_F64_e64 || Op == AMDGPU::V_MUL_F16_e64 ||
           Op == AMDGPU::V_MUL_F16_t16_e64) &&
          MFI->getMode().FP64FP16OutputDenormals))
-      return std::make_pair(nullptr, SIOutMods::NONE);
+      return std::pair(nullptr, SIOutMods::NONE);
 
     const MachineOperand *RegOp = nullptr;
     const MachineOperand *ImmOp = nullptr;
     const MachineOperand *Src0 = TII->getNamedOperand(MI, AMDGPU::OpName::src0);
     const MachineOperand *Src1 = TII->getNamedOperand(MI, AMDGPU::OpName::src1);
     if (Src0->isImm()) {
       ImmOp = Src0;
       RegOp = Src1;
     } else if (Src1->isImm()) {
       ImmOp = Src1;
       RegOp = Src0;
     } else
-      return std::make_pair(nullptr, SIOutMods::NONE);
+      return std::pair(nullptr, SIOutMods::NONE);
 
     int OMod = getOModValue(Op, ImmOp->getImm());
     if (OMod == SIOutMods::NONE ||
         TII->hasModifiersSet(MI, AMDGPU::OpName::src0_modifiers) ||
         TII->hasModifiersSet(MI, AMDGPU::OpName::src1_modifiers) ||
         TII->hasModifiersSet(MI, AMDGPU::OpName::omod) ||
         TII->hasModifiersSet(MI, AMDGPU::OpName::clamp))
-      return std::make_pair(nullptr, SIOutMods::NONE);
+      return std::pair(nullptr, SIOutMods::NONE);
 
-    return std::make_pair(RegOp, OMod);
+    return std::pair(RegOp, OMod);
   }
   case AMDGPU::V_ADD_F64_e64:
   case AMDGPU::V_ADD_F32_e64:
   case AMDGPU::V_ADD_F16_e64:
   case AMDGPU::V_ADD_F16_t16_e64: {
     // If output denormals are enabled, omod is ignored.
     if ((Op == AMDGPU::V_ADD_F32_e64 && MFI->getMode().FP32OutputDenormals) ||
         ((Op == AMDGPU::V_ADD_F64_e64 || Op == AMDGPU::V_ADD_F16_e64 ||
           Op == AMDGPU::V_ADD_F16_t16_e64) &&
          MFI->getMode().FP64FP16OutputDenormals))
-      return std::make_pair(nullptr, SIOutMods::NONE);
+      return std::pair(nullptr, SIOutMods::NONE);
 
     // Look through the DAGCombiner canonicalization fmul x, 2 -> fadd x, x
     const MachineOperand *Src0 = TII->getNamedOperand(MI, AMDGPU::OpName::src0);
     const MachineOperand *Src1 = TII->getNamedOperand(MI, AMDGPU::OpName::src1);
 
     if (Src0->isReg() && Src1->isReg() && Src0->getReg() == Src1->getReg() &&
         Src0->getSubReg() == Src1->getSubReg() &&
         !TII->hasModifiersSet(MI, AMDGPU::OpName::src0_modifiers) &&
         !TII->hasModifiersSet(MI, AMDGPU::OpName::src1_modifiers) &&
         !TII->hasModifiersSet(MI, AMDGPU::OpName::clamp) &&
         !TII->hasModifiersSet(MI, AMDGPU::OpName::omod))
-      return std::make_pair(Src0, SIOutMods::MUL2);
+      return std::pair(Src0, SIOutMods::MUL2);
 
-    return std::make_pair(nullptr, SIOutMods::NONE);
+    return std::pair(nullptr, SIOutMods::NONE);
   }
   default:
-    return std::make_pair(nullptr, SIOutMods::NONE);
+    return std::pair(nullptr, SIOutMods::NONE);
   }
 }
 
 // FIXME: Does this need to check IEEE bit on function?
 bool SIFoldOperands::tryFoldOMod(MachineInstr &MI) {
   const MachineOperand *RegOp;
   int OMod;
   std::tie(RegOp, OMod) = isOMod(MI);

llvm/lib/Target/AMDGPU/SIFormMemoryClauses.cpp

for (const MachineOperand &MO : MI.operands()) {
     LaneBitmask Mask = Reg.isVirtual()
                            ? TRI->getSubRegIndexLaneMask(MO.getSubReg())
                            : LaneBitmask::getAll();
     RegUse &Map = MO.isDef() ? Defs : Uses;
 
     auto Loc = Map.find(Reg);
     unsigned State = getMopState(MO);
     if (Loc == Map.end()) {
-      Map[Reg] = std::make_pair(State, Mask);
+      Map[Reg] = std::pair(State, Mask);
     } else {
       Loc->second.first |= State;
       Loc->second.second |= Mask;
     }
   }
 }
 
 // Check register def/use conflicts, occupancy limits and collect def/use maps.

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

for (unsigned i = 0, e = ArgLocs.size(); i != e; ++i) {
     case CCValAssign::FPExt:
       Arg = DAG.getNode(ISD::FP_EXTEND, DL, VA.getLocVT(), Arg);
       break;
     default:
       llvm_unreachable("Unknown loc info!");
     }
 
     if (VA.isRegLoc()) {
-      RegsToPass.push_back(std::make_pair(VA.getLocReg(), Arg));
+      RegsToPass.push_back(std::pair(VA.getLocReg(), Arg));
     } else {
       assert(VA.isMemLoc());
 
       SDValue DstAddr;
       MachinePointerInfo DstInfo;
 
       unsigned LocMemOffset = VA.getLocMemOffset();
       int32_t Offset = LocMemOffset;
if (InstInLoop) {
     // Move the rest of the block.
     RemainderBB->splice(RemainderBB->begin(), &MBB, Next, MBB.end());
   } else {
     RemainderBB->splice(RemainderBB->begin(), &MBB, I, MBB.end());
   }
 
   MBB.addSuccessor(LoopBB);
 
-  return std::make_pair(LoopBB, RemainderBB);
+  return std::pair(LoopBB, RemainderBB);
 }
 
 /// Insert \p MI into a BUNDLE with an S_WAITCNT 0 immediately following it.
 void SITargetLowering::bundleInstWithWaitcnt(MachineInstr &MI) const {
   MachineBasicBlock *MBB = MI.getParent();
   const SIInstrInfo *TII = getSubtarget()->getInstrInfo();
   auto I = MI.getIterator();
   auto E = std::next(I);
computeIndirectRegAndOffset(const SIRegisterInfo &TRI,
                             const TargetRegisterClass *SuperRC,
                             unsigned VecReg,
                             int Offset) {
   int NumElts = TRI.getRegSizeInBits(*SuperRC) / 32;
 
   // Skip out of bounds offsets, or else we would end up using an undefined
   // register.
   if (Offset >= NumElts || Offset < 0)
-    return std::make_pair(AMDGPU::sub0, Offset);
+    return std::pair(AMDGPU::sub0, Offset);
 
-  return std::make_pair(SIRegisterInfo::getSubRegFromChannel(Offset), 0);
+  return std::pair(SIRegisterInfo::getSubRegFromChannel(Offset), 0);
 }
 
 static void setM0ToIndexFromSGPR(const SIInstrInfo *TII,
                                  MachineRegisterInfo &MRI, MachineInstr &MI,
                                  int Offset) {
   MachineBasicBlock *MBB = MI.getParent();
   const DebugLoc &DL = MI.getDebugLoc();
   MachineBasicBlock::iterator I(&MI);
SDValue SITargetLowering::splitTernaryVectorOp(SDValue Op,
   assert(VT == MVT::v4i16 || VT == MVT::v4f16 || VT == MVT::v8i16 ||
          VT == MVT::v8f16 || VT == MVT::v4f32 || VT == MVT::v16i16 ||
          VT == MVT::v16f16 || VT == MVT::v8f32 || VT == MVT::v16f32 ||
          VT == MVT::v32f32);
 
   SDValue Lo0, Hi0;
   SDValue Op0 = Op.getOperand(0);
   std::tie(Lo0, Hi0) = Op0.getValueType().isVector()
                            ? DAG.SplitVectorOperand(Op.getNode(), 0)
-                         : std::make_pair(Op0, Op0);
+                           : std::pair(Op0, Op0);
   SDValue Lo1, Hi1;
   std::tie(Lo1, Hi1) = DAG.SplitVectorOperand(Op.getNode(), 1);
   SDValue Lo2, Hi2;
   std::tie(Lo2, Hi2) = DAG.SplitVectorOperand(Op.getNode(), 2);
 
   SDLoc SL(Op);
   auto ResVT = DAG.GetSplitDestVTs(VT);
 
case 'r':
         RC = &AMDGPU::SReg_32RegClass;
         break;
       case 64:
         RC = &AMDGPU::SGPR_64RegClass;
         break;
       default:
         RC = SIRegisterInfo::getSGPRClassForBitWidth(BitWidth);
         if (!RC)
-          return std::make_pair(0U, nullptr);
+          return std::pair(0U, nullptr);
         break;
       }
       break;
     case 'v':
       switch (BitWidth) {
       case 16:
         RC = &AMDGPU::VGPR_32RegClass;
         break;
       default:
         RC = TRI->getVGPRClassForBitWidth(BitWidth);
         if (!RC)
-          return std::make_pair(0U, nullptr);
+          return std::pair(0U, nullptr);
         break;
       }
       break;
     case 'a':
       if (!Subtarget->hasMAIInsts())
         break;
       switch (BitWidth) {
       case 16:
         RC = &AMDGPU::AGPR_32RegClass;
         break;
       default:
         RC = TRI->getAGPRClassForBitWidth(BitWidth);
         if (!RC)
-          return std::make_pair(0U, nullptr);
+          return std::pair(0U, nullptr);
         break;
       }
       break;
     }
     // We actually support i128, i16 and f16 as inline parameters
     // even if they are not reported as legal
     if (RC && (isTypeLegal(VT) || VT.SimpleTy == MVT::i128 ||
                VT.SimpleTy == MVT::i16 || VT.SimpleTy == MVT::f16))
-      return std::make_pair(0U, RC);
+      return std::pair(0U, RC);
   }
 
   if (Constraint.startswith("{") && Constraint.endswith("}")) {
     StringRef RegName(Constraint.data() + 1, Constraint.size() - 2);
     if (RegName.consume_front("v")) {
       RC = &AMDGPU::VGPR_32RegClass;
     } else if (RegName.consume_front("s")) {
       RC = &AMDGPU::SGPR_32RegClass;
if (RC) {
           if (SIRegisterInfo::isVGPRClass(RC))
             RC = TRI->getVGPRClassForBitWidth(Width);
           else if (SIRegisterInfo::isSGPRClass(RC))
             RC = TRI->getSGPRClassForBitWidth(Width);
           else if (SIRegisterInfo::isAGPRClass(RC))
             RC = TRI->getAGPRClassForBitWidth(Width);
           if (RC) {
             Reg = TRI->getMatchingSuperReg(Reg, AMDGPU::sub0, RC);
-            return std::make_pair(Reg, RC);
+            return std::pair(Reg, RC);
           }
         }
       } else {
         bool Failed = RegName.getAsInteger(10, Idx);
         if (!Failed && Idx < RC->getNumRegs())
-          return std::make_pair(RC->getRegister(Idx), RC);
+          return std::pair(RC->getRegister(Idx), RC);
       }
     }
   }
 
   auto Ret = TargetLowering::getRegForInlineAsmConstraint(TRI, Constraint, VT);
   if (Ret.first)
     Ret.second = TRI->getPhysRegClass(Ret.first);
 

llvm/lib/Target/AMDGPU/SIInsertWaitcnts.cpp

for (MachineBasicBlock::instr_iterator Iter = Block.instr_begin(),
     }
 
     if (TII->isSMRD(Inst)) {
       for (const MachineMemOperand *Memop : Inst.memoperands()) {
         // No need to handle invariant loads when avoiding WAR conflicts, as
         // there cannot be a vector store to the same memory location.
         if (!Memop->isInvariant()) {
           const Value *Ptr = Memop->getValue();
-          SLoadAddresses.insert(std::make_pair(Ptr, Inst.getParent()));
+          SLoadAddresses.insert(std::pair(Ptr, Inst.getParent()));
         }
       }
       if (ST->hasReadVCCZBug()) {
         // This smem read could complete and clobber vccz at any time.
         VCCZCorrect = false;
       }
     }
 

llvm/lib/Target/AMDGPU/SIInstrInfo.cpp

 std::pair<MachineInstr*, MachineInstr*>
 SIInstrInfo::expandMovDPP64(MachineInstr &MI) const {
   assert (MI.getOpcode() == AMDGPU::V_MOV_B64_DPP_PSEUDO);
 
   if (ST.hasMovB64() &&
       AMDGPU::isLegal64BitDPPControl(
         getNamedOperand(MI, AMDGPU::OpName::dpp_ctrl)->getImm())) {
     MI.setDesc(get(AMDGPU::V_MOV_B64_dpp));
-    return std::make_pair(&MI, nullptr);
+    return std::pair(&MI, nullptr);
   }
 
   MachineBasicBlock &MBB = *MI.getParent();
   DebugLoc DL = MBB.findDebugLoc(MI);
   MachineFunction *MF = MBB.getParent();
   MachineRegisterInfo &MRI = MF->getRegInfo();
   Register Dst = MI.getOperand(0).getReg();
   unsigned Part = 0;
SIInstrInfo::expandMovDPP64(MachineInstr &MI) const {
   if (Dst.isVirtual())
     BuildMI(MBB, MI, DL, get(AMDGPU::REG_SEQUENCE), Dst)
       .addReg(Split[0]->getOperand(0).getReg())
       .addImm(AMDGPU::sub0)
       .addReg(Split[1]->getOperand(0).getReg())
       .addImm(AMDGPU::sub1);
 
   MI.eraseFromParent();
-  return std::make_pair(Split[0], Split[1]);
+  return std::pair(Split[0], Split[1]);
 }
 
 bool SIInstrInfo::swapSourceModifiers(MachineInstr &MI,
                                       MachineOperand &Src0,
                                       unsigned Src0OpName,
                                       MachineOperand &Src1,
                                       unsigned Src1OpName) const {
   MachineOperand *Src0Mods = getNamedOperand(MI, Src0OpName);
extractRsrcPtr(const SIInstrInfo &TII, MachineInstr &MI, MachineOperand &Rsrc) {
   BuildMI(MBB, MI, MI.getDebugLoc(), TII.get(AMDGPU::REG_SEQUENCE), NewSRsrc)
       .addReg(Zero64)
       .addImm(AMDGPU::sub0_sub1)
       .addReg(SRsrcFormatLo)
       .addImm(AMDGPU::sub2)
       .addReg(SRsrcFormatHi)
       .addImm(AMDGPU::sub3);
 
-  return std::make_tuple(RsrcPtr, NewSRsrc);
+  return std::tuple(RsrcPtr, NewSRsrc);
 }
 
 MachineBasicBlock *
 SIInstrInfo::legalizeOperands(MachineInstr &MI,
                               MachineDominatorTree *MDT) const {
   MachineFunction &MF = *MI.getParent()->getParent();
   MachineRegisterInfo &MRI = MF.getRegInfo();
   MachineBasicBlock *CreatedBB = nullptr;
if (ST.hasAddNoCarry()) {
 
     Inst.setDesc(get(NewOpc));
     Inst.addOperand(MachineOperand::CreateImm(0)); // clamp bit
     Inst.addImplicitDefUseOperands(*MBB.getParent());
     MRI.replaceRegWith(OldDstReg, ResultReg);
     MachineBasicBlock *NewBB = legalizeOperands(Inst, MDT);
 
     addUsersToMoveToVALUWorklist(ResultReg, MRI, Worklist);
-    return std::make_pair(true, NewBB);
+    return std::pair(true, NewBB);
   }
 
-  return std::make_pair(false, nullptr);
+  return std::pair(false, nullptr);
 }
 
 void SIInstrInfo::lowerSelect(SetVectorType &Worklist, MachineInstr &Inst,
                               MachineDominatorTree *MDT) const {
 
   MachineBasicBlock &MBB = *Inst.getParent();
   MachineRegisterInfo &MRI = MBB.getParent()->getRegInfo();
   MachineBasicBlock::iterator MII = Inst;
SIInstrInfo::CreateTargetMIHazardRecognizer(const InstrItineraryData *II,
   // track VRegLiveness.
   if (!DAG->hasVRegLiveness())
     return new GCNHazardRecognizer(DAG->MF);
   return TargetInstrInfo::CreateTargetMIHazardRecognizer(II, DAG);
 }
 
 std::pair<unsigned, unsigned>
 SIInstrInfo::decomposeMachineOperandsTargetFlags(unsigned TF) const {
-  return std::make_pair(TF & MO_MASK, TF & ~MO_MASK);
+  return std::pair(TF & MO_MASK, TF & ~MO_MASK);
 }
 
 ArrayRef<std::pair<unsigned, const char *>>
 SIInstrInfo::getSerializableDirectMachineOperandTargetFlags() const {
   static const std::pair<unsigned, const char *> TargetFlags[] = {
     { MO_GOTPCREL, "amdgpu-gotprel" },
     { MO_GOTPCREL32_LO, "amdgpu-gotprel32-lo" },
     { MO_GOTPCREL32_HI, "amdgpu-gotprel32-hi" },

llvm/lib/Target/AMDGPU/SILoadStoreOptimizer.cpp

SILoadStoreOptimizer::getSubRegIdxs(const CombineInfo &CI,
   if (Paired < CI) {
     Idx1 = Idxs[0][Paired.Width - 1];
     Idx0 = Idxs[Paired.Width][CI.Width - 1];
   } else {
     Idx0 = Idxs[0][CI.Width - 1];
     Idx1 = Idxs[CI.Width][Paired.Width - 1];
   }
 
-  return std::make_pair(Idx0, Idx1);
+  return std::pair(Idx0, Idx1);
 }
 
 const TargetRegisterClass *
 SILoadStoreOptimizer::getTargetRegisterClass(const CombineInfo &CI,
                                              const CombineInfo &Paired) {
   if (CI.InstClass == S_BUFFER_LOAD_IMM ||
       CI.InstClass == S_BUFFER_LOAD_SGPR_IMM || CI.InstClass == S_LOAD_IMM) {
     switch (CI.Width + Paired.Width) {
if (Visited.find(&MINext) == Visited.end()) {
       MAddrNext = Visited[&MINext];
 
     if (MAddrNext.Base.LoReg != MAddr.Base.LoReg ||
         MAddrNext.Base.HiReg != MAddr.Base.HiReg ||
         MAddrNext.Base.LoSubReg != MAddr.Base.LoSubReg ||
         MAddrNext.Base.HiSubReg != MAddr.Base.HiSubReg)
       continue;
 
-    InstsWCommonBase.push_back(std::make_pair(&MINext, MAddrNext.Offset));
+    InstsWCommonBase.push_back(std::pair(&MINext, MAddrNext.Offset));
 
     int64_t Dist = MAddr.Offset - MAddrNext.Offset;
     TargetLoweringBase::AddrMode AM;
     AM.HasBaseReg = true;
     AM.BaseOffs = Dist;
     if (TLI->isLegalGlobalAddressingMode(AM) &&
         (uint32_t)std::abs(Dist) > MaxDist) {
       MaxDist = std::abs(Dist);
for (std::list<std::list<CombineInfo>>::iterator I = MergeableInsts.begin(),
     // matches.
     MergeList.sort(
         [] (const CombineInfo &A, const CombineInfo &B) {
           return A.Offset < B.Offset;
         });
     ++I;
   }
 
-  return std::make_pair(BlockI, Modified);
+  return std::pair(BlockI, Modified);
 }
 
 // Scan through looking for adjacent LDS operations with constant offsets from
 // the same base register. We rely on the scheduler to do the hard work of
 // clustering nearby loads, and assume these are all adjacent.
 bool SILoadStoreOptimizer::optimizeBlock(
                        std::list<std::list<CombineInfo> > &MergeableInsts) {
   bool Modified = false;

llvm/lib/Target/AMDGPU/SIMachineScheduler.cpp

if (SuccID == S.first->getID()) {
       if (S.second == SIScheduleBlockLinkKind::NoData &&
           Kind == SIScheduleBlockLinkKind::Data)
         S.second = Kind;
       return;
     }
   }
   if (Succ->isHighLatencyBlock())
     ++NumHighLatencySuccessors;
-  Succs.push_back(std::make_pair(Succ, Kind));
+  Succs.push_back(std::pair(Succ, Kind));
 
   assert(none_of(Preds,
                  [=](SIScheduleBlock *P) { return SuccID == P->getID(); }) &&
          "Loop in the Block Graph!");
 }
 
 #ifndef NDEBUG
 void SIScheduleBlock::printDebug(bool full) {

llvm/lib/Target/AMDGPU/SIMemoryLegalizer.cpp

void SIMemOpAccess::reportUnsupported(const MachineBasicBlock::iterator &MI,
   DiagnosticInfoUnsupported Diag(Func, Msg, MI->getDebugLoc());
   Func.getContext().diagnose(Diag);
 }
 
 std::optional<std::tuple<SIAtomicScope, SIAtomicAddrSpace, bool>>
 SIMemOpAccess::toSIAtomicScope(SyncScope::ID SSID,
                                SIAtomicAddrSpace InstrAddrSpace) const {
   if (SSID == SyncScope::System)
-    return std::make_tuple(SIAtomicScope::SYSTEM,
-                           SIAtomicAddrSpace::ATOMIC,
-                           true);
+    return std::tuple(SIAtomicScope::SYSTEM, SIAtomicAddrSpace::ATOMIC, true);
   if (SSID == MMI->getAgentSSID())
-    return std::make_tuple(SIAtomicScope::AGENT,
-                           SIAtomicAddrSpace::ATOMIC,
-                           true);
+    return std::tuple(SIAtomicScope::AGENT, SIAtomicAddrSpace::ATOMIC, true);
   if (SSID == MMI->getWorkgroupSSID())
-    return std::make_tuple(SIAtomicScope::WORKGROUP,
-                           SIAtomicAddrSpace::ATOMIC,
-                           true);
+    return std::tuple(SIAtomicScope::WORKGROUP, SIAtomicAddrSpace::ATOMIC,
   if (SSID == MMI->getWavefrontSSID())
-    return std::make_tuple(SIAtomicScope::WAVEFRONT,
-                           SIAtomicAddrSpace::ATOMIC,
-                           true);
+    return std::tuple(SIAtomicScope::WAVEFRONT, SIAtomicAddrSpace::ATOMIC,
   if (SSID == SyncScope::SingleThread)
-    return std::make_tuple(SIAtomicScope::SINGLETHREAD,
-                           SIAtomicAddrSpace::ATOMIC,
-                           true);
+    return std::tuple(SIAtomicScope::SINGLETHREAD, SIAtomicAddrSpace::ATOMIC,
   if (SSID == MMI->getSystemOneAddressSpaceSSID())
-    return std::make_tuple(SIAtomicScope::SYSTEM,
-                           SIAtomicAddrSpace::ATOMIC & InstrAddrSpace,
-                           false);
+    return std::tuple(SIAtomicScope::SYSTEM,
+                      SIAtomicAddrSpace::ATOMIC & InstrAddrSpace, false);
   if (SSID == MMI->getAgentOneAddressSpaceSSID())
-    return std::make_tuple(SIAtomicScope::AGENT,
-                           SIAtomicAddrSpace::ATOMIC & InstrAddrSpace,
-                           false);
+    return std::tuple(SIAtomicScope::AGENT,
+                      SIAtomicAddrSpace::ATOMIC & InstrAddrSpace, false);
   if (SSID == MMI->getWorkgroupOneAddressSpaceSSID())
-    return std::make_tuple(SIAtomicScope::WORKGROUP,
-                           SIAtomicAddrSpace::ATOMIC & InstrAddrSpace,
-                           false);
+    return std::tuple(SIAtomicScope::WORKGROUP,
+                      SIAtomicAddrSpace::ATOMIC & InstrAddrSpace, false);
   if (SSID == MMI->getWavefrontOneAddressSpaceSSID())
-    return std::make_tuple(SIAtomicScope::WAVEFRONT,
-                           SIAtomicAddrSpace::ATOMIC & InstrAddrSpace,
-                           false);
+    return std::tuple(SIAtomicScope::WAVEFRONT,
+                      SIAtomicAddrSpace::ATOMIC & InstrAddrSpace, false);
   if (SSID == MMI->getSingleThreadOneAddressSpaceSSID())
-    return std::make_tuple(SIAtomicScope::SINGLETHREAD,
-                           SIAtomicAddrSpace::ATOMIC & InstrAddrSpace,
-                           false);
+    return std::tuple(SIAtomicScope::SINGLETHREAD,
+                      SIAtomicAddrSpace::ATOMIC & InstrAddrSpace, false);
   return std::nullopt;
 }
 
 SIAtomicAddrSpace SIMemOpAccess::toSIAtomicAddrSpace(unsigned AS) const {
   if (AS == AMDGPUAS::FLAT_ADDRESS)
     return SIAtomicAddrSpace::FLAT;
   if (AS == AMDGPUAS::GLOBAL_ADDRESS)
     return SIAtomicAddrSpace::GLOBAL;

llvm/lib/Target/AMDGPU/SIPeepholeSDWA.cpp

auto CheckOROperandsForSDWA =
         MachineInstr *Op1Inst = Op1Def->getParent();
         if (!TII->isSDWA(*Op1Inst))
           return CheckRetType(std::nullopt);
 
         MachineOperand *Op2Def = findSingleRegDef(Op2, MRI);
         if (!Op2Def)
           return CheckRetType(std::nullopt);
 
-        return CheckRetType(std::make_pair(Op1Def, Op2Def));
+        return CheckRetType(std::pair(Op1Def, Op2Def));
       };
 
     MachineOperand *OrSDWA = TII->getNamedOperand(MI, AMDGPU::OpName::src0);
     MachineOperand *OrOther = TII->getNamedOperand(MI, AMDGPU::OpName::src1);
     assert(OrSDWA && OrOther);
     auto Res = CheckOROperandsForSDWA(OrSDWA, OrOther);
     if (!Res) {
       OrSDWA = TII->getNamedOperand(MI, AMDGPU::OpName::src1);

llvm/lib/Target/AMDGPU/SIWholeQuadMode.cpp

public:
 
   explicit PrintState(int State) : State(State) {}
 };
 
 #ifndef NDEBUG
 static raw_ostream &operator<<(raw_ostream &OS, const PrintState &PS) {
 
   static const std::pair<char, const char *> Mapping[] = {
-      std::make_pair(StateWQM, "WQM"),
-      std::make_pair(StateStrictWWM, "StrictWWM"),
-      std::make_pair(StateStrictWQM, "StrictWQM"),
-      std::make_pair(StateExact, "Exact")};
+      std::pair(StateWQM, "WQM"), std::pair(StateStrictWWM, "StrictWWM"),
+      std::pair(StateStrictWQM, "StrictWQM"), std::pair(StateExact, "Exact")};
   char State = PS.State;
   for (auto M : Mapping) {
     if (State & M.first) {
       OS << M.second;
       State &= ~M.first;
 
       if (State)
         OS << '|';

llvm/lib/Target/AMDGPU/Utils/AMDKernelCodeTUtils.cpp

#undef RECORD
   return makeArrayRef(Table);
 }
 
 static StringMap<int> createIndexMap(const ArrayRef<StringRef> &names,
                                      const ArrayRef<StringRef> &altNames) {
   StringMap<int> map;
   assert(names.size() == altNames.size());
   for (unsigned i = 0; i < names.size(); ++i) {
-    map.insert(std::make_pair(names[i], i));
-    map.insert(std::make_pair(altNames[i], i));
+    map.insert(std::pair(names[i], i));
+    map.insert(std::pair(altNames[i], i));
   }
   return map;
 }
 
 static int get_amd_kernel_code_t_FieldIndex(StringRef name) {
   static const auto map = createIndexMap(get_amd_kernel_code_t_FldNames(),
                                          get_amd_kernel_code_t_FldAltNames());
   return map.lookup(name) - 1; // returns -1 if not found