408 lines · cpp
1// RUN: %clang_cc1 -fopenacc -Wno-openacc-deprecated-clause-alias -ast-print %s -o - | FileCheck %s2 3constexpr int get_value() { return 1; }4void foo() {5 int *iPtr;6// CHECK: #pragma acc parallel loop7// CHECK-NEXT: for (int i = 0; i < 5; ++i)8// CHECK-NEXT: ;9#pragma acc parallel loop10 for(int i = 0;i<5;++i);11// CHECK: #pragma acc serial loop12// CHECK-NEXT: for (int i = 0; i < 5; ++i)13// CHECK-NEXT: ;14#pragma acc serial loop15 for(int i = 0;i<5;++i);16// CHECK: #pragma acc kernels loop17// CHECK-NEXT: for (int i = 0; i < 5; ++i)18// CHECK-NEXT: ;19#pragma acc kernels loop20 for(int i = 0;i<5;++i);21// CHECK: #pragma acc parallel loop auto22// CHECK-NEXT: for (int i = 0; i < 5; ++i)23// CHECK-NEXT: ;24#pragma acc parallel loop auto25 for(int i = 0;i<5;++i);26// CHECK: #pragma acc serial loop seq27// CHECK-NEXT: for (int i = 0; i < 5; ++i)28// CHECK-NEXT: ;29#pragma acc serial loop seq30 for(int i = 0;i<5;++i);31// CHECK: #pragma acc kernels loop independent32// CHECK-NEXT: for (int i = 0; i < 5; ++i)33// CHECK-NEXT: ;34#pragma acc kernels loop independent35 for(int i = 0;i<5;++i);36 37 bool SomeB;38 struct SomeStruct{} SomeStructImpl;39 40//CHECK: #pragma acc parallel loop dtype(default)41// CHECK-NEXT: for (int i = 0; i < 5; ++i)42// CHECK-NEXT: ;43#pragma acc parallel loop dtype(default)44 for(int i = 0;i<5;++i);45 46//CHECK: #pragma acc serial loop device_type(radeon, host)47// CHECK-NEXT: for (int i = 0; i < 5; ++i)48// CHECK-NEXT: ;49#pragma acc serial loop device_type(radeon, host)50 for(int i = 0;i<5;++i);51 52//CHECK: #pragma acc kernels loop device_type(nvidia)53// CHECK-NEXT: for (int i = 0; i < 5; ++i)54// CHECK-NEXT: ;55#pragma acc kernels loop device_type(nvidia)56 for(int i = 0;i<5;++i);57 58//CHECK: #pragma acc parallel loop dtype(multicore)59// CHECK-NEXT: for (int i = 0; i < 5; ++i)60// CHECK-NEXT: ;61#pragma acc parallel loop dtype(multicore)62 for(int i = 0;i<5;++i);63 64//CHECK: #pragma acc serial loop device_type(default)65// CHECK-NEXT: for (int i = 0; i < 5; ++i)66// CHECK-NEXT: ;67#pragma acc serial loop device_type (default)68 for(int i = 0;i<5;++i);69 70// CHECK: #pragma acc kernels loop dtype(host)71// CHECK-NEXT: for (int i = 0; i < 5; ++i)72// CHECK-NEXT: ;73#pragma acc kernels loop dtype(host)74 for(int i = 0;i<5;++i);75 76 int i;77 float array[5];78 79// CHECK: #pragma acc parallel loop self(i == 3)80// CHECK-NEXT: for (int i = 0; i < 5; ++i)81// CHECK-NEXT: ;82#pragma acc parallel loop self(i == 3)83 for(int i = 0;i<5;++i);84 85// CHECK: #pragma acc kernels loop if(i == array[1])86// CHECK-NEXT: for (int i = 0; i < 5; ++i)87// CHECK-NEXT: ;88#pragma acc kernels loop if(i == array[1])89 for(int i = 0;i<5;++i);90 91// CHECK: #pragma acc parallel loop default(none)92// CHECK-NEXT: for (int i = 0; i < 5; ++i)93// CHECK-NEXT: ;94#pragma acc parallel loop default(none)95 for(int i = 0;i<5;++i);96// CHECK: #pragma acc serial loop default(present)97// CHECK-NEXT: for (int i = 0; i < 5; ++i)98// CHECK-NEXT: ;99#pragma acc serial loop default(present)100 for(int i = 0;i<5;++i);101 102// CHECK: #pragma acc parallel loop private(i, array[1], array, array[1:2])103// CHECK-NEXT: for (int i = 0; i < 5; ++i)104// CHECK-NEXT: ;105#pragma acc parallel loop private(i, array[1], array, array[1:2])106 for(int i = 0;i<5;++i);107 108// CHECK: #pragma acc serial loop firstprivate(i, array[1], array, array[1:2])109// CHECK-NEXT: for (int i = 0; i < 5; ++i)110// CHECK-NEXT: ;111#pragma acc serial loop firstprivate(i, array[1], array, array[1:2])112 for(int i = 0;i<5;++i);113 114 // CHECK: #pragma acc kernels loop async(*iPtr)115// CHECK-NEXT: for (int i = 0; i < 5; ++i)116// CHECK-NEXT: ;117#pragma acc kernels loop async(*iPtr)118 for(int i = 0;i<5;++i);119 120 // CHECK: #pragma acc kernels loop async121// CHECK-NEXT: for (int i = 0; i < 5; ++i)122// CHECK-NEXT: ;123#pragma acc kernels loop async124 for(int i = 0;i<5;++i);125 126// CHECK: #pragma acc parallel loop present(i, array[1], array, array[1:2])127#pragma acc parallel loop present(i, array[1], array, array[1:2])128 for(int i = 0;i<5;++i);129// CHECK: #pragma acc serial loop present(i, array[1], array, array[1:2])130#pragma acc serial loop present(i, array[1], array, array[1:2])131 for(int i = 0;i<5;++i);132// CHECK: #pragma acc kernels loop present(i, array[1], array, array[1:2])133#pragma acc kernels loop present(i, array[1], array, array[1:2])134 for(int i = 0;i<5;++i);135 136 float *arrayPtr[5];137 138 // CHECK: #pragma acc kernels loop deviceptr(iPtr, arrayPtr[0])139#pragma acc kernels loop deviceptr(iPtr, arrayPtr[0])140 for(int i = 0;i<5;++i);141 142// CHECK: #pragma acc parallel loop wait143#pragma acc parallel loop wait144 for(int i = 0;i<5;++i);145 146// CHECK: #pragma acc parallel loop wait(*iPtr, i)147#pragma acc parallel loop wait(*iPtr, i)148 for(int i = 0;i<5;++i);149 150// CHECK: #pragma acc parallel loop wait(queues: *iPtr, i)151#pragma acc parallel loop wait(queues:*iPtr, i)152 for(int i = 0;i<5;++i);153 154// CHECK: #pragma acc parallel loop wait(devnum: i : *iPtr, i)155#pragma acc parallel loop wait(devnum:i:*iPtr, i)156 for(int i = 0;i<5;++i);157 158// CHECK: #pragma acc parallel loop wait(devnum: i : queues: *iPtr, i)159#pragma acc parallel loop wait(devnum:i:queues:*iPtr, i)160 for(int i = 0;i<5;++i);161 162 // CHECK: #pragma acc serial loop attach(iPtr, arrayPtr[0])163#pragma acc serial loop attach(iPtr, arrayPtr[0])164 for(int i = 0;i<5;++i);165 166// CHECK: #pragma acc parallel loop no_create(i, array[1], array, array[1:2])167#pragma acc parallel loop no_create(i, array[1], array, array[1:2])168 for(int i = 0;i<5;++i);169 170// CHECK: #pragma acc parallel loop no_create(i, array[1], array, array[1:2]) present(i, array[1], array, array[1:2])171#pragma acc parallel loop no_create(i, array[1], array, array[1:2]) present(i, array[1], array, array[1:2])172 for(int i = 0;i<5;++i);173 174// CHECK: #pragma acc parallel loop copy(alwaysin: i, array[1], array, array[1:2]) pcopy(i, array[1], array, array[1:2]) present_or_copy(i, array[1], array, array[1:2])175#pragma acc parallel loop copy(alwaysin: i, array[1], array, array[1:2]) pcopy(i, array[1], array, array[1:2]) present_or_copy(i, array[1], array, array[1:2])176 for(int i = 0;i<5;++i);177 178// CHECK: #pragma acc parallel loop copyin(i, array[1], array, array[1:2]) pcopyin(readonly: i, array[1], array, array[1:2]) present_or_copyin(always, alwaysin: i, array[1], array, array[1:2])179#pragma acc parallel loop copyin(i, array[1], array, array[1:2]) pcopyin(readonly:i, array[1], array, array[1:2]) present_or_copyin(always, alwaysin: i, array[1], array, array[1:2])180 for(int i = 0;i<5;++i);181 182// CHECK: #pragma acc parallel loop copyout(i, array[1], array, array[1:2]) pcopyout(zero: i, array[1], array, array[1:2]) present_or_copyout(always, alwaysout: i, array[1], array, array[1:2])183#pragma acc parallel loop copyout(i, array[1], array, array[1:2]) pcopyout(zero: i, array[1], array, array[1:2]) present_or_copyout(always, alwaysout: i, array[1], array, array[1:2])184 for(int i = 0;i<5;++i);185 186// CHECK: #pragma acc parallel loop create(i, array[1], array, array[1:2]) pcreate(zero: i, array[1], array, array[1:2]) present_or_create(i, array[1], array, array[1:2])187#pragma acc parallel loop create(i, array[1], array, array[1:2]) pcreate(zero: i, array[1], array, array[1:2]) present_or_create(i, array[1], array, array[1:2])188 for(int i = 0;i<5;++i);189 190// CHECK: #pragma acc parallel loop collapse(1)191// CHECK-NEXT: for (int i = 0; i < 5; ++i)192// CHECK-NEXT: ;193#pragma acc parallel loop collapse(1)194 for(int i = 0;i<5;++i);195// CHECK: #pragma acc serial loop collapse(force:1)196// CHECK-NEXT: for (int i = 0; i < 5; ++i)197// CHECK-NEXT: ;198#pragma acc serial loop collapse(force:1)199 for(int i = 0;i<5;++i);200// CHECK: #pragma acc kernels loop collapse(2)201// CHECK-NEXT: for (int i = 0; i < 5; ++i)202// CHECK-NEXT: for (int i = 0; i < 5; ++i)203// CHECK-NEXT: ;204#pragma acc kernels loop collapse(2)205 for(int i = 0;i<5;++i)206 for(int i = 0;i<5;++i);207// CHECK: #pragma acc parallel loop collapse(force:2)208// CHECK-NEXT: for (int i = 0; i < 5; ++i)209// CHECK-NEXT: for (int i = 0; i < 5; ++i)210// CHECK-NEXT: ;211#pragma acc parallel loop collapse(force:2)212 for(int i = 0;i<5;++i)213 for(int i = 0;i<5;++i);214 215// CHECK: #pragma acc serial loop tile(1, 3, *, get_value())216// CHECK-NEXT: for (int i = 0; i < 5; ++i)217// CHECK-NEXT: for (int i = 0; i < 5; ++i)218// CHECK-NEXT: for (int i = 0; i < 5; ++i)219// CHECK-NEXT: for (int i = 0; i < 5; ++i)220// CHECK-NEXT: ;221#pragma acc serial loop tile(1, 3, *, get_value())222 for(int i = 0;i<5;++i)223 for(int i = 0;i<5;++i)224 for(int i = 0;i<5;++i)225 for(int i = 0;i<5;++i);226 227// CHECK: #pragma acc parallel loop num_gangs(i, (int)array[2])228// CHECK-NEXT: for (int i = 0; i < 5; ++i)229// CHECK-NEXT: ;230#pragma acc parallel loop num_gangs(i, (int)array[2])231 for(int i = 0;i<5;++i);232 233// CHECK: #pragma acc parallel loop num_workers(i)234// CHECK-NEXT: for (int i = 0; i < 5; ++i)235// CHECK-NEXT: ;236#pragma acc parallel loop num_workers(i)237 for(int i = 0;i<5;++i);238 239// CHECK: #pragma acc parallel loop vector_length((int)array[1])240// CHECK-NEXT: for (int i = 0; i < 5; ++i)241// CHECK-NEXT: ;242#pragma acc parallel loop vector_length((int)array[1])243 for(int i = 0;i<5;++i);244 245// CHECK: #pragma acc parallel loop gang(dim: 2)246// CHECK-NEXT: for (int i = 0; i < 5; ++i)247// CHECK-NEXT: ;248#pragma acc parallel loop gang(dim:2)249 for(int i = 0;i<5;++i);250 251// CHECK: #pragma acc serial loop gang(static: i)252// CHECK-NEXT: for (int i = 0; i < 5; ++i)253// CHECK-NEXT: ;254#pragma acc serial loop gang(static:i)255 for(int i = 0;i<5;++i);256 257// CHECK: #pragma acc parallel loop gang(static: i) gang(dim: 2)258// CHECK-NEXT: for (int i = 0; i < 5; ++i)259// CHECK-NEXT: ;260#pragma acc parallel loop gang(static:i) gang(dim:2)261 for(int i = 0;i<5;++i);262 263// CHECK: #pragma acc parallel loop gang(static: i, dim: 2)264// CHECK-NEXT: for (int i = 0; i < 5; ++i)265// CHECK-NEXT: ;266#pragma acc parallel loop gang(static:i, dim:2)267 for(int i = 0;i<5;++i);268 269// CHECK: #pragma acc parallel loop gang(dim: 2)270// CHECK-NEXT: for (int i = 0; i < 5; ++i)271// CHECK-NEXT: ;272#pragma acc parallel loop gang(dim:2)273 for(int i = 0;i<5;++i);274 275// CHECK: #pragma acc parallel loop gang(static: i)276// CHECK-NEXT: for (int i = 0; i < 5; ++i)277// CHECK-NEXT: ;278#pragma acc parallel loop gang(static:i)279 for(int i = 0;i<5;++i);280 281// CHECK: #pragma acc parallel loop gang(static: i) gang(dim: 2)282// CHECK-NEXT: for (int i = 0; i < 5; ++i)283// CHECK-NEXT: ;284#pragma acc parallel loop gang(static:i) gang(dim:2)285 for(int i = 0;i<5;++i);286 287// CHECK: #pragma acc parallel loop gang(static: i, dim: 2)288// CHECK-NEXT: for (int i = 0; i < 5; ++i)289// CHECK-NEXT: ;290#pragma acc parallel loop gang(static:i, dim:2)291 for(int i = 0;i<5;++i);292 293// CHECK: #pragma acc kernels loop gang(num: i) gang(static: i)294// CHECK-NEXT: for (int i = 0; i < 5; ++i)295// CHECK-NEXT: ;296#pragma acc kernels loop gang(i) gang(static:i)297 for(int i = 0;i<5;++i);298 299// CHECK: #pragma acc kernels loop gang(num: i) gang(static: i)300// CHECK-NEXT: for (int i = 0; i < 5; ++i)301// CHECK-NEXT: ;302#pragma acc kernels loop gang(num:i) gang(static:i)303 for(int i = 0;i<5;++i);304 305// CHECK: #pragma acc serial loop gang(static: i)306// CHECK-NEXT: for (int i = 0; i < 5; ++i)307// CHECK-NEXT: ;308#pragma acc serial loop gang(static:i)309 for(int i = 0;i<5;++i);310 311// CHECK: #pragma acc serial loop gang(static: *)312// CHECK-NEXT: for (int i = 0; i < 5; ++i)313// CHECK-NEXT: ;314#pragma acc serial loop gang(static:*)315 for(int i = 0;i<5;++i);316 317// CHECK: #pragma acc parallel loop worker318// CHECK-NEXT: for (int i = 0; i < 5; ++i)319// CHECK-NEXT: ;320#pragma acc parallel loop worker321 for(int i = 0;i<5;++i);322 323// CHECK-NEXT: #pragma acc parallel loop worker324// CHECK-NEXT: for (int i = 0; i < 5; ++i)325// CHECK-NEXT: ;326#pragma acc parallel loop worker327 for(int i = 0;i<5;++i);328 329// CHECK-NEXT: #pragma acc serial loop worker330// CHECK-NEXT: for (int i = 0; i < 5; ++i)331// CHECK-NEXT: ;332#pragma acc serial loop worker333 for(int i = 0;i<5;++i);334 335// CHECK-NEXT: #pragma acc kernels loop worker(num: 5)336// CHECK-NEXT: for (int i = 0; i < 5; ++i)337// CHECK-NEXT: ;338#pragma acc kernels loop worker(5)339 for(int i = 0;i<5;++i);340 341// CHECK-NEXT: #pragma acc kernels loop worker(num: 5)342// CHECK-NEXT: for (int i = 0; i < 5; ++i)343// CHECK-NEXT: ;344#pragma acc kernels loop worker(num:5)345 for(int i = 0;i<5;++i);346 347 // CHECK: #pragma acc parallel loop vector348// CHECK-NEXT: for (int i = 0; i < 5; ++i)349// CHECK-NEXT: ;350#pragma acc parallel loop vector351 for(int i = 0;i<5;++i);352 353// CHECK: #pragma acc parallel loop vector(length: 5)354// CHECK-NEXT: for (int i = 0; i < 5; ++i)355// CHECK-NEXT: ;356#pragma acc parallel loop vector(5)357 for(int i = 0;i<5;++i);358 359// CHECK: #pragma acc parallel loop vector(length: 5)360// CHECK-NEXT: for (int i = 0; i < 5; ++i)361// CHECK-NEXT: ;362#pragma acc parallel loop vector(length:5)363 for(int i = 0;i<5;++i);364 365// CHECK-NEXT: #pragma acc kernels loop vector366// CHECK-NEXT: for (int i = 0; i < 5; ++i)367// CHECK-NEXT: ;368#pragma acc kernels loop vector369 for(int i = 0;i<5;++i);370 371// CHECK-NEXT: #pragma acc kernels loop vector(length: 5)372// CHECK-NEXT: for (int i = 0; i < 5; ++i)373// CHECK-NEXT: ;374#pragma acc kernels loop vector(5)375 for(int i = 0;i<5;++i);376 377// CHECK-NEXT: #pragma acc kernels loop vector(length: 5)378// CHECK-NEXT: for (int i = 0; i < 5; ++i)379// CHECK-NEXT: ;380#pragma acc kernels loop vector(length:5)381 for(int i = 0;i<5;++i);382 383// CHECK-NEXT: #pragma acc serial loop vector384// CHECK-NEXT: for (int i = 0; i < 5; ++i)385// CHECK-NEXT: ;386#pragma acc serial loop vector387 for(int i = 0;i<5;++i);388 389//CHECK: #pragma acc serial loop reduction(*: i)390#pragma acc serial loop reduction(*: i)391 for(int i = 0;i<5;++i);392//CHECK: #pragma acc kernels loop reduction(max: SomeB)393#pragma acc kernels loop reduction(max: SomeB)394 for(int i = 0;i<5;++i);395//CHECK: #pragma acc serial loop reduction(&: i)396#pragma acc serial loop reduction(&: i)397 for(int i = 0;i<5;++i);398//CHECK: #pragma acc kernels loop reduction(|: SomeB)399#pragma acc kernels loop reduction(|: SomeB)400 for(int i = 0;i<5;++i);401//CHECK: #pragma acc serial loop reduction(&&: i)402#pragma acc serial loop reduction(&&: i)403 for(int i = 0;i<5;++i);404//CHECK: #pragma acc kernels loop reduction(||: SomeB)405#pragma acc kernels loop reduction(||: SomeB)406 for(int i = 0;i<5;++i);407}408