nvptx_target_cuda_mode_messages.cpp
2.83 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
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
// RUN: %clang_cc1 -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-cuda-mode -fopenmp-host-ir-file-path %t-ppc-host.bc -o -
template <typename tx, typename ty>
struct TT {
tx X;
ty Y;
};
int foo(int n, double *ptr) {
int a = 0;
short aa = 0;
float b[10];
double c[5][10];
TT<long long, char> d;
#pragma omp target firstprivate(a) map(tofrom: b) // expected-note 2 {{defined as threadprivate or thread local}}
{
int c; // expected-note {{defined as threadprivate or thread local}}
#pragma omp parallel shared(a, b, c, aa) // expected-error 3 {{threadprivate or thread local variable cannot be shared}}
b[a] = a;
#pragma omp parallel for
for (int i = 0; i < 10; ++i) // expected-note {{defined as threadprivate or thread local}}
#pragma omp parallel shared(i) // expected-error {{threadprivate or thread local variable cannot be shared}}
++i;
}
#pragma omp target map(aa, b, c, d)
{
int e; // expected-note {{defined as threadprivate or thread local}}
#pragma omp parallel private(b, e) // expected-error {{threadprivate or thread local variable cannot be private}}
{
aa += 1;
b[2] = 1.0;
c[1][2] = 1.0;
d.X = 1;
d.Y = 1;
}
}
#pragma omp target private(ptr)
{
ptr[0]++;
}
return a;
}
template <typename tx>
tx ftemplate(int n) {
tx a = 0;
tx b[10];
#pragma omp target reduction(+ \
: a, b) // expected-note {{defined as threadprivate or thread local}}
{
int e; // expected-note {{defined as threadprivate or thread local}}
#pragma omp parallel shared(a, e) // expected-error 2 {{threadprivate or thread local variable cannot be shared}}
a += 1;
b[2] += 1;
}
return a;
}
static int fstatic(int n) {
int a = 0;
char aaa = 0;
int b[10];
#pragma omp target firstprivate(a, aaa, b)
{
a += 1;
aaa += 1;
b[2] += 1;
}
return a;
}
struct S1 {
double a;
int r1(int n) {
int b = n + 1;
#pragma omp target firstprivate(b) // expected-note {{defined as threadprivate or thread local}}
{
int c; // expected-note {{defined as threadprivate or thread local}}
#pragma omp parallel shared(b, c) // expected-error 2 {{threadprivate or thread local variable cannot be shared}}
this->a = (double)b + 1.5;
}
return (int)b;
}
};
int bar(int n, double *ptr) {
int a = 0;
a += foo(n, ptr);
S1 S;
a += S.r1(n);
a += fstatic(n);
a += ftemplate<int>(n); // expected-note {{in instantiation of function template specialization 'ftemplate<int>' requested here}}
return a;
}