clang/test/OpenMP/nvptx_prohibit_thread_local.cpp

   1 // RUN: %clang_cc1 -fopenmp -x c++ -triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-host.bc
   2 // 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-target-device -fopenmp-host-ir-file-path %t-host.bc -fsyntax-only
   3
   4 thread_local const int prohobit_ns_scope = 0;
   5 thread_local int prohobit_ns_scope2 = 0;
   6 thread_local const int allow_ns_scope = 0;
   7
   8 struct S {
   9   static const thread_local int prohibit_static_member;
  10   static thread_local int prohibit_static_member2;
  11 };
  12
  13 struct T {
  14   static const thread_local int allow_static_member;
  15 };
  16
  17 void foo() {
  18   // expected-error@+1{{thread-local storage is not supported for the current target}}
  19   thread_local const int prohibit_local = 0;
  20   // expected-error@+1{{thread-local storage is not supported for the current target}}
  21   thread_local int prohibit_local2;
  22 }
  23
  24 void bar() { thread_local int allow_local; }
  25
  26 void usage() {
  27   // expected-note@+1 {{called by}}
  28   foo();
  29   // expected-error@+1 {{thread-local storage is not supported for the current target}}
  30   (void)prohobit_ns_scope;
  31   // expected-error@+1 {{thread-local storage is not supported for the current target}}
  32   (void)prohobit_ns_scope2;
  33   // expected-error@+1 {{thread-local storage is not supported for the current target}}
  34   (void)S::prohibit_static_member;
  35   // expected-error@+1 {{thread-local storage is not supported for the current target}}
  36   (void)S::prohibit_static_member2;
  37 }
  38
  39 int main() {
  40   // expected-note@+2 2{{called by}}
  41 #pragma omp target
  42   usage();
  43   return 0;
  44 }