347 lines · cpp
1// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s2 3struct InnerStructTy {4 int Member[5];5};6struct StructTy {7 int scalarMember;8 int arrayMember[5];9 short twoDArrayMember[5][3];10 InnerStructTy iSTy;11 12void InlineFunc() {13 // CHECK: cir.func {{.*}}InlineFunc{{.*}}14 // CHECK-NEXT: %[[THIS:.*]] = cir.alloca !cir.ptr<!rec_StructTy>, !cir.ptr<!cir.ptr<!rec_StructTy>>, ["this", init]15 // CHECK-NEXT: cir.store %[[THIS_ARG:.*]], %[[THIS]] : !cir.ptr<!rec_StructTy>, !cir.ptr<!cir.ptr<!rec_StructTy>>16 // CHECK-NEXT: %[[THIS_LOAD:.*]] = cir.load %[[THIS]] : !cir.ptr<!cir.ptr<!rec_StructTy>>, !cir.ptr<!rec_StructTy>17 18#pragma acc parallel copy(scalarMember)19 ;20 // CHECK-NEXT: %[[GETSCALARMEM:.*]] = cir.get_member %[[THIS_LOAD]][0] {name = "scalarMember"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!s32i>21 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSCALARMEM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "this->scalarMember"}22 // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {23 // CHECK-NEXT: acc.yield24 // CHECK-NEXT: } loc25 // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[GETSCALARMEM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "this->scalarMember"}26 27#pragma acc kernels copy(arrayMember[2])28 ;29 // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i30 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si3231 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i6432 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i6433 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i6434 // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)35 // CHECK-NEXT: %[[GETARRAYMEM:.*]] = cir.get_member %[[THIS_LOAD]][1] {name = "arrayMember"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!cir.array<!s32i x 5>>36 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETARRAYMEM]] : !cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS1]]) -> !cir.ptr<!cir.array<!s32i x 5>> {dataClause = #acc<data_clause acc_copy>, name = "this->arrayMember[2]"}37 // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) {38 // CHECK-NEXT: acc.terminator39 // CHECK-NEXT: } loc40 // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS1]]) to varPtr(%[[GETARRAYMEM]] : !cir.ptr<!cir.array<!s32i x 5>>) {dataClause = #acc<data_clause acc_copy>, name = "this->arrayMember[2]"}41 42#pragma acc kernels copy(twoDArrayMember[1][2])43 ;44 // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i45 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si3246 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i6447 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i6448 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i6449 // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)50 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i51 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si3252 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i6453 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i6454 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i6455 // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)56 // CHECK-NEXT: %[[GET2DARRAYMEM:.*]] = cir.get_member %[[THIS_LOAD]][2] {name = "twoDArrayMember"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>57 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GET2DARRAYMEM]] : !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 = "this->twoDArrayMember[1][2]"}58 // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>) {59 // CHECK-NEXT: acc.terminator60 // CHECK-NEXT: } loc61 // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) to varPtr(%[[GET2DARRAYMEM]] : !cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>) {dataClause = #acc<data_clause acc_copy>, name = "this->twoDArrayMember[1][2]"}62 63#pragma acc kernels copy(iSTy)64 ;65 // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!rec_InnerStructTy>66 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSTRUCTMEM]] : !cir.ptr<!rec_InnerStructTy>) -> !cir.ptr<!rec_InnerStructTy> {dataClause = #acc<data_clause acc_copy>, name = "this->iSTy"}67 // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr<!rec_InnerStructTy>) {68 // CHECK-NEXT: acc.terminator69 // CHECK-NEXT: } loc70 // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!rec_InnerStructTy>) to varPtr(%[[GETSTRUCTMEM]] : !cir.ptr<!rec_InnerStructTy>) {dataClause = #acc<data_clause acc_copy>, name = "this->iSTy"}71 72#pragma acc parallel copy(iSTy.Member)73 ;74 // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!rec_InnerStructTy>75 // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr<!rec_InnerStructTy> -> !cir.ptr<!cir.array<!s32i x 5>>76 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr<!cir.array<!s32i x 5>>) -> !cir.ptr<!cir.array<!s32i x 5>> {dataClause = #acc<data_clause acc_copy>, name = "this->iSTy.Member"}77 // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) {78 // CHECK-NEXT: acc.yield79 // CHECK-NEXT: } loc80 // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) to varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr<!cir.array<!s32i x 5>>) {dataClause = #acc<data_clause acc_copy>, name = "this->iSTy.Member"}81 82#pragma acc serial copy(iSTy.Member[1])83 ;84 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i85 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si3286 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i6487 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i6488 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i6489 // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)90 // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!rec_InnerStructTy>91 // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr<!rec_InnerStructTy> -> !cir.ptr<!cir.array<!s32i x 5>>92 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS1]]) -> !cir.ptr<!cir.array<!s32i x 5>> {dataClause = #acc<data_clause acc_copy>, name = "this->iSTy.Member[1]"}93 // CHECK-NEXT: acc.serial dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) {94 // CHECK-NEXT: acc.yield95 // CHECK-NEXT: } loc96 // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS1]]) to varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr<!cir.array<!s32i x 5>>) {dataClause = #acc<data_clause acc_copy>, name = "this->iSTy.Member[1]"}97 98#pragma acc parallel copy(this->scalarMember)99 ;100 // CHECK-NEXT: %[[GETSCALARMEM:.*]] = cir.get_member %[[THIS_LOAD]][0] {name = "scalarMember"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!s32i>101 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSCALARMEM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "this->scalarMember"}102 // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {103 // CHECK-NEXT: acc.yield104 // CHECK-NEXT: } loc105 // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[GETSCALARMEM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "this->scalarMember"}106 107#pragma acc kernels copy(this->arrayMember[2])108 ;109 // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i110 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32111 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64112 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64113 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64114 // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)115 // CHECK-NEXT: %[[GETARRAYMEM:.*]] = cir.get_member %[[THIS_LOAD]][1] {name = "arrayMember"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!cir.array<!s32i x 5>>116 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETARRAYMEM]] : !cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS1]]) -> !cir.ptr<!cir.array<!s32i x 5>> {dataClause = #acc<data_clause acc_copy>, name = "this->arrayMember[2]"}117 // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) {118 // CHECK-NEXT: acc.terminator119 // CHECK-NEXT: } loc120 // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS1]]) to varPtr(%[[GETARRAYMEM]] : !cir.ptr<!cir.array<!s32i x 5>>) {dataClause = #acc<data_clause acc_copy>, name = "this->arrayMember[2]"}121#pragma acc kernels copy(this->twoDArrayMember[1][2])122 ;123 // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i124 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32125 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64126 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64127 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64128 // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)129 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i130 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32131 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64132 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64133 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64134 // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)135 // CHECK-NEXT: %[[GET2DARRAYMEM:.*]] = cir.get_member %[[THIS_LOAD]][2] {name = "twoDArrayMember"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>136 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GET2DARRAYMEM]] : !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 = "this->twoDArrayMember[1][2]"}137 // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>) {138 // CHECK-NEXT: acc.terminator139 // CHECK-NEXT: } loc140 // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) to varPtr(%[[GET2DARRAYMEM]] : !cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>) {dataClause = #acc<data_clause acc_copy>, name = "this->twoDArrayMember[1][2]"}141 142#pragma acc kernels copy(this->iSTy)143 ;144 // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!rec_InnerStructTy>145 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSTRUCTMEM]] : !cir.ptr<!rec_InnerStructTy>) -> !cir.ptr<!rec_InnerStructTy> {dataClause = #acc<data_clause acc_copy>, name = "this->iSTy"}146 // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr<!rec_InnerStructTy>) {147 // CHECK-NEXT: acc.terminator148 // CHECK-NEXT: } loc149 // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!rec_InnerStructTy>) to varPtr(%[[GETSTRUCTMEM]] : !cir.ptr<!rec_InnerStructTy>) {dataClause = #acc<data_clause acc_copy>, name = "this->iSTy"}150 151#pragma acc parallel copy(this->iSTy.Member)152 ;153 // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!rec_InnerStructTy>154 // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr<!rec_InnerStructTy> -> !cir.ptr<!cir.array<!s32i x 5>>155 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr<!cir.array<!s32i x 5>>) -> !cir.ptr<!cir.array<!s32i x 5>> {dataClause = #acc<data_clause acc_copy>, name = "this->iSTy.Member"}156 // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) {157 // CHECK-NEXT: acc.yield158 // CHECK-NEXT: } loc159 // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) to varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr<!cir.array<!s32i x 5>>) {dataClause = #acc<data_clause acc_copy>, name = "this->iSTy.Member"}160 161#pragma acc serial copy(this->iSTy.Member[1])162 ;163 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i164 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32165 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64166 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64167 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64168 // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)169 // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!rec_InnerStructTy>170 // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr<!rec_InnerStructTy> -> !cir.ptr<!cir.array<!s32i x 5>>171 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS1]]) -> !cir.ptr<!cir.array<!s32i x 5>> {dataClause = #acc<data_clause acc_copy>, name = "this->iSTy.Member[1]"}172 // CHECK-NEXT: acc.serial dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) {173 // CHECK-NEXT: acc.yield174 // CHECK-NEXT: } loc175 // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS1]]) to varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr<!cir.array<!s32i x 5>>) {dataClause = #acc<data_clause acc_copy>, name = "this->iSTy.Member[1]"}176}177 178void OutlineFunc();179};180 181void InlineUse() {182 StructTy s;183 s.InlineFunc();184}185 186void StructTy::OutlineFunc() {187 // CHECK: cir.func {{.*}}OutlineFunc{{.*}}188 // CHECK-NEXT: %[[THIS:.*]] = cir.alloca !cir.ptr<!rec_StructTy>, !cir.ptr<!cir.ptr<!rec_StructTy>>, ["this", init]189 // CHECK-NEXT: cir.store %[[THIS_ARG:.*]], %[[THIS]] : !cir.ptr<!rec_StructTy>, !cir.ptr<!cir.ptr<!rec_StructTy>>190 // CHECK-NEXT: %[[THIS_LOAD:.*]] = cir.load %[[THIS]] : !cir.ptr<!cir.ptr<!rec_StructTy>>, !cir.ptr<!rec_StructTy>191#pragma acc parallel copy(scalarMember)192 ;193 // CHECK-NEXT: %[[GETSCALARMEM:.*]] = cir.get_member %[[THIS_LOAD]][0] {name = "scalarMember"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!s32i>194 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSCALARMEM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "this->scalarMember"}195 // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {196 // CHECK-NEXT: acc.yield197 // CHECK-NEXT: } loc198 // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[GETSCALARMEM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "this->scalarMember"}199#pragma acc kernels copy(arrayMember[2])200 ;201 // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i202 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32203 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64204 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64205 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64206 // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)207 // CHECK-NEXT: %[[GETARRAYMEM:.*]] = cir.get_member %[[THIS_LOAD]][1] {name = "arrayMember"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!cir.array<!s32i x 5>>208 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETARRAYMEM]] : !cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS1]]) -> !cir.ptr<!cir.array<!s32i x 5>> {dataClause = #acc<data_clause acc_copy>, name = "this->arrayMember[2]"}209 // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) {210 // CHECK-NEXT: acc.terminator211 // CHECK-NEXT: } loc212 // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS1]]) to varPtr(%[[GETARRAYMEM]] : !cir.ptr<!cir.array<!s32i x 5>>) {dataClause = #acc<data_clause acc_copy>, name = "this->arrayMember[2]"}213#pragma acc kernels copy(twoDArrayMember[1][2])214 ;215 // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i216 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32217 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64218 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64219 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64220 // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)221 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i222 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32223 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64224 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64225 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64226 // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)227 // CHECK-NEXT: %[[GET2DARRAYMEM:.*]] = cir.get_member %[[THIS_LOAD]][2] {name = "twoDArrayMember"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>228 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GET2DARRAYMEM]] : !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 = "this->twoDArrayMember[1][2]"}229 // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>) {230 // CHECK-NEXT: acc.terminator231 // CHECK-NEXT: } loc232 // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) to varPtr(%[[GET2DARRAYMEM]] : !cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>) {dataClause = #acc<data_clause acc_copy>, name = "this->twoDArrayMember[1][2]"}233#pragma acc kernels copy(iSTy)234 ;235 // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!rec_InnerStructTy>236 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSTRUCTMEM]] : !cir.ptr<!rec_InnerStructTy>) -> !cir.ptr<!rec_InnerStructTy> {dataClause = #acc<data_clause acc_copy>, name = "this->iSTy"}237 // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr<!rec_InnerStructTy>) {238 // CHECK-NEXT: acc.terminator239 // CHECK-NEXT: } loc240 // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!rec_InnerStructTy>) to varPtr(%[[GETSTRUCTMEM]] : !cir.ptr<!rec_InnerStructTy>) {dataClause = #acc<data_clause acc_copy>, name = "this->iSTy"}241 242#pragma acc parallel copy(iSTy.Member)243 ;244 // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!rec_InnerStructTy>245 // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr<!rec_InnerStructTy> -> !cir.ptr<!cir.array<!s32i x 5>>246 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr<!cir.array<!s32i x 5>>) -> !cir.ptr<!cir.array<!s32i x 5>> {dataClause = #acc<data_clause acc_copy>, name = "this->iSTy.Member"}247 // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) {248 // CHECK-NEXT: acc.yield249 // CHECK-NEXT: } loc250 // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) to varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr<!cir.array<!s32i x 5>>) {dataClause = #acc<data_clause acc_copy>, name = "this->iSTy.Member"}251 252#pragma acc serial copy(iSTy.Member[1])253 ;254 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i255 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32256 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64257 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64258 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64259 // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)260 // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!rec_InnerStructTy>261 // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr<!rec_InnerStructTy> -> !cir.ptr<!cir.array<!s32i x 5>>262 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS1]]) -> !cir.ptr<!cir.array<!s32i x 5>> {dataClause = #acc<data_clause acc_copy>, name = "this->iSTy.Member[1]"}263 // CHECK-NEXT: acc.serial dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) {264 // CHECK-NEXT: acc.yield265 // CHECK-NEXT: } loc266 // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS1]]) to varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr<!cir.array<!s32i x 5>>) {dataClause = #acc<data_clause acc_copy>, name = "this->iSTy.Member[1]"}267 268#pragma acc parallel copy(this->scalarMember)269 ;270 // CHECK-NEXT: %[[GETSCALARMEM:.*]] = cir.get_member %[[THIS_LOAD]][0] {name = "scalarMember"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!s32i>271 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSCALARMEM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "this->scalarMember"}272 // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {273 // CHECK-NEXT: acc.yield274 // CHECK-NEXT: } loc275 // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[GETSCALARMEM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "this->scalarMember"}276#pragma acc kernels copy(this->arrayMember[2])277 ;278 // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i279 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32280 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64281 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64282 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64283 // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)284 // CHECK-NEXT: %[[GETARRAYMEM:.*]] = cir.get_member %[[THIS_LOAD]][1] {name = "arrayMember"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!cir.array<!s32i x 5>>285 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETARRAYMEM]] : !cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS1]]) -> !cir.ptr<!cir.array<!s32i x 5>> {dataClause = #acc<data_clause acc_copy>, name = "this->arrayMember[2]"}286 // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) {287 // CHECK-NEXT: acc.terminator288 // CHECK-NEXT: } loc289 // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS1]]) to varPtr(%[[GETARRAYMEM]] : !cir.ptr<!cir.array<!s32i x 5>>) {dataClause = #acc<data_clause acc_copy>, name = "this->arrayMember[2]"}290 291#pragma acc kernels copy(this->twoDArrayMember[1][2])292 ;293 // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i294 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32295 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64296 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64297 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64298 // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)299 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i300 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32301 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64302 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64303 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64304 // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)305 // CHECK-NEXT: %[[GET2DARRAYMEM:.*]] = cir.get_member %[[THIS_LOAD]][2] {name = "twoDArrayMember"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>306 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GET2DARRAYMEM]] : !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 = "this->twoDArrayMember[1][2]"}307 // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>) {308 // CHECK-NEXT: acc.terminator309 // CHECK-NEXT: } loc310 // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) to varPtr(%[[GET2DARRAYMEM]] : !cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>) {dataClause = #acc<data_clause acc_copy>, name = "this->twoDArrayMember[1][2]"}311 312#pragma acc kernels copy(this->iSTy)313 ;314 // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!rec_InnerStructTy>315 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSTRUCTMEM]] : !cir.ptr<!rec_InnerStructTy>) -> !cir.ptr<!rec_InnerStructTy> {dataClause = #acc<data_clause acc_copy>, name = "this->iSTy"}316 // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr<!rec_InnerStructTy>) {317 // CHECK-NEXT: acc.terminator318 // CHECK-NEXT: } loc319 // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!rec_InnerStructTy>) to varPtr(%[[GETSTRUCTMEM]] : !cir.ptr<!rec_InnerStructTy>) {dataClause = #acc<data_clause acc_copy>, name = "this->iSTy"}320 321#pragma acc parallel copy(this->iSTy.Member)322 ;323 // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!rec_InnerStructTy>324 // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr<!rec_InnerStructTy> -> !cir.ptr<!cir.array<!s32i x 5>>325 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr<!cir.array<!s32i x 5>>) -> !cir.ptr<!cir.array<!s32i x 5>> {dataClause = #acc<data_clause acc_copy>, name = "this->iSTy.Member"}326 // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) {327 // CHECK-NEXT: acc.yield328 // CHECK-NEXT: } loc329 // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) to varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr<!cir.array<!s32i x 5>>) {dataClause = #acc<data_clause acc_copy>, name = "this->iSTy.Member"}330 331#pragma acc serial copy(this->iSTy.Member[1])332 ;333 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i334 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32335 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64336 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64337 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64338 // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)339 // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!rec_InnerStructTy>340 // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr<!rec_InnerStructTy> -> !cir.ptr<!cir.array<!s32i x 5>>341 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS1]]) -> !cir.ptr<!cir.array<!s32i x 5>> {dataClause = #acc<data_clause acc_copy>, name = "this->iSTy.Member[1]"}342 // CHECK-NEXT: acc.serial dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) {343 // CHECK-NEXT: acc.yield344 // CHECK-NEXT: } loc345 // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS1]]) to varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr<!cir.array<!s32i x 5>>) {dataClause = #acc<data_clause acc_copy>, name = "this->iSTy.Member[1]"}346}347