1334 lines · cpp
1 2 3// expected-no-diagnostics4#ifndef HEADER5#define HEADER6 7// RUN: %clang_cc1 -DOMP45 -verify -Wno-vla -fopenmp -fopenmp-version=45 -ast-print %s | FileCheck %s --check-prefix=OMP458// RUN: %clang_cc1 -DOMP45 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -emit-pch -o %t %s9// RUN: %clang_cc1 -DOMP45 -fopenmp -fopenmp-version=45 -std=c++11 -include-pch %t -verify -Wno-vla %s -ast-print | FileCheck %s --check-prefix=OMP4510 11// RUN: %clang_cc1 -DOMP45 -verify -Wno-vla -fopenmp-simd -fopenmp-version=45 -ast-print %s | FileCheck %s --check-prefix=OMP4512// RUN: %clang_cc1 -DOMP45 -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -emit-pch -o %t %s13// RUN: %clang_cc1 -DOMP45 -fopenmp-simd -fopenmp-version=45 -std=c++11 -include-pch %t -verify -Wno-vla %s -ast-print | FileCheck %s --check-prefix=OMP4514#ifdef OMP4515 16void foo() {}17 18template <typename T, int C>19T tmain(T argc, T *argv) {20 T i, j, a[20], always, close;21#pragma omp target22 foo();23#pragma omp target if (target:argc > 0)24 foo();25#pragma omp target if (C)26 foo();27#pragma omp target map(i)28 foo();29#pragma omp target map(a[0:10], i)30 foo();31#pragma omp target map(to: i) map(from: j)32 foo();33#pragma omp target map(always,alloc: i)34 foo();35#pragma omp target map(always from: i)36 foo();37#pragma omp target map(always)38 {always++;}39#pragma omp target map(always,i)40 {always++;i++;}41#pragma omp target map(close,alloc: i)42 foo();43#pragma omp target map(close, from: i)44 foo();45#pragma omp target map(close)46 {close++;}47#pragma omp target map(close,i)48 {close++;i++;}49#pragma omp target nowait50 foo();51#pragma omp target depend(in : argc, argv[i:argc], a[:])52 foo();53#pragma omp target defaultmap(tofrom: scalar)54 foo();55 return 0;56}57 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{{$}}61// OMP45-NEXT: foo();62// OMP45-NEXT: #pragma omp target if(target: argc > 0)63// OMP45-NEXT: foo()64// OMP45-NEXT: #pragma omp target if(C)65// OMP45-NEXT: foo()66// OMP45-NEXT: #pragma omp target map(tofrom: i)67// OMP45-NEXT: foo()68// OMP45-NEXT: #pragma omp target map(tofrom: a[0:10],i)69// OMP45-NEXT: foo()70// OMP45-NEXT: #pragma omp target map(to: i) map(from: j)71// OMP45-NEXT: foo()72// OMP45-NEXT: #pragma omp target map(always,alloc: i)73// OMP45-NEXT: foo()74// OMP45-NEXT: #pragma omp target map(always,from: i)75// OMP45-NEXT: foo()76// OMP45-NEXT: #pragma omp target map(tofrom: always)77// OMP45-NEXT: {78// OMP45-NEXT: always++;79// OMP45-NEXT: }80// OMP45-NEXT: #pragma omp target map(tofrom: always,i)81// OMP45-NEXT: {82// OMP45-NEXT: always++;83// OMP45-NEXT: i++;84// OMP45-NEXT: }85// OMP45-NEXT: #pragma omp target map(close,alloc: i)86// OMP45-NEXT: foo()87// OMP45-NEXT: #pragma omp target map(close,from: i)88// OMP45-NEXT: foo()89// OMP45-NEXT: #pragma omp target map(tofrom: close)90// OMP45-NEXT: {91// OMP45-NEXT: close++;92// OMP45-NEXT: }93// OMP45-NEXT: #pragma omp target map(tofrom: close,i)94// OMP45-NEXT: {95// OMP45-NEXT: close++;96// OMP45-NEXT: i++;97// OMP45-NEXT: }98// OMP45-NEXT: #pragma omp target nowait99// OMP45-NEXT: foo()100// OMP45-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])101// OMP45-NEXT: foo()102// OMP45-NEXT: #pragma omp target defaultmap(tofrom: scalar)103// OMP45-NEXT: foo()104// OMP45: template<> int tmain<int, 5>(int argc, int *argv) {105// OMP45-NEXT: int i, j, a[20]106// OMP45-NEXT: #pragma omp target107// OMP45-NEXT: foo();108// OMP45-NEXT: #pragma omp target if(target: argc > 0)109// OMP45-NEXT: foo()110// OMP45-NEXT: #pragma omp target if(5)111// OMP45-NEXT: foo()112// OMP45-NEXT: #pragma omp target map(tofrom: i)113// OMP45-NEXT: foo()114// OMP45-NEXT: #pragma omp target map(tofrom: a[0:10],i)115// OMP45-NEXT: foo()116// OMP45-NEXT: #pragma omp target map(to: i) map(from: j)117// OMP45-NEXT: foo()118// OMP45-NEXT: #pragma omp target map(always,alloc: i)119// OMP45-NEXT: foo()120// OMP45-NEXT: #pragma omp target map(always,from: i)121// OMP45-NEXT: foo()122// OMP45-NEXT: #pragma omp target map(tofrom: always)123// OMP45-NEXT: {124// OMP45-NEXT: always++;125// OMP45-NEXT: }126// OMP45-NEXT: #pragma omp target map(tofrom: always,i)127// OMP45-NEXT: {128// OMP45-NEXT: always++;129// OMP45-NEXT: i++;130// OMP45-NEXT: }131// OMP45-NEXT: #pragma omp target map(close,alloc: i)132// OMP45-NEXT: foo()133// OMP45-NEXT: #pragma omp target map(close,from: i)134// OMP45-NEXT: foo()135// OMP45-NEXT: #pragma omp target map(tofrom: close)136// OMP45-NEXT: {137// OMP45-NEXT: close++;138// OMP45-NEXT: }139// OMP45-NEXT: #pragma omp target map(tofrom: close,i)140// OMP45-NEXT: {141// OMP45-NEXT: close++;142// OMP45-NEXT: i++;143// OMP45-NEXT: }144// OMP45-NEXT: #pragma omp target nowait145// OMP45-NEXT: foo()146// OMP45-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])147// OMP45-NEXT: foo()148// OMP45-NEXT: #pragma omp target defaultmap(tofrom: scalar)149// OMP45-NEXT: foo()150// OMP45: template<> char tmain<char, 1>(char argc, char *argv) {151// OMP45-NEXT: char i, j, a[20]152// OMP45-NEXT: #pragma omp target153// OMP45-NEXT: foo();154// OMP45-NEXT: #pragma omp target if(target: argc > 0)155// OMP45-NEXT: foo()156// OMP45-NEXT: #pragma omp target if(1)157// OMP45-NEXT: foo()158// OMP45-NEXT: #pragma omp target map(tofrom: i)159// OMP45-NEXT: foo()160// OMP45-NEXT: #pragma omp target map(tofrom: a[0:10],i)161// OMP45-NEXT: foo()162// OMP45-NEXT: #pragma omp target map(to: i) map(from: j)163// OMP45-NEXT: foo()164// OMP45-NEXT: #pragma omp target map(always,alloc: i)165// OMP45-NEXT: foo()166// OMP45-NEXT: #pragma omp target map(always,from: i)167// OMP45-NEXT: foo()168// OMP45-NEXT: #pragma omp target map(tofrom: always)169// OMP45-NEXT: {170// OMP45-NEXT: always++;171// OMP45-NEXT: }172// OMP45-NEXT: #pragma omp target map(tofrom: always,i)173// OMP45-NEXT: {174// OMP45-NEXT: always++;175// OMP45-NEXT: i++;176// OMP45-NEXT: }177// OMP45-NEXT: #pragma omp target map(close,alloc: i)178// OMP45-NEXT: foo()179// OMP45-NEXT: #pragma omp target map(close,from: i)180// OMP45-NEXT: foo()181// OMP45-NEXT: #pragma omp target map(tofrom: close)182// OMP45-NEXT: {183// OMP45-NEXT: close++;184// OMP45-NEXT: }185// OMP45-NEXT: #pragma omp target map(tofrom: close,i)186// OMP45-NEXT: {187// OMP45-NEXT: close++;188// OMP45-NEXT: i++;189// OMP45-NEXT: }190// OMP45-NEXT: #pragma omp target nowait191// OMP45-NEXT: foo()192// OMP45-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])193// OMP45-NEXT: foo()194// OMP45-NEXT: #pragma omp target defaultmap(tofrom: scalar)195// OMP45-NEXT: foo()196 197// OMP45-LABEL: class S {198class S {199 void foo() {200// OMP45-NEXT: void foo() {201 int a = 0;202// OMP45-NEXT: int a = 0;203 #pragma omp target map(this[0])204// OMP45-NEXT: #pragma omp target map(tofrom: this[0])205 a++;206// OMP45-NEXT: a++;207 #pragma omp target map(this[:1])208// OMP45-NEXT: #pragma omp target map(tofrom: this[:1])209 a++;210// OMP45-NEXT: a++;211 #pragma omp target map((this)[0])212// OMP45-NEXT: #pragma omp target map(tofrom: (this)[0])213 a++;214// OMP45-NEXT: a++;215 #pragma omp target map(this[:a])216// OMP45-NEXT: #pragma omp target map(tofrom: this[:a])217 a++;218// OMP45-NEXT: a++;219 #pragma omp target map(this[a:1])220// OMP45-NEXT: #pragma omp target map(tofrom: this[a:1])221 a++;222// OMP45-NEXT: a++;223 #pragma omp target map(this[a])224// OMP45-NEXT: #pragma omp target map(tofrom: this[a])225 a++;226// OMP45-NEXT: a++;227 }228// OMP45-NEXT: }229};230// OMP45-NEXT: };231 232// OMP45-LABEL: int main(int argc, char **argv) {233int main (int argc, char **argv) {234 int i, j, a[20], always, close;235// OMP45-NEXT: int i, j, a[20]236#pragma omp target237// OMP45-NEXT: #pragma omp target238 foo();239// OMP45-NEXT: foo();240#pragma omp target if (argc > 0)241// OMP45-NEXT: #pragma omp target if(argc > 0)242 foo();243// OMP45-NEXT: foo();244 245#pragma omp target map(i) if(argc>0)246// OMP45-NEXT: #pragma omp target map(tofrom: i) if(argc > 0)247 foo();248// OMP45-NEXT: foo();249 250#pragma omp target map(i)251// OMP45-NEXT: #pragma omp target map(tofrom: i)252 foo();253// OMP45-NEXT: foo();254 255#pragma omp target map(a[0:10], i)256// OMP45-NEXT: #pragma omp target map(tofrom: a[0:10],i)257 foo();258// OMP45-NEXT: foo();259 260#pragma omp target map(to: i) map(from: j)261// OMP45-NEXT: #pragma omp target map(to: i) map(from: j)262 foo();263// OMP45-NEXT: foo();264 265#pragma omp target map(always,alloc: i)266// OMP45-NEXT: #pragma omp target map(always,alloc: i)267 foo();268// OMP45-NEXT: foo();269 270#pragma omp target map(always from: i)271// OMP45-NEXT: #pragma omp target map(always,from: i)272 foo();273// OMP45-NEXT: foo();274 275#pragma omp target map(always)276// OMP45-NEXT: #pragma omp target map(tofrom: always)277 {always++;}278// OMP45-NEXT: {279// OMP45-NEXT: always++;280// OMP45-NEXT: }281 282#pragma omp target map(always,i)283// OMP45-NEXT: #pragma omp target map(tofrom: always,i)284 {always++;i++;}285// OMP45-NEXT: {286// OMP45-NEXT: always++;287// OMP45-NEXT: i++;288// OMP45-NEXT: }289 290#pragma omp target map(close,alloc: i)291// OMP45-NEXT: #pragma omp target map(close,alloc: i)292 foo();293// OMP45-NEXT: foo();294 295#pragma omp target map(close from: i)296// OMP45-NEXT: #pragma omp target map(close,from: i)297 foo();298// OMP45-NEXT: foo();299 300#pragma omp target map(close)301// OMP45-NEXT: #pragma omp target map(tofrom: close)302 {close++;}303// OMP45-NEXT: {304// OMP45-NEXT: close++;305// OMP45-NEXT: }306 307#pragma omp target map(close,i)308// OMP45-NEXT: #pragma omp target map(tofrom: close,i)309 {close++;i++;}310// OMP45-NEXT: {311// OMP45-NEXT: close++;312// OMP45-NEXT: i++;313// OMP45-NEXT: }314 315#pragma omp target nowait316// OMP45-NEXT: #pragma omp target nowait317 foo();318// OMP45-NEXT: foo();319 320#pragma omp target depend(in : argc, argv[i:argc], a[:])321// OMP45-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])322 foo();323// OMP45-NEXT: foo();324 325#pragma omp target defaultmap(tofrom: scalar)326// OMP45-NEXT: #pragma omp target defaultmap(tofrom: scalar)327 foo();328// OMP45-NEXT: foo();329 330 return tmain<int, 5>(argc, &argc) + tmain<char, 1>(argv[0][0], argv[0]);331}332 333#endif334 335#ifdef OMP5336 337///==========================================================================///338// RUN: %clang_cc1 -DOMP5 -verify -Wno-vla -fopenmp -fopenmp-version=50 -ast-print %s | FileCheck %s --check-prefix OMP5339// RUN: %clang_cc1 -DOMP5 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s340// RUN: %clang_cc1 -DOMP5 -fopenmp -fopenmp-version=50 -std=c++11 -include-pch %t -verify -Wno-vla %s -ast-print | FileCheck %s --check-prefix OMP5341 342// RUN: %clang_cc1 -DOMP5 -verify -Wno-vla -fopenmp-simd -fopenmp-version=50 -ast-print %s | FileCheck %s --check-prefix OMP5343// RUN: %clang_cc1 -DOMP5 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s344// RUN: %clang_cc1 -DOMP5 -fopenmp-simd -fopenmp-version=50 -std=c++11 -include-pch %t -verify -Wno-vla %s -ast-print | FileCheck %s --check-prefix OMP5345 346// RUN: %clang_cc1 -DOMP5 -verify -Wno-vla -fopenmp -fopenmp-version=99 -DOMP99 -ast-print %s | FileCheck %s --check-prefixes=OMP5,REV347// RUN: %clang_cc1 -DOMP5 -fopenmp -fopenmp-version=99 -DOMP99 -x c++ -std=c++11 -emit-pch -o %t %s348// RUN: %clang_cc1 -DOMP5 -fopenmp -fopenmp-version=99 -DOMP99 -std=c++11 -include-pch %t -verify -Wno-vla %s -ast-print | FileCheck %s --check-prefixes=OMP5,REV349 350// RUN: %clang_cc1 -DOMP5 -verify -Wno-vla -fopenmp-simd -fopenmp-version=99 -DOMP99 -ast-print %s | FileCheck %s --check-prefixes=OMP5,REV351// RUN: %clang_cc1 -DOMP5 -fopenmp-simd -fopenmp-version=99 -DOMP99 -x c++ -std=c++11 -emit-pch -o %t %s352// RUN: %clang_cc1 -DOMP5 -fopenmp-simd -fopenmp-version=99 -DOMP99 -std=c++11 -include-pch %t -verify -Wno-vla %s -ast-print | FileCheck %s --check-prefixes=OMP5,REV353 354#ifdef OMP99355#pragma omp requires reverse_offload356#endif357typedef void **omp_allocator_handle_t;358extern const omp_allocator_handle_t omp_null_allocator;359extern const omp_allocator_handle_t omp_default_mem_alloc;360extern const omp_allocator_handle_t omp_large_cap_mem_alloc;361extern const omp_allocator_handle_t omp_const_mem_alloc;362extern const omp_allocator_handle_t omp_high_bw_mem_alloc;363extern const omp_allocator_handle_t omp_low_lat_mem_alloc;364extern const omp_allocator_handle_t omp_cgroup_mem_alloc;365extern const omp_allocator_handle_t omp_pteam_mem_alloc;366extern const omp_allocator_handle_t omp_thread_mem_alloc;367 368void foo() {}369 370#pragma omp begin declare target371void bar() {}372#pragma omp end declare target373 374int a;375#pragma omp declare target link(a)376 377template <typename T, int C>378T tmain(T argc, T *argv) {379 T i, j, a[20], always, close;380#pragma omp target device(argc)381 foo();382#pragma omp target if (target:argc > 0) device(device_num: C)383 foo();384#ifdef OMP99385#pragma omp target if (C) device(ancestor: argc)386 foo();387#endif388#pragma omp target map(i)389 foo();390#pragma omp target map(a[0:10], i)391 foo();392#pragma omp target map(to: i) map(from: j)393 foo();394#pragma omp target map(always,alloc: i)395 foo();396#pragma omp target map(always, from: i)397 foo();398#pragma omp target map(always)399 {always++;}400#pragma omp target map(always,i)401 {always++;i++;}402#pragma omp target map(close,alloc: i)403 foo();404#pragma omp target map(close, from: i)405 foo();406#pragma omp target map(close)407 {close++;}408#pragma omp target map(close,i)409 {close++;i++;}410#pragma omp target nowait411 foo();412#pragma omp target depend(in : argc, argv[i:argc], a[:])413 foo();414#pragma omp target defaultmap(alloc: scalar)415 foo();416#pragma omp target defaultmap(to: scalar)417 foo();418#pragma omp target defaultmap(from: scalar)419 foo();420#pragma omp target defaultmap(tofrom: scalar)421 foo();422#pragma omp target defaultmap(firstprivate: scalar)423 foo();424#pragma omp target defaultmap(none: scalar)425 foo();426#pragma omp target defaultmap(default: scalar)427 foo();428#pragma omp target defaultmap(alloc: aggregate)429 foo();430#pragma omp target defaultmap(to: aggregate)431 foo();432#pragma omp target defaultmap(from: aggregate)433 foo();434#pragma omp target defaultmap(tofrom: aggregate)435 foo();436#pragma omp target defaultmap(firstprivate: aggregate)437 foo();438#pragma omp target defaultmap(none: aggregate)439 foo();440#pragma omp target defaultmap(default: aggregate)441 foo();442#pragma omp target defaultmap(alloc: pointer)443 foo();444#pragma omp target defaultmap(to: pointer)445 foo();446#pragma omp target defaultmap(from: pointer)447 foo();448#pragma omp target defaultmap(tofrom: pointer)449 foo();450#pragma omp target defaultmap(firstprivate: pointer)451 foo();452#pragma omp target defaultmap(none: pointer)453 foo();454#pragma omp target defaultmap(default: pointer)455 foo();456#pragma omp target defaultmap(to: scalar) defaultmap(tofrom: pointer)457 foo();458#pragma omp target defaultmap(from: pointer) defaultmap(none: aggregate)459 foo();460#pragma omp target defaultmap(default: aggregate) defaultmap(alloc: scalar)461 foo();462#pragma omp target defaultmap(alloc: aggregate) defaultmap(firstprivate: scalar) defaultmap(tofrom: pointer)463 foo();464#pragma omp target defaultmap(tofrom: aggregate) defaultmap(to: pointer) defaultmap(alloc: scalar)465 foo();466 467 int *g;468 469#pragma omp target is_device_ptr(g) defaultmap(none: pointer)470 g++;471#pragma omp target private(g) defaultmap(none: pointer)472 g++;473#pragma omp target firstprivate(g) defaultmap(none: pointer)474 g++;475#pragma omp target defaultmap(none: scalar) map(to: i)476 i++;477#pragma omp target defaultmap(none: aggregate) map(to: a)478 a[3]++;479#pragma omp target defaultmap(none: scalar)480 bar();481 482 return 0;483}484 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){{$}}488// OMP5-NEXT: foo();489// OMP5-NEXT: #pragma omp target if(target: argc > 0) device(device_num: C)490// OMP5-NEXT: foo()491// REV: #pragma omp target if(C) device(ancestor: argc)492// REV-NEXT: foo()493// OMP5-NEXT: #pragma omp target map(tofrom: i)494// OMP5-NEXT: foo()495// OMP5-NEXT: #pragma omp target map(tofrom: a[0:10],i)496// OMP5-NEXT: foo()497// OMP5-NEXT: #pragma omp target map(to: i) map(from: j)498// OMP5-NEXT: foo()499// OMP5-NEXT: #pragma omp target map(always,alloc: i)500// OMP5-NEXT: foo()501// OMP5-NEXT: #pragma omp target map(always,from: i)502// OMP5-NEXT: foo()503// OMP5-NEXT: #pragma omp target map(tofrom: always)504// OMP5-NEXT: {505// OMP5-NEXT: always++;506// OMP5-NEXT: }507// OMP5-NEXT: #pragma omp target map(tofrom: always,i)508// OMP5-NEXT: {509// OMP5-NEXT: always++;510// OMP5-NEXT: i++;511// OMP5-NEXT: }512// OMP5-NEXT: #pragma omp target map(close,alloc: i)513// OMP5-NEXT: foo()514// OMP5-NEXT: #pragma omp target map(close,from: i)515// OMP5-NEXT: foo()516// OMP5-NEXT: #pragma omp target map(tofrom: close)517// OMP5-NEXT: {518// OMP5-NEXT: close++;519// OMP5-NEXT: }520// OMP5-NEXT: #pragma omp target map(tofrom: close,i)521// OMP5-NEXT: {522// OMP5-NEXT: close++;523// OMP5-NEXT: i++;524// OMP5-NEXT: }525// OMP5-NEXT: #pragma omp target nowait526// OMP5-NEXT: foo()527// OMP5-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])528// OMP5-NEXT: foo()529// OMP5-NEXT: #pragma omp target defaultmap(alloc: scalar)530// OMP5-NEXT: foo()531// OMP5-NEXT: #pragma omp target defaultmap(to: scalar)532// OMP5-NEXT: foo()533// OMP5-NEXT: #pragma omp target defaultmap(from: scalar)534// OMP5-NEXT: foo()535// OMP5-NEXT: #pragma omp target defaultmap(tofrom: scalar)536// OMP5-NEXT: foo()537// OMP5-NEXT: #pragma omp target defaultmap(firstprivate: scalar)538// OMP5-NEXT: foo()539// OMP5-NEXT: #pragma omp target defaultmap(none: scalar)540// OMP5-NEXT: foo()541// OMP5-NEXT: #pragma omp target defaultmap(default: scalar)542// OMP5-NEXT: foo()543// OMP5-NEXT: #pragma omp target defaultmap(alloc: aggregate)544// OMP5-NEXT: foo()545// OMP5-NEXT: #pragma omp target defaultmap(to: aggregate)546// OMP5-NEXT: foo()547// OMP5-NEXT: #pragma omp target defaultmap(from: aggregate)548// OMP5-NEXT: foo()549// OMP5-NEXT: #pragma omp target defaultmap(tofrom: aggregate)550// OMP5-NEXT: foo()551// OMP5-NEXT: #pragma omp target defaultmap(firstprivate: aggregate)552// OMP5-NEXT: foo()553// OMP5-NEXT: #pragma omp target defaultmap(none: aggregate)554// OMP5-NEXT: foo()555// OMP5-NEXT: #pragma omp target defaultmap(default: aggregate)556// OMP5-NEXT: foo()557// OMP5-NEXT: #pragma omp target defaultmap(alloc: pointer)558// OMP5-NEXT: foo()559// OMP5-NEXT: #pragma omp target defaultmap(to: pointer)560// OMP5-NEXT: foo()561// OMP5-NEXT: #pragma omp target defaultmap(from: pointer)562// OMP5-NEXT: foo()563// OMP5-NEXT: #pragma omp target defaultmap(tofrom: pointer)564// OMP5-NEXT: foo()565// OMP5-NEXT: #pragma omp target defaultmap(firstprivate: pointer)566// OMP5-NEXT: foo()567// OMP5-NEXT: #pragma omp target defaultmap(none: pointer)568// OMP5-NEXT: foo()569// OMP5-NEXT: #pragma omp target defaultmap(default: pointer)570// OMP5-NEXT: foo()571// OMP5-NEXT: #pragma omp target defaultmap(to: scalar) defaultmap(tofrom: pointer)572// OMP5-NEXT: foo()573// OMP5-NEXT: #pragma omp target defaultmap(from: pointer) defaultmap(none: aggregate)574// OMP5-NEXT: foo()575// OMP5-NEXT: #pragma omp target defaultmap(default: aggregate) defaultmap(alloc: scalar)576// OMP5-NEXT: foo()577// OMP5-NEXT: #pragma omp target defaultmap(alloc: aggregate) defaultmap(firstprivate: scalar) defaultmap(tofrom: pointer)578// OMP5-NEXT: foo()579// OMP5-NEXT: #pragma omp target defaultmap(tofrom: aggregate) defaultmap(to: pointer) defaultmap(alloc: scalar)580// OMP5-NEXT: foo()581// OMP5: template<> int tmain<int, 5>(int argc, int *argv) {582// OMP5-NEXT: int i, j, a[20]583// OMP5-NEXT: #pragma omp target584// OMP5-NEXT: foo();585// OMP5-NEXT: #pragma omp target if(target: argc > 0)586// OMP5-NEXT: foo()587// REV: #pragma omp target if(5)588// REV-NEXT: foo()589// OMP5-NEXT: #pragma omp target map(tofrom: i)590// OMP5-NEXT: foo()591// OMP5-NEXT: #pragma omp target map(tofrom: a[0:10],i)592// OMP5-NEXT: foo()593// OMP5-NEXT: #pragma omp target map(to: i) map(from: j)594// OMP5-NEXT: foo()595// OMP5-NEXT: #pragma omp target map(always,alloc: i)596// OMP5-NEXT: foo()597// OMP5-NEXT: #pragma omp target map(always,from: i)598// OMP5-NEXT: foo()599// OMP5-NEXT: #pragma omp target map(tofrom: always)600// OMP5-NEXT: {601// OMP5-NEXT: always++;602// OMP5-NEXT: }603// OMP5-NEXT: #pragma omp target map(tofrom: always,i)604// OMP5-NEXT: {605// OMP5-NEXT: always++;606// OMP5-NEXT: i++;607// OMP5-NEXT: }608// OMP5-NEXT: #pragma omp target map(close,alloc: i)609// OMP5-NEXT: foo()610// OMP5-NEXT: #pragma omp target map(close,from: i)611// OMP5-NEXT: foo()612// OMP5-NEXT: #pragma omp target map(tofrom: close)613// OMP5-NEXT: {614// OMP5-NEXT: close++;615// OMP5-NEXT: }616// OMP5-NEXT: #pragma omp target map(tofrom: close,i)617// OMP5-NEXT: {618// OMP5-NEXT: close++;619// OMP5-NEXT: i++;620// OMP5-NEXT: }621// OMP5-NEXT: #pragma omp target nowait622// OMP5-NEXT: foo()623// OMP5-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])624// OMP5-NEXT: foo()625// OMP5-NEXT: #pragma omp target defaultmap(alloc: scalar)626// OMP5-NEXT: foo()627// OMP5-NEXT: #pragma omp target defaultmap(to: scalar)628// OMP5-NEXT: foo()629// OMP5-NEXT: #pragma omp target defaultmap(from: scalar)630// OMP5-NEXT: foo()631// OMP5-NEXT: #pragma omp target defaultmap(tofrom: scalar)632// OMP5-NEXT: foo()633// OMP5-NEXT: #pragma omp target defaultmap(firstprivate: scalar)634// OMP5-NEXT: foo()635// OMP5-NEXT: #pragma omp target defaultmap(none: scalar)636// OMP5-NEXT: foo()637// OMP5-NEXT: #pragma omp target defaultmap(default: scalar)638// OMP5-NEXT: foo()639// OMP5-NEXT: #pragma omp target defaultmap(alloc: aggregate)640// OMP5-NEXT: foo()641// OMP5-NEXT: #pragma omp target defaultmap(to: aggregate)642// OMP5-NEXT: foo()643// OMP5-NEXT: #pragma omp target defaultmap(from: aggregate)644// OMP5-NEXT: foo()645// OMP5-NEXT: #pragma omp target defaultmap(tofrom: aggregate)646// OMP5-NEXT: foo()647// OMP5-NEXT: #pragma omp target defaultmap(firstprivate: aggregate)648// OMP5-NEXT: foo()649// OMP5-NEXT: #pragma omp target defaultmap(none: aggregate)650// OMP5-NEXT: foo()651// OMP5-NEXT: #pragma omp target defaultmap(default: aggregate)652// OMP5-NEXT: foo()653// OMP5-NEXT: #pragma omp target defaultmap(alloc: pointer)654// OMP5-NEXT: foo()655// OMP5-NEXT: #pragma omp target defaultmap(to: pointer)656// OMP5-NEXT: foo()657// OMP5-NEXT: #pragma omp target defaultmap(from: pointer)658// OMP5-NEXT: foo()659// OMP5-NEXT: #pragma omp target defaultmap(tofrom: pointer)660// OMP5-NEXT: foo()661// OMP5-NEXT: #pragma omp target defaultmap(firstprivate: pointer)662// OMP5-NEXT: foo()663// OMP5-NEXT: #pragma omp target defaultmap(none: pointer)664// OMP5-NEXT: foo()665// OMP5-NEXT: #pragma omp target defaultmap(default: pointer)666// OMP5-NEXT: foo()667// OMP5-NEXT: #pragma omp target defaultmap(to: scalar) defaultmap(tofrom: pointer)668// OMP5-NEXT: foo()669// OMP5-NEXT: #pragma omp target defaultmap(from: pointer) defaultmap(none: aggregate)670// OMP5-NEXT: foo()671// OMP5-NEXT: #pragma omp target defaultmap(default: aggregate) defaultmap(alloc: scalar)672// OMP5-NEXT: foo()673// OMP5-NEXT: #pragma omp target defaultmap(alloc: aggregate) defaultmap(firstprivate: scalar) defaultmap(tofrom: pointer)674// OMP5-NEXT: foo()675// OMP5-NEXT: #pragma omp target defaultmap(tofrom: aggregate) defaultmap(to: pointer) defaultmap(alloc: scalar)676// OMP5-NEXT: foo()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)680// OMP5-NEXT: foo();681// OMP5-NEXT: #pragma omp target if(target: argc > 0) device(device_num: 1)682// OMP5-NEXT: foo()683// REV: #pragma omp target if(1) device(ancestor: argc)684// REV-NEXT: foo()685// OMP5-NEXT: #pragma omp target map(tofrom: i)686// OMP5-NEXT: foo()687// OMP5-NEXT: #pragma omp target map(tofrom: a[0:10],i)688// OMP5-NEXT: foo()689// OMP5-NEXT: #pragma omp target map(to: i) map(from: j)690// OMP5-NEXT: foo()691// OMP5-NEXT: #pragma omp target map(always,alloc: i)692// OMP5-NEXT: foo()693// OMP5-NEXT: #pragma omp target map(always,from: i)694// OMP5-NEXT: foo()695// OMP5-NEXT: #pragma omp target map(tofrom: always)696// OMP5-NEXT: {697// OMP5-NEXT: always++;698// OMP5-NEXT: }699// OMP5-NEXT: #pragma omp target map(tofrom: always,i)700// OMP5-NEXT: {701// OMP5-NEXT: always++;702// OMP5-NEXT: i++;703// OMP5-NEXT: }704// OMP5-NEXT: #pragma omp target map(close,alloc: i)705// OMP5-NEXT: foo()706// OMP5-NEXT: #pragma omp target map(close,from: i)707// OMP5-NEXT: foo()708// OMP5-NEXT: #pragma omp target map(tofrom: close)709// OMP5-NEXT: {710// OMP5-NEXT: close++;711// OMP5-NEXT: }712// OMP5-NEXT: #pragma omp target map(tofrom: close,i)713// OMP5-NEXT: {714// OMP5-NEXT: close++;715// OMP5-NEXT: i++;716// OMP5-NEXT: }717// OMP5-NEXT: #pragma omp target nowait718// OMP5-NEXT: foo()719// OMP5-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])720// OMP5-NEXT: foo()721// OMP5-NEXT: #pragma omp target defaultmap(alloc: scalar)722// OMP5-NEXT: foo()723// OMP5-NEXT: #pragma omp target defaultmap(to: scalar)724// OMP5-NEXT: foo()725// OMP5-NEXT: #pragma omp target defaultmap(from: scalar)726// OMP5-NEXT: foo()727// OMP5-NEXT: #pragma omp target defaultmap(tofrom: scalar)728// OMP5-NEXT: foo()729// OMP5-NEXT: #pragma omp target defaultmap(firstprivate: scalar)730// OMP5-NEXT: foo()731// OMP5-NEXT: #pragma omp target defaultmap(none: scalar)732// OMP5-NEXT: foo()733// OMP5-NEXT: #pragma omp target defaultmap(default: scalar)734// OMP5-NEXT: foo()735// OMP5-NEXT: #pragma omp target defaultmap(alloc: aggregate)736// OMP5-NEXT: foo()737// OMP5-NEXT: #pragma omp target defaultmap(to: aggregate)738// OMP5-NEXT: foo()739// OMP5-NEXT: #pragma omp target defaultmap(from: aggregate)740// OMP5-NEXT: foo()741// OMP5-NEXT: #pragma omp target defaultmap(tofrom: aggregate)742// OMP5-NEXT: foo()743// OMP5-NEXT: #pragma omp target defaultmap(firstprivate: aggregate)744// OMP5-NEXT: foo()745// OMP5-NEXT: #pragma omp target defaultmap(none: aggregate)746// OMP5-NEXT: foo()747// OMP5-NEXT: #pragma omp target defaultmap(default: aggregate)748// OMP5-NEXT: foo()749// OMP5-NEXT: #pragma omp target defaultmap(alloc: pointer)750// OMP5-NEXT: foo()751// OMP5-NEXT: #pragma omp target defaultmap(to: pointer)752// OMP5-NEXT: foo()753// OMP5-NEXT: #pragma omp target defaultmap(from: pointer)754// OMP5-NEXT: foo()755// OMP5-NEXT: #pragma omp target defaultmap(tofrom: pointer)756// OMP5-NEXT: foo()757// OMP5-NEXT: #pragma omp target defaultmap(firstprivate: pointer)758// OMP5-NEXT: foo()759// OMP5-NEXT: #pragma omp target defaultmap(none: pointer)760// OMP5-NEXT: foo()761// OMP5-NEXT: #pragma omp target defaultmap(default: pointer)762// OMP5-NEXT: foo()763// OMP5-NEXT: #pragma omp target defaultmap(to: scalar) defaultmap(tofrom: pointer)764// OMP5-NEXT: foo()765// OMP5-NEXT: #pragma omp target defaultmap(from: pointer) defaultmap(none: aggregate)766// OMP5-NEXT: foo()767// OMP5-NEXT: #pragma omp target defaultmap(default: aggregate) defaultmap(alloc: scalar)768// OMP5-NEXT: foo()769// OMP5-NEXT: #pragma omp target defaultmap(alloc: aggregate) defaultmap(firstprivate: scalar) defaultmap(tofrom: pointer)770// OMP5-NEXT: foo()771// OMP5-NEXT: #pragma omp target defaultmap(tofrom: aggregate) defaultmap(to: pointer) defaultmap(alloc: scalar)772// OMP5-NEXT: foo()773// OMP5-NEXT: int *g;774// OMP5-NEXT: #pragma omp target is_device_ptr(g) defaultmap(none: pointer)775// OMP5-NEXT: g++;776// OMP5-NEXT: #pragma omp target private(g) defaultmap(none: pointer)777// OMP5-NEXT: g++;778// OMP5-NEXT: #pragma omp target firstprivate(g) defaultmap(none: pointer)779// OMP5-NEXT: g++;780// OMP5-NEXT: #pragma omp target defaultmap(none: scalar) map(to: i)781// OMP5-NEXT: 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)785// OMP5-NEXT: bar();786 787// OMP5-LABEL: class S {788class S {789 void foo() {790// OMP5-NEXT: void foo() {791 int a = 0;792// OMP5-NEXT: int a = 0;793 #pragma omp target map(this[0])794// OMP5-NEXT: #pragma omp target map(tofrom: this[0])795 a++;796// OMP5-NEXT: a++;797 #pragma omp target map(this[:1])798// OMP5-NEXT: #pragma omp target map(tofrom: this[:1])799 a++;800// OMP5-NEXT: a++;801 #pragma omp target map((this)[0])802// OMP5-NEXT: #pragma omp target map(tofrom: (this)[0])803 a++;804// OMP5-NEXT: a++;805 #pragma omp target map(this[:a])806// OMP5-NEXT: #pragma omp target map(tofrom: this[:a])807 a++;808// OMP5-NEXT: a++;809 #pragma omp target map(this[a:1])810// OMP5-NEXT: #pragma omp target map(tofrom: this[a:1])811 a++;812// OMP5-NEXT: a++;813 #pragma omp target map(this[a])814// OMP5-NEXT: #pragma omp target map(tofrom: this[a])815 a++;816// OMP5-NEXT: a++;817 }818// OMP5-NEXT: }819};820// OMP5-NEXT: };821 822// OMP5-LABEL: int main(int argc, char **argv) {823int main (int argc, char **argv) {824 int i, j, a[20], always, close;825// OMP5-NEXT: int i, j, a[20]826#pragma omp target827// OMP5-NEXT: #pragma omp target828 foo();829// OMP5-NEXT: foo();830#pragma omp target if (argc > 0)831// OMP5-NEXT: #pragma omp target if(argc > 0)832 foo();833// OMP5-NEXT: foo();834 835#pragma omp target map(i) if(argc>0)836// OMP5-NEXT: #pragma omp target map(tofrom: i) if(argc > 0)837 foo();838// OMP5-NEXT: foo();839 840#pragma omp target map(i)841// OMP5-NEXT: #pragma omp target map(tofrom: i)842 foo();843// OMP5-NEXT: foo();844 845#pragma omp target map(a[0:10], i)846// OMP5-NEXT: #pragma omp target map(tofrom: a[0:10],i)847 foo();848// OMP5-NEXT: foo();849 850#pragma omp target map(to: i) map(from: j)851// OMP5-NEXT: #pragma omp target map(to: i) map(from: j)852 foo();853// OMP5-NEXT: foo();854 855#pragma omp target map(always,alloc: i)856// OMP5-NEXT: #pragma omp target map(always,alloc: i)857 foo();858// OMP5-NEXT: foo();859 860#pragma omp target map(always, from: i)861// OMP5-NEXT: #pragma omp target map(always,from: i)862 foo();863// OMP5-NEXT: foo();864 865#pragma omp target map(always)866// OMP5-NEXT: #pragma omp target map(tofrom: always)867 {always++;}868// OMP5-NEXT: {869// OMP5-NEXT: always++;870// OMP5-NEXT: }871 872#pragma omp target map(always,i)873// OMP5-NEXT: #pragma omp target map(tofrom: always,i)874 {always++;i++;}875// OMP5-NEXT: {876// OMP5-NEXT: always++;877// OMP5-NEXT: i++;878// OMP5-NEXT: }879 880#pragma omp target map(close,alloc: i)881// OMP5-NEXT: #pragma omp target map(close,alloc: i)882 foo();883// OMP5-NEXT: foo();884 885#pragma omp target map(close, from: i)886// OMP5-NEXT: #pragma omp target map(close,from: i)887 foo();888// OMP5-NEXT: foo();889 890#pragma omp target map(close)891// OMP5-NEXT: #pragma omp target map(tofrom: close)892 {close++;}893// OMP5-NEXT: {894// OMP5-NEXT: close++;895// OMP5-NEXT: }896 897#pragma omp target map(close,i)898// OMP5-NEXT: #pragma omp target map(tofrom: close,i)899 {close++;i++;}900// OMP5-NEXT: {901// OMP5-NEXT: close++;902// OMP5-NEXT: i++;903// OMP5-NEXT: }904 905#pragma omp target nowait906// OMP5-NEXT: #pragma omp target nowait907 foo();908// OMP5-NEXT: foo();909 910#pragma omp target depend(in : argc, argv[i:argc], a[:])911// OMP5-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])912 foo();913// OMP5-NEXT: foo();914 915#pragma omp target defaultmap(alloc: scalar)916// OMP5-NEXT: #pragma omp target defaultmap(alloc: scalar)917 foo();918// OMP5-NEXT: foo();919 920#pragma omp target defaultmap(to: scalar)921// OMP5-NEXT: #pragma omp target defaultmap(to: scalar)922 foo();923// OMP5-NEXT: foo();924 925#pragma omp target defaultmap(from: scalar)926// OMP5-NEXT: #pragma omp target defaultmap(from: scalar)927 foo();928// OMP5-NEXT: foo();929 930#pragma omp target defaultmap(tofrom: scalar)931// OMP5-NEXT: #pragma omp target defaultmap(tofrom: scalar)932 foo();933// OMP5-NEXT: foo();934 935#pragma omp target defaultmap(firstprivate: scalar)936// OMP5-NEXT: #pragma omp target defaultmap(firstprivate: scalar)937 foo();938// OMP5-NEXT: foo();939 940#pragma omp target defaultmap(none: scalar)941// OMP5-NEXT: #pragma omp target defaultmap(none: scalar)942 foo();943// OMP5-NEXT: foo();944 945#pragma omp target defaultmap(default: scalar)946// OMP5-NEXT: #pragma omp target defaultmap(default: scalar)947 foo();948// OMP5-NEXT: foo();949 950#pragma omp target defaultmap(alloc: aggregate)951// OMP5-NEXT: #pragma omp target defaultmap(alloc: aggregate)952 foo();953// OMP5-NEXT: foo();954 955#pragma omp target defaultmap(to: aggregate)956// OMP5-NEXT: #pragma omp target defaultmap(to: aggregate)957 foo();958// OMP5-NEXT: foo();959 960#pragma omp target defaultmap(from: aggregate)961// OMP5-NEXT: #pragma omp target defaultmap(from: aggregate)962 foo();963// OMP5-NEXT: foo();964 965#pragma omp target defaultmap(tofrom: aggregate)966// OMP5-NEXT: #pragma omp target defaultmap(tofrom: aggregate)967 foo();968// OMP5-NEXT: foo();969 970#pragma omp target defaultmap(firstprivate: aggregate)971// OMP5-NEXT: #pragma omp target defaultmap(firstprivate: aggregate)972 foo();973// OMP5-NEXT: foo();974 975#pragma omp target defaultmap(none: aggregate)976// OMP5-NEXT: #pragma omp target defaultmap(none: aggregate)977 foo();978// OMP5-NEXT: foo();979 980#pragma omp target defaultmap(default: aggregate)981// OMP5-NEXT: #pragma omp target defaultmap(default: aggregate)982 foo();983// OMP5-NEXT: foo();984 985#pragma omp target defaultmap(alloc: pointer)986// OMP5-NEXT: #pragma omp target defaultmap(alloc: pointer)987 foo();988// OMP5-NEXT: foo();989 990#pragma omp target defaultmap(to: pointer)991// OMP5-NEXT: #pragma omp target defaultmap(to: pointer)992 foo();993// OMP5-NEXT: foo();994 995#pragma omp target defaultmap(from: pointer)996// OMP5-NEXT: #pragma omp target defaultmap(from: pointer)997 foo();998// OMP5-NEXT: foo();999 1000#pragma omp target defaultmap(tofrom: pointer)1001// OMP5-NEXT: #pragma omp target defaultmap(tofrom: pointer)1002 foo();1003// OMP5-NEXT: foo();1004 1005#pragma omp target defaultmap(firstprivate: pointer)1006// OMP5-NEXT: #pragma omp target defaultmap(firstprivate: pointer)1007 foo();1008// OMP5-NEXT: foo();1009 1010#pragma omp target defaultmap(none: pointer)1011// OMP5-NEXT: #pragma omp target defaultmap(none: pointer)1012 foo();1013// OMP5-NEXT: foo();1014 1015#pragma omp target defaultmap(default: pointer)1016// OMP5-NEXT: #pragma omp target defaultmap(default: pointer)1017 foo();1018// OMP5-NEXT: foo();1019 1020#pragma omp target defaultmap(to: scalar) defaultmap(tofrom: pointer)1021// OMP5-NEXT: #pragma omp target defaultmap(to: scalar) defaultmap(tofrom: pointer)1022 foo();1023// OMP5-NEXT: foo();1024 1025#pragma omp target defaultmap(from: pointer) defaultmap(none: aggregate)1026// OMP5-NEXT: #pragma omp target defaultmap(from: pointer) defaultmap(none: aggregate)1027 foo();1028// OMP5-NEXT: foo();1029 1030#pragma omp target defaultmap(default: aggregate) defaultmap(alloc: scalar)1031// OMP5-NEXT: #pragma omp target defaultmap(default: aggregate) defaultmap(alloc: scalar)1032 foo();1033// OMP5-NEXT: foo();1034 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)1037 foo();1038// OMP5-NEXT: foo();1039 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)1042 foo();1043// OMP5-NEXT: foo();1044 1045 int *g;1046// OMP5-NEXT: int *g;1047 1048#pragma omp target is_device_ptr(g) defaultmap(none: pointer)1049// OMP5-NEXT: #pragma omp target is_device_ptr(g) defaultmap(none: pointer)1050 g++;1051// OMP5-NEXT: g++;1052 1053#pragma omp target private(g) defaultmap(none: pointer)1054// OMP5-NEXT: #pragma omp target private(g) defaultmap(none: pointer)1055 g++;1056// OMP5-NEXT: g++;1057 1058#pragma omp target firstprivate(g) defaultmap(none: pointer)1059// OMP5-NEXT: #pragma omp target firstprivate(g) defaultmap(none: pointer)1060 g++;1061// OMP5-NEXT: g++;1062 1063#pragma omp target defaultmap(none: scalar) map(to: i)1064// OMP5-NEXT: #pragma omp target defaultmap(none: scalar) map(to: i)1065 i++;1066// OMP5-NEXT: i++;1067 1068#pragma omp target defaultmap(none: aggregate) map(to: a)1069// OMP5-NEXT: #pragma omp target defaultmap(none: aggregate) map(to: a)1070 a[3]++;1071// OMP5-NEXT: a[3]++;1072 1073#pragma omp target defaultmap(none: scalar)1074// OMP5-NEXT: #pragma omp target defaultmap(none: scalar)1075 bar();1076// OMP5-NEXT: bar();1077#pragma omp target defaultmap(none)1078 // OMP5-NEXT: #pragma omp target defaultmap(none)1079 // OMP5-NEXT: bar();1080 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();1084 bar();1085 return tmain<int, 5>(argc, &argc) + tmain<char, 1>(argv[0][0], argv[0]);1086}1087 1088#endif //OMP51089 1090#ifdef OMP511091 1092///==========================================================================///1093// RUN: %clang_cc1 -DOMP51 -verify -Wno-vla -fopenmp -ast-print %s | FileCheck %s --check-prefix OMP511094// RUN: %clang_cc1 -DOMP51 -fopenmp -x c++ -std=c++11 -emit-pch -o %t %s1095// RUN: %clang_cc1 -DOMP51 -fopenmp -std=c++11 -include-pch %t -verify -Wno-vla %s -ast-print | FileCheck %s --check-prefix OMP511096 1097// RUN: %clang_cc1 -DOMP51 -verify -Wno-vla -fopenmp-simd -ast-print %s | FileCheck %s --check-prefix OMP511098// RUN: %clang_cc1 -DOMP51 -fopenmp-simd -x c++ -std=c++11 -emit-pch -o %t %s1099// RUN: %clang_cc1 -DOMP51 -fopenmp-simd -std=c++11 -include-pch %t -verify -Wno-vla %s -ast-print | FileCheck %s --check-prefix OMP511100 1101void foo() {}1102 1103template <typename T, int C>1104T tmain(T argc, T *argv) {1105 #pragma omp target defaultmap(present: scalar)1106 foo();1107 #pragma omp target defaultmap(present: aggregate)1108 foo();1109 #pragma omp target defaultmap(present: pointer)1110 foo();1111 #pragma omp target thread_limit(C)1112 foo();1113 #pragma omp target defaultmap(present)1114 foo();1115 1116 return 0;1117}1118 1119// OMP51: template <typename T, int C> T tmain(T argc, T *argv) {1120// OMP51-NEXT: #pragma omp target defaultmap(present: scalar)1121// OMP51-NEXT: foo()1122// OMP51-NEXT: #pragma omp target defaultmap(present: aggregate)1123// OMP51-NEXT: foo()1124// OMP51-NEXT: #pragma omp target defaultmap(present: pointer)1125// OMP51-NEXT: foo()1126// OMP51-NEXT: #pragma omp target thread_limit(C)1127// OMP51-NEXT: foo()1128// OMP51-NEXT: #pragma omp target defaultmap(present)1129// OMP51-NEXT: foo()1130 1131// OMP51-LABEL: int main(int argc, char **argv) {1132int main (int argc, char **argv) {1133#pragma omp target defaultmap(present: scalar)1134// OMP51-NEXT: #pragma omp target defaultmap(present: scalar)1135 foo();1136// OMP51-NEXT: foo();1137#pragma omp target defaultmap(present: aggregate)1138// OMP51-NEXT: #pragma omp target defaultmap(present: aggregate)1139 foo();1140// OMP51-NEXT: foo();1141#pragma omp target defaultmap(present: pointer)1142// OMP51-NEXT: #pragma omp target defaultmap(present: pointer)1143 foo();1144// OMP51-NEXT: foo();1145 1146 return tmain<int, 5>(argc, &argc) + tmain<char, 1>(argv[0][0], argv[0]);1147}1148#endif // OMP511149 1150#ifdef OMP521151 1152///==========================================================================///1153// RUN: %clang_cc1 -DOMP52 -verify -Wno-vla -fopenmp -fopenmp-version=52 -ast-print %s | FileCheck %s --check-prefix OMP521154// RUN: %clang_cc1 -DOMP52 -fopenmp -fopenmp-version=52 -x c++ -std=c++11 -emit-pch -o %t %s1155// RUN: %clang_cc1 -DOMP52 -fopenmp -fopenmp-version=52 -std=c++11 -include-pch %t -verify -Wno-vla %s -ast-print | FileCheck %s --check-prefix OMP521156 1157// RUN: %clang_cc1 -DOMP52 -verify -Wno-vla -fopenmp-simd -fopenmp-version=52 -ast-print %s | FileCheck %s --check-prefix OMP521158// RUN: %clang_cc1 -DOMP52 -fopenmp-simd -fopenmp-version=52 -x c++ -std=c++11 -emit-pch -o %t %s1159// RUN: %clang_cc1 -DOMP52 -fopenmp-simd -fopenmp-version=52 -std=c++11 -include-pch %t -verify -Wno-vla %s -ast-print | FileCheck %s --check-prefix OMP521160 1161void foo() {}1162 1163template <typename T, int C>1164T tmain(T argc, T *argv) {1165 int N = 100;1166 int v[N];1167 #pragma omp target map(iterator(it = 0:N:2), to: v[it])1168 foo();1169 #pragma omp target map(iterator(it = 0:N:4), from: v[it])1170 foo();1171 1172 return 0;1173}1174 1175// OMP52: template <typename T, int C> T tmain(T argc, T *argv) {1176// OMP52-NEXT: int N = 100;1177// OMP52-NEXT: int v[N];1178// OMP52-NEXT: #pragma omp target map(iterator(int it = 0:N:2),to: v[it])1179// OMP52-NEXT: foo()1180// OMP52-NEXT: #pragma omp target map(iterator(int it = 0:N:4),from: v[it])1181// OMP52-NEXT: foo()1182 1183// OMP52-LABEL: int main(int argc, char **argv) {1184int main (int argc, char **argv) {1185 int i, j, a[20], always, close;1186// OMP52-NEXT: int i, j, a[20]1187#pragma omp target1188// OMP52-NEXT: #pragma omp target1189 foo();1190// OMP52-NEXT: foo();1191#pragma omp target map(iterator(it = 0:20:2), to: a[it])1192// OMP52-NEXT: #pragma omp target map(iterator(int it = 0:20:2),to: a[it])1193 foo();1194// OMP52-NEXT: foo();1195#pragma omp target map(iterator(it = 0:20:4), from: a[it])1196// OMP52-NEXT: #pragma omp target map(iterator(int it = 0:20:4),from: a[it])1197foo();1198// OMP52-NEXT: foo();1199 1200 return tmain<int, 5>(argc, &argc) + tmain<char, 1>(argv[0][0], argv[0]);1201}1202#endif // OMP521203 1204#ifdef OMP601205 1206///==========================================================================///1207// RUN: %clang_cc1 -DOMP60 -verify -Wno-vla -fopenmp -fopenmp-version=60 -ast-print %s | FileCheck %s --check-prefix OMP601208// RUN: %clang_cc1 -DOMP60 -fopenmp -fopenmp-version=60 -x c++ -std=c++11 -emit-pch -o %t %s1209// RUN: %clang_cc1 -DOMP60 -fopenmp -fopenmp-version=60 -std=c++11 -include-pch %t -verify -Wno-vla %s -ast-print | FileCheck %s --check-prefix OMP601210 1211// RUN: %clang_cc1 -DOMP60 -verify -Wno-vla -fopenmp-simd -fopenmp-version=60 -ast-print %s | FileCheck %s --check-prefix OMP601212// RUN: %clang_cc1 -DOMP60 -fopenmp-simd -fopenmp-version=60 -x c++ -std=c++11 -emit-pch -o %t %s1213// RUN: %clang_cc1 -DOMP60 -fopenmp-simd -fopenmp-version=60 -std=c++11 -include-pch %t -verify -Wno-vla %s -ast-print | FileCheck %s --check-prefix OMP601214 1215void foo() {}1216template <typename T, int C>1217T tmain(T argc, T *argv) {1218 T i;1219#pragma omp target map(from always: i)1220 foo();1221#pragma omp target map(from, close: i)1222 foo();1223#pragma omp target map(always,close: i)1224 foo();1225 return 0;1226}1227//OMP60: template <typename T, int C> T tmain(T argc, T *argv) {1228//OMP60-NEXT: T i;1229//OMP60-NEXT: #pragma omp target map(always,from: i)1230//OMP60-NEXT: foo();1231//OMP60-NEXT: #pragma omp target map(close,from: i)1232//OMP60-NEXT: foo();1233//OMP60-NEXT: #pragma omp target map(always,close,tofrom: i)1234//OMP60-NEXT: foo();1235//OMP60-NEXT: return 0;1236//OMP60-NEXT:}1237//OMP60: template<> int tmain<int, 5>(int argc, int *argv) {1238//OMP60-NEXT: int i;1239//OMP60-NEXT: #pragma omp target map(always,from: i)1240//OMP60-NEXT: foo();1241//OMP60-NEXT: #pragma omp target map(close,from: i)1242//OMP60-NEXT: foo();1243//OMP60-NEXT: #pragma omp target map(always,close,tofrom: i)1244//OMP60-NEXT: foo();1245//OMP60-NEXT: return 0;1246//OMP60-NEXT:}1247//OMP60: template<> char tmain<char, 1>(char argc, char *argv) {1248//OMP60-NEXT: char i;1249//OMP60-NEXT: #pragma omp target map(always,from: i)1250//OMP60-NEXT: foo();1251//OMP60-NEXT: #pragma omp target map(close,from: i)1252//OMP60-NEXT: foo();1253//OMP60-NEXT: #pragma omp target map(always,close,tofrom: i)1254//OMP60-NEXT: foo();1255//OMP60-NEXT: return 0;1256//OMP60-NEXT:}1257int main (int argc, char **argv) {1258 return tmain<int, 5>(argc, &argc) + tmain<char, 1>(argv[0][0], argv[0]);1259}1260#endif // OMP601261 1262#ifdef OMPX1263 1264// RUN: %clang_cc1 -DOMPX -verify -Wno-vla -fopenmp -fopenmp-extensions -ast-print %s | FileCheck %s --check-prefix=OMPX1265// RUN: %clang_cc1 -DOMPX -fopenmp -fopenmp-extensions -x c++ -std=c++11 -emit-pch -o %t %s1266// RUN: %clang_cc1 -DOMPX -fopenmp -fopenmp-extensions -std=c++11 -include-pch %t -verify -Wno-vla %s -ast-print | FileCheck %s --check-prefix=OMPX1267 1268// RUN: %clang_cc1 -DOMPX -verify -Wno-vla -fopenmp-simd -fopenmp-extensions -ast-print %s | FileCheck %s --check-prefix=OMPX1269// RUN: %clang_cc1 -DOMPX -fopenmp-simd -fopenmp-extensions -x c++ -std=c++11 -emit-pch -o %t %s1270// RUN: %clang_cc1 -DOMPX -fopenmp-simd -fopenmp-extensions -std=c++11 -include-pch %t -verify -Wno-vla %s -ast-print | FileCheck %s --check-prefix=OMPX1271 1272void foo() {}1273 1274template <typename T, int C>1275T tmain(T argc, T *argv) {1276 T i, ompx_hold;1277#pragma omp target map(ompx_hold,alloc: i)1278 foo();1279#pragma omp target map(ompx_hold from: i)1280 foo();1281#pragma omp target map(ompx_hold)1282 {ompx_hold++;}1283#pragma omp target map(ompx_hold,i)1284 {ompx_hold++;i++;}1285 return 0;1286}1287 1288// OMPX: template <typename T, int C> T tmain(T argc, T *argv) {1289// OMPX-NEXT: T i, ompx_hold;1290// OMPX-NEXT: #pragma omp target map(ompx_hold,alloc: i)1291// OMPX-NEXT: foo()1292// OMPX-NEXT: #pragma omp target map(ompx_hold,from: i)1293// OMPX-NEXT: foo()1294// OMPX-NEXT: #pragma omp target map(tofrom: ompx_hold)1295// OMPX-NEXT: {1296// OMPX-NEXT: ompx_hold++;1297// OMPX-NEXT: }1298// OMPX-NEXT: #pragma omp target map(tofrom: ompx_hold,i)1299// OMPX-NEXT: {1300// OMPX-NEXT: ompx_hold++;1301// OMPX-NEXT: i++;1302// OMPX-NEXT: }1303 1304// OMPX-LABEL: int main(int argc, char **argv) {1305// OMPX-NEXT: int i, ompx_hold;1306// OMPX-NEXT: #pragma omp target map(ompx_hold,alloc: i)1307// OMPX-NEXT: foo();1308// OMPX-NEXT: #pragma omp target map(ompx_hold,from: i)1309// OMPX-NEXT: foo();1310// OMPX-NEXT: #pragma omp target map(tofrom: ompx_hold)1311// OMPX-NEXT: {1312// OMPX-NEXT: ompx_hold++;1313// OMPX-NEXT: }1314// OMPX-NEXT: #pragma omp target map(tofrom: ompx_hold,i)1315// OMPX-NEXT: {1316// OMPX-NEXT: ompx_hold++;1317// OMPX-NEXT: i++;1318// OMPX-NEXT: }1319int main (int argc, char **argv) {1320 int i, ompx_hold;1321 #pragma omp target map(ompx_hold,alloc: i)1322 foo();1323 #pragma omp target map(ompx_hold from: i)1324 foo();1325 #pragma omp target map(ompx_hold)1326 {ompx_hold++;}1327 #pragma omp target map(ompx_hold,i)1328 {ompx_hold++;i++;}1329 return tmain<int, 5>(argc, &argc) + tmain<char, 1>(argv[0][0], argv[0]);1330}1331 1332#endif1333#endif1334