3 // expected-no-diagnostics
7 // RUN: %clang_cc1 -DOMP45 -verify -Wno-vla -fopenmp -fopenmp-version=45 -ast-print %s | FileCheck %s --check-prefix=OMP45
8 // RUN: %clang_cc1 -DOMP45 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -emit-pch -o %t %s
9 // RUN: %clang_cc1 -DOMP45 -fopenmp -fopenmp-version=45 -std=c++11 -include-pch %t -fsyntax-only -verify -Wno-vla %s -ast-print | FileCheck %s --check-prefix=OMP45
11 // RUN: %clang_cc1 -DOMP45 -verify -Wno-vla -fopenmp-simd -fopenmp-version=45 -ast-print %s | FileCheck %s --check-prefix=OMP45
12 // RUN: %clang_cc1 -DOMP45 -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -emit-pch -o %t %s
13 // RUN: %clang_cc1 -DOMP45 -fopenmp-simd -fopenmp-version=45 -std=c++11 -include-pch %t -fsyntax-only -verify -Wno-vla %s -ast-print | FileCheck %s --check-prefix=OMP45
18 template <typename T
, int C
>
19 T
tmain(T argc
, T
*argv
) {
20 T i
, j
, a
[20], always
, close
;
23 #pragma omp target if (target:argc > 0)
25 #pragma omp target if (C)
27 #pragma omp target map(i)
29 #pragma omp target map(a[0:10], i)
31 #pragma omp target map(to: i) map(from: j)
33 #pragma omp target map(always,alloc: i)
35 #pragma omp target map(always from: i)
37 #pragma omp target map(always)
39 #pragma omp target map(always,i)
41 #pragma omp target map(close,alloc: i)
43 #pragma omp target map(close, from: i)
45 #pragma omp target map(close)
47 #pragma omp target map(close,i)
49 #pragma omp target nowait
51 #pragma omp target depend(in : argc, argv[i:argc], a[:])
53 #pragma omp target defaultmap(tofrom: scalar)
58 // OMP45: template <typename T, int C> T tmain(T argc, T *argv) {
59 // OMP45-NEXT: T i, j, a[20]
60 // OMP45-NEXT: #pragma omp target{{$}}
62 // OMP45-NEXT: #pragma omp target if(target: argc > 0)
64 // OMP45-NEXT: #pragma omp target if(C)
66 // OMP45-NEXT: #pragma omp target map(tofrom: i)
68 // OMP45-NEXT: #pragma omp target map(tofrom: a[0:10],i)
70 // OMP45-NEXT: #pragma omp target map(to: i) map(from: j)
72 // OMP45-NEXT: #pragma omp target map(always,alloc: i)
74 // OMP45-NEXT: #pragma omp target map(always,from: i)
76 // OMP45-NEXT: #pragma omp target map(tofrom: always)
78 // OMP45-NEXT: always++;
80 // OMP45-NEXT: #pragma omp target map(tofrom: always,i)
82 // OMP45-NEXT: always++;
85 // OMP45-NEXT: #pragma omp target map(close,alloc: i)
87 // OMP45-NEXT: #pragma omp target map(close,from: i)
89 // OMP45-NEXT: #pragma omp target map(tofrom: close)
91 // OMP45-NEXT: close++;
93 // OMP45-NEXT: #pragma omp target map(tofrom: close,i)
95 // OMP45-NEXT: close++;
98 // OMP45-NEXT: #pragma omp target nowait
100 // OMP45-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])
102 // OMP45-NEXT: #pragma omp target defaultmap(tofrom: scalar)
104 // OMP45: template<> int tmain<int, 5>(int argc, int *argv) {
105 // OMP45-NEXT: int i, j, a[20]
106 // OMP45-NEXT: #pragma omp target
107 // OMP45-NEXT: foo();
108 // OMP45-NEXT: #pragma omp target if(target: argc > 0)
110 // OMP45-NEXT: #pragma omp target if(5)
112 // OMP45-NEXT: #pragma omp target map(tofrom: i)
114 // OMP45-NEXT: #pragma omp target map(tofrom: a[0:10],i)
116 // OMP45-NEXT: #pragma omp target map(to: i) map(from: j)
118 // OMP45-NEXT: #pragma omp target map(always,alloc: i)
120 // OMP45-NEXT: #pragma omp target map(always,from: i)
122 // OMP45-NEXT: #pragma omp target map(tofrom: always)
124 // OMP45-NEXT: always++;
126 // OMP45-NEXT: #pragma omp target map(tofrom: always,i)
128 // OMP45-NEXT: always++;
131 // OMP45-NEXT: #pragma omp target map(close,alloc: i)
133 // OMP45-NEXT: #pragma omp target map(close,from: i)
135 // OMP45-NEXT: #pragma omp target map(tofrom: close)
137 // OMP45-NEXT: close++;
139 // OMP45-NEXT: #pragma omp target map(tofrom: close,i)
141 // OMP45-NEXT: close++;
144 // OMP45-NEXT: #pragma omp target nowait
146 // OMP45-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])
148 // OMP45-NEXT: #pragma omp target defaultmap(tofrom: scalar)
150 // OMP45: template<> char tmain<char, 1>(char argc, char *argv) {
151 // OMP45-NEXT: char i, j, a[20]
152 // OMP45-NEXT: #pragma omp target
153 // OMP45-NEXT: foo();
154 // OMP45-NEXT: #pragma omp target if(target: argc > 0)
156 // OMP45-NEXT: #pragma omp target if(1)
158 // OMP45-NEXT: #pragma omp target map(tofrom: i)
160 // OMP45-NEXT: #pragma omp target map(tofrom: a[0:10],i)
162 // OMP45-NEXT: #pragma omp target map(to: i) map(from: j)
164 // OMP45-NEXT: #pragma omp target map(always,alloc: i)
166 // OMP45-NEXT: #pragma omp target map(always,from: i)
168 // OMP45-NEXT: #pragma omp target map(tofrom: always)
170 // OMP45-NEXT: always++;
172 // OMP45-NEXT: #pragma omp target map(tofrom: always,i)
174 // OMP45-NEXT: always++;
177 // OMP45-NEXT: #pragma omp target map(close,alloc: i)
179 // OMP45-NEXT: #pragma omp target map(close,from: i)
181 // OMP45-NEXT: #pragma omp target map(tofrom: close)
183 // OMP45-NEXT: close++;
185 // OMP45-NEXT: #pragma omp target map(tofrom: close,i)
187 // OMP45-NEXT: close++;
190 // OMP45-NEXT: #pragma omp target nowait
192 // OMP45-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])
194 // OMP45-NEXT: #pragma omp target defaultmap(tofrom: scalar)
197 // OMP45-LABEL: class S {
200 // OMP45-NEXT: void foo() {
202 // OMP45-NEXT: int a = 0;
203 #pragma omp target map(this[0])
204 // OMP45-NEXT: #pragma omp target map(tofrom: this[0])
207 #pragma omp target map(this[:1])
208 // OMP45-NEXT: #pragma omp target map(tofrom: this[:1])
211 #pragma omp target map((this)[0])
212 // OMP45-NEXT: #pragma omp target map(tofrom: (this)[0])
215 #pragma omp target map(this[:a])
216 // OMP45-NEXT: #pragma omp target map(tofrom: this[:a])
219 #pragma omp target map(this[a:1])
220 // OMP45-NEXT: #pragma omp target map(tofrom: this[a:1])
223 #pragma omp target map(this[a])
224 // OMP45-NEXT: #pragma omp target map(tofrom: this[a])
232 // OMP45-LABEL: int main(int argc, char **argv) {
233 int main (int argc
, char **argv
) {
234 int i
, j
, a
[20], always
, close
;
235 // OMP45-NEXT: int i, j, a[20]
237 // OMP45-NEXT: #pragma omp target
239 // OMP45-NEXT: foo();
240 #pragma omp target if (argc > 0)
241 // OMP45-NEXT: #pragma omp target if(argc > 0)
243 // OMP45-NEXT: foo();
245 #pragma omp target map(i) if(argc>0)
246 // OMP45-NEXT: #pragma omp target map(tofrom: i) if(argc > 0)
248 // OMP45-NEXT: foo();
250 #pragma omp target map(i)
251 // OMP45-NEXT: #pragma omp target map(tofrom: i)
253 // OMP45-NEXT: foo();
255 #pragma omp target map(a[0:10], i)
256 // OMP45-NEXT: #pragma omp target map(tofrom: a[0:10],i)
258 // OMP45-NEXT: foo();
260 #pragma omp target map(to: i) map(from: j)
261 // OMP45-NEXT: #pragma omp target map(to: i) map(from: j)
263 // OMP45-NEXT: foo();
265 #pragma omp target map(always,alloc: i)
266 // OMP45-NEXT: #pragma omp target map(always,alloc: i)
268 // OMP45-NEXT: foo();
270 #pragma omp target map(always from: i)
271 // OMP45-NEXT: #pragma omp target map(always,from: i)
273 // OMP45-NEXT: foo();
275 #pragma omp target map(always)
276 // OMP45-NEXT: #pragma omp target map(tofrom: always)
279 // OMP45-NEXT: always++;
282 #pragma omp target map(always,i)
283 // OMP45-NEXT: #pragma omp target map(tofrom: always,i)
286 // OMP45-NEXT: always++;
290 #pragma omp target map(close,alloc: i)
291 // OMP45-NEXT: #pragma omp target map(close,alloc: i)
293 // OMP45-NEXT: foo();
295 #pragma omp target map(close from: i)
296 // OMP45-NEXT: #pragma omp target map(close,from: i)
298 // OMP45-NEXT: foo();
300 #pragma omp target map(close)
301 // OMP45-NEXT: #pragma omp target map(tofrom: close)
304 // OMP45-NEXT: close++;
307 #pragma omp target map(close,i)
308 // OMP45-NEXT: #pragma omp target map(tofrom: close,i)
311 // OMP45-NEXT: close++;
315 #pragma omp target nowait
316 // OMP45-NEXT: #pragma omp target nowait
318 // OMP45-NEXT: foo();
320 #pragma omp target depend(in : argc, argv[i:argc], a[:])
321 // OMP45-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])
323 // OMP45-NEXT: foo();
325 #pragma omp target defaultmap(tofrom: scalar)
326 // OMP45-NEXT: #pragma omp target defaultmap(tofrom: scalar)
328 // OMP45-NEXT: foo();
330 return tmain
<int, 5>(argc
, &argc
) + tmain
<char, 1>(argv
[0][0], argv
[0]);
337 ///==========================================================================///
338 // RUN: %clang_cc1 -DOMP5 -verify -Wno-vla -fopenmp -fopenmp-version=50 -ast-print %s | FileCheck %s --check-prefix OMP5
339 // RUN: %clang_cc1 -DOMP5 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
340 // RUN: %clang_cc1 -DOMP5 -fopenmp -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify -Wno-vla %s -ast-print | FileCheck %s --check-prefix OMP5
342 // RUN: %clang_cc1 -DOMP5 -verify -Wno-vla -fopenmp-simd -fopenmp-version=50 -ast-print %s | FileCheck %s --check-prefix OMP5
343 // RUN: %clang_cc1 -DOMP5 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
344 // RUN: %clang_cc1 -DOMP5 -fopenmp-simd -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify -Wno-vla %s -ast-print | FileCheck %s --check-prefix OMP5
346 // RUN: %clang_cc1 -DOMP5 -verify -Wno-vla -fopenmp -fopenmp-version=99 -DOMP99 -ast-print %s | FileCheck %s --check-prefixes=OMP5,REV
347 // RUN: %clang_cc1 -DOMP5 -fopenmp -fopenmp-version=99 -DOMP99 -x c++ -std=c++11 -emit-pch -o %t %s
348 // RUN: %clang_cc1 -DOMP5 -fopenmp -fopenmp-version=99 -DOMP99 -std=c++11 -include-pch %t -fsyntax-only -verify -Wno-vla %s -ast-print | FileCheck %s --check-prefixes=OMP5,REV
350 // RUN: %clang_cc1 -DOMP5 -verify -Wno-vla -fopenmp-simd -fopenmp-version=99 -DOMP99 -ast-print %s | FileCheck %s --check-prefixes=OMP5,REV
351 // RUN: %clang_cc1 -DOMP5 -fopenmp-simd -fopenmp-version=99 -DOMP99 -x c++ -std=c++11 -emit-pch -o %t %s
352 // RUN: %clang_cc1 -DOMP5 -fopenmp-simd -fopenmp-version=99 -DOMP99 -std=c++11 -include-pch %t -fsyntax-only -verify -Wno-vla %s -ast-print | FileCheck %s --check-prefixes=OMP5,REV
355 #pragma omp requires reverse_offload
357 typedef void **omp_allocator_handle_t
;
358 extern const omp_allocator_handle_t omp_null_allocator
;
359 extern const omp_allocator_handle_t omp_default_mem_alloc
;
360 extern const omp_allocator_handle_t omp_large_cap_mem_alloc
;
361 extern const omp_allocator_handle_t omp_const_mem_alloc
;
362 extern const omp_allocator_handle_t omp_high_bw_mem_alloc
;
363 extern const omp_allocator_handle_t omp_low_lat_mem_alloc
;
364 extern const omp_allocator_handle_t omp_cgroup_mem_alloc
;
365 extern const omp_allocator_handle_t omp_pteam_mem_alloc
;
366 extern const omp_allocator_handle_t omp_thread_mem_alloc
;
370 #pragma omp declare target
372 #pragma omp end declare target
375 #pragma omp declare target link(a)
377 template <typename T
, int C
>
378 T
tmain(T argc
, T
*argv
) {
379 T i
, j
, a
[20], always
, close
;
380 #pragma omp target device(argc)
382 #pragma omp target if (target:argc > 0) device(device_num: C)
385 #pragma omp target if (C) device(ancestor: argc)
388 #pragma omp target map(i)
390 #pragma omp target map(a[0:10], i)
392 #pragma omp target map(to: i) map(from: j)
394 #pragma omp target map(always,alloc: i)
396 #pragma omp target map(always, from: i)
398 #pragma omp target map(always)
400 #pragma omp target map(always,i)
402 #pragma omp target map(close,alloc: i)
404 #pragma omp target map(close, from: i)
406 #pragma omp target map(close)
408 #pragma omp target map(close,i)
410 #pragma omp target nowait
412 #pragma omp target depend(in : argc, argv[i:argc], a[:])
414 #pragma omp target defaultmap(alloc: scalar)
416 #pragma omp target defaultmap(to: scalar)
418 #pragma omp target defaultmap(from: scalar)
420 #pragma omp target defaultmap(tofrom: scalar)
422 #pragma omp target defaultmap(firstprivate: scalar)
424 #pragma omp target defaultmap(none: scalar)
426 #pragma omp target defaultmap(default: scalar)
428 #pragma omp target defaultmap(alloc: aggregate)
430 #pragma omp target defaultmap(to: aggregate)
432 #pragma omp target defaultmap(from: aggregate)
434 #pragma omp target defaultmap(tofrom: aggregate)
436 #pragma omp target defaultmap(firstprivate: aggregate)
438 #pragma omp target defaultmap(none: aggregate)
440 #pragma omp target defaultmap(default: aggregate)
442 #pragma omp target defaultmap(alloc: pointer)
444 #pragma omp target defaultmap(to: pointer)
446 #pragma omp target defaultmap(from: pointer)
448 #pragma omp target defaultmap(tofrom: pointer)
450 #pragma omp target defaultmap(firstprivate: pointer)
452 #pragma omp target defaultmap(none: pointer)
454 #pragma omp target defaultmap(default: pointer)
456 #pragma omp target defaultmap(to: scalar) defaultmap(tofrom: pointer)
458 #pragma omp target defaultmap(from: pointer) defaultmap(none: aggregate)
460 #pragma omp target defaultmap(default: aggregate) defaultmap(alloc: scalar)
462 #pragma omp target defaultmap(alloc: aggregate) defaultmap(firstprivate: scalar) defaultmap(tofrom: pointer)
464 #pragma omp target defaultmap(tofrom: aggregate) defaultmap(to: pointer) defaultmap(alloc: scalar)
469 #pragma omp target is_device_ptr(g) defaultmap(none: pointer)
471 #pragma omp target private(g) defaultmap(none: pointer)
473 #pragma omp target firstprivate(g) defaultmap(none: pointer)
475 #pragma omp target defaultmap(none: scalar) map(to: i)
477 #pragma omp target defaultmap(none: aggregate) map(to: a)
479 #pragma omp target defaultmap(none: scalar)
485 // OMP5: template <typename T, int C> T tmain(T argc, T *argv) {
486 // OMP5-NEXT: T i, j, a[20]
487 // OMP5-NEXT: #pragma omp target device(argc){{$}}
489 // OMP5-NEXT: #pragma omp target if(target: argc > 0) device(device_num: C)
491 // REV: #pragma omp target if(C) device(ancestor: argc)
493 // OMP5-NEXT: #pragma omp target map(tofrom: i)
495 // OMP5-NEXT: #pragma omp target map(tofrom: a[0:10],i)
497 // OMP5-NEXT: #pragma omp target map(to: i) map(from: j)
499 // OMP5-NEXT: #pragma omp target map(always,alloc: i)
501 // OMP5-NEXT: #pragma omp target map(always,from: i)
503 // OMP5-NEXT: #pragma omp target map(tofrom: always)
505 // OMP5-NEXT: always++;
507 // OMP5-NEXT: #pragma omp target map(tofrom: always,i)
509 // OMP5-NEXT: always++;
512 // OMP5-NEXT: #pragma omp target map(close,alloc: i)
514 // OMP5-NEXT: #pragma omp target map(close,from: i)
516 // OMP5-NEXT: #pragma omp target map(tofrom: close)
518 // OMP5-NEXT: close++;
520 // OMP5-NEXT: #pragma omp target map(tofrom: close,i)
522 // OMP5-NEXT: close++;
525 // OMP5-NEXT: #pragma omp target nowait
527 // OMP5-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])
529 // OMP5-NEXT: #pragma omp target defaultmap(alloc: scalar)
531 // OMP5-NEXT: #pragma omp target defaultmap(to: scalar)
533 // OMP5-NEXT: #pragma omp target defaultmap(from: scalar)
535 // OMP5-NEXT: #pragma omp target defaultmap(tofrom: scalar)
537 // OMP5-NEXT: #pragma omp target defaultmap(firstprivate: scalar)
539 // OMP5-NEXT: #pragma omp target defaultmap(none: scalar)
541 // OMP5-NEXT: #pragma omp target defaultmap(default: scalar)
543 // OMP5-NEXT: #pragma omp target defaultmap(alloc: aggregate)
545 // OMP5-NEXT: #pragma omp target defaultmap(to: aggregate)
547 // OMP5-NEXT: #pragma omp target defaultmap(from: aggregate)
549 // OMP5-NEXT: #pragma omp target defaultmap(tofrom: aggregate)
551 // OMP5-NEXT: #pragma omp target defaultmap(firstprivate: aggregate)
553 // OMP5-NEXT: #pragma omp target defaultmap(none: aggregate)
555 // OMP5-NEXT: #pragma omp target defaultmap(default: aggregate)
557 // OMP5-NEXT: #pragma omp target defaultmap(alloc: pointer)
559 // OMP5-NEXT: #pragma omp target defaultmap(to: pointer)
561 // OMP5-NEXT: #pragma omp target defaultmap(from: pointer)
563 // OMP5-NEXT: #pragma omp target defaultmap(tofrom: pointer)
565 // OMP5-NEXT: #pragma omp target defaultmap(firstprivate: pointer)
567 // OMP5-NEXT: #pragma omp target defaultmap(none: pointer)
569 // OMP5-NEXT: #pragma omp target defaultmap(default: pointer)
571 // OMP5-NEXT: #pragma omp target defaultmap(to: scalar) defaultmap(tofrom: pointer)
573 // OMP5-NEXT: #pragma omp target defaultmap(from: pointer) defaultmap(none: aggregate)
575 // OMP5-NEXT: #pragma omp target defaultmap(default: aggregate) defaultmap(alloc: scalar)
577 // OMP5-NEXT: #pragma omp target defaultmap(alloc: aggregate) defaultmap(firstprivate: scalar) defaultmap(tofrom: pointer)
579 // OMP5-NEXT: #pragma omp target defaultmap(tofrom: aggregate) defaultmap(to: pointer) defaultmap(alloc: scalar)
581 // OMP5: template<> int tmain<int, 5>(int argc, int *argv) {
582 // OMP5-NEXT: int i, j, a[20]
583 // OMP5-NEXT: #pragma omp target
585 // OMP5-NEXT: #pragma omp target if(target: argc > 0)
587 // REV: #pragma omp target if(5)
589 // OMP5-NEXT: #pragma omp target map(tofrom: i)
591 // OMP5-NEXT: #pragma omp target map(tofrom: a[0:10],i)
593 // OMP5-NEXT: #pragma omp target map(to: i) map(from: j)
595 // OMP5-NEXT: #pragma omp target map(always,alloc: i)
597 // OMP5-NEXT: #pragma omp target map(always,from: i)
599 // OMP5-NEXT: #pragma omp target map(tofrom: always)
601 // OMP5-NEXT: always++;
603 // OMP5-NEXT: #pragma omp target map(tofrom: always,i)
605 // OMP5-NEXT: always++;
608 // OMP5-NEXT: #pragma omp target map(close,alloc: i)
610 // OMP5-NEXT: #pragma omp target map(close,from: i)
612 // OMP5-NEXT: #pragma omp target map(tofrom: close)
614 // OMP5-NEXT: close++;
616 // OMP5-NEXT: #pragma omp target map(tofrom: close,i)
618 // OMP5-NEXT: close++;
621 // OMP5-NEXT: #pragma omp target nowait
623 // OMP5-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])
625 // OMP5-NEXT: #pragma omp target defaultmap(alloc: scalar)
627 // OMP5-NEXT: #pragma omp target defaultmap(to: scalar)
629 // OMP5-NEXT: #pragma omp target defaultmap(from: scalar)
631 // OMP5-NEXT: #pragma omp target defaultmap(tofrom: scalar)
633 // OMP5-NEXT: #pragma omp target defaultmap(firstprivate: scalar)
635 // OMP5-NEXT: #pragma omp target defaultmap(none: scalar)
637 // OMP5-NEXT: #pragma omp target defaultmap(default: scalar)
639 // OMP5-NEXT: #pragma omp target defaultmap(alloc: aggregate)
641 // OMP5-NEXT: #pragma omp target defaultmap(to: aggregate)
643 // OMP5-NEXT: #pragma omp target defaultmap(from: aggregate)
645 // OMP5-NEXT: #pragma omp target defaultmap(tofrom: aggregate)
647 // OMP5-NEXT: #pragma omp target defaultmap(firstprivate: aggregate)
649 // OMP5-NEXT: #pragma omp target defaultmap(none: aggregate)
651 // OMP5-NEXT: #pragma omp target defaultmap(default: aggregate)
653 // OMP5-NEXT: #pragma omp target defaultmap(alloc: pointer)
655 // OMP5-NEXT: #pragma omp target defaultmap(to: pointer)
657 // OMP5-NEXT: #pragma omp target defaultmap(from: pointer)
659 // OMP5-NEXT: #pragma omp target defaultmap(tofrom: pointer)
661 // OMP5-NEXT: #pragma omp target defaultmap(firstprivate: pointer)
663 // OMP5-NEXT: #pragma omp target defaultmap(none: pointer)
665 // OMP5-NEXT: #pragma omp target defaultmap(default: pointer)
667 // OMP5-NEXT: #pragma omp target defaultmap(to: scalar) defaultmap(tofrom: pointer)
669 // OMP5-NEXT: #pragma omp target defaultmap(from: pointer) defaultmap(none: aggregate)
671 // OMP5-NEXT: #pragma omp target defaultmap(default: aggregate) defaultmap(alloc: scalar)
673 // OMP5-NEXT: #pragma omp target defaultmap(alloc: aggregate) defaultmap(firstprivate: scalar) defaultmap(tofrom: pointer)
675 // OMP5-NEXT: #pragma omp target defaultmap(tofrom: aggregate) defaultmap(to: pointer) defaultmap(alloc: scalar)
677 // OMP5: template<> char tmain<char, 1>(char argc, char *argv) {
678 // OMP5-NEXT: char i, j, a[20]
679 // OMP5-NEXT: #pragma omp target device(argc)
681 // OMP5-NEXT: #pragma omp target if(target: argc > 0) device(device_num: 1)
683 // REV: #pragma omp target if(1) device(ancestor: argc)
685 // OMP5-NEXT: #pragma omp target map(tofrom: i)
687 // OMP5-NEXT: #pragma omp target map(tofrom: a[0:10],i)
689 // OMP5-NEXT: #pragma omp target map(to: i) map(from: j)
691 // OMP5-NEXT: #pragma omp target map(always,alloc: i)
693 // OMP5-NEXT: #pragma omp target map(always,from: i)
695 // OMP5-NEXT: #pragma omp target map(tofrom: always)
697 // OMP5-NEXT: always++;
699 // OMP5-NEXT: #pragma omp target map(tofrom: always,i)
701 // OMP5-NEXT: always++;
704 // OMP5-NEXT: #pragma omp target map(close,alloc: i)
706 // OMP5-NEXT: #pragma omp target map(close,from: i)
708 // OMP5-NEXT: #pragma omp target map(tofrom: close)
710 // OMP5-NEXT: close++;
712 // OMP5-NEXT: #pragma omp target map(tofrom: close,i)
714 // OMP5-NEXT: close++;
717 // OMP5-NEXT: #pragma omp target nowait
719 // OMP5-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])
721 // OMP5-NEXT: #pragma omp target defaultmap(alloc: scalar)
723 // OMP5-NEXT: #pragma omp target defaultmap(to: scalar)
725 // OMP5-NEXT: #pragma omp target defaultmap(from: scalar)
727 // OMP5-NEXT: #pragma omp target defaultmap(tofrom: scalar)
729 // OMP5-NEXT: #pragma omp target defaultmap(firstprivate: scalar)
731 // OMP5-NEXT: #pragma omp target defaultmap(none: scalar)
733 // OMP5-NEXT: #pragma omp target defaultmap(default: scalar)
735 // OMP5-NEXT: #pragma omp target defaultmap(alloc: aggregate)
737 // OMP5-NEXT: #pragma omp target defaultmap(to: aggregate)
739 // OMP5-NEXT: #pragma omp target defaultmap(from: aggregate)
741 // OMP5-NEXT: #pragma omp target defaultmap(tofrom: aggregate)
743 // OMP5-NEXT: #pragma omp target defaultmap(firstprivate: aggregate)
745 // OMP5-NEXT: #pragma omp target defaultmap(none: aggregate)
747 // OMP5-NEXT: #pragma omp target defaultmap(default: aggregate)
749 // OMP5-NEXT: #pragma omp target defaultmap(alloc: pointer)
751 // OMP5-NEXT: #pragma omp target defaultmap(to: pointer)
753 // OMP5-NEXT: #pragma omp target defaultmap(from: pointer)
755 // OMP5-NEXT: #pragma omp target defaultmap(tofrom: pointer)
757 // OMP5-NEXT: #pragma omp target defaultmap(firstprivate: pointer)
759 // OMP5-NEXT: #pragma omp target defaultmap(none: pointer)
761 // OMP5-NEXT: #pragma omp target defaultmap(default: pointer)
763 // OMP5-NEXT: #pragma omp target defaultmap(to: scalar) defaultmap(tofrom: pointer)
765 // OMP5-NEXT: #pragma omp target defaultmap(from: pointer) defaultmap(none: aggregate)
767 // OMP5-NEXT: #pragma omp target defaultmap(default: aggregate) defaultmap(alloc: scalar)
769 // OMP5-NEXT: #pragma omp target defaultmap(alloc: aggregate) defaultmap(firstprivate: scalar) defaultmap(tofrom: pointer)
771 // OMP5-NEXT: #pragma omp target defaultmap(tofrom: aggregate) defaultmap(to: pointer) defaultmap(alloc: scalar)
773 // OMP5-NEXT: int *g;
774 // OMP5-NEXT: #pragma omp target is_device_ptr(g) defaultmap(none: pointer)
776 // OMP5-NEXT: #pragma omp target private(g) defaultmap(none: pointer)
778 // OMP5-NEXT: #pragma omp target firstprivate(g) defaultmap(none: pointer)
780 // OMP5-NEXT: #pragma omp target defaultmap(none: scalar) map(to: i)
782 // OMP5-NEXT: #pragma omp target defaultmap(none: aggregate) map(to: a)
783 // OMP5-NEXT: a[3]++;
784 // OMP5-NEXT: #pragma omp target defaultmap(none: scalar)
787 // OMP5-LABEL: class S {
790 // OMP5-NEXT: void foo() {
792 // OMP5-NEXT: int a = 0;
793 #pragma omp target map(this[0])
794 // OMP5-NEXT: #pragma omp target map(tofrom: this[0])
797 #pragma omp target map(this[:1])
798 // OMP5-NEXT: #pragma omp target map(tofrom: this[:1])
801 #pragma omp target map((this)[0])
802 // OMP5-NEXT: #pragma omp target map(tofrom: (this)[0])
805 #pragma omp target map(this[:a])
806 // OMP5-NEXT: #pragma omp target map(tofrom: this[:a])
809 #pragma omp target map(this[a:1])
810 // OMP5-NEXT: #pragma omp target map(tofrom: this[a:1])
813 #pragma omp target map(this[a])
814 // OMP5-NEXT: #pragma omp target map(tofrom: this[a])
822 // OMP5-LABEL: int main(int argc, char **argv) {
823 int main (int argc
, char **argv
) {
824 int i
, j
, a
[20], always
, close
;
825 // OMP5-NEXT: int i, j, a[20]
827 // OMP5-NEXT: #pragma omp target
830 #pragma omp target if (argc > 0)
831 // OMP5-NEXT: #pragma omp target if(argc > 0)
835 #pragma omp target map(i) if(argc>0)
836 // OMP5-NEXT: #pragma omp target map(tofrom: i) if(argc > 0)
840 #pragma omp target map(i)
841 // OMP5-NEXT: #pragma omp target map(tofrom: i)
845 #pragma omp target map(a[0:10], i)
846 // OMP5-NEXT: #pragma omp target map(tofrom: a[0:10],i)
850 #pragma omp target map(to: i) map(from: j)
851 // OMP5-NEXT: #pragma omp target map(to: i) map(from: j)
855 #pragma omp target map(always,alloc: i)
856 // OMP5-NEXT: #pragma omp target map(always,alloc: i)
860 #pragma omp target map(always, from: i)
861 // OMP5-NEXT: #pragma omp target map(always,from: i)
865 #pragma omp target map(always)
866 // OMP5-NEXT: #pragma omp target map(tofrom: always)
869 // OMP5-NEXT: always++;
872 #pragma omp target map(always,i)
873 // OMP5-NEXT: #pragma omp target map(tofrom: always,i)
876 // OMP5-NEXT: always++;
880 #pragma omp target map(close,alloc: i)
881 // OMP5-NEXT: #pragma omp target map(close,alloc: i)
885 #pragma omp target map(close, from: i)
886 // OMP5-NEXT: #pragma omp target map(close,from: i)
890 #pragma omp target map(close)
891 // OMP5-NEXT: #pragma omp target map(tofrom: close)
894 // OMP5-NEXT: close++;
897 #pragma omp target map(close,i)
898 // OMP5-NEXT: #pragma omp target map(tofrom: close,i)
901 // OMP5-NEXT: close++;
905 #pragma omp target nowait
906 // OMP5-NEXT: #pragma omp target nowait
910 #pragma omp target depend(in : argc, argv[i:argc], a[:])
911 // OMP5-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])
915 #pragma omp target defaultmap(alloc: scalar)
916 // OMP5-NEXT: #pragma omp target defaultmap(alloc: scalar)
920 #pragma omp target defaultmap(to: scalar)
921 // OMP5-NEXT: #pragma omp target defaultmap(to: scalar)
925 #pragma omp target defaultmap(from: scalar)
926 // OMP5-NEXT: #pragma omp target defaultmap(from: scalar)
930 #pragma omp target defaultmap(tofrom: scalar)
931 // OMP5-NEXT: #pragma omp target defaultmap(tofrom: scalar)
935 #pragma omp target defaultmap(firstprivate: scalar)
936 // OMP5-NEXT: #pragma omp target defaultmap(firstprivate: scalar)
940 #pragma omp target defaultmap(none: scalar)
941 // OMP5-NEXT: #pragma omp target defaultmap(none: scalar)
945 #pragma omp target defaultmap(default: scalar)
946 // OMP5-NEXT: #pragma omp target defaultmap(default: scalar)
950 #pragma omp target defaultmap(alloc: aggregate)
951 // OMP5-NEXT: #pragma omp target defaultmap(alloc: aggregate)
955 #pragma omp target defaultmap(to: aggregate)
956 // OMP5-NEXT: #pragma omp target defaultmap(to: aggregate)
960 #pragma omp target defaultmap(from: aggregate)
961 // OMP5-NEXT: #pragma omp target defaultmap(from: aggregate)
965 #pragma omp target defaultmap(tofrom: aggregate)
966 // OMP5-NEXT: #pragma omp target defaultmap(tofrom: aggregate)
970 #pragma omp target defaultmap(firstprivate: aggregate)
971 // OMP5-NEXT: #pragma omp target defaultmap(firstprivate: aggregate)
975 #pragma omp target defaultmap(none: aggregate)
976 // OMP5-NEXT: #pragma omp target defaultmap(none: aggregate)
980 #pragma omp target defaultmap(default: aggregate)
981 // OMP5-NEXT: #pragma omp target defaultmap(default: aggregate)
985 #pragma omp target defaultmap(alloc: pointer)
986 // OMP5-NEXT: #pragma omp target defaultmap(alloc: pointer)
990 #pragma omp target defaultmap(to: pointer)
991 // OMP5-NEXT: #pragma omp target defaultmap(to: pointer)
995 #pragma omp target defaultmap(from: pointer)
996 // OMP5-NEXT: #pragma omp target defaultmap(from: pointer)
1000 #pragma omp target defaultmap(tofrom: pointer)
1001 // OMP5-NEXT: #pragma omp target defaultmap(tofrom: pointer)
1003 // OMP5-NEXT: foo();
1005 #pragma omp target defaultmap(firstprivate: pointer)
1006 // OMP5-NEXT: #pragma omp target defaultmap(firstprivate: pointer)
1008 // OMP5-NEXT: foo();
1010 #pragma omp target defaultmap(none: pointer)
1011 // OMP5-NEXT: #pragma omp target defaultmap(none: pointer)
1013 // OMP5-NEXT: foo();
1015 #pragma omp target defaultmap(default: pointer)
1016 // OMP5-NEXT: #pragma omp target defaultmap(default: pointer)
1018 // OMP5-NEXT: foo();
1020 #pragma omp target defaultmap(to: scalar) defaultmap(tofrom: pointer)
1021 // OMP5-NEXT: #pragma omp target defaultmap(to: scalar) defaultmap(tofrom: pointer)
1023 // OMP5-NEXT: foo();
1025 #pragma omp target defaultmap(from: pointer) defaultmap(none: aggregate)
1026 // OMP5-NEXT: #pragma omp target defaultmap(from: pointer) defaultmap(none: aggregate)
1028 // OMP5-NEXT: foo();
1030 #pragma omp target defaultmap(default: aggregate) defaultmap(alloc: scalar)
1031 // OMP5-NEXT: #pragma omp target defaultmap(default: aggregate) defaultmap(alloc: scalar)
1033 // OMP5-NEXT: foo();
1035 #pragma omp target defaultmap(alloc: aggregate) defaultmap(firstprivate: scalar) defaultmap(tofrom: pointer)
1036 // OMP5-NEXT: #pragma omp target defaultmap(alloc: aggregate) defaultmap(firstprivate: scalar) defaultmap(tofrom: pointer)
1038 // OMP5-NEXT: foo();
1040 #pragma omp target defaultmap(tofrom: aggregate) defaultmap(to: pointer) defaultmap(alloc: scalar)
1041 // OMP5-NEXT: #pragma omp target defaultmap(tofrom: aggregate) defaultmap(to: pointer) defaultmap(alloc: scalar)
1043 // OMP5-NEXT: foo();
1046 // OMP5-NEXT: int *g;
1048 #pragma omp target is_device_ptr(g) defaultmap(none: pointer)
1049 // OMP5-NEXT: #pragma omp target is_device_ptr(g) defaultmap(none: pointer)
1053 #pragma omp target private(g) defaultmap(none: pointer)
1054 // OMP5-NEXT: #pragma omp target private(g) defaultmap(none: pointer)
1058 #pragma omp target firstprivate(g) defaultmap(none: pointer)
1059 // OMP5-NEXT: #pragma omp target firstprivate(g) defaultmap(none: pointer)
1063 #pragma omp target defaultmap(none: scalar) map(to: i)
1064 // OMP5-NEXT: #pragma omp target defaultmap(none: scalar) map(to: i)
1068 #pragma omp target defaultmap(none: aggregate) map(to: a)
1069 // OMP5-NEXT: #pragma omp target defaultmap(none: aggregate) map(to: a)
1071 // OMP5-NEXT: a[3]++;
1073 #pragma omp target defaultmap(none: scalar)
1074 // OMP5-NEXT: #pragma omp target defaultmap(none: scalar)
1076 // OMP5-NEXT: bar();
1077 #pragma omp target defaultmap(none)
1078 // OMP5-NEXT: #pragma omp target defaultmap(none)
1079 // OMP5-NEXT: bar();
1081 #pragma omp target allocate(omp_default_mem_alloc:argv) uses_allocators(omp_default_mem_alloc,omp_large_cap_mem_alloc) allocate(omp_large_cap_mem_alloc:argc) private(argc, argv)
1082 // OMP5-NEXT: #pragma omp target allocate(omp_default_mem_alloc: argv) uses_allocators(omp_default_mem_alloc,omp_large_cap_mem_alloc) allocate(omp_large_cap_mem_alloc: argc) private(argc,argv)
1083 // OMP5-NEXT: bar();
1085 return tmain
<int, 5>(argc
, &argc
) + tmain
<char, 1>(argv
[0][0], argv
[0]);
1092 ///==========================================================================///
1093 // RUN: %clang_cc1 -DOMP51 -verify -Wno-vla -fopenmp -ast-print %s | FileCheck %s --check-prefix OMP51
1094 // RUN: %clang_cc1 -DOMP51 -fopenmp -x c++ -std=c++11 -emit-pch -o %t %s
1095 // RUN: %clang_cc1 -DOMP51 -fopenmp -std=c++11 -include-pch %t -fsyntax-only -verify -Wno-vla %s -ast-print | FileCheck %s --check-prefix OMP51
1097 // RUN: %clang_cc1 -DOMP51 -verify -Wno-vla -fopenmp-simd -ast-print %s | FileCheck %s --check-prefix OMP51
1098 // RUN: %clang_cc1 -DOMP51 -fopenmp-simd -x c++ -std=c++11 -emit-pch -o %t %s
1099 // RUN: %clang_cc1 -DOMP51 -fopenmp-simd -std=c++11 -include-pch %t -fsyntax-only -verify -Wno-vla %s -ast-print | FileCheck %s --check-prefix OMP51
1103 template <typename T
, int C
>
1104 T
tmain(T argc
, T
*argv
) {
1105 #pragma omp target defaultmap(present: scalar)
1107 #pragma omp target defaultmap(present: aggregate)
1109 #pragma omp target defaultmap(present: pointer)
1111 #pragma omp target thread_limit(C)
1117 // OMP51: template <typename T, int C> T tmain(T argc, T *argv) {
1118 // OMP51-NEXT: #pragma omp target defaultmap(present: scalar)
1119 // OMP51-NEXT: foo()
1120 // OMP51-NEXT: #pragma omp target defaultmap(present: aggregate)
1121 // OMP51-NEXT: foo()
1122 // OMP51-NEXT: #pragma omp target defaultmap(present: pointer)
1123 // OMP51-NEXT: foo()
1124 // OMP51-NEXT: #pragma omp target thread_limit(C)
1125 // OMP51-NEXT: foo()
1127 // OMP51-LABEL: int main(int argc, char **argv) {
1128 int main (int argc
, char **argv
) {
1129 #pragma omp target defaultmap(present: scalar)
1130 // OMP51-NEXT: #pragma omp target defaultmap(present: scalar)
1132 // OMP51-NEXT: foo();
1133 #pragma omp target defaultmap(present: aggregate)
1134 // OMP51-NEXT: #pragma omp target defaultmap(present: aggregate)
1136 // OMP51-NEXT: foo();
1137 #pragma omp target defaultmap(present: pointer)
1138 // OMP51-NEXT: #pragma omp target defaultmap(present: pointer)
1140 // OMP51-NEXT: foo();
1142 return tmain
<int, 5>(argc
, &argc
) + tmain
<char, 1>(argv
[0][0], argv
[0]);
1148 ///==========================================================================///
1149 // RUN: %clang_cc1 -DOMP52 -verify -Wno-vla -fopenmp -fopenmp-version=52 -ast-print %s | FileCheck %s --check-prefix OMP52
1150 // RUN: %clang_cc1 -DOMP52 -fopenmp -fopenmp-version=52 -x c++ -std=c++11 -emit-pch -o %t %s
1151 // RUN: %clang_cc1 -DOMP52 -fopenmp -fopenmp-version=52 -std=c++11 -include-pch %t -fsyntax-only -verify -Wno-vla %s -ast-print | FileCheck %s --check-prefix OMP52
1153 // RUN: %clang_cc1 -DOMP52 -verify -Wno-vla -fopenmp-simd -fopenmp-version=52 -ast-print %s | FileCheck %s --check-prefix OMP52
1154 // RUN: %clang_cc1 -DOMP52 -fopenmp-simd -fopenmp-version=52 -x c++ -std=c++11 -emit-pch -o %t %s
1155 // RUN: %clang_cc1 -DOMP52 -fopenmp-simd -fopenmp-version=52 -std=c++11 -include-pch %t -fsyntax-only -verify -Wno-vla %s -ast-print | FileCheck %s --check-prefix OMP52
1159 template <typename T
, int C
>
1160 T
tmain(T argc
, T
*argv
) {
1163 #pragma omp target map(iterator(it = 0:N:2), to: v[it])
1165 #pragma omp target map(iterator(it = 0:N:4), from: v[it])
1171 // OMP52: template <typename T, int C> T tmain(T argc, T *argv) {
1172 // OMP52-NEXT: int N = 100;
1173 // OMP52-NEXT: int v[N];
1174 // OMP52-NEXT: #pragma omp target map(iterator(int it = 0:N:2),to: v[it])
1175 // OMP52-NEXT: foo()
1176 // OMP52-NEXT: #pragma omp target map(iterator(int it = 0:N:4),from: v[it])
1177 // OMP52-NEXT: foo()
1179 // OMP52-LABEL: int main(int argc, char **argv) {
1180 int main (int argc
, char **argv
) {
1181 int i
, j
, a
[20], always
, close
;
1182 // OMP52-NEXT: int i, j, a[20]
1184 // OMP52-NEXT: #pragma omp target
1186 // OMP52-NEXT: foo();
1187 #pragma omp target map(iterator(it = 0:20:2), to: a[it])
1188 // OMP52-NEXT: #pragma omp target map(iterator(int it = 0:20:2),to: a[it])
1190 // OMP52-NEXT: foo();
1191 #pragma omp target map(iterator(it = 0:20:4), from: a[it])
1192 // OMP52-NEXT: #pragma omp target map(iterator(int it = 0:20:4),from: a[it])
1194 // OMP52-NEXT: foo();
1196 return tmain
<int, 5>(argc
, &argc
) + tmain
<char, 1>(argv
[0][0], argv
[0]);
1202 // RUN: %clang_cc1 -DOMPX -verify -Wno-vla -fopenmp -fopenmp-extensions -ast-print %s | FileCheck %s --check-prefix=OMPX
1203 // RUN: %clang_cc1 -DOMPX -fopenmp -fopenmp-extensions -x c++ -std=c++11 -emit-pch -o %t %s
1204 // RUN: %clang_cc1 -DOMPX -fopenmp -fopenmp-extensions -std=c++11 -include-pch %t -fsyntax-only -verify -Wno-vla %s -ast-print | FileCheck %s --check-prefix=OMPX
1206 // RUN: %clang_cc1 -DOMPX -verify -Wno-vla -fopenmp-simd -fopenmp-extensions -ast-print %s | FileCheck %s --check-prefix=OMPX
1207 // RUN: %clang_cc1 -DOMPX -fopenmp-simd -fopenmp-extensions -x c++ -std=c++11 -emit-pch -o %t %s
1208 // RUN: %clang_cc1 -DOMPX -fopenmp-simd -fopenmp-extensions -std=c++11 -include-pch %t -fsyntax-only -verify -Wno-vla %s -ast-print | FileCheck %s --check-prefix=OMPX
1212 template <typename T
, int C
>
1213 T
tmain(T argc
, T
*argv
) {
1215 #pragma omp target map(ompx_hold,alloc: i)
1217 #pragma omp target map(ompx_hold from: i)
1219 #pragma omp target map(ompx_hold)
1221 #pragma omp target map(ompx_hold,i)
1226 // OMPX: template <typename T, int C> T tmain(T argc, T *argv) {
1227 // OMPX-NEXT: T i, ompx_hold;
1228 // OMPX-NEXT: #pragma omp target map(ompx_hold,alloc: i)
1230 // OMPX-NEXT: #pragma omp target map(ompx_hold,from: i)
1232 // OMPX-NEXT: #pragma omp target map(tofrom: ompx_hold)
1234 // OMPX-NEXT: ompx_hold++;
1236 // OMPX-NEXT: #pragma omp target map(tofrom: ompx_hold,i)
1238 // OMPX-NEXT: ompx_hold++;
1242 // OMPX-LABEL: int main(int argc, char **argv) {
1243 // OMPX-NEXT: int i, ompx_hold;
1244 // OMPX-NEXT: #pragma omp target map(ompx_hold,alloc: i)
1245 // OMPX-NEXT: foo();
1246 // OMPX-NEXT: #pragma omp target map(ompx_hold,from: i)
1247 // OMPX-NEXT: foo();
1248 // OMPX-NEXT: #pragma omp target map(tofrom: ompx_hold)
1250 // OMPX-NEXT: ompx_hold++;
1252 // OMPX-NEXT: #pragma omp target map(tofrom: ompx_hold,i)
1254 // OMPX-NEXT: ompx_hold++;
1257 int main (int argc
, char **argv
) {
1259 #pragma omp target map(ompx_hold,alloc: i)
1261 #pragma omp target map(ompx_hold from: i)
1263 #pragma omp target map(ompx_hold)
1265 #pragma omp target map(ompx_hold,i)
1267 return tmain
<int, 5>(argc
, &argc
) + tmain
<char, 1>(argv
[0][0], argv
[0]);