'num_gangs', 'num_workers', and 'vector_length' are similar to eachother, and are all the same implementation as for compute constructs, so this patch just enables them and adds the necessary testing to ensure they work correctly. These will get more complicated when they get combined with 'gang', 'worker', 'vector' and 'reduction', but those restrictions will be implemented when those clauses are enabled.
246 lines
8.7 KiB
C++
246 lines
8.7 KiB
C++
// RUN: %clang_cc1 -fopenacc -Wno-openacc-deprecated-clause-alias -ast-print %s -o - | FileCheck %s
|
|
|
|
constexpr int get_value() { return 1; }
|
|
void foo() {
|
|
int *iPtr;
|
|
// CHECK: #pragma acc parallel loop
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: ;
|
|
#pragma acc parallel loop
|
|
for(int i = 0;i<5;++i);
|
|
// CHECK: #pragma acc serial loop
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: ;
|
|
#pragma acc serial loop
|
|
for(int i = 0;i<5;++i);
|
|
// CHECK: #pragma acc kernels loop
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: ;
|
|
#pragma acc kernels loop
|
|
for(int i = 0;i<5;++i);
|
|
// CHECK: #pragma acc parallel loop auto
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: ;
|
|
#pragma acc parallel loop auto
|
|
for(int i = 0;i<5;++i);
|
|
// CHECK: #pragma acc serial loop seq
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: ;
|
|
#pragma acc serial loop seq
|
|
for(int i = 0;i<5;++i);
|
|
// CHECK: #pragma acc kernels loop independent
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: ;
|
|
#pragma acc kernels loop independent
|
|
for(int i = 0;i<5;++i);
|
|
|
|
bool SomeB;
|
|
struct SomeStruct{} SomeStructImpl;
|
|
|
|
//CHECK: #pragma acc parallel loop dtype(SomeB)
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: ;
|
|
#pragma acc parallel loop dtype(SomeB)
|
|
for(int i = 0;i<5;++i);
|
|
|
|
//CHECK: #pragma acc serial loop device_type(SomeStruct)
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: ;
|
|
#pragma acc serial loop device_type(SomeStruct)
|
|
for(int i = 0;i<5;++i);
|
|
|
|
//CHECK: #pragma acc kernels loop device_type(int)
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: ;
|
|
#pragma acc kernels loop device_type(int)
|
|
for(int i = 0;i<5;++i);
|
|
|
|
//CHECK: #pragma acc parallel loop dtype(bool)
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: ;
|
|
#pragma acc parallel loop dtype(bool)
|
|
for(int i = 0;i<5;++i);
|
|
|
|
//CHECK: #pragma acc serial loop device_type(SomeStructImpl)
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: ;
|
|
#pragma acc serial loop device_type (SomeStructImpl)
|
|
for(int i = 0;i<5;++i);
|
|
|
|
// CHECK: #pragma acc kernels loop dtype(AnotherIdent)
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: ;
|
|
#pragma acc kernels loop dtype(AnotherIdent)
|
|
for(int i = 0;i<5;++i);
|
|
|
|
int i;
|
|
float array[5];
|
|
|
|
// CHECK: #pragma acc parallel loop self(i == 3)
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: ;
|
|
#pragma acc parallel loop self(i == 3)
|
|
for(int i = 0;i<5;++i);
|
|
|
|
// CHECK: #pragma acc kernels loop if(i == array[1])
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: ;
|
|
#pragma acc kernels loop if(i == array[1])
|
|
for(int i = 0;i<5;++i);
|
|
|
|
// CHECK: #pragma acc parallel loop default(none)
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: ;
|
|
#pragma acc parallel loop default(none)
|
|
for(int i = 0;i<5;++i);
|
|
// CHECK: #pragma acc serial loop default(present)
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: ;
|
|
#pragma acc serial loop default(present)
|
|
for(int i = 0;i<5;++i);
|
|
|
|
// CHECK: #pragma acc parallel loop private(i, array[1], array, array[1:2])
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: ;
|
|
#pragma acc parallel loop private(i, array[1], array, array[1:2])
|
|
for(int i = 0;i<5;++i);
|
|
|
|
// CHECK: #pragma acc serial loop firstprivate(i, array[1], array, array[1:2])
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: ;
|
|
#pragma acc serial loop firstprivate(i, array[1], array, array[1:2])
|
|
for(int i = 0;i<5;++i);
|
|
|
|
// CHECK: #pragma acc kernels loop async(*iPtr)
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: ;
|
|
#pragma acc kernels loop async(*iPtr)
|
|
for(int i = 0;i<5;++i);
|
|
|
|
// CHECK: #pragma acc kernels loop async
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: ;
|
|
#pragma acc kernels loop async
|
|
for(int i = 0;i<5;++i);
|
|
|
|
// CHECK: #pragma acc parallel loop present(i, array[1], array, array[1:2])
|
|
#pragma acc parallel loop present(i, array[1], array, array[1:2])
|
|
for(int i = 0;i<5;++i);
|
|
// CHECK: #pragma acc serial loop present(i, array[1], array, array[1:2])
|
|
#pragma acc serial loop present(i, array[1], array, array[1:2])
|
|
for(int i = 0;i<5;++i);
|
|
// CHECK: #pragma acc kernels loop present(i, array[1], array, array[1:2])
|
|
#pragma acc kernels loop present(i, array[1], array, array[1:2])
|
|
for(int i = 0;i<5;++i);
|
|
|
|
float *arrayPtr[5];
|
|
|
|
// CHECK: #pragma acc kernels loop deviceptr(iPtr, arrayPtr[0])
|
|
#pragma acc kernels loop deviceptr(iPtr, arrayPtr[0])
|
|
for(int i = 0;i<5;++i);
|
|
|
|
// CHECK: #pragma acc parallel loop wait()
|
|
#pragma acc parallel loop wait()
|
|
for(int i = 0;i<5;++i);
|
|
|
|
// CHECK: #pragma acc parallel loop wait(*iPtr, i)
|
|
#pragma acc parallel loop wait(*iPtr, i)
|
|
for(int i = 0;i<5;++i);
|
|
|
|
// CHECK: #pragma acc parallel loop wait(queues: *iPtr, i)
|
|
#pragma acc parallel loop wait(queues:*iPtr, i)
|
|
for(int i = 0;i<5;++i);
|
|
|
|
// CHECK: #pragma acc parallel loop wait(devnum: i : *iPtr, i)
|
|
#pragma acc parallel loop wait(devnum:i:*iPtr, i)
|
|
for(int i = 0;i<5;++i);
|
|
|
|
// CHECK: #pragma acc parallel loop wait(devnum: i : queues: *iPtr, i)
|
|
#pragma acc parallel loop wait(devnum:i:queues:*iPtr, i)
|
|
for(int i = 0;i<5;++i);
|
|
|
|
// CHECK: #pragma acc serial loop attach(iPtr, arrayPtr[0])
|
|
#pragma acc serial loop attach(iPtr, arrayPtr[0])
|
|
for(int i = 0;i<5;++i);
|
|
|
|
// CHECK: #pragma acc parallel loop no_create(i, array[1], array, array[1:2])
|
|
#pragma acc parallel loop no_create(i, array[1], array, array[1:2])
|
|
for(int i = 0;i<5;++i);
|
|
|
|
// CHECK: #pragma acc parallel loop no_create(i, array[1], array, array[1:2]) present(i, array[1], array, array[1:2])
|
|
#pragma acc parallel loop no_create(i, array[1], array, array[1:2]) present(i, array[1], array, array[1:2])
|
|
for(int i = 0;i<5;++i);
|
|
|
|
// 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])
|
|
#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])
|
|
for(int i = 0;i<5;++i);
|
|
|
|
// 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])
|
|
#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])
|
|
for(int i = 0;i<5;++i);
|
|
|
|
// 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])
|
|
#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])
|
|
for(int i = 0;i<5;++i);
|
|
|
|
// 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])
|
|
#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])
|
|
for(int i = 0;i<5;++i);
|
|
|
|
// CHECK: #pragma acc parallel loop collapse(1)
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: ;
|
|
#pragma acc parallel loop collapse(1)
|
|
for(int i = 0;i<5;++i);
|
|
// CHECK: #pragma acc serial loop collapse(force:1)
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: ;
|
|
#pragma acc serial loop collapse(force:1)
|
|
for(int i = 0;i<5;++i);
|
|
// CHECK: #pragma acc kernels loop collapse(2)
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: ;
|
|
#pragma acc kernels loop collapse(2)
|
|
for(int i = 0;i<5;++i)
|
|
for(int i = 0;i<5;++i);
|
|
// CHECK: #pragma acc parallel loop collapse(force:2)
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: ;
|
|
#pragma acc parallel loop collapse(force:2)
|
|
for(int i = 0;i<5;++i)
|
|
for(int i = 0;i<5;++i);
|
|
|
|
// CHECK: #pragma acc serial loop tile(1, 3, *, get_value())
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: ;
|
|
#pragma acc serial loop tile(1, 3, *, get_value())
|
|
for(int i = 0;i<5;++i)
|
|
for(int i = 0;i<5;++i)
|
|
for(int i = 0;i<5;++i)
|
|
for(int i = 0;i<5;++i);
|
|
|
|
// CHECK: #pragma acc parallel loop num_gangs(i, (int)array[2])
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: ;
|
|
#pragma acc parallel loop num_gangs(i, (int)array[2])
|
|
for(int i = 0;i<5;++i);
|
|
|
|
// CHECK: #pragma acc parallel loop num_workers(i)
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: ;
|
|
#pragma acc parallel loop num_workers(i)
|
|
for(int i = 0;i<5;++i);
|
|
|
|
// CHECK: #pragma acc parallel loop vector_length((int)array[1])
|
|
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
|
|
// CHECK-NEXT: ;
|
|
#pragma acc parallel loop vector_length((int)array[1])
|
|
for(int i = 0;i<5;++i);
|
|
|
|
}
|