Fix debug info for NoDebug attr

[yaxunl]

Summary

NoDebug attr does not totally eliminate debug info about a function when
inlining is enabled. This is inconsistent with when inlining is disabled.

This patch fixes that.

[tra]

LGTM. Added @dblaikie as reviewer for debug info expertise.

[dblaikie]

I'm not quite following the test case - could you explain how it's testing the situation?

It might be better/easier to have a more intentional test - also, could you check the commit history for this declaration-emitting functionality (CGDebugInfo::EmitFunctionDecl when used for a non-member function) & pull them into this review too?

[yaxunl]

In D79967#2037675, @dblaikie wrote:

I'm not quite following the test case - could you explain how it's testing the situation?

It might be better/easier to have a more intentional test - also, could you check the commit history for this declaration-emitting functionality (CGDebugInfo::EmitFunctionDecl when used for a non-member function) & pull them into this review too?

This patch is a follow up of https://reviews.llvm.org/D79866, where we marked a compiler generated stub function with NoDebug attribute to eliminate debug info for it. Then we found there is a bug causing debug info not fully eliminated for a function with NoDebug attribute when the function is inlined, therefore I created this patch.

The bug can be demonstrated by

https://godbolt.org/z/z_tA3Z

There is debug info with function foo, which should not. If you change -O3 to -O0, the debug info on function foo disappears, which is the expected behavior.

The following code is supposed to eliminate all debug info for functions with NoDebugAttr, however it is not sufficient when the function gets inlined:

https://github.com/llvm/llvm-project/blob/dad2e92eaf531500af5ab9ee3350e3e186944185/clang/lib/CodeGen/CodeGenFunction.cpp#L1272

I can add a lit test for a function with NoDebugAttr.

[yaxunl]

add a test for nodebug attr.

[dblaikie]

Could you check the commit history for this feature and rope in some folks who added the function declaration work (it's for debug call sites) - maybe @vsk is the right person, or knows who is, to check this is the right fix for it/doesn't adversely affect the feature this code was added to implement.

clang/test/CodeGen/nodebug-attr.c

5–6 ↗

(On Diff #264148)

Does this test fail when the bug is present? I'd have thought not - my understanding was that CHECK-LABEL is found first, then CHECK-NOT is tested between labels/other checks, so it wouldn't find @foo.*!dbg anyway.

I think maybe it'd be better to tighten up the CHECK-LABEL to include "#0 {" at the end and a comment saying how that CHECK part ensures there's no !dbg attached to it.

8–15 ↗

(On Diff #264148)

It looks like this test case currently crashes LLVM:

h$ clang++-tot test.cpp -g -c -O3 && llvm-dwarfdump-tot test.o
PLEASE submit a bug report to https://bugs.llvm.org/ and include the crash backtrace, preprocessed source, and associated run script.
Stack dump:
0. Program arguments: clang++-tot test.cpp -g -c -O3

1. <eof> parser at end of file
2. Code generation
3. Running pass 'Function Pass Manager' on module 'test.cpp'.
4. Running pass 'Debug Variable Analysis' on function '@_Z3fooPi' #0 0x0000000006b49927 llvm::sys::PrintStackTrace(llvm::raw_ostream&) /usr/local/google/home/blaikie/dev/llvm/src/llvm/lib/Support/Unix/Signals.inc:564:11 #1 0x0000000006b49ac9 PrintStackTraceSignalHandler(void*) /usr/local/google/home/blaikie/dev/llvm/src/llvm/lib/Support/Unix/Signals.inc:625:1 #2 0x0000000006b482db llvm::sys::RunSignalHandlers() /usr/local/google/home/blaikie/dev/llvm/src/llvm/lib/Support/Signals.cpp:67:5 #3 0x0000000006b4921e llvm::sys::CleanupOnSignal(unsigned long) /usr/local/google/home/blaikie/dev/llvm/src/llvm/lib/Support/Unix/Signals.inc:362:1 #4 0x0000000006a7cae8 (anonymous namespace)::CrashRecoveryContextImpl::HandleCrash(int, unsigned long) /usr/local/google/home/blaikie/dev/llvm/src/llvm/lib/Support/CrashRecoveryContext.cpp:77:20 #5 0x0000000006a7cd6e CrashRecoverySignalHandler(int) /usr/local/google/home/blaikie/dev/llvm/src/llvm/lib/Support/CrashRecoveryContext.cpp:383:1 #6 0x00007fc197167520 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x13520) #7 0x0000000005a309ac llvm::DICompileUnit::getEmissionKind() const /usr/local/google/home/blaikie/dev/llvm/src/llvm/include/llvm/IR/DebugInfoMetadata.h:1272:31 #8 0x0000000005a4a107 llvm::LexicalScopes::initialize(llvm::MachineFunction const&) /usr/local/google/home/blaikie/dev/llvm/src/llvm/lib/CodeGen/LexicalScopes.cpp:53:70 #9 0x0000000005e128dd (anonymous namespace)::LDVImpl::computeIntervals() /usr/local/google/home/blaikie/dev/llvm/src/llvm/lib/CodeGen/LiveDebugVariables.cpp:972:17

#10 0x0000000005e11284 (anonymous namespace)::LDVImpl::runOnMachineFunction(llvm::MachineFunction&) /usr/local/google/home/blaikie/dev/llvm/src/llvm/lib/CodeGen/LiveDebugVariables.cpp:0:3
#11 0x0000000005e10fe4 llvm::LiveDebugVariables::runOnMachineFunction(llvm::MachineFunction&) /usr/local/google/home/blaikie/dev/llvm/src/llvm/lib/CodeGen/LiveDebugVariables.cpp:1014:3

Is that crash something this patch is intended to address, or unrelated? I guess it's intended to address that problem - because it's a DISubprogram for 'foo' with no DICompileUnit attachment that causes the crash later on.

[yaxunl]

In D79967#2039153, @dblaikie wrote:

Could you check the commit history for this feature and rope in some folks who added the function declaration work (it's for debug call sites) - maybe @vsk is the right person, or knows who is, to check this is the right fix for it/doesn't adversely affect the feature this code was added to implement.

Anders Carlsson introduced support of nodebug attribute by the following two commits:

https://github.com/llvm/llvm-project/commit/76187b4d689a6ce601f2055b3dad5be6a4ab1012
https://github.com/llvm/llvm-project/commit/63784f4e5e125b7a81c92c2cf176797ce67b2e6d

However, it seems he is not in Phabricator.

Based on documentation of nodebug attribute:

https://clang.llvm.org/docs/AttributeReference.html#nodebug

clang should not emit any debug information for functions with nodebug attr.

clang/test/CodeGen/nodebug-attr.c

5–6 ↗

(On Diff #264148)

You are right. The test is not good at detecting the bad IR. Will fix it by removing CHECK-NOT and tighten up CHECK-LABEL.

The test CodeGenCUDA/kernel-dbg-info.cu has similar issue and will also be fixed.

8–15 ↗

(On Diff #264148)

Yes this patch fixes the crash as a side effect. Basically what happens before this patch is:

When clang emits call of foo, it tries to emit DISubprogram for foo as declaration to provide debug info for the call site:

https://github.com/llvm/llvm-project/blob/57d8b8d6f0b91b380ca3b270b4439c338ed67f53/clang/lib/CodeGen/CGExpr.cpp#L5148

Then control flow goes to CGDebugInfo::EmitFuncDeclForCallSite.

https://github.com/llvm/llvm-project/blob/master/clang/lib/CodeGen/CGDebugInfo.cpp#L3877

As we can see for normal function with definitions, clang should have emitted DISubprogram, therefore clang just returns at line 3886. Also clang does not always emit debug info for function decl for call site. e.g. it does not do that for builtin functions, static and inline functions. IMHO it should not do that for functions with nodebug info either. However it does that.

Then control goes to

https://github.com/llvm/llvm-project/blob/57d8b8d6f0b91b380ca3b270b4439c338ed67f53/clang/lib/CodeGen/CGDebugInfo.cpp#L3866

Here clang tries to emit DISubprogram for foo as a declaration instead of definition. It does not set llvm::DISubprogram::SPFlagDefinition, which causes a nullptr Unit in DISubprogram.

When control flow goes to

https://github.com/llvm/llvm-project/blob/5b0502dff5b6f510fbf823059faa042dd1523cef/llvm/lib/CodeGen/LexicalScopes.cpp#L53

Since getUnit() is nullptr, clang crashes.

The situation of a function with nodebug attr is similar to a builtin function. They are actually function definitions since clang emits instructions for them. On the other hand, there is no debug info for these instructions and there is no need to emit any debug info for the function as declarations. For builtin functions, clang does not emit DISubprogram for them as declarations for call sites, the same logic should apply to functions with nodebug attr.

[yaxunl]

Fix the tests and move the logic to CGDebugInfo::EmitFuncDeclForCallSite to make it clearer.

[dblaikie]

In D79967#2041595, @yaxunl wrote:

In D79967#2039153, @dblaikie wrote:

Could you check the commit history for this feature and rope in some folks who added the function declaration work (it's for debug call sites) - maybe @vsk is the right person, or knows who is, to check this is the right fix for it/doesn't adversely affect the feature this code was added to implement.

Anders Carlsson introduced support of nodebug attribute by the following two commits:

https://github.com/llvm/llvm-project/commit/76187b4d689a6ce601f2055b3dad5be6a4ab1012
https://github.com/llvm/llvm-project/commit/63784f4e5e125b7a81c92c2cf176797ce67b2e6d

However, it seems he is not in Phabricator.

Based on documentation of nodebug attribute:

https://clang.llvm.org/docs/AttributeReference.html#nodebug

clang should not emit any debug information for functions with nodebug attr.

Oh, sorry - I didn't mean the nodebug attribute feature I meant the "emitting function declarations" feature - looks like @vsk and @djtodoro have worked on that support & I'd like them to chime in on this.

When control flow goes to

https://github.com/llvm/llvm-project/blob/5b0502dff5b6f510fbf823059faa042dd1523cef/llvm/lib/CodeGen/LexicalScopes.cpp#L53

Since getUnit() is nullptr, clang crashes.

Hmm, this sort of seems like it might lead to other bugs - @vsk wouldn't this assert/fail under LTO too? (we might emit a DISubprogram declaration for a function in one translation unit, but if another translation unit that defines that function is emitted without debug info - would the DISubprogram declaration (with no unit field) get attached to the definition and lead to this sort of crash?)? hmm, nope, because these things are never attached to definitions usually. So then a side question becomes "how did this end up attached to a definition?" (I guess "because there was a definition")

@vsk - we might need to have some broader design discussion. I see that these function declarations end up in the CU's retainedTypes field (I might've been involved in the discussion that got them there? I'm not sure - either way, in retrospect/now that I see it, I have concerns). This means the declarations persist through LTO and don't get deduplicated against the definitions - that's not ideal/will bloat up LTO and ThinLTO bitcode, and persist through other optimizations (eg: if the function gets optimized away/all call sites are dead and optimized out, etc) - it's /probably/ better they get attached to the llvm::Function declaration so it can be cleaned up as needed. That's why DISubprogram moved to be attached to an llvm::Function (previously all DISubprograms were listed on the CU and DISubprograms pointed to llvm::Functions) to aid in LTO (@aprantl and @dexonsmith who were involved in that migration).

But as it pertains to this bug - I /imagine/ the fix proposed right now is about right, but I'd like @vsk or similar to chime in to confirm (& then maybe start another thread on ^)

[djtodoro]

This seems reasonable, so this change looks good to me!

@dblaikie Thanks for pointing out to the potential problems of the usage of the func decl !dbg in the purpose of call sites debug info. It is currently being stored into CU's retainedTypes field.

That's why DISubprogram moved to be attached to an llvm::Function (previously all DISubprograms were listed on the CU and DISubprograms pointed to llvm::Functions) to aid in LTO

Is it only one patch or a set of patches? I'll try to find these changes in the log history.

I am open for the discussion about redesigning it (please wait for a comment from @vsk as well). :)

[vsk]

@yaxunl thanks, this patch lgtm.

@dblaikie I've kicked off a thread on cfe-dev about the topics you brought up ("Design discussion re: DW_TAG_call_site support in clang") and cc'd you.

[dblaikie]

In D79967#2045196, @vsk wrote:

@yaxunl thanks, this patch lgtm.

@dblaikie I've kicked off a thread on cfe-dev about the topics you brought up ("Design discussion re: DW_TAG_call_site support in clang") and cc'd you.

Thanks on both counts!

clang/lib/CodeGen/CGDebugInfo.cpp

 }
 
 void CGDebugInfo::EmitFunctionDecl(GlobalDecl GD, SourceLocation Loc,
                                    QualType FnType, llvm::Function *Fn) {
   StringRef Name;
   StringRef LinkageName;
 
   const Decl *D = GD.getDecl();
-  if (!D)
+  if (!D || D->hasAttr<NoDebugAttr>())
     return;
 
   llvm::TimeTraceScope TimeScope("DebugFunction", [&]() {
     std::string Name;
     llvm::raw_string_ostream OS(Name);
     if (const NamedDecl *ND = dyn_cast<NamedDecl>(D))
       ND->getNameForDiagnostic(OS, getPrintingPolicy(),
                                /*Qualified=*/true);

clang/test/CodeGenCUDA/kernel-dbg-info.cu

 // RUN: echo "GPU binary would be here" > %t
 
 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -O0 \
 // RUN:   -fcuda-include-gpubinary %t -debug-info-kind=limited \
-// RUN:   -o - -x hip | FileCheck %s
+// RUN:   -o - -x hip | FileCheck -check-prefixes=CHECK,O0 %s
 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm %s -O0 \
 // RUN:   -fcuda-include-gpubinary %t -debug-info-kind=limited \
 // RUN:   -o - -x hip -fcuda-is-device | FileCheck -check-prefix=DEV %s
 
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -O0 \
+// RUN:   -fcuda-include-gpubinary %t -debug-info-kind=limited \
+// RUN:   -o - -x hip -debugger-tuning=gdb -dwarf-version=4 \
+// RUN:   | FileCheck -check-prefixes=CHECK,O0 %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm %s -O0 \
+// RUN:   -fcuda-include-gpubinary %t -debug-info-kind=limited \
+// RUN:   -o - -x hip -debugger-tuning=gdb -dwarf-version=4 \
+// RUN:   -fcuda-is-device | FileCheck -check-prefix=DEV %s
+
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -O3 \
+// RUN:   -fcuda-include-gpubinary %t -debug-info-kind=limited \
+// RUN:   -o - -x hip -debugger-tuning=gdb -dwarf-version=4 | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm %s -O3 \
+// RUN:   -fcuda-include-gpubinary %t -debug-info-kind=limited \
+// RUN:   -o - -x hip -debugger-tuning=gdb -dwarf-version=4 \
+// RUN:   -fcuda-is-device | FileCheck -check-prefix=DEV %s
+
 #include "Inputs/cuda.h"
 
 extern "C" __global__ void ckernel(int *a) {
   *a = 1;
 }
 
 // Device side kernel names
 // CHECK: @[[CKERN:[0-9]*]] = {{.*}} c"ckernel\00"
 
 // DEV: define {{.*}}@ckernel{{.*}}!dbg
 // DEV:  store {{.*}}!dbg
 // DEV:  ret {{.*}}!dbg
 
 // CHECK-NOT: define {{.*}}@__device_stub__ckernel{{.*}}!dbg
 // CHECK: define {{.*}}@[[CSTUB:__device_stub__ckernel]]
 // CHECK-NOT: call {{.*}}@hipLaunchByPtr{{.*}}!dbg
 // CHECK: call {{.*}}@hipLaunchByPtr{{.*}}@[[CSTUB]]
 // CHECK-NOT: ret {{.*}}!dbg
 
 // CHECK-LABEL: define {{.*}}@_Z8hostfuncPi{{.*}}!dbg
-// CHECK: call void @[[CSTUB]]{{.*}}!dbg
+// O0: call void @[[CSTUB]]{{.*}}!dbg
 void hostfunc(int *a) {
   ckernel<<<1, 1>>>(a);
 }