[PATCH] D79967: Fix debug info for NoDebug attr

Yaxun Liu via Phabricator via cfe-commits Thu, 14 May 2020 14:44:25 -0700

yaxunl created this revision.
yaxunl added reviewers: tra, rjmccall.
Herald added a subscriber: aprantl.


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.


https://reviews.llvm.org/D79967

Files:
  clang/lib/CodeGen/CGDebugInfo.cpp
  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) {
@@ -27,7 +44,7 @@
 // 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/lib/CodeGen/CGDebugInfo.cpp
===================================================================
--- clang/lib/CodeGen/CGDebugInfo.cpp
+++ clang/lib/CodeGen/CGDebugInfo.cpp
@@ -3820,7 +3820,7 @@
   StringRef LinkageName;
 
   const Decl *D = GD.getDecl();
-  if (!D)
+  if (!D || D->hasAttr<NoDebugAttr>())
     return;
 
   llvm::TimeTraceScope TimeScope("DebugFunction", [&]() {

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) {
@@ -27,7 +44,7 @@
 // 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/lib/CodeGen/CGDebugInfo.cpp
===================================================================
--- clang/lib/CodeGen/CGDebugInfo.cpp
+++ clang/lib/CodeGen/CGDebugInfo.cpp
@@ -3820,7 +3820,7 @@
   StringRef LinkageName;
 
   const Decl *D = GD.getDecl();
-  if (!D)
+  if (!D || D->hasAttr<NoDebugAttr>())
     return;
 
   llvm::TimeTraceScope TimeScope("DebugFunction", [&]() {

_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits

[PATCH] D79967: Fix debug info for NoDebug attr

Reply via email to