[SYCL][OpenMP] Implement thread-local storage restriction

[Fznamznon]

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.

[Fznamznon]

@jdoerfert , @ABataev , if OpenMP needs same diagnostic as well, I can generalize it between SYCL and OpenMP.

[erichkeane]

Give @ABataev and @jdoerfert a day or two to review before committing.

[riccibruno]

clang/lib/Sema/SemaExpr.cpp

216

Nit: The convention is auto *VD.

[Fznamznon]

Fixed code style.

[Fznamznon]

clang/lib/Sema/SemaExpr.cpp

216

Fixed, thanks!

[jdoerfert]

OpenMP has the same restriction (no surprise I guess). Thanks for the ping!

I think we do not emit diagnosis right now: https://godbolt.org/z/srDkXZ
I think we also should diagnose this the same way, though it might be beyond the scope of this patch: https://godbolt.org/z/rRZFi4

[Fznamznon]

Generalized diagnostic between SYCL and OpenMP.

[Fznamznon]

In D81641#2088446, @jdoerfert wrote:

OpenMP has the same restriction (no surprise I guess). Thanks for the ping!

I think we do not emit diagnosis right now: https://godbolt.org/z/srDkXZ
I think we also should diagnose this the same way, though it might be beyond the scope of this patch: https://godbolt.org/z/rRZFi4

Alright, now the first case that you pointed is diagnosed as well.
The second case is not diagnosed at all and I agree that it seems beyond the scope of this patch, because it looks like using of #pragma omp task untied doesn't trigger emission of deferred diagnostics.

[jdoerfert]

In D81641#2096114, @Fznamznon wrote:

In D81641#2088446, @jdoerfert wrote:

OpenMP has the same restriction (no surprise I guess). Thanks for the ping!

I think we do not emit diagnosis right now: https://godbolt.org/z/srDkXZ
I think we also should diagnose this the same way, though it might be beyond the scope of this patch: https://godbolt.org/z/rRZFi4

Alright, now the first case that you pointed is diagnosed as well.

Nice, thanks a lot!

The second case is not diagnosed at all and I agree that it seems beyond the scope of this patch, because it looks like using of #pragma omp task untied doesn't trigger emission of deferred diagnostics.

Yeah, we need to track untied tasks first. Not as part of this though.

LGTM :)

[Fznamznon]

Seems that test/OpenMP/nvptx_target_codegen.cpp is completely not formatted. If I apply suggestion from pre-merge checks, this will look like a big unrelated to this patch change and it will contradict with the whole file style.

[bader]

In D81641#2097433, @Fznamznon wrote:

Seems that test/OpenMP/nvptx_target_codegen.cpp is completely not formatted. If I apply suggestion from pre-merge checks, this will look like a big unrelated to this patch change and it will contradict with the whole file style.

I formatted this file here: https://github.com/llvm/llvm-project/commit/93cd4115799cefa698833ca7a2f1899243d94c77.
Could you rebase your patch, please?

[Fznamznon]

Rebased.

clang/lib/Sema/SemaDecl.cpp

if (DeclSpec::TSCS TSCS = D.getDeclSpec().getThreadStorageClassSpec()) {
     if (NewVD->hasLocalStorage() &&
         (SCSpec != DeclSpec::SCS_unspecified ||
          TSCS != DeclSpec::TSCS_thread_local ||
          !DC->isFunctionOrMethod()))
       Diag(D.getDeclSpec().getThreadStorageClassSpecLoc(),
            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.
         EmitTLSUnsupportedError = true;
         // 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);
if (getLangOpts().OpenCL) {
     deduceOpenCLAddressSpace(NewVD);
 
     diagnoseOpenCLTypes(S, *this, D, DC, NewVD->getType());
   }
 
   // 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 &&
         (NewVD->hasAttr<CUDASharedAttr>() ||
          NewVD->hasAttr<CUDAConstantAttr>())) {
       NewVD->setStorageClass(SC_Static);
     }
   }

clang/lib/Sema/SemaExpr.cpp

 /// referenced), false otherwise.
 ///
 bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef<SourceLocation> Locs,
                              const ObjCInterfaceDecl *UnknownObjCClass,
                              bool ObjCPropertyAccess,
                              bool AvoidPartialAvailabilityChecks,
                              ObjCInterfaceDecl *ClassReceiver) {
   SourceLocation Loc = Locs.front();
   if (getLangOpts().CPlusPlus && isa<FunctionDecl>(D)) {
     // If there were any diagnostics suppressed by template argument deduction,
     // emit them now.
     auto Pos = SuppressedDiagnostics.find(D->getCanonicalDecl());
     if (Pos != SuppressedDiagnostics.end()) {
       for (const PartialDiagnosticAt &Suppressed : Pos->second)
         Diag(Suppressed.first, Suppressed.second);
 
       // Clear out the list of suppressed diagnostics, so that we don't emit
bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef<SourceLocation> Locs,
 
   DiagnoseAvailabilityOfDecl(D, Locs, UnknownObjCClass, ObjCPropertyAccess,
                              AvoidPartialAvailabilityChecks, ClassReceiver);
 
   DiagnoseUnusedOfDecl(*this, D, Loc);
 
   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
     //   A local parameter shall only appear as an unevaluated operand
     //   (Clause 8) within the constraint-expression.
     Diag(Loc, diag::err_requires_expr_parameter_referenced_in_evaluated_context)
         << D;
     Diag(D->getLocation(), diag::note_entity_declared_at) << D;

clang/test/OpenMP/nvptx_prohibit_thread_local.cpp

+// 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;
+}

clang/test/OpenMP/nvptx_target_codegen.cpp

 // Test target codegen - host bc file has to be created first.
 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -disable-llvm-optzns | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64
 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - -disable-llvm-optzns | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
 // RUN: %clang_cc1 -verify -fopenmp -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - -disable-llvm-optzns | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
 // expected-no-diagnostics
 
 #ifndef HEADER
int foo(int n) {
   //
   // CHECK: [[BAR_PARALLEL]]
   // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
   // CHECK: br label {{%?}}[[AWAIT_WORK]]
   //
   // 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*
   // CHECK-DAG: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
   // CHECK-DAG: [[NTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
   // CHECK-DAG: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
   // CHECK-DAG: [[TH_LIMIT:%.+]] = sub nuw i32 [[NTH]], [[WS]]
   // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[TH_LIMIT]]
int foo(int n) {
   // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
   // CHECK: br label {{%?}}[[EXIT]]
   //
   // CHECK: [[EXIT]]
   // CHECK: ret void
   #pragma omp target if(1)
   {
     aa += 1;
-    id = aa;
+    aa += 2;
   }
 
   // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l310}}_worker()
   // CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
   // CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
   // CHECK: store i8* null, i8** [[OMP_WORK_FN]],
   // CHECK: store i8 0, i8* [[OMP_EXEC_STATUS]],
   // CHECK: br label {{%?}}[[AWAIT_WORK:.+]]

clang/test/SemaSYCL/prohibit-thread-local.cpp

+// 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;
+}