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 
30 void foo() {}
31 
32 struct S {
33   S(): a(0) {}
34   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;
43   S7() : a(0) {}
44 
45 public:
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   }
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> {
64   S8() {}
65 
66 public:
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   }
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 
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