230 lines · c
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 5// int6// CHECK: acc.private.recipe @privatization__ZTSi : !cir.ptr<!s32i> init {7// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!s32i> {{.*}}):8// CHECK-NEXT: cir.alloca !s32i, !cir.ptr<!s32i>, ["openacc.private.init"]9// CHECK-NEXT: acc.yield10// CHECK-NEXT: }11//12// float13// CHECK-NEXT: acc.private.recipe @privatization__ZTSf : !cir.ptr<!cir.float> init {14// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.float> {{.*}}):15// CHECK-NEXT: cir.alloca !cir.float, !cir.ptr<!cir.float>, ["openacc.private.init"]16// CHECK-NEXT: acc.yield17// CHECK-NEXT: }18//19// NoCopyConstruct20// CHECK-NEXT: acc.private.recipe @privatization__ZTS15NoCopyConstruct : !cir.ptr<!rec_NoCopyConstruct> init {21// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_NoCopyConstruct> {{.*}}):22// CHECK-NEXT: cir.alloca !rec_NoCopyConstruct, !cir.ptr<!rec_NoCopyConstruct>, ["openacc.private.init"]23// CHECK-NEXT: acc.yield24// CHECK-NEXT: }25//26// int[5] with 1 'bound'27// CHECK-NEXT: acc.private.recipe @privatization__Bcnt1__ZTSA5_i : !cir.ptr<!cir.array<!s32i x 5>> init {28// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!s32i x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}):29// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !cir.array<!s32i x 5>, !cir.ptr<!cir.array<!s32i x 5>>, ["openacc.private.init"]30// CHECK-NEXT: acc.yield31// CHECK-NEXT: }32//33// float[5] with 1 'bound'34// CHECK-NEXT: acc.private.recipe @privatization__Bcnt1__ZTSA5_f : !cir.ptr<!cir.array<!cir.float x 5>> init {35// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!cir.float x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}):36// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !cir.array<!cir.float x 5>, !cir.ptr<!cir.array<!cir.float x 5>>, ["openacc.private.init"]37// CHECK-NEXT: acc.yield38// CHECK-NEXT: }39//40// NoCopyConstruct[5] with 1 'bound'41// CHECK-NEXT: acc.private.recipe @privatization__Bcnt1__ZTSA5_15NoCopyConstruct : !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> init {42// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}):43// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !cir.array<!rec_NoCopyConstruct x 5>, !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>, ["openacc.private.init"]44// CHECK-NEXT: acc.yield45// CHECK-NEXT: }46 47void acc_compute() {48 // CHECK: cir.func{{.*}} @acc_compute()49 50 int someInt;51 // CHECK-NEXT: %[[SOMEINT:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["someInt"]52 float someFloat;53 // CHECK-NEXT: %[[SOMEFLOAT:.*]] = cir.alloca !cir.float, !cir.ptr<!cir.float>, ["someFloat"]54 struct NoCopyConstruct noCopy;55 // CHECK-NEXT: %[[NOCOPY:.*]] = cir.alloca !rec_NoCopyConstruct, !cir.ptr<!rec_NoCopyConstruct>, ["noCopy"]56 int someIntArr[5];57 // CHECK-NEXT: %[[INTARR:.*]] = cir.alloca !cir.array<!s32i x 5>, !cir.ptr<!cir.array<!s32i x 5>>, ["someIntArr"]58 float someFloatArr[5];59 // CHECK-NEXT: %[[FLOATARR:.*]] = cir.alloca !cir.array<!cir.float x 5>, !cir.ptr<!cir.array<!cir.float x 5>>, ["someFloatArr"]60 struct NoCopyConstruct noCopyArr[5];61 // CHECK-NEXT: %[[NOCOPYARR:.*]] = cir.alloca !cir.array<!rec_NoCopyConstruct x 5>, !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>, ["noCopyArr"]62 63#pragma acc parallel private(someInt)64 ;65 // CHECK: %[[PRIVATE:.*]] = acc.private varPtr(%[[SOMEINT]] : !cir.ptr<!s32i>) recipe(@privatization__ZTSi) -> !cir.ptr<!s32i> {name = "someInt"}66 // CHECK-NEXT: acc.parallel private(%[[PRIVATE]] : !cir.ptr<!s32i>)67 // CHECK-NEXT: acc.yield68 // CHECK-NEXT: } loc69#pragma acc serial private(someFloat)70 ;71 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[SOMEFLOAT]] : !cir.ptr<!cir.float>) recipe(@privatization__ZTSf) -> !cir.ptr<!cir.float> {name = "someFloat"}72 // CHECK-NEXT: acc.serial private(%[[PRIVATE]] : !cir.ptr<!cir.float>)73 // CHECK-NEXT: acc.yield74 // CHECK-NEXT: } loc75 76#pragma acc parallel private(noCopy)77 ;78 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[NOCOPY]] : !cir.ptr<!rec_NoCopyConstruct>) recipe(@privatization__ZTS15NoCopyConstruct) -> !cir.ptr<!rec_NoCopyConstruct> {name = "noCopy"}79 // CHECK-NEXT: acc.parallel private(%[[PRIVATE]] : !cir.ptr<!rec_NoCopyConstruct>80 // CHECK-NEXT: acc.yield81 // CHECK-NEXT: } loc82 83#pragma acc parallel private(someInt, someFloat, noCopy)84 ;85 // CHECK: %[[PRIVATE1:.*]] = acc.private varPtr(%[[SOMEINT]] : !cir.ptr<!s32i>) recipe(@privatization__ZTSi) -> !cir.ptr<!s32i> {name = "someInt"}86 // CHECK-NEXT: %[[PRIVATE2:.*]] = acc.private varPtr(%[[SOMEFLOAT]] : !cir.ptr<!cir.float>) recipe(@privatization__ZTSf) -> !cir.ptr<!cir.float> {name = "someFloat"}87 // CHECK-NEXT: %[[PRIVATE3:.*]] = acc.private varPtr(%[[NOCOPY]] : !cir.ptr<!rec_NoCopyConstruct>) recipe(@privatization__ZTS15NoCopyConstruct) -> !cir.ptr<!rec_NoCopyConstruct> {name = "noCopy"}88 // CHECK-NEXT: acc.parallel private(%[[PRIVATE1]], %[[PRIVATE2]], %[[PRIVATE3]] : !cir.ptr<!s32i>,89 // CHECK-SAME: !cir.ptr<!cir.float>,90 // CHECK-SAME: !cir.ptr<!rec_NoCopyConstruct>)91 // CHECK-NEXT: acc.yield92 // CHECK-NEXT: } loc93 94#pragma acc serial private(someIntArr[1])95 ;96 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i97 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si3298 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 199 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0100 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1101 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)102 // 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]"}103 // CHECK-NEXT: acc.serial private(%[[PRIVATE]] : !cir.ptr<!cir.array<!s32i x 5>>)104 // CHECK-NEXT: acc.yield105 // CHECK-NEXT: } loc106#pragma acc parallel private(someFloatArr[1])107 ;108 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i109 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32110 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1111 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0112 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1113 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)114 // 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]"}115 // CHECK-NEXT: acc.parallel private(%[[PRIVATE]] : !cir.ptr<!cir.array<!cir.float x 5>>)116 // CHECK-NEXT: acc.yield117 // CHECK-NEXT: } loc118#pragma acc serial private(noCopyArr[1])119 ;120 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i121 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32122 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1123 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0124 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1125 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)126 // 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]"}127 // CHECK-NEXT: acc.serial private(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>)128 // CHECK-NEXT: acc.yield129 // CHECK-NEXT: } loc130#pragma acc serial private(someIntArr[1], someFloatArr[1], noCopyArr[1])131 ;132 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i133 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32134 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1135 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0136 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1137 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)138 // 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]"}139 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i140 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32141 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1142 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0143 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1144 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)145 // 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]"}146 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i147 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32148 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1149 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0150 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1151 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)152 // 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]"}153 // CHECK-NEXT: acc.serial private(%[[PRIVATE1]], %[[PRIVATE2]], %[[PRIVATE3]] : !cir.ptr<!cir.array<!s32i x 5>>,154 // CHECK-SAME: !cir.ptr<!cir.array<!cir.float x 5>>,155 // CHECK-SAME: !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>)156 // CHECK-NEXT: acc.yield157 // CHECK-NEXT: } loc158 159#pragma acc parallel private(someIntArr[1:1])160 ;161 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i162 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32163 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i164 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32165 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0166 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1167 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)168 // 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]"}169 // CHECK-NEXT: acc.parallel private(%[[PRIVATE]] : !cir.ptr<!cir.array<!s32i x 5>>)170 // CHECK-NEXT: acc.yield171 // CHECK-NEXT: } loc172#pragma acc serial private(someFloatArr[1:1])173 ;174 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i175 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32176 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i177 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32178 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0179 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1180 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)181 // 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]"}182 // CHECK-NEXT: acc.serial private(%[[PRIVATE]] : !cir.ptr<!cir.array<!cir.float x 5>>)183 // CHECK-NEXT: acc.yield184 // CHECK-NEXT: } loc185#pragma acc parallel private(noCopyArr[1:1])186 ;187 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i188 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32189 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i190 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32191 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0192 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1193 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)194 // 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]"}195 // CHECK-NEXT: acc.parallel private(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>)196 // CHECK-NEXT: acc.yield197 // CHECK-NEXT: } loc198#pragma acc parallel private(someIntArr[1:1], someFloatArr[1:1], noCopyArr[1:1])199 ;200 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i201 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32202 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i203 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32204 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0205 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1206 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)207 // 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]"}208 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i209 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32210 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i211 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32212 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0213 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1214 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)215 // 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]"}216 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i217 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32218 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i219 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32220 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0221 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1222 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)223 // 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]"}224 // CHECK-NEXT: acc.parallel private(%[[PRIVATE1]], %[[PRIVATE2]], %[[PRIVATE3]] : !cir.ptr<!cir.array<!s32i x 5>>,225 // CHECK-SAME: !cir.ptr<!cir.array<!cir.float x 5>>,226 // CHECK-SAME: !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>)227 // CHECK-NEXT: acc.yield228 // CHECK-NEXT: } loc229}230