419 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 loop copy(scalarMember)19 for(int i = 0; i < 5; ++i);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 combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {23 // CHECK-NEXT: acc.loop combined(parallel) {24 // CHECK: acc.yield25 // CHECK-NEXT: }26 // CHECK-NEXT: acc.yield27 // CHECK-NEXT: } loc28 // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[GETSCALARMEM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "this->scalarMember"}29 30#pragma acc kernels loop copy(arrayMember[2])31 for(int i = 0; i < 5; ++i);32 // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i33 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si3234 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i6435 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i6436 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i6437 // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)38 // CHECK-NEXT: %[[GETARRAYMEM:.*]] = cir.get_member %[[THIS_LOAD]][1] {name = "arrayMember"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!cir.array<!s32i x 5>>39 // 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]"}40 // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) {41 // CHECK-NEXT: acc.loop combined(kernels) {42 // CHECK: acc.yield43 // CHECK-NEXT: }44 // CHECK-NEXT: acc.terminator45 // CHECK-NEXT: } loc46 // 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]"}47 48#pragma acc kernels loop copy(twoDArrayMember[1][2])49 for(int i = 0; i < 5; ++i);50 // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i51 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !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: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)56 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i57 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si3258 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i6459 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i6460 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i6461 // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)62 // 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>>63 // 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]"}64 // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>) {65 // CHECK-NEXT: acc.loop combined(kernels) {66 // CHECK: acc.yield67 // CHECK-NEXT: }68 // CHECK-NEXT: acc.terminator69 // CHECK-NEXT: } loc70 // 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]"}71 72#pragma acc kernels loop copy(iSTy)73 for(int i = 0; i < 5; ++i);74 // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!rec_InnerStructTy>75 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSTRUCTMEM]] : !cir.ptr<!rec_InnerStructTy>) -> !cir.ptr<!rec_InnerStructTy> {dataClause = #acc<data_clause acc_copy>, name = "this->iSTy"}76 // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!rec_InnerStructTy>) {77 // CHECK-NEXT: acc.loop combined(kernels) {78 // CHECK: acc.yield79 // CHECK-NEXT: }80 // CHECK-NEXT: acc.terminator81 // CHECK-NEXT: } loc82 // 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"}83 84#pragma acc parallel loop copy(iSTy.Member)85 for(int i = 0; i < 5; ++i);86 // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!rec_InnerStructTy>87 // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr<!rec_InnerStructTy> -> !cir.ptr<!cir.array<!s32i x 5>>88 // 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"}89 // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) {90 // CHECK-NEXT: acc.loop combined(parallel) {91 // CHECK: acc.yield92 // CHECK-NEXT: }93 // CHECK-NEXT: acc.yield94 // CHECK-NEXT: } loc95 // 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"}96 97#pragma acc serial loop copy(iSTy.Member[1])98 for(int i = 0; i < 5; ++i);99 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i100 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32101 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64102 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64103 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64104 // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)105 // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!rec_InnerStructTy>106 // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr<!rec_InnerStructTy> -> !cir.ptr<!cir.array<!s32i x 5>>107 // 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]"}108 // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) {109 // CHECK-NEXT: acc.loop combined(serial) {110 // CHECK: acc.yield111 // CHECK-NEXT: }112 // CHECK-NEXT: acc.yield113 // CHECK-NEXT: } loc114 // 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]"}115 116#pragma acc parallel loop copy(this->scalarMember)117 for(int i = 0; i < 5; ++i);118 // CHECK-NEXT: %[[GETSCALARMEM:.*]] = cir.get_member %[[THIS_LOAD]][0] {name = "scalarMember"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!s32i>119 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSCALARMEM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "this->scalarMember"}120 // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {121 // CHECK-NEXT: acc.loop combined(parallel) {122 // CHECK: acc.yield123 // CHECK-NEXT: }124 // CHECK-NEXT: acc.yield125 // CHECK-NEXT: } loc126 // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[GETSCALARMEM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "this->scalarMember"}127 128#pragma acc kernels loop copy(this->arrayMember[2])129 for(int i = 0; i < 5; ++i);130 // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i131 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32132 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64133 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64134 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64135 // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)136 // CHECK-NEXT: %[[GETARRAYMEM:.*]] = cir.get_member %[[THIS_LOAD]][1] {name = "arrayMember"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!cir.array<!s32i x 5>>137 // 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]"}138 // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) {139 // CHECK-NEXT: acc.loop combined(kernels) {140 // CHECK: acc.yield141 // CHECK-NEXT: }142 // CHECK-NEXT: acc.terminator143 // CHECK-NEXT: } loc144 // 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]"}145#pragma acc kernels loop copy(this->twoDArrayMember[1][2])146 for(int i = 0; i < 5; ++i);147 // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i148 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32149 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64150 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64151 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64152 // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)153 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i154 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32155 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64156 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64157 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64158 // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)159 // 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>>160 // 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]"}161 // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>) {162 // CHECK-NEXT: acc.loop combined(kernels) {163 // CHECK: acc.yield164 // CHECK-NEXT: }165 // CHECK-NEXT: acc.terminator166 // CHECK-NEXT: } loc167 // 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]"}168 169#pragma acc kernels loop copy(this->iSTy)170 for(int i = 0; i < 5; ++i);171 // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!rec_InnerStructTy>172 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSTRUCTMEM]] : !cir.ptr<!rec_InnerStructTy>) -> !cir.ptr<!rec_InnerStructTy> {dataClause = #acc<data_clause acc_copy>, name = "this->iSTy"}173 // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!rec_InnerStructTy>) {174 // CHECK-NEXT: acc.loop combined(kernels) {175 // CHECK: acc.yield176 // CHECK-NEXT: }177 // CHECK-NEXT: acc.terminator178 // CHECK-NEXT: } loc179 // 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"}180 181#pragma acc parallel loop copy(this->iSTy.Member)182 for(int i = 0; i < 5; ++i);183 // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!rec_InnerStructTy>184 // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr<!rec_InnerStructTy> -> !cir.ptr<!cir.array<!s32i x 5>>185 // 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"}186 // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) {187 // CHECK-NEXT: acc.loop combined(parallel) {188 // CHECK: acc.yield189 // CHECK-NEXT: }190 // CHECK-NEXT: acc.yield191 // CHECK-NEXT: } loc192 // 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"}193 194#pragma acc serial loop copy(this->iSTy.Member[1])195 for(int i = 0; i < 5; ++i);196 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i197 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32198 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64199 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64200 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64201 // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)202 // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!rec_InnerStructTy>203 // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr<!rec_InnerStructTy> -> !cir.ptr<!cir.array<!s32i x 5>>204 // 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]"}205 // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) {206 // CHECK-NEXT: acc.loop combined(serial) {207 // CHECK: acc.yield208 // CHECK-NEXT: }209 // CHECK-NEXT: acc.yield210 // CHECK-NEXT: } loc211 // 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]"}212}213 214void OutlineFunc();215};216 217void InlineUse() {218 StructTy s;219 s.InlineFunc();220}221 222void StructTy::OutlineFunc() {223 // CHECK: cir.func {{.*}}OutlineFunc{{.*}}224 // CHECK-NEXT: %[[THIS:.*]] = cir.alloca !cir.ptr<!rec_StructTy>, !cir.ptr<!cir.ptr<!rec_StructTy>>, ["this", init]225 // CHECK-NEXT: cir.store %[[THIS_ARG:.*]], %[[THIS]] : !cir.ptr<!rec_StructTy>, !cir.ptr<!cir.ptr<!rec_StructTy>>226 // CHECK-NEXT: %[[THIS_LOAD:.*]] = cir.load %[[THIS]] : !cir.ptr<!cir.ptr<!rec_StructTy>>, !cir.ptr<!rec_StructTy>227#pragma acc parallel loop copy(scalarMember)228 for(int i = 0; i < 5; ++i);229 // CHECK-NEXT: %[[GETSCALARMEM:.*]] = cir.get_member %[[THIS_LOAD]][0] {name = "scalarMember"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!s32i>230 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSCALARMEM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "this->scalarMember"}231 // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {232 // CHECK-NEXT: acc.loop combined(parallel) {233 // CHECK: acc.yield234 // CHECK-NEXT: }235 // CHECK-NEXT: acc.yield236 // CHECK-NEXT: } loc237 // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[GETSCALARMEM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "this->scalarMember"}238#pragma acc kernels loop copy(arrayMember[2])239 for(int i = 0; i < 5; ++i);240 // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i241 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32242 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64243 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64244 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64245 // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)246 // CHECK-NEXT: %[[GETARRAYMEM:.*]] = cir.get_member %[[THIS_LOAD]][1] {name = "arrayMember"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!cir.array<!s32i x 5>>247 // 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]"}248 // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) {249 // CHECK-NEXT: acc.loop combined(kernels) {250 // CHECK: acc.yield251 // CHECK-NEXT: }252 // CHECK-NEXT: acc.terminator253 // CHECK-NEXT: } loc254 // 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]"}255#pragma acc kernels loop copy(twoDArrayMember[1][2])256 for(int i = 0; i < 5; ++i);257 // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i258 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32259 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64260 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64261 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64262 // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)263 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i264 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32265 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64266 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64267 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64268 // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)269 // 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>>270 // 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]"}271 // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>) {272 // CHECK-NEXT: acc.loop combined(kernels) {273 // CHECK: acc.yield274 // CHECK-NEXT: }275 // CHECK-NEXT: acc.terminator276 // CHECK-NEXT: } loc277 // 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]"}278#pragma acc kernels loop copy(iSTy)279 for(int i = 0; i < 5; ++i);280 // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!rec_InnerStructTy>281 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSTRUCTMEM]] : !cir.ptr<!rec_InnerStructTy>) -> !cir.ptr<!rec_InnerStructTy> {dataClause = #acc<data_clause acc_copy>, name = "this->iSTy"}282 // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!rec_InnerStructTy>) {283 // CHECK-NEXT: acc.loop combined(kernels) {284 // CHECK: acc.yield285 // CHECK-NEXT: }286 // CHECK-NEXT: acc.terminator287 // CHECK-NEXT: } loc288 // 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"}289 290#pragma acc parallel loop copy(iSTy.Member)291 for(int i = 0; i < 5; ++i);292 // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!rec_InnerStructTy>293 // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr<!rec_InnerStructTy> -> !cir.ptr<!cir.array<!s32i x 5>>294 // 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"}295 // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) {296 // CHECK-NEXT: acc.loop combined(parallel) {297 // CHECK: acc.yield298 // CHECK-NEXT: }299 // CHECK-NEXT: acc.yield300 // CHECK-NEXT: } loc301 // 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"}302 303#pragma acc serial loop copy(iSTy.Member[1])304 for(int i = 0; i < 5; ++i);305 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i306 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32307 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64308 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64309 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64310 // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)311 // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!rec_InnerStructTy>312 // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr<!rec_InnerStructTy> -> !cir.ptr<!cir.array<!s32i x 5>>313 // 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]"}314 // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) {315 // CHECK-NEXT: acc.loop combined(serial) {316 // CHECK: acc.yield317 // CHECK-NEXT: }318 // CHECK-NEXT: acc.yield319 // CHECK-NEXT: } loc320 // 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]"}321 322#pragma acc parallel loop copy(this->scalarMember)323 for(int i = 0; i < 5; ++i);324 // CHECK-NEXT: %[[GETSCALARMEM:.*]] = cir.get_member %[[THIS_LOAD]][0] {name = "scalarMember"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!s32i>325 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSCALARMEM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "this->scalarMember"}326 // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {327 // CHECK-NEXT: acc.loop combined(parallel) {328 // CHECK: acc.yield329 // CHECK-NEXT: }330 // CHECK-NEXT: acc.yield331 // CHECK-NEXT: } loc332 // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[GETSCALARMEM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "this->scalarMember"}333#pragma acc kernels loop copy(this->arrayMember[2])334 for(int i = 0; i < 5; ++i);335 // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i336 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32337 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64338 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64339 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64340 // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)341 // CHECK-NEXT: %[[GETARRAYMEM:.*]] = cir.get_member %[[THIS_LOAD]][1] {name = "arrayMember"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!cir.array<!s32i x 5>>342 // 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]"}343 // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) {344 // CHECK-NEXT: acc.loop combined(kernels) {345 // CHECK: acc.yield346 // CHECK-NEXT: }347 // CHECK-NEXT: acc.terminator348 // CHECK-NEXT: } loc349 // 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]"}350 351#pragma acc kernels loop copy(this->twoDArrayMember[1][2])352 for(int i = 0; i < 5; ++i);353 // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i354 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32355 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64356 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64357 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64358 // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)359 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i360 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32361 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64362 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64363 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64364 // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)365 // 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>>366 // 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]"}367 // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!cir.array<!s16i x 3> x 5>>) {368 // CHECK-NEXT: acc.loop combined(kernels) {369 // CHECK: acc.yield370 // CHECK-NEXT: }371 // CHECK-NEXT: acc.terminator372 // CHECK-NEXT: } loc373 // 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]"}374 375#pragma acc kernels loop copy(this->iSTy)376 for(int i = 0; i < 5; ++i);377 // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!rec_InnerStructTy>378 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSTRUCTMEM]] : !cir.ptr<!rec_InnerStructTy>) -> !cir.ptr<!rec_InnerStructTy> {dataClause = #acc<data_clause acc_copy>, name = "this->iSTy"}379 // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!rec_InnerStructTy>) {380 // CHECK-NEXT: acc.loop combined(kernels) {381 // CHECK: acc.yield382 // CHECK-NEXT: }383 // CHECK-NEXT: acc.terminator384 // CHECK-NEXT: } loc385 // 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"}386 387#pragma acc parallel loop copy(this->iSTy.Member)388 for(int i = 0; i < 5; ++i);389 // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!rec_InnerStructTy>390 // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr<!rec_InnerStructTy> -> !cir.ptr<!cir.array<!s32i x 5>>391 // 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"}392 // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) {393 // CHECK-NEXT: acc.loop combined(parallel) {394 // CHECK: acc.yield395 // CHECK-NEXT: }396 // CHECK-NEXT: acc.yield397 // CHECK-NEXT: } loc398 // 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"}399 400#pragma acc serial loop copy(this->iSTy.Member[1])401 for(int i = 0; i < 5; ++i);402 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i403 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32404 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64405 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64406 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64407 // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)408 // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr<!rec_StructTy> -> !cir.ptr<!rec_InnerStructTy>409 // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr<!rec_InnerStructTy> -> !cir.ptr<!cir.array<!s32i x 5>>410 // 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]"}411 // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!cir.array<!s32i x 5>>) {412 // CHECK-NEXT: acc.loop combined(serial) {413 // CHECK: acc.yield414 // CHECK-NEXT: }415 // CHECK-NEXT: acc.yield416 // CHECK-NEXT: } loc417 // 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]"}418}419