replica
/
hanchenye-llvm-project
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

Fix debug info for NoDebug attr 

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.

Differential Revision: https://reviews.llvm.org/D79967
This commit is contained in:
Yaxun (Sam) Liu 2020-05-14 23:09:30 -04:00
parent 3411a1a920
commit 361e4f14e3
3 changed files with 42 additions and 10 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

12
clang/lib/CodeGen/CGDebugInfo.cpp
View File

@ -3885,12 +3885,12 @@ void CGDebugInfo::EmitFuncDeclForCallSite(llvm::CallBase *CallOrInvoke,
if (Func->getSubprogram())
return;
// Do not emit a declaration subprogram for a builtin or if call site info
// isn't required. Also, elide declarations for functions with reserved names,
// as call site-related features aren't interesting in this case (& also, the
// compiler may emit calls to these functions without debug locations, which
// makes the verifier complain).
if (CalleeDecl->getBuiltinID() != 0 ||
// Do not emit a declaration subprogram for a builtin, a function with nodebug
// attribute, or if call site info isn't required. Also, elide declarations
// for functions with reserved names, as call site-related features aren't
// interesting in this case (& also, the compiler may emit calls to these
// functions without debug locations, which makes the verifier complain).
if (CalleeDecl->getBuiltinID() != 0 || CalleeDecl->hasAttr<NoDebugAttr>() ||
getCallSiteRelatedAttrs() == llvm::DINode::FlagZero)
return;
if (const auto *Id = CalleeDecl->getIdentifier())

15
clang/test/CodeGen/nodebug-attr.c Normal file
View File

@ -0,0 +1,15 @@
// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -O3 \
// RUN: -debug-info-kind=limited -o - -debugger-tuning=gdb -dwarf-version=4 \
// RUN: | FileCheck %s
// Makes sure there is no !dbg between function attributes and '{'.
// CHECK-LABEL: define void @foo{{.*}} #{{[0-9]+}} {
// CHECK-NOT: ret {{.*}}!dbg
__attribute__((nodebug)) void foo(int *a) {
*a = 1;
}
// CHECK-LABEL: define {{.*}}@bar{{.*}}!dbg
void bar(int *a) {
foo(a);
}

25
clang/test/CodeGenCUDA/kernel-dbg-info.cu
View File

@ -2,11 +2,28 @@
// 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) {
@ -20,14 +37,14 @@ extern "C" __global__ void ckernel(int *a) {
// DEV: store {{.*}}!dbg
// DEV: ret {{.*}}!dbg
// CHECK-NOT: define {{.*}}@__device_stub__ckernel{{.*}}!dbg
// CHECK: define {{.*}}@[[CSTUB:__device_stub__ckernel]]
// Make sure there is no !dbg between function attributes and '{'
// CHECK: define void @[[CSTUB:__device_stub__ckernel]]{{.*}} #{{[0-9]+}} {
// 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);
}