Hi jpienaar, eliben, rnk,
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.
http://reviews.llvm.org/D9269
Files:
include/clang/Sema/SemaInternal.h
lib/Sema/SemaDecl.cpp
lib/Sema/SemaStmtAsm.cpp
test/SemaCUDA/qualifiers.cu
EMAIL PREFERENCES
http://reviews.llvm.org/settings/panel/emailpreferences/
Index: include/clang/Sema/SemaInternal.h
===================================================================
--- include/clang/Sema/SemaInternal.h
+++ include/clang/Sema/SemaInternal.h
@@ -48,6 +48,18 @@
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.
Index: lib/Sema/SemaDecl.cpp
===================================================================
--- lib/Sema/SemaDecl.cpp
+++ lib/Sema/SemaDecl.cpp
@@ -5752,6 +5752,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
@@ -5766,10 +5767,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);
}
@@ -5818,6 +5825,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 &&
Index: lib/Sema/SemaStmtAsm.cpp
===================================================================
--- lib/Sema/SemaStmtAsm.cpp
+++ lib/Sema/SemaStmtAsm.cpp
@@ -124,16 +124,8 @@
// 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];
Index: test/SemaCUDA/qualifiers.cu
===================================================================
--- test/SemaCUDA/qualifiers.cu
+++ test/SemaCUDA/qualifiers.cu
@@ -1,7 +1,23 @@
-// 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;
_______________________________________________
cfe-commits mailing list
[email protected]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
