Var->getAnyInitializer(DefVD) && DefVD->checkInitIsICE();
}
+// Helper function to check whether D's 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 (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
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);
}
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 &&
// 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];
-// 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;