Skip to content

[SYCL] Align TLS diagnosing with community code #1866

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Jun 11, 2020
Merged

[SYCL] Align TLS diagnosing with community code #1866

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
23 changes: 9 additions & 14 deletions clang/lib/Sema/SemaDecl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7071,7 +7071,8 @@ NamedDecl *Sema::ActOnVariableDeclarator(
diag::err_thread_non_global)
<< DeclSpec::getSpecifierName(TSCS);
else if (!Context.getTargetInfo().isTLSSupported()) {
if (getLangOpts().CUDA || getLangOpts().OpenMPIsDevice) {
if (getLangOpts().CUDA || getLangOpts().OpenMPIsDevice ||
getLangOpts().SYCLIsDevice) {
// 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.
Expand All @@ -7080,14 +7081,6 @@ NamedDecl *Sema::ActOnVariableDeclarator(
// proper storage class for other tools to use even if we're not going
// to emit any code for it.
NewVD->setTSCSpec(TSCS);
} else if (getLangOpts().SYCLIsDevice) {
// While SYCL does not support TLS, emitting the diagnostic here
// prevents the compilation of header files with TLS declarations.
// When TLS objects are used in kernel code, they are diagnosed.
// We still need to mark the variable as TLS so it shows up in AST with
// proper storage class for other tools to use even if we're not going
// to emit any code for it.
NewVD->setTSCSpec(TSCS);
} else
Diag(D.getDeclSpec().getThreadStorageClassSpecLoc(),
diag::err_thread_unsupported);
Expand All @@ -7097,13 +7090,10 @@ NamedDecl *Sema::ActOnVariableDeclarator(

// Static variables declared inside SYCL device code must be const or
// constexpr
if (getLangOpts().SYCLIsDevice) {
if (getLangOpts().SYCLIsDevice)
if (SCSpec == DeclSpec::SCS_static && !R.isConstant(Context))
SYCLDiagIfDeviceCode(D.getIdentifierLoc(), diag::err_sycl_restrict)
<< Sema::KernelNonConstStaticDataVariable;
else if (NewVD->getTSCSpec() == DeclSpec::TSCS_thread_local)
SYCLDiagIfDeviceCode(D.getIdentifierLoc(), diag::err_thread_unsupported);
}

switch (D.getDeclSpec().getConstexprSpecifier()) {
case CSK_unspecified:
Expand Down Expand Up @@ -7190,13 +7180,18 @@ NamedDecl *Sema::ActOnVariableDeclarator(
// Handle attributes prior to checking for duplicates in MergeVarDecl
ProcessDeclAttributes(S, NewVD, D);

if (getLangOpts().CUDA || getLangOpts().OpenMPIsDevice) {
if (getLangOpts().CUDA || getLangOpts().OpenMPIsDevice ||
getLangOpts().SYCLIsDevice) {
if (EmitTLSUnsupportedError &&
((getLangOpts().CUDA && DeclAttrsMatchCUDAMode(getLangOpts(), NewVD)) ||
(getLangOpts().OpenMPIsDevice &&
OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(NewVD))))
Diag(D.getDeclSpec().getThreadStorageClassSpecLoc(),
diag::err_thread_unsupported);

if (EmitTLSUnsupportedError && getLangOpts().SYCLIsDevice)
SYCLDiagIfDeviceCode(D.getIdentifierLoc(), diag::err_thread_unsupported);

// CUDA B.2.5: "__shared__ and __constant__ variables have implied static
// storage [duration]."
if (SC == SC_None && S->getFnParent() != nullptr &&
Expand Down
2 changes: 1 addition & 1 deletion clang/test/SemaSYCL/prohibit-thread-local.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
// RUN: %clang_cc1 -fsycl -fsycl-is-device -verify -fsyntax-only %s
// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64 -verify -fsyntax-only %s

thread_local const int prohobit_ns_scope = 0;
thread_local int prohobit_ns_scope2 = 0;
Expand Down