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-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 
foo()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 
bar()24 void bar() { thread_local int allow_local; }
25 
usage()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 
main()39 int main() {
40   // expected-note@+2 2{{called by}}
41 #pragma omp target
42   usage();
43   return 0;
44 }
45