// RUN: not %clang_cc1 -fopenacc -triple x86_64-linux-gnu -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir -triple x86_64-linux-pc %s -o - | FileCheck %s template void do_things(unsigned A, unsigned B) { T *OnePtr; #pragma acc parallel private(OnePtr[A:B]) // CHECK: acc.private.recipe @privatization__Bcnt1__ZTSPi : !cir.ptr> init { // CHECK-NEXT: ^bb0(%arg0: !cir.ptr> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}): // TODO: Add Init here. // CHECK-NEXT: acc.yield // CHECK-NEXT: } ; #pragma acc parallel private(OnePtr[B]) ; #pragma acc parallel private(OnePtr) // CHECK-NEXT: acc.private.recipe @privatization__ZTSPi : !cir.ptr> init { // CHECK-NEXT: ^bb0(%arg0: !cir.ptr> {{.*}}): // CHECK-NEXT: cir.alloca !cir.ptr, !cir.ptr>, ["openacc.private.init"] {alignment = 8 : i64} // CHECK-NEXT: acc.yield // CHECK-NEXT: } ; T **TwoPtr; #pragma acc parallel private(TwoPtr[B][B]) // CHECK-NEXT: acc.private.recipe @privatization__Bcnt2__ZTSPPi : !cir.ptr>> init { // CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}, %[[BOUND2:.*]]: !acc.data_bounds_ty {{.*}}): // TODO: Add Init here. // CHECK-NEXT: acc.yield // CHECK-NEXT: } ; #pragma acc parallel private(TwoPtr[B][A:B]) ; #pragma acc parallel private(TwoPtr[A:B][A:B]) ; #pragma acc parallel private(TwoPtr) // CHECK-NEXT: acc.private.recipe @privatization__ZTSPPi : !cir.ptr>> init { // CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr>> {{.*}}): // CHECK-NEXT: cir.alloca !cir.ptr>, !cir.ptr>>, ["openacc.private.init"] {alignment = 8 : i64} // CHECK-NEXT: acc.yield // CHECK-NEXT: } ; T ***ThreePtr; #pragma acc parallel private(ThreePtr[B][B][B]) // CHECK-NEXT: acc.private.recipe @privatization__Bcnt3__ZTSPPPi : !cir.ptr>>> init { // CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr>>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}, %[[BOUND2:.*]]: !acc.data_bounds_ty {{.*}}, %[[BOUND3:.*]]: !acc.data_bounds_ty {{.*}}): // TODO: Add Init here. // CHECK-NEXT: acc.yield // CHECK-NEXT: } ; #pragma acc parallel private(ThreePtr[B][B][A:B]) ; #pragma acc parallel private(ThreePtr[B][A:B][A:B]) ; #pragma acc parallel private(ThreePtr[A:B][A:B][A:B]) ; #pragma acc parallel private(ThreePtr[B][B]) // CHECK-NEXT: acc.private.recipe @privatization__Bcnt2__ZTSPPPi : !cir.ptr>>> init { // CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr>>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}, %[[BOUND2:.*]]: !acc.data_bounds_ty {{.*}}): // TODO: Add Init here. // CHECK-NEXT: acc.yield // CHECK-NEXT: } ; #pragma acc parallel private(ThreePtr[B][A:B]) ; #pragma acc parallel private(ThreePtr[A:B][A:B]) ; #pragma acc parallel private(ThreePtr) // CHECK-NEXT: acc.private.recipe @privatization__ZTSPPPi : !cir.ptr>>> init { // CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr>>> {{.*}}): // CHECK-NEXT: cir.alloca !cir.ptr>>, !cir.ptr>>>, ["openacc.private.init"] {alignment = 8 : i64} // CHECK-NEXT: acc.yield // CHECK-NEXT: } ; T *ArrayOfPtr[5]; #pragma acc parallel private(ArrayOfPtr[B][A:B]) // CHECK-NEXT: acc.private.recipe @privatization__Bcnt2__ZTSA5_Pi : !cir.ptr x 5>> init { // CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}, %[[BOUND2:.*]]: !acc.data_bounds_ty {{.*}}): // TODO: Add Init here. // CHECK-NEXT: acc.yield // CHECK-NEXT: } ; #pragma acc parallel private(ArrayOfPtr[A:B][A:B]) ; #pragma acc parallel private(ArrayOfPtr[B][B]) ; #pragma acc parallel private(ArrayOfPtr) // CHECK-NEXT: acc.private.recipe @privatization__ZTSA5_Pi : !cir.ptr x 5>> init { // CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr x 5>> {{.*}}): // CHECK-NEXT: cir.alloca !cir.array x 5>, !cir.ptr x 5>>, ["openacc.private.init"] {alignment = 16 : i64} // CHECK-NEXT: acc.yield // CHECK-NEXT: } ; using TArrayTy = T[5]; TArrayTy *PtrToArrays; #pragma acc parallel private(PtrToArrays[B][B]) // CHECK-NEXT: acc.private.recipe @privatization__Bcnt2__ZTSPA5_i : !cir.ptr>> init { // CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}, %[[BOUND2:.*]]: !acc.data_bounds_ty {{.*}}): // TODO: Add Init here. // CHECK-NEXT: acc.yield // CHECK-NEXT: } ; #pragma acc parallel private(PtrToArrays[B][A:B]) ; #pragma acc parallel private(PtrToArrays[A:B][A:B]) ; #pragma acc parallel private(PtrToArrays) // CHECK-NEXT: acc.private.recipe @privatization__ZTSPA5_i : !cir.ptr>> init { // CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr>> {{.*}}): // CHECK-NEXT: cir.alloca !cir.ptr>, !cir.ptr>>, ["openacc.private.init"] {alignment = 8 : i64} // CHECK-NEXT: acc.yield // CHECK-NEXT: } ; T **ArrayOfPtrPtr[5]; #pragma acc parallel private(ArrayOfPtrPtr[B][B][B]) // CHECK-NEXT: acc.private.recipe @privatization__Bcnt3__ZTSA5_PPi : !cir.ptr> x 5>> init { // CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr> x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}, %[[BOUND2:.*]]: !acc.data_bounds_ty {{.*}}, %[[BOUND3:.*]]: !acc.data_bounds_ty {{.*}}): // TODO: Add Init here. // CHECK-NEXT: acc.yield // CHECK-NEXT: } ; #pragma acc parallel private(ArrayOfPtrPtr[B][B][A:B]) ; #pragma acc parallel private(ArrayOfPtrPtr[B][A:B][A:B]) ; #pragma acc parallel private(ArrayOfPtrPtr[A:B][A:B][A:B]) ; #pragma acc parallel private(ArrayOfPtrPtr[B][B]) // CHECK-NEXT: acc.private.recipe @privatization__Bcnt2__ZTSA5_PPi : !cir.ptr> x 5>> init { // CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr> x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}, %[[BOUND2:.*]]: !acc.data_bounds_ty {{.*}}): // TODO: Add Init here. // CHECK-NEXT: acc.yield // CHECK-NEXT: } ; #pragma acc parallel private(ArrayOfPtrPtr[B][A:B]) ; #pragma acc parallel private(ArrayOfPtrPtr[A:B][A:B]) ; #pragma acc parallel private(ArrayOfPtrPtr) // CHECK-NEXT: acc.private.recipe @privatization__ZTSA5_PPi : !cir.ptr> x 5>> init { // CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr> x 5>> {{.*}}): // CHECK-NEXT: cir.alloca !cir.array> x 5>, !cir.ptr> x 5>>, ["openacc.private.init"] {alignment = 16 : i64} // CHECK-NEXT: acc.yield // CHECK-NEXT: } ; TArrayTy **PtrPtrToArray; #pragma acc parallel private(PtrPtrToArray[B][B][B]) // CHECK-NEXT: acc.private.recipe @privatization__Bcnt3__ZTSPPA5_i : !cir.ptr>>> init { // CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr>>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}, %[[BOUND2:.*]]: !acc.data_bounds_ty {{.*}}, %[[BOUND3:.*]]: !acc.data_bounds_ty {{.*}}): // TODO: Add Init here. // CHECK-NEXT: acc.yield // CHECK-NEXT:} ; #pragma acc parallel private(PtrPtrToArray[B][B][A:B]) ; #pragma acc parallel private(PtrPtrToArray[B][A:B][A:B]) ; #pragma acc parallel private(PtrPtrToArray[A:B][A:B][A:B]) ; #pragma acc parallel private(PtrPtrToArray[B][B]) // CHECK-NEXT: acc.private.recipe @privatization__Bcnt2__ZTSPPA5_i : !cir.ptr>>> init { // CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr>>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}, %[[BOUND2:.*]]: !acc.data_bounds_ty {{.*}}): // TODO: Add Init here. // CHECK-NEXT: acc.yield // CHECK-NEXT: } ; #pragma acc parallel private(PtrPtrToArray[B][A:B]) ; #pragma acc parallel private(PtrPtrToArray[A:B][A:B]) ; #pragma acc parallel private(PtrPtrToArray) // CHECK-NEXT: acc.private.recipe @privatization__ZTSPPA5_i : !cir.ptr>>> init { // CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr>>> // CHECK-NEXT: cir.alloca !cir.ptr>>, !cir.ptr>>>, ["openacc.private.init"] {alignment = 8 : i64} // CHECK-NEXT: acc.yield // CHECK-NEXT: } ; using PtrTArrayTy = T*[5]; PtrTArrayTy *PtrArrayPtr; #pragma acc parallel private(PtrArrayPtr[B][B][B]) // CHECK-NEXT: acc.private.recipe @privatization__Bcnt3__ZTSPA5_Pi : !cir.ptr x 5>>> init { // CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr x 5>>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}, %[[BOUND2:.*]]: !acc.data_bounds_ty {{.*}}, %[[BOUND3:.*]]: !acc.data_bounds_ty {{.*}}): // TODO: Add Init here. // CHECK-NEXT: acc.yield // CHECK-NEXT: } ; #pragma acc parallel private(PtrArrayPtr[B][B][A:B]) ; #pragma acc parallel private(PtrArrayPtr[B][A:B][A:B]) ; #pragma acc parallel private(PtrArrayPtr[A:B][A:B][A:B]) ; #pragma acc parallel private(PtrArrayPtr[B][B]) // #pragma acc parallel private(PtrArrayPtr[B][B]) // CHECK-NEXT: acc.private.recipe @privatization__Bcnt2__ZTSPA5_Pi : !cir.ptr x 5>>> init { // CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr x 5>>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}, %[[BOUND2:.*]]: !acc.data_bounds_ty {{.*}}): // TODO: Add Init here. // CHECK-NEXT: acc.yield // CHECK-NEXT: } ; #pragma acc parallel private(PtrArrayPtr[B][A:B]) ; #pragma acc parallel private(PtrArrayPtr[A:B][A:B]) ; #pragma acc parallel private(PtrArrayPtr) // CHECK: acc.private.recipe @privatization__ZTSPA5_Pi : !cir.ptr x 5>>> init { // CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr x 5>>> {{.*}}): // CHECK-NEXT: cir.alloca !cir.ptr x 5>>, !cir.ptr x 5>>>, ["openacc.private.init"] // CHECK-NEXT: acc.yield // CHECK-NEXT: } ; } void use(unsigned A, unsigned B) { do_things(A, B); }