1 // RUN: %clang_cc1 -fopenacc -Wno-openacc-deprecated-clause-alias -ast-print %s -o - | FileCheck %s 2 3 constexpr int get_value() { return 1; } 4 void foo() { 5 int *iPtr; 6 // CHECK: #pragma acc parallel loop 7 // CHECK-NEXT: for (int i = 0; i < 5; ++i) 8 // CHECK-NEXT: ; 9 #pragma acc parallel loop 10 for(int i = 0;i<5;++i); 11 // CHECK: #pragma acc serial loop 12 // CHECK-NEXT: for (int i = 0; i < 5; ++i) 13 // CHECK-NEXT: ; 14 #pragma acc serial loop 15 for(int i = 0;i<5;++i); 16 // CHECK: #pragma acc kernels loop 17 // CHECK-NEXT: for (int i = 0; i < 5; ++i) 18 // CHECK-NEXT: ; 19 #pragma acc kernels loop 20 for(int i = 0;i<5;++i); 21 // CHECK: #pragma acc parallel loop auto 22 // CHECK-NEXT: for (int i = 0; i < 5; ++i) 23 // CHECK-NEXT: ; 24 #pragma acc parallel loop auto 25 for(int i = 0;i<5;++i); 26 // CHECK: #pragma acc serial loop seq 27 // CHECK-NEXT: for (int i = 0; i < 5; ++i) 28 // CHECK-NEXT: ; 29 #pragma acc serial loop seq 30 for(int i = 0;i<5;++i); 31 // CHECK: #pragma acc kernels loop independent 32 // CHECK-NEXT: for (int i = 0; i < 5; ++i) 33 // CHECK-NEXT: ; 34 #pragma acc kernels loop independent 35 for(int i = 0;i<5;++i); 36 37 bool SomeB; 38 struct SomeStruct{} SomeStructImpl; 39 40 //CHECK: #pragma acc parallel loop dtype(SomeB) 41 // CHECK-NEXT: for (int i = 0; i < 5; ++i) 42 // CHECK-NEXT: ; 43 #pragma acc parallel loop dtype(SomeB) 44 for(int i = 0;i<5;++i); 45 46 //CHECK: #pragma acc serial loop device_type(SomeStruct) 47 // CHECK-NEXT: for (int i = 0; i < 5; ++i) 48 // CHECK-NEXT: ; 49 #pragma acc serial loop device_type(SomeStruct) 50 for(int i = 0;i<5;++i); 51 52 //CHECK: #pragma acc kernels loop device_type(int) 53 // CHECK-NEXT: for (int i = 0; i < 5; ++i) 54 // CHECK-NEXT: ; 55 #pragma acc kernels loop device_type(int) 56 for(int i = 0;i<5;++i); 57 58 //CHECK: #pragma acc parallel loop dtype(bool) 59 // CHECK-NEXT: for (int i = 0; i < 5; ++i) 60 // CHECK-NEXT: ; 61 #pragma acc parallel loop dtype(bool) 62 for(int i = 0;i<5;++i); 63 64 //CHECK: #pragma acc serial loop device_type(SomeStructImpl) 65 // CHECK-NEXT: for (int i = 0; i < 5; ++i) 66 // CHECK-NEXT: ; 67 #pragma acc serial loop device_type (SomeStructImpl) 68 for(int i = 0;i<5;++i); 69 70 // CHECK: #pragma acc kernels loop dtype(AnotherIdent) 71 // CHECK-NEXT: for (int i = 0; i < 5; ++i) 72 // CHECK-NEXT: ; 73 #pragma acc kernels loop dtype(AnotherIdent) 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 async 121 // CHECK-NEXT: for (int i = 0; i < 5; ++i) 122 // CHECK-NEXT: ; 123 #pragma acc kernels loop async 124 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 wait() 143 #pragma acc parallel loop wait() 144 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(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(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(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(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(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(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 worker 318 // CHECK-NEXT: for (int i = 0; i < 5; ++i) 319 // CHECK-NEXT: ; 320 #pragma acc parallel loop worker 321 for(int i = 0;i<5;++i); 322 323 // CHECK-NEXT: #pragma acc parallel loop worker 324 // CHECK-NEXT: for (int i = 0; i < 5; ++i) 325 // CHECK-NEXT: ; 326 #pragma acc parallel loop worker 327 for(int i = 0;i<5;++i); 328 329 // CHECK-NEXT: #pragma acc serial loop worker 330 // CHECK-NEXT: for (int i = 0; i < 5; ++i) 331 // CHECK-NEXT: ; 332 #pragma acc serial loop worker 333 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 vector 348 // CHECK-NEXT: for (int i = 0; i < 5; ++i) 349 // CHECK-NEXT: ; 350 #pragma acc parallel loop vector 351 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 vector 366 // CHECK-NEXT: for (int i = 0; i < 5; ++i) 367 // CHECK-NEXT: ; 368 #pragma acc kernels loop vector 369 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 vector 384 // CHECK-NEXT: for (int i = 0; i < 5; ++i) 385 // CHECK-NEXT: ; 386 #pragma acc serial loop vector 387 for(int i = 0;i<5;++i); 388 389 //CHECK: #pragma acc parallel loop reduction(+: iPtr) 390 #pragma acc parallel loop reduction(+: iPtr) 391 for(int i = 0;i<5;++i); 392 //CHECK: #pragma acc serial loop reduction(*: i) 393 #pragma acc serial loop reduction(*: i) 394 for(int i = 0;i<5;++i); 395 //CHECK: #pragma acc kernels loop reduction(max: SomeB) 396 #pragma acc kernels loop reduction(max: SomeB) 397 for(int i = 0;i<5;++i); 398 //CHECK: #pragma acc parallel loop reduction(min: iPtr) 399 #pragma acc parallel loop reduction(min: iPtr) 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 //CHECK: #pragma acc parallel loop reduction(^: iPtr) 408 #pragma acc parallel loop reduction(^: iPtr) 409 for(int i = 0;i<5;++i); 410 //CHECK: #pragma acc serial loop reduction(&&: i) 411 #pragma acc serial loop reduction(&&: i) 412 for(int i = 0;i<5;++i); 413 //CHECK: #pragma acc kernels loop reduction(||: SomeB) 414 #pragma acc kernels loop reduction(||: SomeB) 415 for(int i = 0;i<5;++i); 416 } 417