================
@@ -754,12 +754,351 @@ void acc_compute(int parmVar) {
// CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64
// CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUNDS3:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32)
extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64)
startIdx(%[[ZERO_CONST]] : i64) loc
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[THREEDARRAY]] :
!cir.ptr<!cir.array<!cir.array<!cir.array<!cir.double x 7> x 6> x 5>>)
bounds(%[[BOUNDS]], %[[BOUNDS2]], %[[BOUNDS3]]) ->
!cir.ptr<!cir.array<!cir.array<!cir.array<!cir.double x 7> x 6> x 5>>
{dataClause = #acc<data_clause acc_copy>, name = "threeDArray"} loc
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[THREEDARRAY]] :
!cir.ptr<!cir.array<!cir.array<!cir.array<!cir.double x 7> x 6> x 5>>)
bounds(%[[BOUNDS]], %[[BOUNDS2]], %[[BOUNDS3]]) ->
!cir.ptr<!cir.array<!cir.array<!cir.array<!cir.double x 7> x 6> x 5>>
{dataClause = #acc<data_clause acc_copy>, name = "threeDArray[1:1][2:1][3:1]"}
loc
// CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] :
!cir.ptr<!cir.array<!cir.array<!cir.array<!cir.double x 7> x 6> x 5>>) {
// CHECK-NEXT: acc.loop combined(parallel) {
// CHECK: acc.yield
// CHECK-NEXT: }
// CHECK-NEXT: acc.yield
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] :
!cir.ptr<!cir.array<!cir.array<!cir.array<!cir.double x 7> x 6> x 5>>)
bounds(%[[BOUNDS]], %[[BOUNDS2]], %[[BOUNDS3]]) to varPtr(%[[THREEDARRAY]] :
!cir.ptr<!cir.array<!cir.array<!cir.array<!cir.double x 7> x 6> x 5>>)
{dataClause = #acc<data_clause acc_copy>, name = "threeDArray"} loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] :
!cir.ptr<!cir.array<!cir.array<!cir.array<!cir.double x 7> x 6> x 5>>)
bounds(%[[BOUNDS]], %[[BOUNDS2]], %[[BOUNDS3]]) to varPtr(%[[THREEDARRAY]] :
!cir.ptr<!cir.array<!cir.array<!cir.array<!cir.double x 7> x 6> x 5>>)
{dataClause = #acc<data_clause acc_copy>, name = "threeDArray[1:1][2:1][3:1]"}
loc
+}
+
+typedef struct StructTy {
+ int scalarMember;
+ int arrayMember[5];
+ short twoDArrayMember[5][3];
+ float *ptrArrayMember[5];
+ double **ptrPtrMember;
+} Struct ;
+
+void acc_compute_members() {
+ // CHECK: cir.func @acc_compute_members()
+ Struct localStruct;
+ // CHECK-NEXT: %[[LOCALSTRUCT:.*]] = cir.alloca !rec_StructTy,
!cir.ptr<!rec_StructTy>, ["localStruct"]
+
+#pragma acc parallel loop copy(localStruct)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALSTRUCT]] :
!cir.ptr<!rec_StructTy>) -> !cir.ptr<!rec_StructTy> {dataClause =
#acc<data_clause acc_copy>, name = "localStruct"}
+ // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] :
!cir.ptr<!rec_StructTy>) {
+ // CHECK-NEXT: acc.loop combined(parallel) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: }
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!rec_StructTy>) to
varPtr(%[[LOCALSTRUCT]] : !cir.ptr<!rec_StructTy>) {dataClause =
#acc<data_clause acc_copy>, name = "localStruct"}
+
+#pragma acc serial loop copy(localStruct.scalarMember)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: %[[GETMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][0] {name
= "scalarMember"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETMEMBER]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>,
name = "localStruct.scalarMember"}
+ // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] :
!cir.ptr<!s32i>) {
+ // CHECK-NEXT: acc.loop combined(serial) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: }
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to
varPtr(%[[GETMEMBER]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause
acc_copy>, name = "localStruct.scalarMember"}
+
+#pragma acc kernels loop copy(localStruct.arrayMember)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: %[[GETARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][1]
{name = "arrayMember"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!cir.array<!s32i x
5>>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETARRAYMEMBER]] :
!cir.ptr<!cir.array<!s32i x 5>>) -> !cir.ptr<!cir.array<!s32i x 5>> {dataClause
= #acc<data_clause acc_copy>, name = "localStruct.arrayMember"}
+ // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] :
!cir.ptr<!cir.array<!s32i x 5>>) {
+ // CHECK-NEXT: acc.loop combined(kernels) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: }
+ // CHECK-NEXT: acc.terminator
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x
5>>) to varPtr(%[[GETARRAYMEMBER]] : !cir.ptr<!cir.array<!s32i x 5>>)
{dataClause = #acc<data_clause acc_copy>, name = "localStruct.arrayMember"} loc
+
+#pragma acc parallel loop copy(localStruct.arrayMember[2])
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: %[[TWO_CONST:.*]] = cir.const #cir.int<2> : !s32i
+ // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast
%[[TWO_CONST]] : !s32i to si32
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64
+ // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32)
extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64)
startIdx(%[[ZERO_CONST]] : i64)
+ // CHECK-NEXT: %[[GETARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][1]
{name = "arrayMember"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!cir.array<!s32i x
5>>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETARRAYMEMBER]] :
!cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS]]) ->
!cir.ptr<!cir.array<!s32i x 5>> {dataClause = #acc<data_clause acc_copy>, name
= "localStruct.arrayMember[2]"}
+ // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] :
!cir.ptr<!cir.array<!s32i x 5>>) {
+ // CHECK-NEXT: acc.loop combined(parallel) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: }
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x
5>>) bounds(%[[BOUNDS]]) to varPtr(%[[GETARRAYMEMBER]] :
!cir.ptr<!cir.array<!s32i x 5>>) {dataClause = #acc<data_clause acc_copy>, name
= "localStruct.arrayMember[2]"} loc
+
+#pragma acc serial loop copy(localStruct.arrayMember[1:2])
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast
%[[ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i
+ // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast
%[[TWO]] : !s32i to si32
+ // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32)
extent(%[[TWO_CAST]] : si32) stride(%[[ONE_CONST]] : i64)
startIdx(%[[ZERO_CONST]] : i64)
+ // CHECK-NEXT: %[[GETARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][1]
{name = "arrayMember"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!cir.array<!s32i x
5>>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETARRAYMEMBER]] :
!cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS]]) ->
!cir.ptr<!cir.array<!s32i x 5>> {dataClause = #acc<data_clause acc_copy>, name
= "localStruct.arrayMember[1:2]"}
+ // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] :
!cir.ptr<!cir.array<!s32i x 5>>) {
+ // CHECK-NEXT: acc.loop combined(serial) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: }
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x
5>>) bounds(%[[BOUNDS]]) to varPtr(%[[GETARRAYMEMBER]] :
!cir.ptr<!cir.array<!s32i x 5>>) {dataClause = #acc<data_clause acc_copy>, name
= "localStruct.arrayMember[1:2]"} loc
+
+#pragma acc kernels loop copy(localStruct.arrayMember[1:])
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast
%[[ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[FOUR_CONST:.*]] = arith.constant 4 : i64
+ // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32)
upperbound(%[[FOUR_CONST]] : i64) stride(%[[ONE_CONST]] : i64)
startIdx(%[[ZERO_CONST]] : i64)
+ // CHECK-NEXT: %[[GETARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][1]
{name = "arrayMember"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!cir.array<!s32i x
5>>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETARRAYMEMBER]] :
!cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS]]) ->
!cir.ptr<!cir.array<!s32i x 5>> {dataClause = #acc<data_clause acc_copy>, name
= "localStruct.arrayMember[1:]"}
+ // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] :
!cir.ptr<!cir.array<!s32i x 5>>) {
+ // CHECK-NEXT: acc.loop combined(kernels) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: }
+ // CHECK-NEXT: acc.terminator
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x
5>>) bounds(%[[BOUNDS]]) to varPtr(%[[GETARRAYMEMBER]] :
!cir.ptr<!cir.array<!s32i x 5>>) {dataClause = #acc<data_clause acc_copy>, name
= "localStruct.arrayMember[1:]"} loc
+
+#pragma acc parallel loop copy(localStruct.arrayMember[:2])
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64
+ // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i
+ // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast
%[[TWO]] : !s32i to si32
+ // CHECK-NEXT: %[[ZERO_CONST2:.*]] = arith.constant 0 : i64
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ZERO_CONST]] :
i64) extent(%[[TWO_CAST]] : si32) stride(%[[ONE_CONST]] : i64)
startIdx(%[[ZERO_CONST2]] : i64)
+ // CHECK-NEXT: %[[GETARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][1]
{name = "arrayMember"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!cir.array<!s32i x
5>>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETARRAYMEMBER]] :
!cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS]]) ->
!cir.ptr<!cir.array<!s32i x 5>> {dataClause = #acc<data_clause acc_copy>, name
= "localStruct.arrayMember[:2]"}
+ // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] :
!cir.ptr<!cir.array<!s32i x 5>>) {
+ // CHECK-NEXT: acc.loop combined(parallel) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: }
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x
5>>) bounds(%[[BOUNDS]]) to varPtr(%[[GETARRAYMEMBER]] :
!cir.ptr<!cir.array<!s32i x 5>>) {dataClause = #acc<data_clause acc_copy>, name
= "localStruct.arrayMember[:2]"} loc
+
+#pragma acc serial loop copy(localStruct.twoDArrayMember)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: %[[GET2DARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][2]
{name = "twoDArrayMember"} : !cir.ptr<!rec_StructTy> ->
!cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GET2DARRAYMEMBER]] :
!cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>) ->
!cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>> {dataClause = #acc<data_clause
acc_copy>, name = "localStruct.twoDArrayMember"}
+ // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] :
!cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>) {
+ // CHECK-NEXT: acc.loop combined(serial) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: }
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] :
!cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>) to
varPtr(%[[GET2DARRAYMEMBER]] : !cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>)
{dataClause = #acc<data_clause acc_copy>, name = "localStruct.twoDArrayMember"}
+
+#pragma acc kernels loop copy(localStruct.twoDArrayMember[3][2])
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i
+ // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast
%[[TWO]] : !s32i to si32
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64
+ // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32)
extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64)
startIdx(%[[ZERO_CONST]] : i64)
+ // CHECK-NEXT: %[[THREE:.*]] = cir.const #cir.int<3> : !s32i
+ // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast
%[[THREE]] : !s32i to si32
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64
+ // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[THREE_CAST]] :
si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64)
startIdx(%[[ZERO_CONST]] : i64)
+ // CHECK-NEXT: %[[GET2DARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][2]
{name = "twoDArrayMember"} : !cir.ptr<!rec_StructTy> ->
!cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GET2DARRAYMEMBER]] :
!cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>) bounds(%[[BOUNDS1]],
%[[BOUNDS2]]) -> !cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>> {dataClause =
#acc<data_clause acc_copy>, name = "localStruct.twoDArrayMember[3][2]"}
+ // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] :
!cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>) {
+ // CHECK-NEXT: acc.loop combined(kernels) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: }
+ // CHECK-NEXT: acc.terminator
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] :
!cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>) bounds(%[[BOUNDS1]],
%[[BOUNDS2]]) to varPtr(%[[GET2DARRAYMEMBER]] :
!cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>) {dataClause = #acc<data_clause
acc_copy>, name = "localStruct.twoDArrayMember[3][2]"}
+
+#pragma acc parallel loop copy(localStruct.twoDArrayMember[1:3][1:1])
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast
%[[ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast
%[[ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64
+ // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32)
extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64)
startIdx(%[[ZERO_CONST]] : i64)
+ // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast
%[[ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[THREE:.*]] = cir.const #cir.int<3> : !s32i
+ // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast
%[[THREE]] : !s32i to si32
+ // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64
+ // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32)
extent(%[[THREE_CAST]] : si32) stride(%[[ONE_CONST2]] : i64)
startIdx(%[[ZERO_CONST]] : i64)
+ // CHECK-NEXT: %[[GET2DARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][2]
{name = "twoDArrayMember"} : !cir.ptr<!rec_StructTy> ->
!cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GET2DARRAYMEMBER]] :
!cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>) bounds(%[[BOUNDS1]],
%[[BOUNDS2]]) -> !cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>> {dataClause =
#acc<data_clause acc_copy>, name = "localStruct.twoDArrayMember[1:3][1:1]"}
+ // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] :
!cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>) {
+ // CHECK-NEXT: acc.loop combined(parallel) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: }
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] :
!cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>) bounds(%[[BOUNDS1]],
%[[BOUNDS2]]) to varPtr(%[[GET2DARRAYMEMBER]] :
!cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>) {dataClause = #acc<data_clause
acc_copy>, name = "localStruct.twoDArrayMember[1:3][1:1]"}
+
+#pragma acc serial loop copy(localStruct.ptrArrayMember)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: %[[GETPTRARRAYMEMBER:.*]] = cir.get_member
%[[LOCALSTRUCT]][3] {name = "ptrArrayMember"} : !cir.ptr<!rec_StructTy> ->
!cir.ptr<!cir.array<!cir.ptr<!cir.float> x 5>>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETPTRARRAYMEMBER]] :
!cir.ptr<!cir.array<!cir.ptr<!cir.float> x 5>>) ->
!cir.ptr<!cir.array<!cir.ptr<!cir.float> x 5>> {dataClause = #acc<data_clause
acc_copy>, name = "localStruct.ptrArrayMember"}
+ // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] :
!cir.ptr<!cir.array<!cir.ptr<!cir.float> x 5>>) {
+ // CHECK-NEXT: acc.loop combined(serial) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: }
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] :
!cir.ptr<!cir.array<!cir.ptr<!cir.float> x 5>>) to
varPtr(%[[GETPTRARRAYMEMBER]] : !cir.ptr<!cir.array<!cir.ptr<!cir.float> x 5>>)
{dataClause = #acc<data_clause acc_copy>, name = "localStruct.ptrArrayMember"}
+
+#pragma acc kernels loop copy(localStruct.ptrArrayMember[3][2])
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i
+ // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast
%[[TWO]] : !s32i to si32
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64
+ // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32)
extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64)
startIdx(%[[ZERO_CONST]] : i64)
+ // CHECK-NEXT: %[[THREE:.*]] = cir.const #cir.int<3> : !s32i
+ // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast
%[[THREE]] : !s32i to si32
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64
+ // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[THREE_CAST]] :
si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64)
startIdx(%[[ZERO_CONST]] : i64)
+ // CHECK-NEXT: %[[GETPTRARRAYMEMBER:.*]] = cir.get_member
%[[LOCALSTRUCT]][3] {name = "ptrArrayMember"} : !cir.ptr<!rec_StructTy> ->
!cir.ptr<!cir.array<!cir.ptr<!cir.float> x 5>>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETPTRARRAYMEMBER]] :
!cir.ptr<!cir.array<!cir.ptr<!cir.float> x 5>>) bounds(%[[BOUNDS1]],
%[[BOUNDS2]]) -> !cir.ptr<!cir.array<!cir.ptr<!cir.float> x 5>> {dataClause =
#acc<data_clause acc_copy>, name = "localStruct.ptrArrayMember[3][2]"}
+ // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] :
!cir.ptr<!cir.array<!cir.ptr<!cir.float> x 5>>) {
+ // CHECK-NEXT: acc.loop combined(kernels) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: }
+ // CHECK-NEXT: acc.terminator
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] :
!cir.ptr<!cir.array<!cir.ptr<!cir.float> x 5>>) bounds(%[[BOUNDS1]],
%[[BOUNDS2]]) to varPtr(%[[GETPTRARRAYMEMBER]] :
!cir.ptr<!cir.array<!cir.ptr<!cir.float> x 5>>) {dataClause = #acc<data_clause
acc_copy>, name = "localStruct.ptrArrayMember[3][2]"}
+
+#pragma acc parallel loop copy(localStruct.ptrArrayMember[1:3][1:1])
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast
%[[ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast
%[[ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64
+ // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32)
extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64)
startIdx(%[[ZERO_CONST]] : i64)
+ // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast
%[[ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[THREE:.*]] = cir.const #cir.int<3> : !s32i
+ // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast
%[[THREE]] : !s32i to si32
+ // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64
+ // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32)
extent(%[[THREE_CAST]] : si32) stride(%[[ONE_CONST2]] : i64)
startIdx(%[[ZERO_CONST]] : i64)
+ // CHECK-NEXT: %[[GETPTRARRAYMEMBER:.*]] = cir.get_member
%[[LOCALSTRUCT]][3] {name = "ptrArrayMember"} : !cir.ptr<!rec_StructTy> ->
!cir.ptr<!cir.array<!cir.ptr<!cir.float> x 5>>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETPTRARRAYMEMBER]] :
!cir.ptr<!cir.array<!cir.ptr<!cir.float> x 5>>) bounds(%[[BOUNDS1]],
%[[BOUNDS2]]) -> !cir.ptr<!cir.array<!cir.ptr<!cir.float> x 5>> {dataClause =
#acc<data_clause acc_copy>, name = "localStruct.ptrArrayMember[1:3][1:1]"}
+ // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] :
!cir.ptr<!cir.array<!cir.ptr<!cir.float> x 5>>) {
+ // CHECK-NEXT: acc.loop combined(parallel) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: }
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] :
!cir.ptr<!cir.array<!cir.ptr<!cir.float> x 5>>) bounds(%[[BOUNDS1]],
%[[BOUNDS2]]) to varPtr(%[[GETPTRARRAYMEMBER]] :
!cir.ptr<!cir.array<!cir.ptr<!cir.float> x 5>>) {dataClause = #acc<data_clause
acc_copy>, name = "localStruct.ptrArrayMember[1:3][1:1]"}
+
+#pragma acc serial loop copy(localStruct.ptrPtrMember)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: %[[GETPTRPTRMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][4]
{name = "ptrPtrMember"} : !cir.ptr<!rec_StructTy> ->
!cir.ptr<!cir.ptr<!cir.ptr<!cir.double>>>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETPTRPTRMEMBER]] :
!cir.ptr<!cir.ptr<!cir.ptr<!cir.double>>>) ->
!cir.ptr<!cir.ptr<!cir.ptr<!cir.double>>> {dataClause = #acc<data_clause
acc_copy>, name = "localStruct.ptrPtrMember"}
+ // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] :
!cir.ptr<!cir.ptr<!cir.ptr<!cir.double>>>) {
+ // CHECK-NEXT: acc.loop combined(serial) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: }
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] :
!cir.ptr<!cir.ptr<!cir.ptr<!cir.double>>>) to varPtr(%[[GETPTRPTRMEMBER]] :
!cir.ptr<!cir.ptr<!cir.ptr<!cir.double>>>) {dataClause = #acc<data_clause
acc_copy>, name = "localStruct.ptrPtrMember"}
+
+#pragma acc kernels loop copy(localStruct.ptrPtrMember[3][2])
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i
+ // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast
%[[TWO]] : !s32i to si32
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64
+ // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32)
extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64)
startIdx(%[[ZERO_CONST]] : i64)
+ // CHECK-NEXT: %[[THREE:.*]] = cir.const #cir.int<3> : !s32i
+ // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast
%[[THREE]] : !s32i to si32
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64
+ // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[THREE_CAST]] :
si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64)
startIdx(%[[ZERO_CONST]] : i64)
+ // CHECK-NEXT: %[[GETPTRPTRMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][4]
{name = "ptrPtrMember"} : !cir.ptr<!rec_StructTy> ->
!cir.ptr<!cir.ptr<!cir.ptr<!cir.double>>>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETPTRPTRMEMBER]] :
!cir.ptr<!cir.ptr<!cir.ptr<!cir.double>>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]])
-> !cir.ptr<!cir.ptr<!cir.ptr<!cir.double>>> {dataClause = #acc<data_clause
acc_copy>, name = "localStruct.ptrPtrMember[3][2]"}
+ // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] :
!cir.ptr<!cir.ptr<!cir.ptr<!cir.double>>>) {
+ // CHECK-NEXT: acc.loop combined(kernels) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: }
+ // CHECK-NEXT: acc.terminator
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] :
!cir.ptr<!cir.ptr<!cir.ptr<!cir.double>>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]])
to varPtr(%[[GETPTRPTRMEMBER]] : !cir.ptr<!cir.ptr<!cir.ptr<!cir.double>>>)
{dataClause = #acc<data_clause acc_copy>, name =
"localStruct.ptrPtrMember[3][2]"}
+
+#pragma acc parallel loop copy(localStruct.ptrPtrMember[1:3][1:1])
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast
%[[ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast
%[[ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64
+ // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32)
extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64)
startIdx(%[[ZERO_CONST]] : i64)
+ // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast
%[[ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[THREE:.*]] = cir.const #cir.int<3> : !s32i
+ // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast
%[[THREE]] : !s32i to si32
+ // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64
+ // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32)
extent(%[[THREE_CAST]] : si32) stride(%[[ONE_CONST2]] : i64)
startIdx(%[[ZERO_CONST]] : i64)
+ // CHECK-NEXT: %[[GETPTRPTRMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][4]
{name = "ptrPtrMember"} : !cir.ptr<!rec_StructTy> ->
!cir.ptr<!cir.ptr<!cir.ptr<!cir.double>>>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETPTRPTRMEMBER]] :
!cir.ptr<!cir.ptr<!cir.ptr<!cir.double>>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]])
-> !cir.ptr<!cir.ptr<!cir.ptr<!cir.double>>> {dataClause = #acc<data_clause
acc_copy>, name = "localStruct.ptrPtrMember[1:3][1:1]"}
+ // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] :
!cir.ptr<!cir.ptr<!cir.ptr<!cir.double>>>) {
+ // CHECK-NEXT: acc.loop combined(parallel) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: }
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] :
!cir.ptr<!cir.ptr<!cir.ptr<!cir.double>>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]])
to varPtr(%[[GETPTRPTRMEMBER]] : !cir.ptr<!cir.ptr<!cir.ptr<!cir.double>>>)
{dataClause = #acc<data_clause acc_copy>, name =
"localStruct.ptrPtrMember[1:3][1:1]"}
+}
+
+typedef struct InnerTy {
+ int a;
+ int b;
+} Inner;
+
+typedef struct OuterTy {
+ Inner inner[4];
+} Outer;
+
+void copy_member_of_array_element_member() {
+ // CHECK: cir.func @copy_member_of_array_element_member() {
+ Outer outer;
+ // CHECK-NEXT: %[[OUTER:.*]] = cir.alloca !rec_OuterTy,
!cir.ptr<!rec_OuterTy>, ["outer"]
+
+ #pragma acc parallel loop copy(outer.inner[2].b)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i
+ // CHECK-NEXT: %[[GETINNER:.*]] = cir.get_member %[[OUTER]][0] {name =
"inner"} : !cir.ptr<!rec_OuterTy> -> !cir.ptr<!cir.array<!rec_InnerTy x 4>>
+ // CHECK-NEXT: %[[INNERDECAY:.*]] = cir.cast(array_to_ptrdecay,
%[[GETINNER]] : !cir.ptr<!cir.array<!rec_InnerTy x 4>>), !cir.ptr<!rec_InnerTy>
+ // CHECK-NEXT: %[[STRIDE:.*]] = cir.ptr_stride(%[[INNERDECAY]] :
!cir.ptr<!rec_InnerTy>, %[[TWO]] : !s32i), !cir.ptr<!rec_InnerTy>
+ // CHECK-NEXT: %[[GETB:.*]] = cir.get_member %[[STRIDE]][1] {name = "b"} :
!cir.ptr<!rec_InnerTy> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETB]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>,
name = "outer.inner[2].b"}
----------------
andykaylor wrote:
This looks good. It's finding the member in the way I expected, and the
PointerLikeType handling will work with this.
https://github.com/llvm/llvm-project/pull/142998
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
