https://github.com/pvelesko created https://github.com/llvm/llvm-project/pull/236525
#213153 made collectVarDeclProps scope every static local in the innermost lexical block and assert that LexicalBlockStack is non-empty. In CUDA/HIP device compilation, static `__device__` locals that host code ODR-uses (ASTContext::CUDADeviceVarODRUsedByHost) are emitted from EmitDeferred at the end of the translation unit, outside any function, so with -g clang hits "Region stack mismatch, stack empty!" (and without assertions reads an empty vector). The pattern is the one host-used-device-var.cu already covers without debug info. This falls back to the declaration context, as before #213153, when no lexical block is open. This PR was generated using AI >From c506be9a7171e424265bfa2c38b4449d2972cbd7 Mon Sep 17 00:00:00 2001 From: Paulius Velesko <[email protected]> Date: Wed, 28 Oct 2026 19:15:13 +0200 Subject: [PATCH] [clang][DebugInfo] Don't assume static locals are emitted inside their function collectVarDeclProps scopes a static local in the innermost lexical block on LexicalBlockStack, asserting the stack is non-empty. CUDA/HIP device compilation also emits static __device__ locals that host code ODR-uses (ASTContext::CUDADeviceVarODRUsedByHost) from EmitDeferred at the end of the translation unit, outside any function, so the stack is empty there and clang asserts (or reads past an empty vector without assertions). Fall back to the declaration context, as before #213153, when no lexical block is open. --- clang/lib/CodeGen/CGDebugInfo.cpp | 3 +-- .../debug-info-host-used-static-local.cu | 13 +++++++++++++ 2 files changed, 14 insertions(+), 2 deletions(-) create mode 100644 clang/test/CodeGenCUDA/debug-info-host-used-static-local.cu diff --git a/clang/lib/CodeGen/CGDebugInfo.cpp b/clang/lib/CodeGen/CGDebugInfo.cpp index 96cda4555e49db..51371aea2d2650 100644 --- a/clang/lib/CodeGen/CGDebugInfo.cpp +++ b/clang/lib/CodeGen/CGDebugInfo.cpp @@ -4759,8 +4759,7 @@ void CGDebugInfo::collectVarDeclProps(const VarDecl *VD, llvm::DIFile *&Unit, // Get context for static locals (that are technically globals) the same way // we do for "local" locals -- by using current lexical block. - if (VD->isStaticLocal()) { - assert(!LexicalBlockStack.empty() && "Region stack mismatch, stack empty!"); + if (VD->isStaticLocal() && !LexicalBlockStack.empty()) { VDContext = LexicalBlockStack.back(); return; } diff --git a/clang/test/CodeGenCUDA/debug-info-host-used-static-local.cu b/clang/test/CodeGenCUDA/debug-info-host-used-static-local.cu new file mode 100644 index 00000000000000..68b016b31aba23 --- /dev/null +++ b/clang/test/CodeGenCUDA/debug-info-host-used-static-local.cu @@ -0,0 +1,13 @@ +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s -emit-llvm -debug-info-kind=limited -o - | FileCheck %s + +#include "Inputs/cuda.h" + +// A static __device__ local ODR-used by host code is emitted outside any function. +void use(int *); +void fun1() { + static __device__ int static_var1 = 3; + use(&static_var1); +} + +// CHECK: @_ZZ4fun1vE11static_var1 = addrspace(1) externally_initialized global i32 3 +// CHECK: !DIGlobalVariable(name: "static_var1" _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
