diff --git a/llvm/lib/Target/AMDGPU/AMDGPUAsanInstrumentation.cpp b/llvm/lib/Target/AMDGPU/AMDGPUAsanInstrumentation.cpp
index 4c8ddbd9aabd5a..a6a8597f22a4a2 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUAsanInstrumentation.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUAsanInstrumentation.cpp
@@ -235,7 +235,7 @@ void getInterestingMemoryOperands(
     Interesting.emplace_back(I, XCHG->getPointerOperandIndex(), true,
                              XCHG->getCompareOperand()->getType(),
                              std::nullopt);
-  } else if (auto CI = dyn_cast<CallInst>(I)) {
+  } else if (auto *CI = dyn_cast<CallInst>(I)) {
     switch (CI->getIntrinsicID()) {
     case Intrinsic::masked_load:
     case Intrinsic::masked_store:
@@ -257,7 +257,7 @@ void getInterestingMemoryOperands(
     case Intrinsic::masked_compressstore: {
       bool IsWrite = CI->getIntrinsicID() == Intrinsic::masked_compressstore;
       unsigned OpOffset = IsWrite ? 1 : 0;
-      auto BasePtr = CI->getOperand(OpOffset);
+      auto *BasePtr = CI->getOperand(OpOffset);
       MaybeAlign Alignment = BasePtr->getPointerAlignment(DL);
       Type *Ty = IsWrite ? CI->getArgOperand(0)->getType() : CI->getType();
       IRBuilder<> IB(I);
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp b/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
index 974f1c80d95530..93b6ba0595b70f 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
@@ -332,7 +332,7 @@ void AMDGPUAsmPrinter::emitGlobalVariable(const GlobalVariable *GV) {
 
     emitVisibility(GVSym, GV->getVisibility(), !GV->isDeclaration());
     emitLinkage(GV, GVSym);
-    auto TS = getTargetStreamer();
+    auto *TS = getTargetStreamer();
     TS->emitAMDGPULDS(GVSym, Size, Alignment);
     return;
   }
@@ -1238,8 +1238,8 @@ void AMDGPUAsmPrinter::getSIProgramInfo(SIProgramInfo &ProgInfo,
     // return ((Dst & ~Mask) | (Value << Shift))
     auto SetBits = [&Ctx](const MCExpr *Dst, const MCExpr *Value, uint32_t Mask,
                           uint32_t Shift) {
-      auto Shft = MCConstantExpr::create(Shift, Ctx);
-      auto Msk = MCConstantExpr::create(Mask, Ctx);
+      const auto *Shft = MCConstantExpr::create(Shift, Ctx);
+      const auto *Msk = MCConstantExpr::create(Mask, Ctx);
       Dst = MCBinaryExpr::createAnd(Dst, MCUnaryExpr::createNot(Msk, Ctx), Ctx);
       Dst = MCBinaryExpr::createOr(
           Dst, MCBinaryExpr::createShl(Value, Shft, Ctx), Ctx);
@@ -1414,7 +1414,7 @@ void AMDGPUAsmPrinter::EmitPALMetadata(const MachineFunction &MF,
        const SIProgramInfo &CurrentProgramInfo) {
   const SIMachineFunctionInfo *MFI = MF.getInfo<SIMachineFunctionInfo>();
   auto CC = MF.getFunction().getCallingConv();
-  auto MD = getTargetStreamer()->getPALMetadata();
+  auto *MD = getTargetStreamer()->getPALMetadata();
   auto &Ctx = MF.getContext();
 
   MD->setEntryPoint(CC, MF.getFunction().getName());
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp b/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp
index 25e36dc4b3691f..9807e31c1c7f83 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp
@@ -142,7 +142,7 @@ struct AMDGPUIncomingArgHandler : public CallLowering::IncomingValueHandler {
                             const CCValAssign &VA) override {
     MachineFunction &MF = MIRBuilder.getMF();
 
-    auto MMO = MF.getMachineMemOperand(
+    auto *MMO = MF.getMachineMemOperand(
         MPO, MachineMemOperand::MOLoad | MachineMemOperand::MOInvariant, MemTy,
         inferAlignFromPtrInfo(MF, MPO));
     MIRBuilder.buildLoad(ValVReg, Addr, *MMO);
@@ -244,7 +244,7 @@ struct AMDGPUOutgoingArgHandler : public AMDGPUOutgoingValueHandler {
     uint64_t LocMemOffset = VA.getLocMemOffset();
     const auto &ST = MF.getSubtarget<GCNSubtarget>();
 
-    auto MMO = MF.getMachineMemOperand(
+    auto *MMO = MF.getMachineMemOperand(
         MPO, MachineMemOperand::MOStore, MemTy,
         commonAlignment(ST.getStackAlignment(), LocMemOffset));
     MIRBuilder.buildStore(ValVReg, Addr, *MMO);
@@ -1007,7 +1007,7 @@ bool AMDGPUCallLowering::doCallerAndCalleePassArgsTheSameWay(
   const GCNSubtarget &ST = MF.getSubtarget<GCNSubtarget>();
 
   // Make sure that the caller and callee preserve all of the same registers.
-  auto TRI = ST.getRegisterInfo();
+  const auto *TRI = ST.getRegisterInfo();
 
   const uint32_t *CallerPreserved = TRI->getCallPreservedMask(MF, CallerCC);
   const uint32_t *CalleePreserved = TRI->getCallPreservedMask(MF, CalleeCC);
@@ -1219,7 +1219,7 @@ bool AMDGPUCallLowering::lowerTailCall(
     if (!ExecArg.Ty->isIntegerTy(ST.getWavefrontSize()))
       return false;
 
-    if (auto CI = dyn_cast<ConstantInt>(ExecArg.OrigValue)) {
+    if (const auto *CI = dyn_cast<ConstantInt>(ExecArg.OrigValue)) {
       MIB.addImm(CI->getSExtValue());
     } else {
       MIB.addReg(ExecArg.Regs[0]);
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp b/llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp
index 26116bfa3c2feb..b67d78e450bb82 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp
@@ -163,8 +163,8 @@ std::string MetadataStreamerMsgPackV4::getTypeName(Type *Ty,
   case Type::DoubleTyID:
     return "double";
   case Type::FixedVectorTyID: {
-    auto VecTy = cast<FixedVectorType>(Ty);
-    auto ElTy = VecTy->getElementType();
+    auto *VecTy = cast<FixedVectorType>(Ty);
+    auto *ElTy = VecTy->getElementType();
     auto NumElements = VecTy->getNumElements();
     return (Twine(getTypeName(ElTy, Signed)) + Twine(NumElements)).str();
   }
@@ -199,7 +199,7 @@ void MetadataStreamerMsgPackV4::emitTargetID(
 }
 
 void MetadataStreamerMsgPackV4::emitPrintf(const Module &Mod) {
-  auto Node = Mod.getNamedMetadata("llvm.printf.fmts");
+  auto *Node = Mod.getNamedMetadata("llvm.printf.fmts");
   if (!Node)
     return;
 
@@ -214,10 +214,10 @@ void MetadataStreamerMsgPackV4::emitPrintf(const Module &Mod) {
 void MetadataStreamerMsgPackV4::emitKernelLanguage(const Function &Func,
                                                    msgpack::MapDocNode Kern) {
   // TODO: What about other languages?
-  auto Node = Func.getParent()->getNamedMetadata("opencl.ocl.version");
+  auto *Node = Func.getParent()->getNamedMetadata("opencl.ocl.version");
   if (!Node || !Node->getNumOperands())
     return;
-  auto Op0 = Node->getOperand(0);
+  auto *Op0 = Node->getOperand(0);
   if (Op0->getNumOperands() <= 1)
     return;
 
@@ -233,11 +233,11 @@ void MetadataStreamerMsgPackV4::emitKernelLanguage(const Function &Func,
 void MetadataStreamerMsgPackV4::emitKernelAttrs(const Function &Func,
                                                 msgpack::MapDocNode Kern) {
 
-  if (auto Node = Func.getMetadata("reqd_work_group_size"))
+  if (auto *Node = Func.getMetadata("reqd_work_group_size"))
     Kern[".reqd_workgroup_size"] = getWorkGroupDimensions(Node);
-  if (auto Node = Func.getMetadata("work_group_size_hint"))
+  if (auto *Node = Func.getMetadata("work_group_size_hint"))
     Kern[".workgroup_size_hint"] = getWorkGroupDimensions(Node);
-  if (auto Node = Func.getMetadata("vec_type_hint")) {
+  if (auto *Node = Func.getMetadata("vec_type_hint")) {
     Kern[".vec_type_hint"] = Kern.getDocument()->getNode(
         getTypeName(
             cast<ValueAsMetadata>(Node->getOperand(0))->getType(),
@@ -271,7 +271,7 @@ void MetadataStreamerMsgPackV4::emitKernelArgs(const MachineFunction &MF,
 void MetadataStreamerMsgPackV4::emitKernelArg(const Argument &Arg,
                                               unsigned &Offset,
                                               msgpack::ArrayDocNode Args) {
-  auto Func = Arg.getParent();
+  const auto *Func = Arg.getParent();
   auto ArgNo = Arg.getArgNo();
   const MDNode *Node;
 
@@ -317,7 +317,7 @@ void MetadataStreamerMsgPackV4::emitKernelArg(const Argument &Arg,
   Type *Ty = Arg.hasByRefAttr() ? Arg.getParamByRefType() : Arg.getType();
 
   // FIXME: Need to distinguish in memory alignment from pointer alignment.
-  if (auto PtrTy = dyn_cast<PointerType>(Ty)) {
+  if (auto *PtrTy = dyn_cast<PointerType>(Ty)) {
     if (PtrTy->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS)
       PointeeAlign = Arg.getParamAlign().valueOrOne();
   }
@@ -353,7 +353,7 @@ void MetadataStreamerMsgPackV4::emitKernelArg(
   if (PointeeAlign)
     Arg[".pointee_align"] = Arg.getDocument()->getNode(PointeeAlign->value());
 
-  if (auto PtrTy = dyn_cast<PointerType>(Ty))
+  if (auto *PtrTy = dyn_cast<PointerType>(Ty))
     if (auto Qualifier = getAddressSpaceQualifier(PtrTy->getAddressSpace()))
       // Limiting address space to emit only for a certain ValueKind.
       if (ValueKind == "global_buffer" || ValueKind == "dynamic_shared_pointer")
@@ -393,7 +393,7 @@ void MetadataStreamerMsgPackV4::emitHiddenKernelArgs(
 
   const Module *M = Func.getParent();
   auto &DL = M->getDataLayout();
-  auto Int64Ty = Type::getInt64Ty(Func.getContext());
+  auto *Int64Ty = Type::getInt64Ty(Func.getContext());
 
   Offset = alignTo(Offset, ST.getAlignmentForImplicitArgPtr());
 
@@ -407,7 +407,7 @@ void MetadataStreamerMsgPackV4::emitHiddenKernelArgs(
     emitKernelArg(DL, Int64Ty, Align(8), "hidden_global_offset_z", Offset,
                   Args);
 
-  auto Int8PtrTy =
+  auto *Int8PtrTy =
       PointerType::get(Func.getContext(), AMDGPUAS::GLOBAL_ADDRESS);
 
   if (HiddenArgNumBytes >= 32) {
@@ -592,9 +592,9 @@ void MetadataStreamerMsgPackV5::emitHiddenKernelArgs(
   auto &DL = M->getDataLayout();
   const SIMachineFunctionInfo &MFI = *MF.getInfo<SIMachineFunctionInfo>();
 
-  auto Int64Ty = Type::getInt64Ty(Func.getContext());
-  auto Int32Ty = Type::getInt32Ty(Func.getContext());
-  auto Int16Ty = Type::getInt16Ty(Func.getContext());
+  auto *Int64Ty = Type::getInt64Ty(Func.getContext());
+  auto *Int32Ty = Type::getInt32Ty(Func.getContext());
+  auto *Int16Ty = Type::getInt16Ty(Func.getContext());
 
   Offset = alignTo(Offset, ST.getAlignmentForImplicitArgPtr());
   emitKernelArg(DL, Int32Ty, Align(4), "hidden_block_count_x", Offset, Args);
@@ -621,7 +621,7 @@ void MetadataStreamerMsgPackV5::emitHiddenKernelArgs(
   emitKernelArg(DL, Int16Ty, Align(2), "hidden_grid_dims", Offset, Args);
 
   Offset += 6; // Reserved.
-  auto Int8PtrTy =
+  auto *Int8PtrTy =
       PointerType::get(Func.getContext(), AMDGPUAS::GLOBAL_ADDRESS);
 
   if (M->getNamedMetadata("llvm.printf.fmts")) {
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp b/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp
index aaad2a56de2cd0..a367db70627748 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp
@@ -394,7 +394,7 @@ void PipelineSolver::reset() {
     for (auto &SG : SyncPipeline) {
       SmallVector<SUnit *, 32> TempCollection = SG.Collection;
       SG.Collection.clear();
-      auto SchedBarr = llvm::find_if(TempCollection, [](SUnit *SU) {
+      auto *SchedBarr = llvm::find_if(TempCollection, [](SUnit *SU) {
         return SU->getInstr()->getOpcode() == AMDGPU::SCHED_GROUP_BARRIER;
       });
       if (SchedBarr != TempCollection.end())
@@ -421,7 +421,7 @@ void PipelineSolver::convertSyncMapsToArrays() {
             std::pair(SUsToCandSGs.first, SUsToCandSGs.second));
         continue;
       }
-      auto SortPosition = PipelineInstrs[PipelineIDx].begin();
+      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() &&
@@ -515,7 +515,7 @@ void PipelineSolver::removeEdges(
     SUnit *Pred = PredSuccPair.first;
     SUnit *Succ = PredSuccPair.second;
 
-    auto Match = llvm::find_if(
+    auto *Match = llvm::find_if(
         Succ->Preds, [&Pred](SDep &P) { return P.getSUnit() == Pred; });
     if (Match != Succ->Preds.end()) {
       assert(Match->isArtificial());
@@ -639,8 +639,8 @@ bool PipelineSolver::solveExact() {
              : populateReadyList(ReadyList, CurrSU.second.begin(),
                                  CurrSU.second.end());
 
-  auto I = ReadyList.begin();
-  auto E = ReadyList.end();
+  auto *I = ReadyList.begin();
+  auto *E = ReadyList.end();
   for (; I != E; ++I) {
     // If we are trying SGs in least cost order, and the current SG is cost
     // infeasible, then all subsequent SGs will also be cost infeasible, so we
@@ -942,7 +942,7 @@ class MFMAExpInterleaveOpt final : public IGLPStrategy {
     bool apply(const SUnit *SU, const ArrayRef<SUnit *> Collection,
                SmallVectorImpl<SchedGroup> &SyncPipe) override {
 
-      auto DAG = SyncPipe[0].DAG;
+      auto *DAG = SyncPipe[0].DAG;
 
       if (Cache->empty()) {
         auto I = DAG->SUnits.rbegin();
@@ -976,7 +976,7 @@ class MFMAExpInterleaveOpt final : public IGLPStrategy {
                SmallVectorImpl<SchedGroup> &SyncPipe) override {
       bool FoundTrans = false;
       unsigned Counter = 1;
-      auto DAG = SyncPipe[0].DAG;
+      auto *DAG = SyncPipe[0].DAG;
 
       if (Cache->empty()) {
         SmallVector<SUnit *, 8> Worklist;
@@ -1016,13 +1016,13 @@ class MFMAExpInterleaveOpt final : public IGLPStrategy {
   public:
     bool apply(const SUnit *SU, const ArrayRef<SUnit *> Collection,
                SmallVectorImpl<SchedGroup> &SyncPipe) override {
-      auto DAG = SyncPipe[0].DAG;
+      auto *DAG = SyncPipe[0].DAG;
 
       if (!SU || !TII->isMFMAorWMMA(*ChainSeed->getInstr()))
         return false;
 
       if (Cache->empty()) {
-        auto TempSU = ChainSeed;
+        auto *TempSU = ChainSeed;
         auto Depth = Number;
         while (Depth > 0) {
           --Depth;
@@ -1232,7 +1232,7 @@ class MFMAExpInterleaveOpt final : public IGLPStrategy {
       if (!OtherGroup->Collection.size())
         return true;
 
-      auto DAG = SyncPipe[0].DAG;
+      auto *DAG = SyncPipe[0].DAG;
 
       for (auto &OtherEle : OtherGroup->Collection)
         if (DAG->IsReachable(const_cast<SUnit *>(SU), OtherEle))
@@ -1275,7 +1275,7 @@ class MFMAExpInterleaveOpt final : public IGLPStrategy {
         return false;
 
       if (Cache->empty()) {
-        auto TempSU = ChainSeed;
+        auto *TempSU = ChainSeed;
         auto Depth = Number;
         while (Depth > 0) {
           --Depth;
@@ -1315,7 +1315,7 @@ class MFMAExpInterleaveOpt final : public IGLPStrategy {
                SmallVectorImpl<SchedGroup> &SyncPipe) override {
 
       SmallVector<SUnit *, 12> Worklist;
-      auto DAG = SyncPipe[0].DAG;
+      auto *DAG = SyncPipe[0].DAG;
       if (Cache->empty()) {
         for (auto &SU : DAG->SUnits)
           if (TII->isTRANS(SU.getInstr()->getOpcode())) {
@@ -1509,7 +1509,7 @@ bool MFMAExpInterleaveOpt::analyzeDAG(const SIInstrInfo *TII) {
                       return isBitPack(Opc);
                     });
 
-  auto PackPred =
+  auto *PackPred =
       std::find_if((*TempMFMA)->Preds.begin(), (*TempMFMA)->Preds.end(),
                    [&isBitPack](SDep &Pred) {
                      auto Opc = Pred.getSUnit()->getInstr()->getOpcode();
@@ -1868,7 +1868,7 @@ class MFMASmallGemmSingleWaveOpt final : public IGLPStrategy {
       }
 
       assert(Cache->size());
-      auto DAG = SyncPipe[0].DAG;
+      auto *DAG = SyncPipe[0].DAG;
       for (auto &Elt : *Cache) {
         if (DAG->IsReachable(Elt, const_cast<SUnit *>(SU)))
           return true;
@@ -1886,7 +1886,7 @@ class MFMASmallGemmSingleWaveOpt final : public IGLPStrategy {
   public:
     bool apply(const SUnit *SU, const ArrayRef<SUnit *> Collection,
                SmallVectorImpl<SchedGroup> &SyncPipe) override {
-      auto MI = SU->getInstr();
+      auto *MI = SU->getInstr();
       if (MI->getOpcode() != AMDGPU::V_PERM_B32_e64)
         return false;
 
@@ -1952,7 +1952,7 @@ class MFMASmallGemmSingleWaveOpt final : public IGLPStrategy {
   public:
     bool apply(const SUnit *SU, const ArrayRef<SUnit *> Collection,
                SmallVectorImpl<SchedGroup> &SyncPipe) override {
-      auto MI = SU->getInstr();
+      auto *MI = SU->getInstr();
       if (MI->getOpcode() == TargetOpcode::BUNDLE)
         return false;
       if (!Collection.size())
@@ -2023,7 +2023,7 @@ class MFMASmallGemmSingleWaveOpt final : public IGLPStrategy {
           return false;
       }
 
-      auto DAG = SyncPipe[0].DAG;
+      auto *DAG = SyncPipe[0].DAG;
       // Does the previous DS_WRITE share a V_PERM predecessor with this
       // VMEM_READ
       return llvm::any_of(*Cache, [&SU, &DAG](SUnit *Elt) {
@@ -2070,7 +2070,7 @@ bool MFMASmallGemmSingleWaveOpt::applyIGLPStrategy(
          "DSWCounters should be zero in pre-RA scheduling!");
   SmallVector<SUnit *, 6> DSWithPerms;
   for (auto &SU : DAG->SUnits) {
-    auto I = SU.getInstr();
+    auto *I = SU.getInstr();
     if (TII->isMFMAorWMMA(*I))
       ++MFMACount;
     else if (TII->isDS(*I)) {
@@ -2091,8 +2091,8 @@ bool MFMASmallGemmSingleWaveOpt::applyIGLPStrategy(
 
   if (IsInitial) {
     DSWWithPermCount = DSWithPerms.size();
-    auto I = DSWithPerms.begin();
-    auto E = DSWithPerms.end();
+    auto *I = DSWithPerms.begin();
+    auto *E = DSWithPerms.end();
 
     // Get the count of DS_WRITES with V_PERM predecessors which
     // have loop carried dependencies (WAR) on the same VMEM_READs.
@@ -2113,7 +2113,7 @@ bool MFMASmallGemmSingleWaveOpt::applyIGLPStrategy(
           break;
 
         for (auto &Succ : Pred.getSUnit()->Succs) {
-          auto MI = Succ.getSUnit()->getInstr();
+          auto *MI = Succ.getSUnit()->getInstr();
           if (!TII->isVMEM(*MI) || !MI->mayLoad())
             continue;
 
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp b/llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp
index ff8798edb3cc0f..21fffba14287ef 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp
@@ -1537,7 +1537,7 @@ static bool IsCopyFromSGPR(const SIRegisterInfo &TRI, SDValue Val) {
   auto Reg = cast<RegisterSDNode>(Val.getOperand(1))->getReg();
   if (!Reg.isPhysical())
     return false;
-  auto RC = TRI.getPhysRegBaseClass(Reg);
+  const auto *RC = TRI.getPhysRegBaseClass(Reg);
   return RC && TRI.isSGPRClass(RC);
 }
 
@@ -1855,13 +1855,13 @@ bool AMDGPUDAGToDAGISel::SelectGlobalSAddr(SDNode *N,
 }
 
 static SDValue SelectSAddrFI(SelectionDAG *CurDAG, SDValue SAddr) {
-  if (auto FI = dyn_cast<FrameIndexSDNode>(SAddr)) {
+  if (auto *FI = dyn_cast<FrameIndexSDNode>(SAddr)) {
     SAddr = CurDAG->getTargetFrameIndex(FI->getIndex(), FI->getValueType(0));
   } else if (SAddr.getOpcode() == ISD::ADD &&
              isa<FrameIndexSDNode>(SAddr.getOperand(0))) {
     // Materialize this into a scalar move for scalar address to avoid
     // readfirstlane.
-    auto FI = cast<FrameIndexSDNode>(SAddr.getOperand(0));
+    auto *FI = cast<FrameIndexSDNode>(SAddr.getOperand(0));
     SDValue TFI = CurDAG->getTargetFrameIndex(FI->getIndex(),
                                               FI->getValueType(0));
     SAddr = SDValue(CurDAG->getMachineNode(AMDGPU::S_ADD_I32, SDLoc(SAddr),
@@ -2391,7 +2391,7 @@ bool AMDGPUDAGToDAGISel::isCBranchSCC(const SDNode *N) const {
     return true;
 
   if (VT == MVT::i64) {
-    auto ST = static_cast<const GCNSubtarget *>(Subtarget);
+    const auto *ST = static_cast<const GCNSubtarget *>(Subtarget);
 
     ISD::CondCode CC = cast<CondCodeSDNode>(Cond.getOperand(2))->get();
     return (CC == ISD::SETEQ || CC == ISD::SETNE) && ST->hasScalarCompareEq64();
@@ -3600,7 +3600,7 @@ bool AMDGPUDAGToDAGISel::isVGPRImm(const SDNode * N) const {
 }
 
 bool AMDGPUDAGToDAGISel::isUniformLoad(const SDNode *N) const {
-  auto Ld = cast<LoadSDNode>(N);
+  const auto *Ld = cast<LoadSDNode>(N);
 
   const MachineMemOperand *MMO = Ld->getMemOperand();
   if (N->isDivergent() && !AMDGPUInstrInfo::isUniformMMO(MMO))
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp b/llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp
index 94fdf4effa10a1..cceb89e23f1290 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp
@@ -4202,7 +4202,7 @@ SDValue AMDGPUTargetLowering::performTruncateCombine(
   // integer operation.
   // trunc (srl (bitcast (build_vector x, y))), 16 -> trunc (bitcast y)
   if (Src.getOpcode() == ISD::SRL && !VT.isVector()) {
-    if (auto K = isConstOrConstSplat(Src.getOperand(1))) {
+    if (auto *K = isConstOrConstSplat(Src.getOperand(1))) {
       if (2 * K->getZExtValue() == Src.getValueType().getScalarSizeInBits()) {
         SDValue BV = stripBitcast(Src.getOperand(0));
         if (BV.getOpcode() == ISD::BUILD_VECTOR &&
@@ -5779,7 +5779,7 @@ void AMDGPUTargetLowering::computeKnownBitsForTargetNode(
     break;
   }
   case AMDGPUISD::LDS: {
-    auto GA = cast<GlobalAddressSDNode>(Op.getOperand(0).getNode());
+    auto *GA = cast<GlobalAddressSDNode>(Op.getOperand(0).getNode());
     Align Alignment = GA->getGlobal()->getPointerAlignment(DAG.getDataLayout());
 
     Known.Zero.setHighBits(16);
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUImageIntrinsicOptimizer.cpp b/llvm/lib/Target/AMDGPU/AMDGPUImageIntrinsicOptimizer.cpp
index 88429e3f0e2181..45207c06a788a2 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUImageIntrinsicOptimizer.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUImageIntrinsicOptimizer.cpp
@@ -116,8 +116,8 @@ void addInstToMergeableList(
       Value *Arg = II->getArgOperand(I);
       if (I == ImageDimIntr->VAddrEnd - 1) {
         // Check FragId group.
-        auto FragIdList = cast<ConstantInt>(IIList.front()->getArgOperand(I));
-        auto FragId = cast<ConstantInt>(II->getArgOperand(I));
+        auto *FragIdList = cast<ConstantInt>(IIList.front()->getArgOperand(I));
+        auto *FragId = cast<ConstantInt>(II->getArgOperand(I));
         AllEqual = FragIdList->getValue().udiv(4) == FragId->getValue().udiv(4);
       } else {
         // Check all arguments except FragId.
@@ -219,7 +219,8 @@ bool optimizeSection(ArrayRef<SmallVector<IntrinsicInst *, 4>> MergeableInsts) {
       continue;
 
     const uint8_t FragIdIndex = ImageDimIntr->VAddrEnd - 1;
-    auto FragId = cast<ConstantInt>(IIList.front()->getArgOperand(FragIdIndex));
+    auto *FragId =
+        cast<ConstantInt>(IIList.front()->getArgOperand(FragIdIndex));
     const APInt &NewFragIdVal = FragId->getValue().udiv(4) * 4;
 
     // Create the new instructions.
@@ -251,7 +252,7 @@ bool optimizeSection(ArrayRef<SmallVector<IntrinsicInst *, 4>> MergeableInsts) {
     // Create the new extractelement instructions.
     for (auto &II : IIList) {
       Value *VecOp = nullptr;
-      auto Idx = cast<ConstantInt>(II->getArgOperand(FragIdIndex));
+      auto *Idx = cast<ConstantInt>(II->getArgOperand(FragIdIndex));
       B.SetCurrentDebugLocation(II->getDebugLoc());
       if (NumElts == 1) {
         VecOp = B.CreateExtractElement(NewCalls[0], Idx->getValue().urem(4));
@@ -275,7 +276,7 @@ bool optimizeSection(ArrayRef<SmallVector<IntrinsicInst *, 4>> MergeableInsts) {
     Modified = true;
   }
 
-  for (auto I : InstrsToErase)
+  for (auto *I : InstrsToErase)
     I->eraseFromParent();
 
   return Modified;
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp b/llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp
index e8674c4c775950..ecb4d4fa5d5c39 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp
@@ -144,7 +144,7 @@ static std::optional<Instruction *> modifyIntrinsicCall(
 
   bool RemoveOldIntr = &OldIntr != &InstToReplace;
 
-  auto RetValue = IC.eraseInstFromFunction(InstToReplace);
+  auto *RetValue = IC.eraseInstFromFunction(InstToReplace);
   if (RemoveOldIntr)
     IC.eraseInstFromFunction(OldIntr);
 
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp b/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp
index febf0711c7d574..15bdd9ae293a12 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp
@@ -1474,8 +1474,8 @@ bool AMDGPUInstructionSelector::selectRelocConstant(MachineInstr &I) const {
   Module *M = MF->getFunction().getParent();
   const MDNode *Metadata = I.getOperand(2).getMetadata();
   auto SymbolName = cast<MDString>(Metadata->getOperand(0))->getString();
-  auto RelocSymbol = cast<GlobalVariable>(
-    M->getOrInsertGlobal(SymbolName, Type::getInt32Ty(M->getContext())));
+  auto *RelocSymbol = cast<GlobalVariable>(
+      M->getOrInsertGlobal(SymbolName, Type::getInt32Ty(M->getContext())));
 
   MachineBasicBlock *BB = I.getParent();
   BuildMI(*BB, &I, I.getDebugLoc(),
diff --git a/llvm/lib/Target/AMDGPU/AMDGPULateCodeGenPrepare.cpp b/llvm/lib/Target/AMDGPU/AMDGPULateCodeGenPrepare.cpp
index 3767a6b379f81c..77350dbb6167cb 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPULateCodeGenPrepare.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPULateCodeGenPrepare.cpp
@@ -111,7 +111,7 @@ class LiveRegOptimizer {
     if (!VTy)
       return false;
 
-    auto TLI = ST->getTargetLowering();
+    const auto *TLI = ST->getTargetLowering();
 
     Type *EltTy = VTy->getElementType();
     // If the element size is not less than the convert to scalar size, then we
@@ -454,7 +454,7 @@ bool AMDGPULateCodeGenPrepare::visitLoadInst(LoadInst &LI) {
   IRB.SetCurrentDebugLocation(LI.getDebugLoc());
 
   unsigned LdBits = DL->getTypeStoreSizeInBits(LI.getType());
-  auto IntNTy = Type::getIntNTy(LI.getContext(), LdBits);
+  auto *IntNTy = Type::getIntNTy(LI.getContext(), LdBits);
 
   auto *NewPtr = IRB.CreateConstGEP1_64(
       IRB.getInt8Ty(),
diff --git a/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp b/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp
index 271c8d45fd4a21..c3f751c1a98830 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp
@@ -4802,7 +4802,7 @@ bool AMDGPULegalizerInfo::legalizeFastUnsafeFDIV(MachineInstr &MI,
   bool AllowInaccurateRcp = MI.getFlag(MachineInstr::FmAfn) ||
                             MF.getTarget().Options.UnsafeFPMath;
 
-  if (auto CLHS = getConstantFPVRegVal(LHS, MRI)) {
+  if (const auto *CLHS = getConstantFPVRegVal(LHS, MRI)) {
     if (!AllowInaccurateRcp && ResTy != LLT::scalar(16))
       return false;
 
diff --git a/llvm/lib/Target/AMDGPU/AMDGPULowerKernelAttributes.cpp b/llvm/lib/Target/AMDGPU/AMDGPULowerKernelAttributes.cpp
index e724c978c44b61..7d66d07c9d0fb7 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPULowerKernelAttributes.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPULowerKernelAttributes.cpp
@@ -87,7 +87,7 @@ Function *getBasePtrIntrinsic(Module &M, bool IsV5OrAbove) {
 static bool processUse(CallInst *CI, bool IsV5OrAbove) {
   Function *F = CI->getParent()->getParent();
 
-  auto MD = F->getMetadata("reqd_work_group_size");
+  auto *MD = F->getMetadata("reqd_work_group_size");
   const bool HasReqdWorkGroupSize = MD && MD->getNumOperands() == 3;
 
   const bool HasUniformWorkGroupSize =
diff --git a/llvm/lib/Target/AMDGPU/AMDGPULowerModuleLDSPass.cpp b/llvm/lib/Target/AMDGPU/AMDGPULowerModuleLDSPass.cpp
index a166087a5cdc78..51a5b7702c0093 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPULowerModuleLDSPass.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPULowerModuleLDSPass.cpp
@@ -324,7 +324,7 @@ class AMDGPULowerModuleLDS {
     for (GlobalVariable *GV : Variables) {
       auto ConstantGepIt = LDSVarsToConstantGEP.find(GV);
       if (ConstantGepIt != LDSVarsToConstantGEP.end()) {
-        auto elt = ConstantExpr::getPtrToInt(ConstantGepIt->second, I32);
+        auto *elt = ConstantExpr::getPtrToInt(ConstantGepIt->second, I32);
         Elements.push_back(elt);
       } else {
         Elements.push_back(PoisonValue::get(I32));
@@ -850,7 +850,7 @@ class AMDGPULowerModuleLDS {
     }
 
     assert(func->hasName()); // Checked by caller
-    auto emptyCharArray = ArrayType::get(Type::getInt8Ty(Ctx), 0);
+    auto *emptyCharArray = ArrayType::get(Type::getInt8Ty(Ctx), 0);
     GlobalVariable *N = new GlobalVariable(
         M, emptyCharArray, false, GlobalValue::ExternalLinkage, nullptr,
         Twine("llvm.amdgcn." + func->getName() + ".dynlds"), nullptr, GlobalValue::NotThreadLocal, AMDGPUAS::LOCAL_ADDRESS,
@@ -890,8 +890,8 @@ class AMDGPULowerModuleLDS {
 
           markUsedByKernel(func, N);
 
-          auto emptyCharArray = ArrayType::get(Type::getInt8Ty(Ctx), 0);
-          auto GEP = ConstantExpr::getGetElementPtr(
+          auto *emptyCharArray = ArrayType::get(Type::getInt8Ty(Ctx), 0);
+          auto *GEP = ConstantExpr::getGetElementPtr(
               emptyCharArray, N, ConstantInt::get(I32, 0), true);
           newDynamicLDS.push_back(ConstantExpr::getPtrToInt(GEP, I32));
         } else {
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUMCInstLower.h b/llvm/lib/Target/AMDGPU/AMDGPUMCInstLower.h
index 5c656f158e7146..7176cc5d3439bf 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUMCInstLower.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPUMCInstLower.h
@@ -55,7 +55,7 @@ static inline const MCExpr *lowerAddrSpaceCast(const TargetMachine &TM,
   // Clang generates addrspacecast for null pointers in private and local
   // address space, which needs to be lowered.
   if (CE && CE->getOpcode() == Instruction::AddrSpaceCast) {
-    auto Op = CE->getOperand(0);
+    auto *Op = CE->getOperand(0);
     auto SrcAddr = Op->getType()->getPointerAddressSpace();
     if (Op->isNullValue() && AT.getNullPointerValue(SrcAddr) == 0) {
       auto DstAddr = CE->getType()->getPointerAddressSpace();
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp b/llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp
index 99b4fca20bb2da..1d83d0c4c93372 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp
@@ -31,7 +31,7 @@ getKernelDynLDSGlobalFromFunction(const Function &F) {
 static bool hasLDSKernelArgument(const Function &F) {
   for (const Argument &Arg : F.args()) {
     Type *ArgTy = Arg.getType();
-    if (auto PtrTy = dyn_cast<PointerType>(ArgTy)) {
+    if (auto *PtrTy = dyn_cast<PointerType>(ArgTy)) {
       if (PtrTy->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS)
         return true;
     }
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUPerfHintAnalysis.cpp b/llvm/lib/Target/AMDGPU/AMDGPUPerfHintAnalysis.cpp
index 040e931b82af2f..3a3751892c8b67 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUPerfHintAnalysis.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUPerfHintAnalysis.cpp
@@ -118,15 +118,15 @@ struct AMDGPUPerfHint {
 
 static std::pair<const Value *, const Type *> getMemoryInstrPtrAndType(
     const Instruction *Inst) {
-  if (auto LI = dyn_cast<LoadInst>(Inst))
+  if (const auto *LI = dyn_cast<LoadInst>(Inst))
     return {LI->getPointerOperand(), LI->getType()};
-  if (auto SI = dyn_cast<StoreInst>(Inst))
+  if (const auto *SI = dyn_cast<StoreInst>(Inst))
     return {SI->getPointerOperand(), SI->getValueOperand()->getType()};
-  if (auto AI = dyn_cast<AtomicCmpXchgInst>(Inst))
+  if (const auto *AI = dyn_cast<AtomicCmpXchgInst>(Inst))
     return {AI->getPointerOperand(), AI->getCompareOperand()->getType()};
-  if (auto AI = dyn_cast<AtomicRMWInst>(Inst))
+  if (const auto *AI = dyn_cast<AtomicRMWInst>(Inst))
     return {AI->getPointerOperand(), AI->getValOperand()->getType()};
-  if (auto MI = dyn_cast<AnyMemIntrinsic>(Inst))
+  if (const auto *MI = dyn_cast<AnyMemIntrinsic>(Inst))
     return {MI->getRawDest(), Type::getInt8Ty(MI->getContext())};
 
   return {nullptr, nullptr};
@@ -148,8 +148,8 @@ bool AMDGPUPerfHint::isIndirectAccess(const Instruction *Inst) const {
       continue;
     LLVM_DEBUG(dbgs() << "  check: " << *V << '\n');
 
-    if (auto LD = dyn_cast<LoadInst>(V)) {
-      auto M = LD->getPointerOperand();
+    if (const auto *LD = dyn_cast<LoadInst>(V)) {
+      const auto *M = LD->getPointerOperand();
       if (isGlobalAddr(M)) {
         LLVM_DEBUG(dbgs() << "    is IA\n");
         return true;
@@ -157,32 +157,32 @@ bool AMDGPUPerfHint::isIndirectAccess(const Instruction *Inst) const {
       continue;
     }
 
-    if (auto GEP = dyn_cast<GetElementPtrInst>(V)) {
-      auto P = GEP->getPointerOperand();
+    if (const auto *GEP = dyn_cast<GetElementPtrInst>(V)) {
+      const auto *P = GEP->getPointerOperand();
       WorkSet.insert(P);
       for (unsigned I = 1, E = GEP->getNumIndices() + 1; I != E; ++I)
         WorkSet.insert(GEP->getOperand(I));
       continue;
     }
 
-    if (auto U = dyn_cast<UnaryInstruction>(V)) {
+    if (const auto *U = dyn_cast<UnaryInstruction>(V)) {
       WorkSet.insert(U->getOperand(0));
       continue;
     }
 
-    if (auto BO = dyn_cast<BinaryOperator>(V)) {
+    if (const auto *BO = dyn_cast<BinaryOperator>(V)) {
       WorkSet.insert(BO->getOperand(0));
       WorkSet.insert(BO->getOperand(1));
       continue;
     }
 
-    if (auto S = dyn_cast<SelectInst>(V)) {
+    if (const auto *S = dyn_cast<SelectInst>(V)) {
       WorkSet.insert(S->getFalseValue());
       WorkSet.insert(S->getTrueValue());
       continue;
     }
 
-    if (auto E = dyn_cast<ExtractElementInst>(V)) {
+    if (const auto *E = dyn_cast<ExtractElementInst>(V)) {
       WorkSet.insert(E->getVectorOperand());
       continue;
     }
@@ -331,7 +331,7 @@ bool AMDGPUPerfHint::needLimitWave(const AMDGPUPerfHintAnalysis::FuncInfo &FI) {
 }
 
 bool AMDGPUPerfHint::isGlobalAddr(const Value *V) const {
-  if (auto PT = dyn_cast<PointerType>(V->getType())) {
+  if (auto *PT = dyn_cast<PointerType>(V->getType())) {
     unsigned As = PT->getAddressSpace();
     // Flat likely points to global too.
     return As == AMDGPUAS::GLOBAL_ADDRESS || As == AMDGPUAS::FLAT_ADDRESS;
@@ -340,7 +340,7 @@ bool AMDGPUPerfHint::isGlobalAddr(const Value *V) const {
 }
 
 bool AMDGPUPerfHint::isLocalAddr(const Value *V) const {
-  if (auto PT = dyn_cast<PointerType>(V->getType()))
+  if (auto *PT = dyn_cast<PointerType>(V->getType()))
     return PT->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS;
   return false;
 }
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUPrintfRuntimeBinding.cpp b/llvm/lib/Target/AMDGPU/AMDGPUPrintfRuntimeBinding.cpp
index 02b0d436451a3f..a899805dc46b1a 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUPrintfRuntimeBinding.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUPrintfRuntimeBinding.cpp
@@ -436,7 +436,7 @@ bool AMDGPUPrintfRuntimeBindingImpl::run(Module &M) {
   if (TT.getArch() == Triple::r600)
     return false;
 
-  auto PrintfFunction = M.getFunction("printf");
+  auto *PrintfFunction = M.getFunction("printf");
   if (!PrintfFunction || !PrintfFunction->isDeclaration() ||
       M.getModuleFlag("openmp"))
     return false;
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUSubtarget.cpp b/llvm/lib/Target/AMDGPU/AMDGPUSubtarget.cpp
index 67d8715d3f1c26..9809a289df093b 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUSubtarget.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUSubtarget.cpp
@@ -190,7 +190,7 @@ std::pair<unsigned, unsigned> AMDGPUSubtarget::getWavesPerEU(
 }
 
 static unsigned getReqdWorkGroupSize(const Function &Kernel, unsigned Dim) {
-  auto Node = Kernel.getMetadata("reqd_work_group_size");
+  auto *Node = Kernel.getMetadata("reqd_work_group_size");
   if (Node && Node->getNumOperands() == 3)
     return mdconst::extract<ConstantInt>(Node->getOperand(Dim))->getZExtValue();
   return std::numeric_limits<unsigned>::max();
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUSwLowerLDS.cpp b/llvm/lib/Target/AMDGPU/AMDGPUSwLowerLDS.cpp
index e75b70b00c3ef7..3736e9dbd32198 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUSwLowerLDS.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUSwLowerLDS.cpp
@@ -368,7 +368,7 @@ void AMDGPUSwLowerLDS::buildSwDynLDSGlobal(Function *Func) {
       LDSParams.IndirectAccess.DynamicLDSGlobals.empty())
     return;
   // Create new global pointer variable
-  auto emptyCharArray = ArrayType::get(IRB.getInt8Ty(), 0);
+  auto *emptyCharArray = ArrayType::get(IRB.getInt8Ty(), 0);
   LDSParams.SwDynLDS = new GlobalVariable(
       M, emptyCharArray, false, GlobalValue::ExternalLinkage, nullptr,
       "llvm.amdgcn." + Func->getName() + ".dynlds", nullptr,
@@ -1070,7 +1070,8 @@ void AMDGPUSwLowerLDS::lowerNonKernelLDSAccesses(
       IRB.CreateLoad(IRB.getPtrTy(AMDGPUAS::GLOBAL_ADDRESS), BaseLoad);
 
   for (GlobalVariable *GV : LDSGlobals) {
-    auto GVIt = std::find(OrdereLDSGlobals.begin(), OrdereLDSGlobals.end(), GV);
+    const auto *GVIt =
+        std::find(OrdereLDSGlobals.begin(), OrdereLDSGlobals.end(), GV);
     assert(GVIt != OrdereLDSGlobals.end());
     uint32_t GVOffset = std::distance(OrdereLDSGlobals.begin(), GVIt);
 
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
index 1f2148c2922de9..44ef9e206a8aa0 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
@@ -556,8 +556,8 @@ createGCNMaxILPMachineScheduler(MachineSchedContext *C) {
 static ScheduleDAGInstrs *
 createIterativeGCNMaxOccupancyMachineScheduler(MachineSchedContext *C) {
   const GCNSubtarget &ST = C->MF->getSubtarget<GCNSubtarget>();
-  auto DAG = new GCNIterativeScheduler(C,
-    GCNIterativeScheduler::SCHEDULE_LEGACYMAXOCCUPANCY);
+  auto *DAG = new GCNIterativeScheduler(
+      C, GCNIterativeScheduler::SCHEDULE_LEGACYMAXOCCUPANCY);
   DAG->addMutation(createLoadClusterDAGMutation(DAG->TII, DAG->TRI));
   if (ST.shouldClusterStores())
     DAG->addMutation(createStoreClusterDAGMutation(DAG->TII, DAG->TRI));
@@ -572,8 +572,7 @@ static ScheduleDAGInstrs *createMinRegScheduler(MachineSchedContext *C) {
 static ScheduleDAGInstrs *
 createIterativeILPMachineScheduler(MachineSchedContext *C) {
   const GCNSubtarget &ST = C->MF->getSubtarget<GCNSubtarget>();
-  auto DAG = new GCNIterativeScheduler(C,
-    GCNIterativeScheduler::SCHEDULE_ILP);
+  auto *DAG = new GCNIterativeScheduler(C, GCNIterativeScheduler::SCHEDULE_ILP);
   DAG->addMutation(createLoadClusterDAGMutation(DAG->TII, DAG->TRI));
   if (ST.shouldClusterStores())
     DAG->addMutation(createStoreClusterDAGMutation(DAG->TII, DAG->TRI));
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUTargetTransformInfo.cpp b/llvm/lib/Target/AMDGPU/AMDGPUTargetTransformInfo.cpp
index 4cf7733a260ff0..d348166c2d9a04 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUTargetTransformInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUTargetTransformInfo.cpp
@@ -784,7 +784,7 @@ InstructionCost GCNTTIImpl::getCFInstrCost(unsigned Opcode,
   switch (Opcode) {
   case Instruction::Br: {
     // Branch instruction takes about 4 slots on gfx900.
-    auto BI = dyn_cast_or_null<BranchInst>(I);
+    const auto *BI = dyn_cast_or_null<BranchInst>(I);
     if (BI && BI->isUnconditional())
       return SCost ? 1 : 4;
     // Suppose conditional branch takes additional 3 exec manipulations
@@ -792,7 +792,7 @@ InstructionCost GCNTTIImpl::getCFInstrCost(unsigned Opcode,
     return CBrCost;
   }
   case Instruction::Switch: {
-    auto SI = dyn_cast_or_null<SwitchInst>(I);
+    const auto *SI = dyn_cast_or_null<SwitchInst>(I);
     // Each case (including default) takes 1 cmp + 1 cbr instructions in
     // average.
     return (SI ? (SI->getNumCases() + 1) : 4) * (CBrCost + 1);
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUUnifyMetadata.cpp b/llvm/lib/Target/AMDGPU/AMDGPUUnifyMetadata.cpp
index 327666c2a0f28b..cd7866b86d55b0 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUUnifyMetadata.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUUnifyMetadata.cpp
@@ -53,16 +53,16 @@ namespace {
     /// Keep the largest version as the sole operand if PickFirst is false.
     /// Otherwise pick it from the first value, representing kernel module.
     bool unifyVersionMD(Module &M, StringRef Name, bool PickFirst) {
-      auto NamedMD = M.getNamedMetadata(Name);
+      auto *NamedMD = M.getNamedMetadata(Name);
       if (!NamedMD || NamedMD->getNumOperands() <= 1)
         return false;
       MDNode *MaxMD = nullptr;
       auto MaxVer = 0U;
       for (auto *VersionMD : NamedMD->operands()) {
         assert(VersionMD->getNumOperands() == 2);
-        auto CMajor = mdconst::extract<ConstantInt>(VersionMD->getOperand(0));
+        auto *CMajor = mdconst::extract<ConstantInt>(VersionMD->getOperand(0));
         auto VersionMajor = CMajor->getZExtValue();
-        auto CMinor = mdconst::extract<ConstantInt>(VersionMD->getOperand(1));
+        auto *CMinor = mdconst::extract<ConstantInt>(VersionMD->getOperand(1));
         auto VersionMinor = CMinor->getZExtValue();
         auto Ver = (VersionMajor * 100) + (VersionMinor * 10);
         if (Ver > MaxVer) {
@@ -86,7 +86,7 @@ namespace {
   /// !n2 = !{!"cl_khr_image"}
   /// Combine it into a single list with unique operands.
   bool unifyExtensionMD(Module &M, StringRef Name) {
-    auto NamedMD = M.getNamedMetadata(Name);
+    auto *NamedMD = M.getNamedMetadata(Name);
     if (!NamedMD || NamedMD->getNumOperands() == 1)
       return false;
 
diff --git a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
index e12db4ab058ed6..758f864fd20e6a 100644
--- a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
+++ b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
@@ -6118,7 +6118,7 @@ bool AMDGPUAsmParser::ParseDirectivePALMetadataBegin() {
                           AMDGPU::PALMD::AssemblerDirectiveEnd, String))
     return true;
 
-  auto PALMetadata = getTargetStreamer().getPALMetadata();
+  auto *PALMetadata = getTargetStreamer().getPALMetadata();
   if (!PALMetadata->setFromString(String))
     return Error(getLoc(), "invalid PAL metadata");
   return false;
@@ -6132,7 +6132,7 @@ bool AMDGPUAsmParser::ParseDirectivePALMetadata() {
                  "not available on non-amdpal OSes")).str());
   }
 
-  auto PALMetadata = getTargetStreamer().getPALMetadata();
+  auto *PALMetadata = getTargetStreamer().getPALMetadata();
   PALMetadata->setLegacy();
   for (;;) {
     uint32_t Key, Value;
@@ -7959,7 +7959,7 @@ AMDGPUAsmParser::parseStructuredOpFields(ArrayRef<StructuredOpField *> Fields) {
         !skipToken(AsmToken::Colon, "colon expected"))
       return ParseStatus::Failure;
 
-    auto I =
+    const auto *I =
         find_if(Fields, [Id](StructuredOpField *F) { return F->Id == Id; });
     if (I == Fields.end())
       return Error(IdLoc, "unknown field");
@@ -8798,7 +8798,7 @@ void AMDGPUAsmParser::cvtVOP3(MCInst &Inst, const OperandVector &Operands,
   // we don't allow modifiers for this operand in assembler so src2_modifiers
   // should be 0.
   if (isMAC(Opc)) {
-    auto it = Inst.begin();
+    auto *it = Inst.begin();
     std::advance(it, AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::src2_modifiers));
     it = Inst.insert(it, MCOperand::createImm(0)); // no modifiers for src2
     ++it;
@@ -9627,7 +9627,7 @@ void AMDGPUAsmParser::cvtSDWA(MCInst &Inst, const OperandVector &Operands,
   // it has src2 register operand that is tied to dst operand
   if (Inst.getOpcode() == AMDGPU::V_MAC_F32_sdwa_vi ||
       Inst.getOpcode() == AMDGPU::V_MAC_F16_sdwa_vi)  {
-    auto it = Inst.begin();
+    auto *it = Inst.begin();
     std::advance(
       it, AMDGPU::getNamedOperandIdx(Inst.getOpcode(), AMDGPU::OpName::src2));
     Inst.insert(it, Inst.getOperand(0)); // src2 = dst
diff --git a/llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp b/llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp
index 9eedcc636fd94e..fdef9865b82c06 100644
--- a/llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp
+++ b/llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp
@@ -94,7 +94,7 @@ static int insertNamedMCOperand(MCInst &MI, const MCOperand &Op,
                                 uint16_t NameIdx) {
   int OpIdx = AMDGPU::getNamedOperandIdx(MI.getOpcode(), NameIdx);
   if (OpIdx != -1) {
-    auto I = MI.begin();
+    auto *I = MI.begin();
     std::advance(I, OpIdx);
     MI.insert(I, Op);
   }
@@ -104,7 +104,7 @@ static int insertNamedMCOperand(MCInst &MI, const MCOperand &Op,
 static DecodeStatus decodeSOPPBrTarget(MCInst &Inst, unsigned Imm,
                                        uint64_t Addr,
                                        const MCDisassembler *Decoder) {
-  auto DAsm = static_cast<const AMDGPUDisassembler*>(Decoder);
+  const auto *DAsm = static_cast<const AMDGPUDisassembler *>(Decoder);
 
   // Our branches take a simm16.
   int64_t Offset = SignExtend64<16>(Imm) * 4 + 4 + Addr;
@@ -116,7 +116,7 @@ static DecodeStatus decodeSOPPBrTarget(MCInst &Inst, unsigned Imm,
 
 static DecodeStatus decodeSMEMOffset(MCInst &Inst, unsigned Imm, uint64_t Addr,
                                      const MCDisassembler *Decoder) {
-  auto DAsm = static_cast<const AMDGPUDisassembler*>(Decoder);
+  const auto *DAsm = static_cast<const AMDGPUDisassembler *>(Decoder);
   int64_t Offset;
   if (DAsm->isGFX12Plus()) { // GFX12 supports 24-bit signed offsets.
     Offset = SignExtend64<24>(Imm);
@@ -130,20 +130,20 @@ static DecodeStatus decodeSMEMOffset(MCInst &Inst, unsigned Imm, uint64_t Addr,
 
 static DecodeStatus decodeBoolReg(MCInst &Inst, unsigned Val, uint64_t Addr,
                                   const MCDisassembler *Decoder) {
-  auto DAsm = static_cast<const AMDGPUDisassembler*>(Decoder);
+  const auto *DAsm = static_cast<const AMDGPUDisassembler *>(Decoder);
   return addOperand(Inst, DAsm->decodeBoolReg(Val));
 }
 
 static DecodeStatus decodeSplitBarrier(MCInst &Inst, unsigned Val,
                                        uint64_t Addr,
                                        const MCDisassembler *Decoder) {
-  auto DAsm = static_cast<const AMDGPUDisassembler *>(Decoder);
+  const auto *DAsm = static_cast<const AMDGPUDisassembler *>(Decoder);
   return addOperand(Inst, DAsm->decodeSplitBarrier(Val));
 }
 
 static DecodeStatus decodeDpp8FI(MCInst &Inst, unsigned Val, uint64_t Addr,
                                  const MCDisassembler *Decoder) {
-  auto DAsm = static_cast<const AMDGPUDisassembler *>(Decoder);
+  const auto *DAsm = static_cast<const AMDGPUDisassembler *>(Decoder);
   return addOperand(Inst, DAsm->decodeDpp8FI(Val));
 }
 
@@ -185,7 +185,7 @@ static DecodeStatus decodeSrcOp(MCInst &Inst, unsigned EncSize,
                                 AMDGPU::OperandSemantics Sema,
                                 const MCDisassembler *Decoder) {
   assert(Imm < (1U << EncSize) && "Operand doesn't fit encoding!");
-  auto DAsm = static_cast<const AMDGPUDisassembler *>(Decoder);
+  const auto *DAsm = static_cast<const AMDGPUDisassembler *>(Decoder);
   return addOperand(Inst, DAsm->decodeSrcOp(OpWidth, EncImm, MandatoryLiteral,
                                             ImmWidth, Sema));
 }
@@ -312,7 +312,7 @@ static DecodeStatus DecodeVGPR_16RegisterClass(MCInst &Inst, unsigned Imm,
 
   bool IsHi = Imm & (1 << 9);
   unsigned RegIdx = Imm & 0xff;
-  auto DAsm = static_cast<const AMDGPUDisassembler *>(Decoder);
+  const auto *DAsm = static_cast<const AMDGPUDisassembler *>(Decoder);
   return addOperand(Inst, DAsm->createVGPR16Operand(RegIdx, IsHi));
 }
 
@@ -323,7 +323,7 @@ DecodeVGPR_16_Lo128RegisterClass(MCInst &Inst, unsigned Imm, uint64_t /*Addr*/,
 
   bool IsHi = Imm & (1 << 7);
   unsigned RegIdx = Imm & 0x7f;
-  auto DAsm = static_cast<const AMDGPUDisassembler *>(Decoder);
+  const auto *DAsm = static_cast<const AMDGPUDisassembler *>(Decoder);
   return addOperand(Inst, DAsm->createVGPR16Operand(RegIdx, IsHi));
 }
 
@@ -393,7 +393,7 @@ static bool IsAGPROperand(const MCInst &Inst, int OpIdx,
 static DecodeStatus decodeAVLdSt(MCInst &Inst, unsigned Imm,
                                  AMDGPUDisassembler::OpWidthTy Opw,
                                  const MCDisassembler *Decoder) {
-  auto DAsm = static_cast<const AMDGPUDisassembler*>(Decoder);
+  const auto *DAsm = static_cast<const AMDGPUDisassembler *>(Decoder);
   if (!DAsm->isGFX90A()) {
     Imm &= 511;
   } else {
@@ -435,7 +435,7 @@ static DecodeStatus decodeOperand_VSrc_f64(MCInst &Inst, unsigned Imm,
                                            uint64_t Addr,
                                            const MCDisassembler *Decoder) {
   assert(Imm < (1 << 9) && "9-bit encoding");
-  auto DAsm = static_cast<const AMDGPUDisassembler *>(Decoder);
+  const auto *DAsm = static_cast<const AMDGPUDisassembler *>(Decoder);
   return addOperand(Inst,
                     DAsm->decodeSrcOp(AMDGPUDisassembler::OPW64, Imm, false, 64,
                                       AMDGPU::OperandSemantics::FP64));
@@ -451,7 +451,7 @@ DECODE_SDWA(VopcDst)
 static DecodeStatus decodeVersionImm(MCInst &Inst, unsigned Imm,
                                      uint64_t /* Addr */,
                                      const MCDisassembler *Decoder) {
-  auto DAsm = static_cast<const AMDGPUDisassembler *>(Decoder);
+  const auto *DAsm = static_cast<const AMDGPUDisassembler *>(Decoder);
   return addOperand(Inst, DAsm->decodeVersionImm(Imm));
 }
 
@@ -675,7 +675,7 @@ DecodeStatus AMDGPUDisassembler::getInstruction(MCInst &MI, uint64_t &Size,
     int TFEOpIdx =
         AMDGPU::getNamedOperandIdx(MI.getOpcode(), AMDGPU::OpName::tfe);
     if (TFEOpIdx != -1) {
-      auto TFEIter = MI.begin();
+      auto *TFEIter = MI.begin();
       std::advance(TFEIter, TFEOpIdx);
       MI.insert(TFEIter, MCOperand::createImm(0));
     }
@@ -686,7 +686,7 @@ DecodeStatus AMDGPUDisassembler::getInstruction(MCInst &MI, uint64_t &Size,
     int SWZOpIdx =
         AMDGPU::getNamedOperandIdx(MI.getOpcode(), AMDGPU::OpName::swz);
     if (SWZOpIdx != -1) {
-      auto SWZIter = MI.begin();
+      auto *SWZIter = MI.begin();
       std::advance(SWZIter, SWZOpIdx);
       MI.insert(SWZIter, MCOperand::createImm(0));
     }
@@ -1774,10 +1774,10 @@ MCOperand AMDGPUDisassembler::decodeVersionImm(unsigned Imm) const {
     return MCOperand::createImm(Imm);
 
   const auto &Versions = AMDGPU::UCVersion::getGFXVersions();
-  auto I = find_if(Versions,
-                   [Version = Version](const AMDGPU::UCVersion::GFXVersion &V) {
-                     return V.Code == Version;
-                   });
+  const auto *I = find_if(
+      Versions, [Version = Version](const AMDGPU::UCVersion::GFXVersion &V) {
+        return V.Code == Version;
+      });
   MCContext &Ctx = getContext();
   const MCExpr *E;
   if (I == Versions.end())
diff --git a/llvm/lib/Target/AMDGPU/GCNDPPCombine.cpp b/llvm/lib/Target/AMDGPU/GCNDPPCombine.cpp
index 833cd6dfa4ba2e..cc802b5fbb67c1 100644
--- a/llvm/lib/Target/AMDGPU/GCNDPPCombine.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNDPPCombine.cpp
@@ -501,7 +501,7 @@ MachineInstr *GCNDPPCombine::createDPPInst(
       return nullptr;
     }
     CombOldVGPR = getRegSubRegPair(*Src1);
-    auto MovDst = TII->getNamedOperand(MovMI, AMDGPU::OpName::vdst);
+    auto *MovDst = TII->getNamedOperand(MovMI, AMDGPU::OpName::vdst);
     const TargetRegisterClass *RC = MRI->getRegClass(MovDst->getReg());
     if (!isOfRegClass(CombOldVGPR, *RC, *MRI)) {
       LLVM_DEBUG(dbgs() << "  failed: src1 has wrong register class\n");
diff --git a/llvm/lib/Target/AMDGPU/GCNILPSched.cpp b/llvm/lib/Target/AMDGPU/GCNILPSched.cpp
index 8f15cc1b2b5378..79b4a68c6a44ec 100644
--- a/llvm/lib/Target/AMDGPU/GCNILPSched.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNILPSched.cpp
@@ -240,7 +240,7 @@ GCNILPScheduler::Candidate* GCNILPScheduler::pickCandidate() {
     return nullptr;
   auto Best = AvailQueue.begin();
   for (auto I = std::next(AvailQueue.begin()), E = AvailQueue.end(); I != E; ++I) {
-    auto NewBestSU = pickBest(Best->SU, I->SU);
+    const auto *NewBestSU = pickBest(Best->SU, I->SU);
     if (NewBestSU != Best->SU) {
       assert(NewBestSU == I->SU);
       Best = I;
@@ -272,7 +272,7 @@ void GCNILPScheduler::advanceToCycle(unsigned NextCycle) {
 
 void GCNILPScheduler::releasePredecessors(const SUnit* SU) {
   for (const auto &PredEdge : SU->Preds) {
-    auto PredSU = PredEdge.getSUnit();
+    auto *PredSU = PredEdge.getSUnit();
     if (PredEdge.isWeak())
       continue;
     assert(PredSU->isBoundaryNode() || PredSU->NumSuccsLeft > 0);
@@ -311,7 +311,7 @@ GCNILPScheduler::schedule(ArrayRef<const SUnit*> BotRoots,
   Schedule.reserve(SUnits.size());
   while (true) {
     if (AvailQueue.empty() && !PendingQueue.empty()) {
-      auto EarliestSU =
+      auto *EarliestSU =
           llvm::min_element(PendingQueue, [=](const Candidate &C1,
                                               const Candidate &C2) {
             return C1.SU->getHeight() < C2.SU->getHeight();
@@ -328,10 +328,10 @@ GCNILPScheduler::schedule(ArrayRef<const SUnit*> BotRoots,
                << ' ' << C.SU->NodeNum;
                dbgs() << '\n';);
 
-    auto C = pickCandidate();
+    auto *C = pickCandidate();
     assert(C);
     AvailQueue.remove(*C);
-    auto SU = C->SU;
+    auto *SU = C->SU;
     LLVM_DEBUG(dbgs() << "Selected "; DAG.dumpNode(*SU));
 
     advanceToCycle(SU->getHeight());
diff --git a/llvm/lib/Target/AMDGPU/GCNIterativeScheduler.cpp b/llvm/lib/Target/AMDGPU/GCNIterativeScheduler.cpp
index 061b0515031b1b..13504508e2fb2e 100644
--- a/llvm/lib/Target/AMDGPU/GCNIterativeScheduler.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNIterativeScheduler.cpp
@@ -47,7 +47,7 @@ static void printRegion(raw_ostream &OS,
                         const LiveIntervals *LIS,
                         unsigned MaxInstNum =
                           std::numeric_limits<unsigned>::max()) {
-  auto BB = Begin->getParent();
+  auto *BB = Begin->getParent();
   OS << BB->getParent()->getName() << ":" << printMBBReference(*BB) << ' '
      << BB->getName() << ":\n";
   auto I = Begin;
@@ -76,7 +76,7 @@ static void printLivenessInfo(raw_ostream &OS,
                               MachineBasicBlock::iterator Begin,
                               MachineBasicBlock::iterator End,
                               const LiveIntervals *LIS) {
-  const auto BB = Begin->getParent();
+  auto *const BB = Begin->getParent();
   const auto &MRI = BB->getParent()->getRegInfo();
 
   const auto LiveIns = getLiveRegsBefore(*Begin, *LIS);
@@ -90,7 +90,7 @@ static void printLivenessInfo(raw_ostream &OS,
 LLVM_DUMP_METHOD
 void GCNIterativeScheduler::printRegions(raw_ostream &OS) const {
   const auto &ST = MF.getSubtarget<GCNSubtarget>();
-  for (const auto R : Regions) {
+  for (auto *const R : Regions) {
     OS << "Region to schedule ";
     printRegion(OS, R->Begin, R->End, LIS, 1);
     printLivenessInfo(OS, R->Begin, R->End, LIS);
@@ -127,7 +127,7 @@ class GCNIterativeScheduler::BuildDAG {
 public:
   BuildDAG(const Region &R, GCNIterativeScheduler &_Sch)
     : Sch(_Sch) {
-    auto BB = R.Begin->getParent();
+    auto *BB = R.Begin->getParent();
     Sch.BaseClass::startBlock(BB);
     Sch.BaseClass::enterRegion(BB, R.Begin, R.End, R.NumRegionInstrs);
 
@@ -165,7 +165,7 @@ class GCNIterativeScheduler::OverrideLegacyStrategy {
     , SaveSchedImpl(std::move(_Sch.SchedImpl))
     , SaveMaxRP(R.MaxPressure) {
     Sch.SchedImpl.reset(&OverrideStrategy);
-    auto BB = R.Begin->getParent();
+    auto *BB = R.Begin->getParent();
     Sch.BaseClass::startBlock(BB);
     Sch.BaseClass::enterRegion(BB, R.Begin, R.End, R.NumRegionInstrs);
   }
@@ -355,7 +355,7 @@ void GCNIterativeScheduler::scheduleRegion(Region &R, Range &&Schedule,
 #ifndef NDEBUG
   const auto SchedMaxRP = getSchedulePressure(R, Schedule);
 #endif
-  auto BB = R.Begin->getParent();
+  auto *BB = R.Begin->getParent();
   auto Top = R.Begin;
   for (const auto &I : Schedule) {
     auto MI = getMachineInstr(I);
diff --git a/llvm/lib/Target/AMDGPU/GCNMinRegStrategy.cpp b/llvm/lib/Target/AMDGPU/GCNMinRegStrategy.cpp
index 04a6f2a7f4fd5b..4154f946de8ff9 100644
--- a/llvm/lib/Target/AMDGPU/GCNMinRegStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNMinRegStrategy.cpp
@@ -88,7 +88,7 @@ int GCNMinRegScheduler::getReadySuccessors(const SUnit *SU) const {
   for (auto SDep : SU->Succs) {
     bool wouldBeScheduled = true;
     for (auto PDep : SDep.getSUnit()->Preds) {
-      auto PSU = PDep.getSUnit();
+      auto *PSU = PDep.getSUnit();
       assert(!PSU->isBoundaryNode());
       if (PSU != SU && !isScheduled(PSU)) {
         wouldBeScheduled = false;
@@ -143,7 +143,7 @@ GCNMinRegScheduler::Candidate* GCNMinRegScheduler::pickCandidate() {
     LLVM_DEBUG(dbgs() << "\nSelecting min non-ready producing candidate among "
                       << Num << '\n');
     Num = findMax(Num, [=](const Candidate &C) {
-      auto SU = C.SU;
+      const auto *SU = C.SU;
       int Res = getNotReadySuccessors(SU);
       LLVM_DEBUG(dbgs() << "SU(" << SU->NodeNum << ") would left non-ready "
                         << Res << " successors, metric = " << -Res << '\n');
@@ -154,7 +154,7 @@ GCNMinRegScheduler::Candidate* GCNMinRegScheduler::pickCandidate() {
     LLVM_DEBUG(dbgs() << "\nSelecting most producing candidate among " << Num
                       << '\n');
     Num = findMax(Num, [=](const Candidate &C) {
-      auto SU = C.SU;
+      const auto *SU = C.SU;
       auto Res = getReadySuccessors(SU);
       LLVM_DEBUG(dbgs() << "SU(" << SU->NodeNum << ") would make ready " << Res
                         << " successors, metric = " << Res << '\n');
@@ -181,7 +181,7 @@ void GCNMinRegScheduler::bumpPredsPriority(const SUnit *SchedSU, int Priority) {
         S.getKind() != SDep::Data)
       continue;
     for (const auto &P : S.getSUnit()->Preds) {
-      auto PSU = P.getSUnit();
+      auto *PSU = P.getSUnit();
       assert(!PSU->isBoundaryNode());
       if (PSU != SchedSU && !isScheduled(PSU)) {
         Set.insert(PSU);
@@ -190,7 +190,7 @@ void GCNMinRegScheduler::bumpPredsPriority(const SUnit *SchedSU, int Priority) {
   }
   SmallVector<const SUnit*, 32> Worklist(Set.begin(), Set.end());
   while (!Worklist.empty()) {
-    auto SU = Worklist.pop_back_val();
+    const auto *SU = Worklist.pop_back_val();
     assert(!SU->isBoundaryNode());
     for (const auto &P : SU->Preds) {
       if (!P.getSUnit()->isBoundaryNode() && !isScheduled(P.getSUnit()) &&
@@ -212,7 +212,7 @@ void GCNMinRegScheduler::bumpPredsPriority(const SUnit *SchedSU, int Priority) {
 
 void GCNMinRegScheduler::releaseSuccessors(const SUnit* SU, int Priority) {
   for (const auto &S : SU->Succs) {
-    auto SuccSU = S.getSUnit();
+    auto *SuccSU = S.getSUnit();
     if (S.isWeak())
       continue;
     assert(SuccSU->isBoundaryNode() || getNumPreds(SuccSU) > 0);
@@ -246,10 +246,10 @@ GCNMinRegScheduler::schedule(ArrayRef<const SUnit*> TopRoots,
                << ' ' << C.SU->NodeNum << "(P" << C.Priority << ')';
                dbgs() << '\n';);
 
-    auto C = pickCandidate();
+    auto *C = pickCandidate();
     assert(C);
     RQ.remove(*C);
-    auto SU = C->SU;
+    const auto *SU = C->SU;
     LLVM_DEBUG(dbgs() << "Selected "; DAG.dumpNode(*SU));
 
     releaseSuccessors(SU, StepNo);
diff --git a/llvm/lib/Target/AMDGPU/GCNNSAReassign.cpp b/llvm/lib/Target/AMDGPU/GCNNSAReassign.cpp
index d6395fd75924de..18ffd8820f95e0 100644
--- a/llvm/lib/Target/AMDGPU/GCNNSAReassign.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNNSAReassign.cpp
@@ -328,12 +328,14 @@ bool GCNNSAReassign::runOnMachineFunction(MachineFunction &MF) {
         continue;
     } else {
       // Check we did not make it worse for other instructions.
-      auto I = std::lower_bound(Candidates.begin(), &C, MinInd,
-                                [this](const Candidate &C, SlotIndex I) {
-                                  return LIS->getInstructionIndex(*C.first) < I;
-                                });
-      for (auto E = Candidates.end(); Success && I != E &&
-              LIS->getInstructionIndex(*I->first) < MaxInd; ++I) {
+      auto *I =
+          std::lower_bound(Candidates.begin(), &C, MinInd,
+                           [this](const Candidate &C, SlotIndex I) {
+                             return LIS->getInstructionIndex(*C.first) < I;
+                           });
+      for (auto *E = Candidates.end();
+           Success && I != E && LIS->getInstructionIndex(*I->first) < MaxInd;
+           ++I) {
         if (I->second && CheckNSA(*I->first, true) < NSA_Status::CONTIGUOUS) {
           Success = false;
           LLVM_DEBUG(dbgs() << "\tNSA conversion conflict with " << *I->first);
diff --git a/llvm/lib/Target/AMDGPU/GCNRegPressure.cpp b/llvm/lib/Target/AMDGPU/GCNRegPressure.cpp
index c83af729f501fe..cb0624f11592d2 100644
--- a/llvm/lib/Target/AMDGPU/GCNRegPressure.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNRegPressure.cpp
@@ -38,8 +38,9 @@ bool llvm::isEqual(const GCNRPTracker::LiveRegSet &S1,
 unsigned GCNRegPressure::getRegKind(Register Reg,
                                     const MachineRegisterInfo &MRI) {
   assert(Reg.isVirtual());
-  const auto RC = MRI.getRegClass(Reg);
-  auto STI = static_cast<const SIRegisterInfo*>(MRI.getTargetRegisterInfo());
+  const auto *const RC = MRI.getRegClass(Reg);
+  const auto *STI =
+      static_cast<const SIRegisterInfo *>(MRI.getTargetRegisterInfo());
   return STI->isSGPRClass(RC)
              ? (STI->getRegSizeInBits(*RC) == 32 ? SGPR32 : SGPR_TUPLE)
          : STI->isAGPRClass(RC)
diff --git a/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUMCExpr.cpp b/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUMCExpr.cpp
index 4fbd7d0f889457..d1212ec76f9860 100644
--- a/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUMCExpr.cpp
+++ b/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUMCExpr.cpp
@@ -77,7 +77,7 @@ void AMDGPUMCExpr::printImpl(raw_ostream &OS, const MCAsmInfo *MAI) const {
     OS << "occupancy(";
     break;
   }
-  for (auto It = Args.begin(); It != Args.end(); ++It) {
+  for (const auto *It = Args.begin(); It != Args.end(); ++It) {
     (*It)->print(OS, MAI, /*InParens=*/false);
     if ((It + 1) != Args.end())
       OS << ", ";
diff --git a/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUMCKernelDescriptor.cpp b/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUMCKernelDescriptor.cpp
index 77e7e30ff5281b..14b3cdf37650c0 100644
--- a/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUMCKernelDescriptor.cpp
+++ b/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUMCKernelDescriptor.cpp
@@ -82,8 +82,8 @@ MCKernelDescriptor::getDefaultAmdhsaKernelDescriptor(const MCSubtargetInfo *STI,
 void MCKernelDescriptor::bits_set(const MCExpr *&Dst, const MCExpr *Value,
                                   uint32_t Shift, uint32_t Mask,
                                   MCContext &Ctx) {
-  auto Sft = MCConstantExpr::create(Shift, Ctx);
-  auto Msk = MCConstantExpr::create(Mask, Ctx);
+  const auto *Sft = MCConstantExpr::create(Shift, Ctx);
+  const auto *Msk = MCConstantExpr::create(Mask, Ctx);
   Dst = MCBinaryExpr::createAnd(Dst, MCUnaryExpr::createNot(Msk, Ctx), Ctx);
   Dst = MCBinaryExpr::createOr(Dst, MCBinaryExpr::createShl(Value, Sft, Ctx),
                                Ctx);
@@ -91,8 +91,8 @@ void MCKernelDescriptor::bits_set(const MCExpr *&Dst, const MCExpr *Value,
 
 const MCExpr *MCKernelDescriptor::bits_get(const MCExpr *Src, uint32_t Shift,
                                            uint32_t Mask, MCContext &Ctx) {
-  auto Sft = MCConstantExpr::create(Shift, Ctx);
-  auto Msk = MCConstantExpr::create(Mask, Ctx);
+  const auto *Sft = MCConstantExpr::create(Shift, Ctx);
+  const auto *Msk = MCConstantExpr::create(Mask, Ctx);
   return MCBinaryExpr::createLShr(MCBinaryExpr::createAnd(Src, Msk, Ctx), Sft,
                                   Ctx);
 }
diff --git a/llvm/lib/Target/AMDGPU/R600OpenCLImageTypeLoweringPass.cpp b/llvm/lib/Target/AMDGPU/R600OpenCLImageTypeLoweringPass.cpp
index 604a4cb1bf881b..02fa4cd581b934 100644
--- a/llvm/lib/Target/AMDGPU/R600OpenCLImageTypeLoweringPass.cpp
+++ b/llvm/lib/Target/AMDGPU/R600OpenCLImageTypeLoweringPass.cpp
@@ -83,7 +83,7 @@ GetFunctionFromMDNode(MDNode *Node) {
   if (NumOps != NumKernelArgMDNodes + 1)
     return nullptr;
 
-  auto F = mdconst::dyn_extract<Function>(Node->getOperand(0));
+  auto *F = mdconst::dyn_extract<Function>(Node->getOperand(0));
   if (!F)
     return nullptr;
 
@@ -153,7 +153,7 @@ class R600OpenCLImageTypeLoweringPass : public ModulePass {
     bool Modified = false;
 
     for (auto &Use : ImageArg.uses()) {
-      auto Inst = dyn_cast<CallInst>(Use.getUser());
+      auto *Inst = dyn_cast<CallInst>(Use.getUser());
       if (!Inst) {
         continue;
       }
@@ -186,7 +186,7 @@ class R600OpenCLImageTypeLoweringPass : public ModulePass {
     bool Modified = false;
 
     for (const auto &Use : SamplerArg.uses()) {
-      auto Inst = dyn_cast<CallInst>(Use.getUser());
+      auto *Inst = dyn_cast<CallInst>(Use.getUser());
       if (!Inst) {
         continue;
       }
@@ -218,7 +218,7 @@ class R600OpenCLImageTypeLoweringPass : public ModulePass {
 
     bool Modified = false;
     InstsToErase.clear();
-    for (auto ArgI = F->arg_begin(); ArgI != F->arg_end(); ++ArgI) {
+    for (auto *ArgI = F->arg_begin(); ArgI != F->arg_end(); ++ArgI) {
       Argument &Arg = *ArgI;
       StringRef Type = ArgTypeFromMD(KernelMDNode, Arg.getArgNo());
 
@@ -287,10 +287,10 @@ class R600OpenCLImageTypeLoweringPass : public ModulePass {
     }
 
     // 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());
+    auto *NewFT = FunctionType::get(FT->getReturnType(), ArgTypes, false);
+    auto *NewF = Function::Create(NewFT, F->getLinkage(), F->getName());
     ValueToValueMapTy VMap;
-    auto NewFArgIt = NewF->arg_begin();
+    auto *NewFArgIt = NewF->arg_begin();
     for (auto &Arg: F->args()) {
       auto ArgName = Arg.getName();
       NewFArgIt->setName(ArgName);
diff --git a/llvm/lib/Target/AMDGPU/SIFixSGPRCopies.cpp b/llvm/lib/Target/AMDGPU/SIFixSGPRCopies.cpp
index 46f5097c679fb3..654ae412f39c13 100644
--- a/llvm/lib/Target/AMDGPU/SIFixSGPRCopies.cpp
+++ b/llvm/lib/Target/AMDGPU/SIFixSGPRCopies.cpp
@@ -587,7 +587,7 @@ static bool hoistAndMergeSGPRInits(unsigned Reg,
   for (auto &Init : Inits) {
     auto &Defs = Init.second;
     for (auto *MI : Defs) {
-      auto MBB = MI->getParent();
+      auto *MBB = MI->getParent();
       MachineInstr &BoundaryMI = *getFirstNonPrologue(MBB, TII);
       MachineBasicBlock::reverse_iterator B(BoundaryMI);
       // Check if B should actually be a boundary. If not set the previous
@@ -769,7 +769,7 @@ bool SIFixSGPRCopies::run(MachineFunction &MF) {
   lowerVGPR2SGPRCopies(MF);
   // Postprocessing
   fixSCCCopies(MF);
-  for (auto MI : S2VCopies) {
+  for (auto *MI : S2VCopies) {
     // Check if it is still valid
     if (MI->isCopy()) {
       const TargetRegisterClass *SrcRC, *DstRC;
@@ -778,12 +778,12 @@ bool SIFixSGPRCopies::run(MachineFunction &MF) {
         tryChangeVGPRtoSGPRinCopy(*MI, TRI, TII);
     }
   }
-  for (auto MI : RegSequences) {
+  for (auto *MI : RegSequences) {
     // Check if it is still valid
     if (MI->isRegSequence())
       foldVGPRCopyIntoRegSequence(*MI, TRI, TII, *MRI);
   }
-  for (auto MI : PHINodes) {
+  for (auto *MI : PHINodes) {
     processPHINode(*MI);
   }
   if (MF.getTarget().getOptLevel() > CodeGenOptLevel::None && EnableM0Merge)
@@ -968,7 +968,7 @@ void SIFixSGPRCopies::analyzeVGPRToSGPRCopy(MachineInstr* MI) {
         for (auto &U : MRI->use_instructions(Reg))
           Users.push_back(&U);
     }
-    for (auto U : Users) {
+    for (auto *U : Users) {
       if (TII->isSALU(*U))
         Info.SChain.insert(U);
       AnalysisWorklist.push_back(U);
@@ -996,7 +996,7 @@ bool SIFixSGPRCopies::needToBeConvertedToVALU(V2SCopyInfo *Info) {
   // of the same register.
   SmallSet<std::pair<Register, unsigned>, 4> SrcRegs;
   for (auto J : Info->Siblings) {
-    auto InfoIt = V2SCopies.find(J);
+    auto *InfoIt = V2SCopies.find(J);
     if (InfoIt != V2SCopies.end()) {
       MachineInstr *SiblingCopy = InfoIt->second.Copy;
       if (SiblingCopy->isImplicitDef())
@@ -1031,12 +1031,12 @@ void SIFixSGPRCopies::lowerVGPR2SGPRCopies(MachineFunction &MF) {
 
   while (!LoweringWorklist.empty()) {
     unsigned CurID = LoweringWorklist.pop_back_val();
-    auto CurInfoIt = V2SCopies.find(CurID);
+    auto *CurInfoIt = V2SCopies.find(CurID);
     if (CurInfoIt != V2SCopies.end()) {
       V2SCopyInfo C = CurInfoIt->second;
       LLVM_DEBUG(dbgs() << "Processing ...\n"; C.dump());
       for (auto S : C.Siblings) {
-        auto SibInfoIt = V2SCopies.find(S);
+        auto *SibInfoIt = V2SCopies.find(S);
         if (SibInfoIt != V2SCopies.end()) {
           V2SCopyInfo &SI = SibInfoIt->second;
           LLVM_DEBUG(dbgs() << "Sibling:\n"; SI.dump());
diff --git a/llvm/lib/Target/AMDGPU/SIFrameLowering.cpp b/llvm/lib/Target/AMDGPU/SIFrameLowering.cpp
index 3d1657392884f5..bc162b0953a7be 100644
--- a/llvm/lib/Target/AMDGPU/SIFrameLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/SIFrameLowering.cpp
@@ -749,11 +749,11 @@ void SIFrameLowering::emitEntryFunctionScratchRsrcRegSetup(
     // at offset 0 (or offset 16 for a compute shader).
     MachinePointerInfo PtrInfo(AMDGPUAS::CONSTANT_ADDRESS);
     const MCInstrDesc &LoadDwordX4 = TII->get(AMDGPU::S_LOAD_DWORDX4_IMM);
-    auto MMO = MF.getMachineMemOperand(PtrInfo,
-                                       MachineMemOperand::MOLoad |
-                                           MachineMemOperand::MOInvariant |
-                                           MachineMemOperand::MODereferenceable,
-                                       16, Align(4));
+    auto *MMO = MF.getMachineMemOperand(
+        PtrInfo,
+        MachineMemOperand::MOLoad | MachineMemOperand::MOInvariant |
+            MachineMemOperand::MODereferenceable,
+        16, Align(4));
     unsigned Offset = Fn.getCallingConv() == CallingConv::AMDGPU_CS ? 16 : 0;
     const GCNSubtarget &Subtarget = MF.getSubtarget<GCNSubtarget>();
     unsigned EncodedOffset = AMDGPU::convertSMRDOffsetUnits(Subtarget, Offset);
@@ -800,7 +800,7 @@ void SIFrameLowering::emitEntryFunctionScratchRsrcRegSetup(
         const MCInstrDesc &LoadDwordX2 = TII->get(AMDGPU::S_LOAD_DWORDX2_IMM);
 
         MachinePointerInfo PtrInfo(AMDGPUAS::CONSTANT_ADDRESS);
-        auto MMO = MF.getMachineMemOperand(
+        auto *MMO = MF.getMachineMemOperand(
             PtrInfo,
             MachineMemOperand::MOLoad | MachineMemOperand::MOInvariant |
                 MachineMemOperand::MODereferenceable,
diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
index 5e4cf705cc9e47..44d7804647fa02 100644
--- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
@@ -7287,7 +7287,7 @@ SDValue SITargetLowering::lowerINSERT_VECTOR_ELT(SDValue Op,
 
   // Specially handle the case of v4i16 with static indexing.
   unsigned NumElts = VecVT.getVectorNumElements();
-  auto KIdx = dyn_cast<ConstantSDNode>(Idx);
+  auto *KIdx = dyn_cast<ConstantSDNode>(Idx);
   if (NumElts == 4 && EltSize == 16 && KIdx) {
     SDValue BCVec = DAG.getNode(ISD::BITCAST, SL, MVT::v2i32, Vec);
 
@@ -7879,7 +7879,7 @@ static SDValue constructRetValue(SelectionDAG &DAG, MachineSDNode *Result,
 
 static bool parseTexFail(SDValue TexFailCtrl, SelectionDAG &DAG, SDValue *TFE,
                          SDValue *LWE, bool &IsTexFail) {
-  auto TexFailCtrlConst = cast<ConstantSDNode>(TexFailCtrl.getNode());
+  auto *TexFailCtrlConst = cast<ConstantSDNode>(TexFailCtrl.getNode());
 
   uint64_t Value = TexFailCtrlConst->getZExtValue();
   if (Value) {
@@ -8284,7 +8284,7 @@ SDValue SITargetLowering::lowerImage(SDValue Op,
     return Op;
 
   MachineSDNode *NewNode = DAG.getMachineNode(Opcode, DL, ResultTypes, Ops);
-  if (auto MemOp = dyn_cast<MemSDNode>(Op)) {
+  if (auto *MemOp = dyn_cast<MemSDNode>(Op)) {
     MachineMemOperand *MemRef = MemOp->getMemOperand();
     DAG.setNodeMemRefs(NewNode, {MemRef});
   }
@@ -8435,7 +8435,7 @@ SDValue SITargetLowering::lowerWorkitemID(SelectionDAG &DAG, SDValue Op,
 SDValue SITargetLowering::LowerINTRINSIC_WO_CHAIN(SDValue Op,
                                                   SelectionDAG &DAG) const {
   MachineFunction &MF = DAG.getMachineFunction();
-  auto MFI = MF.getInfo<SIMachineFunctionInfo>();
+  auto *MFI = MF.getInfo<SIMachineFunctionInfo>();
 
   EVT VT = Op.getValueType();
   SDLoc DL(Op);
@@ -8755,7 +8755,7 @@ SDValue SITargetLowering::LowerINTRINSIC_WO_CHAIN(SDValue Op,
     Module *M = const_cast<Module *>(MF.getFunction().getParent());
     const MDNode *Metadata = cast<MDNodeSDNode>(Op.getOperand(1))->getMD();
     auto SymbolName = cast<MDString>(Metadata->getOperand(0))->getString();
-    auto RelocSymbol = cast<GlobalVariable>(
+    auto *RelocSymbol = cast<GlobalVariable>(
         M->getOrInsertGlobal(SymbolName, Type::getInt32Ty(M->getContext())));
     SDValue GA = DAG.getTargetGlobalAddress(RelocSymbol, DL, MVT::i32, 0,
                                             SIInstrInfo::MO_ABS32_LO);
@@ -9382,7 +9382,7 @@ SDValue SITargetLowering::LowerINTRINSIC_W_CHAIN(SDValue Op,
       Ops.push_back(M0Val.getValue(0));
     }
 
-    auto NewMI = DAG.getMachineNode(Opc, DL, Op->getVTList(), Ops);
+    auto *NewMI = DAG.getMachineNode(Opc, DL, Op->getVTList(), Ops);
     return SDValue(NewMI, 0);
   }
   default:
@@ -9829,7 +9829,7 @@ SDValue SITargetLowering::LowerINTRINSIC_VOID(SDValue Op,
         StorePtrI, F | MachineMemOperand::MOStore, sizeof(int32_t),
         LoadMMO->getBaseAlign(), LoadMMO->getAAInfo());
 
-    auto Load = DAG.getMachineNode(Opc, DL, M->getVTList(), Ops);
+    auto *Load = DAG.getMachineNode(Opc, DL, M->getVTList(), Ops);
     DAG.setNodeMemRefs(Load, {LoadMMO, StoreMMO});
 
     return SDValue(Load, 0);
@@ -9907,7 +9907,7 @@ SDValue SITargetLowering::LowerINTRINSIC_VOID(SDValue Op,
         StorePtrI, F | MachineMemOperand::MOStore, sizeof(int32_t), Align(4),
         LoadMMO->getAAInfo());
 
-    auto Load = DAG.getMachineNode(Opc, DL, Op->getVTList(), Ops);
+    auto *Load = DAG.getMachineNode(Opc, DL, Op->getVTList(), Ops);
     DAG.setNodeMemRefs(Load, {LoadMMO, StoreMMO});
 
     return SDValue(Load, 0);
@@ -9984,7 +9984,7 @@ SDValue SITargetLowering::LowerINTRINSIC_VOID(SDValue Op,
       Ops.push_back(copyToM0(DAG, Chain, DL, BarOp).getValue(0));
     }
 
-    auto NewMI = DAG.getMachineNode(Opc, DL, Op->getVTList(), Ops);
+    auto *NewMI = DAG.getMachineNode(Opc, DL, Op->getVTList(), Ops);
     return SDValue(NewMI, 0);
   }
   case Intrinsic::amdgcn_s_prefetch_data: {
@@ -11815,7 +11815,7 @@ calculateSrcByte(const SDValue Op, uint64_t DestByte, uint64_t SrcIndex = 0,
 
   case ISD::SRA:
   case ISD::SRL: {
-    auto ShiftOp = dyn_cast<ConstantSDNode>(Op->getOperand(1));
+    auto *ShiftOp = dyn_cast<ConstantSDNode>(Op->getOperand(1));
     if (!ShiftOp)
       return std::nullopt;
 
@@ -11885,7 +11885,7 @@ calculateByteProvider(const SDValue &Op, unsigned Index, unsigned Depth,
     if (IsVec)
       return std::nullopt;
 
-    auto BitMaskOp = dyn_cast<ConstantSDNode>(Op->getOperand(1));
+    auto *BitMaskOp = dyn_cast<ConstantSDNode>(Op->getOperand(1));
     if (!BitMaskOp)
       return std::nullopt;
 
@@ -11909,7 +11909,7 @@ calculateByteProvider(const SDValue &Op, unsigned Index, unsigned Depth,
       return std::nullopt;
 
     // fshr(X,Y,Z): (X << (BW - (Z % BW))) | (Y >> (Z % BW))
-    auto ShiftOp = dyn_cast<ConstantSDNode>(Op->getOperand(2));
+    auto *ShiftOp = dyn_cast<ConstantSDNode>(Op->getOperand(2));
     if (!ShiftOp || Op.getValueType().isVector())
       return std::nullopt;
 
@@ -11936,7 +11936,7 @@ calculateByteProvider(const SDValue &Op, unsigned Index, unsigned Depth,
     if (IsVec)
       return std::nullopt;
 
-    auto ShiftOp = dyn_cast<ConstantSDNode>(Op->getOperand(1));
+    auto *ShiftOp = dyn_cast<ConstantSDNode>(Op->getOperand(1));
     if (!ShiftOp)
       return std::nullopt;
 
@@ -11964,7 +11964,7 @@ calculateByteProvider(const SDValue &Op, unsigned Index, unsigned Depth,
     if (IsVec)
       return std::nullopt;
 
-    auto ShiftOp = dyn_cast<ConstantSDNode>(Op->getOperand(1));
+    auto *ShiftOp = dyn_cast<ConstantSDNode>(Op->getOperand(1));
     if (!ShiftOp)
       return std::nullopt;
 
@@ -12033,7 +12033,7 @@ calculateByteProvider(const SDValue &Op, unsigned Index, unsigned Depth,
   }
 
   case ISD::LOAD: {
-    auto L = cast<LoadSDNode>(Op.getNode());
+    auto *L = cast<LoadSDNode>(Op.getNode());
 
     unsigned NarrowBitWidth = L->getMemoryVT().getSizeInBits();
     if (NarrowBitWidth % 8 != 0)
@@ -12066,7 +12066,7 @@ calculateByteProvider(const SDValue &Op, unsigned Index, unsigned Depth,
   }
 
   case ISD::EXTRACT_VECTOR_ELT: {
-    auto IdxOp = dyn_cast<ConstantSDNode>(Op->getOperand(1));
+    auto *IdxOp = dyn_cast<ConstantSDNode>(Op->getOperand(1));
     if (!IdxOp)
       return std::nullopt;
     auto VecIdx = IdxOp->getZExtValue();
@@ -12081,7 +12081,7 @@ calculateByteProvider(const SDValue &Op, unsigned Index, unsigned Depth,
     if (IsVec)
       return std::nullopt;
 
-    auto PermMask = dyn_cast<ConstantSDNode>(Op->getOperand(2));
+    auto *PermMask = dyn_cast<ConstantSDNode>(Op->getOperand(2));
     if (!PermMask)
       return std::nullopt;
 
@@ -12380,7 +12380,7 @@ SDValue SITargetLowering::performOrCombine(SDNode *N,
         return true;
 
       // If we have any non-vectorized use, then it is a candidate for v_perm
-      for (auto VUse : OrUse->uses()) {
+      for (auto *VUse : OrUse->uses()) {
         if (!VUse->getValueType(0).isVector())
           return true;
 
@@ -13902,7 +13902,7 @@ static void placeSources(ByteProvider<SDValue> &Src0,
                (IterElt.DWordOffset == (BPP.first.SrcOffset / 4));
       };
 
-      auto Match = llvm::find_if(Srcs, MatchesFirst);
+      auto *Match = llvm::find_if(Srcs, MatchesFirst);
       if (Match != Srcs.end()) {
         Match->PermMask = addPermMasks(FirstMask, Match->PermMask);
         FirstGroup = I;
@@ -13915,7 +13915,7 @@ static void placeSources(ByteProvider<SDValue> &Src0,
         return IterElt.SrcOp == *BPP.second.Src &&
                (IterElt.DWordOffset == (BPP.second.SrcOffset / 4));
       };
-      auto Match = llvm::find_if(Srcs, MatchesSecond);
+      auto *Match = llvm::find_if(Srcs, MatchesSecond);
       if (Match != Srcs.end()) {
         Match->PermMask = addPermMasks(SecondMask, Match->PermMask);
       } else
@@ -13948,7 +13948,7 @@ static SDValue resolveSources(SelectionDAG &DAG, SDLoc SL,
 
   // If we just have one source, just permute it accordingly.
   if (Srcs.size() == 1) {
-    auto Elt = Srcs.begin();
+    auto *Elt = Srcs.begin();
     auto EltOp = getDWordFromOffset(DAG, SL, Elt->SrcOp, Elt->DWordOffset);
 
     // v_perm will produce the original value
@@ -13959,8 +13959,8 @@ static SDValue resolveSources(SelectionDAG &DAG, SDLoc SL,
                        DAG.getConstant(Elt->PermMask, SL, MVT::i32));
   }
 
-  auto FirstElt = Srcs.begin();
-  auto SecondElt = std::next(FirstElt);
+  auto *FirstElt = Srcs.begin();
+  auto *SecondElt = std::next(FirstElt);
 
   SmallVector<SDValue, 2> Perms;
 
@@ -14214,11 +14214,11 @@ SDValue SITargetLowering::performAddCombine(SDNode *N,
       if (UniqueEntries) {
         UseOriginalSrc = true;
 
-        auto FirstElt = Src0s.begin();
+        auto *FirstElt = Src0s.begin();
         auto FirstEltOp =
             getDWordFromOffset(DAG, SL, FirstElt->SrcOp, FirstElt->DWordOffset);
 
-        auto SecondElt = Src1s.begin();
+        auto *SecondElt = Src1s.begin();
         auto SecondEltOp = getDWordFromOffset(DAG, SL, SecondElt->SrcOp,
                                               SecondElt->DWordOffset);
 
@@ -14561,7 +14561,7 @@ SDValue SITargetLowering::performSetCCCombine(SDNode *N,
   EVT VT = LHS.getValueType();
   ISD::CondCode CC = cast<CondCodeSDNode>(N->getOperand(2))->get();
 
-  auto CRHS = dyn_cast<ConstantSDNode>(RHS);
+  auto *CRHS = dyn_cast<ConstantSDNode>(RHS);
   if (!CRHS) {
     CRHS = dyn_cast<ConstantSDNode>(LHS);
     if (CRHS) {
diff --git a/llvm/lib/Target/AMDGPU/SIInsertWaitcnts.cpp b/llvm/lib/Target/AMDGPU/SIInsertWaitcnts.cpp
index 2728db064f5db9..15f4114826e401 100644
--- a/llvm/lib/Target/AMDGPU/SIInsertWaitcnts.cpp
+++ b/llvm/lib/Target/AMDGPU/SIInsertWaitcnts.cpp
@@ -2409,7 +2409,7 @@ bool SIInsertWaitcnts::runOnMachineFunction(MachineFunction &MF) {
   const SIMachineFunctionInfo *MFI = MF.getInfo<SIMachineFunctionInfo>();
   MLI = &getAnalysis<MachineLoopInfoWrapperPass>().getLI();
   PDT = &getAnalysis<MachinePostDominatorTreeWrapperPass>().getPostDomTree();
-  if (auto AAR = getAnalysisIfAvailable<AAResultsWrapperPass>())
+  if (auto *AAR = getAnalysisIfAvailable<AAResultsWrapperPass>())
     AA = &AAR->getAAResults();
 
   AMDGPU::IsaVersion IV = AMDGPU::getIsaVersion(ST->getCPU());
@@ -2538,7 +2538,7 @@ bool SIInsertWaitcnts::runOnMachineFunction(MachineFunction &MF) {
       if (Brackets->hasPendingEvent()) {
         BlockInfo *MoveBracketsToSucc = nullptr;
         for (MachineBasicBlock *Succ : MBB->successors()) {
-          auto SuccBII = BlockInfos.find(Succ);
+          auto *SuccBII = BlockInfos.find(Succ);
           BlockInfo &SuccBI = SuccBII->second;
           if (!SuccBI.Incoming) {
             SuccBI.Dirty = true;
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
index 5c39b2a4fc96aa..b379447dd8d4c8 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
@@ -528,13 +528,13 @@ static bool memOpsHaveSameBasePtr(const MachineInstr &MI1,
   if (!MI1.hasOneMemOperand() || !MI2.hasOneMemOperand())
     return false;
 
-  auto MO1 = *MI1.memoperands_begin();
-  auto MO2 = *MI2.memoperands_begin();
+  auto *MO1 = *MI1.memoperands_begin();
+  auto *MO2 = *MI2.memoperands_begin();
   if (MO1->getAddrSpace() != MO2->getAddrSpace())
     return false;
 
-  auto Base1 = MO1->getValue();
-  auto Base2 = MO2->getValue();
+  const auto *Base1 = MO1->getValue();
+  const auto *Base2 = MO2->getValue();
   if (!Base1 || !Base2)
     return false;
   Base1 = getUnderlyingObject(Base1);
@@ -2008,7 +2008,7 @@ void SIInstrInfo::insertNoops(MachineBasicBlock &MBB,
 }
 
 void SIInstrInfo::insertReturn(MachineBasicBlock &MBB) const {
-  auto MF = MBB.getParent();
+  auto *MF = MBB.getParent();
   SIMachineFunctionInfo *Info = MF->getInfo<SIMachineFunctionInfo>();
 
   assert(Info->isEntryFunction());
@@ -3819,7 +3819,7 @@ MachineInstr *SIInstrInfo::convertToThreeAddress(MachineInstr &MI,
         SlotIndex NewIndex = LIS->getInstructionIndex(*MIB).getRegSlot(true);
         auto &LI = LIS->getInterval(Def.getReg());
         auto UpdateDefIndex = [&](LiveRange &LR) {
-          auto S = LR.find(OldIndex);
+          auto *S = LR.find(OldIndex);
           if (S != LR.end() && S->start == OldIndex) {
             assert(S->valno && S->valno->def == OldIndex);
             S->start = NewIndex;
@@ -5061,7 +5061,7 @@ bool SIInstrInfo::verifyInstruction(const MachineInstr &MI,
   }
 
   if (isSOPK(MI)) {
-    auto Op = getNamedOperand(MI, AMDGPU::OpName::simm16);
+    const auto *Op = getNamedOperand(MI, AMDGPU::OpName::simm16);
     if (Desc.isBranch()) {
       if (!Op->isMBB()) {
         ErrInfo = "invalid branch target for SOPK instruction";
@@ -6402,7 +6402,7 @@ static void emitLoadScalarOpsFromVGPRLoop(
         }
       } // End for loop.
 
-      auto SScalarOpRC =
+      const auto *SScalarOpRC =
           TRI->getEquivalentSGPRClass(MRI.getRegClass(VScalarOp));
       Register SScalarOp = MRI.createVirtualRegister(SScalarOpRC);
 
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.h b/llvm/lib/Target/AMDGPU/SIInstrInfo.h
index f7554906a9c98b..7041b59964645a 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.h
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.h
@@ -52,12 +52,12 @@ struct SIInstrWorklist {
   void insert(MachineInstr *MI);
 
   MachineInstr *top() const {
-    auto iter = InstrList.begin();
+    const auto *iter = InstrList.begin();
     return *iter;
   }
 
   void erase_top() {
-    auto iter = InstrList.begin();
+    const auto *iter = InstrList.begin();
     InstrList.erase(iter);
   }
 
diff --git a/llvm/lib/Target/AMDGPU/SILateBranchLowering.cpp b/llvm/lib/Target/AMDGPU/SILateBranchLowering.cpp
index afc6353ec81167..62cda9e7ef2f3e 100644
--- a/llvm/lib/Target/AMDGPU/SILateBranchLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/SILateBranchLowering.cpp
@@ -226,7 +226,7 @@ bool SILateBranchLowering::runOnMachineFunction(MachineFunction &MF) {
     }
 
     for (auto *MI : EpilogInstrs) {
-      auto MBB = MI->getParent();
+      auto *MBB = MI->getParent();
       if (MBB == &MF.back() && MI == &MBB->back())
         continue;
 
diff --git a/llvm/lib/Target/AMDGPU/SILoadStoreOptimizer.cpp b/llvm/lib/Target/AMDGPU/SILoadStoreOptimizer.cpp
index 23d04fae420150..904321a344db11 100644
--- a/llvm/lib/Target/AMDGPU/SILoadStoreOptimizer.cpp
+++ b/llvm/lib/Target/AMDGPU/SILoadStoreOptimizer.cpp
@@ -2056,7 +2056,7 @@ Register SILoadStoreOptimizer::computeBase(MachineInstr &MI,
 void SILoadStoreOptimizer::updateBaseAndOffset(MachineInstr &MI,
                                                Register NewBase,
                                                int32_t NewOffset) const {
-  auto Base = TII->getNamedOperand(MI, AMDGPU::OpName::vaddr);
+  auto *Base = TII->getNamedOperand(MI, AMDGPU::OpName::vaddr);
   Base->setReg(NewBase);
   Base->setIsKill(false);
   TII->getNamedOperand(MI, AMDGPU::OpName::offset)->setImm(NewOffset);
diff --git a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp
index 8be9a082a7fd08..e59dd724b94f8b 100644
--- a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp
@@ -347,7 +347,7 @@ void SIMachineFunctionInfo::shiftWwmVGPRsToLowestRange(
 
     // Replace the register in SpillPhysVGPRs. This is needed to look for free
     // lanes while spilling special SGPRs like FP, BP, etc. during PEI.
-    auto RegItr = std::find(SpillPhysVGPRs.begin(), SpillPhysVGPRs.end(), Reg);
+    auto *RegItr = std::find(SpillPhysVGPRs.begin(), SpillPhysVGPRs.end(), Reg);
     if (RegItr != SpillPhysVGPRs.end()) {
       unsigned Idx = std::distance(SpillPhysVGPRs.begin(), RegItr);
       SpillPhysVGPRs[Idx] = NewReg;
diff --git a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h
index 669f98dd865d61..c8c305e24c7101 100644
--- a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h
+++ b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h
@@ -639,15 +639,17 @@ class SIMachineFunctionInfo final : public AMDGPUMachineFunction,
   // Check if an entry created for \p Reg in PrologEpilogSGPRSpills. Return true
   // on success and false otherwise.
   bool hasPrologEpilogSGPRSpillEntry(Register Reg) const {
-    auto I = find_if(PrologEpilogSGPRSpills,
-                     [&Reg](const auto &Spill) { return Spill.first == Reg; });
+    const auto *I = find_if(PrologEpilogSGPRSpills, [&Reg](const auto &Spill) {
+      return Spill.first == Reg;
+    });
     return I != PrologEpilogSGPRSpills.end();
   }
 
   // Get the scratch SGPR if allocated to save/restore \p Reg.
   Register getScratchSGPRCopyDstReg(Register Reg) const {
-    auto I = find_if(PrologEpilogSGPRSpills,
-                     [&Reg](const auto &Spill) { return Spill.first == Reg; });
+    const auto *I = find_if(PrologEpilogSGPRSpills, [&Reg](const auto &Spill) {
+      return Spill.first == Reg;
+    });
     if (I != PrologEpilogSGPRSpills.end() &&
         I->second.getKind() == SGPRSaveKind::COPY_TO_SCRATCH_SGPR)
       return I->second.getReg();
@@ -676,8 +678,9 @@ class SIMachineFunctionInfo final : public AMDGPUMachineFunction,
 
   const PrologEpilogSGPRSaveRestoreInfo &
   getPrologEpilogSGPRSaveRestoreInfo(Register Reg) const {
-    auto I = find_if(PrologEpilogSGPRSpills,
-                     [&Reg](const auto &Spill) { return Spill.first == Reg; });
+    const auto *I = find_if(PrologEpilogSGPRSpills, [&Reg](const auto &Spill) {
+      return Spill.first == Reg;
+    });
     assert(I != PrologEpilogSGPRSpills.end());
 
     return I->second;
@@ -888,7 +891,7 @@ class SIMachineFunctionInfo final : public AMDGPUMachineFunction,
   }
 
   MCRegister getPreloadedReg(AMDGPUFunctionArgInfo::PreloadedValue Value) const {
-    auto Arg = std::get<0>(ArgInfo.getPreloadedValue(Value));
+    const auto *Arg = std::get<0>(ArgInfo.getPreloadedValue(Value));
     return Arg ? Arg->getRegister() : MCRegister();
   }
 
diff --git a/llvm/lib/Target/AMDGPU/SIOptimizeExecMaskingPreRA.cpp b/llvm/lib/Target/AMDGPU/SIOptimizeExecMaskingPreRA.cpp
index 494b2341f0e0e4..31f65d82a4d2bc 100644
--- a/llvm/lib/Target/AMDGPU/SIOptimizeExecMaskingPreRA.cpp
+++ b/llvm/lib/Target/AMDGPU/SIOptimizeExecMaskingPreRA.cpp
@@ -392,7 +392,7 @@ bool SIOptimizeExecMaskingPreRA::runOnMachineFunction(MachineFunction &MF) {
       SmallVector<MachineBasicBlock*, 4> Blocks({&MBB});
 
       while (!Blocks.empty()) {
-        auto CurBB = Blocks.pop_back_val();
+        auto *CurBB = Blocks.pop_back_val();
         auto I = CurBB->rbegin(), E = CurBB->rend();
         if (I != E) {
           if (I->isUnconditionalBranch() || I->getOpcode() == AMDGPU::S_ENDPGM)
diff --git a/llvm/lib/Target/AMDGPU/SIPeepholeSDWA.cpp b/llvm/lib/Target/AMDGPU/SIPeepholeSDWA.cpp
index 86cb0e6944ed75..467f042892cebe 100644
--- a/llvm/lib/Target/AMDGPU/SIPeepholeSDWA.cpp
+++ b/llvm/lib/Target/AMDGPU/SIPeepholeSDWA.cpp
@@ -710,7 +710,7 @@ SIPeepholeSDWA::matchSDWAOperand(MachineInstr &MI) {
 
     MachineOperand *Src0 = TII->getNamedOperand(MI, AMDGPU::OpName::src0);
     MachineOperand *Src1 = TII->getNamedOperand(MI, AMDGPU::OpName::src1);
-    auto ValSrc = Src1;
+    auto *ValSrc = Src1;
     auto Imm = foldToImm(*Src0);
 
     if (!Imm) {
@@ -1151,7 +1151,7 @@ bool SIPeepholeSDWA::convertToSDWA(MachineInstr &MI,
   }
 
   // Check for a preserved register that needs to be copied.
-  auto DstUnused = TII->getNamedOperand(MI, AMDGPU::OpName::dst_unused);
+  auto *DstUnused = TII->getNamedOperand(MI, AMDGPU::OpName::dst_unused);
   if (DstUnused &&
       DstUnused->getImm() == AMDGPU::SDWA::DstUnused::UNUSED_PRESERVE) {
     // We expect, if we are here, that the instruction was already in it's SDWA form,
diff --git a/llvm/lib/Target/AMDGPU/SIShrinkInstructions.cpp b/llvm/lib/Target/AMDGPU/SIShrinkInstructions.cpp
index 9967a65244413a..78267d402b6c9e 100644
--- a/llvm/lib/Target/AMDGPU/SIShrinkInstructions.cpp
+++ b/llvm/lib/Target/AMDGPU/SIShrinkInstructions.cpp
@@ -752,13 +752,13 @@ MachineInstr *SIShrinkInstructions::matchSwap(MachineInstr &MovT) const {
     MachineBasicBlock &MBB = *MovT.getParent();
     SmallVector<MachineInstr *, 4> Swaps;
     if (Size == 2) {
-      auto MIB = BuildMI(MBB, MovX->getIterator(), MovT.getDebugLoc(),
-                         TII->get(AMDGPU::V_SWAP_B16))
-                     .addDef(X)
-                     .addDef(Y)
-                     .addReg(Y)
-                     .addReg(X)
-                     .getInstr();
+      auto *MIB = BuildMI(MBB, MovX->getIterator(), MovT.getDebugLoc(),
+                          TII->get(AMDGPU::V_SWAP_B16))
+                      .addDef(X)
+                      .addDef(Y)
+                      .addReg(Y)
+                      .addReg(X)
+                      .getInstr();
       Swaps.push_back(MIB);
     } else {
       assert(Size > 0 && Size % 4 == 0);
@@ -766,13 +766,13 @@ MachineInstr *SIShrinkInstructions::matchSwap(MachineInstr &MovT) const {
         TargetInstrInfo::RegSubRegPair X1, Y1;
         X1 = getSubRegForIndex(X, Xsub, I);
         Y1 = getSubRegForIndex(Y, Ysub, I);
-        auto MIB = BuildMI(MBB, MovX->getIterator(), MovT.getDebugLoc(),
-                           TII->get(AMDGPU::V_SWAP_B32))
-                       .addDef(X1.Reg, 0, X1.SubReg)
-                       .addDef(Y1.Reg, 0, Y1.SubReg)
-                       .addReg(Y1.Reg, 0, Y1.SubReg)
-                       .addReg(X1.Reg, 0, X1.SubReg)
-                       .getInstr();
+        auto *MIB = BuildMI(MBB, MovX->getIterator(), MovT.getDebugLoc(),
+                            TII->get(AMDGPU::V_SWAP_B32))
+                        .addDef(X1.Reg, 0, X1.SubReg)
+                        .addDef(Y1.Reg, 0, Y1.SubReg)
+                        .addReg(Y1.Reg, 0, Y1.SubReg)
+                        .addReg(X1.Reg, 0, X1.SubReg)
+                        .getInstr();
         Swaps.push_back(MIB);
       }
     }
diff --git a/llvm/lib/Target/AMDGPU/SIWholeQuadMode.cpp b/llvm/lib/Target/AMDGPU/SIWholeQuadMode.cpp
index ef6c92dfa9b9f2..fba1d0c0269589 100644
--- a/llvm/lib/Target/AMDGPU/SIWholeQuadMode.cpp
+++ b/llvm/lib/Target/AMDGPU/SIWholeQuadMode.cpp
@@ -366,8 +366,8 @@ void SIWholeQuadMode::markDefs(const MachineInstr &UseMI, LiveRange &LR,
 
       // Find next predecessor to process
       unsigned Idx = NextPredIdx;
-      auto PI = MBB->pred_begin() + Idx;
-      auto PE = MBB->pred_end();
+      const auto *PI = MBB->pred_begin() + Idx;
+      const auto *PE = MBB->pred_end();
       for (; PI != PE && !NextValue; ++PI, ++Idx) {
         if (const VNInfo *VN = LR.getVNInfoBefore(LIS->getMBBEndIdx(*PI))) {
           if (!Visited.count(VisitKey(VN, DefinedLanes)))
@@ -1036,7 +1036,7 @@ MachineInstr *SIWholeQuadMode::lowerKillI1(MachineBasicBlock &MBB,
 // This can only happen once all the live mask registers have been created
 // and the execute state (WQM/StrictWWM/Exact) of instructions is known.
 void SIWholeQuadMode::lowerBlock(MachineBasicBlock &MBB) {
-  auto BII = Blocks.find(&MBB);
+  auto *BII = Blocks.find(&MBB);
   if (BII == Blocks.end())
     return;
 
@@ -1261,7 +1261,7 @@ void SIWholeQuadMode::fromStrictMode(MachineBasicBlock &MBB,
 }
 
 void SIWholeQuadMode::processBlock(MachineBasicBlock &MBB, bool IsEntry) {
-  auto BII = Blocks.find(&MBB);
+  auto *BII = Blocks.find(&MBB);
   if (BII == Blocks.end())
     return;
 
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
index feec7b47ae294b..20a81a3135f0b2 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
@@ -171,7 +171,7 @@ bool isHsaAbi(const MCSubtargetInfo &STI) {
 }
 
 unsigned getAMDHSACodeObjectVersion(const Module &M) {
-  if (auto Ver = mdconst::extract_or_null<ConstantInt>(
+  if (auto *Ver = mdconst::extract_or_null<ConstantInt>(
           M.getModuleFlag("amdhsa_code_object_version"))) {
     return (unsigned)Ver->getZExtValue() / 100;
   }
@@ -649,8 +649,8 @@ int getVOPDFull(unsigned OpX, unsigned OpY, unsigned EncodingFamily) {
 std::pair<unsigned, unsigned> getVOPDComponents(unsigned VOPDOpcode) {
   const VOPDInfo *Info = getVOPDOpcodeHelper(VOPDOpcode);
   assert(Info);
-  auto OpX = getVOPDBaseFromComponent(Info->OpX);
-  auto OpY = getVOPDBaseFromComponent(Info->OpY);
+  const auto *OpX = getVOPDBaseFromComponent(Info->OpX);
+  const auto *OpY = getVOPDBaseFromComponent(Info->OpY);
   assert(OpX && OpY);
   return {OpX->BaseVOP, OpY->BaseVOP};
 }
@@ -1789,7 +1789,7 @@ static StringLiteral const *getNfmtLookupTable(const MCSubtargetInfo &STI) {
 }
 
 int64_t getNfmt(const StringRef Name, const MCSubtargetInfo &STI) {
-  auto lookupTable = getNfmtLookupTable(STI);
+  const auto *lookupTable = getNfmtLookupTable(STI);
   for (int Id = NFMT_MIN; Id <= NFMT_MAX; ++Id) {
     if (Name == lookupTable[Id])
       return Id;
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUPALMetadata.cpp b/llvm/lib/Target/AMDGPU/Utils/AMDGPUPALMetadata.cpp
index 92d09b3afa77d7..4ad26ee895c7dd 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUPALMetadata.cpp
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUPALMetadata.cpp
@@ -29,14 +29,14 @@ using namespace llvm::AMDGPU;
 
 // Read the PAL metadata from IR metadata, where it was put by the frontend.
 void AMDGPUPALMetadata::readFromIR(Module &M) {
-  auto NamedMD = M.getNamedMetadata("amdgpu.pal.metadata.msgpack");
+  auto *NamedMD = M.getNamedMetadata("amdgpu.pal.metadata.msgpack");
   if (NamedMD && NamedMD->getNumOperands()) {
     // This is the new msgpack format for metadata. It is a NamedMD containing
     // an MDTuple containing an MDString containing the msgpack data.
     BlobType = ELF::NT_AMDGPU_METADATA;
-    auto MDN = dyn_cast<MDTuple>(NamedMD->getOperand(0));
+    auto *MDN = dyn_cast<MDTuple>(NamedMD->getOperand(0));
     if (MDN && MDN->getNumOperands()) {
-      if (auto MDS = dyn_cast<MDString>(MDN->getOperand(0)))
+      if (auto *MDS = dyn_cast<MDString>(MDN->getOperand(0)))
         setFromMsgPackBlob(MDS->getString());
     }
     return;
@@ -52,12 +52,12 @@ void AMDGPUPALMetadata::readFromIR(Module &M) {
   // containing an MDTuple containing a number of MDNodes each of which is an
   // integer value, and each two integer values forms a key=value pair that we
   // store as Registers[key]=value in the map.
-  auto Tuple = dyn_cast<MDTuple>(NamedMD->getOperand(0));
+  auto *Tuple = dyn_cast<MDTuple>(NamedMD->getOperand(0));
   if (!Tuple)
     return;
   for (unsigned I = 0, E = Tuple->getNumOperands() & -2; I != E; I += 2) {
-    auto Key = mdconst::dyn_extract<ConstantInt>(Tuple->getOperand(I));
-    auto Val = mdconst::dyn_extract<ConstantInt>(Tuple->getOperand(I + 1));
+    auto *Key = mdconst::dyn_extract<ConstantInt>(Tuple->getOperand(I));
+    auto *Val = mdconst::dyn_extract<ConstantInt>(Tuple->getOperand(I + 1));
     if (!Key || !Val)
       continue;
     setRegister(Key->getZExtValue(), Val->getZExtValue());
@@ -76,7 +76,7 @@ bool AMDGPUPALMetadata::setFromBlob(unsigned Type, StringRef Blob) {
 
 // Set PAL metadata from legacy (array of key=value pairs) blob.
 bool AMDGPUPALMetadata::setFromLegacyBlob(StringRef Blob) {
-  auto Data = reinterpret_cast<const uint32_t *>(Blob.data());
+  const auto *Data = reinterpret_cast<const uint32_t *>(Blob.data());
   for (unsigned I = 0; I != Blob.size() / sizeof(uint32_t) / 2; ++I)
     setRegister(Data[I * 2], Data[I * 2 + 1]);
   return true;
@@ -751,7 +751,7 @@ static const char *getRegisterName(unsigned RegNum) {
       {0x2c75, "SPI_SHADER_USER_ACCUM_VS_3"},
 
       {0, nullptr}};
-  auto Entry = RegInfoTable;
+  const auto *Entry = RegInfoTable;
   for (; Entry->Num && Entry->Num != RegNum; ++Entry)
     ;
   return Entry->Name;
@@ -1070,13 +1070,13 @@ msgpack::DocNode *AMDGPUPALMetadata::refComputeRegister(StringRef field) {
 }
 
 bool AMDGPUPALMetadata::checkComputeRegisters(StringRef field, unsigned Val) {
-  if (auto N = refComputeRegister(field))
+  if (auto *N = refComputeRegister(field))
     return N->getUInt() == Val;
   return false;
 }
 
 bool AMDGPUPALMetadata::checkComputeRegisters(StringRef field, bool Val) {
-  if (auto N = refComputeRegister(field))
+  if (auto *N = refComputeRegister(field))
     return N->getBool() == Val;
   return false;
 }