From: Mariya Podchishchaeva Date: Wed, 17 Jun 2020 11:31:38 +0000 (+0300) Subject: [SYCL][OpenMP] Implement thread-local storage restriction X-Git-Tag: llvmorg-12-init~2814 X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=0bdcd95bf20f159a2512aff1ef032bec52039bf6;p=platform%2Fupstream%2Fllvm.git [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 --- diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 2bf16d1..80469e3 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 66a2ec1..ffc7214 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -355,10 +355,16 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef 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(D)) checkDeviceDecl(VD, Loc); + if (!Context.getTargetInfo().isTLSSupported()) + if (const auto *VD = dyn_cast(D)) + if (VD->getTLSKind() != VarDecl::TLS_None) + targetDiag(*Locs.begin(), diag::err_thread_unsupported); + } + if (isa(D) && isa(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 0000000..b84918e --- /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@+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; +} diff --git a/clang/test/OpenMP/nvptx_target_codegen.cpp b/clang/test/OpenMP/nvptx_target_codegen.cpp index d615b85..91f3118 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 0000000..4fd1136 --- /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@+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 +__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([]() { usage(); }); + return 0; +}