506 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_loop() {161 // CHECK: cir.func{{.*}} @acc_loop()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 loop private(someInt)190 for(int i = 0; i < 5; ++i);191 // CHECK: %[[PRIVATE:.*]] = acc.private varPtr(%[[SOMEINT]] : !cir.ptr<!s32i>) recipe(@privatization__ZTSi) -> !cir.ptr<!s32i> {name = "someInt"}192 // CHECK-NEXT: acc.loop private(%[[PRIVATE]] : !cir.ptr<!s32i>)193 // CHECK: acc.yield194 // CHECK-NEXT: } loc195#pragma acc loop private(someFloat)196 for(int i = 0; i < 5; ++i);197 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[SOMEFLOAT]] : !cir.ptr<!cir.float>) recipe(@privatization__ZTSf) -> !cir.ptr<!cir.float> {name = "someFloat"}198 // CHECK-NEXT: acc.loop private(%[[PRIVATE]] : !cir.ptr<!cir.float>)199 // CHECK: acc.yield200 // CHECK-NEXT: } loc201 202#pragma acc loop private(noCopy)203 for(int i = 0; i < 5; ++i);204 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[NOCOPY]] : !cir.ptr<!rec_NoCopyConstruct>) recipe(@privatization__ZTS15NoCopyConstruct) -> !cir.ptr<!rec_NoCopyConstruct> {name = "noCopy"}205 // CHECK-NEXT: acc.loop private(%[[PRIVATE]] : !cir.ptr<!rec_NoCopyConstruct>206 // CHECK: acc.yield207 // CHECK-NEXT: } loc208#pragma acc loop private(hasCopy)209 for(int i = 0; i < 5; ++i);210 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[HASCOPY]] : !cir.ptr<!rec_CopyConstruct>) recipe(@privatization__ZTS13CopyConstruct) -> !cir.ptr<!rec_CopyConstruct> {name = "hasCopy"}211 // CHECK-NEXT: acc.loop private(%[[PRIVATE]] : !cir.ptr<!rec_CopyConstruct>)212 // CHECK: acc.yield213 // CHECK-NEXT: } loc214#pragma acc loop private(notDefCtor)215 for(int i = 0; i < 5; ++i);216 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[NOTDEFCTOR]] : !cir.ptr<!rec_NonDefaultCtor>) recipe(@privatization__ZTS14NonDefaultCtor) -> !cir.ptr<!rec_NonDefaultCtor> {name = "notDefCtor"}217 // CHECK-NEXT: acc.loop private(%[[PRIVATE]] : !cir.ptr<!rec_NonDefaultCtor>)218 // CHECK: acc.yield219 // CHECK-NEXT: } loc220#pragma acc loop private(dtor)221 for(int i = 0; i < 5; ++i);222 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[DTOR]] : !cir.ptr<!rec_HasDtor>) recipe(@privatization__ZTS7HasDtor) -> !cir.ptr<!rec_HasDtor> {name = "dtor"}223 // CHECK-NEXT: acc.loop private(%[[PRIVATE]] : !cir.ptr<!rec_HasDtor>)224 // CHECK: acc.yield225 // CHECK-NEXT: } loc226 227#pragma acc loop private(someInt, someFloat, noCopy, hasCopy, notDefCtor, dtor)228 for(int i = 0; i < 5; ++i);229 // CHECK: %[[PRIVATE1:.*]] = acc.private varPtr(%[[SOMEINT]] : !cir.ptr<!s32i>) recipe(@privatization__ZTSi) -> !cir.ptr<!s32i> {name = "someInt"}230 // CHECK-NEXT: %[[PRIVATE2:.*]] = acc.private varPtr(%[[SOMEFLOAT]] : !cir.ptr<!cir.float>) recipe(@privatization__ZTSf) -> !cir.ptr<!cir.float> {name = "someFloat"}231 // CHECK-NEXT: %[[PRIVATE3:.*]] = acc.private varPtr(%[[NOCOPY]] : !cir.ptr<!rec_NoCopyConstruct>) recipe(@privatization__ZTS15NoCopyConstruct) -> !cir.ptr<!rec_NoCopyConstruct> {name = "noCopy"}232 // CHECK-NEXT: %[[PRIVATE4:.*]] = acc.private varPtr(%[[HASCOPY]] : !cir.ptr<!rec_CopyConstruct>) recipe(@privatization__ZTS13CopyConstruct) -> !cir.ptr<!rec_CopyConstruct> {name = "hasCopy"}233 // CHECK-NEXT: %[[PRIVATE5:.*]] = acc.private varPtr(%[[NOTDEFCTOR]] : !cir.ptr<!rec_NonDefaultCtor>) recipe(@privatization__ZTS14NonDefaultCtor) -> !cir.ptr<!rec_NonDefaultCtor> {name = "notDefCtor"}234 // CHECK-NEXT: %[[PRIVATE6:.*]] = acc.private varPtr(%[[DTOR]] : !cir.ptr<!rec_HasDtor>) recipe(@privatization__ZTS7HasDtor) -> !cir.ptr<!rec_HasDtor> {name = "dtor"}235 // CHECK-NEXT: acc.loop private(%[[PRIVATE1]], %[[PRIVATE2]], %[[PRIVATE3]], %[[PRIVATE4]], %[[PRIVATE5]], %[[PRIVATE6]] : !cir.ptr<!s32i>,236 // CHECK-SAME: !cir.ptr<!cir.float>,237 // CHECK-SAME: !cir.ptr<!rec_NoCopyConstruct>,238 // CHECK-SAME: !cir.ptr<!rec_CopyConstruct>,239 // CHECK-SAME: !cir.ptr<!rec_NonDefaultCtor>,240 // CHECK-SAME: !cir.ptr<!rec_HasDtor>)241 // CHECK: acc.yield242 // CHECK-NEXT: } loc243 244#pragma acc loop private(someIntArr[1])245 for(int i = 0; i < 5; ++i);246 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i247 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32248 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1249 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0250 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1251 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)252 // 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]"}253 // CHECK-NEXT: acc.loop private(%[[PRIVATE]] : !cir.ptr<!cir.array<!s32i x 5>>)254 // CHECK: acc.yield255 // CHECK-NEXT: } loc256#pragma acc loop private(someFloatArr[1])257 for(int i = 0; i < 5; ++i);258 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i259 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32260 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1261 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0262 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1263 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)264 // 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]"}265 // CHECK-NEXT: acc.loop private(%[[PRIVATE]] : !cir.ptr<!cir.array<!cir.float x 5>>)266 // CHECK: acc.yield267 // CHECK-NEXT: } loc268#pragma acc loop private(noCopyArr[1])269 for(int i = 0; i < 5; ++i);270 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i271 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32272 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1273 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0274 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1275 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)276 // 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]"}277 // CHECK-NEXT: acc.loop private(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>)278 // CHECK: acc.yield279 // CHECK-NEXT: } loc280#pragma acc loop private(hasCopyArr[1])281 for(int i = 0; i < 5; ++i);282 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i283 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32284 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1285 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0286 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1287 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)288 // 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]"}289 // CHECK-NEXT: acc.loop private(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>)290 // CHECK: acc.yield291 // CHECK-NEXT: } loc292#pragma acc loop private(notDefCtorArr[1])293 for(int i = 0; i < 5; ++i);294 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i295 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32296 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1297 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0298 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1299 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)300 // 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]"}301 // CHECK-NEXT: acc.loop private(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>)302 // CHECK: acc.yield303 // CHECK-NEXT: } loc304#pragma acc loop private(dtorArr[1])305 for(int i = 0; i < 5; ++i);306 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i307 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32308 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1309 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0310 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1311 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)312 // 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]"}313 // CHECK-NEXT: acc.loop private(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_HasDtor x 5>>)314 // CHECK: acc.yield315 // CHECK-NEXT: } loc316#pragma acc loop private(someIntArr[1], someFloatArr[1], noCopyArr[1], hasCopyArr[1], notDefCtorArr[1], dtorArr[1])317 for(int i = 0; i < 5; ++i);318 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i319 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32320 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1321 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0322 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1323 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)324 // 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]"}325 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i326 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32327 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1328 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0329 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1330 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)331 // 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]"}332 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i333 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32334 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1335 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0336 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1337 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)338 // 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]"}339 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i340 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32341 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1342 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0343 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1344 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)345 // 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]"}346 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i347 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32348 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1349 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0350 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1351 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)352 // 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]"}353 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i354 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32355 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1356 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0357 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1358 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)359 // 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]"}360 // CHECK-NEXT: acc.loop private(%[[PRIVATE1]], %[[PRIVATE2]], %[[PRIVATE3]], %[[PRIVATE4]], %[[PRIVATE5]], %[[PRIVATE6]] : !cir.ptr<!cir.array<!s32i x 5>>,361 // CHECK-SAME: !cir.ptr<!cir.array<!cir.float x 5>>,362 // CHECK-SAME: !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>,363 // CHECK-SAME: !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>,364 // CHECK-SAME: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>,365 // CHECK-SAME: !cir.ptr<!cir.array<!rec_HasDtor x 5>>)366 // CHECK: acc.yield367 // CHECK-NEXT: } loc368 369#pragma acc loop private(someIntArr[1:1])370 for(int i = 0; i < 5; ++i);371 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i372 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32373 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i374 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32375 // 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_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)378 // 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]"}379 // CHECK-NEXT: acc.loop private(%[[PRIVATE]] : !cir.ptr<!cir.array<!s32i x 5>>)380 // CHECK: acc.yield381 // CHECK-NEXT: } loc382#pragma acc loop private(someFloatArr[1:1])383 for(int i = 0; i < 5; ++i);384 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i385 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32386 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i387 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32388 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0389 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1390 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)391 // 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]"}392 // CHECK-NEXT: acc.loop private(%[[PRIVATE]] : !cir.ptr<!cir.array<!cir.float x 5>>)393 // CHECK: acc.yield394 // CHECK-NEXT: } loc395#pragma acc loop private(noCopyArr[1:1])396 for(int i = 0; i < 5; ++i);397 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i398 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32399 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i400 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32401 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0402 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1403 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)404 // 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]"}405 // CHECK-NEXT: acc.loop private(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>)406 // CHECK: acc.yield407 // CHECK-NEXT: } loc408#pragma acc loop private(hasCopyArr[1:1])409 for(int i = 0; i < 5; ++i);410 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i411 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32412 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i413 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32414 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0415 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1416 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)417 // 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]"}418 // CHECK-NEXT: acc.loop private(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>)419 // CHECK: acc.yield420 // CHECK-NEXT: } loc421#pragma acc loop private(notDefCtorArr[1:1])422 for(int i = 0; i < 5; ++i);423 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i424 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32425 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i426 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32427 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0428 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1429 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)430 // 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]"}431 // CHECK-NEXT: acc.loop private(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>)432 // CHECK: acc.yield433 // CHECK-NEXT: } loc434#pragma acc loop private(dtorArr[1:1])435 for(int i = 0; i < 5; ++i);436 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i437 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32438 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i439 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32440 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0441 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1442 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)443 // 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]"}444 // CHECK-NEXT: acc.loop private(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_HasDtor x 5>>)445 // CHECK: acc.yield446 // CHECK-NEXT: } loc447#pragma acc loop private(someIntArr[1:1], someFloatArr[1:1], noCopyArr[1:1], hasCopyArr[1:1], notDefCtorArr[1:1], dtorArr[1:1])448 for(int i = 0; i < 5; ++i);449 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i450 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32451 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i452 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32453 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0454 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1455 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)456 // 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]"}457 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i458 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32459 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i460 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32461 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0462 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1463 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)464 // 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]"}465 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i466 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32467 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i468 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32469 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0470 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1471 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)472 // 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]"}473 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i474 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32475 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i476 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32477 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0478 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1479 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)480 // 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]"}481 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i482 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32483 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i484 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32485 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0486 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1487 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)488 // 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]"}489 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i490 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32491 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i492 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32493 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0494 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1495 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)496 // 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]"}497 // CHECK-NEXT: acc.loop private(%[[PRIVATE1]], %[[PRIVATE2]], %[[PRIVATE3]], %[[PRIVATE4]], %[[PRIVATE5]], %[[PRIVATE6]] : !cir.ptr<!cir.array<!s32i x 5>>,498 // CHECK-SAME: !cir.ptr<!cir.array<!cir.float x 5>>,499 // CHECK-SAME: !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>,500 // CHECK-SAME: !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>,501 // CHECK-SAME: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>,502 // CHECK-SAME: !cir.ptr<!cir.array<!rec_HasDtor x 5>>)503 // CHECK: acc.yield504 // CHECK-NEXT: } loc505}506