diff options
Diffstat (limited to 'clang/lib/Sema')
| -rw-r--r-- | clang/lib/Sema/SemaDecl.cpp | 18 | ||||
| -rw-r--r-- | clang/lib/Sema/SemaStmtAsm.cpp | 12 |
2 files changed, 16 insertions, 14 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 && diff --git a/clang/lib/Sema/SemaStmtAsm.cpp b/clang/lib/Sema/SemaStmtAsm.cpp index 179e207e76a..9f48616ea5b 100644 --- a/clang/lib/Sema/SemaStmtAsm.cpp +++ b/clang/lib/Sema/SemaStmtAsm.cpp @@ -124,16 +124,8 @@ StmtResult Sema::ActOnGCCAsmStmt(SourceLocation AsmLoc, bool IsSimple, // 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]; |
