[clang] 0bdcd95 - [SYCL][OpenMP] Implement thread-local storage restriction
Alexey Bader via cfe-commits
cfe-commits at lists.llvm.org
Wed Jun 17 04:38:05 PDT 2020
Author: Mariya Podchishchaeva
Date: 2020-06-17T14:36:00+03:00
New Revision: 0bdcd95bf20f159a2512aff1ef032bec52039bf6
URL: https://github.com/llvm/llvm-project/commit/0bdcd95bf20f159a2512aff1ef032bec52039bf6
DIFF: https://github.com/llvm/llvm-project/commit/0bdcd95bf20f159a2512aff1ef032bec52039bf6.diff
LOG: [SYCL][OpenMP] Implement thread-local storage restriction
Summary:
SYCL and OpenMP prohibits thread local storage in device code,
so this commit ensures that error is emitted for device code and not
emitted for host code when host target supports it.
Reviewers: jdoerfert, erichkeane, bader
Reviewed By: jdoerfert, erichkeane
Subscribers: guansong, riccibruno, ABataev, yaxunl, ebevhan, Anastasia, sstefan1, cfe-commits
Tags: #clang
Differential Revision: https://reviews.llvm.org/D81641
Added:
clang/test/OpenMP/nvptx_prohibit_thread_local.cpp
clang/test/SemaSYCL/prohibit-thread-local.cpp
Modified:
clang/lib/Sema/SemaDecl.cpp
clang/lib/Sema/SemaExpr.cpp
clang/test/OpenMP/nvptx_target_codegen.cpp
Removed:
################################################################################
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 2bf16d138d5a..80469e3bedbe 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -7077,7 +7077,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.
@@ -7179,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 &&
+ (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 &&
diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp
index 66a2ec1fe9dc..ffc72140dcf4 100644
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/clang/lib/Sema/SemaExpr.cpp
@@ -355,10 +355,16 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef<SourceLocation> Locs,
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
diff --git a/clang/test/OpenMP/nvptx_prohibit_thread_local.cpp b/clang/test/OpenMP/nvptx_prohibit_thread_local.cpp
new file mode 100644
index 000000000000..b84918e528cb
--- /dev/null
+++ b/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 at +1{{thread-local storage is not supported for the current target}}
+ thread_local const int prohibit_local = 0;
+ // expected-error at +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 at +1 {{called by}}
+ foo();
+ // expected-error at +1 {{thread-local storage is not supported for the current target}}
+ (void)prohobit_ns_scope;
+ // expected-error at +1 {{thread-local storage is not supported for the current target}}
+ (void)prohobit_ns_scope2;
+ // expected-error at +1 {{thread-local storage is not supported for the current target}}
+ (void)S::prohibit_static_member;
+ // expected-error at +1 {{thread-local storage is not supported for the current target}}
+ (void)S::prohibit_static_member2;
+}
+
+int main() {
+ // expected-note at +2 2{{called by}}
+#pragma omp target
+ usage();
+ return 0;
+}
diff --git a/clang/test/OpenMP/nvptx_target_codegen.cpp b/clang/test/OpenMP/nvptx_target_codegen.cpp
index d615b8536c48..91f31185d8c1 100644
--- a/clang/test/OpenMP/nvptx_target_codegen.cpp
+++ b/clang/test/OpenMP/nvptx_target_codegen.cpp
@@ -160,7 +160,7 @@ int foo(int n) {
// 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 @@ int foo(int n) {
#pragma omp target if (1)
{
aa += 1;
- id = aa;
+ aa += 2;
}
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l310}}_worker()
diff --git a/clang/test/SemaSYCL/prohibit-thread-local.cpp b/clang/test/SemaSYCL/prohibit-thread-local.cpp
new file mode 100644
index 000000000000..4fd113626ea7
--- /dev/null
+++ b/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 at +1{{thread-local storage is not supported for the current target}}
+ thread_local const int prohibit_local = 0;
+ // expected-error at +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 at +1 {{called by}}
+ foo();
+ // expected-error at +1 {{thread-local storage is not supported for the current target}}
+ (void)prohobit_ns_scope;
+ // expected-error at +1 {{thread-local storage is not supported for the current target}}
+ (void)prohobit_ns_scope2;
+ // expected-error at +1 {{thread-local storage is not supported for the current target}}
+ (void)S::prohibit_static_member;
+ // expected-error at +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 at +2 2{{called by}}
+void
+kernel_single_task(Func kernelFunc) { kernelFunc(); }
+
+int main() {
+ // expected-note at +1 2{{called by}}
+ kernel_single_task<class fake_kernel>([]() { usage(); });
+ return 0;
+}
More information about the cfe-commits
mailing list