[llvm] 8d13e7b - [AMDGPU] Qualify auto. NFC. (#110878)
via llvm-commits
llvm-commits at lists.llvm.org
Thu Oct 3 05:08:00 PDT 2024
Author: Jay Foad
Date: 2024-10-03T13:07:54+01:00
New Revision: 8d13e7b8c382499c1cf0c2a3184b483e760f266b
URL: https://github.com/llvm/llvm-project/commit/8d13e7b8c382499c1cf0c2a3184b483e760f266b
DIFF: https://github.com/llvm/llvm-project/commit/8d13e7b8c382499c1cf0c2a3184b483e760f266b.diff
LOG: [AMDGPU] Qualify auto. NFC. (#110878)
Generated automatically with:
$ clang-tidy -fix -checks=-*,llvm-qualified-auto $(find
lib/Target/AMDGPU/ -type f)
Added:
Modified:
llvm/lib/Target/AMDGPU/AMDGPUAsanInstrumentation.cpp
llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp
llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp
llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp
llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp
llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp
llvm/lib/Target/AMDGPU/AMDGPUImageIntrinsicOptimizer.cpp
llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp
llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp
llvm/lib/Target/AMDGPU/AMDGPULateCodeGenPrepare.cpp
llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp
llvm/lib/Target/AMDGPU/AMDGPULowerKernelAttributes.cpp
llvm/lib/Target/AMDGPU/AMDGPULowerModuleLDSPass.cpp
llvm/lib/Target/AMDGPU/AMDGPUMCInstLower.h
llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp
llvm/lib/Target/AMDGPU/AMDGPUPerfHintAnalysis.cpp
llvm/lib/Target/AMDGPU/AMDGPUPrintfRuntimeBinding.cpp
llvm/lib/Target/AMDGPU/AMDGPUSubtarget.cpp
llvm/lib/Target/AMDGPU/AMDGPUSwLowerLDS.cpp
llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
llvm/lib/Target/AMDGPU/AMDGPUTargetTransformInfo.cpp
llvm/lib/Target/AMDGPU/AMDGPUUnifyMetadata.cpp
llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp
llvm/lib/Target/AMDGPU/GCNDPPCombine.cpp
llvm/lib/Target/AMDGPU/GCNILPSched.cpp
llvm/lib/Target/AMDGPU/GCNIterativeScheduler.cpp
llvm/lib/Target/AMDGPU/GCNMinRegStrategy.cpp
llvm/lib/Target/AMDGPU/GCNNSAReassign.cpp
llvm/lib/Target/AMDGPU/GCNRegPressure.cpp
llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUMCExpr.cpp
llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUMCKernelDescriptor.cpp
llvm/lib/Target/AMDGPU/R600OpenCLImageTypeLoweringPass.cpp
llvm/lib/Target/AMDGPU/SIFixSGPRCopies.cpp
llvm/lib/Target/AMDGPU/SIFrameLowering.cpp
llvm/lib/Target/AMDGPU/SIISelLowering.cpp
llvm/lib/Target/AMDGPU/SIInsertWaitcnts.cpp
llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
llvm/lib/Target/AMDGPU/SIInstrInfo.h
llvm/lib/Target/AMDGPU/SILateBranchLowering.cpp
llvm/lib/Target/AMDGPU/SILoadStoreOptimizer.cpp
llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp
llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h
llvm/lib/Target/AMDGPU/SIOptimizeExecMaskingPreRA.cpp
llvm/lib/Target/AMDGPU/SIPeepholeSDWA.cpp
llvm/lib/Target/AMDGPU/SIShrinkInstructions.cpp
llvm/lib/Target/AMDGPU/SIWholeQuadMode.cpp
llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
llvm/lib/Target/AMDGPU/Utils/AMDGPUPALMetadata.cpp
Removed:
################################################################################
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 2d48ecd2e257ba..2c84cdac76d027 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;
}
More information about the llvm-commits
mailing list