1 // RUN: %clang_cc1 -verify -fopenmp -std=c++11 -ast-print %s | FileCheck %s
2 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -emit-pch -o %t %s
3 // RUN: %clang_cc1 -fopenmp -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
4
5 // RUN: %clang_cc1 -verify -fopenmp-simd -std=c++11 -ast-print %s | FileCheck %s
6 // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -emit-pch -o %t %s
7 // RUN: %clang_cc1 -fopenmp-simd -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
8 // expected-no-diagnostics
9
10 #ifndef HEADER
11 #define HEADER
12
foo()13 void foo() {}
14
15 struct ST {
16 int *a;
17 };
18 typedef int arr[10];
19 typedef ST STarr[10];
20 struct SA {
21 const int da[5] = { 0 };
22 ST g[10];
23 STarr &rg = g;
24 int i;
25 int &j = i;
26 int *k = &j;
27 int *&z = k;
28 int aa[10];
29 arr &raa = aa;
funcSA30 void func(int arg) {
31 #pragma omp target parallel for is_device_ptr(k)
32 for (int i=0; i<100; i++) foo();
33
34 #pragma omp target parallel for is_device_ptr(z)
35 for (int i=0; i<100; i++) foo();
36
37 #pragma omp target parallel for is_device_ptr(aa) // OK
38 for (int i=0; i<100; i++) foo();
39
40 #pragma omp target parallel for is_device_ptr(raa) // OK
41 for (int i=0; i<100; i++) foo();
42
43 #pragma omp target parallel for is_device_ptr(g) // OK
44 for (int i=0; i<100; i++) foo();
45
46 #pragma omp target parallel for is_device_ptr(rg) // OK
47 for (int i=0; i<100; i++) foo();
48
49 #pragma omp target parallel for is_device_ptr(da) // OK
50 for (int i=0; i<100; i++) foo();
51
52 return;
53 }
54 };
55 // CHECK: struct SA
56 // CHECK-NEXT: const int da[5] = {0};
57 // CHECK-NEXT: ST g[10];
58 // CHECK-NEXT: STarr &rg = this->g;
59 // CHECK-NEXT: int i;
60 // CHECK-NEXT: int &j = this->i;
61 // CHECK-NEXT: int *k = &this->j;
62 // CHECK-NEXT: int *&z = this->k;
63 // CHECK-NEXT: int aa[10];
64 // CHECK-NEXT: arr &raa = this->aa;
65 // CHECK-NEXT: func(
66 // CHECK-NEXT: #pragma omp target parallel for is_device_ptr(this->k){{$}}
67 // CHECK-NEXT: for (int i = 0; i < 100; i++)
68 // CHECK-NEXT: foo();
69 // CHECK-NEXT: #pragma omp target parallel for is_device_ptr(this->z)
70 // CHECK-NEXT: for (int i = 0; i < 100; i++)
71 // CHECK-NEXT: foo();
72 // CHECK-NEXT: #pragma omp target parallel for is_device_ptr(this->aa)
73 // CHECK-NEXT: for (int i = 0; i < 100; i++)
74 // CHECK-NEXT: foo();
75 // CHECK-NEXT: #pragma omp target parallel for is_device_ptr(this->raa)
76 // CHECK-NEXT: for (int i = 0; i < 100; i++)
77 // CHECK-NEXT: foo();
78 // CHECK-NEXT: #pragma omp target parallel for is_device_ptr(this->g)
79 // CHECK-NEXT: for (int i = 0; i < 100; i++)
80 // CHECK-NEXT: foo();
81 // CHECK-NEXT: #pragma omp target parallel for is_device_ptr(this->rg)
82 // CHECK-NEXT: for (int i = 0; i < 100; i++)
83 // CHECK-NEXT: foo();
84 // CHECK-NEXT: #pragma omp target parallel for is_device_ptr(this->da)
85 // CHECK-NEXT: for (int i = 0; i < 100; i++)
86 // CHECK-NEXT: foo();
87
88 struct SB {
89 unsigned A;
90 unsigned B;
91 float Arr[100];
92 float *Ptr;
fooSB93 float *foo() {
94 return &Arr[0];
95 }
96 };
97
98 struct SC {
99 unsigned A : 2;
100 unsigned B : 3;
101 unsigned C;
102 unsigned D;
103 float Arr[100];
104 SB S;
105 SB ArrS[100];
106 SB *PtrS;
107 SB *&RPtrS;
108 float *Ptr;
109
SCSC110 SC(SB *&_RPtrS) : RPtrS(_RPtrS) {}
111 };
112
113 union SD {
114 unsigned A;
115 float B;
116 };
117
118 struct S1;
119 extern S1 a;
120 class S2 {
121 mutable int a;
122 public:
S2()123 S2():a(0) { }
S2(S2 & s2)124 S2(S2 &s2):a(s2.a) { }
125 static float S2s;
126 static const float S2sc;
127 };
128 const float S2::S2sc = 0;
129 const S2 b;
130 const S2 ba[5];
131 class S3 {
132 int a;
133 public:
S3()134 S3():a(0) { }
S3(S3 & s3)135 S3(S3 &s3):a(s3.a) { }
136 };
137 const S3 c;
138 const S3 ca[5];
139 extern const int f;
140 class S4 {
141 int a;
142 S4();
143 S4(const S4 &s4);
144 public:
S4(int v)145 S4(int v):a(v) { }
146 };
147 class S5 {
148 int a;
S5()149 S5():a(0) {}
S5(const S5 & s5)150 S5(const S5 &s5):a(s5.a) { }
151 public:
S5(int v)152 S5(int v):a(v) { }
153 };
154
155 S3 h;
156 #pragma omp threadprivate(h)
157
158 typedef struct {
159 int a;
160 } S6;
161
162 template <typename T>
tmain(T argc)163 T tmain(T argc) {
164 const T da[5] = { 0 };
165 S6 h[10];
166 auto &rh = h;
167 T i;
168 T &j = i;
169 T *k = &j;
170 T *&z = k;
171 T aa[10];
172 auto &raa = aa;
173 #pragma omp target parallel for is_device_ptr(k)
174 for (int i=0; i<100; i++) foo();
175 #pragma omp target parallel for is_device_ptr(z)
176 for (int i=0; i<100; i++) foo();
177 #pragma omp target parallel for is_device_ptr(aa)
178 for (int i=0; i<100; i++) foo();
179 #pragma omp target parallel for is_device_ptr(raa)
180 for (int i=0; i<100; i++) foo();
181 #pragma omp target parallel for is_device_ptr(h)
182 for (int i=0; i<100; i++) foo();
183 #pragma omp target parallel for is_device_ptr(rh)
184 for (int i=0; i<100; i++) foo();
185 #pragma omp target parallel for is_device_ptr(da)
186 for (int i=0; i<100; i++) foo();
187 return 0;
188 }
189
190 // CHECK: template<> int tmain<int>(int argc) {
191 // CHECK-NEXT: const int da[5] = {0};
192 // CHECK-NEXT: S6 h[10];
193 // CHECK-NEXT: auto &rh = h;
194 // CHECK-NEXT: int i;
195 // CHECK-NEXT: int &j = i;
196 // CHECK-NEXT: int *k = &j;
197 // CHECK-NEXT: int *&z = k;
198 // CHECK-NEXT: int aa[10];
199 // CHECK-NEXT: auto &raa = aa;
200 // CHECK-NEXT: #pragma omp target parallel for is_device_ptr(k)
201 // CHECK-NEXT: for (int i = 0; i < 100; i++)
202 // CHECK-NEXT: foo();
203 // CHECK-NEXT: #pragma omp target parallel for is_device_ptr(z)
204 // CHECK-NEXT: for (int i = 0; i < 100; i++)
205 // CHECK-NEXT: foo();
206 // CHECK-NEXT: #pragma omp target parallel for is_device_ptr(aa)
207 // CHECK-NEXT: for (int i = 0; i < 100; i++)
208 // CHECK-NEXT: foo();
209 // CHECK-NEXT: #pragma omp target parallel for is_device_ptr(raa)
210 // CHECK-NEXT: for (int i = 0; i < 100; i++)
211 // CHECK-NEXT: foo();
212 // CHECK-NEXT: #pragma omp target parallel for is_device_ptr(h)
213 // CHECK-NEXT: for (int i = 0; i < 100; i++)
214 // CHECK-NEXT: foo();
215 // CHECK-NEXT: #pragma omp target parallel for is_device_ptr(rh)
216 // CHECK-NEXT: for (int i = 0; i < 100; i++)
217 // CHECK-NEXT: foo();
218 // CHECK-NEXT: #pragma omp target parallel for is_device_ptr(da)
219 // CHECK-NEXT: for (int i = 0; i < 100; i++)
220 // CHECK-NEXT: foo();
221
222 // CHECK: template<> int *tmain<int *>(int *argc) {
223 // CHECK-NEXT: int *const da[5] = {0};
224 // CHECK-NEXT: S6 h[10];
225 // CHECK-NEXT: auto &rh = h;
226 // CHECK-NEXT: int *i;
227 // CHECK-NEXT: int *&j = i;
228 // CHECK-NEXT: int **k = &j;
229 // CHECK-NEXT: int **&z = k;
230 // CHECK-NEXT: int *aa[10];
231 // CHECK-NEXT: auto &raa = aa;
232 // CHECK-NEXT: #pragma omp target parallel for is_device_ptr(k)
233 // CHECK-NEXT: for (int i = 0; i < 100; i++)
234 // CHECK-NEXT: foo();
235 // CHECK-NEXT: #pragma omp target parallel for is_device_ptr(z)
236 // CHECK-NEXT: for (int i = 0; i < 100; i++)
237 // CHECK-NEXT: foo();
238 // CHECK-NEXT: #pragma omp target parallel for is_device_ptr(aa)
239 // CHECK-NEXT: for (int i = 0; i < 100; i++)
240 // CHECK-NEXT: foo();
241 // CHECK-NEXT: #pragma omp target parallel for is_device_ptr(raa)
242 // CHECK-NEXT: for (int i = 0; i < 100; i++)
243 // CHECK-NEXT: foo();
244 // CHECK-NEXT: #pragma omp target parallel for is_device_ptr(h)
245 // CHECK-NEXT: for (int i = 0; i < 100; i++)
246 // CHECK-NEXT: foo();
247 // CHECK-NEXT: #pragma omp target parallel for is_device_ptr(rh)
248 // CHECK-NEXT: for (int i = 0; i < 100; i++)
249 // CHECK-NEXT: foo();
250 // CHECK-NEXT: #pragma omp target parallel for is_device_ptr(da)
251 // CHECK-NEXT: for (int i = 0; i < 100; i++)
252 // CHECK-NEXT: foo();
253
254 // CHECK-LABEL: int main(int argc, char **argv) {
main(int argc,char ** argv)255 int main(int argc, char **argv) {
256 const int da[5] = { 0 };
257 S6 h[10];
258 auto &rh = h;
259 int i;
260 int &j = i;
261 int *k = &j;
262 int *&z = k;
263 int aa[10];
264 auto &raa = aa;
265 // CHECK-NEXT: const int da[5] = {0};
266 // CHECK-NEXT: S6 h[10];
267 // CHECK-NEXT: auto &rh = h;
268 // CHECK-NEXT: int i;
269 // CHECK-NEXT: int &j = i;
270 // CHECK-NEXT: int *k = &j;
271 // CHECK-NEXT: int *&z = k;
272 // CHECK-NEXT: int aa[10];
273 // CHECK-NEXT: auto &raa = aa;
274 #pragma omp target parallel for is_device_ptr(k)
275 for (int i=0; i<100; i++) foo();
276 // CHECK-NEXT: #pragma omp target parallel for is_device_ptr(k)
277 // CHECK-NEXT: for (int i = 0; i < 100; i++)
278 // CHECK-NEXT: foo();
279
280 #pragma omp target parallel for is_device_ptr(z)
281 for (int i=0; i<100; i++) foo();
282 // CHECK-NEXT: #pragma omp target parallel for is_device_ptr(z)
283 // CHECK-NEXT: for (int i = 0; i < 100; i++)
284 // CHECK-NEXT: foo();
285
286 #pragma omp target parallel for is_device_ptr(aa)
287 for (int i=0; i<100; i++) foo();
288 // CHECK-NEXT: #pragma omp target parallel for is_device_ptr(aa)
289 // CHECK-NEXT: for (int i = 0; i < 100; i++)
290 // CHECK-NEXT: foo();
291
292 #pragma omp target parallel for is_device_ptr(raa)
293 for (int i=0; i<100; i++) foo();
294 // CHECK-NEXT: #pragma omp target parallel for is_device_ptr(raa)
295 // CHECK-NEXT: for (int i = 0; i < 100; i++)
296 // CHECK-NEXT: foo();
297
298 #pragma omp target parallel for is_device_ptr(h)
299 for (int i=0; i<100; i++) foo();
300 // CHECK-NEXT: #pragma omp target parallel for is_device_ptr(h)
301 // CHECK-NEXT: for (int i = 0; i < 100; i++)
302 // CHECK-NEXT: foo();
303
304 #pragma omp target parallel for is_device_ptr(rh)
305 for (int i=0; i<100; i++) foo();
306 // CHECK-NEXT: #pragma omp target parallel for is_device_ptr(rh)
307 // CHECK-NEXT: for (int i = 0; i < 100; i++)
308 // CHECK-NEXT: foo();
309
310 #pragma omp target parallel for is_device_ptr(da)
311 for (int i=0; i<100; i++) foo();
312 // CHECK-NEXT: #pragma omp target parallel for is_device_ptr(da)
313 // CHECK-NEXT: for (int i = 0; i < 100; i++)
314 // CHECK-NEXT: foo();
315 return tmain<int>(argc) + *tmain<int *>(&argc);
316 }
317
318
319 #endif
320