diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index d2be8141a3a4b..55648963df36a 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -1441,6 +1441,7 @@ llvm::Value *CGOpenMPRuntime::getThreadID(CodeGenFunction &CGF,
     setLocThreadIdInsertPt(CGF);
   CGBuilderTy::InsertPointGuard IPG(CGF.Builder);
   CGF.Builder.SetInsertPoint(Elem.second.ServiceInsertPt);
+  auto DL = ApplyDebugLocation::CreateDefaultArtificial(CGF, Loc);
   llvm::CallInst *Call = CGF.Builder.CreateCall(
       OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(),
                                             OMPRTL___kmpc_global_thread_num),
diff --git a/clang/test/OpenMP/debug-info-kmpc.cpp b/clang/test/OpenMP/debug-info-kmpc.cpp
new file mode 100644
index 0000000000000..ee4e0108469c8
--- /dev/null
+++ b/clang/test/OpenMP/debug-info-kmpc.cpp
@@ -0,0 +1,17 @@
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple amdgcn-amd-amdhsa -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-target-device "-debug-info-kind=constructor" -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s
+
+// Check that we properly attach debug info to the __kmpc_global_thread_num call
+// CHECK: call {{.*}} @__kmpc_global_thread_num{{.*}}!dbg
+
+extern int bar();
+void foo() {
+#pragma omp target teams
+  {
+#pragma omp parallel
+    {
+      bar();
+    }
+  }
+}
+