diff --git a/clang/lib/CodeGen/CGDebugInfo.cpp b/clang/lib/CodeGen/CGDebugInfo.cpp
--- a/clang/lib/CodeGen/CGDebugInfo.cpp
+++ b/clang/lib/CodeGen/CGDebugInfo.cpp
@@ -3885,12 +3885,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())
diff --git a/clang/test/CodeGen/nodebug-attr.c b/clang/test/CodeGen/nodebug-attr.c
new file mode 100644
--- /dev/null
+++ b/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);
+}
diff --git a/clang/test/CodeGenCUDA/kernel-dbg-info.cu b/clang/test/CodeGenCUDA/kernel-dbg-info.cu
--- a/clang/test/CodeGenCUDA/kernel-dbg-info.cu
+++ b/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);
 }