Skip to content

Commit 8d13e7b

Browse files
jayfoadjayfoad
authored
[AMDGPU] Qualify auto. NFC. (#110878)
Generated automatically with: $ clang-tidy -fix -checks=-*,llvm-qualified-auto $(find lib/Target/AMDGPU/ -type f)
1 parent 5901463 commit 8d13e7b

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

50 files changed

+269
-262
lines changed

‎llvm/lib/Target/AMDGPU/AMDGPUAsanInstrumentation.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -235,7 +235,7 @@ void getInterestingMemoryOperands(
235235
Interesting.emplace_back(I, XCHG->getPointerOperandIndex(), true,
236236
XCHG->getCompareOperand()->getType(),
237237
std::nullopt);
238-
} else if (auto CI = dyn_cast<CallInst>(I)) {
238+
} else if (auto *CI = dyn_cast<CallInst>(I)) {
239239
switch (CI->getIntrinsicID()) {
240240
case Intrinsic::masked_load:
241241
case Intrinsic::masked_store:
@@ -257,7 +257,7 @@ void getInterestingMemoryOperands(
257257
case Intrinsic::masked_compressstore: {
258258
bool IsWrite = CI->getIntrinsicID() == Intrinsic::masked_compressstore;
259259
unsigned OpOffset = IsWrite ? 1 : 0;
260-
auto BasePtr = CI->getOperand(OpOffset);
260+
auto *BasePtr = CI->getOperand(OpOffset);
261261
MaybeAlign Alignment = BasePtr->getPointerAlignment(DL);
262262
Type *Ty = IsWrite ? CI->getArgOperand(0)->getType() : CI->getType();
263263
IRBuilder<> IB(I);

‎llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -332,7 +332,7 @@ void AMDGPUAsmPrinter::emitGlobalVariable(const GlobalVariable *GV) {
332332

333333
emitVisibility(GVSym, GV->getVisibility(), !GV->isDeclaration());
334334
emitLinkage(GV, GVSym);
335-
auto TS = getTargetStreamer();
335+
auto *TS = getTargetStreamer();
336336
TS->emitAMDGPULDS(GVSym, Size, Alignment);
337337
return;
338338
}
@@ -1238,8 +1238,8 @@ void AMDGPUAsmPrinter::getSIProgramInfo(SIProgramInfo &ProgInfo,
12381238
// return ((Dst & ~Mask) | (Value << Shift))
12391239
auto SetBits = [&Ctx](const MCExpr *Dst, const MCExpr *Value, uint32_t Mask,
12401240
uint32_t Shift) {
1241-
auto Shft = MCConstantExpr::create(Shift, Ctx);
1242-
auto Msk = MCConstantExpr::create(Mask, Ctx);
1241+
const auto *Shft = MCConstantExpr::create(Shift, Ctx);
1242+
const auto *Msk = MCConstantExpr::create(Mask, Ctx);
12431243
Dst = MCBinaryExpr::createAnd(Dst, MCUnaryExpr::createNot(Msk, Ctx), Ctx);
12441244
Dst = MCBinaryExpr::createOr(
12451245
Dst, MCBinaryExpr::createShl(Value, Shft, Ctx), Ctx);
@@ -1414,7 +1414,7 @@ void AMDGPUAsmPrinter::EmitPALMetadata(const MachineFunction &MF,
14141414
const SIProgramInfo &CurrentProgramInfo) {
14151415
const SIMachineFunctionInfo *MFI = MF.getInfo<SIMachineFunctionInfo>();
14161416
auto CC = MF.getFunction().getCallingConv();
1417-
auto MD = getTargetStreamer()->getPALMetadata();
1417+
auto *MD = getTargetStreamer()->getPALMetadata();
14181418
auto &Ctx = MF.getContext();
14191419

14201420
MD->setEntryPoint(CC, MF.getFunction().getName());

‎llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -142,7 +142,7 @@ struct AMDGPUIncomingArgHandler : public CallLowering::IncomingValueHandler {
142142
const CCValAssign &VA) override {
143143
MachineFunction &MF = MIRBuilder.getMF();
144144

145-
auto MMO = MF.getMachineMemOperand(
145+
auto *MMO = MF.getMachineMemOperand(
146146
MPO, MachineMemOperand::MOLoad | MachineMemOperand::MOInvariant, MemTy,
147147
inferAlignFromPtrInfo(MF, MPO));
148148
MIRBuilder.buildLoad(ValVReg, Addr, *MMO);
@@ -244,7 +244,7 @@ struct AMDGPUOutgoingArgHandler : public AMDGPUOutgoingValueHandler {
244244
uint64_t LocMemOffset = VA.getLocMemOffset();
245245
const auto &ST = MF.getSubtarget<GCNSubtarget>();
246246

247-
auto MMO = MF.getMachineMemOperand(
247+
auto *MMO = MF.getMachineMemOperand(
248248
MPO, MachineMemOperand::MOStore, MemTy,
249249
commonAlignment(ST.getStackAlignment(), LocMemOffset));
250250
MIRBuilder.buildStore(ValVReg, Addr, *MMO);
@@ -1007,7 +1007,7 @@ bool AMDGPUCallLowering::doCallerAndCalleePassArgsTheSameWay(
10071007
const GCNSubtarget &ST = MF.getSubtarget<GCNSubtarget>();
10081008

10091009
// Make sure that the caller and callee preserve all of the same registers.
1010-
auto TRI = ST.getRegisterInfo();
1010+
const auto *TRI = ST.getRegisterInfo();
10111011

10121012
const uint32_t *CallerPreserved = TRI->getCallPreservedMask(MF, CallerCC);
10131013
const uint32_t *CalleePreserved = TRI->getCallPreservedMask(MF, CalleeCC);
@@ -1219,7 +1219,7 @@ bool AMDGPUCallLowering::lowerTailCall(
12191219
if (!ExecArg.Ty->isIntegerTy(ST.getWavefrontSize()))
12201220
return false;
12211221

1222-
if (auto CI = dyn_cast<ConstantInt>(ExecArg.OrigValue)) {
1222+
if (const auto *CI = dyn_cast<ConstantInt>(ExecArg.OrigValue)) {
12231223
MIB.addImm(CI->getSExtValue());
12241224
} else {
12251225
MIB.addReg(ExecArg.Regs[0]);

‎llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp

Lines changed: 17 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -163,8 +163,8 @@ std::string MetadataStreamerMsgPackV4::getTypeName(Type *Ty,
163163
case Type::DoubleTyID:
164164
return "double";
165165
case Type::FixedVectorTyID: {
166-
auto VecTy = cast<FixedVectorType>(Ty);
167-
auto ElTy = VecTy->getElementType();
166+
auto *VecTy = cast<FixedVectorType>(Ty);
167+
auto *ElTy = VecTy->getElementType();
168168
auto NumElements = VecTy->getNumElements();
169169
return (Twine(getTypeName(ElTy, Signed)) + Twine(NumElements)).str();
170170
}
@@ -199,7 +199,7 @@ void MetadataStreamerMsgPackV4::emitTargetID(
199199
}
200200

201201
void MetadataStreamerMsgPackV4::emitPrintf(const Module &Mod) {
202-
auto Node = Mod.getNamedMetadata("llvm.printf.fmts");
202+
auto *Node = Mod.getNamedMetadata("llvm.printf.fmts");
203203
if (!Node)
204204
return;
205205

@@ -214,10 +214,10 @@ void MetadataStreamerMsgPackV4::emitPrintf(const Module &Mod) {
214214
void MetadataStreamerMsgPackV4::emitKernelLanguage(const Function &Func,
215215
msgpack::MapDocNode Kern) {
216216
// TODO: What about other languages?
217-
auto Node = Func.getParent()->getNamedMetadata("opencl.ocl.version");
217+
auto *Node = Func.getParent()->getNamedMetadata("opencl.ocl.version");
218218
if (!Node || !Node->getNumOperands())
219219
return;
220-
auto Op0 = Node->getOperand(0);
220+
auto *Op0 = Node->getOperand(0);
221221
if (Op0->getNumOperands() <= 1)
222222
return;
223223

@@ -233,11 +233,11 @@ void MetadataStreamerMsgPackV4::emitKernelLanguage(const Function &Func,
233233
void MetadataStreamerMsgPackV4::emitKernelAttrs(const Function &Func,
234234
msgpack::MapDocNode Kern) {
235235

236-
if (auto Node = Func.getMetadata("reqd_work_group_size"))
236+
if (auto *Node = Func.getMetadata("reqd_work_group_size"))
237237
Kern[".reqd_workgroup_size"] = getWorkGroupDimensions(Node);
238-
if (auto Node = Func.getMetadata("work_group_size_hint"))
238+
if (auto *Node = Func.getMetadata("work_group_size_hint"))
239239
Kern[".workgroup_size_hint"] = getWorkGroupDimensions(Node);
240-
if (auto Node = Func.getMetadata("vec_type_hint")) {
240+
if (auto *Node = Func.getMetadata("vec_type_hint")) {
241241
Kern[".vec_type_hint"] = Kern.getDocument()->getNode(
242242
getTypeName(
243243
cast<ValueAsMetadata>(Node->getOperand(0))->getType(),
@@ -271,7 +271,7 @@ void MetadataStreamerMsgPackV4::emitKernelArgs(const MachineFunction &MF,
271271
void MetadataStreamerMsgPackV4::emitKernelArg(const Argument &Arg,
272272
unsigned &Offset,
273273
msgpack::ArrayDocNode Args) {
274-
auto Func = Arg.getParent();
274+
const auto *Func = Arg.getParent();
275275
auto ArgNo = Arg.getArgNo();
276276
const MDNode *Node;
277277

@@ -317,7 +317,7 @@ void MetadataStreamerMsgPackV4::emitKernelArg(const Argument &Arg,
317317
Type *Ty = Arg.hasByRefAttr() ? Arg.getParamByRefType() : Arg.getType();
318318

319319
// FIXME: Need to distinguish in memory alignment from pointer alignment.
320-
if (auto PtrTy = dyn_cast<PointerType>(Ty)) {
320+
if (auto *PtrTy = dyn_cast<PointerType>(Ty)) {
321321
if (PtrTy->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS)
322322
PointeeAlign = Arg.getParamAlign().valueOrOne();
323323
}
@@ -353,7 +353,7 @@ void MetadataStreamerMsgPackV4::emitKernelArg(
353353
if (PointeeAlign)
354354
Arg[".pointee_align"] = Arg.getDocument()->getNode(PointeeAlign->value());
355355

356-
if (auto PtrTy = dyn_cast<PointerType>(Ty))
356+
if (auto *PtrTy = dyn_cast<PointerType>(Ty))
357357
if (auto Qualifier = getAddressSpaceQualifier(PtrTy->getAddressSpace()))
358358
// Limiting address space to emit only for a certain ValueKind.
359359
if (ValueKind == "global_buffer" || ValueKind == "dynamic_shared_pointer")
@@ -393,7 +393,7 @@ void MetadataStreamerMsgPackV4::emitHiddenKernelArgs(
393393

394394
const Module *M = Func.getParent();
395395
auto &DL = M->getDataLayout();
396-
auto Int64Ty = Type::getInt64Ty(Func.getContext());
396+
auto *Int64Ty = Type::getInt64Ty(Func.getContext());
397397

398398
Offset = alignTo(Offset, ST.getAlignmentForImplicitArgPtr());
399399

@@ -407,7 +407,7 @@ void MetadataStreamerMsgPackV4::emitHiddenKernelArgs(
407407
emitKernelArg(DL, Int64Ty, Align(8), "hidden_global_offset_z", Offset,
408408
Args);
409409

410-
auto Int8PtrTy =
410+
auto *Int8PtrTy =
411411
PointerType::get(Func.getContext(), AMDGPUAS::GLOBAL_ADDRESS);
412412

413413
if (HiddenArgNumBytes >= 32) {
@@ -592,9 +592,9 @@ void MetadataStreamerMsgPackV5::emitHiddenKernelArgs(
592592
auto &DL = M->getDataLayout();
593593
const SIMachineFunctionInfo &MFI = *MF.getInfo<SIMachineFunctionInfo>();
594594

595-
auto Int64Ty = Type::getInt64Ty(Func.getContext());
596-
auto Int32Ty = Type::getInt32Ty(Func.getContext());
597-
auto Int16Ty = Type::getInt16Ty(Func.getContext());
595+
auto *Int64Ty = Type::getInt64Ty(Func.getContext());
596+
auto *Int32Ty = Type::getInt32Ty(Func.getContext());
597+
auto *Int16Ty = Type::getInt16Ty(Func.getContext());
598598

599599
Offset = alignTo(Offset, ST.getAlignmentForImplicitArgPtr());
600600
emitKernelArg(DL, Int32Ty, Align(4), "hidden_block_count_x", Offset, Args);
@@ -621,7 +621,7 @@ void MetadataStreamerMsgPackV5::emitHiddenKernelArgs(
621621
emitKernelArg(DL, Int16Ty, Align(2), "hidden_grid_dims", Offset, Args);
622622

623623
Offset += 6; // Reserved.
624-
auto Int8PtrTy =
624+
auto *Int8PtrTy =
625625
PointerType::get(Func.getContext(), AMDGPUAS::GLOBAL_ADDRESS);
626626

627627
if (M->getNamedMetadata("llvm.printf.fmts")) {

‎llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp

Lines changed: 21 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -394,7 +394,7 @@ void PipelineSolver::reset() {
394394
for (auto &SG : SyncPipeline) {
395395
SmallVector<SUnit *, 32> TempCollection = SG.Collection;
396396
SG.Collection.clear();
397-
auto SchedBarr = llvm::find_if(TempCollection, [](SUnit *SU) {
397+
auto *SchedBarr = llvm::find_if(TempCollection, [](SUnit *SU) {
398398
return SU->getInstr()->getOpcode() == AMDGPU::SCHED_GROUP_BARRIER;
399399
});
400400
if (SchedBarr != TempCollection.end())
@@ -421,7 +421,7 @@ void PipelineSolver::convertSyncMapsToArrays() {
421421
std::pair(SUsToCandSGs.first, SUsToCandSGs.second));
422422
continue;
423423
}
424-
auto SortPosition = PipelineInstrs[PipelineIDx].begin();
424+
auto *SortPosition = PipelineInstrs[PipelineIDx].begin();
425425
// Insert them in sorted order -- this allows for good parsing order in
426426
// the greedy algorithm
427427
while (SortPosition != PipelineInstrs[PipelineIDx].end() &&
@@ -515,7 +515,7 @@ void PipelineSolver::removeEdges(
515515
SUnit *Pred = PredSuccPair.first;
516516
SUnit *Succ = PredSuccPair.second;
517517

518-
auto Match = llvm::find_if(
518+
auto *Match = llvm::find_if(
519519
Succ->Preds, [&Pred](SDep &P) { return P.getSUnit() == Pred; });
520520
if (Match != Succ->Preds.end()) {
521521
assert(Match->isArtificial());
@@ -639,8 +639,8 @@ bool PipelineSolver::solveExact() {
639639
: populateReadyList(ReadyList, CurrSU.second.begin(),
640640
CurrSU.second.end());
641641

642-
auto I = ReadyList.begin();
643-
auto E = ReadyList.end();
642+
auto *I = ReadyList.begin();
643+
auto *E = ReadyList.end();
644644
for (; I != E; ++I) {
645645
// If we are trying SGs in least cost order, and the current SG is cost
646646
// infeasible, then all subsequent SGs will also be cost infeasible, so we
@@ -942,7 +942,7 @@ class MFMAExpInterleaveOpt final : public IGLPStrategy {
942942
bool apply(const SUnit *SU, const ArrayRef<SUnit *> Collection,
943943
SmallVectorImpl<SchedGroup> &SyncPipe) override {
944944

945-
auto DAG = SyncPipe[0].DAG;
945+
auto *DAG = SyncPipe[0].DAG;
946946

947947
if (Cache->empty()) {
948948
auto I = DAG->SUnits.rbegin();
@@ -976,7 +976,7 @@ class MFMAExpInterleaveOpt final : public IGLPStrategy {
976976
SmallVectorImpl<SchedGroup> &SyncPipe) override {
977977
bool FoundTrans = false;
978978
unsigned Counter = 1;
979-
auto DAG = SyncPipe[0].DAG;
979+
auto *DAG = SyncPipe[0].DAG;
980980

981981
if (Cache->empty()) {
982982
SmallVector<SUnit *, 8> Worklist;
@@ -1016,13 +1016,13 @@ class MFMAExpInterleaveOpt final : public IGLPStrategy {
10161016
public:
10171017
bool apply(const SUnit *SU, const ArrayRef<SUnit *> Collection,
10181018
SmallVectorImpl<SchedGroup> &SyncPipe) override {
1019-
auto DAG = SyncPipe[0].DAG;
1019+
auto *DAG = SyncPipe[0].DAG;
10201020

10211021
if (!SU || !TII->isMFMAorWMMA(*ChainSeed->getInstr()))
10221022
return false;
10231023

10241024
if (Cache->empty()) {
1025-
auto TempSU = ChainSeed;
1025+
auto *TempSU = ChainSeed;
10261026
auto Depth = Number;
10271027
while (Depth > 0) {
10281028
--Depth;
@@ -1232,7 +1232,7 @@ class MFMAExpInterleaveOpt final : public IGLPStrategy {
12321232
if (!OtherGroup->Collection.size())
12331233
return true;
12341234

1235-
auto DAG = SyncPipe[0].DAG;
1235+
auto *DAG = SyncPipe[0].DAG;
12361236

12371237
for (auto &OtherEle : OtherGroup->Collection)
12381238
if (DAG->IsReachable(const_cast<SUnit *>(SU), OtherEle))
@@ -1275,7 +1275,7 @@ class MFMAExpInterleaveOpt final : public IGLPStrategy {
12751275
return false;
12761276

12771277
if (Cache->empty()) {
1278-
auto TempSU = ChainSeed;
1278+
auto *TempSU = ChainSeed;
12791279
auto Depth = Number;
12801280
while (Depth > 0) {
12811281
--Depth;
@@ -1315,7 +1315,7 @@ class MFMAExpInterleaveOpt final : public IGLPStrategy {
13151315
SmallVectorImpl<SchedGroup> &SyncPipe) override {
13161316

13171317
SmallVector<SUnit *, 12> Worklist;
1318-
auto DAG = SyncPipe[0].DAG;
1318+
auto *DAG = SyncPipe[0].DAG;
13191319
if (Cache->empty()) {
13201320
for (auto &SU : DAG->SUnits)
13211321
if (TII->isTRANS(SU.getInstr()->getOpcode())) {
@@ -1509,7 +1509,7 @@ bool MFMAExpInterleaveOpt::analyzeDAG(const SIInstrInfo *TII) {
15091509
return isBitPack(Opc);
15101510
});
15111511

1512-
auto PackPred =
1512+
auto *PackPred =
15131513
std::find_if((*TempMFMA)->Preds.begin(), (*TempMFMA)->Preds.end(),
15141514
[&isBitPack](SDep &Pred) {
15151515
auto Opc = Pred.getSUnit()->getInstr()->getOpcode();
@@ -1868,7 +1868,7 @@ class MFMASmallGemmSingleWaveOpt final : public IGLPStrategy {
18681868
}
18691869

18701870
assert(Cache->size());
1871-
auto DAG = SyncPipe[0].DAG;
1871+
auto *DAG = SyncPipe[0].DAG;
18721872
for (auto &Elt : *Cache) {
18731873
if (DAG->IsReachable(Elt, const_cast<SUnit *>(SU)))
18741874
return true;
@@ -1886,7 +1886,7 @@ class MFMASmallGemmSingleWaveOpt final : public IGLPStrategy {
18861886
public:
18871887
bool apply(const SUnit *SU, const ArrayRef<SUnit *> Collection,
18881888
SmallVectorImpl<SchedGroup> &SyncPipe) override {
1889-
auto MI = SU->getInstr();
1889+
auto *MI = SU->getInstr();
18901890
if (MI->getOpcode() != AMDGPU::V_PERM_B32_e64)
18911891
return false;
18921892

@@ -1952,7 +1952,7 @@ class MFMASmallGemmSingleWaveOpt final : public IGLPStrategy {
19521952
public:
19531953
bool apply(const SUnit *SU, const ArrayRef<SUnit *> Collection,
19541954
SmallVectorImpl<SchedGroup> &SyncPipe) override {
1955-
auto MI = SU->getInstr();
1955+
auto *MI = SU->getInstr();
19561956
if (MI->getOpcode() == TargetOpcode::BUNDLE)
19571957
return false;
19581958
if (!Collection.size())
@@ -2023,7 +2023,7 @@ class MFMASmallGemmSingleWaveOpt final : public IGLPStrategy {
20232023
return false;
20242024
}
20252025

2026-
auto DAG = SyncPipe[0].DAG;
2026+
auto *DAG = SyncPipe[0].DAG;
20272027
// Does the previous DS_WRITE share a V_PERM predecessor with this
20282028
// VMEM_READ
20292029
return llvm::any_of(*Cache, [&SU, &DAG](SUnit *Elt) {
@@ -2070,7 +2070,7 @@ bool MFMASmallGemmSingleWaveOpt::applyIGLPStrategy(
20702070
"DSWCounters should be zero in pre-RA scheduling!");
20712071
SmallVector<SUnit *, 6> DSWithPerms;
20722072
for (auto &SU : DAG->SUnits) {
2073-
auto I = SU.getInstr();
2073+
auto *I = SU.getInstr();
20742074
if (TII->isMFMAorWMMA(*I))
20752075
++MFMACount;
20762076
else if (TII->isDS(*I)) {
@@ -2091,8 +2091,8 @@ bool MFMASmallGemmSingleWaveOpt::applyIGLPStrategy(
20912091

20922092
if (IsInitial) {
20932093
DSWWithPermCount = DSWithPerms.size();
2094-
auto I = DSWithPerms.begin();
2095-
auto E = DSWithPerms.end();
2094+
auto *I = DSWithPerms.begin();
2095+
auto *E = DSWithPerms.end();
20962096

20972097
// Get the count of DS_WRITES with V_PERM predecessors which
20982098
// have loop carried dependencies (WAR) on the same VMEM_READs.
@@ -2113,7 +2113,7 @@ bool MFMASmallGemmSingleWaveOpt::applyIGLPStrategy(
21132113
break;
21142114

21152115
for (auto &Succ : Pred.getSUnit()->Succs) {
2116-
auto MI = Succ.getSUnit()->getInstr();
2116+
auto *MI = Succ.getSUnit()->getInstr();
21172117
if (!TII->isVMEM(*MI) || !MI->mayLoad())
21182118
continue;
21192119

0 commit comments

Comments
 (0)