llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: Paulius Velesko (pvelesko) <details> <summary>Changes</summary> #<!-- -->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 --- Full diff: https://github.com/llvm/llvm-project/pull/236525.diff 2 Files Affected: - (modified) clang/lib/CodeGen/CGDebugInfo.cpp (+1-2) - (added) clang/test/CodeGenCUDA/debug-info-host-used-static-local.cu (+13) ``````````diff 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" `````````` </details> https://github.com/llvm/llvm-project/pull/236525 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
