commit fa62ad40872b0d85f9b9a55f0b806c910edab257
author Artem Belevich <tra@google.com> Mon Apr 27 19:37:53 2015 +0000
committer Artem Belevich <tra@google.com> Mon Apr 27 19:37:53 2015 +0000
tree a84ff38123dfd0635d6cf3f876b2245558998be0
parent 0eafe5df712a0394feb8d49d629dc8bace0f292d

[cuda] Ignore "TLS unsupported by target" errors for host variables during device compilation.

During device-side CUDA compilation clang currently complains about
all TLS variables, regardless of whether they are __host__ or
__device__.

This patch suppresses "TLS unsupported" errors for host variables
during device compilation and for device variables during host
compilation.

Differential Revision: http://reviews.llvm.org/D9269

llvm-svn: 235907

clang/lib/Sema/SemaDecl.cpp

diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 07dbdaf..58e7838 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -5753,6 +5753,7 @@
   if (IsLocalExternDecl)
     NewVD->setLocalExternDecl();
 
+  bool EmitTLSUnsupportedError = false;
   if (DeclSpec::TSCS TSCS = D.getDeclSpec().getThreadStorageClassSpec()) {
     // C++11 [dcl.stc]p4:
     //   When thread_local is applied to a variable of block scope the
@@ -5767,10 +5768,16 @@
       Diag(D.getDeclSpec().getThreadStorageClassSpecLoc(),
            diag::err_thread_non_global)
         << DeclSpec::getSpecifierName(TSCS);
-    else if (!Context.getTargetInfo().isTLSSupported())
-      Diag(D.getDeclSpec().getThreadStorageClassSpecLoc(),
-           diag::err_thread_unsupported);
-    else
+    else if (!Context.getTargetInfo().isTLSSupported()) {
+      if (getLangOpts().CUDA)
+        // 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.
+        EmitTLSUnsupportedError = true;
+      else
+        Diag(D.getDeclSpec().getThreadStorageClassSpecLoc(),
+             diag::err_thread_unsupported);
+    } else
       NewVD->setTSCSpec(TSCS);
   }
 
@@ -5819,6 +5826,9 @@
   ProcessDeclAttributes(S, NewVD, D);
 
   if (getLangOpts().CUDA) {
+    if (EmitTLSUnsupportedError && DeclAttrsMatchCUDAMode(getLangOpts(), NewVD))
+      Diag(D.getDeclSpec().getThreadStorageClassSpecLoc(),
+           diag::err_thread_unsupported);
     // CUDA B.2.5: "__shared__ and __constant__ variables have implied static
     // storage [duration]."
     if (SC == SC_None && S->getFnParent() != nullptr &&