https://github.com/mgcarrasco updated https://github.com/llvm/llvm-project/pull/222011
>From 1f115a0e2e7aa83f040fb762db8018259879d770 Mon Sep 17 00:00:00 2001 From: Manuel Carrasco <[email protected]> Date: Tue, 8 Sep 2026 09:14:20 -0500 Subject: [PATCH 1/2] [SPIRV] Emit NonSemantic DebugDeclare, DebugExpression and DebugOperation. Add support for [DebugDeclare](https://github.khronos.org/SPIRV-Registry/nonsemantic/NonSemantic.Shader.DebugInfo.html#DebugDeclare), [DebugExpression](https://github.khronos.org/SPIRV-Registry/nonsemantic/NonSemantic.Shader.DebugInfo.html#DebugExpression) and [DebugOperation](https://github.khronos.org/SPIRV-Registry/nonsemantic/NonSemantic.Shader.DebugInfo.html#DebugOperation). Changes: - Emit DebugDeclare from an indirect DBG_VALUE whose location register is defined by OpVariable, which is what the spec requires of the Variable operand. #dbg_declare does not survive as its own opcode in MIR, IRTranslator lowers it to an indirect DBG_VALUE. - Skip every other shape instead of emitting an invalid instruction: an access chain into a field, an OpFunctionParameter for a byval argument, a constant address, a register left without a def after dead storage was erased, a variadic #dbg_value, an expression using an unmapped operation, and a variable whose DebugLocalVariable was not emitted. - Collect expressions from MIR, so DebugExpression and DebugOperation are emitted at module scope and shared by every declare using the same DIExpression. - Map the nine operations the instruction set defines, and reject an expression whole when an argument does not fit in 32 bits (as per spec). - Cache DebugLocalVariable result ids, which emitDebugLocalVariable already returned but nothing stored. - Add tests for each emitted and skipped case, for the interaction with DebugLine and DebugScope, and for the expression operations. spirv-val and the spec disagree in two places, in opposite directions: - DebugDeclare's Variable operand: spirv-val accepts an OpFunctionParameter, while the spec names only OpVariable. This patch follows the spec, so the byval case is skipped even though spirv-val would take it, see debug-declare-function-parameter.ll. - DebugGlobalVariable's Variable operand: spirv-val rejects a DebugExpression there although the spec permits one, KhronosGroup/SPIRV-Tools#6469, see debug-global-variable-constant-value.ll. --- .../SPIRV/SPIRVNonSemanticDebugHandler.cpp | 192 +++++++++++++++++- .../SPIRV/SPIRVNonSemanticDebugHandler.h | 53 ++++- .../debug-info/debug-declare-access-chain.ll | 36 ++++ .../debug-declare-dbg-value-variadic.ll | 36 ++++ .../debug-info/debug-declare-dead-alloca.ll | 32 +++ .../debug-declare-expression-unsupported.ll | 31 +++ .../debug-declare-expression-xderef.ll | 60 ++++++ .../debug-declare-function-parameter.ll | 39 ++++ .../debug-info/debug-declare-int-storage.ll | 34 ++++ .../debug-info/debug-declare-line-scope.ll | 98 +++++++++ .../debug-declare-module-scope-variable.ll | 45 ++++ .../debug-info/debug-declare-null-storage.ll | 30 +++ .../debug-declare-poison-storage.ll | 28 +++ .../debug-info/debug-declare-skip-type.ll | 33 +++ .../CodeGen/SPIRV/debug-info/debug-declare.ll | 55 +++++ .../debug-info/debug-expression-bit-piece.ll | 40 ++++ .../debug-info/debug-expression-operations.ll | 54 +++++ .../debug-expression-out-of-range.ll | 76 +++++++ .../debug-global-variable-constant-value.ll | 76 +++++++ .../debug-global-variable-init-expr.ll | 40 +++- .../debug-local-variable-dbg-value.ll | 7 +- 21 files changed, 1074 insertions(+), 21 deletions(-) create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-declare-access-chain.ll create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-declare-dbg-value-variadic.ll create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-declare-dead-alloca.ll create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-declare-expression-unsupported.ll create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-declare-expression-xderef.ll create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-declare-function-parameter.ll create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-declare-int-storage.ll create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-declare-line-scope.ll create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-declare-module-scope-variable.ll create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-declare-null-storage.ll create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-declare-poison-storage.ll create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-declare-skip-type.ll create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-declare.ll create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-expression-bit-piece.ll create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-expression-operations.ll create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-expression-out-of-range.ll create mode 100644 llvm/test/CodeGen/SPIRV/debug-info/debug-global-variable-constant-value.ll diff --git a/llvm/lib/Target/SPIRV/SPIRVNonSemanticDebugHandler.cpp b/llvm/lib/Target/SPIRV/SPIRVNonSemanticDebugHandler.cpp index af553d75f6c04..8088555404562 100644 --- a/llvm/lib/Target/SPIRV/SPIRVNonSemanticDebugHandler.cpp +++ b/llvm/lib/Target/SPIRV/SPIRVNonSemanticDebugHandler.cpp @@ -18,6 +18,7 @@ #include "llvm/CodeGen/AsmPrinter.h" #include "llvm/CodeGen/MachineFunction.h" #include "llvm/CodeGen/MachineInstr.h" +#include "llvm/CodeGen/MachineModuleInfo.h" #include "llvm/IR/DebugInfo.h" #include "llvm/IR/DebugInfoMetadata.h" #include "llvm/IR/DebugProgramInstruction.h" @@ -28,6 +29,7 @@ #include "llvm/MC/MCInst.h" #include "llvm/MC/MCStreamer.h" #include "llvm/Support/ErrorHandling.h" +#include "llvm/Support/MathExtras.h" #include "llvm/Support/Path.h" #include <cassert> @@ -320,6 +322,8 @@ void SPIRVNonSemanticDebugHandler::beginModule(Module *M) { UniqueDebugLocations.clear(); GlobalVariableDebugInfoMap.clear(); LocalVariables.clear(); + DebugLocalVariableRegs.clear(); + DebugExpressionRegs.clear(); LexicalBlocks.clear(); DebugScopeRegs.clear(); DebugInlinedAtRegs.clear(); @@ -842,10 +846,82 @@ std::optional<MCRegister> SPIRVNonSemanticDebugHandler::mapDISignatureTypeToReg( return lookupOptReg(DebugScopeRegs, Ty); } -// Unimplemented no-op; see emitDebugExpression declaration. +// NonSemantic.Shader.DebugInfo.100 debug operation encodings +// (section 4.5, "Debug Operations"). +namespace NonSemanticDebugOp { +enum : uint32_t { + Deref = 0, + Plus = 1, + Minus = 2, + PlusUconst = 3, + BitPiece = 4, + Swap = 5, + Xderef = 6, + StackValue = 7, + Constu = 8, + Fragment = 9 +}; +} // namespace NonSemanticDebugOp + +static std::optional<uint32_t> mapDwarfOpToNonSemanticOp(uint64_t DwarfOp) { + switch (DwarfOp) { + case dwarf::DW_OP_deref: + return NonSemanticDebugOp::Deref; + case dwarf::DW_OP_plus: + return NonSemanticDebugOp::Plus; + case dwarf::DW_OP_minus: + return NonSemanticDebugOp::Minus; + case dwarf::DW_OP_plus_uconst: + return NonSemanticDebugOp::PlusUconst; + case dwarf::DW_OP_bit_piece: + return NonSemanticDebugOp::BitPiece; + case dwarf::DW_OP_swap: + return NonSemanticDebugOp::Swap; + case dwarf::DW_OP_xderef: + return NonSemanticDebugOp::Xderef; + case dwarf::DW_OP_stack_value: + return NonSemanticDebugOp::StackValue; + case dwarf::DW_OP_constu: + return NonSemanticDebugOp::Constu; + case dwarf::DW_OP_LLVM_fragment: + return NonSemanticDebugOp::Fragment; + default: + return std::nullopt; + } +} + std::optional<MCRegister> SPIRVNonSemanticDebugHandler::emitDebugExpression( - const DIExpression *, MCRegister, MCRegister, SPIRV::ModuleAnalysisInfo &) { - return std::nullopt; + const DIExpression *Expr, MCRegister VoidTypeReg, MCRegister I32TypeReg, + MCRegister ExtInstSetReg, SPIRV::ModuleAnalysisInfo &MAI) { + assert(Expr && "Expr must not be null in emitDebugExpression"); + + // Check the whole expression before emitting anything, and bail out if + // unsupported. Verify that the operation is supported, and that each argument + // is a 32-bit constant (as per spec). + for (const DIExpression::ExprOperand &Op : Expr->expr_ops()) { + if (!mapDwarfOpToNonSemanticOp(Op.getOp())) + return std::nullopt; + for (unsigned I = 0, E = Op.getNumArgs(); I != E; ++I) + if (!isUInt<32>(Op.getArg(I))) + return std::nullopt; + } + + SmallVector<MCRegister, 4> OperationRegs; + for (const DIExpression::ExprOperand &Op : Expr->expr_ops()) { + SmallVector<MCRegister, 3> Operands{emitOpConstantI32( + *mapDwarfOpToNonSemanticOp(Op.getOp()), I32TypeReg, MAI)}; + for (unsigned I = 0, E = Op.getNumArgs(); I != E; ++I) + // Operands are truncated to 32 bits but we already checked that they are + // in range. + Operands.push_back(emitOpConstantI32(static_cast<uint32_t>(Op.getArg(I)), + I32TypeReg, MAI)); + OperationRegs.push_back( + emitExtInst(SPIRV::NonSemanticExtInst::DebugOperation, VoidTypeReg, + ExtInstSetReg, Operands, MAI)); + } + + return emitExtInst(SPIRV::NonSemanticExtInst::DebugExpression, VoidTypeReg, + ExtInstSetReg, OperationRegs, MAI); } std::optional<MCRegister> SPIRVNonSemanticDebugHandler::emitDebugGlobalVariable( @@ -894,15 +970,18 @@ std::optional<MCRegister> SPIRVNonSemanticDebugHandler::emitDebugGlobalVariable( MCRegister ColReg = emitOpConstantI32(0, I32TypeReg, MAI); // Variable: @g OpVariable id when !dbg matches; else a DebugExpression for - // the GVE init value when no @g exists; else DebugInfoNone. + // the GVE init value when no @g exists and the expression is non-empty; else + // DebugInfoNone. As per spec, the DebugExpression must contains the constant + // value of the variable that was optimized out. An empty expression contains + // no value, so we emit DebugInfoNone instead. MCRegister VariableReg = CachedDebugInfoNoneReg; if (const GlobalVariable *LLVMGV = Info.LLVMGV) { MCRegister GVReg = MAI.getGlobalObjReg(LLVMGV); if (GVReg.isValid()) VariableReg = GVReg; - } else if (Info.Expr) { - if (auto ExprReg = - emitDebugExpression(Info.Expr, VoidTypeReg, ExtInstSetReg, MAI)) + } else if (Info.Expr && Info.Expr->getNumElements() != 0) { + if (auto ExprReg = emitDebugExpression(Info.Expr, VoidTypeReg, I32TypeReg, + ExtInstSetReg, MAI)) VariableReg = *ExprReg; } @@ -1300,6 +1379,72 @@ void SPIRVNonSemanticDebugHandler::beginInstruction(const MachineInstr *MI) { emitDebugScopeForInstruction(*Target); emitDebugLineForInstruction(*Target); + + emitDebugDeclare(MI); +} + +// The register that holds the variable's address in \p MI, or std::nullopt +// when \p MI is not a declare this backend can describe. +// +// The spec requires DebugDeclare's Variable operand to be "the <id> of an +// OpVariable instruction that defines the local variable". MIR has no +// DBG_DECLARE, so what this looks for is an indirect DBG_VALUE whose location +// register an OpVariable defines. +static std::optional<Register> +getDebugDeclareStorageReg(const MachineInstr &MI) { + // #dbg_declare is an indirect DBG_VALUE in MIR; #dbg_value is normally a + // direct one except for the variadic case. + if (!MI.isIndirectDebugValue()) + return std::nullopt; + + // A variadic #dbg_value becomes DBG_VALUE $noreg, 0, ... which is indirect + // too, and $noreg is not virtual. + Register LocReg = MI.getDebugOperand(0).getReg(); + if (!LocReg.isVirtual()) + return std::nullopt; + + // DebugDeclare can only encode the address of an OpVariable. + // Other legitimate #dbg_declare cannot be encoded. + // Examples: an access chain for a field, an OpFunctionParameter for a byval + // argument, or a module-scope constant for a null or fixed address. + + // LocReg may also have no def at all: erasing dead storage leaves the + // DBG_VALUE pointing at an undefined register. MachineVerifier permits that + // because LiveDebugVariables normally clears it, but this pipeline has no + // register allocation, so LiveDebugVariables never runs. + const MachineInstr *Def = MI.getMF()->getRegInfo().getUniqueVRegDef(LocReg); + if (!Def || Def->getOpcode() != SPIRV::OpVariable) + return std::nullopt; + + return LocReg; +} + +void SPIRVNonSemanticDebugHandler::emitDebugDeclare(const MachineInstr *MI) { + assert(DebugFunctionDefinitionEmitted && + "DebugFunctionDefinition must be emitted"); + assert(CurrentMAI && "CurrentMAI must be set"); + + std::optional<Register> LocReg = getDebugDeclareStorageReg(*MI); + if (!LocReg) + return; + + auto VarRegOpt = lookupOptReg(DebugLocalVariableRegs, MI->getDebugVariable()); + if (!VarRegOpt) + return; + + auto ExprRegOpt = lookupOptReg(DebugExpressionRegs, MI->getDebugExpression()); + if (!ExprRegOpt) + return; + + SPIRV::ModuleAnalysisInfo &MAI = *CurrentMAI; + MCRegister StorageReg = MAI.getRegisterAlias(MI->getMF(), *LocReg); + if (!StorageReg.isValid()) + return; + + MCRegister VoidTypeReg = getOrEmitOpTypeVoidReg(MAI); + MCRegister ExtInstSetReg = MAI.getExtInstSetReg(NSSet); + emitExtInst(SPIRV::NonSemanticExtInst::DebugDeclare, VoidTypeReg, + ExtInstSetReg, {*VarRegOpt, StorageReg, *ExprRegOpt}, MAI); } static bool isMergeInstruction(unsigned Opcode) { @@ -1519,6 +1664,22 @@ void SPIRVNonSemanticDebugHandler::notifyEntryLabelEmitted( tryEmitDebugFunctionDefinition(*CurrentMAI); } +void SPIRVNonSemanticDebugHandler::collectDebugExpressions( + SetVector<const DIExpression *> &Out) const { + MachineModuleInfo *ModuleInfo = Asm->MMI; + assert(ModuleInfo && "MachineModuleInfo must be set during module output"); + + for (const Function &F : *ModuleInfo->getModule()) { + const MachineFunction *MF = ModuleInfo->getMachineFunction(F); + if (!MF) + continue; + for (const MachineBasicBlock &MBB : *MF) + for (const MachineInstr &MI : MBB) + if (MI.isDebugValueLike()) + Out.insert(MI.getDebugExpression()); + } +} + void SPIRVNonSemanticDebugHandler::emitNonSemanticGlobalDebugInfo( SPIRV::ModuleAnalysisInfo &MAI) { if (GlobalDIEmitted) @@ -1729,9 +1890,22 @@ void SPIRVNonSemanticDebugHandler::emitNonSemanticGlobalDebugInfo( } // Emit DebugLocalVariable after DebugFunction and their lexical blocks so the - // Parent operand can resolve. + // Parent operand can resolve. Record the ids for DebugDeclare. for (const DILocalVariable *LV : LocalVariables) - emitDebugLocalVariable(LV, VoidTypeReg, I32TypeReg, ExtInstSetReg, MAI); + if (auto LVReg = emitDebugLocalVariable(LV, VoidTypeReg, I32TypeReg, + ExtInstSetReg, MAI)) + DebugLocalVariableRegs[LV] = *LVReg; + + // Opcodes like DebugDeclare are part of the function body, but + // DebugExpression cannot. For such opcodes, we collect the expressions + // directly from the MIR to avoid inconsistencies with those in the LLVM IR + // module. + SetVector<const DIExpression *> Expressions; + collectDebugExpressions(Expressions); + for (const DIExpression *Expr : Expressions) + if (auto ExprReg = emitDebugExpression(Expr, VoidTypeReg, I32TypeReg, + ExtInstSetReg, MAI)) + DebugExpressionRegs[Expr] = *ExprReg; // Emit DebugGlobalVariable for each collected DIGlobalVariable. for (const auto &[GV, Info] : GlobalVariableDebugInfoMap) diff --git a/llvm/lib/Target/SPIRV/SPIRVNonSemanticDebugHandler.h b/llvm/lib/Target/SPIRV/SPIRVNonSemanticDebugHandler.h index 2650402d7c796..fbfb7bfce1666 100644 --- a/llvm/lib/Target/SPIRV/SPIRVNonSemanticDebugHandler.h +++ b/llvm/lib/Target/SPIRV/SPIRVNonSemanticDebugHandler.h @@ -108,6 +108,16 @@ class SPIRVNonSemanticDebugHandler : public DebugHandlerBase { // records, dbg intrinsics, and subprogram retained nodes. SetVector<const DILocalVariable *> LocalVariables; + // DebugLocalVariable result id per variable that module-scope emission + // actually emitted. DebugDeclare needs it for its Local Variable operand; a + // variable missing here (skipped type or scope) gets no declare. + DenseMap<const DILocalVariable *, MCRegister> DebugLocalVariableRegs; + + // DebugExpression result id per DIExpression that could be lowered. An + // expression missing here uses operations with no NonSemantic counterpart, + // so declares referencing it are skipped rather than described wrongly. + DenseMap<const DIExpression *, MCRegister> DebugExpressionRegs; + // Distinct DILexicalBlock and DINamespace scopes, parent-before-child // order, collected in beginModule() for DebugLexicalBlock emission. SetVector<const DIScope *> LexicalBlocks; @@ -394,14 +404,51 @@ class SPIRVNonSemanticDebugHandler : public DebugHandlerBase { MCRegister VoidTypeReg, MCRegister I32TypeReg, MCRegister ExtInstSetReg, SPIRV::ModuleAnalysisInfo &MAI); - /// Emit \c DebugExpression for \p Expr. Unimplemented: defined as a no-op - /// (\returns \c std::nullopt, emits nothing) so \c emitDebugGlobalVariable - /// can complete Variable-operand resolution for the opcodes we support today. + /// Collect the \c DIExpression of every debug value in the module + /// (\c DBG_VALUE, \c DBG_VALUE_LIST, \c DBG_INSTR_REF), in MIR order. + /// + /// Reads MIR rather than IR because only MIR shows which debug values + /// survived codegen and in what form, and because an expression synthesized + /// during lowering never appears in the IR at all. Must be called from + /// module-scope emission, which is where the resulting \c DebugExpression + /// instructions have to be emitted; every \c MachineFunction is still + /// reachable at that point through \c MachineModuleInfo. + /// + /// Deliberately independent of what the consumers can currently emit, so + /// that adding an instruction that needs an expression (\c DebugValue) needs + /// no change here. The cost is a \c DebugExpression that nothing references + /// yet, for a debug value no instruction is emitted for. + void collectDebugExpressions(SetVector<const DIExpression *> &Out) const; + + /// Emit one \c DebugOperation per element of \p Expr followed by the + /// \c DebugExpression that lists them. An empty \p Expr yields a + /// \c DebugExpression with no operands, which is what a plain + /// \c !DIExpression() means. + /// + /// Must be called from module-scope emission only: \c DebugExpression and + /// \c DebugOperation are not in the spec's list of instructions allowed + /// inside a function, and forward references were removed in Rev 2. + /// + /// \returns The result id register on success. Returns \c std::nullopt and + /// emits nothing if any element has no NonSemantic counterpart, or carries an + /// argument too large for the 32-bit \c OpConstant operands this set + /// requires. std::optional<MCRegister> emitDebugExpression(const DIExpression *Expr, MCRegister VoidTypeReg, + MCRegister I32TypeReg, MCRegister ExtInstSetReg, SPIRV::ModuleAnalysisInfo &MAI); + /// Emit \c DebugDeclare for \p MI when it is an indirect \c DBG_VALUE whose + /// location register is defined by \c OpVariable, which is the shape + /// \c IRTranslator gives a \c #dbg_declare on storage the backend kept. + /// + /// Emits nothing when \p MI is not such a declare, when the variable has no + /// \c DebugLocalVariable, when the expression was not lowered, or when the + /// storage is anything other than an \c OpVariable (an access chain, a + /// constant, a function parameter, or a dead alloca with no def at all). + void emitDebugDeclare(const MachineInstr *MI); + /// Emit \c DebugTypeVector for the vector composite type \p VT. /// /// \returns The result id register on success. Returns \c std::nullopt and diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-access-chain.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-access-chain.ll new file mode 100644 index 0000000000000..9c051a6ac9622 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-access-chain.ll @@ -0,0 +1,36 @@ +; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s --implicit-check-not=DebugDeclare +; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; A declare whose address is a field of an aggregate. The location register is +; defined by OpInBoundsPtrAccessChain, not OpVariable, so the declare is +; dropped. +; clang seems to emit dbg_declare on the aggregate's own alloca but not in the GEP. + +; CHECK: OpExtInst {{.*}} DebugLocalVariable + +target triple = "spirv64-unknown-unknown" + +%struct.S = type { i32, i32 } + +define spir_func void @f() !dbg !5 { +entry: + %s = alloca %struct.S, align 4 + %b = getelementptr inbounds %struct.S, ptr %s, i32 0, i32 1 + #dbg_declare(ptr %b, !9, !DIExpression(), !10) + store i32 1, ptr %b, align 4, !dbg !10 + ret void, !dbg !10 +} + +!llvm.dbg.cu = !{!0} +!llvm.module.flags = !{!2, !3} + +!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None) +!1 = !DIFile(filename: "debug-declare-access-chain.c", directory: "/src") +!2 = !{i32 7, !"Dwarf Version", i32 5} +!3 = !{i32 2, !"Debug Info Version", i32 3} +!4 = !DISubroutineType(types: !6) +!6 = !{null} +!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed) +!5 = distinct !DISubprogram(name: "f", linkageName: "f", scope: !1, file: !1, line: 1, type: !4, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0) +!9 = !DILocalVariable(name: "b", scope: !5, file: !1, line: 2, type: !7) +!10 = !DILocation(line: 3, column: 1, scope: !5) diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-dbg-value-variadic.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-dbg-value-variadic.ll new file mode 100644 index 0000000000000..ee22e024a2397 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-dbg-value-variadic.ll @@ -0,0 +1,36 @@ +; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s --implicit-check-not=DebugDeclare --implicit-check-not=DebugExpression +; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; A variadic #dbg_value. IRTranslator cannot lower the DIArgList. +; It emits DBG_VALUE $noreg, 0, which is an indirect DBG_VALUE just like a +; declare is. +; The test checks that it must not become a DebugDeclare. + +; Clang seem to emit no DIArgList. The optimizer does, +; when it deletes a dead binary operation and rewrites the variable as an +; expression over the two operands. + +; CHECK: OpExtInst {{.*}} DebugLocalVariable + +target triple = "spirv64-unknown-unknown" + +define spir_func i32 @sum(i32 %a, i32 %b) !dbg !5 { +entry: + %add = add nsw i32 %a, %b, !dbg !11 + #dbg_value(!DIArgList(i32 %a, i32 %b), !9, !DIExpression(DW_OP_LLVM_arg, 0, DW_OP_LLVM_arg, 1, DW_OP_plus, DW_OP_stack_value), !11) + ret i32 %add, !dbg !11 +} + +!llvm.dbg.cu = !{!0} +!llvm.module.flags = !{!2, !3} + +!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None) +!1 = !DIFile(filename: "debug-declare-dbg-value-variadic.c", directory: "/src") +!2 = !{i32 7, !"Dwarf Version", i32 5} +!3 = !{i32 2, !"Debug Info Version", i32 3} +!4 = !DISubroutineType(types: !6) +!6 = !{!7, !7, !7} +!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed) +!5 = distinct !DISubprogram(name: "sum", linkageName: "sum", scope: !1, file: !1, line: 1, type: !4, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0) +!9 = !DILocalVariable(name: "total", scope: !5, file: !1, line: 2, type: !7) +!11 = !DILocation(line: 2, column: 7, scope: !5) diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-dead-alloca.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-dead-alloca.ll new file mode 100644 index 0000000000000..0c7cde2bcaeb5 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-dead-alloca.ll @@ -0,0 +1,32 @@ +; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s --implicit-check-not=DebugDeclare +; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; An alloca whose only use is the declare, which is what clang emits at -O0 for +; an unused local. Nothing keeps the alloca alive, so no OpVariable is left to +; point at and the declare is dropped. The DebugLocalVariable still describes +; the variable. + +; CHECK: OpExtInst {{.*}} DebugLocalVariable + +target triple = "spirv64-unknown-unknown" + +define spir_func void @f() !dbg !5 { +entry: + %unused = alloca i32, align 4 + #dbg_declare(ptr %unused, !9, !DIExpression(), !10) + ret void, !dbg !10 +} + +!llvm.dbg.cu = !{!0} +!llvm.module.flags = !{!2, !3} + +!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None) +!1 = !DIFile(filename: "debug-declare-dead-alloca.c", directory: "/src") +!2 = !{i32 7, !"Dwarf Version", i32 5} +!3 = !{i32 2, !"Debug Info Version", i32 3} +!4 = !DISubroutineType(types: !6) +!6 = !{null} +!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed) +!5 = distinct !DISubprogram(name: "f", linkageName: "f", scope: !1, file: !1, line: 1, type: !4, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0) +!9 = !DILocalVariable(name: "unused", scope: !5, file: !1, line: 2, type: !7) +!10 = !DILocation(line: 3, column: 1, scope: !5) diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-expression-unsupported.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-expression-unsupported.ll new file mode 100644 index 0000000000000..a816590c2604b --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-expression-unsupported.ll @@ -0,0 +1,31 @@ +; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s --implicit-check-not=DebugDeclare --implicit-check-not=DebugExpression --implicit-check-not=DebugOperation +; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; A declare whose expression contains an operation with no counterpart in +; NonSemantic.Shader.DebugInfo.100. No DebugExpression is built for it. + +; CHECK: OpExtInst {{.*}} DebugLocalVariable + +target triple = "spirv64-unknown-unknown" + +define spir_func void @f() !dbg !5 { +entry: + %x = alloca i32, align 4 + #dbg_declare(ptr %x, !9, !DIExpression(DW_OP_LLVM_convert, 32, DW_ATE_signed), !10) + store i32 1, ptr %x, align 4, !dbg !10 + ret void, !dbg !10 +} + +!llvm.dbg.cu = !{!0} +!llvm.module.flags = !{!2, !3} + +!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None) +!1 = !DIFile(filename: "debug-declare-expression-unsupported.c", directory: "/src") +!2 = !{i32 7, !"Dwarf Version", i32 5} +!3 = !{i32 2, !"Debug Info Version", i32 3} +!4 = !DISubroutineType(types: !6) +!6 = !{null} +!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed) +!5 = distinct !DISubprogram(name: "f", linkageName: "f", scope: !1, file: !1, line: 1, type: !4, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0) +!9 = !DILocalVariable(name: "x", scope: !5, file: !1, line: 2, type: !7) +!10 = !DILocation(line: 3, column: 1, scope: !5) diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-expression-xderef.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-expression-xderef.ll new file mode 100644 index 0000000000000..d00075a58bae3 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-expression-xderef.ll @@ -0,0 +1,60 @@ +; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s +; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; The address-space expression clang puts on every declare for a SPIR-V target. + +; CHECK-DAG: [[EXT:%[0-9]+]] = OpExtInstImport "NonSemantic.Shader.DebugInfo.100" +; CHECK-DAG: [[VOID:%[0-9]+]] = OpTypeVoid +; CHECK-DAG: [[I32T:%[0-9]+]] = OpTypeInt 32 0 +; CHECK-DAG: [[X:%[0-9]+]] = OpString "x" +; CHECK-DAG: [[Y:%[0-9]+]] = OpString "y" +; The trailing anchors keep e.g. [[C8]] from binding to "OpConstant %3 80". +; CHECK-DAG: [[C0:%[0-9]+]] = OpConstant [[I32T]] 0{{ *$}} +; CHECK-DAG: [[C5:%[0-9]+]] = OpConstant [[I32T]] 5{{ *$}} +; CHECK-DAG: [[C6:%[0-9]+]] = OpConstant [[I32T]] 6{{ *$}} +; CHECK-DAG: [[C8:%[0-9]+]] = OpConstant [[I32T]] 8{{ *$}} +; CHECK-DAG: [[XVAR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugLocalVariable [[X]] +; CHECK-DAG: [[YVAR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugLocalVariable [[Y]] + +; CHECK-DAG: [[CONSTU:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C8]] [[C0]]{{ *$}} +; CHECK-DAG: [[SWAP:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C5]]{{ *$}} +; CHECK-DAG: [[XDEREF:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C6]]{{ *$}} +; CHECK-DAG: [[EXPR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugExpression [[CONSTU]] [[SWAP]] [[XDEREF]]{{ *$}} + +; CHECK: [[XADDR:%[0-9]+]] = OpVariable {{%[0-9]+}} Function +; CHECK: [[YADDR:%[0-9]+]] = OpVariable {{%[0-9]+}} Function +; CHECK: OpExtInst [[VOID]] [[EXT]] DebugDeclare [[XVAR]] [[XADDR]] [[EXPR]] +; CHECK: OpExtInst [[VOID]] [[EXT]] DebugDeclare [[YVAR]] [[YADDR]] [[EXPR]] + +target triple = "spirv64-unknown-unknown" + +define spir_func i32 @sum(i32 noundef %x) !dbg !5 { +entry: + %x.addr = alloca i32, align 4 + %y = alloca i32, align 4 + store i32 %x, ptr %x.addr, align 4 + #dbg_declare(ptr %x.addr, !9, !DIExpression(DW_OP_constu, 0, DW_OP_swap, DW_OP_xderef), !11) + #dbg_declare(ptr %y, !10, !DIExpression(DW_OP_constu, 0, DW_OP_swap, DW_OP_xderef), !12) + %0 = load i32, ptr %x.addr, align 4, !dbg !12 + store i32 %0, ptr %y, align 4, !dbg !12 + %1 = load i32, ptr %y, align 4, !dbg !12 + ret i32 %1, !dbg !12 +} + +!llvm.dbg.cu = !{!0} +!llvm.module.flags = !{!2, !3} + +!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None) +!1 = !DIFile(filename: "debug-declare-expression-xderef.c", directory: "/src") +!2 = !{i32 7, !"Dwarf Version", i32 5} +!3 = !{i32 2, !"Debug Info Version", i32 3} + +!4 = !DISubroutineType(types: !6) +!6 = !{!7, !7} +!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed) + +!5 = distinct !DISubprogram(name: "sum", linkageName: "sum", scope: !1, file: !1, line: 1, type: !4, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0) +!9 = !DILocalVariable(name: "x", arg: 1, scope: !5, file: !1, line: 1, type: !7) +!10 = !DILocalVariable(name: "y", scope: !5, file: !1, line: 2, type: !7) +!11 = !DILocation(line: 1, column: 13, scope: !5) +!12 = !DILocation(line: 2, column: 7, scope: !5) diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-function-parameter.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-function-parameter.ll new file mode 100644 index 0000000000000..1de2ea6f1947b --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-function-parameter.ll @@ -0,0 +1,39 @@ +; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s --implicit-check-not=DebugDeclare +; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; A declare whose address is an incoming argument rather than an alloca, which +; is the shape clang produces for a byval parameter on this triple. The +; location register is defined by OpFunctionParameter, so the declare is +; dropped as per spec. + +; At the time of writing, spirv-val accepts it, since its rule allows OpVariable or +; OpFunctionParameter, but the spec restricts the Variable operand to +; OpVariable. +; +; spirv64-amd-amdhsa never reaches this shape: its ABI passes the struct byref, +; so the declare lands on the callee's own copy, an alloca. + +; CHECK: OpExtInst {{.*}} DebugLocalVariable + +target triple = "spirv64-unknown-unknown" + +define spir_func void @f(ptr %p) !dbg !5 { +entry: + #dbg_declare(ptr %p, !9, !DIExpression(), !10) + ret void, !dbg !10 +} + +!llvm.dbg.cu = !{!0} +!llvm.module.flags = !{!2, !3} + +!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None) +!1 = !DIFile(filename: "debug-declare-function-parameter.c", directory: "/src") +!2 = !{i32 7, !"Dwarf Version", i32 5} +!3 = !{i32 2, !"Debug Info Version", i32 3} +!4 = !DISubroutineType(types: !6) +!6 = !{null, !8} +!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed) +!8 = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !7, size: 64, dwarfAddressSpace: 4) +!5 = distinct !DISubprogram(name: "f", linkageName: "f", scope: !1, file: !1, line: 1, type: !4, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0) +!9 = !DILocalVariable(name: "p", arg: 1, scope: !5, file: !1, line: 1, type: !8) +!10 = !DILocation(line: 1, column: 20, scope: !5) diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-int-storage.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-int-storage.ll new file mode 100644 index 0000000000000..5dd2010ccfb5b --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-int-storage.ll @@ -0,0 +1,34 @@ +; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s --implicit-check-not=DebugDeclare +; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; A declare on a fixed numeric address, which the IR verifier accepts. The +; address is a constant rather than an OpVariable, and nothing defines the +; location register inside the function, so the declare is dropped. +; +; Clang seems not to emit declares on constants. The verifier permits an +; integer location only because deleting an inttoptr rewrites the declare onto +; the integer itself. + +; CHECK: OpExtInst {{.*}} DebugLocalVariable + +target triple = "spirv64-unknown-unknown" + +define spir_func void @f() !dbg !5 { +entry: + #dbg_declare(ptr inttoptr (i64 4096 to ptr), !9, !DIExpression(), !10) + ret void, !dbg !10 +} + +!llvm.dbg.cu = !{!0} +!llvm.module.flags = !{!2, !3} + +!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None) +!1 = !DIFile(filename: "debug-declare-int-storage.c", directory: "/src") +!2 = !{i32 7, !"Dwarf Version", i32 5} +!3 = !{i32 2, !"Debug Info Version", i32 3} +!4 = !DISubroutineType(types: !6) +!6 = !{null} +!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed) +!5 = distinct !DISubprogram(name: "f", linkageName: "f", scope: !1, file: !1, line: 1, type: !4, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0) +!9 = !DILocalVariable(name: "mapped", scope: !5, file: !1, line: 2, type: !7) +!10 = !DILocation(line: 3, column: 1, scope: !5) diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-line-scope.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-line-scope.ll new file mode 100644 index 0000000000000..eb03ddee8869f --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-line-scope.ll @@ -0,0 +1,98 @@ +; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s +; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; A declare has a location of its own, so it takes part in DebugLine and +; DebugScope tracking like a real instruction. The body interleaves declares +; and instructions whose locations disagree on purpose: +; +; #dbg_declare(ptr %x, !9, !DIExpression(), !20) ; !20 is line 20 +; store i32 1, ptr %x, align 4, !dbg !21 ; !21 is line 5 +; + +; CHECK-DAG: [[EXT:%[0-9]+]] = OpExtInstImport "NonSemantic.Shader.DebugInfo.100" +; CHECK-DAG: [[VOID:%[0-9]+]] = OpTypeVoid +; CHECK-DAG: [[I32:%[0-9]+]] = OpTypeInt 32 0 +; CHECK-DAG: [[PATH:%[0-9]+]] = OpString "{{[/\\]}}src{{[/\\]}}debug-declare-line-scope.c" +; CHECK-DAG: [[X:%[0-9]+]] = OpString "x" +; CHECK-DAG: [[Y:%[0-9]+]] = OpString "y" +; CHECK-DAG: [[DS:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugSource [[PATH]] +; CHECK-DAG: [[DF:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugFunction {{.*}} +; CHECK-DAG: [[LB:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugLexicalBlock [[DS]] {{.*}} [[DF]] +; CHECK-DAG: [[XVAR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugLocalVariable [[X]] +; CHECK-DAG: [[YVAR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugLocalVariable [[Y]] +; CHECK-DAG: [[EXPR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugExpression{{ *$}} +; CHECK-DAG: [[V3:%[0-9]+]] = OpConstant [[I32]] 3{{$}} +; CHECK-DAG: [[V4:%[0-9]+]] = OpConstant [[I32]] 4{{$}} +; CHECK-DAG: [[V5:%[0-9]+]] = OpConstant [[I32]] 5{{$}} +; CHECK-DAG: [[V6:%[0-9]+]] = OpConstant [[I32]] 6{{$}} +; CHECK-DAG: [[V7:%[0-9]+]] = OpConstant [[I32]] 7{{$}} +; CHECK-DAG: [[V8:%[0-9]+]] = OpConstant [[I32]] 8{{$}} +; CHECK-DAG: [[V9:%[0-9]+]] = OpConstant [[I32]] 9{{$}} +; CHECK-DAG: [[V10:%[0-9]+]] = OpConstant [[I32]] 10{{$}} +; CHECK-DAG: [[V20:%[0-9]+]] = OpConstant [[I32]] 20{{$}} +; CHECK-DAG: [[V31:%[0-9]+]] = OpConstant [[I32]] 31{{$}} + +; CHECK: [[FN:%[0-9]+]] = OpFunction +; CHECK-NEXT: OpLabel +; CHECK-NEXT: [[XADDR:%[0-9]+]] = OpVariable {{%[0-9]+}} Function +; CHECK-NEXT: [[YADDR:%[0-9]+]] = OpVariable {{%[0-9]+}} Function +; CHECK-NEXT: OpExtInst [[VOID]] [[EXT]] DebugFunctionDefinition [[DF]] [[FN]] + +; The declare's own location, line 20, not the line 5 of the store below it. +; CHECK-NEXT: OpExtInst [[VOID]] [[EXT]] DebugScope [[DF]] +; CHECK-NEXT: OpExtInst [[VOID]] [[EXT]] DebugLine [[DS]] [[V20]] [[V20]] [[V7]] [[V8]] +; CHECK-NEXT: OpExtInst [[VOID]] [[EXT]] DebugDeclare [[XVAR]] [[XADDR]] [[EXPR]] +; CHECK-NEXT: ;DEBUG_VALUE: + +; CHECK-NEXT: OpExtInst [[VOID]] [[EXT]] DebugLine [[DS]] [[V5]] [[V5]] [[V3]] [[V4]] +; CHECK-NEXT: OpStore [[XADDR]] + +; A declare in a lexical block moves the scope, on the declare alone. +; CHECK-NEXT: OpExtInst [[VOID]] [[EXT]] DebugScope [[LB]] +; CHECK-NEXT: OpExtInst [[VOID]] [[EXT]] DebugLine [[DS]] [[V31]] [[V31]] [[V9]] [[V10]] +; CHECK-NEXT: OpExtInst [[VOID]] [[EXT]] DebugDeclare [[YVAR]] [[YADDR]] [[EXPR]] +; CHECK-NEXT: ;DEBUG_VALUE: + +; The unlocated store drops both, rather than keeping the declare's line. +; CHECK-NEXT: OpExtInst [[VOID]] [[EXT]] DebugNoScope +; CHECK-NEXT: OpExtInst [[VOID]] [[EXT]] DebugNoLine +; CHECK-NEXT: OpStore [[YADDR]] + +; CHECK-NEXT: OpExtInst [[VOID]] [[EXT]] DebugScope [[DF]] +; CHECK-NEXT: OpExtInst [[VOID]] [[EXT]] DebugLine [[DS]] [[V6]] [[V6]] [[V3]] [[V4]] +; CHECK-NEXT: OpReturn +; CHECK-NEXT: OpFunctionEnd + +target triple = "spirv64-unknown-unknown" + +define spir_func void @f() !dbg !5 { +entry: + %x = alloca i32, align 4 + %y = alloca i32, align 4 + #dbg_declare(ptr %x, !9, !DIExpression(), !20) ; function scope, line 20 + store i32 1, ptr %x, align 4, !dbg !21 ; function scope, line 5 + #dbg_declare(ptr %y, !10, !DIExpression(), !22) ; lexical block, line 31 + store i32 2, ptr %y, align 4 ; no debug location + ret void, !dbg !23 ; function scope, line 6 +} + +!llvm.dbg.cu = !{!0} +!llvm.module.flags = !{!2, !3} + +!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None) +!1 = !DIFile(filename: "debug-declare-line-scope.c", directory: "/src") +!2 = !{i32 7, !"Dwarf Version", i32 5} +!3 = !{i32 2, !"Debug Info Version", i32 3} + +!4 = !DISubroutineType(types: !6) +!6 = !{null} +!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed) + +!5 = distinct !DISubprogram(name: "f", linkageName: "f", scope: !1, file: !1, line: 1, type: !4, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0) +!8 = distinct !DILexicalBlock(scope: !5, file: !1, line: 30, column: 3) +!9 = !DILocalVariable(name: "x", scope: !5, file: !1, line: 20, type: !7) +!10 = !DILocalVariable(name: "y", scope: !8, file: !1, line: 31, type: !7) +!20 = !DILocation(line: 20, column: 7, scope: !5) +!21 = !DILocation(line: 5, column: 3, scope: !5) +!22 = !DILocation(line: 31, column: 9, scope: !8) +!23 = !DILocation(line: 6, column: 3, scope: !5) diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-module-scope-variable.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-module-scope-variable.ll new file mode 100644 index 0000000000000..710f3ef0221f5 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-module-scope-variable.ll @@ -0,0 +1,45 @@ +; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s +; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; A #dbg_declare whose address is a module-scope global rather than an alloca. +; The G_GLOBAL_VALUE is selected to the Workgroup OpVariable, so the storage +; operand resolves and the declare is emitted from inside the function. +; +; This is a synthetic case, clang doesn't seem to emit this. +; clang describes an OpenCL __local or a HIP __shared__ variable with a +; DIGlobalVariable, which takes the DebugGlobalVariable path instead. + +; CHECK-DAG: [[EXT:%[0-9]+]] = OpExtInstImport "NonSemantic.Shader.DebugInfo.100" +; CHECK-DAG: [[VOID:%[0-9]+]] = OpTypeVoid +; CHECK-DAG: [[NAME:%[0-9]+]] = OpString "shared" +; CHECK-DAG: [[STORAGE:%[0-9]+]] = OpVariable {{%[0-9]+}} Workgroup +; CHECK-DAG: [[VAR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugLocalVariable [[NAME]] +; CHECK-DAG: [[EXPR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugExpression{{ *$}} + +; CHECK: OpExtInst [[VOID]] [[EXT]] DebugFunctionDefinition +; CHECK: OpExtInst [[VOID]] [[EXT]] DebugDeclare [[VAR]] [[STORAGE]] [[EXPR]] + +target triple = "spirv64-unknown-unknown" + +@shared = internal addrspace(3) global i32 undef, align 4 + +define spir_func void @k() !dbg !5 { +entry: + #dbg_declare(ptr addrspace(3) @shared, !9, !DIExpression(), !10) + store i32 1, ptr addrspace(3) @shared, align 4, !dbg !10 + ret void, !dbg !10 +} + +!llvm.dbg.cu = !{!0} +!llvm.module.flags = !{!2, !3} + +!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None) +!1 = !DIFile(filename: "debug-declare-module-scope-variable.c", directory: "/src") +!2 = !{i32 7, !"Dwarf Version", i32 5} +!3 = !{i32 2, !"Debug Info Version", i32 3} +!4 = !DISubroutineType(types: !6) +!6 = !{null} +!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed) +!5 = distinct !DISubprogram(name: "k", linkageName: "k", scope: !1, file: !1, line: 1, type: !4, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0) +!9 = !DILocalVariable(name: "shared", scope: !5, file: !1, line: 2, type: !7) +!10 = !DILocation(line: 3, column: 3, scope: !5) diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-null-storage.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-null-storage.ll new file mode 100644 index 0000000000000..615b35e98408a --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-null-storage.ll @@ -0,0 +1,30 @@ +; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s --implicit-check-not=DebugDeclare +; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; A declare on a null address. The pointer is a module-scope OpConstantNull, so +; nothing defines the location register inside the function and the declare is +; dropped. The register still has a SPIR-V id. + +; CHECK: OpExtInst {{.*}} DebugLocalVariable + +target triple = "spirv64-unknown-unknown" + +define spir_func void @f() !dbg !5 { +entry: + #dbg_declare(ptr null, !9, !DIExpression(), !10) + ret void, !dbg !10 +} + +!llvm.dbg.cu = !{!0} +!llvm.module.flags = !{!2, !3} + +!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None) +!1 = !DIFile(filename: "debug-declare-null-storage.c", directory: "/src") +!2 = !{i32 7, !"Dwarf Version", i32 5} +!3 = !{i32 2, !"Debug Info Version", i32 3} +!4 = !DISubroutineType(types: !6) +!6 = !{null} +!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed) +!5 = distinct !DISubprogram(name: "f", linkageName: "f", scope: !1, file: !1, line: 1, type: !4, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0) +!9 = !DILocalVariable(name: "x", scope: !5, file: !1, line: 2, type: !7) +!10 = !DILocation(line: 3, column: 1, scope: !5) diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-poison-storage.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-poison-storage.ll new file mode 100644 index 0000000000000..c5961ec21a697 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-poison-storage.ll @@ -0,0 +1,28 @@ +; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s --implicit-check-not=DebugDeclare --implicit-check-not=DebugExpression +; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; A declare on poison, which is how a dropped variable location is treated. + +; CHECK: OpExtInst {{.*}} DebugLocalVariable + +target triple = "spirv64-unknown-unknown" + +define spir_func void @f() !dbg !5 { +entry: + #dbg_declare(ptr poison, !9, !DIExpression(), !10) + ret void, !dbg !10 +} + +!llvm.dbg.cu = !{!0} +!llvm.module.flags = !{!2, !3} + +!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None) +!1 = !DIFile(filename: "debug-declare-poison-storage.c", directory: "/src") +!2 = !{i32 7, !"Dwarf Version", i32 5} +!3 = !{i32 2, !"Debug Info Version", i32 3} +!4 = !DISubroutineType(types: !6) +!6 = !{null} +!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed) +!5 = distinct !DISubprogram(name: "f", linkageName: "f", scope: !1, file: !1, line: 1, type: !4, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0) +!9 = !DILocalVariable(name: "dropped", scope: !5, file: !1, line: 2, type: !7) +!10 = !DILocation(line: 3, column: 1, scope: !5) diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-skip-type.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-skip-type.ll new file mode 100644 index 0000000000000..bee90184131ef --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-skip-type.ll @@ -0,0 +1,33 @@ +; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s --implicit-check-not=DebugLocalVariable --implicit-check-not=DebugDeclare +; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; A declare on live storage, so the OpVariable is there, but the variable's +; type has no DWARF address space and its DebugTypePointer is skipped. + +; CHECK: OpVariable {{%[0-9]+}} Function +; CHECK: OpExtInst {{.*}} DebugFunctionDefinition + +target triple = "spirv64-unknown-unknown" + +define spir_func void @f() !dbg !5 { +entry: + %p = alloca ptr, align 8 + #dbg_declare(ptr %p, !9, !DIExpression(), !10) + store ptr null, ptr %p, align 8, !dbg !10 + ret void, !dbg !10 +} + +!llvm.dbg.cu = !{!0} +!llvm.module.flags = !{!2, !3} + +!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None) +!1 = !DIFile(filename: "debug-declare-skip-type.c", directory: "/src") +!2 = !{i32 7, !"Dwarf Version", i32 5} +!3 = !{i32 2, !"Debug Info Version", i32 3} +!4 = !DISubroutineType(types: !6) +!6 = !{null} +!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed) +!8 = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !7, size: 64) +!5 = distinct !DISubprogram(name: "f", linkageName: "f", scope: !1, file: !1, line: 1, type: !4, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0) +!9 = !DILocalVariable(name: "p", scope: !5, file: !1, line: 2, type: !8) +!10 = !DILocation(line: 3, column: 1, scope: !5) diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare.ll new file mode 100644 index 0000000000000..fd8fa157c2466 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare.ll @@ -0,0 +1,55 @@ +; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s +; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; A #dbg_declare on an alloca that survives to MIR: the parameter copy and a +; local. Both use an empty DIExpression, so a single DebugExpression with no +; operations is shared by both declares. + +; CHECK-DAG: [[EXT:%[0-9]+]] = OpExtInstImport "NonSemantic.Shader.DebugInfo.100" +; CHECK-DAG: [[VOID:%[0-9]+]] = OpTypeVoid +; CHECK-DAG: [[VALUE:%[0-9]+]] = OpString "value" +; CHECK-DAG: [[RESULT:%[0-9]+]] = OpString "result" +; CHECK-DAG: [[VALUEVAR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugLocalVariable [[VALUE]] +; CHECK-DAG: [[RESULTVAR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugLocalVariable [[RESULT]] +; An empty DIExpression lowers to a DebugExpression with no operands. +; CHECK-DAG: [[EXPR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugExpression{{ *$}} + +; CHECK: [[VALUEADDR:%[0-9]+]] = OpVariable {{%[0-9]+}} Function +; CHECK: [[RESULTADDR:%[0-9]+]] = OpVariable {{%[0-9]+}} Function +; CHECK: OpExtInst [[VOID]] [[EXT]] DebugFunctionDefinition +; CHECK: OpExtInst [[VOID]] [[EXT]] DebugDeclare [[VALUEVAR]] [[VALUEADDR]] [[EXPR]] +; CHECK: OpExtInst [[VOID]] [[EXT]] DebugDeclare [[RESULTVAR]] [[RESULTADDR]] [[EXPR]] + +target triple = "spirv64-unknown-unknown" + +define spir_func i32 @add_one(i32 %value) !dbg !5 { +entry: + %value.addr = alloca i32, align 4 + %result = alloca i32, align 4 + store i32 %value, ptr %value.addr, align 4 + #dbg_declare(ptr %value.addr, !9, !DIExpression(), !11) + %0 = load i32, ptr %value.addr, align 4, !dbg !11 + %add = add nsw i32 %0, %0, !dbg !11 + store i32 %add, ptr %result, align 4, !dbg !11 + #dbg_declare(ptr %result, !10, !DIExpression(), !12) + %1 = load i32, ptr %result, align 4, !dbg !12 + ret i32 %1, !dbg !12 +} + +!llvm.dbg.cu = !{!0} +!llvm.module.flags = !{!2, !3} + +!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None) +!1 = !DIFile(filename: "debug-declare.c", directory: "/src") +!2 = !{i32 7, !"Dwarf Version", i32 5} +!3 = !{i32 2, !"Debug Info Version", i32 3} + +!4 = !DISubroutineType(types: !6) +!6 = !{!7, !7} +!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed) + +!5 = distinct !DISubprogram(name: "add_one", linkageName: "add_one", scope: !1, file: !1, line: 7, type: !4, scopeLine: 7, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0) +!9 = !DILocalVariable(name: "value", arg: 1, scope: !5, file: !1, line: 7, type: !7) +!10 = !DILocalVariable(name: "result", scope: !5, file: !1, line: 11, type: !7) +!11 = !DILocation(line: 7, column: 20, scope: !5) +!12 = !DILocation(line: 11, column: 30, scope: !5) diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-expression-bit-piece.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-expression-bit-piece.ll new file mode 100644 index 0000000000000..cfda50da35b0c --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-expression-bit-piece.ll @@ -0,0 +1,40 @@ +; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o /dev/null 2>&1 | FileCheck %s --check-prefix=VERIFY +; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s --implicit-check-not=Debug +; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; DW_OP_bit_piece is the DWARF counterpart of NonSemantic.Shader.DebugInfo.100 +; BitPiece (encoding 4). The IR verifier currently rejects it. +; If rejected, the debug information is dropped, and llc continues the execution. + +; Future implementation must be careful with the following: +; DW_OP_bit_piece is (size, offset); NSDI BitPiece is (offset, size). + +; VERIFY: invalid expression +; VERIFY: !DIExpression(157, 32, 8) +; VERIFY: warning: ignoring invalid debug info + +; CHECK: OpFunction + +target triple = "spirv64-unknown-unknown" + +define spir_func void @f() !dbg !5 { +entry: + %x = alloca i32, align 4 + #dbg_declare(ptr %x, !9, !DIExpression(DW_OP_bit_piece, 32, 8), !10) + store i32 1, ptr %x, align 4, !dbg !10 + ret void, !dbg !10 +} + +!llvm.dbg.cu = !{!0} +!llvm.module.flags = !{!2, !3} + +!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None) +!1 = !DIFile(filename: "debug-expression-bit-piece.c", directory: "/src") +!2 = !{i32 7, !"Dwarf Version", i32 5} +!3 = !{i32 2, !"Debug Info Version", i32 3} +!4 = !DISubroutineType(types: !6) +!6 = !{null} +!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed) +!5 = distinct !DISubprogram(name: "f", linkageName: "f", scope: !1, file: !1, line: 1, type: !4, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0) +!9 = !DILocalVariable(name: "x", scope: !5, file: !1, line: 2, type: !7) +!10 = !DILocation(line: 3, column: 1, scope: !5) diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-expression-operations.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-expression-operations.ll new file mode 100644 index 0000000000000..84e5c766afc41 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-expression-operations.ll @@ -0,0 +1,54 @@ +; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s +; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; One expression using the operations no other test reaches: Deref (0), +; Plus (1), Minus (2), PlusUconst (3), StackValue (7) and Fragment (9). + +; CHECK-DAG: [[EXT:%[0-9]+]] = OpExtInstImport "NonSemantic.Shader.DebugInfo.100" +; CHECK-DAG: [[VOID:%[0-9]+]] = OpTypeVoid +; CHECK-DAG: [[I32T:%[0-9]+]] = OpTypeInt 32 0 +; CHECK-DAG: [[C0:%[0-9]+]] = OpConstant [[I32T]] 0{{ *$}} +; CHECK-DAG: [[C1:%[0-9]+]] = OpConstant [[I32T]] 1{{ *$}} +; CHECK-DAG: [[C2:%[0-9]+]] = OpConstant [[I32T]] 2{{ *$}} +; CHECK-DAG: [[C3:%[0-9]+]] = OpConstant [[I32T]] 3{{ *$}} +; CHECK-DAG: [[C4:%[0-9]+]] = OpConstant [[I32T]] 4{{ *$}} +; CHECK-DAG: [[C7:%[0-9]+]] = OpConstant [[I32T]] 7{{ *$}} +; CHECK-DAG: [[C8:%[0-9]+]] = OpConstant [[I32T]] 8{{ *$}} +; CHECK-DAG: [[C9:%[0-9]+]] = OpConstant [[I32T]] 9{{ *$}} + +; CHECK-DAG: [[DEREF:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C0]]{{ *$}} +; CHECK-DAG: [[PLUSU:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C3]] [[C4]]{{ *$}} +; CHECK-DAG: [[CONSTU8:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C8]] [[C8]]{{ *$}} +; CHECK-DAG: [[MINUS:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C2]]{{ *$}} +; CHECK-DAG: [[CONSTU2:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C8]] [[C2]]{{ *$}} +; CHECK-DAG: [[PLUS:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C1]]{{ *$}} +; CHECK-DAG: [[STACK:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C7]]{{ *$}} +; CHECK-DAG: [[FRAG:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C9]] [[C0]] [[C8]]{{ *$}} +; CHECK-DAG: [[EXPR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugExpression [[DEREF]] [[PLUSU]] [[CONSTU8]] [[MINUS]] [[CONSTU2]] [[PLUS]] [[STACK]] [[FRAG]]{{ *$}} + +; CHECK: OpExtInst [[VOID]] [[EXT]] DebugDeclare {{%[0-9]+}} {{%[0-9]+}} [[EXPR]] + +target triple = "spirv64-unknown-unknown" + +define spir_func void @f() !dbg !5 { +entry: + %x = alloca i32, align 4 + #dbg_declare(ptr %x, !9, !DIExpression(DW_OP_deref, DW_OP_plus_uconst, 4, DW_OP_constu, 8, DW_OP_minus, DW_OP_constu, 2, DW_OP_plus, DW_OP_stack_value, DW_OP_LLVM_fragment, 0, 8), !10) + ; 99 keeps this constant clear of the operation encodings captured above. + store i32 99, ptr %x, align 4, !dbg !10 + ret void, !dbg !10 +} + +!llvm.dbg.cu = !{!0} +!llvm.module.flags = !{!2, !3} + +!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None) +!1 = !DIFile(filename: "debug-expression-operations.c", directory: "/src") +!2 = !{i32 7, !"Dwarf Version", i32 5} +!3 = !{i32 2, !"Debug Info Version", i32 3} +!4 = !DISubroutineType(types: !6) +!6 = !{null} +!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed) +!5 = distinct !DISubprogram(name: "f", linkageName: "f", scope: !1, file: !1, line: 1, type: !4, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0) +!9 = !DILocalVariable(name: "x", scope: !5, file: !1, line: 2, type: !7) +!10 = !DILocation(line: 3, column: 1, scope: !5) diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-expression-out-of-range.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-expression-out-of-range.ll new file mode 100644 index 0000000000000..108b1f7338084 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-expression-out-of-range.ll @@ -0,0 +1,76 @@ +; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s +; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s --check-prefix=ONE --implicit-check-not=DebugOperation --implicit-check-not=DebugExpression --implicit-check-not=DebugDeclare +; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; NonSemantic sets take no literals, so every DIExpression argument becomes a +; 32-bit OpConstant. An argument that does not fit is dropped. + +; CHECK-DAG: [[EXT:%[0-9]+]] = OpExtInstImport "NonSemantic.Shader.DebugInfo.100" +; CHECK-DAG: [[VOID:%[0-9]+]] = OpTypeVoid +; CHECK-DAG: [[I32T:%[0-9]+]] = OpTypeInt 32 0 +; CHECK-DAG: [[OK:%[0-9]+]] = OpString "ok" +; CHECK-DAG: [[BIG:%[0-9]+]] = OpString "big" +; CHECK-DAG: [[GBIG:%[0-9]+]] = OpString "gbig" +; The trailing anchors keep e.g. [[C4]] from binding to "OpConstant %3 40". +; CHECK-DAG: [[C0:%[0-9]+]] = OpConstant [[I32T]] 0{{ *$}} +; CHECK-DAG: [[C4:%[0-9]+]] = OpConstant [[I32T]] 4{{ *$}} +; CHECK-DAG: [[C5:%[0-9]+]] = OpConstant [[I32T]] 5{{ *$}} +; CHECK-DAG: [[C6:%[0-9]+]] = OpConstant [[I32T]] 6{{ *$}} +; CHECK-DAG: [[C8:%[0-9]+]] = OpConstant [[I32T]] 8{{ *$}} +; CHECK-DAG: [[NONE:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugInfoNone + +; Both variables are collected and emitted; only the expressions differ. +; CHECK-DAG: [[OKVAR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugLocalVariable [[OK]] +; CHECK-DAG: [[BIGVAR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugLocalVariable [[BIG]] + +; CHECK-DAG: [[CONSTU:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C8]] [[C4]]{{ *$}} +; CHECK-DAG: [[SWAP:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C5]]{{ *$}} +; CHECK-DAG: [[XDEREF:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C6]]{{ *$}} +; CHECK-DAG: [[EXPR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugExpression [[CONSTU]] [[SWAP]] [[XDEREF]]{{ *$}} + +; CHECK-DAG: OpExtInst [[VOID]] [[EXT]] DebugGlobalVariable [[GBIG]] {{.*}} [[GBIG]] [[NONE]] + +; Both allocas still become OpVariable, in declaration order; only the declare +; for the first one can be emitted. +; CHECK: [[OKADDR:%[0-9]+]] = OpVariable {{%[0-9]+}} Function +; CHECK: [[BIGADDR:%[0-9]+]] = OpVariable {{%[0-9]+}} Function +; CHECK: OpExtInst [[VOID]] [[EXT]] DebugDeclare [[OKVAR]] [[OKADDR]] [[EXPR]] + +; ONE: OpExtInst {{.*}} DebugOperation +; ONE: OpExtInst {{.*}} DebugOperation +; ONE: OpExtInst {{.*}} DebugOperation +; ONE: OpExtInst {{.*}} DebugExpression +; ONE: OpExtInst {{.*}} DebugDeclare + +target triple = "spirv64-unknown-unknown" + +define spir_func i32 @f(i32 noundef %x) !dbg !10 { +entry: + %ok = alloca i32, align 4 + %big = alloca i32, align 4 + store i32 %x, ptr %ok, align 4 + #dbg_declare(ptr %ok, !11, !DIExpression(DW_OP_constu, 4, DW_OP_swap, DW_OP_xderef), !13) + #dbg_declare(ptr %big, !12, !DIExpression(DW_OP_constu, 4294967296, DW_OP_swap, DW_OP_xderef), !13) + %0 = load i32, ptr %ok, align 4, !dbg !13 + store i32 %0, ptr %big, align 4, !dbg !13 + %1 = load i32, ptr %big, align 4, !dbg !13 + ret i32 %1, !dbg !13 +} + +!llvm.dbg.cu = !{!0} +!llvm.module.flags = !{!1, !2} + +!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !3, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, globals: !4, splitDebugInlining: false, nameTableKind: None) +!1 = !{i32 7, !"Dwarf Version", i32 5} +!2 = !{i32 2, !"Debug Info Version", i32 3} +!3 = !DIFile(filename: "t.c", directory: "/src") +!4 = !{!5} +!5 = !DIGlobalVariableExpression(var: !6, expr: !DIExpression(DW_OP_constu, 4294967296, DW_OP_swap, DW_OP_xderef)) +!6 = distinct !DIGlobalVariable(name: "gbig", linkageName: "gbig", scope: !0, file: !3, line: 1, type: !7, isLocal: false, isDefinition: true) +!7 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed) +!8 = !DISubroutineType(cc: DW_CC_LLVM_SpirFunction, types: !9) +!9 = !{!7, !7} +!10 = distinct !DISubprogram(name: "f", scope: !3, file: !3, line: 3, type: !8, scopeLine: 3, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0) +!11 = !DILocalVariable(name: "ok", scope: !10, file: !3, line: 4, type: !7) +!12 = !DILocalVariable(name: "big", scope: !10, file: !3, line: 5, type: !7) +!13 = !DILocation(line: 4, column: 7, scope: !10) diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-global-variable-constant-value.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-global-variable-constant-value.ll new file mode 100644 index 0000000000000..b5f73c39bedb4 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-global-variable-constant-value.ll @@ -0,0 +1,76 @@ +; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s +; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s --check-prefix=ONE --implicit-check-not=DebugOperation --implicit-check-not=DebugExpression +; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | not spirv-val 2>&1 | FileCheck %s --check-prefix=VAL %} + +; This HIP code compiles to the metadata below with no optimisations as it is the case for amdgcnspirv: +; +; constexpr unsigned long long Hash = 0xff51afd7ed558ccdULL; +; constexpr double Scale = 1.5; +; constexpr unsigned Small = 42u; +; + +; DebugOperation requires int32 operands and DI for Hash and Scale cannot be encoded. + +; CHECK-DAG: [[EXT:%[0-9]+]] = OpExtInstImport "NonSemantic.Shader.DebugInfo.100" +; CHECK-DAG: [[VOID:%[0-9]+]] = OpTypeVoid +; CHECK-DAG: [[I32T:%[0-9]+]] = OpTypeInt 32 0 +; CHECK-DAG: [[HASH:%[0-9]+]] = OpString "Hash" +; CHECK-DAG: [[SCALE:%[0-9]+]] = OpString "Scale" +; CHECK-DAG: [[SMALL:%[0-9]+]] = OpString "Small" +; The trailing anchors keep e.g. [[C7]] from binding to "OpConstant %3 72". +; CHECK-DAG: [[C7:%[0-9]+]] = OpConstant [[I32T]] 7{{ *$}} +; CHECK-DAG: [[C8:%[0-9]+]] = OpConstant [[I32T]] 8{{ *$}} +; CHECK-DAG: [[C42:%[0-9]+]] = OpConstant [[I32T]] 42{{ *$}} +; CHECK-DAG: [[NONE:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugInfoNone + +; CHECK-DAG: [[CONSTU:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C8]] [[C42]]{{ *$}} +; CHECK-DAG: [[SV:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C7]]{{ *$}} +; CHECK-DAG: [[EXPR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugExpression [[CONSTU]] [[SV]]{{ *$}} +; CHECK-DAG: OpExtInst [[VOID]] [[EXT]] DebugGlobalVariable [[SMALL]] {{.*}} [[SMALL]] [[EXPR]] + +; CHECK-DAG: OpExtInst [[VOID]] [[EXT]] DebugGlobalVariable [[HASH]] {{.*}} [[HASH]] [[NONE]] +; CHECK-DAG: OpExtInst [[VOID]] [[EXT]] DebugGlobalVariable [[SCALE]] {{.*}} [[SCALE]] [[NONE]] + +; Small's two operations and one expression are the only ones emitted. +; ONE: OpExtInst {{.*}} DebugOperation +; ONE: OpExtInst {{.*}} DebugOperation +; ONE: OpExtInst {{.*}} DebugExpression + +; spirv-val is broken and is rejecting an expression in a DebugGlobalVariable (KhronosGroup/SPIRV-Tools#6469). +; VAL: DebugGlobalVariable: expected operand Variable must be a result id of + +target triple = "spirv64-unknown-unknown" + +define spir_func void @f() !dbg !20 { +entry: + ret void, !dbg !21 +} + +!llvm.dbg.cu = !{!0} +!llvm.module.flags = !{!1, !2} + +!0 = distinct !DICompileUnit(language: DW_LANG_HIP, file: !3, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, globals: !4, splitDebugInlining: false, nameTableKind: None) +!1 = !{i32 7, !"Dwarf Version", i32 5} +!2 = !{i32 2, !"Debug Info Version", i32 3} +!3 = !DIFile(filename: "fe.hip", directory: "/src") +!4 = !{!5, !8, !11} + +; constexpr unsigned long long Hash = 0xff51afd7ed558ccdULL; +!5 = !DIGlobalVariableExpression(var: !6, expr: !DIExpression(DW_OP_constu, 18397679294719823053, DW_OP_stack_value)) +!6 = distinct !DIGlobalVariable(name: "Hash", linkageName: "Hash", scope: !0, file: !3, line: 2, type: !7, isLocal: true, isDefinition: true) +!7 = !DIBasicType(name: "unsigned long long", size: 64, encoding: DW_ATE_unsigned) + +; constexpr double Scale = 1.5; (0x3FF8000000000000) +!8 = !DIGlobalVariableExpression(var: !9, expr: !DIExpression(DW_OP_constu, 4609434218613702656, DW_OP_stack_value)) +!9 = distinct !DIGlobalVariable(name: "Scale", linkageName: "Scale", scope: !0, file: !3, line: 3, type: !10, isLocal: true, isDefinition: true) +!10 = !DIBasicType(name: "double", size: 64, encoding: DW_ATE_float) + +; constexpr unsigned Small = 42u; +!11 = !DIGlobalVariableExpression(var: !12, expr: !DIExpression(DW_OP_constu, 42, DW_OP_stack_value)) +!12 = distinct !DIGlobalVariable(name: "Small", linkageName: "Small", scope: !0, file: !3, line: 4, type: !13, isLocal: true, isDefinition: true) +!13 = !DIBasicType(name: "unsigned int", size: 32, encoding: DW_ATE_unsigned) + +!18 = !DISubroutineType(cc: DW_CC_LLVM_SpirFunction, types: !19) +!19 = !{null} +!20 = distinct !DISubprogram(name: "f", scope: !3, file: !3, line: 6, type: !18, scopeLine: 6, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0) +!21 = !DILocation(line: 7, column: 1, scope: !20) diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-global-variable-init-expr.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-global-variable-init-expr.ll index fb83a6350ce03..a62fe822555d6 100644 --- a/llvm/test/CodeGen/SPIRV/debug-info/debug-global-variable-init-expr.ll +++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-global-variable-init-expr.ll @@ -1,22 +1,48 @@ ; RUN: llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s -; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} +; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | not spirv-val 2>&1 | FileCheck %s --check-prefix=VAL %} ; A DIGlobalVariable with no backing llvm::GlobalVariable but whose ; DIGlobalVariableExpression carries a non-empty DIExpression (a constant -; initializer). Since DebugExpression emission is not implemented, the Variable -; operand falls back to DebugInfoNone. Flags encode IsLocal|IsDefinition (12). +; initializer). Both operations map, so the Variable operand is the resulting +; DebugExpression rather than DebugInfoNone: Constu (8) carrying 42, then +; StackValue (7) to say the operand stack holds the value itself. +; Flags encode IsLocal|IsDefinition (12). +; +; spirv-val rejects that, so the second RUN line expects it to fail. The +; extension permits it: "If the variable is optimized out, this operand can be +; the <id> of a DebugExpression instruction that contains the constant value of +; the variable that was optimized out." The validator instead shares one rule +; with OpenCL.DebugInfo.100, whose wording stops at DebugInfoNone, and checks +; the operand against a fixed list of OpVariable and constant opcodes. +; +; This is KhronosGroup/SPIRV-Tools#6469, open, where the maintainers agree the +; validator is at fault twice over: it should accept a DebugExpression for +; NonSemantic.Shader.DebugInfo.100, and it should not be accepting the OpConstant +; variants that only the OpenCL wording allows. Every version tried rejects it, +; from 2022.2 to 2026.2, and SPIRV-LLVM-Translator emits the same thing, so this +; is not a stale-binary problem. When the fix lands this RUN line will start +; failing, which is the signal to restore a plain spirv-val invocation. ; CHECK-DAG: [[EXT:%[0-9]+]] = OpExtInstImport "NonSemantic.Shader.DebugInfo.100" ; CHECK-DAG: [[VOID:%[0-9]+]] = OpTypeVoid ; CHECK-DAG: [[I32T:%[0-9]+]] = OpTypeInt 32 0 ; CHECK-DAG: [[NAME:%[0-9]+]] = OpString "constg" ; CHECK-DAG: [[STR_INT:%[0-9]+]] = OpString "int" -; CHECK-DAG: [[C42:%[0-9]+]] = OpConstant [[I32T]] 42 -; CHECK-DAG: [[C12:%[0-9]+]] = OpConstant [[I32T]] 12 -; CHECK-DAG: [[NONE:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugInfoNone +; The trailing anchors keep e.g. [[C7]] from binding to "OpConstant %3 72". +; CHECK-DAG: [[C7:%[0-9]+]] = OpConstant [[I32T]] 7{{ *$}} +; CHECK-DAG: [[C8:%[0-9]+]] = OpConstant [[I32T]] 8{{ *$}} +; CHECK-DAG: [[C12:%[0-9]+]] = OpConstant [[I32T]] 12{{ *$}} +; CHECK-DAG: [[C42:%[0-9]+]] = OpConstant [[I32T]] 42{{ *$}} ; CHECK-DAG: [[DS:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugSource ; CHECK-DAG: [[DTI:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugTypeBasic [[STR_INT]] -; CHECK-DAG: OpExtInst [[VOID]] [[EXT]] DebugGlobalVariable [[NAME]] [[DTI]] [[DS]] [[C42]] {{%[0-9]+}} {{%[0-9]+}} [[NAME]] [[NONE]] [[C12]] +; CHECK-DAG: [[CONSTU:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C8]] [[C42]]{{ *$}} +; CHECK-DAG: [[SV:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugOperation [[C7]]{{ *$}} +; CHECK-DAG: [[EXPR:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugExpression [[CONSTU]] [[SV]]{{ *$}} +; CHECK-DAG: OpExtInst [[VOID]] [[EXT]] DebugGlobalVariable [[NAME]] [[DTI]] [[DS]] [[C42]] {{%[0-9]+}} {{%[0-9]+}} [[NAME]] [[EXPR]] [[C12]] + +; The operand list the validator will accept has varied across releases, so +; match only up to it. +; VAL: DebugGlobalVariable: expected operand Variable must be a result id of target triple = "spirv64-unknown-unknown" diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-local-variable-dbg-value.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-local-variable-dbg-value.ll index 83483ac8d5805..91699d80cbe54 100644 --- a/llvm/test/CodeGen/SPIRV/debug-info/debug-local-variable-dbg-value.ll +++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-local-variable-dbg-value.ll @@ -1,7 +1,10 @@ -; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s +; RUN: llc --verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_non_semantic_info %s -o - | FileCheck %s --implicit-check-not=DebugDeclare ; RUN: %if spirv-tools %{ llc --verify-machineinstrs --spirv-ext=+SPV_KHR_non_semantic_info -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} ; Collect DILocalVariable from a #dbg_value record. No retainedNodes. +; +; The record is a plain #dbg_value, so it lowers to a direct DBG_VALUE and gets +; no DebugDeclare. ; CHECK-DAG: [[EXT:%[0-9]+]] = OpExtInstImport "NonSemantic.Shader.DebugInfo.100" ; CHECK-DAG: [[VOID:%[0-9]+]] = OpTypeVoid @@ -16,7 +19,7 @@ ; CHECK-DAG: [[DS:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugSource [[PATH]] ; CHECK-DAG: [[INT:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugTypeBasic [[INTNAME]] {{.*}} [[C0]] ; CHECK-DAG: [[DF:%[0-9]+]] = OpExtInst [[VOID]] [[EXT]] DebugFunction {{.*}} -; CHECK-DAG: OpExtInst [[VOID]] [[EXT]] DebugLocalVariable [[XNAME]] [[INT]] [[DS]] [[C8]] [[C0]] [[DF]] [[C0]] [[C1]] +; CHECK: OpExtInst [[VOID]] [[EXT]] DebugLocalVariable [[XNAME]] [[INT]] [[DS]] [[C8]] [[C0]] [[DF]] [[C0]] [[C1]] target triple = "spirv64-unknown-unknown" >From bab4f38c78425872c58c993ead7d1abd4e771af4 Mon Sep 17 00:00:00 2001 From: Manuel Carrasco <[email protected]> Date: Tue, 8 Sep 2026 09:43:42 -0500 Subject: [PATCH 2/2] Fix linter. --- .../SPIRV/debug-info/debug-declare-module-scope-variable.ll | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-module-scope-variable.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-module-scope-variable.ll index 710f3ef0221f5..a8c33b59d29e0 100644 --- a/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-module-scope-variable.ll +++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-declare-module-scope-variable.ll @@ -21,7 +21,7 @@ target triple = "spirv64-unknown-unknown" -@shared = internal addrspace(3) global i32 undef, align 4 +@shared = internal addrspace(3) global i32 poison, align 4 define spir_func void @k() !dbg !5 { entry: _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
