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

[tra]

Summary

During device-side CUDA compilation clang currently complains about any TLS variable, regardless of whether it's a host or device one.
This patch suppresses "TLS unsupported" errors for host variables during device compilation and for device variables during host compilation.

[jpienaar]

Looks good to me.

[eliben]

LGTM

include/clang/Sema/SemaInternal.h

51

D's attributes

include/clang/Sema/SemaInternal.h

 // to not be value dependent.
 inline bool IsVariableAConstantExpression(VarDecl *Var, ASTContext &Context) {
   const VarDecl *DefVD = nullptr;
   return !isa<ParmVarDecl>(Var) &&
     Var->isUsableInConstantExpressions(Context) &&
     Var->getAnyInitializer(DefVD) && DefVD->checkInitIsICE(); 
 }
 
+// Helper function to check whether D attributes match current CUDA mode.  Decls
+// with mismatched attributes and related diagnostics may have to be
+// ignored during this CUDA compilation pass.
+inline bool DeclAttrsMatchCUDAMode(const LangOptions &LangOpts, Decl *D) {
+  if (!LangOpts.CUDA || !D)
+    return true;
+  bool isDeviceSideDecl = D->hasAttr<CUDADeviceAttr>() ||
+                          D->hasAttr<CUDASharedAttr>() ||
+                          D->hasAttr<CUDAGlobalAttr>();
+  return isDeviceSideDecl == LangOpts.CUDAIsDevice;
+}
+
 // Directly mark a variable odr-used. Given a choice, prefer to use 
 // MarkVariableReferenced since it does additional checks and then 
 // calls MarkVarDeclODRUsed.
 // If the variable must be captured:
 //  - if FunctionScopeIndexToStopAt is null, capture it in the CurContext
 //  - else capture it in the DeclContext that maps to the 
 //    *FunctionScopeIndexToStopAt on the FunctionScopeInfo stack.  
 inline void MarkVarDeclODRUsed(VarDecl *Var,

lib/Sema/SemaDecl.cpp

Sema::ActOnVariableDeclarator(Scope *S, Declarator &D, DeclContext *DC,
   // lexical context will be different from the semantic context.
   NewVD->setLexicalDeclContext(CurContext);
   if (NewTemplate)
     NewTemplate->setLexicalDeclContext(CurContext);
 
   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
     //   storage-class-specifier static is implied if it does not appear
     //   explicitly.
     // Core issue: 'static' is not implied if the variable is declared
     //   'extern'.
     if (NewVD->hasLocalStorage() &&
         (SCSpec != DeclSpec::SCS_unspecified ||
          TSCS != DeclSpec::TSCS_thread_local ||
          !DC->isFunctionOrMethod()))
       Diag(D.getDeclSpec().getThreadStorageClassSpecLoc(),
            diag::err_thread_non_global)
         << DeclSpec::getSpecifierName(TSCS);
-    else if (!Context.getTargetInfo().isTLSSupported())
+    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
+    } else
       NewVD->setTSCSpec(TSCS);
   }
 
   // C99 6.7.4p3
   //   An inline definition of a function with external linkage shall
   //   not contain a definition of a modifiable object with static or
   //   thread storage duration...
   // We only apply this when the function is required to be defined
else {
         NewTemplate->setModulePrivate();
     }
   }
 
   // Handle attributes prior to checking for duplicates in MergeVarDecl
   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 &&
         (NewVD->hasAttr<CUDASharedAttr>() ||
          NewVD->hasAttr<CUDAConstantAttr>())) {
       NewVD->setStorageClass(SC_Static);
     }
   }

lib/Sema/SemaStmtAsm.cpp

StmtResult Sema::ActOnGCCAsmStmt(SourceLocation AsmLoc, bool IsSimple,
   StringLiteral *AsmString = cast<StringLiteral>(asmString);
   StringLiteral **Clobbers = reinterpret_cast<StringLiteral**>(clobbers.data());
 
   SmallVector<TargetInfo::ConstraintInfo, 4> OutputConstraintInfos;
 
   // The parser verifies that there is a string literal here.
   assert(AsmString->isAscii());
 
-  bool ValidateConstraints = true;
-  if (getLangOpts().CUDA) {
-    // In CUDA mode don't verify asm constraints in device functions during host
-    // compilation and vice versa.
-    bool InDeviceMode = getLangOpts().CUDAIsDevice;
-    FunctionDecl *FD = getCurFunctionDecl();
-    bool IsDeviceFunction =
-        FD && (FD->hasAttr<CUDADeviceAttr>() || FD->hasAttr<CUDAGlobalAttr>());
-    ValidateConstraints = IsDeviceFunction == InDeviceMode;
-  }
+  bool ValidateConstraints =
+      DeclAttrsMatchCUDAMode(getLangOpts(), getCurFunctionDecl());
 
   for (unsigned i = 0; i != NumOutputs; i++) {
     StringLiteral *Literal = Constraints[i];
     assert(Literal->isAscii());
 
     StringRef OutputName;
     if (Names[i])
       OutputName = Names[i]->getName();

test/SemaCUDA/qualifiers.cu

-// RUN: %clang_cc1 -fsyntax-only -verify %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s
+// RUN: %clang_cc1 -triple nvptx-unknown-cuda -fsyntax-only -verify -fcuda-is-device %s
 
 #include "Inputs/cuda.h"
 
+// Host (x86) supports TLS and device-side compilation should ignore
+// host variables. No errors in either case.
+int __thread host_tls_var;
+
+#if defined(__CUDA_ARCH__)
+// NVPTX does not support TLS
+__device__ int __thread device_tls_var; // expected-error {{thread-local storage is not supported for the current target}}
+__shared__ int __thread shared_tls_var; // expected-error {{thread-local storage is not supported for the current target}}
+#else
+// Device-side vars should not produce any errors during host-side
+// compilation.
+__device__ int __thread device_tls_var;
+__shared__ int __thread shared_tls_var;
+#endif
+
 __global__ void g1(int x) {}
 __global__ int g2(int x) { // expected-error {{must have void return type}}
   return 1;
 }