[llvm] 0012dea - [AsmPrinter][debug] Materialize abstract lexical blocks referenced by global vars (#193986)
via llvm-commits
llvm-commits at lists.llvm.org
Tue Apr 28 04:21:17 PDT 2026
Author: Matt Davis
Date: 2026-04-28T07:21:12-04:00
New Revision: 0012deaa1b90f7dea5b2f05c7c92902291e46583
URL: https://github.com/llvm/llvm-project/commit/0012deaa1b90f7dea5b2f05c7c92902291e46583
DIFF: https://github.com/llvm/llvm-project/commit/0012deaa1b90f7dea5b2f05c7c92902291e46583.diff
LOG: [AsmPrinter][debug] Materialize abstract lexical blocks referenced by global vars (#193986)
After 63074da2 moved global variable emission from beginModule() to
endModule(), getOrCreateContextDIE() is called after abstract scope DIEs
have been built. However, skipLexicalScope() may have elided an abstract
lexical block that has no local variables. When a global variable (e.g.,
a function-scope static) is scoped to such a block,
getOrCreateContextDIE() returns nullptr causing an assertion. This PR
keeps abstract lexical blocks when they are the scope target of a global
variable in the compile unit.
Derived via LLM usage.
Added:
llvm/test/DebugInfo/Generic/global-var-in-abstract-lexical-block.ll
Modified:
llvm/lib/CodeGen/AsmPrinter/DwarfCompileUnit.cpp
llvm/lib/CodeGen/AsmPrinter/DwarfCompileUnit.h
Removed:
################################################################################
diff --git a/llvm/lib/CodeGen/AsmPrinter/DwarfCompileUnit.cpp b/llvm/lib/CodeGen/AsmPrinter/DwarfCompileUnit.cpp
index d7ed467809225..aad8f3bce0d6a 100644
--- a/llvm/lib/CodeGen/AsmPrinter/DwarfCompileUnit.cpp
+++ b/llvm/lib/CodeGen/AsmPrinter/DwarfCompileUnit.cpp
@@ -81,6 +81,10 @@ DwarfCompileUnit::DwarfCompileUnit(unsigned UID, const DICompileUnit *Node,
: DwarfUnit(GetCompileUnitType(Kind, DW), Node, A, DW, DWU, UID) {
insertDIE(Node, &getUnitDie());
MacroLabelBegin = Asm->createTempSymbol("cu_macro_begin");
+ assert(CUNode);
+ for (auto *GVE : CUNode->getGlobalVariables())
+ if (auto *GV = GVE->getVariable())
+ GlobalVarScopes.insert(GV->getScope());
}
/// addLabelAddress - Add a dwarf label attribute data and value using
@@ -1222,6 +1226,10 @@ DIE &DwarfCompileUnit::constructSubprogramScopeDIE(const DISubprogram *Sub,
return ScopeDIE;
}
+bool DwarfCompileUnit::hasGlobalVariableInScope(const DILocalScope *ScopeNode) {
+ return GlobalVarScopes.contains(ScopeNode);
+}
+
DIE *DwarfCompileUnit::createAndAddScopeChildren(LexicalScope *Scope,
DIE &ScopeDIE) {
DIE *ObjectPointer = nullptr;
@@ -1252,6 +1260,12 @@ DIE *DwarfCompileUnit::createAndAddScopeChildren(LexicalScope *Scope,
auto skipLexicalScope = [this](LexicalScope *S) -> bool {
if (isa<DISubprogram>(S->getScopeNode()))
return false;
+ // Don't skip abstract lexical blocks that are scope targets for global
+ // variables (e.g., function-scope statics). Those globals are emitted
+ // later in endModule() and need to find the block via
+ // getOrCreateContextDIE().
+ if (S->isAbstractScope() && hasGlobalVariableInScope(S->getScopeNode()))
+ return false;
auto Vars = DU->getScopeVariables().lookup(S);
if (!Vars.Args.empty() || !Vars.Locals.empty())
return false;
@@ -1259,8 +1273,9 @@ DIE *DwarfCompileUnit::createAndAddScopeChildren(LexicalScope *Scope,
DD->getLocalDeclsForScope(S->getScopeNode()).empty();
};
for (LexicalScope *LS : Scope->getChildren()) {
- // If the lexical block doesn't have non-scope children, skip
- // its emission and put its children directly to the parent scope.
+ // If the lexical block doesn't have non-scope children or global
+ // variables scoped to it, skip its emission and put its children directly
+ // to the parent scope.
if (skipLexicalScope(LS))
createAndAddScopeChildren(LS, ScopeDIE);
else
diff --git a/llvm/lib/CodeGen/AsmPrinter/DwarfCompileUnit.h b/llvm/lib/CodeGen/AsmPrinter/DwarfCompileUnit.h
index 5e3a2f4b5f75f..fefd8adb3c932 100644
--- a/llvm/lib/CodeGen/AsmPrinter/DwarfCompileUnit.h
+++ b/llvm/lib/CodeGen/AsmPrinter/DwarfCompileUnit.h
@@ -94,6 +94,9 @@ class DwarfCompileUnit final : public DwarfUnit {
/// multiple pointer variables reference the same constant.
DenseMap<std::pair<const DIType *, int64_t>, DIE *> ImplicitPointerDIEs;
+ // Set of scope nodes referenced by global variables in this CU.
+ SmallPtrSet<const MDNode *, 4> GlobalVarScopes;
+
/// DWO ID for correlating skeleton and split units.
uint64_t DWOId = 0;
@@ -157,6 +160,9 @@ class DwarfCompileUnit final : public DwarfUnit {
return DU->getFinalizedAbstractSubprograms();
}
+ /// \returns true if \ref ScopeNode contains a GlobalVariable.
+ bool hasGlobalVariableInScope(const DILocalScope *ScopeNode);
+
void finishNonUnitTypeDIE(DIE& D, const DICompositeType *CTy) override;
/// Add info for Wasm-global-based relocation.
diff --git a/llvm/test/DebugInfo/Generic/global-var-in-abstract-lexical-block.ll b/llvm/test/DebugInfo/Generic/global-var-in-abstract-lexical-block.ll
new file mode 100644
index 0000000000000..e31fd295a4c94
--- /dev/null
+++ b/llvm/test/DebugInfo/Generic/global-var-in-abstract-lexical-block.ll
@@ -0,0 +1,63 @@
+; RUN: %llc_dwarf -O0 -filetype=obj < %s | llvm-dwarfdump -debug-info --name=static_local --show-parents - | FileCheck %s --implicit-check-not=DW_TAG
+
+;; Derived from a CUDA program like:
+;;
+;; __forceinline__ __device__ void foo(int *p) {
+;; { // lexical block -- scope of the static local
+;; static __shared__ int static_local;
+;; *p = static_local;
+;; }
+;; }
+;;
+;; __global__ void bar(int *q) {
+;; {
+;; foo(q); // foo is inlined here
+;; }
+;; }
+;;
+;; __forceinline__ causes foo to be inlined, so an abstract subprogram
+;; origin is created for it. The `static __shared__` variable compiles to an
+;; addrspace(3) global (not a local variable) whose DI scope is the
+;; DILexicalBlock inside foo. When the abstract scope tree is built,
+;; skipLexicalScope() may elide that lexical block because it has no local
+;; variables -- but the global still needs the block as its context DIE.
+
+;; The abstract "foo" subprogram must contain a DW_TAG_lexical_block holding
+;; the "static_local" variable.
+
+; CHECK: DW_TAG_compile_unit
+; CHECK: DW_TAG_subprogram
+; CHECK: DW_AT_name ("foo")
+; CHECK: DW_AT_inline (DW_INL_inlined)
+; CHECK: DW_TAG_lexical_block
+; CHECK: DW_TAG_variable
+; CHECK: DW_AT_name ("static_local")
+
+ at static_local = internal addrspace(3) global i32 0, align 4, !dbg !0
+
+define void @_Z3barPi(ptr %q) !dbg !14 {
+entry:
+ ret void, !dbg !17
+}
+
+!llvm.dbg.cu = !{!10}
+!llvm.module.flags = !{!13}
+
+!0 = !DIGlobalVariableExpression(var: !1, expr: !DIExpression())
+!1 = distinct !DIGlobalVariable(name: "static_local", linkageName: "_ZZ3fooPiE12static_local", scope: !2, file: !3, line: 3, type: !9, isLocal: true, isDefinition: true)
+!2 = distinct !DILexicalBlock(scope: !4, file: !3, line: 2, column: 1)
+!3 = !DIFile(filename: "test.cu", directory: "/")
+!4 = distinct !DISubprogram(name: "foo", linkageName: "_Z3fooPi", scope: !3, file: !3, line: 1, type: !5, scopeLine: 1, spFlags: DISPFlagLocalToUnit | DISPFlagDefinition, unit: !10, retainedNodes: !12)
+!5 = !DISubroutineType(types: !6)
+!6 = !{!7, !8}
+!7 = !DIBasicType(tag: DW_TAG_unspecified_type, name: "void")
+!8 = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !9, size: 64, align: 64, dwarfAddressSpace: 12)
+!9 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed)
+!10 = distinct !DICompileUnit(language: DW_LANG_C_plus_plus, file: !3, producer: "lgenfe: EDG 6.8", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, globals: !11)
+!11 = !{!0}
+!12 = !{}
+!13 = !{i32 1, !"Debug Info Version", i32 3}
+!14 = distinct !DISubprogram(name: "bar", linkageName: "_Z3barPi", scope: !3, file: !3, line: 9, type: !5, scopeLine: 9, spFlags: DISPFlagDefinition, unit: !10, retainedNodes: !12)
+!15 = distinct !DILocation(line: 11, column: 3, scope: !16)
+!16 = distinct !DILexicalBlock(scope: !14, file: !3, line: 10, column: 1)
+!17 = !DILocation(line: 4, column: 3, scope: !2, inlinedAt: !15)
More information about the llvm-commits
mailing list