nvptx_prohibit_thread_local.cpp
1.59 KB
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
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;
}