[llvm] [RFC][NFCI][IR] Extract AMDGPU-specific verification logic into `VerifierAMDGPU.cpp` (PR #204284)
Shilei Tian via llvm-commits
llvm-commits at lists.llvm.org
Mon Jun 22 08:39:14 PDT 2026
https://github.com/shiltian updated https://github.com/llvm/llvm-project/pull/204284
>From 93d04aba867462ad013d21d6788285d63584785d Mon Sep 17 00:00:00 2001
From: Shilei Tian <i at tianshilei.me>
Date: Wed, 17 Jun 2026 00:55:20 -0400
Subject: [PATCH] [RFC][IR] Extract AMDGPU-specific verification logic into
`VerifierAMDGPU.cpp`
`Verifier.cpp` is large and already mixes generic IR verification with
target-specific checks. We also have a growing amount of AMDGPU verifier logic
downstream, which would all end up in the same file if we don't address this,
and that is not ideal.
This patch extracts AMDGPU-specific verification logic into a separate
`VerifierAMDGPU.cpp` file, with shared infrastructure (`VerifierSupport`) moved
into `VerifierInternal.h`.
This is purely a code organization change, not a target-dependent IR verifier.
All checks remain compiled and linked into `LLVMCore` regardless of the target
triple. The extracted functions are called unconditionally at well-defined
extension points in `Verifier.cpp`, and each function internally gates on
target-specific conditions (for example, triple checks or intrinsic IDs) as
needed. The file is strictly limited to AMDGPU-specific IR constructs (amdgcn
intrinsics, AMDGPU module flags, etc.), and does not contain generic IR rules
that vary by target.
This PR deliberately avoided introducing polymorphism, since this is not a
target-dependent verification framework. Instead, this follows a pattern similar
to TargetParser (for example `AMDGPUTargetParser.cpp`): flat file layout, free
functions, no registration, and unconditional compilation.
Other targets that want similar separation can follow the same pattern.
Open to suggestions on whether there is a better long-term way to structure
this.
---
llvm/lib/IR/CMakeLists.txt | 1 +
llvm/lib/IR/Verifier.cpp | 555 +------------------
llvm/lib/IR/VerifierAMDGPU.cpp | 406 ++++++++++++++
llvm/lib/IR/VerifierInternal.h | 233 ++++++++
llvm/test/Verifier/callbr-intrinsic.ll | 12 +-
llvm/utils/gn/secondary/llvm/lib/IR/BUILD.gn | 1 +
6 files changed, 671 insertions(+), 537 deletions(-)
create mode 100644 llvm/lib/IR/VerifierAMDGPU.cpp
create mode 100644 llvm/lib/IR/VerifierInternal.h
diff --git a/llvm/lib/IR/CMakeLists.txt b/llvm/lib/IR/CMakeLists.txt
index 9cc45ef0e1773..3037f01083308 100644
--- a/llvm/lib/IR/CMakeLists.txt
+++ b/llvm/lib/IR/CMakeLists.txt
@@ -79,6 +79,7 @@ add_llvm_component_library(LLVMCore
ValueSymbolTable.cpp
VectorTypeUtils.cpp
Verifier.cpp
+ VerifierAMDGPU.cpp
VFABIDemangler.cpp
RuntimeLibcalls.cpp
diff --git a/llvm/lib/IR/Verifier.cpp b/llvm/lib/IR/Verifier.cpp
index 8207e60857eba..2a0892d1af11a 100644
--- a/llvm/lib/IR/Verifier.cpp
+++ b/llvm/lib/IR/Verifier.cpp
@@ -48,6 +48,7 @@
//===----------------------------------------------------------------------===//
#include "llvm/IR/Verifier.h"
+#include "VerifierInternal.h"
#include "llvm/ADT/APFloat.h"
#include "llvm/ADT/APInt.h"
#include "llvm/ADT/ArrayRef.h"
@@ -95,7 +96,6 @@
#include "llvm/IR/IntrinsicInst.h"
#include "llvm/IR/Intrinsics.h"
#include "llvm/IR/IntrinsicsAArch64.h"
-#include "llvm/IR/IntrinsicsAMDGPU.h"
#include "llvm/IR/IntrinsicsARM.h"
#include "llvm/IR/IntrinsicsNVPTX.h"
#include "llvm/IR/IntrinsicsWebAssembly.h"
@@ -115,7 +115,6 @@
#include "llvm/InitializePasses.h"
#include "llvm/Pass.h"
#include "llvm/ProfileData/InstrProf.h"
-#include "llvm/Support/AMDGPUAddrSpace.h"
#include "llvm/Support/AtomicOrdering.h"
#include "llvm/Support/Casting.h"
#include "llvm/Support/CommandLine.h"
@@ -143,189 +142,6 @@ static cl::opt<bool> VerifyNoAliasScopeDomination(
cl::desc("Ensure that llvm.experimental.noalias.scope.decl for identical "
"scopes are not dominating"));
-struct llvm::VerifierSupport {
- raw_ostream *OS;
- const Module &M;
- ModuleSlotTracker MST;
- const Triple &TT;
- const DataLayout &DL;
- LLVMContext &Context;
-
- /// Track the brokenness of the module while recursively visiting.
- bool Broken = false;
- /// Broken debug info can be "recovered" from by stripping the debug info.
- bool BrokenDebugInfo = false;
- /// Whether to treat broken debug info as an error.
- bool TreatBrokenDebugInfoAsError = true;
-
- explicit VerifierSupport(raw_ostream *OS, const Module &M)
- : OS(OS), M(M), MST(&M), TT(M.getTargetTriple()), DL(M.getDataLayout()),
- Context(M.getContext()) {}
-
-private:
- void Write(const Module *M) {
- *OS << "; ModuleID = '" << M->getModuleIdentifier() << "'\n";
- }
-
- void Write(const Value *V) {
- if (V)
- Write(*V);
- }
-
- void Write(const Value &V) {
- if (isa<Instruction>(V)) {
- V.print(*OS, MST);
- *OS << '\n';
- } else {
- V.printAsOperand(*OS, true, MST);
- *OS << '\n';
- }
- }
-
- void Write(const DbgRecord *DR) {
- if (DR) {
- DR->print(*OS, MST, false);
- *OS << '\n';
- }
- }
-
- void Write(DbgVariableRecord::LocationType Type) {
- switch (Type) {
- case DbgVariableRecord::LocationType::Value:
- *OS << "value";
- break;
- case DbgVariableRecord::LocationType::Declare:
- *OS << "declare";
- break;
- case DbgVariableRecord::LocationType::DeclareValue:
- *OS << "declare_value";
- break;
- case DbgVariableRecord::LocationType::Assign:
- *OS << "assign";
- break;
- case DbgVariableRecord::LocationType::End:
- *OS << "end";
- break;
- case DbgVariableRecord::LocationType::Any:
- *OS << "any";
- break;
- };
- }
-
- void Write(const Metadata *MD) {
- if (!MD)
- return;
- MD->print(*OS, MST, &M);
- *OS << '\n';
- }
-
- template <class T> void Write(const MDTupleTypedArrayWrapper<T> &MD) {
- Write(MD.get());
- }
-
- void Write(const NamedMDNode *NMD) {
- if (!NMD)
- return;
- NMD->print(*OS, MST);
- *OS << '\n';
- }
-
- void Write(Type *T) {
- if (!T)
- return;
- *OS << ' ' << *T;
- }
-
- void Write(const Comdat *C) {
- if (!C)
- return;
- *OS << *C;
- }
-
- void Write(const APInt *AI) {
- if (!AI)
- return;
- *OS << *AI << '\n';
- }
-
- void Write(const unsigned i) { *OS << i << '\n'; }
-
- // NOLINTNEXTLINE(readability-identifier-naming)
- void Write(const Attribute *A) {
- if (!A)
- return;
- *OS << A->getAsString() << '\n';
- }
-
- // NOLINTNEXTLINE(readability-identifier-naming)
- void Write(const AttributeSet *AS) {
- if (!AS)
- return;
- *OS << AS->getAsString() << '\n';
- }
-
- // NOLINTNEXTLINE(readability-identifier-naming)
- void Write(const AttributeList *AL) {
- if (!AL)
- return;
- AL->print(*OS);
- }
-
- void Write(Printable P) { *OS << P << '\n'; }
-
- template <typename T> void Write(ArrayRef<T> Vs) {
- for (const T &V : Vs)
- Write(V);
- }
-
- template <typename T1, typename... Ts>
- void WriteTs(const T1 &V1, const Ts &... Vs) {
- Write(V1);
- WriteTs(Vs...);
- }
-
- template <typename... Ts> void WriteTs() {}
-
-public:
- /// A check failed, so printout out the condition and the message.
- ///
- /// This provides a nice place to put a breakpoint if you want to see why
- /// something is not correct.
- void CheckFailed(const Twine &Message) {
- if (OS)
- *OS << Message << '\n';
- Broken = true;
- }
-
- /// A check failed (with values to print).
- ///
- /// This calls the Message-only version so that the above is easier to set a
- /// breakpoint on.
- template <typename T1, typename... Ts>
- void CheckFailed(const Twine &Message, const T1 &V1, const Ts &... Vs) {
- CheckFailed(Message);
- if (OS)
- WriteTs(V1, Vs...);
- }
-
- /// A debug info check failed.
- void DebugInfoCheckFailed(const Twine &Message) {
- if (OS)
- *OS << Message << '\n';
- Broken |= TreatBrokenDebugInfoAsError;
- BrokenDebugInfo = true;
- }
-
- /// A debug info check failed (with values to print).
- template <typename T1, typename... Ts>
- void DebugInfoCheckFailed(const Twine &Message, const T1 &V1,
- const Ts &... Vs) {
- DebugInfoCheckFailed(Message);
- if (OS)
- WriteTs(V1, Vs...);
- }
-};
-
namespace {
class Verifier : public InstVisitor<Verifier>, VerifierSupport {
@@ -647,7 +463,6 @@ class Verifier : public InstVisitor<Verifier>, VerifierSupport {
void verifyFunctionAttrs(FunctionType *FT, AttributeList Attrs,
const Value *V, bool IsIntrinsic, bool IsInlineAsm);
void verifyFunctionMetadata(ArrayRef<std::pair<unsigned, MDNode *>> MDs);
- void verifyAMDGPUReqdWorkGroupSize(const Function &F);
void verifyUnknownProfileMetadata(MDNode *MD);
void visitConstantExprsRecursively(const Constant *EntryC);
void visitConstantExpr(const ConstantExpr *CE);
@@ -2147,25 +1962,13 @@ Verifier::visitModuleFlag(const MDNode *Op,
"SemanticInterposition metadata requires constant integer argument");
}
- if (ID->getString() == "amdgpu.buffer.oob.mode" ||
- ID->getString() == "amdgpu.tbuffer.oob.mode") {
- Check(MFB == Module::Max,
- "'" + ID->getString() +
- "' module flag must use 'max' merge behaviour");
- ConstantInt *Value =
- mdconst::dyn_extract_or_null<ConstantInt>(Op->getOperand(2));
- Check(Value, "'" + ID->getString() +
- "' module flag must have a constant integer value");
- if (Value) {
- Check(Value->getZExtValue() <= 2,
- "'" + ID->getString() + "' module flag must be 0, 1, or 2");
- }
- }
-
if (ID->getString() == "CG Profile") {
for (const MDOperand &MDO : cast<MDNode>(Op->getOperand(2))->operands())
visitModuleFlagCGProfileEntry(MDO);
}
+
+ // Target-specific module flag checks.
+ verifyAMDGPUModuleFlag(*this, ID, MFB, Op);
}
void Verifier::visitModuleFlagCGProfileEntry(const MDOperand &MDO) {
@@ -2858,57 +2661,6 @@ void Verifier::verifyFunctionMetadata(
}
}
-void Verifier::verifyAMDGPUReqdWorkGroupSize(const Function &F) {
- if (!TT.isAMDGPU())
- return;
-
- MDNode *ReqdWorkGroupSize = F.getMetadata("reqd_work_group_size");
- if (!ReqdWorkGroupSize || ReqdWorkGroupSize->getNumOperands() != 3)
- return;
-
- uint64_t Product = 1;
- for (const MDOperand &Op : ReqdWorkGroupSize->operands()) {
- ConstantInt *C = mdconst::dyn_extract<ConstantInt>(Op);
- if (!C || C->getValue().getActiveBits() > 64)
- return;
- uint64_t Dim = C->getZExtValue();
- if (Dim != 0 && Product > std::numeric_limits<uint64_t>::max() / Dim)
- return;
- Product *= Dim;
- }
-
- Attribute FlatWorkGroupSize = F.getFnAttribute("amdgpu-flat-work-group-size");
- if (!FlatWorkGroupSize.isValid()) {
- CheckFailed("reqd_work_group_size requires amdgpu-flat-work-group-size", &F,
- ReqdWorkGroupSize);
- return;
- }
-
- if (!FlatWorkGroupSize.isStringAttribute()) {
- CheckFailed("amdgpu-flat-work-group-size must be a string attribute", &F);
- return;
- }
-
- StringRef AttrValue = FlatWorkGroupSize.getValueAsString();
- std::pair<StringRef, StringRef> Values = AttrValue.split(',');
- uint64_t Min = 0;
- uint64_t Max = 0;
- bool Parsed = !Values.second.contains(',') &&
- llvm::to_integer(Values.first.trim(), Min) &&
- llvm::to_integer(Values.second.trim(), Max);
- if (!Parsed) {
- CheckFailed("amdgpu-flat-work-group-size must be a pair of unsigned "
- "integers",
- &F);
- return;
- }
-
- Check(Min == Product && Max == Product,
- "amdgpu-flat-work-group-size must equal the product of "
- "reqd_work_group_size operands",
- &F, ReqdWorkGroupSize);
-}
-
void Verifier::visitConstantExprsRecursively(const Constant *EntryC) {
if (EntryC->getNumOperands() == 0)
return;
@@ -3379,7 +3131,9 @@ void Verifier::visitFunction(const Function &F) {
F.getAllMetadata(MDs);
assert(F.hasMetadata() != MDs.empty() && "Bit out-of-sync");
verifyFunctionMetadata(MDs);
- verifyAMDGPUReqdWorkGroupSize(F);
+
+ // Target-specific function metadata checks.
+ verifyAMDGPUFunctionMetadata(*this, F);
// Check validity of the personality function
if (F.hasPersonalityFn()) {
@@ -3678,27 +3432,22 @@ void Verifier::visitIndirectBrInst(IndirectBrInst &BI) {
visitTerminator(BI);
}
+static bool isSupportedCallBrIntrinsic(Intrinsic::ID ID) {
+ // Currently we only support callbr for amdgcn.kill. Add more checks here as
+ // needed.
+ return isAMDGPUCallBrIntrinsic(ID);
+}
+
void Verifier::visitCallBrInst(CallBrInst &CBI) {
if (!CBI.isInlineAsm()) {
Check(CBI.getCalledFunction(),
- "Callbr: indirect function / invalid signature");
+ "callbr: indirect function / invalid signature");
Check(!CBI.hasOperandBundles(),
- "Callbr for intrinsics currently doesn't support operand bundles");
-
- switch (CBI.getIntrinsicID()) {
- case Intrinsic::amdgcn_kill: {
- Check(CBI.getNumIndirectDests() == 1,
- "Callbr amdgcn_kill only supports one indirect dest");
- bool Unreachable = isa<UnreachableInst>(CBI.getIndirectDest(0)->begin());
- CallInst *Call = dyn_cast<CallInst>(CBI.getIndirectDest(0)->begin());
- Check(Unreachable || (Call && Call->getIntrinsicID() ==
- Intrinsic::amdgcn_unreachable),
- "Callbr amdgcn_kill indirect dest needs to be unreachable");
- break;
- }
- default:
+ "callbr for intrinsics currently doesn't support operand bundles");
+
+ if (!isSupportedCallBrIntrinsic(CBI.getIntrinsicID())) {
CheckFailed(
- "Callbr currently only supports asm-goto and selected intrinsics");
+ "callbr currently only supports asm-goto and selected intrinsics");
}
visitIntrinsicCall(CBI.getIntrinsicID(), CBI);
} else {
@@ -4886,12 +4635,10 @@ void Verifier::visitAllocaInst(AllocaInst &AI) {
verifySwiftErrorValue(&AI);
}
- if (TT.isAMDGPU()) {
- Check(AI.getAddressSpace() == AMDGPUAS::PRIVATE_ADDRESS,
- "alloca on amdgpu must be in addrspace(5)", &AI);
- }
-
visitInstruction(AI);
+
+ // Target-specific alloca checks.
+ verifyAMDGPUAlloca(*this, AI);
}
void Verifier::visitAtomicCmpXchgInst(AtomicCmpXchgInst &CXI) {
@@ -7163,263 +6910,6 @@ void Verifier::visitIntrinsicCall(Intrinsic::ID ID, CallBase &Call) {
"@llvm.structured.alloca calls require elementtype attribute.",
&Call);
break;
- case Intrinsic::amdgcn_cs_chain: {
- auto CallerCC = Call.getCaller()->getCallingConv();
- switch (CallerCC) {
- case CallingConv::AMDGPU_CS:
- case CallingConv::AMDGPU_CS_Chain:
- case CallingConv::AMDGPU_CS_ChainPreserve:
- case CallingConv::AMDGPU_ES:
- case CallingConv::AMDGPU_GS:
- case CallingConv::AMDGPU_HS:
- case CallingConv::AMDGPU_LS:
- case CallingConv::AMDGPU_VS:
- break;
- default:
- CheckFailed("Intrinsic cannot be called from functions with this "
- "calling convention",
- &Call);
- break;
- }
-
- Check(Call.paramHasAttr(2, Attribute::InReg),
- "SGPR arguments must have the `inreg` attribute", &Call);
- Check(!Call.paramHasAttr(3, Attribute::InReg),
- "VGPR arguments must not have the `inreg` attribute", &Call);
-
- auto *FlagsArg = cast<ConstantInt>(Call.getArgOperand(4));
- Check(FlagsArg->getValue().ult(2),
- "flags must be 0 or 1 for llvm.amdgcn.cs.chain", &Call);
-
- auto *Next = Call.getNextNode();
- bool IsAMDUnreachable = Next && isa<IntrinsicInst>(Next) &&
- cast<IntrinsicInst>(Next)->getIntrinsicID() ==
- Intrinsic::amdgcn_unreachable;
- Check(Next && (isa<UnreachableInst>(Next) || IsAMDUnreachable),
- "llvm.amdgcn.cs.chain must be followed by unreachable", &Call);
- break;
- }
- case Intrinsic::amdgcn_init_exec_from_input: {
- const Argument *Arg = dyn_cast<Argument>(Call.getOperand(0));
- Check(Arg && Arg->hasInRegAttr(),
- "only inreg arguments to the parent function are valid as inputs to "
- "this intrinsic",
- &Call);
- break;
- }
- case Intrinsic::amdgcn_set_inactive_chain_arg: {
- auto CallerCC = Call.getCaller()->getCallingConv();
- switch (CallerCC) {
- case CallingConv::AMDGPU_CS_Chain:
- case CallingConv::AMDGPU_CS_ChainPreserve:
- break;
- default:
- CheckFailed("Intrinsic can only be used from functions with the "
- "amdgpu_cs_chain or amdgpu_cs_chain_preserve "
- "calling conventions",
- &Call);
- break;
- }
-
- unsigned InactiveIdx = 1;
- Check(!Call.paramHasAttr(InactiveIdx, Attribute::InReg),
- "Value for inactive lanes must not have the `inreg` attribute",
- &Call);
- Check(isa<Argument>(Call.getArgOperand(InactiveIdx)),
- "Value for inactive lanes must be a function argument", &Call);
- Check(!cast<Argument>(Call.getArgOperand(InactiveIdx))->hasInRegAttr(),
- "Value for inactive lanes must be a VGPR function argument", &Call);
- break;
- }
- case Intrinsic::amdgcn_call_whole_wave: {
- auto F = dyn_cast<Function>(Call.getArgOperand(0));
- Check(F, "Indirect whole wave calls are not allowed", &Call);
-
- CallingConv::ID CC = F->getCallingConv();
- Check(CC == CallingConv::AMDGPU_Gfx_WholeWave,
- "Callee must have the amdgpu_gfx_whole_wave calling convention",
- &Call);
-
- Check(!F->isVarArg(), "Variadic whole wave calls are not allowed", &Call);
-
- Check(Call.arg_size() == F->arg_size(),
- "Call argument count must match callee argument count", &Call);
-
- // The first argument of the call is the callee, and the first argument of
- // the callee is the active mask. The rest of the arguments must match.
- Check(F->arg_begin()->getType()->isIntegerTy(1),
- "Callee must have i1 as its first argument", &Call);
- for (auto [CallArg, FuncArg] :
- drop_begin(zip_equal(Call.args(), F->args()))) {
- Check(CallArg->getType() == FuncArg.getType(),
- "Argument types must match", &Call);
-
- // Check that inreg attributes match between call site and function
- Check(Call.paramHasAttr(FuncArg.getArgNo(), Attribute::InReg) ==
- FuncArg.hasInRegAttr(),
- "Argument inreg attributes must match", &Call);
- }
- break;
- }
- case Intrinsic::amdgcn_s_prefetch_data: {
- Check(
- AMDGPU::isFlatGlobalAddrSpace(
- Call.getArgOperand(0)->getType()->getPointerAddressSpace()),
- "llvm.amdgcn.s.prefetch.data only supports global or constant memory");
- break;
- }
- case Intrinsic::amdgcn_load_to_lds:
- case Intrinsic::amdgcn_load_async_to_lds:
- case Intrinsic::amdgcn_global_load_lds:
- case Intrinsic::amdgcn_global_load_async_lds:
- case Intrinsic::amdgcn_raw_buffer_load_lds:
- case Intrinsic::amdgcn_raw_buffer_load_async_lds:
- case Intrinsic::amdgcn_raw_ptr_buffer_load_lds:
- case Intrinsic::amdgcn_raw_ptr_buffer_load_async_lds:
- case Intrinsic::amdgcn_struct_buffer_load_lds:
- case Intrinsic::amdgcn_struct_buffer_load_async_lds:
- case Intrinsic::amdgcn_struct_ptr_buffer_load_lds:
- case Intrinsic::amdgcn_struct_ptr_buffer_load_async_lds: {
- // The data byte size immarg is operand 2 for every load-to-LDS intrinsic.
- uint64_t Size = cast<ConstantInt>(Call.getArgOperand(2))->getZExtValue();
- Check(Size == 1 || Size == 2 || Size == 4 || Size == 12 || Size == 16,
- "invalid data size for load-to-LDS intrinsic; must be 1, 2, 4, 12, "
- "or 16",
- &Call);
- break;
- }
- case Intrinsic::amdgcn_mfma_scale_f32_16x16x128_f8f6f4:
- case Intrinsic::amdgcn_mfma_scale_f32_32x32x64_f8f6f4: {
- Value *Src0 = Call.getArgOperand(0);
- Value *Src1 = Call.getArgOperand(1);
-
- uint64_t CBSZ = cast<ConstantInt>(Call.getArgOperand(3))->getZExtValue();
- uint64_t BLGP = cast<ConstantInt>(Call.getArgOperand(4))->getZExtValue();
- Check(CBSZ <= 4, "invalid value for cbsz format", Call,
- Call.getArgOperand(3));
- Check(BLGP <= 4, "invalid value for blgp format", Call,
- Call.getArgOperand(4));
-
- // AMDGPU::MFMAScaleFormats values
- auto getFormatNumRegs = [](unsigned FormatVal) {
- switch (FormatVal) {
- case 0:
- case 1:
- return 8u;
- case 2:
- case 3:
- return 6u;
- case 4:
- return 4u;
- default:
- llvm_unreachable("invalid format value");
- }
- };
-
- auto isValidSrcASrcBVector = [](FixedVectorType *Ty) {
- if (!Ty || !Ty->getElementType()->isIntegerTy(32))
- return false;
- unsigned NumElts = Ty->getNumElements();
- return NumElts == 4 || NumElts == 6 || NumElts == 8;
- };
-
- auto *Src0Ty = dyn_cast<FixedVectorType>(Src0->getType());
- auto *Src1Ty = dyn_cast<FixedVectorType>(Src1->getType());
- Check(isValidSrcASrcBVector(Src0Ty),
- "operand 0 must be 4, 6 or 8 element i32 vector", &Call, Src0);
- Check(isValidSrcASrcBVector(Src1Ty),
- "operand 1 must be 4, 6 or 8 element i32 vector", &Call, Src1);
-
- // Permit excess registers for the format.
- Check(Src0Ty->getNumElements() >= getFormatNumRegs(CBSZ),
- "invalid vector type for format", &Call, Src0, Call.getArgOperand(3));
- Check(Src1Ty->getNumElements() >= getFormatNumRegs(BLGP),
- "invalid vector type for format", &Call, Src1, Call.getArgOperand(5));
- break;
- }
- case Intrinsic::amdgcn_wmma_f32_16x16x128_f8f6f4:
- case Intrinsic::amdgcn_wmma_scale_f32_16x16x128_f8f6f4:
- case Intrinsic::amdgcn_wmma_scale16_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::amdgcn_cooperative_atomic_load_32x4B:
- case Intrinsic::amdgcn_cooperative_atomic_load_16x8B:
- case Intrinsic::amdgcn_cooperative_atomic_load_8x16B:
- case Intrinsic::amdgcn_cooperative_atomic_store_32x4B:
- case Intrinsic::amdgcn_cooperative_atomic_store_16x8B:
- case Intrinsic::amdgcn_cooperative_atomic_store_8x16B: {
- // Check we only use this intrinsic on the FLAT or GLOBAL address spaces.
- Value *PtrArg = Call.getArgOperand(0);
- const unsigned AS = PtrArg->getType()->getPointerAddressSpace();
- Check(AS == AMDGPUAS::FLAT_ADDRESS || AS == AMDGPUAS::GLOBAL_ADDRESS,
- "cooperative atomic intrinsics require a generic or global pointer",
- &Call, PtrArg);
-
- // Last argument must be a MD string
- auto *Op = cast<MetadataAsValue>(Call.getArgOperand(Call.arg_size() - 1));
- MDNode *MD = cast<MDNode>(Op->getMetadata());
- Check((MD->getNumOperands() == 1) && isa<MDString>(MD->getOperand(0)),
- "cooperative atomic intrinsics require that the last argument is a "
- "metadata string",
- &Call, Op);
- break;
- }
- case Intrinsic::amdgcn_av_load_b128:
- case Intrinsic::amdgcn_av_store_b128: {
- // Last argument must be a MD string
- auto *Op = cast<MetadataAsValue>(Call.getArgOperand(Call.arg_size() - 1));
- auto *MD = dyn_cast<MDNode>(Op->getMetadata());
- Check(MD && (MD->getNumOperands() == 1) && isa<MDString>(MD->getOperand(0)),
- "the last argument to av load/store intrinsics must be a "
- "metadata string",
- &Call, Op);
- break;
- }
case Intrinsic::nvvm_setmaxnreg_inc_sync_aligned_u32:
case Intrinsic::nvvm_setmaxnreg_dec_sync_aligned_u32: {
Value *V = Call.getArgOperand(0);
@@ -7531,6 +7021,9 @@ void Verifier::visitIntrinsicCall(Intrinsic::ID ID, CallBase &Call) {
Check(HasToken, "Missing funclet token on intrinsic call", &Call);
}
}
+
+ // Target-specific intrinsic call checks.
+ verifyAMDGPUIntrinsicCall(*this, ID, Call);
}
/// Carefully grab the subprogram from a local scope.
diff --git a/llvm/lib/IR/VerifierAMDGPU.cpp b/llvm/lib/IR/VerifierAMDGPU.cpp
new file mode 100644
index 0000000000000..04cb214ef2520
--- /dev/null
+++ b/llvm/lib/IR/VerifierAMDGPU.cpp
@@ -0,0 +1,406 @@
+//===-- VerifierAMDGPU.cpp - AMDGPU-specific IR verification ---------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file contains AMDGPU-specific IR verification logic that was extracted
+// from Verifier.cpp for code organization purposes only. These checks are
+// always compiled and linked as part of LLVMCore — this is not a target-
+// dependent IR verifier, which would require a different design.
+//
+// This file should only contain checks for AMDGPU-specific IR constructs
+// (e.g. amdgcn intrinsics, AMDGPU address spaces). It must not contain
+// checks for generic IR that might behave differently under AMDGPU.
+//
+//===----------------------------------------------------------------------===//
+
+#include "VerifierInternal.h"
+#include "llvm/ADT/StringExtras.h"
+#include "llvm/IR/CallingConv.h"
+#include "llvm/IR/Constants.h"
+#include "llvm/IR/DerivedTypes.h"
+#include "llvm/IR/Function.h"
+#include "llvm/IR/IntrinsicInst.h"
+#include "llvm/IR/IntrinsicsAMDGPU.h"
+#include "llvm/Support/AMDGPUAddrSpace.h"
+
+using namespace llvm;
+
+#define Check(C, ...) \
+ do { \
+ if (!(C)) { \
+ VS.CheckFailed(__VA_ARGS__); \
+ return; \
+ } \
+ } while (false)
+
+void llvm::verifyAMDGPUModuleFlag(VerifierSupport &VS, const MDString *ID,
+ Module::ModFlagBehavior MFB,
+ const MDNode *Op) {
+ if (ID->getString() != "amdgpu.buffer.oob.mode" &&
+ ID->getString() != "amdgpu.tbuffer.oob.mode")
+ return;
+
+ Check(MFB == Module::Max,
+ "'" + ID->getString() + "' module flag must use 'max' merge behaviour");
+ ConstantInt *Value =
+ mdconst::dyn_extract_or_null<ConstantInt>(Op->getOperand(2));
+ Check(Value, "'" + ID->getString() +
+ "' module flag must have a constant integer value");
+ Check(Value->getZExtValue() <= 2,
+ "'" + ID->getString() + "' module flag must be 0, 1, or 2");
+}
+
+// Verify that when a function has !reqd_work_group_size metadata, it also has
+// an amdgpu-flat-work-group-size attribute that matches the product of the
+// reqd_work_group_size operands.
+static void verifyAMDGPUReqdWorkGroupSize(VerifierSupport &VS,
+ const Function &F) {
+ // This is not required for other targets so we only check for AMDGPU.
+ if (!VS.TT.isAMDGPU())
+ return;
+
+ MDNode *ReqdWorkGroupSize = F.getMetadata("reqd_work_group_size");
+ if (!ReqdWorkGroupSize || ReqdWorkGroupSize->getNumOperands() != 3)
+ return;
+
+ uint64_t Product = 1;
+ for (const MDOperand &Op : ReqdWorkGroupSize->operands()) {
+ ConstantInt *C = mdconst::dyn_extract<ConstantInt>(Op);
+ if (!C || C->getValue().getActiveBits() > 64)
+ return;
+ uint64_t Dim = C->getZExtValue();
+ if (Dim != 0 && Product > std::numeric_limits<uint64_t>::max() / Dim)
+ return;
+ Product *= Dim;
+ }
+
+ Attribute FlatWorkGroupSize = F.getFnAttribute("amdgpu-flat-work-group-size");
+ if (!FlatWorkGroupSize.isValid()) {
+ VS.CheckFailed("reqd_work_group_size requires amdgpu-flat-work-group-size",
+ &F, ReqdWorkGroupSize);
+ return;
+ }
+
+ if (!FlatWorkGroupSize.isStringAttribute()) {
+ VS.CheckFailed("amdgpu-flat-work-group-size must be a string attribute",
+ &F);
+ return;
+ }
+
+ StringRef AttrValue = FlatWorkGroupSize.getValueAsString();
+ std::pair<StringRef, StringRef> Values = AttrValue.split(',');
+ uint64_t Min = 0;
+ uint64_t Max = 0;
+ bool Parsed = !Values.second.contains(',') &&
+ llvm::to_integer(Values.first.trim(), Min) &&
+ llvm::to_integer(Values.second.trim(), Max);
+ if (!Parsed) {
+ VS.CheckFailed("amdgpu-flat-work-group-size must be a pair of unsigned "
+ "integers",
+ &F);
+ return;
+ }
+
+ if (Min != Product || Max != Product) {
+ VS.CheckFailed("amdgpu-flat-work-group-size must equal the product of "
+ "reqd_work_group_size operands",
+ &F, ReqdWorkGroupSize);
+ }
+}
+
+void llvm::verifyAMDGPUFunctionMetadata(VerifierSupport &VS,
+ const Function &F) {
+ verifyAMDGPUReqdWorkGroupSize(VS, F);
+}
+
+void llvm::verifyAMDGPUAlloca(VerifierSupport &VS, const AllocaInst &AI) {
+ // This is not required for other targets so we only check for AMDGPU.
+ if (!VS.TT.isAMDGPU())
+ return;
+
+ if (AI.getAddressSpace() != AMDGPUAS::PRIVATE_ADDRESS)
+ VS.CheckFailed("alloca on amdgpu must be in addrspace(5)", &AI);
+}
+
+bool llvm::isAMDGPUCallBrIntrinsic(Intrinsic::ID ID) {
+ switch (ID) {
+ default:
+ return false;
+ case Intrinsic::amdgcn_kill:
+ return true;
+ }
+}
+
+void llvm::verifyAMDGPUIntrinsicCall(VerifierSupport &VS, Intrinsic::ID ID,
+ CallBase &Call) {
+ switch (ID) {
+ default:
+ return;
+ case Intrinsic::amdgcn_kill: {
+ if (auto *CBI = dyn_cast<CallBrInst>(&Call)) {
+ Check(CBI->getNumIndirectDests() == 1,
+ "callbr amdgcn_kill only supports one indirect dest");
+ bool Unreachable = isa<UnreachableInst>(CBI->getIndirectDest(0)->begin());
+ CallInst *CI = dyn_cast<CallInst>(CBI->getIndirectDest(0)->begin());
+ Check(Unreachable ||
+ (CI && CI->getIntrinsicID() == Intrinsic::amdgcn_unreachable),
+ "callbr amdgcn_kill indirect dest needs to be unreachable");
+ }
+ break;
+ }
+ case Intrinsic::amdgcn_cs_chain: {
+ CallingConv::ID CallerCC = Call.getCaller()->getCallingConv();
+ switch (CallerCC) {
+ case CallingConv::AMDGPU_CS:
+ case CallingConv::AMDGPU_CS_Chain:
+ case CallingConv::AMDGPU_CS_ChainPreserve:
+ case CallingConv::AMDGPU_ES:
+ case CallingConv::AMDGPU_GS:
+ case CallingConv::AMDGPU_HS:
+ case CallingConv::AMDGPU_LS:
+ case CallingConv::AMDGPU_VS:
+ break;
+ default:
+ VS.CheckFailed("Intrinsic cannot be called from functions with this "
+ "calling convention",
+ &Call);
+ break;
+ }
+
+ Check(Call.paramHasAttr(2, Attribute::InReg),
+ "SGPR arguments must have the `inreg` attribute", &Call);
+ Check(!Call.paramHasAttr(3, Attribute::InReg),
+ "VGPR arguments must not have the `inreg` attribute", &Call);
+
+ ConstantInt *FlagsArg = cast<ConstantInt>(Call.getArgOperand(4));
+ Check(FlagsArg->getValue().ult(2),
+ "flags must be 0 or 1 for llvm.amdgcn.cs.chain", &Call);
+
+ Instruction *Next = Call.getNextNode();
+ bool IsAMDUnreachable = isa_and_nonnull<IntrinsicInst>(Next) &&
+ cast<IntrinsicInst>(Next)->getIntrinsicID() ==
+ Intrinsic::amdgcn_unreachable;
+ Check(Next && (isa<UnreachableInst>(Next) || IsAMDUnreachable),
+ "llvm.amdgcn.cs.chain must be followed by unreachable", &Call);
+ break;
+ }
+ case Intrinsic::amdgcn_init_exec_from_input: {
+ const Argument *Arg = dyn_cast<Argument>(Call.getOperand(0));
+ Check(Arg && Arg->hasInRegAttr(),
+ "only inreg arguments to the parent function are valid as inputs to "
+ "this intrinsic",
+ &Call);
+ break;
+ }
+ case Intrinsic::amdgcn_set_inactive_chain_arg: {
+ CallingConv::ID CallerCC = Call.getCaller()->getCallingConv();
+ switch (CallerCC) {
+ case CallingConv::AMDGPU_CS_Chain:
+ case CallingConv::AMDGPU_CS_ChainPreserve:
+ break;
+ default:
+ VS.CheckFailed("Intrinsic can only be used from functions with the "
+ "amdgpu_cs_chain or amdgpu_cs_chain_preserve "
+ "calling conventions",
+ &Call);
+ break;
+ }
+
+ unsigned InactiveIdx = 1;
+ Check(!Call.paramHasAttr(InactiveIdx, Attribute::InReg),
+ "Value for inactive lanes must not have the `inreg` attribute",
+ &Call);
+ Check(isa<Argument>(Call.getArgOperand(InactiveIdx)),
+ "Value for inactive lanes must be a function argument", &Call);
+ Check(!cast<Argument>(Call.getArgOperand(InactiveIdx))->hasInRegAttr(),
+ "Value for inactive lanes must be a VGPR function argument", &Call);
+ break;
+ }
+ case Intrinsic::amdgcn_call_whole_wave: {
+ Function *F = dyn_cast<Function>(Call.getArgOperand(0));
+ Check(F, "Indirect whole wave calls are not allowed", &Call);
+
+ CallingConv::ID CC = F->getCallingConv();
+ Check(CC == CallingConv::AMDGPU_Gfx_WholeWave,
+ "Callee must have the amdgpu_gfx_whole_wave calling convention",
+ &Call);
+
+ Check(!F->isVarArg(), "Variadic whole wave calls are not allowed", &Call);
+
+ Check(Call.arg_size() == F->arg_size(),
+ "Call argument count must match callee argument count", &Call);
+
+ Check(F->arg_begin()->getType()->isIntegerTy(1),
+ "Callee must have i1 as its first argument", &Call);
+ for (auto [CallArg, FuncArg] :
+ drop_begin(zip_equal(Call.args(), F->args()))) {
+ Check(CallArg->getType() == FuncArg.getType(),
+ "Argument types must match", &Call);
+
+ Check(Call.paramHasAttr(FuncArg.getArgNo(), Attribute::InReg) ==
+ FuncArg.hasInRegAttr(),
+ "Argument inreg attributes must match", &Call);
+ }
+ break;
+ }
+ case Intrinsic::amdgcn_s_prefetch_data: {
+ Check(
+ AMDGPU::isFlatGlobalAddrSpace(
+ Call.getArgOperand(0)->getType()->getPointerAddressSpace()),
+ "llvm.amdgcn.s.prefetch.data only supports global or constant memory");
+ break;
+ }
+ case Intrinsic::amdgcn_load_to_lds:
+ case Intrinsic::amdgcn_load_async_to_lds:
+ case Intrinsic::amdgcn_global_load_lds:
+ case Intrinsic::amdgcn_global_load_async_lds:
+ case Intrinsic::amdgcn_raw_buffer_load_lds:
+ case Intrinsic::amdgcn_raw_buffer_load_async_lds:
+ case Intrinsic::amdgcn_raw_ptr_buffer_load_lds:
+ case Intrinsic::amdgcn_raw_ptr_buffer_load_async_lds:
+ case Intrinsic::amdgcn_struct_buffer_load_lds:
+ case Intrinsic::amdgcn_struct_buffer_load_async_lds:
+ case Intrinsic::amdgcn_struct_ptr_buffer_load_lds:
+ case Intrinsic::amdgcn_struct_ptr_buffer_load_async_lds: {
+ uint64_t Size = cast<ConstantInt>(Call.getArgOperand(2))->getZExtValue();
+ Check(Size == 1 || Size == 2 || Size == 4 || Size == 12 || Size == 16,
+ "invalid data size for load-to-LDS intrinsic; must be 1, 2, 4, 12, "
+ "or 16",
+ &Call);
+ break;
+ }
+ case Intrinsic::amdgcn_mfma_scale_f32_16x16x128_f8f6f4:
+ case Intrinsic::amdgcn_mfma_scale_f32_32x32x64_f8f6f4: {
+ Value *Src0 = Call.getArgOperand(0);
+ Value *Src1 = Call.getArgOperand(1);
+
+ uint64_t CBSZ = cast<ConstantInt>(Call.getArgOperand(3))->getZExtValue();
+ uint64_t BLGP = cast<ConstantInt>(Call.getArgOperand(4))->getZExtValue();
+ Check(CBSZ <= 4, "invalid value for cbsz format", Call,
+ Call.getArgOperand(3));
+ Check(BLGP <= 4, "invalid value for blgp format", Call,
+ Call.getArgOperand(4));
+
+ auto GetFormatNumRegs = [](unsigned FormatVal) {
+ switch (FormatVal) {
+ case 0:
+ case 1:
+ return 8u;
+ case 2:
+ case 3:
+ return 6u;
+ case 4:
+ return 4u;
+ default:
+ llvm_unreachable("invalid format value");
+ }
+ };
+
+ auto IsValidSrcASrcBVector = [](FixedVectorType *Ty) {
+ if (!Ty || !Ty->getElementType()->isIntegerTy(32))
+ return false;
+ unsigned NumElts = Ty->getNumElements();
+ return NumElts == 4 || NumElts == 6 || NumElts == 8;
+ };
+
+ FixedVectorType *Src0Ty = dyn_cast<FixedVectorType>(Src0->getType());
+ FixedVectorType *Src1Ty = dyn_cast<FixedVectorType>(Src1->getType());
+ Check(IsValidSrcASrcBVector(Src0Ty),
+ "operand 0 must be 4, 6 or 8 element i32 vector", &Call, Src0);
+ Check(IsValidSrcASrcBVector(Src1Ty),
+ "operand 1 must be 4, 6 or 8 element i32 vector", &Call, Src1);
+
+ Check(Src0Ty->getNumElements() >= GetFormatNumRegs(CBSZ),
+ "invalid vector type for format", &Call, Src0, Call.getArgOperand(3));
+ Check(Src1Ty->getNumElements() >= GetFormatNumRegs(BLGP),
+ "invalid vector type for format", &Call, Src1, Call.getArgOperand(5));
+ break;
+ }
+ case Intrinsic::amdgcn_wmma_f32_16x16x128_f8f6f4:
+ case Intrinsic::amdgcn_wmma_scale_f32_16x16x128_f8f6f4:
+ case Intrinsic::amdgcn_wmma_scale16_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));
+
+ 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;
+ };
+
+ FixedVectorType *Src0Ty = dyn_cast<FixedVectorType>(Src0->getType());
+ FixedVectorType *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);
+
+ 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::amdgcn_cooperative_atomic_load_32x4B:
+ case Intrinsic::amdgcn_cooperative_atomic_load_16x8B:
+ case Intrinsic::amdgcn_cooperative_atomic_load_8x16B:
+ case Intrinsic::amdgcn_cooperative_atomic_store_32x4B:
+ case Intrinsic::amdgcn_cooperative_atomic_store_16x8B:
+ case Intrinsic::amdgcn_cooperative_atomic_store_8x16B: {
+ Value *PtrArg = Call.getArgOperand(0);
+ const unsigned AS = PtrArg->getType()->getPointerAddressSpace();
+ Check(AS == AMDGPUAS::FLAT_ADDRESS || AS == AMDGPUAS::GLOBAL_ADDRESS,
+ "cooperative atomic intrinsics require a generic or global pointer",
+ &Call, PtrArg);
+
+ MetadataAsValue *Op =
+ cast<MetadataAsValue>(Call.getArgOperand(Call.arg_size() - 1));
+ MDNode *MD = cast<MDNode>(Op->getMetadata());
+ Check((MD->getNumOperands() == 1) && isa<MDString>(MD->getOperand(0)),
+ "cooperative atomic intrinsics require that the last argument is a "
+ "metadata string",
+ &Call, Op);
+ break;
+ }
+ case Intrinsic::amdgcn_av_load_b128:
+ case Intrinsic::amdgcn_av_store_b128: {
+ MetadataAsValue *Op =
+ cast<MetadataAsValue>(Call.getArgOperand(Call.arg_size() - 1));
+ MDNode *MD = dyn_cast<MDNode>(Op->getMetadata());
+ Check(MD && (MD->getNumOperands() == 1) && isa<MDString>(MD->getOperand(0)),
+ "the last argument to av load/store intrinsics must be a "
+ "metadata string",
+ &Call, Op);
+ break;
+ }
+ }
+}
+
+#undef Check
diff --git a/llvm/lib/IR/VerifierInternal.h b/llvm/lib/IR/VerifierInternal.h
new file mode 100644
index 0000000000000..922385230179b
--- /dev/null
+++ b/llvm/lib/IR/VerifierInternal.h
@@ -0,0 +1,233 @@
+//===-- VerifierInternal.h - Internal verifier infrastructure --------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// Shared definitions used by the verifier implementation files.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_LIB_IR_VERIFIERINTERNAL_H
+#define LLVM_LIB_IR_VERIFIERINTERNAL_H
+
+#include "llvm/ADT/Twine.h"
+#include "llvm/IR/Attributes.h"
+#include "llvm/IR/DebugInfoMetadata.h"
+#include "llvm/IR/Instructions.h"
+#include "llvm/IR/Metadata.h"
+#include "llvm/IR/Module.h"
+#include "llvm/IR/ModuleSlotTracker.h"
+#include "llvm/IR/Value.h"
+#include "llvm/Support/Casting.h"
+#include "llvm/Support/Printable.h"
+#include "llvm/Support/raw_ostream.h"
+#include "llvm/TargetParser/Triple.h"
+
+namespace llvm {
+
+struct VerifierSupport {
+ raw_ostream *OS;
+ const Module &M;
+ ModuleSlotTracker MST;
+ const Triple &TT;
+ const DataLayout &DL;
+ LLVMContext &Context;
+
+ /// Track the brokenness of the module while recursively visiting.
+ bool Broken = false;
+ /// Broken debug info can be "recovered" from by stripping the debug info.
+ bool BrokenDebugInfo = false;
+ /// Whether to treat broken debug info as an error.
+ bool TreatBrokenDebugInfoAsError = true;
+
+ explicit VerifierSupport(raw_ostream *OS, const Module &M)
+ : OS(OS), M(M), MST(&M), TT(M.getTargetTriple()), DL(M.getDataLayout()),
+ Context(M.getContext()) {}
+
+private:
+ void Write(const Module *M) {
+ *OS << "; ModuleID = '" << M->getModuleIdentifier() << "'\n";
+ }
+
+ void Write(const Value *V) {
+ if (V)
+ Write(*V);
+ }
+
+ void Write(const Value &V) {
+ if (isa<Instruction>(V)) {
+ V.print(*OS, MST);
+ *OS << '\n';
+ } else {
+ V.printAsOperand(*OS, true, MST);
+ *OS << '\n';
+ }
+ }
+
+ void Write(const DbgRecord *DR) {
+ if (DR) {
+ DR->print(*OS, MST, false);
+ *OS << '\n';
+ }
+ }
+
+ void Write(DbgVariableRecord::LocationType Type) {
+ switch (Type) {
+ case DbgVariableRecord::LocationType::Value:
+ *OS << "value";
+ break;
+ case DbgVariableRecord::LocationType::Declare:
+ *OS << "declare";
+ break;
+ case DbgVariableRecord::LocationType::DeclareValue:
+ *OS << "declare_value";
+ break;
+ case DbgVariableRecord::LocationType::Assign:
+ *OS << "assign";
+ break;
+ case DbgVariableRecord::LocationType::End:
+ *OS << "end";
+ break;
+ case DbgVariableRecord::LocationType::Any:
+ *OS << "any";
+ break;
+ };
+ }
+
+ void Write(const Metadata *MD) {
+ if (!MD)
+ return;
+ MD->print(*OS, MST, &M);
+ *OS << '\n';
+ }
+
+ template <class T> void Write(const MDTupleTypedArrayWrapper<T> &MD) {
+ Write(MD.get());
+ }
+
+ void Write(const NamedMDNode *NMD) {
+ if (!NMD)
+ return;
+ NMD->print(*OS, MST);
+ *OS << '\n';
+ }
+
+ void Write(Type *T) {
+ if (!T)
+ return;
+ *OS << ' ' << *T;
+ }
+
+ void Write(const Comdat *C) {
+ if (!C)
+ return;
+ *OS << *C;
+ }
+
+ void Write(const APInt *AI) {
+ if (!AI)
+ return;
+ *OS << *AI << '\n';
+ }
+
+ void Write(const unsigned i) { *OS << i << '\n'; }
+
+ // NOLINTNEXTLINE(readability-identifier-naming)
+ void Write(const Attribute *A) {
+ if (!A)
+ return;
+ *OS << A->getAsString() << '\n';
+ }
+
+ // NOLINTNEXTLINE(readability-identifier-naming)
+ void Write(const AttributeSet *AS) {
+ if (!AS)
+ return;
+ *OS << AS->getAsString() << '\n';
+ }
+
+ // NOLINTNEXTLINE(readability-identifier-naming)
+ void Write(const AttributeList *AL) {
+ if (!AL)
+ return;
+ AL->print(*OS);
+ }
+
+ void Write(Printable P) { *OS << P << '\n'; }
+
+ template <typename T> void Write(ArrayRef<T> Vs) {
+ for (const T &V : Vs)
+ Write(V);
+ }
+
+ template <typename T1, typename... Ts>
+ void WriteTs(const T1 &V1, const Ts &...Vs) {
+ Write(V1);
+ WriteTs(Vs...);
+ }
+
+ template <typename... Ts> void WriteTs() {}
+
+public:
+ /// A check failed, so printout out the condition and the message.
+ ///
+ /// This provides a nice place to put a breakpoint if you want to see why
+ /// something is not correct.
+ void CheckFailed(const Twine &Message) {
+ if (OS)
+ *OS << Message << '\n';
+ Broken = true;
+ }
+
+ /// A check failed (with values to print).
+ ///
+ /// This calls the Message-only version so that the above is easier to set a
+ /// breakpoint on.
+ template <typename T1, typename... Ts>
+ void CheckFailed(const Twine &Message, const T1 &V1, const Ts &...Vs) {
+ CheckFailed(Message);
+ if (OS)
+ WriteTs(V1, Vs...);
+ }
+
+ /// A debug info check failed.
+ void DebugInfoCheckFailed(const Twine &Message) {
+ if (OS)
+ *OS << Message << '\n';
+ Broken |= TreatBrokenDebugInfoAsError;
+ BrokenDebugInfo = true;
+ }
+
+ /// A debug info check failed (with values to print).
+ template <typename T1, typename... Ts>
+ void DebugInfoCheckFailed(const Twine &Message, const T1 &V1,
+ const Ts &...Vs) {
+ DebugInfoCheckFailed(Message);
+ if (OS)
+ WriteTs(V1, Vs...);
+ }
+};
+
+//==============================================================================
+// AMDGPU-specific verification functions
+
+void verifyAMDGPUModuleFlag(VerifierSupport &VS, const MDString *ID,
+ Module::ModFlagBehavior MFB, const MDNode *Op);
+
+void verifyAMDGPUFunctionMetadata(VerifierSupport &VS, const Function &F);
+
+void verifyAMDGPUAlloca(VerifierSupport &VS, const AllocaInst &AI);
+
+void verifyAMDGPUIntrinsicCall(VerifierSupport &VS, Intrinsic::ID ID,
+ CallBase &Call);
+
+bool isAMDGPUCallBrIntrinsic(Intrinsic::ID ID);
+
+//==============================================================================
+
+} // namespace llvm
+
+#endif // LLVM_LIB_IR_VERIFIERINTERNAL_H
diff --git a/llvm/test/Verifier/callbr-intrinsic.ll b/llvm/test/Verifier/callbr-intrinsic.ll
index 5cc9e7034e2bc..eec73fb796698 100644
--- a/llvm/test/Verifier/callbr-intrinsic.ll
+++ b/llvm/test/Verifier/callbr-intrinsic.ll
@@ -2,7 +2,7 @@
declare void @llvm.amdgcn.kill(i1)
-; CHECK: Callbr amdgcn_kill only supports one indirect dest
+; CHECK: callbr amdgcn_kill only supports one indirect dest
define void @test_callbr_intrinsic_indirect0(i1 %c) {
callbr void @llvm.amdgcn.kill(i1 %c) to label %cont []
kill:
@@ -11,7 +11,7 @@ cont:
ret void
}
-; CHECK-NEXT: Callbr amdgcn_kill only supports one indirect dest
+; CHECK-NEXT: callbr amdgcn_kill only supports one indirect dest
define void @test_callbr_intrinsic_indirect2(i1 %c) {
callbr void @llvm.amdgcn.kill(i1 %c) to label %cont [label %kill1, label %kill2]
kill1:
@@ -22,7 +22,7 @@ cont:
ret void
}
-; CHECK-NEXT: Callbr amdgcn_kill indirect dest needs to be unreachable
+; CHECK-NEXT: callbr amdgcn_kill indirect dest needs to be unreachable
define void @test_callbr_intrinsic_no_unreachable(i1 %c) {
callbr void @llvm.amdgcn.kill(i1 %c) to label %cont [label %kill]
kill:
@@ -31,7 +31,7 @@ cont:
ret void
}
-; CHECK-NEXT: Callbr currently only supports asm-goto and selected intrinsics
+; CHECK-NEXT: callbr currently only supports asm-goto and selected intrinsics
declare i32 @llvm.amdgcn.workitem.id.x()
define void @test_callbr_intrinsic_unsupported() {
callbr i32 @llvm.amdgcn.workitem.id.x() to label %cont []
@@ -39,7 +39,7 @@ cont:
ret void
}
-; CHECK-NEXT: Callbr: indirect function / invalid signature
+; CHECK-NEXT: callbr: indirect function / invalid signature
define void @test_callbr_intrinsic_wrong_signature(ptr %ptr) {
%func = load ptr, ptr %ptr, align 8
callbr void %func() to label %cont []
@@ -47,7 +47,7 @@ cont:
ret void
}
-; CHECK-NEXT: Callbr for intrinsics currently doesn't support operand bundles
+; CHECK-NEXT: callbr for intrinsics currently doesn't support operand bundles
define void @test_callbr_intrinsic_no_operand_bundles(i1 %c) {
callbr void @llvm.amdgcn.kill(i1 %c) [ "foo"(i1 %c) ] to label %cont [label %kill]
kill:
diff --git a/llvm/utils/gn/secondary/llvm/lib/IR/BUILD.gn b/llvm/utils/gn/secondary/llvm/lib/IR/BUILD.gn
index b91e7f72cf712..1135884bdfd68 100644
--- a/llvm/utils/gn/secondary/llvm/lib/IR/BUILD.gn
+++ b/llvm/utils/gn/secondary/llvm/lib/IR/BUILD.gn
@@ -95,5 +95,6 @@ static_library("IR") {
"ValueSymbolTable.cpp",
"VectorTypeUtils.cpp",
"Verifier.cpp",
+ "VerifierAMDGPU.cpp",
]
}
More information about the llvm-commits
mailing list