This revision was automatically updated to reflect the committed changes.
Closed by commit rG0bdcd95bf20f: [SYCL][OpenMP] Implement thread-local storage
restriction (authored by Fznamznon, committed by bader).
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/OpenMP/nvptx_prohibit_thread_local.cpp
clang/test/OpenMP/nvptx_target_codegen.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/test/OpenMP/nvptx_target_codegen.cpp
===================================================================
--- clang/test/OpenMP/nvptx_target_codegen.cpp
+++ clang/test/OpenMP/nvptx_target_codegen.cpp
@@ -160,7 +160,7 @@
// CHECK: [[EXIT]]
// CHECK: ret void
-// CHECK: define {{.*}}void [[T2:@__omp_offloading_.+foo.+l200]](i[[SZ:32|64]] [[ARG1:%[a-zA-Z_]+]], i[[SZ:32|64]] [[ID:%[a-zA-Z_]+]])
+// CHECK: define {{.*}}void [[T2:@__omp_offloading_.+foo.+l200]](i[[SZ:32|64]] [[ARG1:%[a-zA-Z_]+]])
// CHECK: [[AA_ADDR:%.+]] = alloca i[[SZ]],
// CHECK: store i[[SZ]] [[ARG1]], i[[SZ]]* [[AA_ADDR]],
// CHECK: [[AA_CADDR:%.+]] = bitcast i[[SZ]]* [[AA_ADDR]] to i16*
@@ -200,7 +200,7 @@
#pragma omp target if (1)
{
aa += 1;
- id = aa;
+ aa += 2;
}
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l310}}_worker()
Index: clang/test/OpenMP/nvptx_prohibit_thread_local.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/nvptx_prohibit_thread_local.cpp
@@ -0,0 +1,44 @@
+// RUN: %clang_cc1 -fopenmp -x c++ -triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -fsyntax-only
+
+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;
+}
+
+int main() {
+ // expected-note@+2 2{{called by}}
+#pragma omp target
+ usage();
+ return 0;
+}
Index: clang/lib/Sema/SemaExpr.cpp
===================================================================
--- clang/lib/Sema/SemaExpr.cpp
+++ clang/lib/Sema/SemaExpr.cpp
@@ -355,10 +355,16 @@
diagnoseUseOfInternalDeclInInlineFunction(*this, D, Loc);
- if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice))
+ if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice)) {
if (const auto *VD = dyn_cast<ValueDecl>(D))
checkDeviceDecl(VD, Loc);
+ if (!Context.getTargetInfo().isTLSSupported())
+ if (const auto *VD = dyn_cast<VarDecl>(D))
+ if (VD->getTLSKind() != VarDecl::TLS_None)
+ targetDiag(*Locs.begin(), diag::err_thread_unsupported);
+ }
+
if (isa<ParmVarDecl>(D) && isa<RequiresExprBodyDecl>(D->getDeclContext()) &&
!isUnevaluatedContext()) {
// C++ [expr.prim.req.nested] p3
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,18 @@
// 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 &&
+ (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice)))
+ targetDiag(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
