[PATCH] D79967: Fix debug info for NoDebug attr

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl updated this revision to Diff 264622.
yaxunl marked 2 inline comments as done.
yaxunl added a comment.

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


CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D79967/new/

https://reviews.llvm.org/D79967

Files:
  clang/lib/CodeGen/CGDebugInfo.cpp
  clang/test/CodeGen/nodebug-attr.c
  clang/test/CodeGenCUDA/kernel-dbg-info.cu


Index: clang/test/CodeGenCUDA/kernel-dbg-info.cu
===================================================================
--- clang/test/CodeGenCUDA/kernel-dbg-info.cu
+++ clang/test/CodeGenCUDA/kernel-dbg-info.cu
@@ -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 @@
 // 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);
 }
Index: clang/test/CodeGen/nodebug-attr.c
===================================================================
--- /dev/null
+++ clang/test/CodeGen/nodebug-attr.c
@@ -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);
+}
Index: clang/lib/CodeGen/CGDebugInfo.cpp
===================================================================
--- clang/lib/CodeGen/CGDebugInfo.cpp
+++ clang/lib/CodeGen/CGDebugInfo.cpp
@@ -3844,12 +3844,12 @@
   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())


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D79967.264622.patch
Type: text/x-patch
Size: 4272 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20200518/9ddc6c6e/attachment-0001.bin>