diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -7077,7 +7077,8 @@
            diag::err_thread_non_global)
         << DeclSpec::getSpecifierName(TSCS);
     else if (!Context.getTargetInfo().isTLSSupported()) {
-      if (getLangOpts().CUDA || getLangOpts().OpenMPIsDevice) {
+      if (getLangOpts().CUDA || getLangOpts().OpenMPIsDevice ||
+          getLangOpts().SYCLIsDevice) {
         // Postpone error emission until we've collected attributes required to
         // figure out whether it's a host or device variable and whether the
         // error should be ignored.
@@ -7179,13 +7180,18 @@
   // Handle attributes prior to checking for duplicates in MergeVarDecl
   ProcessDeclAttributes(S, NewVD, D);
 
-  if (getLangOpts().CUDA || getLangOpts().OpenMPIsDevice) {
+  if (getLangOpts().CUDA || getLangOpts().OpenMPIsDevice ||
+      getLangOpts().SYCLIsDevice) {
     if (EmitTLSUnsupportedError &&
         ((getLangOpts().CUDA && DeclAttrsMatchCUDAMode(getLangOpts(), NewVD)) ||
          (getLangOpts().OpenMPIsDevice &&
           OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(NewVD))))
       Diag(D.getDeclSpec().getThreadStorageClassSpecLoc(),
            diag::err_thread_unsupported);
+
+    if (EmitTLSUnsupportedError &&
+        (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice)))
+      targetDiag(D.getIdentifierLoc(), diag::err_thread_unsupported);
     // CUDA B.2.5: "__shared__ and __constant__ variables have implied static
     // storage [duration]."
     if (SC == SC_None && S->getFnParent() != nullptr &&
diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/clang/lib/Sema/SemaExpr.cpp
@@ -355,10 +355,16 @@
 
   diagnoseUseOfInternalDeclInInlineFunction(*this, D, Loc);
 
-  if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice))
+  if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice)) {
     if (const auto *VD = dyn_cast<ValueDecl>(D))
       checkDeviceDecl(VD, Loc);
 
+    if (!Context.getTargetInfo().isTLSSupported())
+      if (const auto *VD = dyn_cast<VarDecl>(D))
+        if (VD->getTLSKind() != VarDecl::TLS_None)
+          targetDiag(*Locs.begin(), diag::err_thread_unsupported);
+  }
+
   if (isa<ParmVarDecl>(D) && isa<RequiresExprBodyDecl>(D->getDeclContext()) &&
       !isUnevaluatedContext()) {
     // C++ [expr.prim.req.nested] p3
diff --git a/clang/test/OpenMP/nvptx_prohibit_thread_local.cpp b/clang/test/OpenMP/nvptx_prohibit_thread_local.cpp
new file mode 100644
--- /dev/null
+++ b/clang/test/OpenMP/nvptx_prohibit_thread_local.cpp
@@ -0,0 +1,44 @@
+// RUN: %clang_cc1 -fopenmp -x c++ -triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -fsyntax-only
+
+thread_local const int prohobit_ns_scope = 0;
+thread_local int prohobit_ns_scope2 = 0;
+thread_local const int allow_ns_scope = 0;
+
+struct S {
+  static const thread_local int prohibit_static_member;
+  static thread_local int prohibit_static_member2;
+};
+
+struct T {
+  static const thread_local int allow_static_member;
+};
+
+void foo() {
+  // expected-error@+1{{thread-local storage is not supported for the current target}}
+  thread_local const int prohibit_local = 0;
+  // expected-error@+1{{thread-local storage is not supported for the current target}}
+  thread_local int prohibit_local2;
+}
+
+void bar() { thread_local int allow_local; }
+
+void usage() {
+  // expected-note@+1 {{called by}}
+  foo();
+  // expected-error@+1 {{thread-local storage is not supported for the current target}}
+  (void)prohobit_ns_scope;
+  // expected-error@+1 {{thread-local storage is not supported for the current target}}
+  (void)prohobit_ns_scope2;
+  // expected-error@+1 {{thread-local storage is not supported for the current target}}
+  (void)S::prohibit_static_member;
+  // expected-error@+1 {{thread-local storage is not supported for the current target}}
+  (void)S::prohibit_static_member2;
+}
+
+int main() {
+  // expected-note@+2 2{{called by}}
+#pragma omp target
+  usage();
+  return 0;
+}
diff --git a/clang/test/OpenMP/nvptx_target_codegen.cpp b/clang/test/OpenMP/nvptx_target_codegen.cpp
--- a/clang/test/OpenMP/nvptx_target_codegen.cpp
+++ b/clang/test/OpenMP/nvptx_target_codegen.cpp
@@ -160,7 +160,7 @@
 // CHECK: [[EXIT]]
 // CHECK: ret void
 
-// CHECK: define {{.*}}void [[T2:@__omp_offloading_.+foo.+l200]](i[[SZ:32|64]] [[ARG1:%[a-zA-Z_]+]], i[[SZ:32|64]] [[ID:%[a-zA-Z_]+]])
+// CHECK: define {{.*}}void [[T2:@__omp_offloading_.+foo.+l200]](i[[SZ:32|64]] [[ARG1:%[a-zA-Z_]+]])
 // CHECK: [[AA_ADDR:%.+]] = alloca i[[SZ]],
 // CHECK: store i[[SZ]] [[ARG1]], i[[SZ]]* [[AA_ADDR]],
 // CHECK: [[AA_CADDR:%.+]] = bitcast i[[SZ]]* [[AA_ADDR]] to i16*
@@ -200,7 +200,7 @@
 #pragma omp target if (1)
   {
     aa += 1;
-    id = aa;
+    aa += 2;
   }
 
 // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l310}}_worker()
diff --git a/clang/test/SemaSYCL/prohibit-thread-local.cpp b/clang/test/SemaSYCL/prohibit-thread-local.cpp
new file mode 100644
--- /dev/null
+++ b/clang/test/SemaSYCL/prohibit-thread-local.cpp
@@ -0,0 +1,48 @@
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64 -verify -fsyntax-only %s
+
+thread_local const int prohobit_ns_scope = 0;
+thread_local int prohobit_ns_scope2 = 0;
+thread_local const int allow_ns_scope = 0;
+
+struct S {
+  static const thread_local int prohibit_static_member;
+  static thread_local int prohibit_static_member2;
+};
+
+struct T {
+  static const thread_local int allow_static_member;
+};
+
+void foo() {
+  // expected-error@+1{{thread-local storage is not supported for the current target}}
+  thread_local const int prohibit_local = 0;
+  // expected-error@+1{{thread-local storage is not supported for the current target}}
+  thread_local int prohibit_local2;
+}
+
+void bar() { thread_local int allow_local; }
+
+void usage() {
+  // expected-note@+1 {{called by}}
+  foo();
+  // expected-error@+1 {{thread-local storage is not supported for the current target}}
+  (void)prohobit_ns_scope;
+  // expected-error@+1 {{thread-local storage is not supported for the current target}}
+  (void)prohobit_ns_scope2;
+  // expected-error@+1 {{thread-local storage is not supported for the current target}}
+  (void)S::prohibit_static_member;
+  // expected-error@+1 {{thread-local storage is not supported for the current target}}
+  (void)S::prohibit_static_member2;
+}
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel))
+// expected-note@+2 2{{called by}}
+void
+kernel_single_task(Func kernelFunc) { kernelFunc(); }
+
+int main() {
+  // expected-note@+1 2{{called by}}
+  kernel_single_task<class fake_kernel>([]() { usage(); });
+  return 0;
+}