Index: clang/lib/CodeGen/CGDebugInfo.cpp
===================================================================
--- clang/lib/CodeGen/CGDebugInfo.cpp
+++ clang/lib/CodeGen/CGDebugInfo.cpp
@@ -3779,7 +3779,7 @@
   StringRef LinkageName;
 
   const Decl *D = GD.getDecl();
-  if (!D)
+  if (!D || D->hasAttr<NoDebugAttr>())
     return;
 
   llvm::TimeTraceScope TimeScope("DebugFunction", [&]() {
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
+
+// CHECK-NOT: define {{.*}}@foo{{.*}}!dbg
+// CHECK-LABEL: define {{.*}}@foo
+// 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/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);
 }