diff options
Diffstat (limited to 'clang/lib/Sema/SemaDecl.cpp')
| -rw-r--r-- | clang/lib/Sema/SemaDecl.cpp | 18 |
1 files changed, 14 insertions, 4 deletions
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 07dbdaf9968..58e78387fd8 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -5753,6 +5753,7 @@ Sema::ActOnVariableDeclarator(Scope *S, Declarator &D, DeclContext *DC, 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 @@ Sema::ActOnVariableDeclarator(Scope *S, Declarator &D, DeclContext *DC, 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 @@ Sema::ActOnVariableDeclarator(Scope *S, Declarator &D, DeclContext *DC, 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 && |
