1 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -ast-print %s -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP45
2 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -emit-pch -o %t %s -Wno-openmp-mapping
3 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP45
4 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -DOMP5 -ast-print %s -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP50
5 // RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -DOMP5 -x c++ -std=c++11 -emit-pch -o %t %s -Wno-openmp-mapping
6 // RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -DOMP5 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP50
7
8 // RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -ast-print %s -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP45
9 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -emit-pch -o %t %s -Wno-openmp-mapping
10 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP45
11 // RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=50 -DOMP5 -ast-print %s -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP50
12 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -DOMP5 -x c++ -std=c++11 -emit-pch -o %t %s -Wno-openmp-mapping
13 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -DOMP5 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP50
14 // expected-no-diagnostics
15
16 #ifndef HEADER
17 #define HEADER
18
19 typedef void **omp_allocator_handle_t;
20 extern const omp_allocator_handle_t omp_null_allocator;
21 extern const omp_allocator_handle_t omp_default_mem_alloc;
22 extern const omp_allocator_handle_t omp_large_cap_mem_alloc;
23 extern const omp_allocator_handle_t omp_const_mem_alloc;
24 extern const omp_allocator_handle_t omp_high_bw_mem_alloc;
25 extern const omp_allocator_handle_t omp_low_lat_mem_alloc;
26 extern const omp_allocator_handle_t omp_cgroup_mem_alloc;
27 extern const omp_allocator_handle_t omp_pteam_mem_alloc;
28 extern const omp_allocator_handle_t omp_thread_mem_alloc;
29
foo()30 void foo() {}
31
32 struct S {
SS33 S(): a(0) {}
SS34 S(int v) : a(v) {}
35 int a;
36 typedef int type;
37 };
38
39 template <typename T>
40 class S7 : public T {
41 protected:
42 T a;
S7()43 S7() : a(0) {}
44
45 public:
S7(typename T::type v)46 S7(typename T::type v) : a(v) {
47 #pragma omp target simd private(a) private(this->a) private(T::a)
48 for (int k = 0; k < a.a; ++k)
49 ++this->a.a;
50 }
operator =(S7 & s)51 S7 &operator=(S7 &s) {
52 #pragma omp target simd private(a) private(this->a)
53 for (int k = 0; k < s.a.a; ++k)
54 ++s.a.a;
55 return *this;
56 }
57 };
58
59 // CHECK: #pragma omp target simd private(this->a) private(this->a) private(T::a){{$}}
60 // CHECK: #pragma omp target simd private(this->a) private(this->a)
61 // CHECK: #pragma omp target simd private(this->a) private(this->a) private(this->S::a)
62
63 class S8 : public S7<S> {
S8()64 S8() {}
65
66 public:
S8(int v)67 S8(int v) : S7<S>(v){
68 #pragma omp target simd private(a) private(this->a) private(S7<S>::a)
69 for (int k = 0; k < a.a; ++k)
70 ++this->a.a;
71 }
operator =(S8 & s)72 S8 &operator=(S8 &s) {
73 #pragma omp target simd private(a) private(this->a)
74 for (int k = 0; k < s.a.a; ++k)
75 ++s.a.a;
76 return *this;
77 }
78 };
79
80 // CHECK: #pragma omp target simd private(this->a) private(this->a) private(this->S7<S>::a)
81 // CHECK: #pragma omp target simd private(this->a) private(this->a)
82
83 template <class T, int N>
84 T tmain(T argc, T *argv) {
85 T b = argc, c, d, e, f, h;
86 T arr[N][10], arr1[N];
87 T i, j;
88 T s;
89 static T a;
90 // CHECK: static T a;
91 static T g;
92 const T clen = 5;
93 // CHECK: T clen = 5;
94 T *p;
95 #pragma omp threadprivate(g)
96 #pragma omp target simd linear(a) allocate(a)
97 // CHECK: #pragma omp target simd linear(a) allocate(a)
98 for (T i = 0; i < 2; ++i)
99 a = 2;
100 // CHECK-NEXT: for (T i = 0; i < 2; ++i)
101 // CHECK-NEXT: a = 2;
102 #pragma omp target simd allocate(f) private(argc, b), firstprivate(c, d), lastprivate(d, f) collapse(N) if (target :argc) reduction(+ : h)
103 for (int i = 0; i < 2; ++i)
104 for (int j = 0; j < 2; ++j)
105 for (int j = 0; j < 2; ++j)
106 for (int j = 0; j < 2; ++j)
107 for (int j = 0; j < 2; ++j)
108 foo();
109 // CHECK-NEXT: #pragma omp target simd allocate(f) private(argc,b) firstprivate(c,d) lastprivate(d,f) collapse(N) if(target: argc) reduction(+: h)
110 // CHECK-NEXT: for (int i = 0; i < 2; ++i)
111 // CHECK-NEXT: for (int j = 0; j < 2; ++j)
112 // CHECK-NEXT: for (int j = 0; j < 2; ++j)
113 // CHECK-NEXT: for (int j = 0; j < 2; ++j)
114 // CHECK-NEXT: for (int j = 0; j < 2; ++j)
115 // CHECK-NEXT: foo();
116
117 #pragma omp target simd private(argc,b) firstprivate(argv) if(target:argc > 0) reduction(+:c, arr1[argc]) reduction(max:e, arr[:N][0:10])
118 for (T i = 0; i < 2; ++i) {}
119 // CHECK: #pragma omp target simd private(argc,b) firstprivate(argv) if(target: argc > 0) reduction(+: c,arr1[argc]) reduction(max: e,arr[:N][0:10])
120 // CHECK-NEXT: for (T i = 0; i < 2; ++i) {
121 // CHECK-NEXT: }
122
123 #pragma omp target simd if(N) reduction(^:e, f, arr[0:N][:argc]) reduction(&& : h)
124 for (T i = 0; i < 2; ++i) {}
125 // CHECK: #pragma omp target simd if(N) reduction(^: e,f,arr[0:N][:argc]) reduction(&&: h)
126 // CHECK-NEXT: for (T i = 0; i < 2; ++i) {
127 // CHECK-NEXT: }
128
129 #ifdef OMP5
130 #pragma omp target simd if(target:argc > 0) if (simd:argc) allocate(omp_default_mem_alloc:f) private(f) uses_allocators(omp_default_mem_alloc)
131 #else
132 #pragma omp target simd if(target:argc > 0)
133 #endif // OMP5
134 for (T i = 0; i < 2; ++i) {}
135 // OMP45: #pragma omp target simd if(target: argc > 0)
136 // OMP50: #pragma omp target simd if(target: argc > 0) if(simd: argc) allocate(omp_default_mem_alloc: f) private(f) uses_allocators(omp_default_mem_alloc)
137 // CHECK-NEXT: for (T i = 0; i < 2; ++i) {
138 // CHECK-NEXT: }
139
140 #pragma omp target simd if(N)
141 for (T i = 0; i < 2; ++i) {}
142 // CHECK: #pragma omp target simd if(N)
143 // CHECK-NEXT: for (T i = 0; i < 2; ++i) {
144 // CHECK-NEXT: }
145
146 #pragma omp target simd map(i)
147 for (T i = 0; i < 2; ++i) {}
148 // CHECK: #pragma omp target simd map(tofrom: i)
149 // CHECK-NEXT: for (T i = 0; i < 2; ++i) {
150 // CHECK-NEXT: }
151
152 #pragma omp target simd map(arr1[0:10], i)
153 for (T i = 0; i < 2; ++i) {}
154 // CHECK: #pragma omp target simd map(tofrom: arr1[0:10],i)
155 // CHECK-NEXT: for (T i = 0; i < 2; ++i) {
156 // CHECK-NEXT: }
157
158 #pragma omp target simd map(to: i) map(from: j)
159 for (T i = 0; i < 2; ++i) {}
160 // CHECK: #pragma omp target simd map(to: i) map(from: j)
161 // CHECK-NEXT: for (T i = 0; i < 2; ++i) {
162 // CHECK-NEXT: }
163
164 #pragma omp target simd map(always,alloc: i)
165 for (T i = 0; i < 2; ++i) {}
166 // CHECK: #pragma omp target simd map(always,alloc: i)
167 // CHECK-NEXT: for (T i = 0; i < 2; ++i) {
168 // CHECK-NEXT: }
169
170 #pragma omp target simd nowait
171 for (T i = 0; i < 2; ++i) {}
172 // CHECK: #pragma omp target simd nowait
173 // CHECK-NEXT: for (T i = 0; i < 2; ++i) {
174 // CHECK-NEXT: }
175
176 #pragma omp target simd depend(in : argc, arr[i:argc], arr1[:])
177 for (T i = 0; i < 2; ++i) {}
178 // CHECK: #pragma omp target simd depend(in : argc,arr[i:argc],arr1[:])
179 // CHECK-NEXT: for (T i = 0; i < 2; ++i) {
180 // CHECK-NEXT: }
181
182 #pragma omp target simd defaultmap(tofrom: scalar)
183 for (T i = 0; i < 2; ++i) {}
184 // CHECK: #pragma omp target simd defaultmap(tofrom: scalar)
185 // CHECK-NEXT: for (T i = 0; i < 2; ++i) {
186 // CHECK-NEXT: }
187
188 #pragma omp target simd safelen(clen-1)
189 for (T i = 0; i < 2; ++i) {}
190 // CHECK: #pragma omp target simd safelen(clen - 1)
191 // CHECK-NEXT: for (T i = 0; i < 2; ++i) {
192 // CHECK-NEXT: }
193
194 #pragma omp target simd simdlen(clen-1)
195 for (T i = 0; i < 2; ++i) {}
196 // CHECK: #pragma omp target simd simdlen(clen - 1)
197 // CHECK-NEXT: for (T i = 0; i < 2; ++i) {
198 // CHECK-NEXT: }
199
200 #pragma omp target simd aligned(arr1:N-1)
201 for (T i = 0; i < N; ++i) {}
202 // CHECK: #pragma omp target simd aligned(arr1: N - 1)
203 // CHECK-NEXT: for (T i = 0; i < N; ++i) {
204 // CHECK-NEXT: }
205
206 #pragma omp target simd is_device_ptr(p)
207 for (T i = 0; i < N; ++i) {}
208 // CHECK: #pragma omp target simd is_device_ptr(p)
209 // CHECK-NEXT: for (T i = 0; i < N; ++i) {
210 // CHECK-NEXT: }
211
212 return T();
213 }
214
main(int argc,char ** argv)215 int main(int argc, char **argv) {
216 int b = argc, c, d, e, f, h;
217 int arr[5][10], arr1[5];
218 int i, j;
219 int s;
220 static int a;
221 // CHECK: static int a;
222 const int clen = 5;
223 // CHECK: int clen = 5;
224 static float g;
225 #pragma omp threadprivate(g)
226 int *p;
227 #pragma omp target simd linear(a)
228 // CHECK: #pragma omp target simd linear(a)
229 for (int i = 0; i < 2; ++i)
230 a = 2;
231 // CHECK-NEXT: for (int i = 0; i < 2; ++i)
232 // CHECK-NEXT: a = 2;
233
234 #pragma omp target simd private(argc, b), firstprivate(argv, c), lastprivate(d, f) collapse(2) if (target: argc) reduction(+ : h) linear(a:-5)
235 for (int i = 0; i < 10; ++i)
236 for (int j = 0; j < 10; ++j)
237 foo();
238 // CHECK: #pragma omp target simd private(argc,b) firstprivate(argv,c) lastprivate(d,f) collapse(2) if(target: argc) reduction(+: h) linear(a: -5)
239 // CHECK-NEXT: for (int i = 0; i < 10; ++i)
240 // CHECK-NEXT: for (int j = 0; j < 10; ++j)
241 // CHECK-NEXT: foo();
242
243 #pragma omp target simd private(argc,b) firstprivate(argv) if (argc > 0) reduction(+:c, arr1[argc]) reduction(max:e, arr[:5][0:10])
244 for (int i = 0; i < 2; ++i) {}
245 // CHECK: #pragma omp target simd private(argc,b) firstprivate(argv) if(argc > 0) reduction(+: c,arr1[argc]) reduction(max: e,arr[:5][0:10])
246 // CHECK-NEXT: for (int i = 0; i < 2; ++i) {
247 // CHECK-NEXT: }
248
249 #pragma omp target simd if (5) reduction(^:e, f, arr[0:5][:argc]) reduction(&& : h)
250 for (int i = 0; i < 2; ++i) {}
251 // CHECK: #pragma omp target simd if(5) reduction(^: e,f,arr[0:5][:argc]) reduction(&&: h)
252 // CHECK-NEXT: for (int i = 0; i < 2; ++i) {
253 // CHECK-NEXT: }
254
255 #ifdef OMP5
256 #pragma omp target simd if (target:argc > 0) if(simd:argc) nontemporal(argc, c, d) lastprivate(conditional: d, f) order(concurrent)
257 #else
258 #pragma omp target simd if (target:argc > 0)
259 #endif // OMP5
260 for (int i = 0; i < 2; ++i) {}
261 // OMP45: #pragma omp target simd if(target: argc > 0)
262 // OMP50: #pragma omp target simd if(target: argc > 0) if(simd: argc) nontemporal(argc,c,d) lastprivate(conditional: d,f) order(concurrent)
263 // CHECK-NEXT: for (int i = 0; i < 2; ++i) {
264 // CHECK-NEXT: }
265
266 #pragma omp target simd if (5)
267 for (int i = 0; i < 2; ++i) {}
268 // CHECK: #pragma omp target simd if(5)
269 // CHECK-NEXT: for (int i = 0; i < 2; ++i) {
270 // CHECK-NEXT: }
271
272 #pragma omp target simd map(i)
273 for (int i = 0; i < 2; ++i) {}
274 // CHECK: #pragma omp target simd map(tofrom: i)
275 // CHECK-NEXT: for (int i = 0; i < 2; ++i) {
276 // CHECK-NEXT: }
277
278 #pragma omp target simd map(arr1[0:10], i)
279 for (int i = 0; i < 2; ++i) {}
280 // CHECK: #pragma omp target simd map(tofrom: arr1[0:10],i)
281 // CHECK-NEXT: for (int i = 0; i < 2; ++i) {
282 // CHECK-NEXT: }
283
284 #pragma omp target simd map(to: i) map(from: j)
285 for (int i = 0; i < 2; ++i) {}
286 // CHECK: #pragma omp target simd map(to: i) map(from: j)
287 // CHECK-NEXT: for (int i = 0; i < 2; ++i) {
288 // CHECK-NEXT: }
289
290 #pragma omp target simd map(always,alloc: i)
291 for (int i = 0; i < 2; ++i) {}
292 // CHECK: #pragma omp target simd map(always,alloc: i)
293 // CHECK-NEXT: for (int i = 0; i < 2; ++i) {
294 // CHECK-NEXT: }
295
296 #pragma omp target simd nowait
297 for (int i = 0; i < 2; ++i) {}
298 // CHECK: #pragma omp target simd nowait
299 // CHECK-NEXT: for (int i = 0; i < 2; ++i) {
300 // CHECK-NEXT: }
301
302 #pragma omp target simd depend(in : argc, arr[i:argc], arr1[:])
303 for (int i = 0; i < 2; ++i) {}
304 // CHECK: #pragma omp target simd depend(in : argc,arr[i:argc],arr1[:])
305 // CHECK-NEXT: for (int i = 0; i < 2; ++i) {
306 // CHECK-NEXT: }
307
308 #pragma omp target simd defaultmap(tofrom: scalar)
309 for (int i = 0; i < 2; ++i) {}
310 // CHECK: #pragma omp target simd defaultmap(tofrom: scalar)
311 // CHECK-NEXT: for (int i = 0; i < 2; ++i) {
312 // CHECK-NEXT: }
313
314 #pragma omp target simd safelen(clen-1)
315 for (int i = 0; i < 2; ++i) {}
316 // CHECK: #pragma omp target simd safelen(clen - 1)
317 // CHECK-NEXT: for (int i = 0; i < 2; ++i) {
318 // CHECK-NEXT: }
319
320 #pragma omp target simd simdlen(clen-1)
321 for (int i = 0; i < 2; ++i) {}
322 // CHECK: #pragma omp target simd simdlen(clen - 1)
323 // CHECK-NEXT: for (int i = 0; i < 2; ++i) {
324 // CHECK-NEXT: }
325
326 #pragma omp target simd aligned(arr1:4)
327 for (int i = 0; i < 2; ++i) {}
328 // CHECK: #pragma omp target simd aligned(arr1: 4)
329 // CHECK-NEXT: for (int i = 0; i < 2; ++i) {
330 // CHECK-NEXT: }
331
332 #pragma omp target simd is_device_ptr(p)
333 for (int i = 0; i < 2; ++i) {}
334 // CHECK: #pragma omp target simd is_device_ptr(p)
335 // CHECK-NEXT: for (int i = 0; i < 2; ++i) {
336 // CHECK-NEXT: }
337
338 return (tmain<int, 5>(argc, &argc));
339 }
340
341 #endif
342