Index: clang/include/clang/AST/Decl.h
===================================================================
--- clang/include/clang/AST/Decl.h
+++ clang/include/clang/AST/Decl.h
@@ -1633,6 +1633,9 @@
     /// Parameter for captured context
     CapturedContext,
 
+    /// Parameter for Thread private variable
+    ThreadPrivateVar,
+
     /// Other implicit parameter
     Other,
   };
Index: clang/lib/CodeGen/CGDecl.cpp
===================================================================
--- clang/lib/CodeGen/CGDecl.cpp
+++ clang/lib/CodeGen/CGDecl.cpp
@@ -2442,6 +2442,7 @@
 /// for the specified parameter and set up LocalDeclMap.
 void CodeGenFunction::EmitParmDecl(const VarDecl &D, ParamValue Arg,
                                    unsigned ArgNo) {
+  bool NoDebugInfo = false;
   // FIXME: Why isn't ImplicitParamDecl a ParmVarDecl?
   assert((isa<ParmVarDecl>(D) || isa<ImplicitParamDecl>(D)) &&
          "Invalid argument to EmitParmDecl");
@@ -2461,6 +2462,10 @@
       setBlockContextParameter(IPD, ArgNo, V);
       return;
     }
+    // Suppressing debug info for ThreadPrivateVar parameters, else it hides
+    // debug info of TLS variables.
+    NoDebugInfo =
+        (IPD->getParameterKind() == ImplicitParamDecl::ThreadPrivateVar);
   }
 
   Address DeclPtr = Address::invalid();
@@ -2591,7 +2596,8 @@
 
   // Emit debug info for param declarations in non-thunk functions.
   if (CGDebugInfo *DI = getDebugInfo()) {
-    if (CGM.getCodeGenOpts().hasReducedDebugInfo() && !CurFuncIsThunk) {
+    if (CGM.getCodeGenOpts().hasReducedDebugInfo() && !CurFuncIsThunk &&
+        !NoDebugInfo) {
       llvm::DILocalVariable *DILocalVar = DI->EmitDeclareOfArgVariable(
           &D, AllocaPtr.getPointer(), ArgNo, Builder);
       if (const auto *Var = dyn_cast_or_null<ParmVarDecl>(&D))
Index: clang/lib/CodeGen/CGStmtOpenMP.cpp
===================================================================
--- clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -481,7 +481,11 @@
     if (ArgType->isVariablyModifiedType())
       ArgType = getCanonicalParamType(Ctx, ArgType);
     VarDecl *Arg;
-    if (DebugFunctionDecl && (CapVar || I->capturesThis())) {
+    if (CapVar && (CapVar->getTLSKind() != clang::VarDecl::TLS_None)) {
+      Arg = ImplicitParamDecl::Create(Ctx, /*DC=*/nullptr, FD->getLocation(),
+                                      II, ArgType,
+                                      ImplicitParamDecl::ThreadPrivateVar);
+    } else if (DebugFunctionDecl && (CapVar || I->capturesThis())) {
       Arg = ParmVarDecl::Create(
           Ctx, DebugFunctionDecl,
           CapVar ? CapVar->getBeginLoc() : FD->getBeginLoc(),
Index: clang/test/OpenMP/debug_threadprivate_copyin.c
===================================================================
--- /dev/null
+++ clang/test/OpenMP/debug_threadprivate_copyin.c
@@ -0,0 +1,59 @@
+// This testcase checks emission of debug info for threadprivate variables
+// present in any clause of OpenMP construct.
+
+// REQUIRES: x86_64-linux
+
+// RUN: %clang_cc1 -debug-info-kind=constructor -x c -verify -triple x86_64-pc-linux-gnu -fopenmp -emit-llvm %s -o - | FileCheck %s
+// expected-no-diagnostics
+
+// CHECK: define internal void @.omp_outlined._debug__(
+// CHECK: call void @llvm.dbg.declare(metadata ptr %.global_tid..addr,
+// CHECK: call void @llvm.dbg.declare(metadata ptr %.bound_tid..addr,
+// CHECK: call void @llvm.dbg.declare(metadata ptr %nt.addr
+// CHECK: store ptr %gbl_dynamic_int, ptr %gbl_dynamic_int.addr, align 8
+// CHECK-NOT: call void @llvm.dbg.declare(metadata ptr %gbl_dynamic_int.addr
+// CHECK-NOT: call void @llvm.dbg.declare(metadata ptr %gbl_static_int.addr
+
+extern int printf(const char *, ...);
+extern void omp_set_num_threads(int);
+extern int omp_get_num_threads(void);
+extern int omp_get_thread_num(void);
+
+int gbl_dynamic_int;
+__thread int gbl_static_int;
+
+#pragma omp threadprivate(gbl_dynamic_int)
+
+int main() {
+  int nt = 0;
+  int offset = 10;
+  gbl_dynamic_int = 55;
+  gbl_static_int = 77;
+
+  omp_set_num_threads(4);
+#pragma omp parallel copyin(gbl_dynamic_int, gbl_static_int)
+  {
+    int data;
+    int tid;
+    nt = omp_get_num_threads();
+    tid = omp_get_thread_num();
+    data = gbl_dynamic_int + gbl_static_int;
+    gbl_dynamic_int += 10;
+    gbl_static_int += 20;
+#pragma omp barrier
+    if (tid == 0)
+      printf("In parallel region total threads = %d, thread id = %d data=%d gbl_dyn_addr = %p, gbl_static_addr = %p\n",
+             nt, tid, data, &gbl_dynamic_int, &gbl_static_int);
+    if (tid == 1)
+      printf("In parallel region total threads = %d, thread id = %d data=%d gbl_dyn_addr = %p, gbl_static_addr = %p\n",
+             nt, tid, data, &gbl_dynamic_int, &gbl_static_int);
+    if (tid == 2)
+      printf("In parallel region total threads = %d, thread id = %d data=%d gbl_dyn_addr = %p, gbl_static_addr = %p\n",
+             nt, tid, data, &gbl_dynamic_int, &gbl_static_int);
+    if (tid == 3)
+      printf("In parallel region total threads = %d, thread id = %d data=%d gbl_dyn_addr = %p, gbl_static_addr = %p\n",
+             nt, tid, data, &gbl_dynamic_int, &gbl_static_int);
+  }
+
+  return 0;
+}