494 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// CHECK: acc.private.recipe @privatization__ZTSi : !cir.ptr<!s32i> init {19// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!s32i> {{.*}}):20// CHECK-NEXT: cir.alloca !s32i, !cir.ptr<!s32i>, ["openacc.private.init"]21// CHECK-NEXT: acc.yield22// CHECK-NEXT: }23//24// CHECK-NEXT: acc.private.recipe @privatization__ZTSf : !cir.ptr<!cir.float> init {25// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.float> {{.*}}):26// CHECK-NEXT: cir.alloca !cir.float, !cir.ptr<!cir.float>, ["openacc.private.init"]27// CHECK-NEXT: acc.yield28// CHECK-NEXT: }29//30// CHECK-NEXT: acc.private.recipe @privatization__ZTS15NoCopyConstruct : !cir.ptr<!rec_NoCopyConstruct> init {31// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_NoCopyConstruct> {{.*}}):32// CHECK-NEXT: cir.alloca !rec_NoCopyConstruct, !cir.ptr<!rec_NoCopyConstruct>, ["openacc.private.init"]33// CHECK-NEXT: acc.yield34// CHECK-NEXT: }35//36// CHECK-NEXT: acc.private.recipe @privatization__ZTS13CopyConstruct : !cir.ptr<!rec_CopyConstruct> init {37// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_CopyConstruct> {{.*}}):38// CHECK-NEXT: cir.alloca !rec_CopyConstruct, !cir.ptr<!rec_CopyConstruct>, ["openacc.private.init"]39// CHECK-NEXT: acc.yield40// CHECK-NEXT: }41//42// CHECK-NEXT: acc.private.recipe @privatization__ZTS14NonDefaultCtor : !cir.ptr<!rec_NonDefaultCtor> init {43// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_NonDefaultCtor> {{.*}}):44// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !rec_NonDefaultCtor, !cir.ptr<!rec_NonDefaultCtor>, ["openacc.private.init", init]45// CHECK-NEXT: cir.call @_ZN14NonDefaultCtorC1Ev(%[[ALLOCA]]) : (!cir.ptr<!rec_NonDefaultCtor>) -> ()46// CHECK-NEXT: acc.yield47// CHECK-NEXT: }48//49// CHECK-NEXT: acc.private.recipe @privatization__ZTS7HasDtor : !cir.ptr<!rec_HasDtor> init {50// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_HasDtor> {{.*}}):51// CHECK-NEXT: cir.alloca !rec_HasDtor, !cir.ptr<!rec_HasDtor>, ["openacc.private.init"]52// CHECK-NEXT: acc.yield53// CHECK-NEXT: } destroy {54// CHECK-NEXT: ^bb0(%[[ORIG:.*]]: !cir.ptr<!rec_HasDtor> {{.*}}, %[[ARG:.*]]: !cir.ptr<!rec_HasDtor> {{.*}}):55// CHECK-NEXT: cir.call @_ZN7HasDtorD1Ev(%[[ARG]]) nothrow : (!cir.ptr<!rec_HasDtor>) -> ()56// CHECK-NEXT: acc.yield57// CHECK-NEXT: }58//59// CHECK-NEXT: acc.private.recipe @privatization__Bcnt1__ZTSA5_i : !cir.ptr<!cir.array<!s32i x 5>> init {60// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!s32i x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}):61// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !cir.array<!s32i x 5>, !cir.ptr<!cir.array<!s32i x 5>>, ["openacc.private.init"]62// CHECK-NEXT: acc.yield63// CHECK-NEXT: }64//65// CHECK-NEXT: acc.private.recipe @privatization__Bcnt1__ZTSA5_f : !cir.ptr<!cir.array<!cir.float x 5>> init {66// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!cir.float x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}):67// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !cir.array<!cir.float x 5>, !cir.ptr<!cir.array<!cir.float x 5>>, ["openacc.private.init"]68// CHECK-NEXT: acc.yield69// CHECK-NEXT: }70//71// CHECK-NEXT: acc.private.recipe @privatization__Bcnt1__ZTSA5_15NoCopyConstruct : !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> init {72// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}):73// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !cir.array<!rec_NoCopyConstruct x 5>, !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>, ["openacc.private.init"]74// CHECK-NEXT: acc.yield75// CHECK-NEXT: }76//77// CHECK-NEXT: acc.private.recipe @privatization__Bcnt1__ZTSA5_13CopyConstruct : !cir.ptr<!cir.array<!rec_CopyConstruct x 5>> init {78// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!rec_CopyConstruct x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}):79// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !cir.array<!rec_CopyConstruct x 5>, !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>, ["openacc.private.init"]80// CHECK-NEXT: acc.yield81// CHECK-NEXT: }82//83// CHECK-NEXT: acc.private.recipe @privatization__Bcnt1__ZTSA5_14NonDefaultCtor : !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> init {84// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}):85// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !cir.array<!rec_NonDefaultCtor x 5>, !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>, ["openacc.private.init"]86// CHECK-NEXT: cir.scope {87// CHECK-NEXT: %[[LB:.*]] = acc.get_lowerbound %[[BOUND1]] : (!acc.data_bounds_ty) -> index88// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]] : index to !u64i89// CHECK-NEXT: %[[UB:.*]] = acc.get_upperbound %[[BOUND1]] : (!acc.data_bounds_ty) -> index90// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]] : index to !u64i91// CHECK-NEXT: %[[ITR:.*]] = cir.alloca !u64i, !cir.ptr<!u64i>, ["iter"] {alignment = 8 : i64}92// CHECK-NEXT: cir.store %[[LB_CAST]], %[[ITR]] : !u64i, !cir.ptr<!u64i>93// CHECK-NEXT: cir.for : cond {94// CHECK-NEXT: %[[ITR_LOAD:.*]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i95// CHECK-NEXT: %[[COND:.*]] = cir.cmp(lt, %[[ITR_LOAD]], %[[UB_CAST]]) : !u64i, !cir.bool96// CHECK-NEXT: cir.condition(%[[COND]])97// CHECK-NEXT: } body {98// CHECK-NEXT: %[[ITR_LOAD:.*]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i99// CHECK-NEXT: %[[DECAY:.*]] = cir.cast array_to_ptrdecay %[[ALLOCA]] : !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> -> !cir.ptr<!rec_NonDefaultCtor>100// CHECK-NEXT: %[[STRIDE:.*]] = cir.ptr_stride %[[DECAY]], %[[ITR_LOAD]] : (!cir.ptr<!rec_NonDefaultCtor>, !u64i) -> !cir.ptr<!rec_NonDefaultCtor>101// CHECK-NEXT: cir.call @_ZN14NonDefaultCtorC1Ev(%[[STRIDE]]) : (!cir.ptr<!rec_NonDefaultCtor>) -> ()102// CHECK-NEXT: cir.yield103// CHECK-NEXT: } step {104// CHECK-NEXT: %[[ITR_LOAD]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i105// CHECK-NEXT: %[[INC:.*]] = cir.unary(inc, %[[ITR_LOAD]]) : !u64i, !u64i106// CHECK-NEXT: cir.store %[[INC]], %[[ITR]] : !u64i, !cir.ptr<!u64i>107// CHECK-NEXT: cir.yield108// CHECK-NEXT: }109// CHECK-NEXT: }110// CHECK-NEXT: acc.yield111// CHECK-NEXT: }112//113// CHECK: acc.private.recipe @privatization__Bcnt1__ZTSA5_7HasDtor : !cir.ptr<!cir.array<!rec_HasDtor x 5>> init {114// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!rec_HasDtor x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}):115// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !cir.array<!rec_HasDtor x 5>, !cir.ptr<!cir.array<!rec_HasDtor x 5>>, ["openacc.private.init"]116// CHECK-NEXT: acc.yield117// CHECK-NEXT: } destroy {118// 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 {{.*}}):119// CHECK-NEXT: cir.scope {120// CHECK-NEXT: %[[LB1:.*]] = acc.get_lowerbound %[[BOUND1]] : (!acc.data_bounds_ty) -> index121// CHECK-NEXT: %[[LB1_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB1]] : index to !u64i122// CHECK-NEXT: %[[UB1:.*]] = acc.get_upperbound %[[BOUND1]] : (!acc.data_bounds_ty) -> index123// CHECK-NEXT: %[[UB1_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB1]] : index to !u64i124// CHECK-NEXT: %[[ITR:.*]] = cir.alloca !u64i, !cir.ptr<!u64i>, ["iter"] {alignment = 8 : i64}125// CHECK-NEXT: %[[CONST_ONE:.*]] = cir.const #cir.int<1> : !u64i126// CHECK-NEXT: %[[ONE_BELOW_UB:.*]] = cir.binop(sub, %[[UB1_CAST]], %[[CONST_ONE]]) : !u64i127// CHECK-NEXT: cir.store %[[ONE_BELOW_UB]], %[[ITR]] : !u64i, !cir.ptr<!u64i>128// CHECK-NEXT: cir.for : cond {129// CHECK-NEXT: %[[ITR_LOAD:.*]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i130// CHECK-NEXT: %[[COND:.*]] = cir.cmp(ge, %[[ITR_LOAD]], %[[LB1_CAST]]) : !u64i, !cir.bool131// CHECK-NEXT: cir.condition(%[[COND]])132// CHECK-NEXT: } body {133// CHECK-NEXT: %[[ITR_LOAD:.*]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i134// CHECK-NEXT: %[[DECAY:.*]] = cir.cast array_to_ptrdecay %[[ARG]] : !cir.ptr<!cir.array<!rec_HasDtor x 5>> -> !cir.ptr<!rec_HasDtor>135// CHECK-NEXT: %[[STRIDE:.*]] = cir.ptr_stride %[[DECAY]], %[[ITR_LOAD]] : (!cir.ptr<!rec_HasDtor>, !u64i) -> !cir.ptr<!rec_HasDtor>136// CHECK-NEXT: cir.call @_ZN7HasDtorD1Ev(%[[STRIDE]]) nothrow : (!cir.ptr<!rec_HasDtor>) -> ()137// CHECK-NEXT: cir.yield138// CHECK-NEXT: } step {139// CHECK-NEXT: %[[ITR_LOAD]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i140// CHECK-NEXT: %[[DEC:.*]] = cir.unary(dec, %[[ITR_LOAD]]) : !u64i, !u64i141// CHECK-NEXT: cir.store %[[DEC]], %[[ITR]] : !u64i, !cir.ptr<!u64i>142// CHECK-NEXT: cir.yield143// CHECK-NEXT: }144// CHECK-NEXT: }145// CHECK-NEXT: acc.yield146// CHECK-NEXT: }147 148extern "C" void acc_compute() {149 // CHECK: cir.func{{.*}} @acc_compute()150 151 int someInt;152 // CHECK-NEXT: %[[SOMEINT:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["someInt"]153 float someFloat;154 // CHECK-NEXT: %[[SOMEFLOAT:.*]] = cir.alloca !cir.float, !cir.ptr<!cir.float>, ["someFloat"]155 NoCopyConstruct noCopy;156 // CHECK-NEXT: %[[NOCOPY:.*]] = cir.alloca !rec_NoCopyConstruct, !cir.ptr<!rec_NoCopyConstruct>, ["noCopy"]157 CopyConstruct hasCopy;158 // CHECK-NEXT: %[[HASCOPY:.*]] = cir.alloca !rec_CopyConstruct, !cir.ptr<!rec_CopyConstruct>, ["hasCopy"]159 NonDefaultCtor notDefCtor;160 // CHECK-NEXT: %[[NOTDEFCTOR:.*]] = cir.alloca !rec_NonDefaultCtor, !cir.ptr<!rec_NonDefaultCtor>, ["notDefCtor", init]161 HasDtor dtor;162 // CHECK-NEXT: %[[DTOR:.*]] = cir.alloca !rec_HasDtor, !cir.ptr<!rec_HasDtor>, ["dtor"]163 int someIntArr[5];164 // CHECK-NEXT: %[[INTARR:.*]] = cir.alloca !cir.array<!s32i x 5>, !cir.ptr<!cir.array<!s32i x 5>>, ["someIntArr"]165 float someFloatArr[5];166 // CHECK-NEXT: %[[FLOATARR:.*]] = cir.alloca !cir.array<!cir.float x 5>, !cir.ptr<!cir.array<!cir.float x 5>>, ["someFloatArr"]167 NoCopyConstruct noCopyArr[5];168 // CHECK-NEXT: %[[NOCOPYARR:.*]] = cir.alloca !cir.array<!rec_NoCopyConstruct x 5>, !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>, ["noCopyArr"]169 CopyConstruct hasCopyArr[5];170 // CHECK-NEXT: %[[HASCOPYARR:.*]] = cir.alloca !cir.array<!rec_CopyConstruct x 5>, !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>, ["hasCopyArr"]171 NonDefaultCtor notDefCtorArr[5];172 // CHECK-NEXT: %[[NOTDEFCTORARR:.*]] = cir.alloca !cir.array<!rec_NonDefaultCtor x 5>, !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>, ["notDefCtorArr", init]173 HasDtor dtorArr[5];174 // CHECK-NEXT: %[[DTORARR:.*]] = cir.alloca !cir.array<!rec_HasDtor x 5>, !cir.ptr<!cir.array<!rec_HasDtor x 5>>, ["dtorArr"]175 // CHECK-NEXT: cir.call @_ZN14NonDefaultCtorC1Ev(%[[NOTDEFCTOR]]) : (!cir.ptr<!rec_NonDefaultCtor>) -> ()176 177#pragma acc parallel private(someInt)178 ;179 // CHECK: %[[PRIVATE:.*]] = acc.private varPtr(%[[SOMEINT]] : !cir.ptr<!s32i>) recipe(@privatization__ZTSi) -> !cir.ptr<!s32i> {name = "someInt"}180 // CHECK-NEXT: acc.parallel private(%[[PRIVATE]] : !cir.ptr<!s32i>)181 // CHECK-NEXT: acc.yield182 // CHECK-NEXT: } loc183#pragma acc serial private(someFloat)184 ;185 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[SOMEFLOAT]] : !cir.ptr<!cir.float>) recipe(@privatization__ZTSf) -> !cir.ptr<!cir.float> {name = "someFloat"}186 // CHECK-NEXT: acc.serial private(%[[PRIVATE]] : !cir.ptr<!cir.float>)187 // CHECK-NEXT: acc.yield188 // CHECK-NEXT: } loc189 190#pragma acc parallel private(noCopy)191 ;192 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[NOCOPY]] : !cir.ptr<!rec_NoCopyConstruct>) recipe(@privatization__ZTS15NoCopyConstruct) -> !cir.ptr<!rec_NoCopyConstruct> {name = "noCopy"}193 // CHECK-NEXT: acc.parallel private(%[[PRIVATE]] : !cir.ptr<!rec_NoCopyConstruct>194 // CHECK-NEXT: acc.yield195 // CHECK-NEXT: } loc196#pragma acc serial private(hasCopy)197 ;198 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[HASCOPY]] : !cir.ptr<!rec_CopyConstruct>) recipe(@privatization__ZTS13CopyConstruct) -> !cir.ptr<!rec_CopyConstruct> {name = "hasCopy"}199 // CHECK-NEXT: acc.serial private(%[[PRIVATE]] : !cir.ptr<!rec_CopyConstruct>)200 // CHECK-NEXT: acc.yield201 // CHECK-NEXT: } loc202#pragma acc serial private(notDefCtor)203 ;204 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[NOTDEFCTOR]] : !cir.ptr<!rec_NonDefaultCtor>) recipe(@privatization__ZTS14NonDefaultCtor) -> !cir.ptr<!rec_NonDefaultCtor> {name = "notDefCtor"}205 // CHECK-NEXT: acc.serial private(%[[PRIVATE]] : !cir.ptr<!rec_NonDefaultCtor>)206 // CHECK-NEXT: acc.yield207 // CHECK-NEXT: } loc208#pragma acc serial private(dtor)209 ;210 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[DTOR]] : !cir.ptr<!rec_HasDtor>) recipe(@privatization__ZTS7HasDtor) -> !cir.ptr<!rec_HasDtor> {name = "dtor"}211 // CHECK-NEXT: acc.serial private(%[[PRIVATE]] : !cir.ptr<!rec_HasDtor>)212 // CHECK-NEXT: acc.yield213 // CHECK-NEXT: } loc214 215#pragma acc parallel private(someInt, someFloat, noCopy, hasCopy, notDefCtor, dtor)216 ;217 // CHECK: %[[PRIVATE1:.*]] = acc.private varPtr(%[[SOMEINT]] : !cir.ptr<!s32i>) recipe(@privatization__ZTSi) -> !cir.ptr<!s32i> {name = "someInt"}218 // CHECK-NEXT: %[[PRIVATE2:.*]] = acc.private varPtr(%[[SOMEFLOAT]] : !cir.ptr<!cir.float>) recipe(@privatization__ZTSf) -> !cir.ptr<!cir.float> {name = "someFloat"}219 // CHECK-NEXT: %[[PRIVATE3:.*]] = acc.private varPtr(%[[NOCOPY]] : !cir.ptr<!rec_NoCopyConstruct>) recipe(@privatization__ZTS15NoCopyConstruct) -> !cir.ptr<!rec_NoCopyConstruct> {name = "noCopy"}220 // CHECK-NEXT: %[[PRIVATE4:.*]] = acc.private varPtr(%[[HASCOPY]] : !cir.ptr<!rec_CopyConstruct>) recipe(@privatization__ZTS13CopyConstruct) -> !cir.ptr<!rec_CopyConstruct> {name = "hasCopy"}221 // CHECK-NEXT: %[[PRIVATE5:.*]] = acc.private varPtr(%[[NOTDEFCTOR]] : !cir.ptr<!rec_NonDefaultCtor>) recipe(@privatization__ZTS14NonDefaultCtor) -> !cir.ptr<!rec_NonDefaultCtor> {name = "notDefCtor"}222 // CHECK-NEXT: %[[PRIVATE6:.*]] = acc.private varPtr(%[[DTOR]] : !cir.ptr<!rec_HasDtor>) recipe(@privatization__ZTS7HasDtor) -> !cir.ptr<!rec_HasDtor> {name = "dtor"}223 // CHECK-NEXT: acc.parallel private(%[[PRIVATE1]], %[[PRIVATE2]], %[[PRIVATE3]], %[[PRIVATE4]], %[[PRIVATE5]], %[[PRIVATE6]] : !cir.ptr<!s32i>,224 // CHECK-SAME: !cir.ptr<!cir.float>,225 // CHECK-SAME: !cir.ptr<!rec_NoCopyConstruct>,226 // CHECK-SAME: !cir.ptr<!rec_CopyConstruct>,227 // CHECK-SAME: !cir.ptr<!rec_NonDefaultCtor>,228 // CHECK-SAME: !cir.ptr<!rec_HasDtor>)229 // CHECK-NEXT: acc.yield230 // CHECK-NEXT: } loc231 232#pragma acc serial private(someIntArr[1])233 ;234 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i235 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32236 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1237 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0238 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1239 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)240 // 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]"}241 // CHECK-NEXT: acc.serial private(%[[PRIVATE]] : !cir.ptr<!cir.array<!s32i x 5>>)242 // CHECK-NEXT: acc.yield243 // CHECK-NEXT: } loc244#pragma acc parallel private(someFloatArr[1])245 ;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(%[[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]"}253 // CHECK-NEXT: acc.parallel private(%[[PRIVATE]] : !cir.ptr<!cir.array<!cir.float x 5>>)254 // CHECK-NEXT: acc.yield255 // CHECK-NEXT: } loc256#pragma acc serial private(noCopyArr[1])257 ;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(%[[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]"}265 // CHECK-NEXT: acc.serial private(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>)266 // CHECK-NEXT: acc.yield267 // CHECK-NEXT: } loc268#pragma acc parallel private(hasCopyArr[1])269 ;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(%[[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]"}277 // CHECK-NEXT: acc.parallel private(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>)278 // CHECK-NEXT: acc.yield279 // CHECK-NEXT: } loc280#pragma acc parallel private(notDefCtorArr[1])281 ;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(%[[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]"}289 // CHECK-NEXT: acc.parallel private(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>)290 // CHECK-NEXT: acc.yield291 // CHECK-NEXT: } loc292#pragma acc parallel private(dtorArr[1])293 ;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(%[[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]"}301 // CHECK-NEXT: acc.parallel private(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_HasDtor x 5>>)302 // CHECK-NEXT: acc.yield303 // CHECK-NEXT: } loc304#pragma acc serial private(someIntArr[1], someFloatArr[1], noCopyArr[1], hasCopyArr[1], notDefCtorArr[1], dtorArr[1])305 ;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: %[[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]"}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: %[[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]"}320 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i321 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32322 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1323 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0324 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1325 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)326 // 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]"}327 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i328 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32329 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1330 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0331 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1332 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)333 // 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]"}334 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i335 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32336 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1337 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0338 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1339 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)340 // 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]"}341 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i342 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32343 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1344 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0345 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1346 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)347 // 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]"}348 // CHECK-NEXT: acc.serial private(%[[PRIVATE1]], %[[PRIVATE2]], %[[PRIVATE3]], %[[PRIVATE4]], %[[PRIVATE5]], %[[PRIVATE6]] : !cir.ptr<!cir.array<!s32i x 5>>,349 // CHECK-SAME: !cir.ptr<!cir.array<!cir.float x 5>>,350 // CHECK-SAME: !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>,351 // CHECK-SAME: !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>,352 // CHECK-SAME: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>,353 // CHECK-SAME: !cir.ptr<!cir.array<!rec_HasDtor x 5>>)354 // CHECK-NEXT: acc.yield355 // CHECK-NEXT: } loc356 357#pragma acc parallel private(someIntArr[1:1])358 ;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:.*]] = cir.const #cir.int<1> : !s32i362 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32363 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0364 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1365 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)366 // 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]"}367 // CHECK-NEXT: acc.parallel private(%[[PRIVATE]] : !cir.ptr<!cir.array<!s32i x 5>>)368 // CHECK-NEXT: acc.yield369 // CHECK-NEXT: } loc370#pragma acc serial private(someFloatArr[1:1])371 ;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:.*]] = cir.const #cir.int<1> : !s32i375 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32376 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0377 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1378 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)379 // 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]"}380 // CHECK-NEXT: acc.serial private(%[[PRIVATE]] : !cir.ptr<!cir.array<!cir.float x 5>>)381 // CHECK-NEXT: acc.yield382 // CHECK-NEXT: } loc383#pragma acc parallel private(noCopyArr[1:1])384 ;385 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i386 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32387 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i388 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32389 // 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_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)392 // 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]"}393 // CHECK-NEXT: acc.parallel private(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>)394 // CHECK-NEXT: acc.yield395 // CHECK-NEXT: } loc396#pragma acc serial private(hasCopyArr[1:1])397 ;398 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i399 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32400 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i401 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32402 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0403 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1404 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)405 // 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]"}406 // CHECK-NEXT: acc.serial private(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>)407 // CHECK-NEXT: acc.yield408 // CHECK-NEXT: } loc409#pragma acc parallel private(notDefCtorArr[1:1])410 ;411 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i412 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32413 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i414 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32415 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0416 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1417 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)418 // 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]"}419 // CHECK-NEXT: acc.parallel private(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>)420 // CHECK-NEXT: acc.yield421 // CHECK-NEXT: } loc422#pragma acc parallel private(dtorArr[1:1])423 ;424 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i425 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32426 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i427 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32428 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0429 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1430 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)431 // 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]"}432 // CHECK-NEXT: acc.parallel private(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_HasDtor x 5>>)433 // CHECK-NEXT: acc.yield434 // CHECK-NEXT: } loc435#pragma acc parallel private(someIntArr[1:1], someFloatArr[1:1], noCopyArr[1:1], hasCopyArr[1:1], notDefCtorArr[1:1], dtorArr[1:1])436 ;437 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i438 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32439 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i440 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32441 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0442 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1443 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)444 // 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]"}445 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i446 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32447 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i448 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32449 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0450 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1451 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)452 // 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]"}453 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i454 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32455 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i456 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32457 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0458 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1459 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)460 // 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]"}461 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i462 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32463 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i464 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32465 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0466 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1467 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)468 // 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]"}469 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i470 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32471 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i472 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32473 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0474 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1475 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)476 // 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]"}477 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i478 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32479 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i480 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32481 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0482 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1483 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)484 // 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]"}485 // CHECK-NEXT: acc.parallel private(%[[PRIVATE1]], %[[PRIVATE2]], %[[PRIVATE3]], %[[PRIVATE4]], %[[PRIVATE5]], %[[PRIVATE6]] : !cir.ptr<!cir.array<!s32i x 5>>,486 // CHECK-SAME: !cir.ptr<!cir.array<!cir.float x 5>>,487 // CHECK-SAME: !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>,488 // CHECK-SAME: !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>,489 // CHECK-SAME: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>,490 // CHECK-SAME: !cir.ptr<!cir.array<!rec_HasDtor x 5>>)491 // CHECK-NEXT: acc.yield492 // CHECK-NEXT: } loc493}494