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.cpp87
-rw-r--r--llvm/lib/IR/Function.cpp1
-rw-r--r--llvm/lib/IR/Value.cpp9
-rw-r--r--llvm/lib/IR/Verifier.cpp63
6 files changed, 127 insertions, 85 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 8fb33c3..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;) {
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/Value.cpp b/llvm/lib/IR/Value.cpp
index 02c16e2..5928c89 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)
diff --git a/llvm/lib/IR/Verifier.cpp b/llvm/lib/IR/Verifier.cpp
index 9bd573e..3ff9895 100644
--- a/llvm/lib/IR/Verifier.cpp
+++ b/llvm/lib/IR/Verifier.cpp
@@ -2979,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...
@@ -6658,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);
@@ -6710,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-22 04:39:08 +0000