569 lines · cpp
1// RUN: %clang_cc1 -fopenacc -triple x86_64-linux-gnu -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir -triple x86_64-linux-pc %s -o - | FileCheck %s2 3struct NoCopyConstruct {};4 5struct CopyConstruct {6 CopyConstruct() = default;7 CopyConstruct(const CopyConstruct&);8};9 10struct NonDefaultCtor {11 NonDefaultCtor();12};13 14struct HasDtor {15 ~HasDtor();16};17 18// Int19// CHECK: acc.private.recipe @privatization__ZTSi : !cir.ptr<!s32i> init {20// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!s32i> {{.*}}):21// CHECK-NEXT: cir.alloca !s32i, !cir.ptr<!s32i>, ["openacc.private.init"]22// CHECK-NEXT: acc.yield23// CHECK-NEXT: }24//25// Float26// CHECK-NEXT: acc.private.recipe @privatization__ZTSf : !cir.ptr<!cir.float> init {27// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.float> {{.*}}):28// CHECK-NEXT: cir.alloca !cir.float, !cir.ptr<!cir.float>, ["openacc.private.init"]29// CHECK-NEXT: acc.yield30// CHECK-NEXT: }31//32// NoCopyConstruct33// CHECK-NEXT: acc.private.recipe @privatization__ZTS15NoCopyConstruct : !cir.ptr<!rec_NoCopyConstruct> init {34// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_NoCopyConstruct> {{.*}}):35// CHECK-NEXT: cir.alloca !rec_NoCopyConstruct, !cir.ptr<!rec_NoCopyConstruct>, ["openacc.private.init"]36// CHECK-NEXT: acc.yield37// CHECK-NEXT: }38//39// CopyConstruct40// CHECK-NEXT: acc.private.recipe @privatization__ZTS13CopyConstruct : !cir.ptr<!rec_CopyConstruct> init {41// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_CopyConstruct> {{.*}}):42// CHECK-NEXT: cir.alloca !rec_CopyConstruct, !cir.ptr<!rec_CopyConstruct>, ["openacc.private.init"]43// CHECK-NEXT: acc.yield44// CHECK-NEXT: }45//46// NonDefaultCtor47// CHECK-NEXT: acc.private.recipe @privatization__ZTS14NonDefaultCtor : !cir.ptr<!rec_NonDefaultCtor> init {48// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_NonDefaultCtor> {{.*}}):49// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !rec_NonDefaultCtor, !cir.ptr<!rec_NonDefaultCtor>, ["openacc.private.init", init]50// CHECK-NEXT: cir.call @_ZN14NonDefaultCtorC1Ev(%[[ALLOCA]]) : (!cir.ptr<!rec_NonDefaultCtor>) -> ()51// CHECK-NEXT: acc.yield52// CHECK-NEXT: }53//54// HasDtor55// CHECK-NEXT: acc.private.recipe @privatization__ZTS7HasDtor : !cir.ptr<!rec_HasDtor> init {56// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_HasDtor> {{.*}}):57// CHECK-NEXT: cir.alloca !rec_HasDtor, !cir.ptr<!rec_HasDtor>, ["openacc.private.init"]58// CHECK-NEXT: acc.yield59// CHECK-NEXT: } destroy {60// CHECK-NEXT: ^bb0(%[[ORIG:.*]]: !cir.ptr<!rec_HasDtor> {{.*}}, %[[ARG:.*]]: !cir.ptr<!rec_HasDtor> {{.*}}):61// CHECK-NEXT: cir.call @_ZN7HasDtorD1Ev(%[[ARG]]) nothrow : (!cir.ptr<!rec_HasDtor>) -> ()62// CHECK-NEXT: acc.yield63// CHECK-NEXT: }64//65// int[5] with 1 'bound'66// CHECK-NEXT: acc.private.recipe @privatization__Bcnt1__ZTSA5_i : !cir.ptr<!cir.array<!s32i x 5>> init {67// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!s32i x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}):68// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !cir.array<!s32i x 5>, !cir.ptr<!cir.array<!s32i x 5>>, ["openacc.private.init"]69// CHECK-NEXT: acc.yield70// CHECK-NEXT: }71//72// float[5] with 1 'bound'73// CHECK-NEXT: acc.private.recipe @privatization__Bcnt1__ZTSA5_f : !cir.ptr<!cir.array<!cir.float x 5>> init {74// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!cir.float x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}):75// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !cir.array<!cir.float x 5>, !cir.ptr<!cir.array<!cir.float x 5>>, ["openacc.private.init"]76// CHECK-NEXT: acc.yield77// CHECK-NEXT: }78//79// NoCopyConstruct[5] with 1 'bound'80// CHECK-NEXT: acc.private.recipe @privatization__Bcnt1__ZTSA5_15NoCopyConstruct : !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> init {81// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}):82// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !cir.array<!rec_NoCopyConstruct x 5>, !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>, ["openacc.private.init"]83// CHECK-NEXT: acc.yield84// CHECK-NEXT: }85//86// CopyConstruct[5] with 1 'bound'87// CHECK-NEXT: acc.private.recipe @privatization__Bcnt1__ZTSA5_13CopyConstruct : !cir.ptr<!cir.array<!rec_CopyConstruct x 5>> init {88// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!rec_CopyConstruct x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}):89// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !cir.array<!rec_CopyConstruct x 5>, !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>, ["openacc.private.init"]90// CHECK-NEXT: acc.yield91// CHECK-NEXT: }92//93// NonDefaultCtor[5] with 1 'bound'94// CHECK-NEXT: acc.private.recipe @privatization__Bcnt1__ZTSA5_14NonDefaultCtor : !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> init {95// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}):96// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !cir.array<!rec_NonDefaultCtor x 5>, !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>, ["openacc.private.init"]97// CHECK-NEXT: cir.scope {98// CHECK-NEXT: %[[LB:.*]] = acc.get_lowerbound %[[BOUND1]] : (!acc.data_bounds_ty) -> index99// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]] : index to !u64i100// CHECK-NEXT: %[[UB:.*]] = acc.get_upperbound %[[BOUND1]] : (!acc.data_bounds_ty) -> index101// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]] : index to !u64i102// CHECK-NEXT: %[[ITR:.*]] = cir.alloca !u64i, !cir.ptr<!u64i>, ["iter"] {alignment = 8 : i64}103// CHECK-NEXT: cir.store %[[LB_CAST]], %[[ITR]] : !u64i, !cir.ptr<!u64i>104// CHECK-NEXT: cir.for : cond {105// CHECK-NEXT: %[[ITR_LOAD:.*]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i106// CHECK-NEXT: %[[COND:.*]] = cir.cmp(lt, %[[ITR_LOAD]], %[[UB_CAST]]) : !u64i, !cir.bool107// CHECK-NEXT: cir.condition(%[[COND]])108// CHECK-NEXT: } body {109// CHECK-NEXT: %[[ITR_LOAD:.*]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i110// CHECK-NEXT: %[[DECAY:.*]] = cir.cast array_to_ptrdecay %[[ALLOCA]] : !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> -> !cir.ptr<!rec_NonDefaultCtor>111// CHECK-NEXT: %[[STRIDE:.*]] = cir.ptr_stride %[[DECAY]], %[[ITR_LOAD]] : (!cir.ptr<!rec_NonDefaultCtor>, !u64i) -> !cir.ptr<!rec_NonDefaultCtor>112// CHECK-NEXT: cir.call @_ZN14NonDefaultCtorC1Ev(%[[STRIDE]]) : (!cir.ptr<!rec_NonDefaultCtor>) -> ()113// CHECK-NEXT: cir.yield114// CHECK-NEXT: } step {115// CHECK-NEXT: %[[ITR_LOAD]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i116// CHECK-NEXT: %[[INC:.*]] = cir.unary(inc, %[[ITR_LOAD]]) : !u64i, !u64i117// CHECK-NEXT: cir.store %[[INC]], %[[ITR]] : !u64i, !cir.ptr<!u64i>118// CHECK-NEXT: cir.yield119// CHECK-NEXT: }120// CHECK-NEXT: }121// CHECK-NEXT: acc.yield122// CHECK-NEXT: }123//124// HasDtor[5] with 1 'bound'125// CHECK-NEXT: acc.private.recipe @privatization__Bcnt1__ZTSA5_7HasDtor : !cir.ptr<!cir.array<!rec_HasDtor x 5>> init {126// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!rec_HasDtor x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}):127// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !cir.array<!rec_HasDtor x 5>, !cir.ptr<!cir.array<!rec_HasDtor x 5>>, ["openacc.private.init"]128// CHECK-NEXT: acc.yield129// CHECK-NEXT: } destroy {130// CHECK-NEXT: ^bb0(%[[ORIG:.*]]: !cir.ptr<!cir.array<!rec_HasDtor x 5>> {{.*}}, %[[ARG:.*]]: !cir.ptr<!cir.array<!rec_HasDtor x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}):131// CHECK-NEXT: cir.scope {132// CHECK-NEXT: %[[LB1:.*]] = acc.get_lowerbound %[[BOUND1]] : (!acc.data_bounds_ty) -> index133// CHECK-NEXT: %[[LB1_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB1]] : index to !u64i134// CHECK-NEXT: %[[UB1:.*]] = acc.get_upperbound %[[BOUND1]] : (!acc.data_bounds_ty) -> index135// CHECK-NEXT: %[[UB1_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB1]] : index to !u64i136// CHECK-NEXT: %[[ITR:.*]] = cir.alloca !u64i, !cir.ptr<!u64i>, ["iter"] {alignment = 8 : i64}137// CHECK-NEXT: %[[CONST_ONE:.*]] = cir.const #cir.int<1> : !u64i138// CHECK-NEXT: %[[ONE_BELOW_UB:.*]] = cir.binop(sub, %[[UB1_CAST]], %[[CONST_ONE]]) : !u64i139// CHECK-NEXT: cir.store %[[ONE_BELOW_UB]], %[[ITR]] : !u64i, !cir.ptr<!u64i>140// CHECK-NEXT: cir.for : cond {141// CHECK-NEXT: %[[ITR_LOAD:.*]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i142// CHECK-NEXT: %[[COND:.*]] = cir.cmp(ge, %[[ITR_LOAD]], %[[LB1_CAST]]) : !u64i, !cir.bool143// CHECK-NEXT: cir.condition(%[[COND]])144// CHECK-NEXT: } body {145// CHECK-NEXT: %[[ITR_LOAD:.*]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i146// CHECK-NEXT: %[[DECAY:.*]] = cir.cast array_to_ptrdecay %[[ARG]] : !cir.ptr<!cir.array<!rec_HasDtor x 5>> -> !cir.ptr<!rec_HasDtor>147// CHECK-NEXT: %[[STRIDE:.*]] = cir.ptr_stride %[[DECAY]], %[[ITR_LOAD]] : (!cir.ptr<!rec_HasDtor>, !u64i) -> !cir.ptr<!rec_HasDtor>148// CHECK-NEXT: cir.call @_ZN7HasDtorD1Ev(%[[STRIDE]]) nothrow : (!cir.ptr<!rec_HasDtor>) -> ()149// CHECK-NEXT: cir.yield150// CHECK-NEXT: } step {151// CHECK-NEXT: %[[ITR_LOAD]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i152// CHECK-NEXT: %[[DEC:.*]] = cir.unary(dec, %[[ITR_LOAD]]) : !u64i, !u64i153// CHECK-NEXT: cir.store %[[DEC]], %[[ITR]] : !u64i, !cir.ptr<!u64i>154// CHECK-NEXT: cir.yield155// CHECK-NEXT: }156// CHECK-NEXT: }157// CHECK-NEXT: acc.yield158// CHECK-NEXT: }159 160extern "C" void acc_combined() {161 // CHECK: cir.func{{.*}} @acc_combined()162 163 int someInt;164 // CHECK-NEXT: %[[SOMEINT:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["someInt"]165 float someFloat;166 // CHECK-NEXT: %[[SOMEFLOAT:.*]] = cir.alloca !cir.float, !cir.ptr<!cir.float>, ["someFloat"]167 NoCopyConstruct noCopy;168 // CHECK-NEXT: %[[NOCOPY:.*]] = cir.alloca !rec_NoCopyConstruct, !cir.ptr<!rec_NoCopyConstruct>, ["noCopy"]169 CopyConstruct hasCopy;170 // CHECK-NEXT: %[[HASCOPY:.*]] = cir.alloca !rec_CopyConstruct, !cir.ptr<!rec_CopyConstruct>, ["hasCopy"]171 NonDefaultCtor notDefCtor;172 // CHECK-NEXT: %[[NOTDEFCTOR:.*]] = cir.alloca !rec_NonDefaultCtor, !cir.ptr<!rec_NonDefaultCtor>, ["notDefCtor", init]173 HasDtor dtor;174 // CHECK-NEXT: %[[DTOR:.*]] = cir.alloca !rec_HasDtor, !cir.ptr<!rec_HasDtor>, ["dtor"]175 int someIntArr[5];176 // CHECK-NEXT: %[[INTARR:.*]] = cir.alloca !cir.array<!s32i x 5>, !cir.ptr<!cir.array<!s32i x 5>>, ["someIntArr"]177 float someFloatArr[5];178 // CHECK-NEXT: %[[FLOATARR:.*]] = cir.alloca !cir.array<!cir.float x 5>, !cir.ptr<!cir.array<!cir.float x 5>>, ["someFloatArr"]179 NoCopyConstruct noCopyArr[5];180 // CHECK-NEXT: %[[NOCOPYARR:.*]] = cir.alloca !cir.array<!rec_NoCopyConstruct x 5>, !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>, ["noCopyArr"]181 CopyConstruct hasCopyArr[5];182 // CHECK-NEXT: %[[HASCOPYARR:.*]] = cir.alloca !cir.array<!rec_CopyConstruct x 5>, !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>, ["hasCopyArr"]183 NonDefaultCtor notDefCtorArr[5];184 // CHECK-NEXT: %[[NOTDEFCTORARR:.*]] = cir.alloca !cir.array<!rec_NonDefaultCtor x 5>, !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>, ["notDefCtorArr", init]185 HasDtor dtorArr[5];186 // CHECK-NEXT: %[[DTORARR:.*]] = cir.alloca !cir.array<!rec_HasDtor x 5>, !cir.ptr<!cir.array<!rec_HasDtor x 5>>, ["dtorArr"]187 // CHECK-NEXT: cir.call @_ZN14NonDefaultCtorC1Ev(%[[NOTDEFCTOR]]) : (!cir.ptr<!rec_NonDefaultCtor>) -> ()188 189#pragma acc parallel loop private(someInt)190 for(int i = 0; i < 5; ++i);191 // CHECK: acc.parallel combined(loop) {192 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[SOMEINT]] : !cir.ptr<!s32i>) recipe(@privatization__ZTSi) -> !cir.ptr<!s32i> {name = "someInt"}193 // CHECK-NEXT: acc.loop combined(parallel) private(%[[PRIVATE]] : !cir.ptr<!s32i>)194 // CHECK: acc.yield195 // CHECK-NEXT: } loc196 // CHECK-NEXT: acc.yield197 // CHECK-NEXT: } loc198#pragma acc serial loop private(someFloat)199 for(int i = 0; i < 5; ++i);200 // CHECK-NEXT: acc.serial combined(loop) {201 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[SOMEFLOAT]] : !cir.ptr<!cir.float>) recipe(@privatization__ZTSf) -> !cir.ptr<!cir.float> {name = "someFloat"}202 // CHECK-NEXT: acc.loop combined(serial) private(%[[PRIVATE]] : !cir.ptr<!cir.float>)203 // CHECK: acc.yield204 // CHECK-NEXT: } loc205 // CHECK-NEXT: acc.yield206 // CHECK-NEXT: } loc207 208#pragma acc parallel loop private(noCopy)209 for(int i = 0; i < 5; ++i);210 // CHECK-NEXT: acc.parallel combined(loop) {211 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[NOCOPY]] : !cir.ptr<!rec_NoCopyConstruct>) recipe(@privatization__ZTS15NoCopyConstruct) -> !cir.ptr<!rec_NoCopyConstruct> {name = "noCopy"}212 // CHECK-NEXT: acc.loop combined(parallel) private(%[[PRIVATE]] : !cir.ptr<!rec_NoCopyConstruct>213 // CHECK: acc.yield214 // CHECK-NEXT: } loc215 // CHECK-NEXT: acc.yield216 // CHECK-NEXT: } loc217#pragma acc serial loop private(hasCopy)218 for(int i = 0; i < 5; ++i);219 // CHECK-NEXT: acc.serial combined(loop) {220 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[HASCOPY]] : !cir.ptr<!rec_CopyConstruct>) recipe(@privatization__ZTS13CopyConstruct) -> !cir.ptr<!rec_CopyConstruct> {name = "hasCopy"}221 // CHECK-NEXT: acc.loop combined(serial) private(%[[PRIVATE]] : !cir.ptr<!rec_CopyConstruct>)222 // CHECK: acc.yield223 // CHECK-NEXT: } loc224 // CHECK-NEXT: acc.yield225 // CHECK-NEXT: } loc226#pragma acc serial loop private(notDefCtor)227 for(int i = 0; i < 5; ++i);228 // CHECK-NEXT: acc.serial combined(loop) {229 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[NOTDEFCTOR]] : !cir.ptr<!rec_NonDefaultCtor>) recipe(@privatization__ZTS14NonDefaultCtor) -> !cir.ptr<!rec_NonDefaultCtor> {name = "notDefCtor"}230 // CHECK-NEXT: acc.loop combined(serial) private(%[[PRIVATE]] : !cir.ptr<!rec_NonDefaultCtor>)231 // CHECK: acc.yield232 // CHECK-NEXT: } loc233 // CHECK-NEXT: acc.yield234 // CHECK-NEXT: } loc235#pragma acc serial loop private(dtor)236 for(int i = 0; i < 5; ++i);237 // CHECK-NEXT: acc.serial combined(loop) {238 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[DTOR]] : !cir.ptr<!rec_HasDtor>) recipe(@privatization__ZTS7HasDtor) -> !cir.ptr<!rec_HasDtor> {name = "dtor"}239 // CHECK-NEXT: acc.loop combined(serial) private(%[[PRIVATE]] : !cir.ptr<!rec_HasDtor>)240 // CHECK: acc.yield241 // CHECK-NEXT: } loc242 // CHECK-NEXT: acc.yield243 // CHECK-NEXT: } loc244 245#pragma acc parallel loop private(someInt, someFloat, noCopy, hasCopy, notDefCtor, dtor)246 for(int i = 0; i < 5; ++i);247 // CHECK: acc.parallel combined(loop) {248 // CHECK: %[[PRIVATE1:.*]] = acc.private varPtr(%[[SOMEINT]] : !cir.ptr<!s32i>) recipe(@privatization__ZTSi) -> !cir.ptr<!s32i> {name = "someInt"}249 // CHECK-NEXT: %[[PRIVATE2:.*]] = acc.private varPtr(%[[SOMEFLOAT]] : !cir.ptr<!cir.float>) recipe(@privatization__ZTSf) -> !cir.ptr<!cir.float> {name = "someFloat"}250 // CHECK-NEXT: %[[PRIVATE3:.*]] = acc.private varPtr(%[[NOCOPY]] : !cir.ptr<!rec_NoCopyConstruct>) recipe(@privatization__ZTS15NoCopyConstruct) -> !cir.ptr<!rec_NoCopyConstruct> {name = "noCopy"}251 // CHECK-NEXT: %[[PRIVATE4:.*]] = acc.private varPtr(%[[HASCOPY]] : !cir.ptr<!rec_CopyConstruct>) recipe(@privatization__ZTS13CopyConstruct) -> !cir.ptr<!rec_CopyConstruct> {name = "hasCopy"}252 // CHECK-NEXT: %[[PRIVATE5:.*]] = acc.private varPtr(%[[NOTDEFCTOR]] : !cir.ptr<!rec_NonDefaultCtor>) recipe(@privatization__ZTS14NonDefaultCtor) -> !cir.ptr<!rec_NonDefaultCtor> {name = "notDefCtor"}253 // CHECK-NEXT: %[[PRIVATE6:.*]] = acc.private varPtr(%[[DTOR]] : !cir.ptr<!rec_HasDtor>) recipe(@privatization__ZTS7HasDtor) -> !cir.ptr<!rec_HasDtor> {name = "dtor"}254 // CHECK-NEXT: acc.loop combined(parallel) private(%[[PRIVATE1]], %[[PRIVATE2]], %[[PRIVATE3]], %[[PRIVATE4]], %[[PRIVATE5]], %[[PRIVATE6]] : !cir.ptr<!s32i>,255 // CHECK-SAME: !cir.ptr<!cir.float>,256 // CHECK-SAME: !cir.ptr<!rec_NoCopyConstruct>,257 // CHECK-SAME: !cir.ptr<!rec_CopyConstruct>,258 // CHECK-SAME: !cir.ptr<!rec_NonDefaultCtor>,259 // CHECK-SAME: !cir.ptr<!rec_HasDtor>)260 // CHECK: acc.yield261 // CHECK-NEXT: } loc262 // CHECK-NEXT: acc.yield263 // CHECK-NEXT: } loc264 265#pragma acc serial loop private(someIntArr[1])266 for(int i = 0; i < 5; ++i);267 // CHECK-NEXT: acc.serial combined(loop) {268 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i269 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32270 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1271 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0272 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1273 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)274 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[INTARR]] : !cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS]]) recipe(@privatization__Bcnt1__ZTSA5_i) -> !cir.ptr<!cir.array<!s32i x 5>> {name = "someIntArr[1]"}275 // CHECK-NEXT: acc.loop combined(serial) private(%[[PRIVATE]] : !cir.ptr<!cir.array<!s32i x 5>>)276 // CHECK: acc.yield277 // CHECK-NEXT: } loc278 // CHECK-NEXT: acc.yield279 // CHECK-NEXT: } loc280#pragma acc parallel loop private(someFloatArr[1])281 for(int i = 0; i < 5; ++i);282 // CHECK-NEXT: acc.parallel combined(loop) {283 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i284 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32285 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1286 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0287 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1288 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)289 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[FLOATARR]] : !cir.ptr<!cir.array<!cir.float x 5>>) bounds(%[[BOUNDS]]) recipe(@privatization__Bcnt1__ZTSA5_f) -> !cir.ptr<!cir.array<!cir.float x 5>> {name = "someFloatArr[1]"}290 // CHECK-NEXT: acc.loop combined(parallel) private(%[[PRIVATE]] : !cir.ptr<!cir.array<!cir.float x 5>>)291 // CHECK: acc.yield292 // CHECK-NEXT: } loc293 // CHECK-NEXT: acc.yield294 // CHECK-NEXT: } loc295#pragma acc serial loop private(noCopyArr[1])296 for(int i = 0; i < 5; ++i);297 // CHECK-NEXT: acc.serial combined(loop) {298 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i299 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32300 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1301 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0302 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1303 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)304 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[NOCOPYARR]] : !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>) bounds(%[[BOUNDS]]) recipe(@privatization__Bcnt1__ZTSA5_15NoCopyConstruct) -> !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> {name = "noCopyArr[1]"}305 // CHECK-NEXT: acc.loop combined(serial) private(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>)306 // CHECK: acc.yield307 // CHECK-NEXT: } loc308 // CHECK-NEXT: acc.yield309 // CHECK-NEXT: } loc310#pragma acc parallel loop private(hasCopyArr[1])311 for(int i = 0; i < 5; ++i);312 // CHECK-NEXT: acc.parallel combined(loop) {313 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i314 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32315 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1316 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0317 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1318 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)319 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[HASCOPYARR]] : !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>) bounds(%[[BOUNDS]]) recipe(@privatization__Bcnt1__ZTSA5_13CopyConstruct) -> !cir.ptr<!cir.array<!rec_CopyConstruct x 5>> {name = "hasCopyArr[1]"}320 // CHECK-NEXT: acc.loop combined(parallel) private(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>)321 // CHECK: acc.yield322 // CHECK-NEXT: } loc323 // CHECK-NEXT: acc.yield324 // CHECK-NEXT: } loc325#pragma acc parallel loop private(notDefCtorArr[1])326 for(int i = 0; i < 5; ++i);327 // CHECK-NEXT: acc.parallel combined(loop) {328 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i329 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32330 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1331 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0332 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1333 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)334 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[NOTDEFCTORARR]] : !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>) bounds(%[[BOUNDS]]) recipe(@privatization__Bcnt1__ZTSA5_14NonDefaultCtor) -> !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> {name = "notDefCtorArr[1]"}335 // CHECK-NEXT: acc.loop combined(parallel) private(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>)336 // CHECK: acc.yield337 // CHECK-NEXT: } loc338 // CHECK-NEXT: acc.yield339 // CHECK-NEXT: } loc340#pragma acc parallel loop private(dtorArr[1])341 for(int i = 0; i < 5; ++i);342 // CHECK-NEXT: acc.parallel combined(loop) {343 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i344 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32345 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1346 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0347 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1348 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)349 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[DTORARR]] : !cir.ptr<!cir.array<!rec_HasDtor x 5>>) bounds(%[[BOUNDS]]) recipe(@privatization__Bcnt1__ZTSA5_7HasDtor) -> !cir.ptr<!cir.array<!rec_HasDtor x 5>> {name = "dtorArr[1]"}350 // CHECK-NEXT: acc.loop combined(parallel) private(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_HasDtor x 5>>)351 // CHECK: acc.yield352 // CHECK-NEXT: } loc353 // CHECK-NEXT: acc.yield354 // CHECK-NEXT: } loc355#pragma acc serial loop private(someIntArr[1], someFloatArr[1], noCopyArr[1], hasCopyArr[1], notDefCtorArr[1], dtorArr[1])356 for(int i = 0; i < 5; ++i);357 // CHECK-NEXT: acc.serial combined(loop) {358 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i359 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32360 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1361 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0362 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1363 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)364 // CHECK-NEXT: %[[PRIVATE1:.*]] = acc.private varPtr(%[[INTARR]] : !cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS]]) recipe(@privatization__Bcnt1__ZTSA5_i) -> !cir.ptr<!cir.array<!s32i x 5>> {name = "someIntArr[1]"}365 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i366 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32367 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1368 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0369 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1370 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)371 // CHECK-NEXT: %[[PRIVATE2:.*]] = acc.private varPtr(%[[FLOATARR]] : !cir.ptr<!cir.array<!cir.float x 5>>) bounds(%[[BOUNDS]]) recipe(@privatization__Bcnt1__ZTSA5_f) -> !cir.ptr<!cir.array<!cir.float x 5>> {name = "someFloatArr[1]"}372 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i373 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32374 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1375 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0376 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1377 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)378 // CHECK-NEXT: %[[PRIVATE3:.*]] = acc.private varPtr(%[[NOCOPYARR]] : !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>) bounds(%[[BOUNDS]]) recipe(@privatization__Bcnt1__ZTSA5_15NoCopyConstruct) -> !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> {name = "noCopyArr[1]"}379 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i380 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32381 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1382 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0383 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1384 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)385 // CHECK-NEXT: %[[PRIVATE4:.*]] = acc.private varPtr(%[[HASCOPYARR]] : !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>) bounds(%[[BOUNDS]]) recipe(@privatization__Bcnt1__ZTSA5_13CopyConstruct) -> !cir.ptr<!cir.array<!rec_CopyConstruct x 5>> {name = "hasCopyArr[1]"}386 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i387 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32388 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1389 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0390 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1391 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)392 // CHECK-NEXT: %[[PRIVATE5:.*]] = acc.private varPtr(%[[NOTDEFCTORARR]] : !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>) bounds(%[[BOUNDS]]) recipe(@privatization__Bcnt1__ZTSA5_14NonDefaultCtor) -> !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> {name = "notDefCtorArr[1]"}393 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i394 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32395 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1396 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0397 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1398 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)399 // CHECK-NEXT: %[[PRIVATE6:.*]] = acc.private varPtr(%[[DTORARR]] : !cir.ptr<!cir.array<!rec_HasDtor x 5>>) bounds(%[[BOUNDS]]) recipe(@privatization__Bcnt1__ZTSA5_7HasDtor) -> !cir.ptr<!cir.array<!rec_HasDtor x 5>> {name = "dtorArr[1]"}400 // CHECK-NEXT: acc.loop combined(serial) private(%[[PRIVATE1]], %[[PRIVATE2]], %[[PRIVATE3]], %[[PRIVATE4]], %[[PRIVATE5]], %[[PRIVATE6]] : !cir.ptr<!cir.array<!s32i x 5>>,401 // CHECK-SAME: !cir.ptr<!cir.array<!cir.float x 5>>,402 // CHECK-SAME: !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>,403 // CHECK-SAME: !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>,404 // CHECK-SAME: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>,405 // CHECK-SAME: !cir.ptr<!cir.array<!rec_HasDtor x 5>>)406 // CHECK: acc.yield407 // CHECK-NEXT: } loc408 // CHECK-NEXT: acc.yield409 // CHECK-NEXT: } loc410 411#pragma acc parallel loop private(someIntArr[1:1])412 for(int i = 0; i < 5; ++i);413 // CHECK-NEXT: acc.parallel combined(loop) {414 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i415 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32416 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i417 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32418 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0419 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1420 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)421 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[INTARR]] : !cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS]]) recipe(@privatization__Bcnt1__ZTSA5_i) -> !cir.ptr<!cir.array<!s32i x 5>> {name = "someIntArr[1:1]"}422 // CHECK-NEXT: acc.loop combined(parallel) private(%[[PRIVATE]] : !cir.ptr<!cir.array<!s32i x 5>>)423 // CHECK: acc.yield424 // CHECK-NEXT: } loc425 // CHECK-NEXT: acc.yield426 // CHECK-NEXT: } loc427#pragma acc serial loop private(someFloatArr[1:1])428 for(int i = 0; i < 5; ++i);429 // CHECK-NEXT: acc.serial combined(loop) {430 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i431 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32432 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i433 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32434 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0435 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1436 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)437 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[FLOATARR]] : !cir.ptr<!cir.array<!cir.float x 5>>) bounds(%[[BOUNDS]]) recipe(@privatization__Bcnt1__ZTSA5_f) -> !cir.ptr<!cir.array<!cir.float x 5>> {name = "someFloatArr[1:1]"}438 // CHECK-NEXT: acc.loop combined(serial) private(%[[PRIVATE]] : !cir.ptr<!cir.array<!cir.float x 5>>)439 // CHECK: acc.yield440 // CHECK-NEXT: } loc441 // CHECK-NEXT: acc.yield442 // CHECK-NEXT: } loc443#pragma acc parallel loop private(noCopyArr[1:1])444 for(int i = 0; i < 5; ++i);445 // CHECK-NEXT: acc.parallel combined(loop) {446 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i447 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32448 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i449 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32450 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0451 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1452 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)453 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[NOCOPYARR]] : !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>) bounds(%[[BOUNDS]]) recipe(@privatization__Bcnt1__ZTSA5_15NoCopyConstruct) -> !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> {name = "noCopyArr[1:1]"}454 // CHECK-NEXT: acc.loop combined(parallel) private(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>)455 // CHECK: acc.yield456 // CHECK-NEXT: } loc457 // CHECK-NEXT: acc.yield458 // CHECK-NEXT: } loc459#pragma acc serial loop private(hasCopyArr[1:1])460 for(int i = 0; i < 5; ++i);461 // CHECK-NEXT: acc.serial combined(loop) {462 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i463 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32464 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i465 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32466 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0467 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1468 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)469 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[HASCOPYARR]] : !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>) bounds(%[[BOUNDS]]) recipe(@privatization__Bcnt1__ZTSA5_13CopyConstruct) -> !cir.ptr<!cir.array<!rec_CopyConstruct x 5>> {name = "hasCopyArr[1:1]"}470 // CHECK-NEXT: acc.loop combined(serial) private(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>)471 // CHECK: acc.yield472 // CHECK-NEXT: } loc473 // CHECK-NEXT: acc.yield474 // CHECK-NEXT: } loc475#pragma acc parallel loop private(notDefCtorArr[1:1])476 for(int i = 0; i < 5; ++i);477 // CHECK-NEXT: acc.parallel combined(loop) {478 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i479 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32480 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i481 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32482 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0483 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1484 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)485 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[NOTDEFCTORARR]] : !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>) bounds(%[[BOUNDS]]) recipe(@privatization__Bcnt1__ZTSA5_14NonDefaultCtor) -> !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> {name = "notDefCtorArr[1:1]"}486 // CHECK-NEXT: acc.loop combined(parallel) private(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>)487 // CHECK: acc.yield488 // CHECK-NEXT: } loc489 // CHECK-NEXT: acc.yield490 // CHECK-NEXT: } loc491#pragma acc parallel loop private(dtorArr[1:1])492 for(int i = 0; i < 5; ++i);493 // CHECK-NEXT: acc.parallel combined(loop) {494 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i495 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32496 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i497 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32498 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0499 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1500 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)501 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[DTORARR]] : !cir.ptr<!cir.array<!rec_HasDtor x 5>>) bounds(%[[BOUNDS]]) recipe(@privatization__Bcnt1__ZTSA5_7HasDtor) -> !cir.ptr<!cir.array<!rec_HasDtor x 5>> {name = "dtorArr[1:1]"}502 // CHECK-NEXT: acc.loop combined(parallel) private(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_HasDtor x 5>>)503 // CHECK: acc.yield504 // CHECK-NEXT: } loc505 // CHECK-NEXT: acc.yield506 // CHECK-NEXT: } loc507#pragma acc parallel loop private(someIntArr[1:1], someFloatArr[1:1], noCopyArr[1:1], hasCopyArr[1:1], notDefCtorArr[1:1], dtorArr[1:1])508 for(int i = 0; i < 5; ++i);509 // CHECK-NEXT: acc.parallel combined(loop) {510 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i511 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32512 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i513 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32514 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0515 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1516 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)517 // CHECK-NEXT: %[[PRIVATE1:.*]] = acc.private varPtr(%[[INTARR]] : !cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS]]) recipe(@privatization__Bcnt1__ZTSA5_i) -> !cir.ptr<!cir.array<!s32i x 5>> {name = "someIntArr[1:1]"}518 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i519 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32520 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i521 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32522 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0523 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1524 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)525 // CHECK-NEXT: %[[PRIVATE2:.*]] = acc.private varPtr(%[[FLOATARR]] : !cir.ptr<!cir.array<!cir.float x 5>>) bounds(%[[BOUNDS]]) recipe(@privatization__Bcnt1__ZTSA5_f) -> !cir.ptr<!cir.array<!cir.float x 5>> {name = "someFloatArr[1:1]"}526 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i527 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32528 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i529 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32530 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0531 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1532 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)533 // CHECK-NEXT: %[[PRIVATE3:.*]] = acc.private varPtr(%[[NOCOPYARR]] : !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>) bounds(%[[BOUNDS]]) recipe(@privatization__Bcnt1__ZTSA5_15NoCopyConstruct) -> !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> {name = "noCopyArr[1:1]"}534 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i535 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32536 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i537 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32538 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0539 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1540 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)541 // CHECK-NEXT: %[[PRIVATE4:.*]] = acc.private varPtr(%[[HASCOPYARR]] : !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>) bounds(%[[BOUNDS]]) recipe(@privatization__Bcnt1__ZTSA5_13CopyConstruct) -> !cir.ptr<!cir.array<!rec_CopyConstruct x 5>> {name = "hasCopyArr[1:1]"}542 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i543 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32544 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i545 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32546 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0547 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1548 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)549 // CHECK-NEXT: %[[PRIVATE5:.*]] = acc.private varPtr(%[[NOTDEFCTORARR]] : !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>) bounds(%[[BOUNDS]]) recipe(@privatization__Bcnt1__ZTSA5_14NonDefaultCtor) -> !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> {name = "notDefCtorArr[1:1]"}550 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i551 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32552 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i553 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32554 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0555 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1556 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)557 // CHECK-NEXT: %[[PRIVATE6:.*]] = acc.private varPtr(%[[DTORARR]] : !cir.ptr<!cir.array<!rec_HasDtor x 5>>) bounds(%[[BOUNDS]]) recipe(@privatization__Bcnt1__ZTSA5_7HasDtor) -> !cir.ptr<!cir.array<!rec_HasDtor x 5>> {name = "dtorArr[1:1]"}558 // CHECK-NEXT: acc.loop combined(parallel) private(%[[PRIVATE1]], %[[PRIVATE2]], %[[PRIVATE3]], %[[PRIVATE4]], %[[PRIVATE5]], %[[PRIVATE6]] : !cir.ptr<!cir.array<!s32i x 5>>,559 // CHECK-SAME: !cir.ptr<!cir.array<!cir.float x 5>>,560 // CHECK-SAME: !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>,561 // CHECK-SAME: !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>,562 // CHECK-SAME: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>,563 // CHECK-SAME: !cir.ptr<!cir.array<!rec_HasDtor x 5>>)564 // CHECK: acc.yield565 // CHECK-NEXT: } loc566 // CHECK-NEXT: acc.yield567 // CHECK-NEXT: } loc568}569