735 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.firstprivate.recipe @firstprivatization__ZTSi : !cir.ptr<!s32i> init {19// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!s32i> {{.*}}):20// CHECK-NEXT: cir.alloca !s32i, !cir.ptr<!s32i>, ["openacc.firstprivate.init"]21// CHECK-NEXT: acc.yield22// CHECK-NEXT: } copy {23// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!s32i> {{.*}}, %[[ARG_TO:.*]]: !cir.ptr<!s32i> {{.*}}):24// CHECK-NEXT: %[[FROM_LOAD:.*]] = cir.load {{.*}}%[[ARG_FROM]] : !cir.ptr<!s32i>, !s32i25// CHECK-NEXT: cir.store {{.*}} %[[FROM_LOAD]], %[[ARG_TO]] : !s32i, !cir.ptr<!s32i>26// CHECK-NEXT: acc.yield27// CHECK-NEXT: }28//29// CHECK-NEXT: acc.firstprivate.recipe @firstprivatization__ZTSf : !cir.ptr<!cir.float> init {30// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.float> {{.*}}):31// CHECK-NEXT: cir.alloca !cir.float, !cir.ptr<!cir.float>, ["openacc.firstprivate.init"]32// CHECK-NEXT: acc.yield33// CHECK-NEXT: } copy {34// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!cir.float> {{.*}}, %[[ARG_TO:.*]]: !cir.ptr<!cir.float> {{.*}}):35// CHECK-NEXT: %[[FROM_LOAD:.*]] = cir.load {{.*}}%[[ARG_FROM]] : !cir.ptr<!cir.float>, !cir.float36// CHECK-NEXT: cir.store {{.*}} %[[FROM_LOAD]], %[[ARG_TO]] : !cir.float, !cir.ptr<!cir.float>37// CHECK-NEXT: acc.yield38// CHECK-NEXT: }39//40// CHECK-NEXT: acc.firstprivate.recipe @firstprivatization__ZTS15NoCopyConstruct : !cir.ptr<!rec_NoCopyConstruct> init {41// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_NoCopyConstruct> {{.*}}):42// CHECK-NEXT: cir.alloca !rec_NoCopyConstruct, !cir.ptr<!rec_NoCopyConstruct>, ["openacc.firstprivate.init"]43// CHECK-NEXT: acc.yield44// CHECK-NEXT: } copy {45// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!rec_NoCopyConstruct> {{.*}}, %[[ARG_TO:.*]]: !cir.ptr<!rec_NoCopyConstruct> {{.*}}):46// CHECK-NEXT: cir.call @_ZN15NoCopyConstructC1ERKS_(%[[ARG_TO]], %[[ARG_FROM]]) nothrow : (!cir.ptr<!rec_NoCopyConstruct>, !cir.ptr<!rec_NoCopyConstruct>) -> ()47// CHECK-NEXT: acc.yield48// CHECK-NEXT: }49//50// CHECK-NEXT: acc.firstprivate.recipe @firstprivatization__ZTS13CopyConstruct : !cir.ptr<!rec_CopyConstruct> init {51// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_CopyConstruct> {{.*}}):52// CHECK-NEXT: cir.alloca !rec_CopyConstruct, !cir.ptr<!rec_CopyConstruct>, ["openacc.firstprivate.init"]53// CHECK-NEXT: acc.yield54// CHECK-NEXT: } copy {55// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!rec_CopyConstruct> {{.*}}, %[[ARG_TO:.*]]: !cir.ptr<!rec_CopyConstruct> {{.*}}):56// CHECK-NEXT: cir.call @_ZN13CopyConstructC1ERKS_(%[[ARG_TO]], %[[ARG_FROM]]) : (!cir.ptr<!rec_CopyConstruct>, !cir.ptr<!rec_CopyConstruct>) -> ()57// CHECK-NEXT: acc.yield58// CHECK-NEXT: }59//60// CHECK-NEXT: acc.firstprivate.recipe @firstprivatization__ZTS14NonDefaultCtor : !cir.ptr<!rec_NonDefaultCtor> init {61// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_NonDefaultCtor> {{.*}}):62// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !rec_NonDefaultCtor, !cir.ptr<!rec_NonDefaultCtor>, ["openacc.firstprivate.init"]63// CHECK-NEXT: acc.yield64// CHECK-NEXT: } copy {65// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!rec_NonDefaultCtor> {{.*}}, %[[ARG_TO:.*]]: !cir.ptr<!rec_NonDefaultCtor> {{.*}}):66// CHECK-NEXT: cir.call @_ZN14NonDefaultCtorC1ERKS_(%[[ARG_TO]], %[[ARG_FROM]]) nothrow : (!cir.ptr<!rec_NonDefaultCtor>, !cir.ptr<!rec_NonDefaultCtor>) -> ()67// CHECK-NEXT: acc.yield68// CHECK-NEXT: }69//70// CHECK-NEXT: acc.firstprivate.recipe @firstprivatization__ZTS7HasDtor : !cir.ptr<!rec_HasDtor> init {71// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_HasDtor> {{.*}}):72// CHECK-NEXT: cir.alloca !rec_HasDtor, !cir.ptr<!rec_HasDtor>, ["openacc.firstprivate.init"]73// CHECK-NEXT: acc.yield74// CHECK-NEXT: } copy {75// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!rec_HasDtor> {{.*}}, %[[ARG_TO:.*]]: !cir.ptr<!rec_HasDtor> {{.*}}):76// CHECK-NEXT: cir.call @_ZN7HasDtorC1ERKS_(%[[ARG_TO]], %[[ARG_FROM]]) nothrow : (!cir.ptr<!rec_HasDtor>, !cir.ptr<!rec_HasDtor>) -> ()77// CHECK-NEXT: acc.yield78// CHECK-NEXT: } destroy {79// CHECK-NEXT: ^bb0(%[[ORIG:.*]]: !cir.ptr<!rec_HasDtor> {{.*}}, %[[ARG:.*]]: !cir.ptr<!rec_HasDtor> {{.*}}):80// CHECK-NEXT: cir.call @_ZN7HasDtorD1Ev(%[[ARG]]) nothrow : (!cir.ptr<!rec_HasDtor>) -> ()81// CHECK-NEXT: acc.yield82// CHECK-NEXT: }83//84// CHECK-NEXT: acc.firstprivate.recipe @firstprivatization__Bcnt1__ZTSA5_i : !cir.ptr<!cir.array<!s32i x 5>> init {85// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!s32i x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty{{.*}}):86// CHECK-NEXT: cir.alloca !cir.array<!s32i x 5>, !cir.ptr<!cir.array<!s32i x 5>>, ["openacc.firstprivate.init"]87// CHECK-NEXT: acc.yield88// CHECK-NEXT: } copy {89// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!cir.array<!s32i x 5>> {{.*}}, %[[ARG_TO:.*]]: !cir.ptr<!cir.array<!s32i x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty{{.*}}):90// CHECK-NEXT: cir.scope {91// CHECK-NEXT: %[[LB:.*]] = acc.get_lowerbound %[[BOUND1]] : (!acc.data_bounds_ty) -> index92// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]] : index to !u64i93// CHECK-NEXT: %[[UB:.*]] = acc.get_upperbound %[[BOUND1]] : (!acc.data_bounds_ty) -> index94// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]] : index to !u64i95// CHECK-NEXT: %[[ITR:.*]] = cir.alloca !u64i, !cir.ptr<!u64i>, ["iter"] {alignment = 8 : i64}96// CHECK-NEXT: cir.store %[[LB_CAST]], %[[ITR]] : !u64i, !cir.ptr<!u64i>97// CHECK-NEXT: cir.for : cond {98// CHECK-NEXT: %[[ITR_LOAD:.*]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i99// CHECK-NEXT: %[[COND:.*]] = cir.cmp(lt, %[[ITR_LOAD]], %[[UB_CAST]]) : !u64i, !cir.bool100// CHECK-NEXT: cir.condition(%[[COND]])101// CHECK-NEXT: } body {102// CHECK-NEXT: %[[ITR_LOAD:.*]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i103// CHECK-NEXT: %[[DECAY_FROM:.*]] = cir.cast array_to_ptrdecay %[[ARG_FROM]] : !cir.ptr<!cir.array<!s32i x 5>> -> !cir.ptr<!s32i>104// CHECK-NEXT: %[[STRIDE_FROM:.*]] = cir.ptr_stride %[[DECAY_FROM]], %[[ITR_LOAD]] : (!cir.ptr<!s32i>, !u64i) -> !cir.ptr<!s32i>105// CHECK-NEXT: %[[DECAY_TO:.*]] = cir.cast array_to_ptrdecay %[[ARG_TO]] : !cir.ptr<!cir.array<!s32i x 5>> -> !cir.ptr<!s32i>106// CHECK-NEXT: %[[STRIDE_TO:.*]] = cir.ptr_stride %[[DECAY_TO]], %[[ITR_LOAD]] : (!cir.ptr<!s32i>, !u64i) -> !cir.ptr<!s32i>107// CHECK-NEXT: %[[FROM_LOAD:.*]] = cir.load{{.*}} %[[STRIDE_FROM]] : !cir.ptr<!s32i>, !s32i108// CHECK-NEXT: cir.store{{.*}} %[[FROM_LOAD]], %[[STRIDE_TO]] : !s32i, !cir.ptr<!s32i>109// CHECK-NEXT: cir.yield110// CHECK-NEXT: } step {111// CHECK-NEXT: %[[ITR_LOAD]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i112// CHECK-NEXT: %[[INC:.*]] = cir.unary(inc, %[[ITR_LOAD]]) : !u64i, !u64i113// CHECK-NEXT: cir.store %[[INC]], %[[ITR]] : !u64i, !cir.ptr<!u64i>114// CHECK-NEXT: cir.yield115// CHECK-NEXT: }116// CHECK-NEXT: }117// CHECK-NEXT: acc.yield118// CHECK-NEXT: }119//120// CHECK-NEXT: acc.firstprivate.recipe @firstprivatization__Bcnt1__ZTSA5_f : !cir.ptr<!cir.array<!cir.float x 5>> init {121// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!cir.float x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty{{.*}}):122// CHECK-NEXT: cir.alloca !cir.array<!cir.float x 5>, !cir.ptr<!cir.array<!cir.float x 5>>, ["openacc.firstprivate.init"]123// CHECK-NEXT: acc.yield124// CHECK-NEXT: } copy {125// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!cir.array<!cir.float x 5>> {{.*}}, %[[ARG_TO:.*]]: !cir.ptr<!cir.array<!cir.float x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty{{.*}}):126// CHECK-NEXT: cir.scope {127// CHECK-NEXT: %[[LB:.*]] = acc.get_lowerbound %[[BOUND1]] : (!acc.data_bounds_ty) -> index128// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]] : index to !u64i129// CHECK-NEXT: %[[UB:.*]] = acc.get_upperbound %[[BOUND1]] : (!acc.data_bounds_ty) -> index130// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]] : index to !u64i131// CHECK-NEXT: %[[ITR:.*]] = cir.alloca !u64i, !cir.ptr<!u64i>, ["iter"] {alignment = 8 : i64}132// CHECK-NEXT: cir.store %[[LB_CAST]], %[[ITR]] : !u64i, !cir.ptr<!u64i>133// CHECK-NEXT: cir.for : cond {134// CHECK-NEXT: %[[ITR_LOAD:.*]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i135// CHECK-NEXT: %[[COND:.*]] = cir.cmp(lt, %[[ITR_LOAD]], %[[UB_CAST]]) : !u64i, !cir.bool136// CHECK-NEXT: cir.condition(%[[COND]])137// CHECK-NEXT: } body {138// CHECK-NEXT: %[[ITR_LOAD:.*]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i139// CHECK-NEXT: %[[DECAY_FROM:.*]] = cir.cast array_to_ptrdecay %[[ARG_FROM]] : !cir.ptr<!cir.array<!cir.float x 5>> -> !cir.ptr<!cir.float>140// CHECK-NEXT: %[[STRIDE_FROM:.*]] = cir.ptr_stride %[[DECAY_FROM]], %[[ITR_LOAD]] : (!cir.ptr<!cir.float>, !u64i) -> !cir.ptr<!cir.float>141// CHECK-NEXT: %[[DECAY_TO:.*]] = cir.cast array_to_ptrdecay %[[ARG_TO]] : !cir.ptr<!cir.array<!cir.float x 5>> -> !cir.ptr<!cir.float>142// CHECK-NEXT: %[[STRIDE_TO:.*]] = cir.ptr_stride %[[DECAY_TO]], %[[ITR_LOAD]] : (!cir.ptr<!cir.float>, !u64i) -> !cir.ptr<!cir.float>143// CHECK-NEXT: %[[FROM_LOAD:.*]] = cir.load{{.*}} %[[STRIDE_FROM]] : !cir.ptr<!cir.float>, !cir.float144// CHECK-NEXT: cir.store{{.*}} %[[FROM_LOAD]], %[[STRIDE_TO]] : !cir.float, !cir.ptr<!cir.float>145// CHECK-NEXT: cir.yield146// CHECK-NEXT: } step {147// CHECK-NEXT: %[[ITR_LOAD]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i148// CHECK-NEXT: %[[INC:.*]] = cir.unary(inc, %[[ITR_LOAD]]) : !u64i, !u64i149// CHECK-NEXT: cir.store %[[INC]], %[[ITR]] : !u64i, !cir.ptr<!u64i>150// CHECK-NEXT: cir.yield151// CHECK-NEXT: }152// CHECK-NEXT: }153// CHECK-NEXT: acc.yield154// CHECK-NEXT: }155//156// CHECK-NEXT: acc.firstprivate.recipe @firstprivatization__Bcnt1__ZTSA5_15NoCopyConstruct : !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> init {157// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty{{.*}}):158// CHECK-NEXT: cir.alloca !cir.array<!rec_NoCopyConstruct x 5>, !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>, ["openacc.firstprivate.init"]159// CHECK-NEXT: acc.yield160// CHECK-NEXT: } copy {161// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> {{.*}}, %[[ARG_TO:.*]]: !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty{{.*}}):162// CHECK-NEXT: cir.scope {163// CHECK-NEXT: %[[LB:.*]] = acc.get_lowerbound %[[BOUND1]] : (!acc.data_bounds_ty) -> index164// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]] : index to !u64i165// CHECK-NEXT: %[[UB:.*]] = acc.get_upperbound %[[BOUND1]] : (!acc.data_bounds_ty) -> index166// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]] : index to !u64i167// CHECK-NEXT: %[[ITR:.*]] = cir.alloca !u64i, !cir.ptr<!u64i>, ["iter"] {alignment = 8 : i64}168// CHECK-NEXT: cir.store %[[LB_CAST]], %[[ITR]] : !u64i, !cir.ptr<!u64i>169// CHECK-NEXT: cir.for : cond {170// CHECK-NEXT: %[[ITR_LOAD:.*]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i171// CHECK-NEXT: %[[COND:.*]] = cir.cmp(lt, %[[ITR_LOAD]], %[[UB_CAST]]) : !u64i, !cir.bool172// CHECK-NEXT: cir.condition(%[[COND]])173// CHECK-NEXT: } body {174// CHECK-NEXT: %[[ITR_LOAD:.*]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i175// CHECK-NEXT: %[[DECAY_FROM:.*]] = cir.cast array_to_ptrdecay %[[ARG_FROM]] : !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> -> !cir.ptr<!rec_NoCopyConstruct>176// CHECK-NEXT: %[[STRIDE_FROM:.*]] = cir.ptr_stride %[[DECAY_FROM]], %[[ITR_LOAD]] : (!cir.ptr<!rec_NoCopyConstruct>, !u64i) -> !cir.ptr<!rec_NoCopyConstruct>177// CHECK-NEXT: %[[DECAY_TO:.*]] = cir.cast array_to_ptrdecay %[[ARG_TO]] : !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> -> !cir.ptr<!rec_NoCopyConstruct>178// CHECK-NEXT: %[[STRIDE_TO:.*]] = cir.ptr_stride %[[DECAY_TO]], %[[ITR_LOAD]] : (!cir.ptr<!rec_NoCopyConstruct>, !u64i) -> !cir.ptr<!rec_NoCopyConstruct>179// CHECK-NEXT: cir.call @_ZN15NoCopyConstructC1ERKS_(%[[STRIDE_TO]], %[[STRIDE_FROM]]) nothrow : (!cir.ptr<!rec_NoCopyConstruct>, !cir.ptr<!rec_NoCopyConstruct>) -> ()180// CHECK-NEXT: cir.yield181// CHECK-NEXT: } step {182// CHECK-NEXT: %[[ITR_LOAD]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i183// CHECK-NEXT: %[[INC:.*]] = cir.unary(inc, %[[ITR_LOAD]]) : !u64i, !u64i184// CHECK-NEXT: cir.store %[[INC]], %[[ITR]] : !u64i, !cir.ptr<!u64i>185// CHECK-NEXT: cir.yield186// CHECK-NEXT: }187// CHECK-NEXT: }188// CHECK-NEXT: acc.yield189// CHECK-NEXT: }190//191// CHECK-NEXT: acc.firstprivate.recipe @firstprivatization__Bcnt1__ZTSA5_13CopyConstruct : !cir.ptr<!cir.array<!rec_CopyConstruct x 5>> init {192// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!rec_CopyConstruct x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty{{.*}}):193// CHECK-NEXT: cir.alloca !cir.array<!rec_CopyConstruct x 5>, !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>, ["openacc.firstprivate.init"]194// CHECK-NEXT: acc.yield195// CHECK-NEXT: } copy {196// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!cir.array<!rec_CopyConstruct x 5>> {{.*}}, %[[ARG_TO:.*]]: !cir.ptr<!cir.array<!rec_CopyConstruct x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty{{.*}}):197// CHECK-NEXT: cir.scope {198// CHECK-NEXT: %[[LB:.*]] = acc.get_lowerbound %[[BOUND1]] : (!acc.data_bounds_ty) -> index199// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]] : index to !u64i200// CHECK-NEXT: %[[UB:.*]] = acc.get_upperbound %[[BOUND1]] : (!acc.data_bounds_ty) -> index201// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]] : index to !u64i202// CHECK-NEXT: %[[ITR:.*]] = cir.alloca !u64i, !cir.ptr<!u64i>, ["iter"] {alignment = 8 : i64}203// CHECK-NEXT: cir.store %[[LB_CAST]], %[[ITR]] : !u64i, !cir.ptr<!u64i>204// CHECK-NEXT: cir.for : cond {205// CHECK-NEXT: %[[ITR_LOAD:.*]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i206// CHECK-NEXT: %[[COND:.*]] = cir.cmp(lt, %[[ITR_LOAD]], %[[UB_CAST]]) : !u64i, !cir.bool207// CHECK-NEXT: cir.condition(%[[COND]])208// CHECK-NEXT: } body {209// CHECK-NEXT: %[[ITR_LOAD:.*]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i210// CHECK-NEXT: %[[DECAY_FROM:.*]] = cir.cast array_to_ptrdecay %[[ARG_FROM]] : !cir.ptr<!cir.array<!rec_CopyConstruct x 5>> -> !cir.ptr<!rec_CopyConstruct>211// CHECK-NEXT: %[[STRIDE_FROM:.*]] = cir.ptr_stride %[[DECAY_FROM]], %[[ITR_LOAD]] : (!cir.ptr<!rec_CopyConstruct>, !u64i) -> !cir.ptr<!rec_CopyConstruct>212// CHECK-NEXT: %[[DECAY_TO:.*]] = cir.cast array_to_ptrdecay %[[ARG_TO]] : !cir.ptr<!cir.array<!rec_CopyConstruct x 5>> -> !cir.ptr<!rec_CopyConstruct>213// CHECK-NEXT: %[[STRIDE_TO:.*]] = cir.ptr_stride %[[DECAY_TO]], %[[ITR_LOAD]] : (!cir.ptr<!rec_CopyConstruct>, !u64i) -> !cir.ptr<!rec_CopyConstruct>214// CHECK-NEXT: cir.call @_ZN13CopyConstructC1ERKS_(%[[STRIDE_TO]], %[[STRIDE_FROM]]) : (!cir.ptr<!rec_CopyConstruct>, !cir.ptr<!rec_CopyConstruct>) -> ()215// CHECK-NEXT: cir.yield216// CHECK-NEXT: } step {217// CHECK-NEXT: %[[ITR_LOAD]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i218// CHECK-NEXT: %[[INC:.*]] = cir.unary(inc, %[[ITR_LOAD]]) : !u64i, !u64i219// CHECK-NEXT: cir.store %[[INC]], %[[ITR]] : !u64i, !cir.ptr<!u64i>220// CHECK-NEXT: cir.yield221// CHECK-NEXT: }222// CHECK-NEXT: }223// CHECK-NEXT: acc.yield224// CHECK-NEXT: }225//226// CHECK-NEXT: acc.firstprivate.recipe @firstprivatization__Bcnt1__ZTSA5_14NonDefaultCtor : !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> init {227// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty{{.*}}):228// CHECK-NEXT: cir.alloca !cir.array<!rec_NonDefaultCtor x 5>, !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>, ["openacc.firstprivate.init"]229// CHECK-NEXT: acc.yield230// CHECK-NEXT: } copy {231// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> {{.*}}, %[[ARG_TO:.*]]: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty{{.*}}):232// CHECK-NEXT: cir.scope {233// CHECK-NEXT: %[[LB:.*]] = acc.get_lowerbound %[[BOUND1]] : (!acc.data_bounds_ty) -> index234// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]] : index to !u64i235// CHECK-NEXT: %[[UB:.*]] = acc.get_upperbound %[[BOUND1]] : (!acc.data_bounds_ty) -> index236// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]] : index to !u64i237// CHECK-NEXT: %[[ITR:.*]] = cir.alloca !u64i, !cir.ptr<!u64i>, ["iter"] {alignment = 8 : i64}238// CHECK-NEXT: cir.store %[[LB_CAST]], %[[ITR]] : !u64i, !cir.ptr<!u64i>239// CHECK-NEXT: cir.for : cond {240// CHECK-NEXT: %[[ITR_LOAD:.*]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i241// CHECK-NEXT: %[[COND:.*]] = cir.cmp(lt, %[[ITR_LOAD]], %[[UB_CAST]]) : !u64i, !cir.bool242// CHECK-NEXT: cir.condition(%[[COND]])243// CHECK-NEXT: } body {244// CHECK-NEXT: %[[ITR_LOAD:.*]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i245// CHECK-NEXT: %[[DECAY_FROM:.*]] = cir.cast array_to_ptrdecay %[[ARG_FROM]] : !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> -> !cir.ptr<!rec_NonDefaultCtor>246// CHECK-NEXT: %[[STRIDE_FROM:.*]] = cir.ptr_stride %[[DECAY_FROM]], %[[ITR_LOAD]] : (!cir.ptr<!rec_NonDefaultCtor>, !u64i) -> !cir.ptr<!rec_NonDefaultCtor>247// CHECK-NEXT: %[[DECAY_TO:.*]] = cir.cast array_to_ptrdecay %[[ARG_TO]] : !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> -> !cir.ptr<!rec_NonDefaultCtor>248// CHECK-NEXT: %[[STRIDE_TO:.*]] = cir.ptr_stride %[[DECAY_TO]], %[[ITR_LOAD]] : (!cir.ptr<!rec_NonDefaultCtor>, !u64i) -> !cir.ptr<!rec_NonDefaultCtor>249// CHECK-NEXT: cir.call @_ZN14NonDefaultCtorC1ERKS_(%[[STRIDE_TO]], %[[STRIDE_FROM]]) nothrow : (!cir.ptr<!rec_NonDefaultCtor>, !cir.ptr<!rec_NonDefaultCtor>) -> ()250// CHECK-NEXT: cir.yield251// CHECK-NEXT: } step {252// CHECK-NEXT: %[[ITR_LOAD]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i253// CHECK-NEXT: %[[INC:.*]] = cir.unary(inc, %[[ITR_LOAD]]) : !u64i, !u64i254// CHECK-NEXT: cir.store %[[INC]], %[[ITR]] : !u64i, !cir.ptr<!u64i>255// CHECK-NEXT: cir.yield256// CHECK-NEXT: }257// CHECK-NEXT: }258// CHECK-NEXT: acc.yield259// CHECK-NEXT: }260//261// CHECK-NEXT: acc.firstprivate.recipe @firstprivatization__Bcnt1__ZTSA5_7HasDtor : !cir.ptr<!cir.array<!rec_HasDtor x 5>> init {262// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!rec_HasDtor x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty{{.*}}):263// CHECK-NEXT: cir.alloca !cir.array<!rec_HasDtor x 5>, !cir.ptr<!cir.array<!rec_HasDtor x 5>>, ["openacc.firstprivate.init"]264// CHECK-NEXT: acc.yield265// CHECK-NEXT: } copy {266// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!cir.array<!rec_HasDtor x 5>> {{.*}}, %[[ARG_TO:.*]]: !cir.ptr<!cir.array<!rec_HasDtor x 5>> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty{{.*}}):267// CHECK-NEXT: cir.scope {268// CHECK-NEXT: %[[LB:.*]] = acc.get_lowerbound %[[BOUND1]] : (!acc.data_bounds_ty) -> index269// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]] : index to !u64i270// CHECK-NEXT: %[[UB:.*]] = acc.get_upperbound %[[BOUND1]] : (!acc.data_bounds_ty) -> index271// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]] : index to !u64i272// CHECK-NEXT: %[[ITR:.*]] = cir.alloca !u64i, !cir.ptr<!u64i>, ["iter"] {alignment = 8 : i64}273// CHECK-NEXT: cir.store %[[LB_CAST]], %[[ITR]] : !u64i, !cir.ptr<!u64i>274// CHECK-NEXT: cir.for : cond {275// CHECK-NEXT: %[[ITR_LOAD:.*]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i276// CHECK-NEXT: %[[COND:.*]] = cir.cmp(lt, %[[ITR_LOAD]], %[[UB_CAST]]) : !u64i, !cir.bool277// CHECK-NEXT: cir.condition(%[[COND]])278// CHECK-NEXT: } body {279// CHECK-NEXT: %[[ITR_LOAD:.*]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i280// CHECK-NEXT: %[[DECAY_FROM:.*]] = cir.cast array_to_ptrdecay %[[ARG_FROM]] : !cir.ptr<!cir.array<!rec_HasDtor x 5>> -> !cir.ptr<!rec_HasDtor>281// CHECK-NEXT: %[[STRIDE_FROM:.*]] = cir.ptr_stride %[[DECAY_FROM]], %[[ITR_LOAD]] : (!cir.ptr<!rec_HasDtor>, !u64i) -> !cir.ptr<!rec_HasDtor>282// CHECK-NEXT: %[[DECAY_TO:.*]] = cir.cast array_to_ptrdecay %[[ARG_TO]] : !cir.ptr<!cir.array<!rec_HasDtor x 5>> -> !cir.ptr<!rec_HasDtor>283// CHECK-NEXT: %[[STRIDE_TO:.*]] = cir.ptr_stride %[[DECAY_TO]], %[[ITR_LOAD]] : (!cir.ptr<!rec_HasDtor>, !u64i) -> !cir.ptr<!rec_HasDtor>284// CHECK-NEXT: cir.call @_ZN7HasDtorC1ERKS_(%[[STRIDE_TO]], %[[STRIDE_FROM]]) nothrow : (!cir.ptr<!rec_HasDtor>, !cir.ptr<!rec_HasDtor>) -> ()285// CHECK-NEXT: cir.yield286// CHECK-NEXT: } step {287// CHECK-NEXT: %[[ITR_LOAD]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i288// CHECK-NEXT: %[[INC:.*]] = cir.unary(inc, %[[ITR_LOAD]]) : !u64i, !u64i289// CHECK-NEXT: cir.store %[[INC]], %[[ITR]] : !u64i, !cir.ptr<!u64i>290// CHECK-NEXT: cir.yield291// CHECK-NEXT: }292// CHECK-NEXT: }293// CHECK-NEXT: acc.yield294// CHECK-NEXT: } destroy {295// 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{{.*}}):296// CHECK-NEXT: cir.scope {297// CHECK-NEXT: %[[LB:.*]] = acc.get_lowerbound %[[BOUND1]] : (!acc.data_bounds_ty) -> index298// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]] : index to !u64i299// CHECK-NEXT: %[[UB:.*]] = acc.get_upperbound %[[BOUND1]] : (!acc.data_bounds_ty) -> index300// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]] : index to !u64i301// CHECK-NEXT: %[[ITR:.*]] = cir.alloca !u64i, !cir.ptr<!u64i>, ["iter"] {alignment = 8 : i64}302// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !u64i303// CHECK-NEXT: %[[LAST_SUB_ONE:.*]] = cir.binop(sub, %[[UB_CAST]], %[[ONE]]) : !u64i304// CHECK-NEXT: cir.store %[[LAST_SUB_ONE]], %[[ITR]] : !u64i, !cir.ptr<!u64i>305// CHECK-NEXT: cir.for : cond {306// CHECK-NEXT: %[[ITR_LOAD:.*]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i307// CHECK-NEXT: %[[COND:.*]] = cir.cmp(ge, %[[ITR_LOAD]], %[[LB_CAST]]) : !u64i, !cir.bool308// CHECK-NEXT: cir.condition(%[[COND]])309// CHECK-NEXT: } body {310// CHECK-NEXT: %[[ITR_LOAD:.*]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i311// CHECK-NEXT: %[[DECAY:.*]] = cir.cast array_to_ptrdecay %[[ARG]] : !cir.ptr<!cir.array<!rec_HasDtor x 5>> -> !cir.ptr<!rec_HasDtor>312// CHECK-NEXT: %[[STRIDE:.*]] = cir.ptr_stride %[[DECAY]], %[[ITR_LOAD]] : (!cir.ptr<!rec_HasDtor>, !u64i) -> !cir.ptr<!rec_HasDtor>313// CHECK-NEXT: cir.call @_ZN7HasDtorD1Ev(%[[STRIDE]]) nothrow : (!cir.ptr<!rec_HasDtor>) -> ()314// CHECK-NEXT: cir.yield315// CHECK-NEXT: } step {316// CHECK-NEXT: %[[ITR_LOAD]] = cir.load %[[ITR]] : !cir.ptr<!u64i>, !u64i317// CHECK-NEXT: %[[DEC:.*]] = cir.unary(dec, %[[ITR_LOAD]]) : !u64i, !u64i318// CHECK-NEXT: cir.store %[[DEC]], %[[ITR]] : !u64i, !cir.ptr<!u64i>319// CHECK-NEXT: cir.yield320// CHECK-NEXT: }321// CHECK-NEXT: }322// CHECK-NEXT: acc.yield323// CHECK-NEXT: }324//325 326extern "C" void acc_combined() {327 // CHECK: cir.func{{.*}} @acc_combined()328 329 int someInt;330 // CHECK-NEXT: %[[SOMEINT:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["someInt"]331 float someFloat;332 // CHECK-NEXT: %[[SOMEFLOAT:.*]] = cir.alloca !cir.float, !cir.ptr<!cir.float>, ["someFloat"]333 NoCopyConstruct noCopy;334 // CHECK-NEXT: %[[NOCOPY:.*]] = cir.alloca !rec_NoCopyConstruct, !cir.ptr<!rec_NoCopyConstruct>, ["noCopy"]335 CopyConstruct hasCopy;336 // CHECK-NEXT: %[[HASCOPY:.*]] = cir.alloca !rec_CopyConstruct, !cir.ptr<!rec_CopyConstruct>, ["hasCopy"]337 NonDefaultCtor notDefCtor;338 // CHECK-NEXT: %[[NOTDEFCTOR:.*]] = cir.alloca !rec_NonDefaultCtor, !cir.ptr<!rec_NonDefaultCtor>, ["notDefCtor", init]339 HasDtor dtor;340 // CHECK-NEXT: %[[DTOR:.*]] = cir.alloca !rec_HasDtor, !cir.ptr<!rec_HasDtor>, ["dtor"]341 int someIntArr[5];342 // CHECK-NEXT: %[[INTARR:.*]] = cir.alloca !cir.array<!s32i x 5>, !cir.ptr<!cir.array<!s32i x 5>>, ["someIntArr"]343 float someFloatArr[5];344 // CHECK-NEXT: %[[FLOATARR:.*]] = cir.alloca !cir.array<!cir.float x 5>, !cir.ptr<!cir.array<!cir.float x 5>>, ["someFloatArr"]345 NoCopyConstruct noCopyArr[5];346 // CHECK-NEXT: %[[NOCOPYARR:.*]] = cir.alloca !cir.array<!rec_NoCopyConstruct x 5>, !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>, ["noCopyArr"]347 CopyConstruct hasCopyArr[5];348 // CHECK-NEXT: %[[HASCOPYARR:.*]] = cir.alloca !cir.array<!rec_CopyConstruct x 5>, !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>, ["hasCopyArr"]349 NonDefaultCtor notDefCtorArr[5];350 // CHECK-NEXT: %[[NOTDEFCTORARR:.*]] = cir.alloca !cir.array<!rec_NonDefaultCtor x 5>, !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>, ["notDefCtorArr", init]351 HasDtor dtorArr[5];352 // CHECK-NEXT: %[[DTORARR:.*]] = cir.alloca !cir.array<!rec_HasDtor x 5>, !cir.ptr<!cir.array<!rec_HasDtor x 5>>, ["dtorArr"]353 // CHECK-NEXT: cir.call @_ZN14NonDefaultCtorC1Ev(%[[NOTDEFCTOR]]) : (!cir.ptr<!rec_NonDefaultCtor>) -> ()354 355#pragma acc parallel loop firstprivate(someInt)356 for(int i = 0; i < 5; ++i);357 // CHECK: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[SOMEINT]] : !cir.ptr<!s32i>) recipe(@firstprivatization__ZTSi) -> !cir.ptr<!s32i> {name = "someInt"}358 // CHECK-NEXT: acc.parallel combined(loop) firstprivate(%[[PRIVATE]] : !cir.ptr<!s32i>) {359 // CHECK-NEXT: acc.loop combined(parallel)360 // CHECK: acc.yield361 // CHECK-NEXT: } loc362 // CHECK-NEXT: acc.yield363 // CHECK-NEXT: } loc364#pragma acc serial loop firstprivate(someFloat)365 for(int i = 0; i < 5; ++i);366 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[SOMEFLOAT]] : !cir.ptr<!cir.float>) recipe(@firstprivatization__ZTSf) -> !cir.ptr<!cir.float> {name = "someFloat"}367 // CHECK-NEXT: acc.serial combined(loop) firstprivate(%[[PRIVATE]] : !cir.ptr<!cir.float>) {368 // CHECK-NEXT: acc.loop combined(serial)369 // CHECK: acc.yield370 // CHECK-NEXT: } loc371 // CHECK-NEXT: acc.yield372 // CHECK-NEXT: } loc373 374#pragma acc parallel loop firstprivate(noCopy)375 for(int i = 0; i < 5; ++i);376 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[NOCOPY]] : !cir.ptr<!rec_NoCopyConstruct>) recipe(@firstprivatization__ZTS15NoCopyConstruct) -> !cir.ptr<!rec_NoCopyConstruct> {name = "noCopy"}377 // CHECK-NEXT: acc.parallel combined(loop) firstprivate(%[[PRIVATE]] : !cir.ptr<!rec_NoCopyConstruct>) {378 // CHECK-NEXT: acc.loop combined(parallel)379 // CHECK: acc.yield380 // CHECK-NEXT: } loc381 // CHECK-NEXT: acc.yield382 // CHECK-NEXT: } loc383#pragma acc serial loop firstprivate(hasCopy)384 for(int i = 0; i < 5; ++i);385 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[HASCOPY]] : !cir.ptr<!rec_CopyConstruct>) recipe(@firstprivatization__ZTS13CopyConstruct) -> !cir.ptr<!rec_CopyConstruct> {name = "hasCopy"}386 // CHECK-NEXT: acc.serial combined(loop) firstprivate(%[[PRIVATE]] : !cir.ptr<!rec_CopyConstruct>) {387 // CHECK-NEXT: acc.loop combined(serial)388 // CHECK: acc.yield389 // CHECK-NEXT: } loc390 // CHECK-NEXT: acc.yield391 // CHECK-NEXT: } loc392#pragma acc serial loop firstprivate(notDefCtor)393 for(int i = 0; i < 5; ++i);394 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[NOTDEFCTOR]] : !cir.ptr<!rec_NonDefaultCtor>) recipe(@firstprivatization__ZTS14NonDefaultCtor) -> !cir.ptr<!rec_NonDefaultCtor> {name = "notDefCtor"}395 // CHECK-NEXT: acc.serial combined(loop) firstprivate(%[[PRIVATE]] : !cir.ptr<!rec_NonDefaultCtor>) {396 // CHECK-NEXT: acc.loop combined(serial)397 // CHECK: acc.yield398 // CHECK-NEXT: } loc399 // CHECK-NEXT: acc.yield400 // CHECK-NEXT: } loc401#pragma acc serial loop firstprivate(dtor)402 for(int i = 0; i < 5; ++i);403 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[DTOR]] : !cir.ptr<!rec_HasDtor>) recipe(@firstprivatization__ZTS7HasDtor) -> !cir.ptr<!rec_HasDtor> {name = "dtor"}404 // CHECK-NEXT: acc.serial combined(loop) firstprivate(%[[PRIVATE]] : !cir.ptr<!rec_HasDtor>) {405 // CHECK-NEXT: acc.loop combined(serial)406 // CHECK: acc.yield407 // CHECK-NEXT: } loc408 // CHECK-NEXT: acc.yield409 // CHECK-NEXT: } loc410 411#pragma acc parallel loop firstprivate(someInt, someFloat, noCopy, hasCopy, notDefCtor, dtor)412 for(int i = 0; i < 5; ++i);413 // CHECK: %[[PRIVATE1:.*]] = acc.firstprivate varPtr(%[[SOMEINT]] : !cir.ptr<!s32i>) recipe(@firstprivatization__ZTSi) -> !cir.ptr<!s32i> {name = "someInt"}414 // CHECK-NEXT: %[[PRIVATE2:.*]] = acc.firstprivate varPtr(%[[SOMEFLOAT]] : !cir.ptr<!cir.float>) recipe(@firstprivatization__ZTSf) -> !cir.ptr<!cir.float> {name = "someFloat"}415 // CHECK-NEXT: %[[PRIVATE3:.*]] = acc.firstprivate varPtr(%[[NOCOPY]] : !cir.ptr<!rec_NoCopyConstruct>) recipe(@firstprivatization__ZTS15NoCopyConstruct) -> !cir.ptr<!rec_NoCopyConstruct> {name = "noCopy"}416 // CHECK-NEXT: %[[PRIVATE4:.*]] = acc.firstprivate varPtr(%[[HASCOPY]] : !cir.ptr<!rec_CopyConstruct>) recipe(@firstprivatization__ZTS13CopyConstruct) -> !cir.ptr<!rec_CopyConstruct> {name = "hasCopy"}417 // CHECK-NEXT: %[[PRIVATE5:.*]] = acc.firstprivate varPtr(%[[NOTDEFCTOR]] : !cir.ptr<!rec_NonDefaultCtor>) recipe(@firstprivatization__ZTS14NonDefaultCtor) -> !cir.ptr<!rec_NonDefaultCtor> {name = "notDefCtor"}418 // CHECK-NEXT: %[[PRIVATE6:.*]] = acc.firstprivate varPtr(%[[DTOR]] : !cir.ptr<!rec_HasDtor>) recipe(@firstprivatization__ZTS7HasDtor) -> !cir.ptr<!rec_HasDtor> {name = "dtor"}419 // CHECK: acc.parallel combined(loop) firstprivate(%[[PRIVATE1]], %[[PRIVATE2]], %[[PRIVATE3]], %[[PRIVATE4]], %[[PRIVATE5]], %[[PRIVATE6]] : !cir.ptr<!s32i>,420 // CHECK-SAME: !cir.ptr<!cir.float>,421 // CHECK-SAME: !cir.ptr<!rec_NoCopyConstruct>,422 // CHECK-SAME: !cir.ptr<!rec_CopyConstruct>,423 // CHECK-SAME: !cir.ptr<!rec_NonDefaultCtor>,424 // CHECK-SAME: !cir.ptr<!rec_HasDtor>)425 // CHECK-NEXT: acc.loop combined(parallel)426 // CHECK: acc.yield427 // CHECK-NEXT: } loc428 // CHECK-NEXT: acc.yield429 // CHECK-NEXT: } loc430 431#pragma acc serial loop firstprivate(someIntArr[1])432 for(int i = 0; i < 5; ++i);433 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i434 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32435 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1436 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0437 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1438 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)439 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[INTARR]] : !cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS]]) recipe(@firstprivatization__Bcnt1__ZTSA5_i) -> !cir.ptr<!cir.array<!s32i x 5>> {name = "someIntArr[1]"}440 // CHECK-NEXT: acc.serial combined(loop) firstprivate(%[[PRIVATE]] : !cir.ptr<!cir.array<!s32i x 5>>) {441 // CHECK-NEXT: acc.loop combined(serial)442 // CHECK: acc.yield443 // CHECK-NEXT: } loc444 // CHECK-NEXT: acc.yield445 // CHECK-NEXT: } loc446#pragma acc parallel loop firstprivate(someFloatArr[1])447 for(int i = 0; i < 5; ++i);448 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i449 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32450 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1451 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0452 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1453 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)454 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[FLOATARR]] : !cir.ptr<!cir.array<!cir.float x 5>>) bounds(%[[BOUNDS]]) recipe(@firstprivatization__Bcnt1__ZTSA5_f) -> !cir.ptr<!cir.array<!cir.float x 5>> {name = "someFloatArr[1]"}455 // CHECK-NEXT: acc.parallel combined(loop) firstprivate(%[[PRIVATE]] : !cir.ptr<!cir.array<!cir.float x 5>>) {456 // CHECK-NEXT: acc.loop combined(parallel)457 // CHECK: acc.yield458 // CHECK-NEXT: } loc459 // CHECK-NEXT: acc.yield460 // CHECK-NEXT: } loc461#pragma acc serial loop firstprivate(noCopyArr[1])462 for(int i = 0; i < 5; ++i);463 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i464 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32465 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1466 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0467 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1468 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)469 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[NOCOPYARR]] : !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>) bounds(%[[BOUNDS]]) recipe(@firstprivatization__Bcnt1__ZTSA5_15NoCopyConstruct) -> !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> {name = "noCopyArr[1]"}470 // CHECK-NEXT: acc.serial combined(loop) firstprivate(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>) {471 // CHECK-NEXT: acc.loop combined(serial)472 // CHECK: acc.yield473 // CHECK-NEXT: } loc474 // CHECK-NEXT: acc.yield475 // CHECK-NEXT: } loc476#pragma acc parallel loop firstprivate(hasCopyArr[1])477 for(int i = 0; i < 5; ++i);478 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i479 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32480 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1481 // 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_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)484 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[HASCOPYARR]] : !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>) bounds(%[[BOUNDS]]) recipe(@firstprivatization__Bcnt1__ZTSA5_13CopyConstruct) -> !cir.ptr<!cir.array<!rec_CopyConstruct x 5>> {name = "hasCopyArr[1]"}485 // CHECK-NEXT: acc.parallel combined(loop) firstprivate(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>) {486 // CHECK-NEXT: acc.loop combined(parallel)487 // CHECK: acc.yield488 // CHECK-NEXT: } loc489 // CHECK-NEXT: acc.yield490 // CHECK-NEXT: } loc491#pragma acc parallel loop firstprivate(notDefCtorArr[1])492 for(int i = 0; i < 5; ++i);493 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i494 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32495 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1496 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0497 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1498 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)499 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[NOTDEFCTORARR]] : !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>) bounds(%[[BOUNDS]]) recipe(@firstprivatization__Bcnt1__ZTSA5_14NonDefaultCtor) -> !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> {name = "notDefCtorArr[1]"}500 // CHECK-NEXT: acc.parallel combined(loop) firstprivate(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>) {501 // CHECK-NEXT: acc.loop combined(parallel)502 // CHECK: acc.yield503 // CHECK-NEXT: } loc504 // CHECK-NEXT: acc.yield505 // CHECK-NEXT: } loc506#pragma acc parallel loop firstprivate(dtorArr[1])507 for(int i = 0; i < 5; ++i);508 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i509 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32510 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1511 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0512 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1513 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)514 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[DTORARR]] : !cir.ptr<!cir.array<!rec_HasDtor x 5>>) bounds(%[[BOUNDS]]) recipe(@firstprivatization__Bcnt1__ZTSA5_7HasDtor) -> !cir.ptr<!cir.array<!rec_HasDtor x 5>> {name = "dtorArr[1]"}515 // CHECK-NEXT: acc.parallel combined(loop) firstprivate(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_HasDtor x 5>>) {516 // CHECK-NEXT: acc.loop combined(parallel)517 // CHECK: acc.yield518 // CHECK-NEXT: } loc519 // CHECK-NEXT: acc.yield520 // CHECK-NEXT: } loc521#pragma acc serial loop firstprivate(someIntArr[1], someFloatArr[1], noCopyArr[1], hasCopyArr[1], notDefCtorArr[1], dtorArr[1])522 for(int i = 0; i < 5; ++i);523 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i524 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32525 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1526 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0527 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1528 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)529 // CHECK-NEXT: %[[PRIVATE1:.*]] = acc.firstprivate varPtr(%[[INTARR]] : !cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS]]) recipe(@firstprivatization__Bcnt1__ZTSA5_i) -> !cir.ptr<!cir.array<!s32i x 5>> {name = "someIntArr[1]"}530 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i531 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32532 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1533 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0534 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1535 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)536 // CHECK-NEXT: %[[PRIVATE2:.*]] = acc.firstprivate varPtr(%[[FLOATARR]] : !cir.ptr<!cir.array<!cir.float x 5>>) bounds(%[[BOUNDS]]) recipe(@firstprivatization__Bcnt1__ZTSA5_f) -> !cir.ptr<!cir.array<!cir.float x 5>> {name = "someFloatArr[1]"}537 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i538 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32539 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1540 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0541 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1542 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)543 // CHECK-NEXT: %[[PRIVATE3:.*]] = acc.firstprivate varPtr(%[[NOCOPYARR]] : !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>) bounds(%[[BOUNDS]]) recipe(@firstprivatization__Bcnt1__ZTSA5_15NoCopyConstruct) -> !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> {name = "noCopyArr[1]"}544 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i545 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32546 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1547 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0548 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1549 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)550 // CHECK-NEXT: %[[PRIVATE4:.*]] = acc.firstprivate varPtr(%[[HASCOPYARR]] : !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>) bounds(%[[BOUNDS]]) recipe(@firstprivatization__Bcnt1__ZTSA5_13CopyConstruct) -> !cir.ptr<!cir.array<!rec_CopyConstruct x 5>> {name = "hasCopyArr[1]"}551 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i552 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32553 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1554 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0555 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1556 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)557 // CHECK-NEXT: %[[PRIVATE5:.*]] = acc.firstprivate varPtr(%[[NOTDEFCTORARR]] : !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>) bounds(%[[BOUNDS]]) recipe(@firstprivatization__Bcnt1__ZTSA5_14NonDefaultCtor) -> !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> {name = "notDefCtorArr[1]"}558 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i559 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32560 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1561 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0562 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1563 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)564 // CHECK-NEXT: %[[PRIVATE6:.*]] = acc.firstprivate varPtr(%[[DTORARR]] : !cir.ptr<!cir.array<!rec_HasDtor x 5>>) bounds(%[[BOUNDS]]) recipe(@firstprivatization__Bcnt1__ZTSA5_7HasDtor) -> !cir.ptr<!cir.array<!rec_HasDtor x 5>> {name = "dtorArr[1]"}565 // CHECK-NEXT: acc.serial combined(loop) firstprivate(%[[PRIVATE1]], %[[PRIVATE2]], %[[PRIVATE3]], %[[PRIVATE4]], %[[PRIVATE5]], %[[PRIVATE6]] : !cir.ptr<!cir.array<!s32i x 5>>,566 // CHECK-SAME: !cir.ptr<!cir.array<!cir.float x 5>>,567 // CHECK-SAME: !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>,568 // CHECK-SAME: !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>,569 // CHECK-SAME: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>,570 // CHECK-SAME: !cir.ptr<!cir.array<!rec_HasDtor x 5>>)571 // CHECK-NEXT: acc.loop combined(serial)572 // CHECK: acc.yield573 // CHECK-NEXT: } loc574 // CHECK-NEXT: acc.yield575 // CHECK-NEXT: } loc576 577#pragma acc parallel loop firstprivate(someIntArr[1:1])578 for(int i = 0; i < 5; ++i);579 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i580 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32581 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i582 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32583 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0584 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1585 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)586 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[INTARR]] : !cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS]]) recipe(@firstprivatization__Bcnt1__ZTSA5_i) -> !cir.ptr<!cir.array<!s32i x 5>> {name = "someIntArr[1:1]"}587 // CHECK-NEXT: acc.parallel combined(loop) firstprivate(%[[PRIVATE]] : !cir.ptr<!cir.array<!s32i x 5>>) {588 // CHECK-NEXT: acc.loop combined(parallel)589 // CHECK: acc.yield590 // CHECK-NEXT: } loc591 // CHECK-NEXT: acc.yield592 // CHECK-NEXT: } loc593#pragma acc serial loop firstprivate(someFloatArr[1:1])594 for(int i = 0; i < 5; ++i);595 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i596 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32597 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i598 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32599 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0600 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1601 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)602 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[FLOATARR]] : !cir.ptr<!cir.array<!cir.float x 5>>) bounds(%[[BOUNDS]]) recipe(@firstprivatization__Bcnt1__ZTSA5_f) -> !cir.ptr<!cir.array<!cir.float x 5>> {name = "someFloatArr[1:1]"}603 // CHECK-NEXT: acc.serial combined(loop) firstprivate(%[[PRIVATE]] : !cir.ptr<!cir.array<!cir.float x 5>>) {604 // CHECK-NEXT: acc.loop combined(serial)605 // CHECK: acc.yield606 // CHECK-NEXT: } loc607 // CHECK-NEXT: acc.yield608 // CHECK-NEXT: } loc609#pragma acc parallel loop firstprivate(noCopyArr[1:1])610 for(int i = 0; i < 5; ++i);611 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i612 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32613 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i614 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32615 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0616 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1617 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)618 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[NOCOPYARR]] : !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>) bounds(%[[BOUNDS]]) recipe(@firstprivatization__Bcnt1__ZTSA5_15NoCopyConstruct) -> !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> {name = "noCopyArr[1:1]"}619 // CHECK-NEXT: acc.parallel combined(loop) firstprivate(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>) {620 // CHECK-NEXT: acc.loop combined(parallel)621 // CHECK: acc.yield622 // CHECK-NEXT: } loc623 // CHECK-NEXT: acc.yield624 // CHECK-NEXT: } loc625#pragma acc serial loop firstprivate(hasCopyArr[1:1])626 for(int i = 0; i < 5; ++i);627 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i628 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32629 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i630 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32631 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0632 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1633 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)634 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[HASCOPYARR]] : !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>) bounds(%[[BOUNDS]]) recipe(@firstprivatization__Bcnt1__ZTSA5_13CopyConstruct) -> !cir.ptr<!cir.array<!rec_CopyConstruct x 5>> {name = "hasCopyArr[1:1]"}635 // CHECK-NEXT: acc.serial combined(loop) firstprivate(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>) {636 // CHECK-NEXT: acc.loop combined(serial)637 // CHECK: acc.yield638 // CHECK-NEXT: } loc639 // CHECK-NEXT: acc.yield640 // CHECK-NEXT: } loc641#pragma acc parallel loop firstprivate(notDefCtorArr[1:1])642 for(int i = 0; i < 5; ++i);643 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i644 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32645 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i646 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32647 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0648 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1649 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)650 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[NOTDEFCTORARR]] : !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>) bounds(%[[BOUNDS]]) recipe(@firstprivatization__Bcnt1__ZTSA5_14NonDefaultCtor) -> !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> {name = "notDefCtorArr[1:1]"}651 // CHECK-NEXT: acc.parallel combined(loop) firstprivate(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>) {652 // CHECK-NEXT: acc.loop combined(parallel)653 // CHECK: acc.yield654 // CHECK-NEXT: } loc655 // CHECK-NEXT: acc.yield656 // CHECK-NEXT: } loc657#pragma acc parallel loop firstprivate(dtorArr[1:1])658 for(int i = 0; i < 5; ++i);659 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i660 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32661 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i662 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32663 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0664 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1665 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)666 // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[DTORARR]] : !cir.ptr<!cir.array<!rec_HasDtor x 5>>) bounds(%[[BOUNDS]]) recipe(@firstprivatization__Bcnt1__ZTSA5_7HasDtor) -> !cir.ptr<!cir.array<!rec_HasDtor x 5>> {name = "dtorArr[1:1]"}667 // CHECK-NEXT: acc.parallel combined(loop) firstprivate(%[[PRIVATE]] : !cir.ptr<!cir.array<!rec_HasDtor x 5>>) {668 // CHECK-NEXT: acc.loop combined(parallel)669 // CHECK: acc.yield670 // CHECK-NEXT: } loc671 // CHECK-NEXT: acc.yield672 // CHECK-NEXT: } loc673#pragma acc parallel loop firstprivate(someIntArr[1:1], someFloatArr[1:1], noCopyArr[1:1], hasCopyArr[1:1], notDefCtorArr[1:1], dtorArr[1:1])674 for(int i = 0; i < 5; ++i);675 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i676 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32677 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i678 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32679 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0680 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1681 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)682 // CHECK-NEXT: %[[PRIVATE1:.*]] = acc.firstprivate varPtr(%[[INTARR]] : !cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS]]) recipe(@firstprivatization__Bcnt1__ZTSA5_i) -> !cir.ptr<!cir.array<!s32i x 5>> {name = "someIntArr[1:1]"}683 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i684 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32685 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i686 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32687 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0688 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1689 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)690 // CHECK-NEXT: %[[PRIVATE2:.*]] = acc.firstprivate varPtr(%[[FLOATARR]] : !cir.ptr<!cir.array<!cir.float x 5>>) bounds(%[[BOUNDS]]) recipe(@firstprivatization__Bcnt1__ZTSA5_f) -> !cir.ptr<!cir.array<!cir.float x 5>> {name = "someFloatArr[1:1]"}691 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i692 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32693 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i694 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32695 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0696 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1697 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)698 // CHECK-NEXT: %[[PRIVATE3:.*]] = acc.firstprivate varPtr(%[[NOCOPYARR]] : !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>) bounds(%[[BOUNDS]]) recipe(@firstprivatization__Bcnt1__ZTSA5_15NoCopyConstruct) -> !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> {name = "noCopyArr[1:1]"}699 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i700 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32701 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i702 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32703 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0704 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1705 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)706 // CHECK-NEXT: %[[PRIVATE4:.*]] = acc.firstprivate varPtr(%[[HASCOPYARR]] : !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>) bounds(%[[BOUNDS]]) recipe(@firstprivatization__Bcnt1__ZTSA5_13CopyConstruct) -> !cir.ptr<!cir.array<!rec_CopyConstruct x 5>> {name = "hasCopyArr[1:1]"}707 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i708 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32709 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i710 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32711 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0712 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1713 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)714 // CHECK-NEXT: %[[PRIVATE5:.*]] = acc.firstprivate varPtr(%[[NOTDEFCTORARR]] : !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>) bounds(%[[BOUNDS]]) recipe(@firstprivatization__Bcnt1__ZTSA5_14NonDefaultCtor) -> !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> {name = "notDefCtorArr[1:1]"}715 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i716 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32717 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i718 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32719 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0720 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1721 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)722 // CHECK-NEXT: %[[PRIVATE6:.*]] = acc.firstprivate varPtr(%[[DTORARR]] : !cir.ptr<!cir.array<!rec_HasDtor x 5>>) bounds(%[[BOUNDS]]) recipe(@firstprivatization__Bcnt1__ZTSA5_7HasDtor) -> !cir.ptr<!cir.array<!rec_HasDtor x 5>> {name = "dtorArr[1:1]"}723 // CHECK-NEXT: acc.parallel combined(loop) firstprivate(%[[PRIVATE1]], %[[PRIVATE2]], %[[PRIVATE3]], %[[PRIVATE4]], %[[PRIVATE5]], %[[PRIVATE6]] : !cir.ptr<!cir.array<!s32i x 5>>,724 // CHECK-SAME: !cir.ptr<!cir.array<!cir.float x 5>>,725 // CHECK-SAME: !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>,726 // CHECK-SAME: !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>,727 // CHECK-SAME: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>,728 // CHECK-SAME: !cir.ptr<!cir.array<!rec_HasDtor x 5>>)729 // CHECK-NEXT: acc.loop combined(parallel)730 // CHECK: acc.yield731 // CHECK-NEXT: } loc732 // CHECK-NEXT: acc.yield733 // CHECK-NEXT: } loc734}735