aboutsummaryrefslogtreecommitdiff
path: root/llvm/lib/IR
diff options
context:
space:
mode:
Diffstat (limited to 'llvm/lib/IR')
-rw-r--r--llvm/lib/IR/AsmWriter.cpp8
-rw-r--r--llvm/lib/IR/AutoUpgrade.cpp44
-rw-r--r--llvm/lib/IR/DebugInfo.cpp112
-rw-r--r--llvm/lib/IR/DebugInfoMetadata.cpp9
-rw-r--r--llvm/lib/IR/Function.cpp1
-rw-r--r--llvm/lib/IR/Metadata.cpp18
-rw-r--r--llvm/lib/IR/PassInstrumentation.cpp6
-rw-r--r--llvm/lib/IR/Value.cpp55
-rw-r--r--llvm/lib/IR/Verifier.cpp94
9 files changed, 216 insertions, 131 deletions
diff --git a/llvm/lib/IR/AsmWriter.cpp b/llvm/lib/IR/AsmWriter.cpp
index 145ef10..e5a4e1e 100644
--- a/llvm/lib/IR/AsmWriter.cpp
+++ b/llvm/lib/IR/AsmWriter.cpp
@@ -404,6 +404,9 @@ static void PrintCallingConv(unsigned cc, raw_ostream &Out) {
break;
case CallingConv::AMDGPU_KERNEL: Out << "amdgpu_kernel"; break;
case CallingConv::AMDGPU_Gfx: Out << "amdgpu_gfx"; break;
+ case CallingConv::AMDGPU_Gfx_WholeWave:
+ Out << "amdgpu_gfx_whole_wave";
+ break;
case CallingConv::M68k_RTD: Out << "m68k_rtdcc"; break;
case CallingConv::RISCV_VectorCall:
Out << "riscv_vector_cc";
@@ -2398,8 +2401,9 @@ static void writeDIFile(raw_ostream &Out, const DIFile *N, AsmWriterContext &) {
// Print all values for checksum together, or not at all.
if (N->getChecksum())
Printer.printChecksum(*N->getChecksum());
- Printer.printString("source", N->getSource().value_or(StringRef()),
- /* ShouldSkipEmpty */ true);
+ if (N->getSource())
+ Printer.printString("source", *N->getSource(),
+ /* ShouldSkipEmpty */ false);
Out << ")";
}
diff --git a/llvm/lib/IR/AutoUpgrade.cpp b/llvm/lib/IR/AutoUpgrade.cpp
index 86285a0..28ed1e5 100644
--- a/llvm/lib/IR/AutoUpgrade.cpp
+++ b/llvm/lib/IR/AutoUpgrade.cpp
@@ -1310,6 +1310,18 @@ static bool upgradeIntrinsicFunction1(Function *F, Function *&NewFn,
return true;
}
break;
+ case 'l':
+ if (Name.starts_with("lifetime.start") ||
+ Name.starts_with("lifetime.end")) {
+ // Unless remangling is required, do not upgrade the function declaration,
+ // but do upgrade the calls.
+ if (auto Result = llvm::Intrinsic::remangleIntrinsicFunction(F))
+ NewFn = *Result;
+ else
+ NewFn = F;
+ return true;
+ }
+ break;
case 'm': {
// Updating the memory intrinsics (memcpy/memmove/memset) that have an
// alignment parameter to embedding the alignment as an attribute of
@@ -1629,7 +1641,6 @@ bool llvm::UpgradeIntrinsicFunction(Function *F, Function *&NewFn,
NewFn = nullptr;
bool Upgraded =
upgradeIntrinsicFunction1(F, NewFn, CanUpgradeDebugIntrinsicsToRecords);
- assert(F != NewFn && "Intrinsic function upgraded to the same function");
// Upgrade intrinsic attributes. This does not change the function.
if (NewFn)
@@ -4570,6 +4581,9 @@ void llvm::UpgradeIntrinsicCall(CallBase *CI, Function *NewFn) {
}
const auto &DefaultCase = [&]() -> void {
+ if (F == NewFn)
+ return;
+
if (CI->getFunctionType() == NewFn->getFunctionType()) {
// Handle generic mangling change.
assert(
@@ -5109,6 +5123,31 @@ void llvm::UpgradeIntrinsicCall(CallBase *CI, Function *NewFn) {
MTI->setSourceAlignment(Align->getMaybeAlignValue());
break;
}
+
+ case Intrinsic::lifetime_start:
+ case Intrinsic::lifetime_end: {
+ Value *Size = CI->getArgOperand(0);
+ Value *Ptr = CI->getArgOperand(1);
+ if (isa<AllocaInst>(Ptr)) {
+ DefaultCase();
+ return;
+ }
+
+ // Try to strip pointer casts, such that the lifetime works on an alloca.
+ Ptr = Ptr->stripPointerCasts();
+ if (isa<AllocaInst>(Ptr)) {
+ // Don't use NewFn, as we might have looked through an addrspacecast.
+ if (NewFn->getIntrinsicID() == Intrinsic::lifetime_start)
+ NewCall = Builder.CreateLifetimeStart(Ptr, cast<ConstantInt>(Size));
+ else
+ NewCall = Builder.CreateLifetimeEnd(Ptr, cast<ConstantInt>(Size));
+ break;
+ }
+
+ // Otherwise remove the lifetime marker.
+ CI->eraseFromParent();
+ return;
+ }
}
assert(NewCall && "Should have either set this variable or returned through "
"the default case");
@@ -5131,7 +5170,8 @@ void llvm::UpgradeCallsToIntrinsic(Function *F) {
UpgradeIntrinsicCall(CB, NewFn);
// Remove old function, no longer used, from the module.
- F->eraseFromParent();
+ if (F != NewFn)
+ F->eraseFromParent();
}
}
diff --git a/llvm/lib/IR/DebugInfo.cpp b/llvm/lib/IR/DebugInfo.cpp
index 84a56058..ab8ecee 100644
--- a/llvm/lib/IR/DebugInfo.cpp
+++ b/llvm/lib/IR/DebugInfo.cpp
@@ -45,25 +45,6 @@ using namespace llvm;
using namespace llvm::at;
using namespace llvm::dwarf;
-TinyPtrVector<DbgDeclareInst *> llvm::findDbgDeclares(Value *V) {
- // This function is hot. Check whether the value has any metadata to avoid a
- // DenseMap lookup. This check is a bitfield datamember lookup.
- if (!V->isUsedByMetadata())
- return {};
- auto *L = ValueAsMetadata::getIfExists(V);
- if (!L)
- return {};
- auto *MDV = MetadataAsValue::getIfExists(V->getContext(), L);
- if (!MDV)
- return {};
-
- TinyPtrVector<DbgDeclareInst *> Declares;
- for (User *U : MDV->users())
- if (auto *DDI = dyn_cast<DbgDeclareInst>(U))
- Declares.push_back(DDI);
-
- return Declares;
-}
TinyPtrVector<DbgVariableRecord *> llvm::findDVRDeclares(Value *V) {
// This function is hot. Check whether the value has any metadata to avoid a
// DenseMap lookup. This check is a bitfield datamember lookup.
@@ -98,42 +79,31 @@ TinyPtrVector<DbgVariableRecord *> llvm::findDVRValues(Value *V) {
return Values;
}
-template <typename IntrinsicT, bool DbgAssignAndValuesOnly>
+template <bool DbgAssignAndValuesOnly>
static void
-findDbgIntrinsics(SmallVectorImpl<IntrinsicT *> &Result, Value *V,
- SmallVectorImpl<DbgVariableRecord *> *DbgVariableRecords) {
+findDbgIntrinsics(Value *V,
+ SmallVectorImpl<DbgVariableRecord *> &DbgVariableRecords) {
// This function is hot. Check whether the value has any metadata to avoid a
// DenseMap lookup.
if (!V->isUsedByMetadata())
return;
- LLVMContext &Ctx = V->getContext();
// TODO: If this value appears multiple times in a DIArgList, we should still
- // only add the owning DbgValueInst once; use this set to track ArgListUsers.
+ // only add the owning dbg.value once; use this set to track ArgListUsers.
// This behaviour can be removed when we can automatically remove duplicates.
// V will also appear twice in a dbg.assign if its used in the both the value
// and address components.
- SmallPtrSet<IntrinsicT *, 4> EncounteredIntrinsics;
SmallPtrSet<DbgVariableRecord *, 4> EncounteredDbgVariableRecords;
- /// Append IntrinsicT users of MetadataAsValue(MD).
- auto AppendUsers = [&Ctx, &EncounteredIntrinsics,
- &EncounteredDbgVariableRecords, &Result,
- DbgVariableRecords](Metadata *MD) {
- if (auto *MDV = MetadataAsValue::getIfExists(Ctx, MD)) {
- for (User *U : MDV->users())
- if (IntrinsicT *DVI = dyn_cast<IntrinsicT>(U))
- if (EncounteredIntrinsics.insert(DVI).second)
- Result.push_back(DVI);
- }
- if (!DbgVariableRecords)
- return;
+ /// Append users of MetadataAsValue(MD).
+ auto AppendUsers = [&EncounteredDbgVariableRecords,
+ &DbgVariableRecords](Metadata *MD) {
// Get DbgVariableRecords that use this as a single value.
if (LocalAsMetadata *L = dyn_cast<LocalAsMetadata>(MD)) {
for (DbgVariableRecord *DVR : L->getAllDbgVariableRecordUsers()) {
if (!DbgAssignAndValuesOnly || DVR->isDbgValue() || DVR->isDbgAssign())
if (EncounteredDbgVariableRecords.insert(DVR).second)
- DbgVariableRecords->push_back(DVR);
+ DbgVariableRecords.push_back(DVR);
}
}
};
@@ -142,29 +112,23 @@ findDbgIntrinsics(SmallVectorImpl<IntrinsicT *> &Result, Value *V,
AppendUsers(L);
for (Metadata *AL : L->getAllArgListUsers()) {
AppendUsers(AL);
- if (!DbgVariableRecords)
- continue;
DIArgList *DI = cast<DIArgList>(AL);
for (DbgVariableRecord *DVR : DI->getAllDbgVariableRecordUsers())
if (!DbgAssignAndValuesOnly || DVR->isDbgValue() || DVR->isDbgAssign())
if (EncounteredDbgVariableRecords.insert(DVR).second)
- DbgVariableRecords->push_back(DVR);
+ DbgVariableRecords.push_back(DVR);
}
}
}
void llvm::findDbgValues(
- SmallVectorImpl<DbgValueInst *> &DbgValues, Value *V,
- SmallVectorImpl<DbgVariableRecord *> *DbgVariableRecords) {
- findDbgIntrinsics<DbgValueInst, /*DbgAssignAndValuesOnly=*/true>(
- DbgValues, V, DbgVariableRecords);
+ Value *V, SmallVectorImpl<DbgVariableRecord *> &DbgVariableRecords) {
+ findDbgIntrinsics</*DbgAssignAndValuesOnly=*/true>(V, DbgVariableRecords);
}
void llvm::findDbgUsers(
- SmallVectorImpl<DbgVariableIntrinsic *> &DbgUsers, Value *V,
- SmallVectorImpl<DbgVariableRecord *> *DbgVariableRecords) {
- findDbgIntrinsics<DbgVariableIntrinsic, /*DbgAssignAndValuesOnly=*/false>(
- DbgUsers, V, DbgVariableRecords);
+ Value *V, SmallVectorImpl<DbgVariableRecord *> &DbgVariableRecords) {
+ findDbgIntrinsics</*DbgAssignAndValuesOnly=*/false>(V, DbgVariableRecords);
}
DISubprogram *llvm::getDISubprogram(const MDNode *Scope) {
@@ -173,18 +137,6 @@ DISubprogram *llvm::getDISubprogram(const MDNode *Scope) {
return nullptr;
}
-DebugLoc llvm::getDebugValueLoc(DbgVariableIntrinsic *DII) {
- // Original dbg.declare must have a location.
- const DebugLoc &DeclareLoc = DII->getDebugLoc();
- MDNode *Scope = DeclareLoc.getScope();
- DILocation *InlinedAt = DeclareLoc.getInlinedAt();
- // Because no machine insts can come from debug intrinsics, only the scope
- // and inlinedAt is significant. Zero line numbers are used in case this
- // DebugLoc leaks into any adjacent instructions. Produce an unknown location
- // with the correct scope / inlinedAt fields.
- return DILocation::get(DII->getContext(), 0, 0, Scope, InlinedAt);
-}
-
DebugLoc llvm::getDebugValueLoc(DbgVariableRecord *DVR) {
// Original dbg.declare must have a location.
const DebugLoc &DeclareLoc = DVR->getDebugLoc();
@@ -852,19 +804,6 @@ void DebugTypeInfoRemoval::traverse(MDNode *N) {
bool llvm::stripNonLineTableDebugInfo(Module &M) {
bool Changed = false;
- // First off, delete the debug intrinsics.
- auto RemoveUses = [&](StringRef Name) {
- if (auto *DbgVal = M.getFunction(Name)) {
- while (!DbgVal->use_empty())
- cast<Instruction>(DbgVal->user_back())->eraseFromParent();
- DbgVal->eraseFromParent();
- Changed = true;
- }
- };
- RemoveUses("llvm.dbg.declare");
- RemoveUses("llvm.dbg.label");
- RemoveUses("llvm.dbg.value");
-
// Delete non-CU debug info named metadata nodes.
for (auto NMI = M.named_metadata_begin(), NME = M.named_metadata_end();
NMI != NME;) {
@@ -2288,39 +2227,36 @@ bool AssignmentTrackingPass::runOnFunction(Function &F) {
// Collect a map of {backing storage : dbg.declares} (currently "backing
// storage" is limited to Allocas). We'll use this to find dbg.declares to
// delete after running `trackAssignments`.
- DenseMap<const AllocaInst *, SmallPtrSet<DbgDeclareInst *, 2>> DbgDeclares;
DenseMap<const AllocaInst *, SmallPtrSet<DbgVariableRecord *, 2>> DVRDeclares;
// Create another similar map of {storage : variables} that we'll pass to
// trackAssignments.
StorageToVarsMap Vars;
- auto ProcessDeclare = [&](auto *Declare, auto &DeclareList) {
+ auto ProcessDeclare = [&](DbgVariableRecord &Declare) {
// FIXME: trackAssignments doesn't let you specify any modifiers to the
// variable (e.g. fragment) or location (e.g. offset), so we have to
// leave dbg.declares with non-empty expressions in place.
- if (Declare->getExpression()->getNumElements() != 0)
+ if (Declare.getExpression()->getNumElements() != 0)
return;
- if (!Declare->getAddress())
+ if (!Declare.getAddress())
return;
if (AllocaInst *Alloca =
- dyn_cast<AllocaInst>(Declare->getAddress()->stripPointerCasts())) {
+ dyn_cast<AllocaInst>(Declare.getAddress()->stripPointerCasts())) {
// FIXME: Skip VLAs for now (let these variables use dbg.declares).
if (!Alloca->isStaticAlloca())
return;
// Similarly, skip scalable vectors (use dbg.declares instead).
if (auto Sz = Alloca->getAllocationSize(*DL); Sz && Sz->isScalable())
return;
- DeclareList[Alloca].insert(Declare);
- Vars[Alloca].insert(VarRecord(Declare));
+ DVRDeclares[Alloca].insert(&Declare);
+ Vars[Alloca].insert(VarRecord(&Declare));
}
};
for (auto &BB : F) {
for (auto &I : BB) {
for (DbgVariableRecord &DVR : filterDbgVars(I.getDbgRecordRange())) {
if (DVR.isDbgDeclare())
- ProcessDeclare(&DVR, DVRDeclares);
+ ProcessDeclare(DVR);
}
- if (DbgDeclareInst *DDI = dyn_cast<DbgDeclareInst>(&I))
- ProcessDeclare(DDI, DbgDeclares);
}
}
@@ -2336,8 +2272,8 @@ bool AssignmentTrackingPass::runOnFunction(Function &F) {
trackAssignments(F.begin(), F.end(), Vars, *DL);
// Delete dbg.declares for variables now tracked with assignment tracking.
- auto DeleteSubsumedDeclare = [&](const auto &Markers, auto &Declares) {
- (void)Markers;
+ for (auto &[Insts, Declares] : DVRDeclares) {
+ auto Markers = at::getDVRAssignmentMarkers(Insts);
for (auto *Declare : Declares) {
// Assert that the alloca that Declare uses is now linked to a dbg.assign
// describing the same variable (i.e. check that this dbg.declare has
@@ -2356,10 +2292,6 @@ bool AssignmentTrackingPass::runOnFunction(Function &F) {
Changed = true;
}
};
- for (auto &P : DbgDeclares)
- DeleteSubsumedDeclare(at::getAssignmentMarkers(P.first), P.second);
- for (auto &P : DVRDeclares)
- DeleteSubsumedDeclare(at::getDVRAssignmentMarkers(P.first), P.second);
return Changed;
}
diff --git a/llvm/lib/IR/DebugInfoMetadata.cpp b/llvm/lib/IR/DebugInfoMetadata.cpp
index 2270923..f16963d 100644
--- a/llvm/lib/IR/DebugInfoMetadata.cpp
+++ b/llvm/lib/IR/DebugInfoMetadata.cpp
@@ -49,20 +49,11 @@ uint32_t DIType::getAlignInBits() const {
const DIExpression::FragmentInfo DebugVariable::DefaultFragment = {
std::numeric_limits<uint64_t>::max(), std::numeric_limits<uint64_t>::min()};
-DebugVariable::DebugVariable(const DbgVariableIntrinsic *DII)
- : Variable(DII->getVariable()),
- Fragment(DII->getExpression()->getFragmentInfo()),
- InlinedAt(DII->getDebugLoc().getInlinedAt()) {}
-
DebugVariable::DebugVariable(const DbgVariableRecord *DVR)
: Variable(DVR->getVariable()),
Fragment(DVR->getExpression()->getFragmentInfo()),
InlinedAt(DVR->getDebugLoc().getInlinedAt()) {}
-DebugVariableAggregate::DebugVariableAggregate(const DbgVariableIntrinsic *DVI)
- : DebugVariable(DVI->getVariable(), std::nullopt,
- DVI->getDebugLoc()->getInlinedAt()) {}
-
DILocation::DILocation(LLVMContext &C, StorageType Storage, unsigned Line,
unsigned Column, uint64_t AtomGroup, uint8_t AtomRank,
ArrayRef<Metadata *> MDs, bool ImplicitCode)
diff --git a/llvm/lib/IR/Function.cpp b/llvm/lib/IR/Function.cpp
index 7a03663..fc06745 100644
--- a/llvm/lib/IR/Function.cpp
+++ b/llvm/lib/IR/Function.cpp
@@ -1232,6 +1232,7 @@ bool llvm::CallingConv::supportsNonVoidReturnType(CallingConv::ID CC) {
case CallingConv::AArch64_SVE_VectorCall:
case CallingConv::WASM_EmscriptenInvoke:
case CallingConv::AMDGPU_Gfx:
+ case CallingConv::AMDGPU_Gfx_WholeWave:
case CallingConv::M68k_INTR:
case CallingConv::AArch64_SME_ABI_Support_Routines_PreserveMost_From_X0:
case CallingConv::AArch64_SME_ABI_Support_Routines_PreserveMost_From_X2:
diff --git a/llvm/lib/IR/Metadata.cpp b/llvm/lib/IR/Metadata.cpp
index f0448b0..0dbd07f 100644
--- a/llvm/lib/IR/Metadata.cpp
+++ b/llvm/lib/IR/Metadata.cpp
@@ -1303,6 +1303,24 @@ static void addRange(SmallVectorImpl<ConstantInt *> &EndPoints,
EndPoints.push_back(High);
}
+MDNode *MDNode::getMergedCalleeTypeMetadata(const MDNode *A, const MDNode *B) {
+ // Drop the callee_type metadata if either of the call instructions do not
+ // have it.
+ if (!A || !B)
+ return nullptr;
+ SmallVector<Metadata *, 8> AB;
+ SmallPtrSet<Metadata *, 8> MergedCallees;
+ auto AddUniqueCallees = [&AB, &MergedCallees](const MDNode *N) {
+ for (Metadata *MD : N->operands()) {
+ if (MergedCallees.insert(MD).second)
+ AB.push_back(MD);
+ }
+ };
+ AddUniqueCallees(A);
+ AddUniqueCallees(B);
+ return MDNode::get(A->getContext(), AB);
+}
+
MDNode *MDNode::getMostGenericRange(MDNode *A, MDNode *B) {
// Given two ranges, we want to compute the union of the ranges. This
// is slightly complicated by having to combine the intervals and merge
diff --git a/llvm/lib/IR/PassInstrumentation.cpp b/llvm/lib/IR/PassInstrumentation.cpp
index 94ad124..70bbe8f 100644
--- a/llvm/lib/IR/PassInstrumentation.cpp
+++ b/llvm/lib/IR/PassInstrumentation.cpp
@@ -23,6 +23,7 @@ template struct LLVM_EXPORT_TEMPLATE Any::TypeId<const Loop *>;
void PassInstrumentationCallbacks::addClassToPassName(StringRef ClassName,
StringRef PassName) {
+ assert(!PassName.empty() && "PassName can't be empty!");
ClassToPassName.try_emplace(ClassName, PassName.str());
}
@@ -33,7 +34,10 @@ PassInstrumentationCallbacks::getPassNameForClassName(StringRef ClassName) {
Fn();
ClassToPassNameCallbacks.clear();
}
- return ClassToPassName[ClassName];
+ auto PassNameIter = ClassToPassName.find(ClassName);
+ if (PassNameIter != ClassToPassName.end())
+ return PassNameIter->second;
+ return {};
}
AnalysisKey PassInstrumentationAnalysis::Key;
diff --git a/llvm/lib/IR/Value.cpp b/llvm/lib/IR/Value.cpp
index 02c16e2..129ca4a 100644
--- a/llvm/lib/IR/Value.cpp
+++ b/llvm/lib/IR/Value.cpp
@@ -582,16 +582,11 @@ void Value::replaceUsesWithIf(Value *New,
}
}
-/// Replace llvm.dbg.* uses of MetadataAsValue(ValueAsMetadata(V)) outside BB
+/// Replace debug record uses of MetadataAsValue(ValueAsMetadata(V)) outside BB
/// with New.
static void replaceDbgUsesOutsideBlock(Value *V, Value *New, BasicBlock *BB) {
- SmallVector<DbgVariableIntrinsic *> DbgUsers;
SmallVector<DbgVariableRecord *> DPUsers;
- findDbgUsers(DbgUsers, V, &DPUsers);
- for (auto *DVI : DbgUsers) {
- if (DVI->getParent() != BB)
- DVI->replaceVariableLocationOp(V, New);
- }
+ findDbgUsers(V, DPUsers);
for (auto *DVR : DPUsers) {
DbgMarker *Marker = DVR->getMarker();
if (Marker->getParent() != BB)
@@ -752,28 +747,34 @@ const Value *Value::stripAndAccumulateConstantOffsets(
// means when we construct GEPOffset, we need to use the size
// of GEP's pointer type rather than the size of the original
// pointer type.
- APInt GEPOffset(DL.getIndexTypeSizeInBits(V->getType()), 0);
- if (!GEP->accumulateConstantOffset(DL, GEPOffset, ExternalAnalysis))
- return V;
-
- // Stop traversal if the pointer offset wouldn't fit in the bit-width
- // provided by the Offset argument. This can happen due to AddrSpaceCast
- // stripping.
- if (GEPOffset.getSignificantBits() > BitWidth)
- return V;
-
- // External Analysis can return a result higher/lower than the value
- // represents. We need to detect overflow/underflow.
- APInt GEPOffsetST = GEPOffset.sextOrTrunc(BitWidth);
- if (!ExternalAnalysis) {
- Offset += GEPOffsetST;
+ unsigned CurBitWidth = DL.getIndexTypeSizeInBits(V->getType());
+ if (CurBitWidth == BitWidth) {
+ if (!GEP->accumulateConstantOffset(DL, Offset, ExternalAnalysis))
+ return V;
} else {
- bool Overflow = false;
- APInt OldOffset = Offset;
- Offset = Offset.sadd_ov(GEPOffsetST, Overflow);
- if (Overflow) {
- Offset = OldOffset;
+ APInt GEPOffset(CurBitWidth, 0);
+ if (!GEP->accumulateConstantOffset(DL, GEPOffset, ExternalAnalysis))
+ return V;
+
+ // Stop traversal if the pointer offset wouldn't fit in the bit-width
+ // provided by the Offset argument. This can happen due to AddrSpaceCast
+ // stripping.
+ if (GEPOffset.getSignificantBits() > BitWidth)
return V;
+
+ // External Analysis can return a result higher/lower than the value
+ // represents. We need to detect overflow/underflow.
+ APInt GEPOffsetST = GEPOffset.sextOrTrunc(BitWidth);
+ if (!ExternalAnalysis) {
+ Offset += GEPOffsetST;
+ } else {
+ bool Overflow = false;
+ APInt OldOffset = Offset;
+ Offset = Offset.sadd_ov(GEPOffsetST, Overflow);
+ if (Overflow) {
+ Offset = OldOffset;
+ return V;
+ }
}
}
V = GEP->getPointerOperand();
diff --git a/llvm/lib/IR/Verifier.cpp b/llvm/lib/IR/Verifier.cpp
index 8c8ed3c..3ff9895 100644
--- a/llvm/lib/IR/Verifier.cpp
+++ b/llvm/lib/IR/Verifier.cpp
@@ -531,6 +531,7 @@ private:
void visitCallStackMetadata(MDNode *MD);
void visitMemProfMetadata(Instruction &I, MDNode *MD);
void visitCallsiteMetadata(Instruction &I, MDNode *MD);
+ void visitCalleeTypeMetadata(Instruction &I, MDNode *MD);
void visitDIAssignIDMetadata(Instruction &I, MDNode *MD);
void visitMMRAMetadata(Instruction &I, MDNode *MD);
void visitAnnotationMetadata(MDNode *Annotation);
@@ -2978,6 +2979,16 @@ void Verifier::visitFunction(const Function &F) {
"perfect forwarding!",
&F);
break;
+ case CallingConv::AMDGPU_Gfx_WholeWave:
+ Check(!F.arg_empty() && F.arg_begin()->getType()->isIntegerTy(1),
+ "Calling convention requires first argument to be i1", &F);
+ Check(!F.arg_begin()->hasInRegAttr(),
+ "Calling convention requires first argument to not be inreg", &F);
+ Check(!F.isVarArg(),
+ "Calling convention does not support varargs or "
+ "perfect forwarding!",
+ &F);
+ break;
}
// Check that the argument values match the function type for this function...
@@ -5193,6 +5204,33 @@ void Verifier::visitCallsiteMetadata(Instruction &I, MDNode *MD) {
visitCallStackMetadata(MD);
}
+static inline bool isConstantIntMetadataOperand(const Metadata *MD) {
+ if (auto *VAL = dyn_cast<ValueAsMetadata>(MD))
+ return isa<ConstantInt>(VAL->getValue());
+ return false;
+}
+
+void Verifier::visitCalleeTypeMetadata(Instruction &I, MDNode *MD) {
+ Check(isa<CallBase>(I), "!callee_type metadata should only exist on calls",
+ &I);
+ for (Metadata *Op : MD->operands()) {
+ Check(isa<MDNode>(Op),
+ "The callee_type metadata must be a list of type metadata nodes", Op);
+ auto *TypeMD = cast<MDNode>(Op);
+ Check(TypeMD->getNumOperands() == 2,
+ "Well-formed generalized type metadata must contain exactly two "
+ "operands",
+ Op);
+ Check(isConstantIntMetadataOperand(TypeMD->getOperand(0)) &&
+ mdconst::extract<ConstantInt>(TypeMD->getOperand(0))->isZero(),
+ "The first operand of type metadata for functions must be zero", Op);
+ Check(TypeMD->hasGeneralizedMDString(),
+ "Only generalized type metadata can be part of the callee_type "
+ "metadata list",
+ Op);
+ }
+}
+
void Verifier::visitAnnotationMetadata(MDNode *Annotation) {
Check(isa<MDTuple>(Annotation), "annotation must be a tuple");
Check(Annotation->getNumOperands() >= 1,
@@ -5470,6 +5508,9 @@ void Verifier::visitInstruction(Instruction &I) {
if (MDNode *MD = I.getMetadata(LLVMContext::MD_callsite))
visitCallsiteMetadata(I, MD);
+ if (MDNode *MD = I.getMetadata(LLVMContext::MD_callee_type))
+ visitCalleeTypeMetadata(I, MD);
+
if (MDNode *MD = I.getMetadata(LLVMContext::MD_DIAssignID))
visitDIAssignIDMetadata(I, MD);
@@ -6627,6 +6668,54 @@ void Verifier::visitIntrinsicCall(Intrinsic::ID ID, CallBase &Call) {
"invalid vector type for format", &Call, Src1, Call.getArgOperand(5));
break;
}
+ case Intrinsic::amdgcn_wmma_f32_16x16x128_f8f6f4: {
+ Value *Src0 = Call.getArgOperand(1);
+ Value *Src1 = Call.getArgOperand(3);
+
+ unsigned FmtA = cast<ConstantInt>(Call.getArgOperand(0))->getZExtValue();
+ unsigned FmtB = cast<ConstantInt>(Call.getArgOperand(2))->getZExtValue();
+ Check(FmtA <= 4, "invalid value for matrix format", Call,
+ Call.getArgOperand(0));
+ Check(FmtB <= 4, "invalid value for matrix format", Call,
+ Call.getArgOperand(2));
+
+ // AMDGPU::MatrixFMT values
+ auto getFormatNumRegs = [](unsigned FormatVal) {
+ switch (FormatVal) {
+ case 0:
+ case 1:
+ return 16u;
+ case 2:
+ case 3:
+ return 12u;
+ case 4:
+ return 8u;
+ default:
+ llvm_unreachable("invalid format value");
+ }
+ };
+
+ auto isValidSrcASrcBVector = [](FixedVectorType *Ty) {
+ if (!Ty || !Ty->getElementType()->isIntegerTy(32))
+ return false;
+ unsigned NumElts = Ty->getNumElements();
+ return NumElts == 16 || NumElts == 12 || NumElts == 8;
+ };
+
+ auto *Src0Ty = dyn_cast<FixedVectorType>(Src0->getType());
+ auto *Src1Ty = dyn_cast<FixedVectorType>(Src1->getType());
+ Check(isValidSrcASrcBVector(Src0Ty),
+ "operand 1 must be 8, 12 or 16 element i32 vector", &Call, Src0);
+ Check(isValidSrcASrcBVector(Src1Ty),
+ "operand 3 must be 8, 12 or 16 element i32 vector", &Call, Src1);
+
+ // Permit excess registers for the format.
+ Check(Src0Ty->getNumElements() >= getFormatNumRegs(FmtA),
+ "invalid vector type for format", &Call, Src0, Call.getArgOperand(0));
+ Check(Src1Ty->getNumElements() >= getFormatNumRegs(FmtB),
+ "invalid vector type for format", &Call, Src1, Call.getArgOperand(2));
+ break;
+ }
case Intrinsic::nvvm_setmaxnreg_inc_sync_aligned_u32:
case Intrinsic::nvvm_setmaxnreg_dec_sync_aligned_u32: {
Value *V = Call.getArgOperand(0);
@@ -6679,6 +6768,11 @@ void Verifier::visitIntrinsicCall(Intrinsic::ID ID, CallBase &Call) {
"llvm.threadlocal.address operand isThreadLocal() must be true");
break;
}
+ case Intrinsic::lifetime_start:
+ case Intrinsic::lifetime_end:
+ Check(isa<AllocaInst>(Call.getArgOperand(1)),
+ "llvm.lifetime.start/end can only be used on alloca", &Call);
+ break;
};
// Verify that there aren't any unmediated control transfers between funclets.
generated by cgit v1.1 at 2025-07-23 14:18:35 +0000