Fznamznon updated this revision to Diff 270135.
Fznamznon added a comment.
Fixed code style.
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D81641/new/
https://reviews.llvm.org/D81641
Files:
clang/lib/Sema/SemaDecl.cpp
clang/lib/Sema/SemaExpr.cpp
clang/test/SemaSYCL/prohibit-thread-local.cpp
Index: clang/test/SemaSYCL/prohibit-thread-local.cpp
===================================================================
--- /dev/null
+++ clang/test/SemaSYCL/prohibit-thread-local.cpp
@@ -0,0 +1,48 @@
+// 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;
+thread_local const int allow_ns_scope = 0;
+
+struct S {
+ static const thread_local int prohibit_static_member;
+ static thread_local int prohibit_static_member2;
+};
+
+struct T {
+ static const thread_local int allow_static_member;
+};
+
+void foo() {
+ // expected-error@+1{{thread-local storage is not supported for the current target}}
+ thread_local const int prohibit_local = 0;
+ // expected-error@+1{{thread-local storage is not supported for the current target}}
+ thread_local int prohibit_local2;
+}
+
+void bar() { thread_local int allow_local; }
+
+void usage() {
+ // expected-note@+1 {{called by}}
+ foo();
+ // expected-error@+1 {{thread-local storage is not supported for the current target}}
+ (void)prohobit_ns_scope;
+ // expected-error@+1 {{thread-local storage is not supported for the current target}}
+ (void)prohobit_ns_scope2;
+ // expected-error@+1 {{thread-local storage is not supported for the current target}}
+ (void)S::prohibit_static_member;
+ // expected-error@+1 {{thread-local storage is not supported for the current target}}
+ (void)S::prohibit_static_member2;
+}
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel))
+// expected-note@+2 2{{called by}}
+void
+kernel_single_task(Func kernelFunc) { kernelFunc(); }
+
+int main() {
+ // expected-note@+1 2{{called by}}
+ kernel_single_task<class fake_kernel>([]() { usage(); });
+ return 0;
+}
Index: clang/lib/Sema/SemaExpr.cpp
===================================================================
--- clang/lib/Sema/SemaExpr.cpp
+++ clang/lib/Sema/SemaExpr.cpp
@@ -212,6 +212,11 @@
bool ObjCPropertyAccess,
bool AvoidPartialAvailabilityChecks,
ObjCInterfaceDecl *ClassReceiver) {
+ if (getLangOpts().SYCLIsDevice)
+ if (auto *VD = dyn_cast<VarDecl>(D))
+ if (VD->getTLSKind() != VarDecl::TLS_None)
+ SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_thread_unsupported);
+
SourceLocation Loc = Locs.front();
if (getLangOpts().CPlusPlus && isa<FunctionDecl>(D)) {
// If there were any diagnostics suppressed by template argument deduction,
Index: clang/lib/Sema/SemaDecl.cpp
===================================================================
--- clang/lib/Sema/SemaDecl.cpp
+++ clang/lib/Sema/SemaDecl.cpp
@@ -7077,7 +7077,8 @@
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.
@@ -7179,13 +7180,17 @@
// 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 &&
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
