Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[AMDGPU] Qualify auto. NFC. #110878

Merged
merged 2 commits into from
Oct 3, 2024
Merged

[AMDGPU] Qualify auto. NFC. #110878

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
4d8c2d2
[AMDGPU] Qualify auto. NFC.
jayfoad Oct 2, 2024
90a7a60
clang-format
jayfoad Oct 2, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions llvm/lib/Target/AMDGPU/AMDGPUAsanInstrumentation.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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:
Expand All @@ -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);
Expand Down
8 changes: 4 additions & 4 deletions llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
}
Expand Down Expand Up @@ -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);
Expand Down Expand Up @@ -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());
Expand Down
8 changes: 4 additions & 4 deletions llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down Expand Up @@ -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);
Expand Down Expand Up @@ -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);
Expand Down Expand Up @@ -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]);
Expand Down
34 changes: 17 additions & 17 deletions llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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();
}
Expand Down Expand Up @@ -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;

Expand All @@ -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;

Expand All @@ -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(),
Expand Down Expand Up @@ -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;

Expand Down Expand Up @@ -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();
}
Expand Down Expand Up @@ -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")
Expand Down Expand Up @@ -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());

Expand All @@ -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) {
Expand Down Expand Up @@ -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);
Expand All @@ -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")) {
Expand Down
42 changes: 21 additions & 21 deletions llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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())
Expand All @@ -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() &&
Expand Down Expand Up @@ -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());
Expand Down Expand Up @@ -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
Expand Down Expand Up @@ -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();
Expand Down Expand Up @@ -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;
Expand Down Expand Up @@ -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;
Expand Down Expand Up @@ -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))
Expand Down Expand Up @@ -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;
Expand Down Expand Up @@ -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())) {
Expand Down Expand Up @@ -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();
Expand Down Expand Up @@ -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;
Expand All @@ -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;

Expand Down Expand Up @@ -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())
Expand Down Expand Up @@ -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) {
Expand Down Expand Up @@ -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)) {
Expand All @@ -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.
Expand All @@ -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;

Expand Down
Loading
Loading