Index: clang/lib/Sema/SemaDeclAttr.cpp
===================================================================
--- clang/lib/Sema/SemaDeclAttr.cpp
+++ clang/lib/Sema/SemaDeclAttr.cpp
@@ -4361,6 +4361,12 @@
     S.Diag(FD->getBeginLoc(), diag::warn_kern_is_inline) << FD;
 
   D->addAttr(::new (S.Context) CUDAGlobalAttr(S.Context, AL));
+  // In host compilation the kernel is emitted as a stub function, which is
+  // a helper function for launching the kernel. The instructions in the helper
+  // function has nothing to do with the source code of the kernel. Do not emit
+  // debug info for the stub function to avoid confusing the debugger.
+  if (!S.LangOpts.CUDAIsDevice)
+    D->addAttr(NoDebugAttr::CreateImplicit(S.Context));
 }
 
 static void handleGNUInlineAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
Index: clang/test/CodeGenCUDA/kernel-dbg-info.cu
===================================================================
--- /dev/null
+++ clang/test/CodeGenCUDA/kernel-dbg-info.cu
@@ -0,0 +1,33 @@
+// RUN: echo "GPU binary would be here" > %t
+
+// 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: %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
+
+#include "Inputs/cuda.h"
+
+extern "C" __global__ void ckernel(int *a) {
+  *a = 1;
+}
+
+// Device side kernel names
+// CHECK: @[[CKERN:[0-9]*]] = {{.*}} c"ckernel\00"
+
+// DEV: define {{.*}}@ckernel{{.*}}!dbg
+// DEV:  store {{.*}}!dbg
+// DEV:  ret {{.*}}!dbg
+
+// CHECK-NOT: define {{.*}}@__device_stub__ckernel{{.*}}!dbg
+// CHECK: define {{.*}}@[[CSTUB:__device_stub__ckernel]]
+// CHECK-NOT: call {{.*}}@hipLaunchByPtr{{.*}}!dbg
+// CHECK: call {{.*}}@hipLaunchByPtr{{.*}}@[[CSTUB]]
+// CHECK-NOT: ret {{.*}}!dbg
+
+// CHECK-LABEL: define {{.*}}@_Z8hostfuncPi{{.*}}!dbg
+// CHECK: call void @[[CSTUB]]{{.*}}!dbg
+void hostfunc(int *a) {
+  ckernel<<<1, 1>>>(a);
+}