================
@@ -754,12 +754,320 @@ 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]"}
----------------
erichkeane wrote:
None that I can think of... The OpenACC dialect defines these 'bounds', and the
'stride' of arrays/etc in C++ are always 1, based on type.
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
